diff --git a/clang/lib/DPCT/RuleInfra/CallExprRewriter.h b/clang/lib/DPCT/RuleInfra/CallExprRewriter.h index 8142ccd3def9..2d1f28675812 100644 --- a/clang/lib/DPCT/RuleInfra/CallExprRewriter.h +++ b/clang/lib/DPCT/RuleInfra/CallExprRewriter.h @@ -702,10 +702,6 @@ template void printMemberOp(StreamT &Stream, bool IsArrow) { Stream << "."; } -template void printDisambiguator(StreamT &Stream) { - Stream << "template "; -} - template void printCapture(StreamT &Stream, bool IsCaptureRef) { if (IsCaptureRef) @@ -941,7 +937,7 @@ class ArgsPrinter template void printBase(StreamT &Stream, std::pair P, - bool IsArrow, bool NeedDisambiguator) { + bool IsArrow) { { std::unique_ptr> Paren; if (needExtraParensInMemberExpr(P.second)) @@ -949,13 +945,10 @@ void printBase(StreamT &Stream, std::pair P, print(Stream, P); } printMemberOp(Stream, IsArrow); - if (NeedDisambiguator) - printDisambiguator(Stream); } template -void printBase(StreamT &Stream, const Expr *E, bool IsArrow, - bool NeedDisambiguator) { +void printBase(StreamT &Stream, const Expr *E, bool IsArrow) { { std::unique_ptr> Paren; if (needExtraParensInMemberExpr(E)) @@ -963,23 +956,15 @@ void printBase(StreamT &Stream, const Expr *E, bool IsArrow, print(Stream, E); } printMemberOp(Stream, IsArrow); - if (NeedDisambiguator) - printDisambiguator(Stream); } template -void printBase(StreamT &Stream, const DerefExpr &D, bool, - bool NeedDisambiguator) { +void printBase(StreamT &Stream, const DerefExpr &D, bool) { D.printMemberBase(Stream); - if (NeedDisambiguator) - printDisambiguator(Stream); } template -void printBase(StreamT &Stream, const T &Val, bool IsArrow, - bool NeedDisambiguator) { +void printBase(StreamT &Stream, const T &Val, bool IsArrow) { print(Stream, Val); printMemberOp(Stream, IsArrow); - if (NeedDisambiguator) - printDisambiguator(Stream); } template class CallExprPrinter { @@ -1038,24 +1023,17 @@ template class TypeNamePrinter { } }; -template -class MemberExprPrinter { +template class MemberExprPrinter { BaseT Base; bool IsArrow; MemberT MemberName; - bool IsBaseDependentType = false; public: MemberExprPrinter(const BaseT &Base, bool IsArrow, MemberT MemberName) - : Base(Base), IsArrow(IsArrow), MemberName(MemberName) { - if constexpr (std::is_same_v) { - IsBaseDependentType = Base->getType()->isDependentType(); - } - } + : Base(Base), IsArrow(IsArrow), MemberName(MemberName) {} template void print(StreamT &Stream) const { - printBase(Stream, Base, IsArrow, - HasExplicitTemplateArg && IsBaseDependentType); + printBase(Stream, Base, IsArrow); dpct::print(Stream, MemberName); } }; @@ -1074,19 +1052,19 @@ template class StaticMemberExprPrinter { } }; -template class MemberCallPrinter : public CallExprPrinter< - MemberExprPrinter, + MemberExprPrinter, CallArgsT...> { public: MemberCallPrinter(const BaseT &Base, bool IsArrow, MemberT MemberName, CallArgsT &&...Args) : CallExprPrinter< - MemberExprPrinter, + MemberExprPrinter, CallArgsT...>( - MemberExprPrinter( + MemberExprPrinter( std::move(Base), IsArrow, std::move(MemberName)), std::forward(Args)...) {} }; @@ -1441,35 +1419,32 @@ class TemplatedCallExprRewriter template class MemberExprRewriter - : public PrinterRewriter> { + : public PrinterRewriter> { public: MemberExprRewriter( const CallExpr *C, StringRef Source, const std::function &BaseCreator, bool IsArrow, const std::function &MemberCreator) - : PrinterRewriter>( + : PrinterRewriter>( C, Source, BaseCreator(C), IsArrow, MemberCreator(C)) {} }; -template +template class MemberCallExprRewriter - : public PrinterRewriter> { + : public PrinterRewriter> { public: MemberCallExprRewriter( const CallExpr *C, StringRef Source, const std::function &BaseCreator, bool IsArrow, StringRef Member, const std::function &...ArgsCreator) - : PrinterRewriter>( + : PrinterRewriter>( C, Source, BaseCreator(C), IsArrow, Member, ArgsCreator(C)...) {} MemberCallExprRewriter( const CallExpr *C, StringRef Source, const BaseT &BaseCreator, bool IsArrow, StringRef Member, const std::function &...ArgsCreator) - : PrinterRewriter>( + : PrinterRewriter>( C, Source, BaseCreator, IsArrow, Member, ArgsCreator(C)...) {} }; diff --git a/clang/lib/DPCT/RuleInfra/CallExprRewriterCommon.h b/clang/lib/DPCT/RuleInfra/CallExprRewriterCommon.h index 32818f06fff7..b2bb14371633 100644 --- a/clang/lib/DPCT/RuleInfra/CallExprRewriterCommon.h +++ b/clang/lib/DPCT/RuleInfra/CallExprRewriterCommon.h @@ -439,32 +439,31 @@ inline std::function makeDeviceStr() { }; } -template -using MemberCallPrinterCreator = PrinterCreator< - MemberCallPrinter, - std::function, bool, std::string, - std::function...>; - -template -inline std::function(const CallExpr *)> +template +using MemberCallPrinterCreator = + PrinterCreator, + std::function, bool, std::string, + std::function...>; + +template +inline std::function< + MemberCallPrinter(const CallExpr *)> makeMemberCallCreator(std::function BaseFunc, bool IsArrow, std::string Member, std::function... Args) { - return MemberCallPrinterCreator( - BaseFunc, IsArrow, Member, Args...); + return MemberCallPrinterCreator(BaseFunc, IsArrow, + Member, Args...); } -template -inline std::function< - MemberCallPrinter(const CallExpr *)> +template +inline std::function(const CallExpr *)> makeMemberCallCreator(std::function BaseFunc, bool IsArrow, std::function Member) { - return PrinterCreator< - MemberCallPrinter, - std::function, bool, - std::function>(BaseFunc, IsArrow, Member); + return PrinterCreator, + std::function, bool, + std::function>( + BaseFunc, IsArrow, Member); } template @@ -691,10 +690,10 @@ makeArgWithAddressSpaceCast(int ArgIdx) { } template -inline std::function(const CallExpr *)> +inline std::function(const CallExpr *)> makeMemberExprCreator(std::function Base, bool IsArrow, std::function Member) { - return PrinterCreator, + return PrinterCreator, std::function, bool, std::function>(Base, IsArrow, Member); @@ -1344,7 +1343,7 @@ createTemplatedCallExprRewriterFactory( /// \p BaseCreator use to get base expr from original call expr. /// \p IsArrow the member operator is arrow or dot as default. /// \p ArgsCreator use to get call args from original call expr. -template +template inline std::shared_ptr createMemberCallExprRewriterFactory( const std::string &SourceName, @@ -1352,7 +1351,7 @@ createMemberCallExprRewriterFactory( std::string MemberName, std::function... ArgsCreator) { return std::make_shared, + MemberCallExprRewriter, std::function, bool, std::string, std::function...>>( SourceName, @@ -1361,7 +1360,7 @@ createMemberCallExprRewriterFactory( std::forward>(ArgsCreator)...); } -template +template inline std::shared_ptr, CallExprRewriterFactoryBase>> createMemberCallExprRewriterFactory( @@ -1369,8 +1368,8 @@ createMemberCallExprRewriterFactory( std::string MemberName, std::function... ArgsCreator) { return std::make_shared, BaseT, - bool, std::string, std::function...>>( + MemberCallExprRewriter, BaseT, bool, std::string, + std::function...>>( SourceName, BaseCreator, IsArrow, MemberName, std::forward>(ArgsCreator)...); } @@ -1635,19 +1634,19 @@ createBindTextureRewriterFactory(const std::string &Source) { return std::make_shared( makePointerChecker(StartIdx + 0), - createMemberCallExprRewriterFactory( + createMemberCallExprRewriterFactory( Source, makeDerefExprCreator(StartIdx + 0), true, "attach", makeCallArgCreator(StartIdx + 1), makeCallArgCreator(StartIdx + Idx + 1)..., makeDerefExprCreator(StartIdx + 2)), std::make_shared( TypeChecker, - createMemberCallExprRewriterFactory( + createMemberCallExprRewriterFactory( Source, makeCallArgCreatorWithCall(StartIdx + 0), false, "attach", makeCallArgCreatorWithCall(StartIdx + 1), makeCallArgCreatorWithCall(StartIdx + Idx + 1)..., makeCallArgCreatorWithCall(StartIdx + 2)), - createMemberCallExprRewriterFactory( + createMemberCallExprRewriterFactory( Source, makeCallArgCreatorWithCall(StartIdx + 0), false, "attach", makeCallArgCreatorWithCall(StartIdx + 1), makeCallArgCreatorWithCall(StartIdx + Idx)...))); @@ -2184,9 +2183,7 @@ const std::string MipmapNeedBindlessImage = #define UO(Op, E) makeUnaryOperatorCreator(E) #define BO(Op, L, R) makeBinaryOperatorCreator(L, R) #define PAREN(E) makeParenExprCreator(E) -#define MEMBER_CALL(...) makeMemberCallCreator(__VA_ARGS__) -#define MEMBER_CALL_HAS_EXPLICIT_TEMP_ARG(...) \ - makeMemberCallCreator(__VA_ARGS__) +#define MEMBER_CALL(...) makeMemberCallCreator(__VA_ARGS__) #define MEMBER_EXPR(...) makeMemberExprCreator(__VA_ARGS__) #define STATIC_MEMBER_EXPR(...) makeStaticMemberExprCreator(__VA_ARGS__) #define LAMBDA(...) makeLambdaCreator(__VA_ARGS__) @@ -2250,11 +2247,8 @@ const std::string MipmapNeedBindlessImage = #define CALL_FACTORY_ENTRY(FuncName, C) \ std::make_pair(FuncName, createCallExprRewriterFactory(FuncName, C)), #define MEMBER_CALL_FACTORY_ENTRY(FuncName, ...) \ - std::make_pair(FuncName, createMemberCallExprRewriterFactory( \ - FuncName, __VA_ARGS__)), -#define MEMBER_CALL_HAS_EXPLICIT_TEMP_ARG_FACTORY_ENTRY(FuncName, ...) \ - std::make_pair(FuncName, createMemberCallExprRewriterFactory( \ - FuncName, __VA_ARGS__)), + std::make_pair(FuncName, \ + createMemberCallExprRewriterFactory(FuncName, __VA_ARGS__)), #define ARRAYSUBSCRIPT_EXPR_FACTORY_ENTRY(FuncName, ...) \ std::make_pair(FuncName, createArraySubscriptExprRewriterFactory( \ FuncName, __VA_ARGS__)), diff --git a/clang/lib/DPCT/RuleInfra/MemberExprRewriter.cpp b/clang/lib/DPCT/RuleInfra/MemberExprRewriter.cpp index 4b24cf85bda4..8a4fcfb1f4d2 100644 --- a/clang/lib/DPCT/RuleInfra/MemberExprRewriter.cpp +++ b/clang/lib/DPCT/RuleInfra/MemberExprRewriter.cpp @@ -40,7 +40,7 @@ template class MEMemberExprPrinter { : Base(Base), IsArrow(IsArrow), MemberName(MemberName) {} template void print(StreamT &Stream) const { - printBase(Stream, Base, IsArrow, false); + printBase(Stream, Base, IsArrow); dpct::print(Stream, MemberName); } }; diff --git a/clang/lib/DPCT/RulesLang/APINamesComplex.inc b/clang/lib/DPCT/RulesLang/APINamesComplex.inc index 7152711b6146..054c6d752872 100644 --- a/clang/lib/DPCT/RulesLang/APINamesComplex.inc +++ b/clang/lib/DPCT/RulesLang/APINamesComplex.inc @@ -95,5 +95,7 @@ BINARY_OP_FACTORY_ENTRY("cuCfmaf", BinaryOperatorKind::BO_Add, makeCallArgCreatorWithCall(1)), makeCallArgCreatorWithCall(2)) -MEMBER_CALL_FACTORY_ENTRY("cuComplexDoubleToFloat", ARG(0), false, "convert") -MEMBER_CALL_FACTORY_ENTRY("cuComplexFloatToDouble", ARG(0), false, "convert") +MEMBER_CALL_FACTORY_ENTRY("cuComplexDoubleToFloat", ARG(0), false, + "template convert") +MEMBER_CALL_FACTORY_ENTRY("cuComplexFloatToDouble", ARG(0), false, + "template convert") diff --git a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp index 139887d37f3b..c00ffc8a067b 100644 --- a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp +++ b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp @@ -23,7 +23,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { template std::shared_ptr createRewriter(const CallExpr *C, bool RetAssign, BaseT Base) const { - using ReaderPrinter = decltype(makeMemberCallCreator( + using ReaderPrinter = decltype(makeMemberCallCreator( std::declval>(), false, TargetName, makeCallArgCreatorWithCall(Idx)...)(C)); if (RetAssign) { diff --git a/clang/lib/DPCT/RulesLang/Math/CallExprRewriterMath.cpp b/clang/lib/DPCT/RulesLang/Math/CallExprRewriterMath.cpp index 94e58fc5c6df..f38c48bde7a7 100644 --- a/clang/lib/DPCT/RulesLang/Math/CallExprRewriterMath.cpp +++ b/clang/lib/DPCT/RulesLang/Math/CallExprRewriterMath.cpp @@ -288,75 +288,24 @@ std::optional MathTypeCastRewriter::rewrite() { const StringRef &FuncName = SourceCalleeName; std::string ReplStr; llvm::raw_string_ostream OS(ReplStr); - auto MigratedArg0 = getMigratedArgWithExtraParens(0); - if (FuncName == "__float22half2_rn") { - OS << MigratedArg0 - << ".convert<" + MapNames::getClNamespace() + "half, " + - MapNames::getClNamespace() + "rounding_mode::rte>()"; - } else if (FuncName == "__float2half2_rn") { - OS << MapNames::getClNamespace() + "float2{" << MigratedArg0 << "," - << MigratedArg0 - << "}.convert<" + MapNames::getClNamespace() + "half, " + - MapNames::getClNamespace() + "rounding_mode::rte>()"; - } else if (FuncName == "__floats2half2_rn") { - auto MigratedArg1 = getMigratedArg(1); - OS << MapNames::getClNamespace() + "float2{" << MigratedArg0 << "," - << MigratedArg1 - << "}.convert<" + MapNames::getClNamespace() + "half, " + - MapNames::getClNamespace() + "rounding_mode::rte>()"; - } else if (FuncName == "__half22float2") { - OS << MigratedArg0 - << ".convert()"; - } else if (FuncName == "__half2half2") { - OS << MapNames::getClNamespace() + "half2{" << MigratedArg0 << "," - << MigratedArg0 << "}"; - } else if (FuncName == "__halves2half2") { - auto MigratedArg1 = getMigratedArg(1); - OS << MapNames::getClNamespace() + "half2{" << MigratedArg0 << "," - << MigratedArg1 << "}"; - } else if (FuncName == "__high2half") { - OS << MigratedArg0 << "[0]"; - } else if (FuncName == "__high2half2") { - OS << MapNames::getClNamespace() + "half2{" << MigratedArg0 << "[0], " - << MigratedArg0 << "[0]}"; - } else if (FuncName == "__highs2half2") { - auto MigratedArg1 = getMigratedArgWithExtraParens(1); - OS << MapNames::getClNamespace() + "half2{" << MigratedArg0 << "[0], " - << MigratedArg1 << "[0]}"; - } else if (FuncName == "__low2half") { - OS << MigratedArg0 << "[1]"; - } else if (FuncName == "__low2half2") { - OS << MapNames::getClNamespace() + "half2{" << MigratedArg0 << "[1], " - << MigratedArg0 << "[1]}"; - } else if (FuncName == "__lowhigh2highlow") { - OS << MapNames::getClNamespace() + "half2{" << MigratedArg0 << "[1], " - << MigratedArg0 << "[0]}"; - } else if (FuncName == "__lows2half2") { - auto MigratedArg1 = getMigratedArgWithExtraParens(1); - OS << MapNames::getClNamespace() + "half2{" << MigratedArg0 << "[1], " - << MigratedArg1 << "[1]}"; - } else { - //__half2short_rd and __half2float - static SSMap TypeMap{{"ll", "long long"}, - {"ull", "unsigned long long"}, - {"ushort", "unsigned short"}, - {"uint", "unsigned int"}, - {"half", MapNames::getClNamespace() + "half"}}; - std::string RoundingMode; - if (FuncName[FuncName.size() - 3] == '_') - RoundingMode = FuncName.substr(FuncName.size() - 2).str(); - auto FN = FuncName.substr(2, FuncName.find('_', 2) - 2).str(); - auto Types = split(FN, '2'); - assert(Types.size() == 2); - MapNames::replaceName(TypeMap, Types[0]); - MapNames::replaceName(TypeMap, Types[1]); - OS << MapNames::getClNamespace() + "vec<" << Types[0] << ", 1>{" - << MigratedArg0 << "}.convert<" << Types[1] - << ", " + MapNames::getClNamespace() + "rounding_mode::" - << RoundingModeMap[RoundingMode] << ">()[0]"; - } + static SSMap TypeMap{{"ll", "long long"}, + {"ull", "unsigned long long"}, + {"ushort", "unsigned short"}, + {"uint", "unsigned int"}, + {"half", MapNames::getClNamespace() + "half"}}; + std::string RoundingMode; + if (FuncName[FuncName.size() - 3] == '_') + RoundingMode = FuncName.substr(FuncName.size() - 2).str(); + auto FN = FuncName.substr(2, FuncName.find('_', 2) - 2).str(); + auto Types = split(FN, '2'); + assert(Types.size() == 2); + MapNames::replaceName(TypeMap, Types[0]); + MapNames::replaceName(TypeMap, Types[1]); + OS << MapNames::getClNamespace() + "vec<" << Types[0] << ", 1>{" + << MigratedArg0 << "}.template convert<" << Types[1] + << ", " + MapNames::getClNamespace() + "rounding_mode::" + << RoundingModeMap[RoundingMode] << ">()[0]"; OS.flush(); return ReplStr; } diff --git a/clang/lib/DPCT/RulesLang/Math/RewriterBfloat16PrecisionConversionAndDataMovement.cpp b/clang/lib/DPCT/RulesLang/Math/RewriterBfloat16PrecisionConversionAndDataMovement.cpp index 26852b834d5d..d84aae3aa6a6 100644 --- a/clang/lib/DPCT/RulesLang/Math/RewriterBfloat16PrecisionConversionAndDataMovement.cpp +++ b/clang/lib/DPCT/RulesLang/Math/RewriterBfloat16PrecisionConversionAndDataMovement.cpp @@ -67,7 +67,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), LITERAL("0"))) // __bfloat162int_rn @@ -84,7 +84,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), LITERAL("0"))) // __bfloat162int_ru @@ -101,7 +101,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), LITERAL("0"))) // __bfloat162int_rz @@ -118,7 +118,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), LITERAL("0"))) // __bfloat162ll_rd @@ -135,7 +135,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), LITERAL("0"))) // __bfloat162ll_rn @@ -152,7 +152,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), LITERAL("0"))) // __bfloat162ll_ru @@ -169,7 +169,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), LITERAL("0"))) // __bfloat162ll_rz @@ -186,7 +186,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), LITERAL("0"))) // __bfloat162short_rd @@ -203,7 +203,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), LITERAL("0"))) // __bfloat162short_rn @@ -220,7 +220,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), LITERAL("0"))) // __bfloat162short_ru @@ -237,7 +237,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), LITERAL("0"))) // __bfloat162short_rz @@ -254,7 +254,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), LITERAL("0"))) // __bfloat162uint_rd @@ -271,7 +271,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), LITERAL("0"))) // __bfloat162uint_rn @@ -288,7 +288,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), LITERAL("0"))) // __bfloat162uint_ru @@ -305,7 +305,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), LITERAL("0"))) // __bfloat162uint_rz @@ -322,7 +322,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), LITERAL("0"))) // __bfloat162ull_rd @@ -339,8 +339,8 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), + "template convert"), LITERAL("0"))) // __bfloat162ull_rn CONDITIONAL_FACTORY_ENTRY( @@ -356,8 +356,8 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), + "template convert"), LITERAL("0"))) // __bfloat162ull_ru CONDITIONAL_FACTORY_ENTRY( @@ -373,8 +373,8 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), + "template convert"), LITERAL("0"))) // __bfloat162ull_rz CONDITIONAL_FACTORY_ENTRY( @@ -390,8 +390,8 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), + "template convert"), LITERAL("0"))) // __bfloat162ushort_rd CONDITIONAL_FACTORY_ENTRY( @@ -408,8 +408,8 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), + "template convert"), LITERAL("0"))) // __bfloat162ushort_rn CONDITIONAL_FACTORY_ENTRY( @@ -426,8 +426,8 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), + "template convert"), LITERAL("0"))) // __bfloat162ushort_ru CONDITIONAL_FACTORY_ENTRY( @@ -444,8 +444,8 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), + "template convert"), LITERAL("0"))) // __bfloat162ushort_rz CONDITIONAL_FACTORY_ENTRY( @@ -462,8 +462,8 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap() MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert"), + "template convert"), LITERAL("0"))) // __bfloat16_as_short CONDITIONAL_FACTORY_ENTRY( diff --git a/clang/lib/DPCT/RulesLang/Math/RewriterHalfPrecisionConversionAndDataMovement.cpp b/clang/lib/DPCT/RulesLang/Math/RewriterHalfPrecisionConversionAndDataMovement.cpp index 8d6607742957..4125ef1a1eed 100644 --- a/clang/lib/DPCT/RulesLang/Math/RewriterHalfPrecisionConversionAndDataMovement.cpp +++ b/clang/lib/DPCT/RulesLang/Math/RewriterHalfPrecisionConversionAndDataMovement.cpp @@ -30,9 +30,9 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { CALL(MapNames::getClNamespace() + "ext::intel::math::float2half_rn", MEMBER_CALL(ARG(0), false, "y"))))), - MEMBER_CALL_HAS_EXPLICIT_TEMP_ARG_FACTORY_ENTRY( + MEMBER_CALL_FACTORY_ENTRY( "__float22half2_rn", ARG(0), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rte>")) // __float2half CONDITIONAL_FACTORY_ENTRY( @@ -48,7 +48,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::automatic>"), LITERAL("0"))) // __float2half2_rn @@ -65,7 +65,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL_FACTORY_ENTRY( "__float2half2_rn", CALL(MapNames::getClNamespace() + "float2", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rte>")) // __float2half_rd CONDITIONAL_FACTORY_ENTRY( @@ -81,7 +81,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rtn>"), LITERAL("0"))) // __float2half_rn @@ -98,7 +98,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rte>"), LITERAL("0"))) // __float2half_ru @@ -115,7 +115,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rtp>"), LITERAL("0"))) // __float2half_rz @@ -132,7 +132,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rtz>"), LITERAL("0"))) // __floats2half2_rn @@ -153,7 +153,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { "__floats2half2_rn", CALL(MapNames::getClNamespace() + "float2", ARG(0), ARG(1)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rte>")) // __half22float2 CONDITIONAL_FACTORY_ENTRY( @@ -168,10 +168,10 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { CALL(MapNames::getClNamespace() + "ext::intel::math::half2float", MEMBER_CALL(ARG(0), false, "y"))))), - MEMBER_CALL_HAS_EXPLICIT_TEMP_ARG_FACTORY_ENTRY( - "__half22float2", ARG(0), false, - "convert")) + MEMBER_CALL_FACTORY_ENTRY("__half22float2", ARG(0), false, + "template convert")) // __half2float CONDITIONAL_FACTORY_ENTRY( math::UseIntelDeviceMath, @@ -187,7 +187,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2half2 @@ -208,7 +209,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2int_rn @@ -226,7 +228,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2int_ru @@ -244,7 +247,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2int_rz @@ -262,7 +266,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2ll_rd @@ -280,7 +285,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2ll_rn @@ -298,7 +304,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2ll_ru @@ -316,7 +323,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2ll_rz @@ -334,7 +342,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2short_rd @@ -352,7 +361,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2short_rn @@ -370,7 +380,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2short_ru @@ -388,7 +399,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2short_rz @@ -406,7 +418,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2uint_rd @@ -424,7 +437,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2uint_rn @@ -442,7 +456,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2uint_ru @@ -460,7 +475,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2uint_rz @@ -478,7 +494,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) // __half2ull_rd @@ -496,7 +513,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) @@ -515,7 +532,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) @@ -534,7 +551,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) @@ -553,7 +570,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) @@ -572,7 +589,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) @@ -591,7 +608,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) @@ -610,7 +627,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) @@ -629,7 +646,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MapNames::getClNamespace() + "half, 1>", ARG(0)), false, - "convert"), LITERAL("0"))) @@ -675,7 +692,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rtn>"), LITERAL("0"))) // __int2half_rn @@ -692,7 +709,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rte>"), LITERAL("0"))) // __int2half_ru @@ -709,7 +726,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rtp>"), LITERAL("0"))) // __int2half_rz @@ -726,7 +743,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rtz>"), LITERAL("0"))) // __funnelshift_l @@ -863,8 +880,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL(CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + - MapNames::getClNamespace() + + "template convert<" + MapNames::getClNamespace() + + "half, " + MapNames::getClNamespace() + "rounding_mode::rtn>"), LITERAL("0"))) // __ll2half_rn @@ -881,8 +898,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL(CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + - MapNames::getClNamespace() + + "template convert<" + MapNames::getClNamespace() + + "half, " + MapNames::getClNamespace() + "rounding_mode::rte>"), LITERAL("0"))) // __ll2half_ru @@ -899,8 +916,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL(CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + - MapNames::getClNamespace() + + "template convert<" + MapNames::getClNamespace() + + "half, " + MapNames::getClNamespace() + "rounding_mode::rtp>"), LITERAL("0"))) // __ll2half_rz @@ -917,8 +934,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL(CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + - MapNames::getClNamespace() + + "template convert<" + MapNames::getClNamespace() + + "half, " + MapNames::getClNamespace() + "rounding_mode::rtz>"), LITERAL("0"))) // __low2float @@ -953,7 +970,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rtn>"), LITERAL("0"))) // __short2half_rn @@ -970,7 +987,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rte>"), LITERAL("0"))) // __short2half_ru @@ -987,7 +1004,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rtp>"), LITERAL("0"))) // __short2half_rz @@ -1004,7 +1021,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rtz>"), LITERAL("0"))) // __short_as_half @@ -1070,7 +1087,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rtn>"), LITERAL("0"))) // __uint2half_rn @@ -1087,7 +1104,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rte>"), LITERAL("0"))) // __uint2half_ru @@ -1104,7 +1121,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rtp>"), LITERAL("0"))) // __uint2half_rz @@ -1121,7 +1138,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { MEMBER_CALL( CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rtz>"), LITERAL("0"))) // __ull2half_rd @@ -1139,8 +1156,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + - MapNames::getClNamespace() + + "template convert<" + MapNames::getClNamespace() + + "half, " + MapNames::getClNamespace() + "rounding_mode::rtn>"), LITERAL("0"))) // __ull2half_rn @@ -1158,8 +1175,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + - MapNames::getClNamespace() + + "template convert<" + MapNames::getClNamespace() + + "half, " + MapNames::getClNamespace() + "rounding_mode::rte>"), LITERAL("0"))) // __ull2half_ru @@ -1177,8 +1194,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + - MapNames::getClNamespace() + + "template convert<" + MapNames::getClNamespace() + + "half, " + MapNames::getClNamespace() + "rounding_mode::rtp>"), LITERAL("0"))) // __ull2half_rz @@ -1196,8 +1213,8 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + - MapNames::getClNamespace() + + "template convert<" + MapNames::getClNamespace() + + "half, " + MapNames::getClNamespace() + "rounding_mode::rtz>"), LITERAL("0"))) // __ushort2half_rd @@ -1215,7 +1232,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rtn>"), LITERAL("0"))) // __ushort2half_rn @@ -1233,7 +1250,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rte>"), LITERAL("0"))) // __ushort2half_ru @@ -1251,7 +1268,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rtp>"), LITERAL("0"))) // __ushort2half_rz @@ -1269,7 +1286,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() { CALL(MapNames::getClNamespace() + "vec", ARG(0)), false, - "convert<" + MapNames::getClNamespace() + "half, " + + "template convert<" + MapNames::getClNamespace() + "half, " + MapNames::getClNamespace() + "rounding_mode::rtz>"), LITERAL("0"))) // __ushort_as_half diff --git a/clang/lib/DPCT/RulesLang/RewriterSYCLcompat.cpp b/clang/lib/DPCT/RulesLang/RewriterSYCLcompat.cpp index fa01b1b1f99f..57dc8e99fced 100644 --- a/clang/lib/DPCT/RulesLang/RewriterSYCLcompat.cpp +++ b/clang/lib/DPCT/RulesLang/RewriterSYCLcompat.cpp @@ -139,9 +139,8 @@ void dpct::initRewriterMethodMapCooperativeGroupsSYCLcompat( auto SyclId2Dim3 = [](const std::string &SourceMember, const std::string &MemberCallName) { auto GetID = [Member = MemberCallName](std::string Dimension) { - return makeMemberCallCreator(MemberExprBase(), false, - std::move(Member), - makeLiteral(std::move(Dimension))); + return makeMemberCallCreator(MemberExprBase(), false, std::move(Member), + makeLiteral(std::move(Dimension))); }; auto Name = "cooperative_groups::__v1::thread_block." + SourceMember; return std::make_pair( diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 859bf795c812..cb13a4ca8f44 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -2770,8 +2770,8 @@ EventQueryTraversal::buildCallReplacement(const CallExpr *Call) { static std::string MemberName = "get_info<" + MapNames::getClNamespace() + "info::event::command_execution_status>"; std::string ReplStr; - MemberCallPrinter Printer(Call->getArg(0), - true, MemberName); + MemberCallPrinter Printer(Call->getArg(0), true, + MemberName); llvm::raw_string_ostream OS(ReplStr); Printer.print(OS); return new ReplaceStmt(Call, std::move(OS.str())); diff --git a/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp b/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp index 23a634b3595e..bb161c17d53a 100644 --- a/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp @@ -950,7 +950,7 @@ void TextureRule::runRule(const MatchFinder::MatchResult &Result) { std::shared_ptr Rewriter = std::make_shared( CE, std::make_shared>>( + const Expr *, RenameWithSuffix, StringRef>>>( CE, Name, CE->getArg(0), true, RenameWithSuffix("set", MethodName), Value)); std::optional Result = Rewriter->rewrite(); @@ -1308,7 +1308,7 @@ bool TextureRule::SettersMerger::applyResult() { std::string ReplacedText; llvm::raw_string_ostream OS(ReplacedText); - MemberCallPrinter> + MemberCallPrinter> Printer(D->getName(), IsArrow, "set", std::move(ArgsList)); Printer.print(OS); diff --git a/clang/test/dpct/compat_with_clang_4.cu b/clang/test/dpct/compat_with_clang_4.cu index a0053e8da80c..7491840a7380 100644 --- a/clang/test/dpct/compat_with_clang_4.cu +++ b/clang/test/dpct/compat_with_clang_4.cu @@ -7,7 +7,7 @@ #include "cuda_fp16.h" // CHECK: inline void foo1(sycl::half2 *array, sycl::half a) { -// CHECK-NEXT: array[dpct::reverse_bits(123)] = {a, sycl::vec(2.3f).convert()[0]}; +// CHECK-NEXT: array[dpct::reverse_bits(123)] = {a, sycl::vec(2.3f).template convert()[0]}; // CHECK-NEXT: } __device__ inline void foo1(__half2 *array, __half a) { array[__brev(123)] = {a, __float2half(2.3f)}; diff --git a/clang/test/dpct/complex.cu b/clang/test/dpct/complex.cu index 2d4542e4020b..4c62aebe29e2 100644 --- a/clang/test/dpct/complex.cu +++ b/clang/test/dpct/complex.cu @@ -290,9 +290,9 @@ int main() { auto a24 = COMPLEX_F_FMA(f1, f2, f3); r = r && check(a24, expect, index); - // CHECK: f1 = d1.convert(); + // CHECK: f1 = d1.template convert(); f1 = cuComplexDoubleToFloat(d1); - // CHECK: d1 = f1.convert(); + // CHECK: d1 = f1.template convert(); d1 = cuComplexFloatToDouble(f1); int *result = nullptr; diff --git a/clang/test/dpct/math/bfloat16/bfloat16.cu b/clang/test/dpct/math/bfloat16/bfloat16.cu index c66395795d72..df5466eadf61 100644 --- a/clang/test/dpct/math/bfloat16/bfloat16.cu +++ b/clang/test/dpct/math/bfloat16/bfloat16.cu @@ -233,53 +233,53 @@ __global__ void test_conversions_device(__nv_bfloat16 *deviceArrayBFloat16) { bf162 = __bfloat162bfloat162(bf16); // CHECK: f = static_cast(bf16); f = __bfloat162float(bf16); - // CHECK: i = sycl::vec(bf16).convert()[0]; + // CHECK: i = sycl::vec(bf16).template convert()[0]; i = __bfloat162int_rd(bf16); - // CHECK: i = sycl::vec(bf16).convert()[0]; + // CHECK: i = sycl::vec(bf16).template convert()[0]; i = __bfloat162int_rn(bf16); - // CHECK: i = sycl::vec(bf16).convert()[0]; + // CHECK: i = sycl::vec(bf16).template convert()[0]; i = __bfloat162int_ru(bf16); - // CHECK: i = sycl::vec(bf16).convert()[0]; + // CHECK: i = sycl::vec(bf16).template convert()[0]; i = __bfloat162int_rz(bf16); - // CHECK: ll = sycl::vec(bf16).convert()[0]; + // CHECK: ll = sycl::vec(bf16).template convert()[0]; ll = __bfloat162ll_rd(bf16); - // CHECK: ll = sycl::vec(bf16).convert()[0]; + // CHECK: ll = sycl::vec(bf16).template convert()[0]; ll = __bfloat162ll_rn(bf16); - // CHECK: ll = sycl::vec(bf16).convert()[0]; + // CHECK: ll = sycl::vec(bf16).template convert()[0]; ll = __bfloat162ll_ru(bf16); - // CHECK: ll = sycl::vec(bf16).convert()[0]; + // CHECK: ll = sycl::vec(bf16).template convert()[0]; ll = __bfloat162ll_rz(bf16); - // CHECK: s = sycl::vec(bf16).convert()[0]; + // CHECK: s = sycl::vec(bf16).template convert()[0]; s = __bfloat162short_rd(bf16); - // CHECK: s = sycl::vec(bf16).convert()[0]; + // CHECK: s = sycl::vec(bf16).template convert()[0]; s = __bfloat162short_rn(bf16); - // CHECK: s = sycl::vec(bf16).convert()[0]; + // CHECK: s = sycl::vec(bf16).template convert()[0]; s = __bfloat162short_ru(bf16); - // CHECK: s = sycl::vec(bf16).convert()[0]; + // CHECK: s = sycl::vec(bf16).template convert()[0]; s = __bfloat162short_rz(bf16); - // CHECK: u = sycl::vec(bf16).convert()[0]; + // CHECK: u = sycl::vec(bf16).template convert()[0]; u = __bfloat162uint_rd(bf16); - // CHECK: u = sycl::vec(bf16).convert()[0]; + // CHECK: u = sycl::vec(bf16).template convert()[0]; u = __bfloat162uint_rn(bf16); - // CHECK: u = sycl::vec(bf16).convert()[0]; + // CHECK: u = sycl::vec(bf16).template convert()[0]; u = __bfloat162uint_ru(bf16); - // CHECK: u = sycl::vec(bf16).convert()[0]; + // CHECK: u = sycl::vec(bf16).template convert()[0]; u = __bfloat162uint_rz(bf16); - // CHECK: ull = sycl::vec(bf16).convert()[0]; + // CHECK: ull = sycl::vec(bf16).template convert()[0]; ull = __bfloat162ull_rd(bf16); - // CHECK: ull = sycl::vec(bf16).convert()[0]; + // CHECK: ull = sycl::vec(bf16).template convert()[0]; ull = __bfloat162ull_rn(bf16); - // CHECK: ull = sycl::vec(bf16).convert()[0]; + // CHECK: ull = sycl::vec(bf16).template convert()[0]; ull = __bfloat162ull_ru(bf16); - // CHECK: ull = sycl::vec(bf16).convert()[0]; + // CHECK: ull = sycl::vec(bf16).template convert()[0]; ull = __bfloat162ull_rz(bf16); - // CHECK: us = sycl::vec(bf16).convert()[0]; + // CHECK: us = sycl::vec(bf16).template convert()[0]; us = __bfloat162ushort_rd(bf16); - // CHECK: us = sycl::vec(bf16).convert()[0]; + // CHECK: us = sycl::vec(bf16).template convert()[0]; us = __bfloat162ushort_rn(bf16); - // CHECK: us = sycl::vec(bf16).convert()[0]; + // CHECK: us = sycl::vec(bf16).template convert()[0]; us = __bfloat162ushort_ru(bf16); - // CHECK: us = sycl::vec(bf16).convert()[0]; + // CHECK: us = sycl::vec(bf16).template convert()[0]; us = __bfloat162ushort_rz(bf16); // CHECK: /* // CHECK-NEXT: DPCT1007:{{[0-9]+}}: Migration of __bfloat16_as_short is not supported. diff --git a/clang/test/dpct/math/cuda-math-intrinsics.cu b/clang/test/dpct/math/cuda-math-intrinsics.cu index ca94098a8895..dea24b5554e9 100644 --- a/clang/test/dpct/math/cuda-math-intrinsics.cu +++ b/clang/test/dpct/math/cuda-math-intrinsics.cu @@ -1607,115 +1607,115 @@ __global__ void kernelFuncTypecasts() { double d; double2 d2; - // CHECK: f = sycl::vec{d}.convert()[0]; + // CHECK: f = sycl::vec{d}.template convert()[0]; f = __double2float_rd(d); - // CHECK: f = sycl::vec{d}.convert()[0]; + // CHECK: f = sycl::vec{d}.template convert()[0]; f = __double2float_rn(d); - // CHECK: f = sycl::vec{d}.convert()[0]; + // CHECK: f = sycl::vec{d}.template convert()[0]; f = __double2float_ru(d); - // CHECK: f = sycl::vec{d}.convert()[0]; + // CHECK: f = sycl::vec{d}.template convert()[0]; f = __double2float_rz(d); - // CHECK: i = sycl::vec{d}.convert()[0]; + // CHECK: i = sycl::vec{d}.template convert()[0]; i = __double2int_rd(d); - // CHECK: i = sycl::vec{d}.convert()[0]; + // CHECK: i = sycl::vec{d}.template convert()[0]; i = __double2int_rn(d); - // CHECK: i = sycl::vec{d}.convert()[0]; + // CHECK: i = sycl::vec{d}.template convert()[0]; i = __double2int_ru(d); - // CHECK: i = sycl::vec{d}.convert()[0]; + // CHECK: i = sycl::vec{d}.template convert()[0]; i = __double2int_rz(d); - // CHECK: ll = sycl::vec{d}.convert()[0]; + // CHECK: ll = sycl::vec{d}.template convert()[0]; ll = __double2ll_rd(d); - // CHECK: ll = sycl::vec{d}.convert()[0]; + // CHECK: ll = sycl::vec{d}.template convert()[0]; ll = __double2ll_rn(d); - // CHECK: ll = sycl::vec{d}.convert()[0]; + // CHECK: ll = sycl::vec{d}.template convert()[0]; ll = __double2ll_ru(d); - // CHECK: ll = sycl::vec{d}.convert()[0]; + // CHECK: ll = sycl::vec{d}.template convert()[0]; ll = __double2ll_rz(d); - // CHECK: ui = sycl::vec{d}.convert()[0]; + // CHECK: ui = sycl::vec{d}.template convert()[0]; ui = __double2uint_rd(d); - // CHECK:ui = sycl::vec{d}.convert()[0]; + // CHECK:ui = sycl::vec{d}.template convert()[0]; ui = __double2uint_rn(d); - // CHECK: ui = sycl::vec{d}.convert()[0]; + // CHECK: ui = sycl::vec{d}.template convert()[0]; ui = __double2uint_ru(d); - // CHECK: ui = sycl::vec{d}.convert()[0]; + // CHECK: ui = sycl::vec{d}.template convert()[0]; ui = __double2uint_rz(d); - // CHECK: ull = sycl::vec{d}.convert()[0]; + // CHECK: ull = sycl::vec{d}.template convert()[0]; ull = __double2ull_rd(d); - // CHECK: ull = sycl::vec{d}.convert()[0]; + // CHECK: ull = sycl::vec{d}.template convert()[0]; ull = __double2ull_rn(d); - // CHECK: ull = sycl::vec{d}.convert()[0]; + // CHECK: ull = sycl::vec{d}.template convert()[0]; ull = __double2ull_ru(d); - // CHECK: ull = sycl::vec{d}.convert()[0]; + // CHECK: ull = sycl::vec{d}.template convert()[0]; ull = __double2ull_rz(d); // CHECK: ll = sycl::bit_cast(d); ll = __double_as_longlong(d); - // CHECK: i = sycl::vec{f}.convert()[0]; + // CHECK: i = sycl::vec{f}.template convert()[0]; i = __float2int_rd(f); - // CHECK: i = sycl::vec{f}.convert()[0]; + // CHECK: i = sycl::vec{f}.template convert()[0]; i = __float2int_rn(f); - // CHECK: i = sycl::vec{f}.convert()[0]; + // CHECK: i = sycl::vec{f}.template convert()[0]; i = __float2int_ru(f); - // CHECK: i = sycl::vec{f}.convert()[0]; + // CHECK: i = sycl::vec{f}.template convert()[0]; i = __float2int_rz(f); - // CHECK: ll = sycl::vec{f}.convert()[0]; + // CHECK: ll = sycl::vec{f}.template convert()[0]; ll = __float2ll_rd(f); - // CHECK: ll = sycl::vec{f}.convert()[0]; + // CHECK: ll = sycl::vec{f}.template convert()[0]; ll = __float2ll_rn(f); - // CHECK: ll = sycl::vec{f}.convert()[0]; + // CHECK: ll = sycl::vec{f}.template convert()[0]; ll = __float2ll_ru(f); - // CHECK: ll = sycl::vec{f}.convert()[0]; + // CHECK: ll = sycl::vec{f}.template convert()[0]; ll = __float2ll_rz(f); - // CHECK: ui = sycl::vec{f}.convert()[0]; + // CHECK: ui = sycl::vec{f}.template convert()[0]; ui = __float2uint_rd(f); - // CHECK: ui = sycl::vec{f}.convert()[0]; + // CHECK: ui = sycl::vec{f}.template convert()[0]; ui = __float2uint_rn(f); - // CHECK: ui = sycl::vec{f}.convert()[0]; + // CHECK: ui = sycl::vec{f}.template convert()[0]; ui = __float2uint_ru(f); - // CHECK: ui = sycl::vec{f}.convert()[0]; + // CHECK: ui = sycl::vec{f}.template convert()[0]; ui = __float2uint_rz(f); - // CHECK: ull = sycl::vec{f}.convert()[0]; + // CHECK: ull = sycl::vec{f}.template convert()[0]; ull = __float2ull_rd(f); - // CHECK: ull = sycl::vec{f}.convert()[0]; + // CHECK: ull = sycl::vec{f}.template convert()[0]; ull = __float2ull_rn(f); - // CHECK: ull = sycl::vec{f}.convert()[0]; + // CHECK: ull = sycl::vec{f}.template convert()[0]; ull = __float2ull_ru(f); - // CHECK: ull = sycl::vec{f}.convert()[0]; + // CHECK: ull = sycl::vec{f}.template convert()[0]; ull = __float2ull_rz(f); // CHECK: i = sycl::bit_cast(f); @@ -1724,91 +1724,91 @@ __global__ void kernelFuncTypecasts() { // CHECK: ui = sycl::bit_cast(f); ui = __float_as_uint(f); - // CHECK: d = sycl::vec{i}.convert()[0]; + // CHECK: d = sycl::vec{i}.template convert()[0]; d = __int2double_rn(i); - // CHECK: d = sycl::vec{i}.convert()[0]; + // CHECK: d = sycl::vec{i}.template convert()[0]; d = __int2float_rd(i); - // CHECK: d = sycl::vec{i}.convert()[0]; + // CHECK: d = sycl::vec{i}.template convert()[0]; d = __int2float_rn(i); - // CHECK: d = sycl::vec{i}.convert()[0]; + // CHECK: d = sycl::vec{i}.template convert()[0]; d = __int2float_ru(i); - // CHECK: d = sycl::vec{i}.convert()[0]; + // CHECK: d = sycl::vec{i}.template convert()[0]; d = __int2float_rz(i); // CHECK: f = sycl::bit_cast(i); f = __int_as_float(i); - // CHECK: d = sycl::vec{ll}.convert()[0]; + // CHECK: d = sycl::vec{ll}.template convert()[0]; d = __ll2double_rd(ll); - // CHECK: d = sycl::vec{ll}.convert()[0]; + // CHECK: d = sycl::vec{ll}.template convert()[0]; d = __ll2double_rn(ll); - // CHECK: d = sycl::vec{ll}.convert()[0]; + // CHECK: d = sycl::vec{ll}.template convert()[0]; d = __ll2double_ru(ll); - // CHECK: d = sycl::vec{ll}.convert()[0]; + // CHECK: d = sycl::vec{ll}.template convert()[0]; d = __ll2double_rz(ll); - // CHECK: f = sycl::vec{ll}.convert()[0]; + // CHECK: f = sycl::vec{ll}.template convert()[0]; f = __ll2float_rd(ll); - // CHECK: f = sycl::vec{ll}.convert()[0]; + // CHECK: f = sycl::vec{ll}.template convert()[0]; f = __ll2float_rn(ll); - // CHECK: f = sycl::vec{ll}.convert()[0]; + // CHECK: f = sycl::vec{ll}.template convert()[0]; f = __ll2float_ru(ll); - // CHECK: f = sycl::vec{ll}.convert()[0]; + // CHECK: f = sycl::vec{ll}.template convert()[0]; f = __ll2float_rz(ll); // CHECK: d = sycl::bit_cast(ll); d = __longlong_as_double(ll); - // CHECK: d = sycl::vec{ui}.convert()[0]; + // CHECK: d = sycl::vec{ui}.template convert()[0]; d = __uint2double_rn(ui); - // CHECK: f = sycl::vec{ui}.convert()[0]; + // CHECK: f = sycl::vec{ui}.template convert()[0]; f = __uint2float_rd(ui); - // CHECK: f = sycl::vec{ui}.convert()[0]; + // CHECK: f = sycl::vec{ui}.template convert()[0]; f = __uint2float_rn(ui); - // CHECK: f = sycl::vec{ui}.convert()[0]; + // CHECK: f = sycl::vec{ui}.template convert()[0]; f = __uint2float_ru(ui); - // CHECK: f = sycl::vec{ui}.convert()[0]; + // CHECK: f = sycl::vec{ui}.template convert()[0]; f = __uint2float_rz(ui); // CHECK: f = sycl::bit_cast(ui); f = __uint_as_float(ui); - // CHECK: d = sycl::vec{ull}.convert()[0]; + // CHECK: d = sycl::vec{ull}.template convert()[0]; d = __ull2double_rd(ull); - // CHECK: d = sycl::vec{ull}.convert()[0]; + // CHECK: d = sycl::vec{ull}.template convert()[0]; d = __ull2double_rn(ull); - // CHECK: d = sycl::vec{ull}.convert()[0]; + // CHECK: d = sycl::vec{ull}.template convert()[0]; d = __ull2double_ru(ull); - // CHECK: d = sycl::vec{ull}.convert()[0]; + // CHECK: d = sycl::vec{ull}.template convert()[0]; d = __ull2double_rz(ull); - // CHECK: f = sycl::vec{ull}.convert()[0]; + // CHECK: f = sycl::vec{ull}.template convert()[0]; f = __ull2float_rd(ull); - // CHECK: f = sycl::vec{ull}.convert()[0]; + // CHECK: f = sycl::vec{ull}.template convert()[0]; f = __ull2float_rn(ull); - // CHECK: f = sycl::vec{ull}.convert()[0]; + // CHECK: f = sycl::vec{ull}.template convert()[0]; f = __ull2float_ru(ull); - // CHECK: f = sycl::vec{ull}.convert()[0]; + // CHECK: f = sycl::vec{ull}.template convert()[0]; f = __ull2float_rz(ull); } diff --git a/clang/test/dpct/math/cuda-math-need-paren.cu b/clang/test/dpct/math/cuda-math-need-paren.cu index f3342fa2bb86..bcbb8a2becb4 100644 --- a/clang/test/dpct/math/cuda-math-need-paren.cu +++ b/clang/test/dpct/math/cuda-math-need-paren.cu @@ -8,7 +8,7 @@ using namespace std; void __global__ kernel() { half2 h2; - // CHECK: (h2 + h2).convert(); + // CHECK: (h2 + h2).template convert(); __half22float2(__hadd2(h2, h2)); } diff --git a/clang/test/dpct/math/half/half.cu b/clang/test/dpct/math/half/half.cu index ab2d3f989543..7c8fcf98c56c 100644 --- a/clang/test/dpct/math/half/half.cu +++ b/clang/test/dpct/math/half/half.cu @@ -15,75 +15,75 @@ __global__ void kernelFuncHalfConversion() { unsigned u; unsigned long long ull; unsigned short us; - // CHECK: h2 = f2.convert(); + // CHECK: h2 = f2.template convert(); h2 = __float22half2_rn(f2); - // CHECK: h = sycl::vec(f).convert()[0]; + // CHECK: h = sycl::vec(f).template convert()[0]; h = __float2half(f); - // CHECK: h2 = sycl::float2(f).convert(); + // CHECK: h2 = sycl::float2(f).template convert(); h2 = __float2half2_rn(f); - // CHECK: h = sycl::vec(f).convert()[0]; + // CHECK: h = sycl::vec(f).template convert()[0]; h = __float2half_rd(f); - // sycl::vec(f).convert()[0]; + // sycl::vec(f).template convert()[0]; __float2half_rn(f); - // CHECK: h = sycl::vec(f).convert()[0]; + // CHECK: h = sycl::vec(f).template convert()[0]; h = __float2half_ru(f); - // CHECK: h = sycl::vec(f).convert()[0]; + // CHECK: h = sycl::vec(f).template convert()[0]; h = __float2half_rz(f); - // CHECK: h2 = sycl::float2(f, f).convert(); + // CHECK: h2 = sycl::float2(f, f).template convert(); h2 = __floats2half2_rn(f, f); - // CHECK: f2 = h2.convert(); + // CHECK: f2 = h2.template convert(); f2 = __half22float2(h2); - // CHECK: f = sycl::vec(h).convert()[0]; + // CHECK: f = sycl::vec(h).template convert()[0]; f = __half2float(h); // CHECK: h2 = sycl::half2(h); h2 = __half2half2(h); - // CHECK: i = sycl::vec(h).convert()[0]; + // CHECK: i = sycl::vec(h).template convert()[0]; i = __half2int_rd(h); - // CHECK: i = sycl::vec(h).convert()[0]; + // CHECK: i = sycl::vec(h).template convert()[0]; i = __half2int_rn(h); - // CHECK: i = sycl::vec(h).convert()[0]; + // CHECK: i = sycl::vec(h).template convert()[0]; i = __half2int_ru(h); - // CHECK: i = sycl::vec(h).convert()[0]; + // CHECK: i = sycl::vec(h).template convert()[0]; i = __half2int_rz(h); - // CHECK: ll = sycl::vec(h).convert()[0]; + // CHECK: ll = sycl::vec(h).template convert()[0]; ll = __half2ll_rd(h); - // CHECK: ll = sycl::vec(h).convert()[0]; + // CHECK: ll = sycl::vec(h).template convert()[0]; ll = __half2ll_rn(h); - // CHECK: ll = sycl::vec(h).convert()[0]; + // CHECK: ll = sycl::vec(h).template convert()[0]; ll = __half2ll_ru(h); - // CHECK: ll = sycl::vec(h).convert()[0]; + // CHECK: ll = sycl::vec(h).template convert()[0]; ll = __half2ll_rz(h); - // CHECK: s = sycl::vec(h).convert()[0]; + // CHECK: s = sycl::vec(h).template convert()[0]; s = __half2short_rd(h); - // CHECK: s = sycl::vec(h).convert()[0]; + // CHECK: s = sycl::vec(h).template convert()[0]; s = __half2short_rn(h); - // CHECK: s = sycl::vec(h).convert()[0]; + // CHECK: s = sycl::vec(h).template convert()[0]; s = __half2short_ru(h); - // CHECK: s = sycl::vec(h).convert()[0]; + // CHECK: s = sycl::vec(h).template convert()[0]; s = __half2short_rz(h); - // CHECK: u = sycl::vec(h).convert()[0]; + // CHECK: u = sycl::vec(h).template convert()[0]; u = __half2uint_rd(h); - // CHECK: u = sycl::vec(h).convert()[0]; + // CHECK: u = sycl::vec(h).template convert()[0]; u = __half2uint_rn(h); - // CHECK: u = sycl::vec(h).convert()[0]; + // CHECK: u = sycl::vec(h).template convert()[0]; u = __half2uint_ru(h); - // CHECK: u = sycl::vec(h).convert()[0]; + // CHECK: u = sycl::vec(h).template convert()[0]; u = __half2uint_rz(h); - // CHECK: ull = sycl::vec(h).convert()[0]; + // CHECK: ull = sycl::vec(h).template convert()[0]; ull = __half2ull_rd(h); - // CHECK: ull = sycl::vec(h).convert()[0]; + // CHECK: ull = sycl::vec(h).template convert()[0]; ull = __half2ull_rn(h); - // CHECK: ull = sycl::vec(h).convert()[0]; + // CHECK: ull = sycl::vec(h).template convert()[0]; ull = __half2ull_ru(h); - // CHECK: ull = sycl::vec(h).convert()[0]; + // CHECK: ull = sycl::vec(h).template convert()[0]; ull = __half2ull_rz(h); - // CHECK: us = sycl::vec(h).convert()[0]; + // CHECK: us = sycl::vec(h).template convert()[0]; us = __half2ushort_rd(h); - // CHECK: us = sycl::vec(h).convert()[0]; + // CHECK: us = sycl::vec(h).template convert()[0]; us = __half2ushort_rn(h); - // CHECK: us = sycl::vec(h).convert()[0]; + // CHECK: us = sycl::vec(h).template convert()[0]; us = __half2ushort_ru(h); - // CHECK: us = sycl::vec(h).convert()[0]; + // CHECK: us = sycl::vec(h).template convert()[0]; us = __half2ushort_rz(h); // CHECK: s = sycl::bit_cast(h); s = __half_as_short(h); @@ -99,21 +99,21 @@ __global__ void kernelFuncHalfConversion() { h2 = __high2half2(h2); // CHECK: h2 = sycl::half2(h2.y(), h2.y()); h2 = __highs2half2(h2, h2); - // CHECK: h = sycl::vec(i).convert()[0]; + // CHECK: h = sycl::vec(i).template convert()[0]; h = __int2half_rd(i); - // CHECK: h = sycl::vec(i).convert()[0]; + // CHECK: h = sycl::vec(i).template convert()[0]; h = __int2half_rn(i); - // CHECK: h = sycl::vec(i).convert()[0]; + // CHECK: h = sycl::vec(i).template convert()[0]; h = __int2half_ru(i); - // CHECK: h = sycl::vec(i).convert()[0]; + // CHECK: h = sycl::vec(i).template convert()[0]; h = __int2half_rz(i); - // CHECK: h = sycl::vec(ll).convert()[0]; + // CHECK: h = sycl::vec(ll).template convert()[0]; h = __ll2half_rd(ll); - // CHECK: h = sycl::vec(ll).convert()[0]; + // CHECK: h = sycl::vec(ll).template convert()[0]; h = __ll2half_rn(ll); - // CHECK: h = sycl::vec(ll).convert()[0]; + // CHECK: h = sycl::vec(ll).template convert()[0]; h = __ll2half_ru(ll); - // CHECK: h = sycl::vec(ll).convert()[0]; + // CHECK: h = sycl::vec(ll).template convert()[0]; h = __ll2half_rz(ll); // CHECK: f = h2[0]; f = __low2float(h2); @@ -127,39 +127,39 @@ __global__ void kernelFuncHalfConversion() { h2 = __lowhigh2highlow(h2); // CHECK: h2 = sycl::half2(h2.x(), h2.x()); h2 = __lows2half2(h2, h2); - // CHECK: h = sycl::vec(s).convert()[0]; + // CHECK: h = sycl::vec(s).template convert()[0]; h = __short2half_rd(s); - // CHECK: h = sycl::vec(s).convert()[0]; + // CHECK: h = sycl::vec(s).template convert()[0]; h = __short2half_rn(s); - // CHECK: h = sycl::vec(s).convert()[0]; + // CHECK: h = sycl::vec(s).template convert()[0]; h = __short2half_ru(s); - // CHECK: h = sycl::vec(s).convert()[0]; + // CHECK: h = sycl::vec(s).template convert()[0]; h = __short2half_rz(s); // CHECK: h = sycl::bit_cast(s); h = __short_as_half(s); - // CHECK: h = sycl::vec(u).convert()[0]; + // CHECK: h = sycl::vec(u).template convert()[0]; h = __uint2half_rd(u); - // CHECK: h = sycl::vec(u).convert()[0]; + // CHECK: h = sycl::vec(u).template convert()[0]; h = __uint2half_rn(u); - // CHECK: h = sycl::vec(u).convert()[0]; + // CHECK: h = sycl::vec(u).template convert()[0]; h = __uint2half_ru(u); - // CHECK: h = sycl::vec(u).convert()[0]; + // CHECK: h = sycl::vec(u).template convert()[0]; h = __uint2half_rz(u); - // CHECK: h = sycl::vec(ull).convert()[0]; + // CHECK: h = sycl::vec(ull).template convert()[0]; h = __ull2half_rd(ull); - // CHECK: h = sycl::vec(ull).convert()[0]; + // CHECK: h = sycl::vec(ull).template convert()[0]; h = __ull2half_rn(ull); - // CHECK: h = sycl::vec(ull).convert()[0]; + // CHECK: h = sycl::vec(ull).template convert()[0]; h = __ull2half_ru(ull); - // CHECK: h = sycl::vec(ull).convert()[0]; + // CHECK: h = sycl::vec(ull).template convert()[0]; h = __ull2half_rz(ull); - // CHECK: h = sycl::vec(us).convert()[0]; + // CHECK: h = sycl::vec(us).template convert()[0]; h = __ushort2half_rd(us); - // CHECK: h = sycl::vec(us).convert()[0]; + // CHECK: h = sycl::vec(us).template convert()[0]; h = __ushort2half_rn(us); - // CHECK: h = sycl::vec(us).convert()[0]; + // CHECK: h = sycl::vec(us).template convert()[0]; h = __ushort2half_ru(us); - // CHECK: h = sycl::vec(us).convert()[0]; + // CHECK: h = sycl::vec(us).template convert()[0]; h = __ushort2half_rz(us); // CHECK: h = sycl::bit_cast(us); h = __ushort_as_half(us); diff --git a/clang/test/dpct/query_api_mapping/Math/test_type_casting_intrinsics.cu b/clang/test/dpct/query_api_mapping/Math/test_type_casting_intrinsics.cu index f6c0ecc76dc7..ab2bd66a2d8c 100644 --- a/clang/test/dpct/query_api_mapping/Math/test_type_casting_intrinsics.cu +++ b/clang/test/dpct/query_api_mapping/Math/test_type_casting_intrinsics.cu @@ -4,25 +4,25 @@ // __DOUBLE2FLOAT_RD: CUDA API: // __DOUBLE2FLOAT_RD-NEXT: __double2float_rd(d /*double*/); // __DOUBLE2FLOAT_RD-NEXT: Is migrated to: -// __DOUBLE2FLOAT_RD-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2FLOAT_RD-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2float_rn | FileCheck %s -check-prefix=__DOUBLE2FLOAT_RN // __DOUBLE2FLOAT_RN: CUDA API: // __DOUBLE2FLOAT_RN-NEXT: __double2float_rn(d /*double*/); // __DOUBLE2FLOAT_RN-NEXT: Is migrated to: -// __DOUBLE2FLOAT_RN-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2FLOAT_RN-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2float_ru | FileCheck %s -check-prefix=__DOUBLE2FLOAT_RU // __DOUBLE2FLOAT_RU: CUDA API: // __DOUBLE2FLOAT_RU-NEXT: __double2float_ru(d /*double*/); // __DOUBLE2FLOAT_RU-NEXT: Is migrated to: -// __DOUBLE2FLOAT_RU-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2FLOAT_RU-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2float_rz | FileCheck %s -check-prefix=__DOUBLE2FLOAT_RZ // __DOUBLE2FLOAT_RZ: CUDA API: // __DOUBLE2FLOAT_RZ-NEXT: __double2float_rz(d /*double*/); // __DOUBLE2FLOAT_RZ-NEXT: Is migrated to: -// __DOUBLE2FLOAT_RZ-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2FLOAT_RZ-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2hiint | FileCheck %s -check-prefix=__DOUBLE2HIINT // __DOUBLE2HIINT: CUDA API: @@ -34,49 +34,49 @@ // __DOUBLE2INT_RD: CUDA API: // __DOUBLE2INT_RD-NEXT: __double2int_rd(d /*double*/); // __DOUBLE2INT_RD-NEXT: Is migrated to: -// __DOUBLE2INT_RD-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2INT_RD-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2int_rn | FileCheck %s -check-prefix=__DOUBLE2INT_RN // __DOUBLE2INT_RN: CUDA API: // __DOUBLE2INT_RN-NEXT: __double2int_rn(d /*double*/); // __DOUBLE2INT_RN-NEXT: Is migrated to: -// __DOUBLE2INT_RN-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2INT_RN-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2int_ru | FileCheck %s -check-prefix=__DOUBLE2INT_RU // __DOUBLE2INT_RU: CUDA API: // __DOUBLE2INT_RU-NEXT: __double2int_ru(d /*double*/); // __DOUBLE2INT_RU-NEXT: Is migrated to: -// __DOUBLE2INT_RU-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2INT_RU-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2int_rz | FileCheck %s -check-prefix=__DOUBLE2INT_RZ // __DOUBLE2INT_RZ: CUDA API: // __DOUBLE2INT_RZ-NEXT: __double2int_rz(d /*double*/); // __DOUBLE2INT_RZ-NEXT: Is migrated to: -// __DOUBLE2INT_RZ-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2INT_RZ-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2ll_rd | FileCheck %s -check-prefix=__DOUBLE2LL_RD // __DOUBLE2LL_RD: CUDA API: // __DOUBLE2LL_RD-NEXT: __double2ll_rd(d /*double*/); // __DOUBLE2LL_RD-NEXT: Is migrated to: -// __DOUBLE2LL_RD-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2LL_RD-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2ll_rn | FileCheck %s -check-prefix=__DOUBLE2LL_RN // __DOUBLE2LL_RN: CUDA API: // __DOUBLE2LL_RN-NEXT: __double2ll_rn(d /*double*/); // __DOUBLE2LL_RN-NEXT: Is migrated to: -// __DOUBLE2LL_RN-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2LL_RN-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2ll_ru | FileCheck %s -check-prefix=__DOUBLE2LL_RU // __DOUBLE2LL_RU: CUDA API: // __DOUBLE2LL_RU-NEXT: __double2ll_ru(d /*double*/); // __DOUBLE2LL_RU-NEXT: Is migrated to: -// __DOUBLE2LL_RU-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2LL_RU-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2ll_rz | FileCheck %s -check-prefix=__DOUBLE2LL_RZ // __DOUBLE2LL_RZ: CUDA API: // __DOUBLE2LL_RZ-NEXT: __double2ll_rz(d /*double*/); // __DOUBLE2LL_RZ-NEXT: Is migrated to: -// __DOUBLE2LL_RZ-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2LL_RZ-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2loint | FileCheck %s -check-prefix=__DOUBLE2LOINT // __DOUBLE2LOINT: CUDA API: @@ -88,49 +88,49 @@ // __DOUBLE2UINT_RD: CUDA API: // __DOUBLE2UINT_RD-NEXT: __double2uint_rd(d /*double*/); // __DOUBLE2UINT_RD-NEXT: Is migrated to: -// __DOUBLE2UINT_RD-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2UINT_RD-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2uint_rn | FileCheck %s -check-prefix=__DOUBLE2UINT_RN // __DOUBLE2UINT_RN: CUDA API: // __DOUBLE2UINT_RN-NEXT: __double2uint_rn(d /*double*/); // __DOUBLE2UINT_RN-NEXT: Is migrated to: -// __DOUBLE2UINT_RN-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2UINT_RN-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2uint_ru | FileCheck %s -check-prefix=__DOUBLE2UINT_RU // __DOUBLE2UINT_RU: CUDA API: // __DOUBLE2UINT_RU-NEXT: __double2uint_ru(d /*double*/); // __DOUBLE2UINT_RU-NEXT: Is migrated to: -// __DOUBLE2UINT_RU-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2UINT_RU-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2uint_rz | FileCheck %s -check-prefix=__DOUBLE2UINT_RZ // __DOUBLE2UINT_RZ: CUDA API: // __DOUBLE2UINT_RZ-NEXT: __double2uint_rz(d /*double*/); // __DOUBLE2UINT_RZ-NEXT: Is migrated to: -// __DOUBLE2UINT_RZ-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2UINT_RZ-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2ull_rd | FileCheck %s -check-prefix=__DOUBLE2ULL_RD // __DOUBLE2ULL_RD: CUDA API: // __DOUBLE2ULL_RD-NEXT: __double2ull_rd(d /*double*/); // __DOUBLE2ULL_RD-NEXT: Is migrated to: -// __DOUBLE2ULL_RD-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2ULL_RD-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2ull_rn | FileCheck %s -check-prefix=__DOUBLE2ULL_RN // __DOUBLE2ULL_RN: CUDA API: // __DOUBLE2ULL_RN-NEXT: __double2ull_rn(d /*double*/); // __DOUBLE2ULL_RN-NEXT: Is migrated to: -// __DOUBLE2ULL_RN-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2ULL_RN-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2ull_ru | FileCheck %s -check-prefix=__DOUBLE2ULL_RU // __DOUBLE2ULL_RU: CUDA API: // __DOUBLE2ULL_RU-NEXT: __double2ull_ru(d /*double*/); // __DOUBLE2ULL_RU-NEXT: Is migrated to: -// __DOUBLE2ULL_RU-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2ULL_RU-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double2ull_rz | FileCheck %s -check-prefix=__DOUBLE2ULL_RZ // __DOUBLE2ULL_RZ: CUDA API: // __DOUBLE2ULL_RZ-NEXT: __double2ull_rz(d /*double*/); // __DOUBLE2ULL_RZ-NEXT: Is migrated to: -// __DOUBLE2ULL_RZ-NEXT: sycl::vec{d}.convert()[0]; +// __DOUBLE2ULL_RZ-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__double_as_longlong | FileCheck %s -check-prefix=__DOUBLE_AS_LONGLONG // __DOUBLE_AS_LONGLONG: CUDA API: @@ -142,97 +142,97 @@ // __FLOAT2INT_RD: CUDA API: // __FLOAT2INT_RD-NEXT: __float2int_rd(d /*float*/); // __FLOAT2INT_RD-NEXT: Is migrated to: -// __FLOAT2INT_RD-NEXT: sycl::vec{d}.convert()[0]; +// __FLOAT2INT_RD-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__float2int_rn | FileCheck %s -check-prefix=__FLOAT2INT_RN // __FLOAT2INT_RN: CUDA API: // __FLOAT2INT_RN-NEXT: __float2int_rn(d /*float*/); // __FLOAT2INT_RN-NEXT: Is migrated to: -// __FLOAT2INT_RN-NEXT: sycl::vec{d}.convert()[0]; +// __FLOAT2INT_RN-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__float2int_ru | FileCheck %s -check-prefix=__FLOAT2INT_RU // __FLOAT2INT_RU: CUDA API: // __FLOAT2INT_RU-NEXT: __float2int_ru(d /*float*/); // __FLOAT2INT_RU-NEXT: Is migrated to: -// __FLOAT2INT_RU-NEXT: sycl::vec{d}.convert()[0]; +// __FLOAT2INT_RU-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__float2int_rz | FileCheck %s -check-prefix=__FLOAT2INT_RZ // __FLOAT2INT_RZ: CUDA API: // __FLOAT2INT_RZ-NEXT: __float2int_rz(d /*float*/); // __FLOAT2INT_RZ-NEXT: Is migrated to: -// __FLOAT2INT_RZ-NEXT: sycl::vec{d}.convert()[0]; +// __FLOAT2INT_RZ-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__float2ll_rd | FileCheck %s -check-prefix=__FLOAT2LL_RD // __FLOAT2LL_RD: CUDA API: // __FLOAT2LL_RD-NEXT: __float2ll_rd(d /*float*/); // __FLOAT2LL_RD-NEXT: Is migrated to: -// __FLOAT2LL_RD-NEXT: sycl::vec{d}.convert()[0]; +// __FLOAT2LL_RD-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__float2ll_rn | FileCheck %s -check-prefix=__FLOAT2LL_RN // __FLOAT2LL_RN: CUDA API: // __FLOAT2LL_RN-NEXT: __float2ll_rn(d /*float*/); // __FLOAT2LL_RN-NEXT: Is migrated to: -// __FLOAT2LL_RN-NEXT: sycl::vec{d}.convert()[0]; +// __FLOAT2LL_RN-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__float2ll_ru | FileCheck %s -check-prefix=__FLOAT2LL_RU // __FLOAT2LL_RU: CUDA API: // __FLOAT2LL_RU-NEXT: __float2ll_ru(d /*float*/); // __FLOAT2LL_RU-NEXT: Is migrated to: -// __FLOAT2LL_RU-NEXT: sycl::vec{d}.convert()[0]; +// __FLOAT2LL_RU-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__float2ll_rz | FileCheck %s -check-prefix=__FLOAT2LL_RZ // __FLOAT2LL_RZ: CUDA API: // __FLOAT2LL_RZ-NEXT: __float2ll_rz(d /*float*/); // __FLOAT2LL_RZ-NEXT: Is migrated to: -// __FLOAT2LL_RZ-NEXT: sycl::vec{d}.convert()[0]; +// __FLOAT2LL_RZ-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__float2uint_rd | FileCheck %s -check-prefix=__FLOAT2UINT_RD // __FLOAT2UINT_RD: CUDA API: // __FLOAT2UINT_RD-NEXT: __float2uint_rd(d /*float*/); // __FLOAT2UINT_RD-NEXT: Is migrated to: -// __FLOAT2UINT_RD-NEXT: sycl::vec{d}.convert()[0]; +// __FLOAT2UINT_RD-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__float2uint_rn | FileCheck %s -check-prefix=__FLOAT2UINT_RN // __FLOAT2UINT_RN: CUDA API: // __FLOAT2UINT_RN-NEXT: __float2uint_rn(d /*float*/); // __FLOAT2UINT_RN-NEXT: Is migrated to: -// __FLOAT2UINT_RN-NEXT: sycl::vec{d}.convert()[0]; +// __FLOAT2UINT_RN-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__float2uint_ru | FileCheck %s -check-prefix=__FLOAT2UINT_RU // __FLOAT2UINT_RU: CUDA API: // __FLOAT2UINT_RU-NEXT: __float2uint_ru(d /*float*/); // __FLOAT2UINT_RU-NEXT: Is migrated to: -// __FLOAT2UINT_RU-NEXT: sycl::vec{d}.convert()[0]; +// __FLOAT2UINT_RU-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__float2uint_rz | FileCheck %s -check-prefix=__FLOAT2UINT_RZ // __FLOAT2UINT_RZ: CUDA API: // __FLOAT2UINT_RZ-NEXT: __float2uint_rz(d /*float*/); // __FLOAT2UINT_RZ-NEXT: Is migrated to: -// __FLOAT2UINT_RZ-NEXT: sycl::vec{d}.convert()[0]; +// __FLOAT2UINT_RZ-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__float2ull_rd | FileCheck %s -check-prefix=__FLOAT2ULL_RD // __FLOAT2ULL_RD: CUDA API: // __FLOAT2ULL_RD-NEXT: __float2ull_rd(d /*float*/); // __FLOAT2ULL_RD-NEXT: Is migrated to: -// __FLOAT2ULL_RD-NEXT: sycl::vec{d}.convert()[0]; +// __FLOAT2ULL_RD-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__float2ull_rn | FileCheck %s -check-prefix=__FLOAT2ULL_RN // __FLOAT2ULL_RN: CUDA API: // __FLOAT2ULL_RN-NEXT: __float2ull_rn(d /*float*/); // __FLOAT2ULL_RN-NEXT: Is migrated to: -// __FLOAT2ULL_RN-NEXT: sycl::vec{d}.convert()[0]; +// __FLOAT2ULL_RN-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__float2ull_ru | FileCheck %s -check-prefix=__FLOAT2ULL_RU // __FLOAT2ULL_RU: CUDA API: // __FLOAT2ULL_RU-NEXT: __float2ull_ru(d /*float*/); // __FLOAT2ULL_RU-NEXT: Is migrated to: -// __FLOAT2ULL_RU-NEXT: sycl::vec{d}.convert()[0]; +// __FLOAT2ULL_RU-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__float2ull_rz | FileCheck %s -check-prefix=__FLOAT2ULL_RZ // __FLOAT2ULL_RZ: CUDA API: // __FLOAT2ULL_RZ-NEXT: __float2ull_rz(d /*float*/); // __FLOAT2ULL_RZ-NEXT: Is migrated to: -// __FLOAT2ULL_RZ-NEXT: sycl::vec{d}.convert()[0]; +// __FLOAT2ULL_RZ-NEXT: sycl::vec{d}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__float_as_int | FileCheck %s -check-prefix=__FLOAT_AS_INT // __FLOAT_AS_INT: CUDA API: @@ -256,31 +256,31 @@ // __INT2DOUBLE_RN: CUDA API: // __INT2DOUBLE_RN-NEXT: __int2double_rn(i /*int*/); // __INT2DOUBLE_RN-NEXT: Is migrated to: -// __INT2DOUBLE_RN-NEXT: sycl::vec{i}.convert()[0]; +// __INT2DOUBLE_RN-NEXT: sycl::vec{i}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__int2float_rd | FileCheck %s -check-prefix=__INT2FLOAT_RD // __INT2FLOAT_RD: CUDA API: // __INT2FLOAT_RD-NEXT: __int2float_rd(i /*int*/); // __INT2FLOAT_RD-NEXT: Is migrated to: -// __INT2FLOAT_RD-NEXT: sycl::vec{i}.convert()[0]; +// __INT2FLOAT_RD-NEXT: sycl::vec{i}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__int2float_rn | FileCheck %s -check-prefix=__INT2FLOAT_RN // __INT2FLOAT_RN: CUDA API: // __INT2FLOAT_RN-NEXT: __int2float_rn(i /*int*/); // __INT2FLOAT_RN-NEXT: Is migrated to: -// __INT2FLOAT_RN-NEXT: sycl::vec{i}.convert()[0]; +// __INT2FLOAT_RN-NEXT: sycl::vec{i}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__int2float_ru | FileCheck %s -check-prefix=__INT2FLOAT_RU // __INT2FLOAT_RU: CUDA API: // __INT2FLOAT_RU-NEXT: __int2float_ru(i /*int*/); // __INT2FLOAT_RU-NEXT: Is migrated to: -// __INT2FLOAT_RU-NEXT: sycl::vec{i}.convert()[0]; +// __INT2FLOAT_RU-NEXT: sycl::vec{i}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__int2float_rz | FileCheck %s -check-prefix=__INT2FLOAT_RZ // __INT2FLOAT_RZ: CUDA API: // __INT2FLOAT_RZ-NEXT: __int2float_rz(i /*int*/); // __INT2FLOAT_RZ-NEXT: Is migrated to: -// __INT2FLOAT_RZ-NEXT: sycl::vec{i}.convert()[0]; +// __INT2FLOAT_RZ-NEXT: sycl::vec{i}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__int_as_float | FileCheck %s -check-prefix=__INT_AS_FLOAT // __INT_AS_FLOAT: CUDA API: @@ -292,49 +292,49 @@ // __LL2DOUBLE_RD: CUDA API: // __LL2DOUBLE_RD-NEXT: __ll2double_rd(ll /*long long int*/); // __LL2DOUBLE_RD-NEXT: Is migrated to: -// __LL2DOUBLE_RD-NEXT: sycl::vec{ll}.convert()[0]; +// __LL2DOUBLE_RD-NEXT: sycl::vec{ll}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__ll2double_rn | FileCheck %s -check-prefix=__LL2DOUBLE_RN // __LL2DOUBLE_RN: CUDA API: // __LL2DOUBLE_RN-NEXT: __ll2double_rn(ll /*long long int*/); // __LL2DOUBLE_RN-NEXT: Is migrated to: -// __LL2DOUBLE_RN-NEXT: sycl::vec{ll}.convert()[0]; +// __LL2DOUBLE_RN-NEXT: sycl::vec{ll}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__ll2double_ru | FileCheck %s -check-prefix=__LL2DOUBLE_RU // __LL2DOUBLE_RU: CUDA API: // __LL2DOUBLE_RU-NEXT: __ll2double_ru(ll /*long long int*/); // __LL2DOUBLE_RU-NEXT: Is migrated to: -// __LL2DOUBLE_RU-NEXT: sycl::vec{ll}.convert()[0]; +// __LL2DOUBLE_RU-NEXT: sycl::vec{ll}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__ll2double_rz | FileCheck %s -check-prefix=__LL2DOUBLE_RZ // __LL2DOUBLE_RZ: CUDA API: // __LL2DOUBLE_RZ-NEXT: __ll2double_rz(ll /*long long int*/); // __LL2DOUBLE_RZ-NEXT: Is migrated to: -// __LL2DOUBLE_RZ-NEXT: sycl::vec{ll}.convert()[0]; +// __LL2DOUBLE_RZ-NEXT: sycl::vec{ll}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__ll2float_rd | FileCheck %s -check-prefix=__LL2FLOAT_RD // __LL2FLOAT_RD: CUDA API: // __LL2FLOAT_RD-NEXT: __ll2float_rd(ll /*long long int*/); // __LL2FLOAT_RD-NEXT: Is migrated to: -// __LL2FLOAT_RD-NEXT: sycl::vec{ll}.convert()[0]; +// __LL2FLOAT_RD-NEXT: sycl::vec{ll}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__ll2float_rn | FileCheck %s -check-prefix=__LL2FLOAT_RN // __LL2FLOAT_RN: CUDA API: // __LL2FLOAT_RN-NEXT: __ll2float_rn(ll /*long long int*/); // __LL2FLOAT_RN-NEXT: Is migrated to: -// __LL2FLOAT_RN-NEXT: sycl::vec{ll}.convert()[0]; +// __LL2FLOAT_RN-NEXT: sycl::vec{ll}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__ll2float_ru | FileCheck %s -check-prefix=__LL2FLOAT_RU // __LL2FLOAT_RU: CUDA API: // __LL2FLOAT_RU-NEXT: __ll2float_ru(ll /*long long int*/); // __LL2FLOAT_RU-NEXT: Is migrated to: -// __LL2FLOAT_RU-NEXT: sycl::vec{ll}.convert()[0]; +// __LL2FLOAT_RU-NEXT: sycl::vec{ll}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__ll2float_rz | FileCheck %s -check-prefix=__LL2FLOAT_RZ // __LL2FLOAT_RZ: CUDA API: // __LL2FLOAT_RZ-NEXT: __ll2float_rz(ll /*long long int*/); // __LL2FLOAT_RZ-NEXT: Is migrated to: -// __LL2FLOAT_RZ-NEXT: sycl::vec{ll}.convert()[0]; +// __LL2FLOAT_RZ-NEXT: sycl::vec{ll}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__longlong_as_double | FileCheck %s -check-prefix=__LONGLONG_AS_DOUBLE // __LONGLONG_AS_DOUBLE: CUDA API: @@ -346,31 +346,31 @@ // __UINT2DOUBLE_RN: CUDA API: // __UINT2DOUBLE_RN-NEXT: __uint2double_rn(u /*unsigned int*/); // __UINT2DOUBLE_RN-NEXT: Is migrated to: -// __UINT2DOUBLE_RN-NEXT: sycl::vec{u}.convert()[0]; +// __UINT2DOUBLE_RN-NEXT: sycl::vec{u}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__uint2float_rd | FileCheck %s -check-prefix=__UINT2FLOAT_RD // __UINT2FLOAT_RD: CUDA API: // __UINT2FLOAT_RD-NEXT: __uint2float_rd(u /*unsigned int*/); // __UINT2FLOAT_RD-NEXT: Is migrated to: -// __UINT2FLOAT_RD-NEXT: sycl::vec{u}.convert()[0]; +// __UINT2FLOAT_RD-NEXT: sycl::vec{u}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__uint2float_rn | FileCheck %s -check-prefix=__UINT2FLOAT_RN // __UINT2FLOAT_RN: CUDA API: // __UINT2FLOAT_RN-NEXT: __uint2float_rn(u /*unsigned int*/); // __UINT2FLOAT_RN-NEXT: Is migrated to: -// __UINT2FLOAT_RN-NEXT: sycl::vec{u}.convert()[0]; +// __UINT2FLOAT_RN-NEXT: sycl::vec{u}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__uint2float_ru | FileCheck %s -check-prefix=__UINT2FLOAT_RU // __UINT2FLOAT_RU: CUDA API: // __UINT2FLOAT_RU-NEXT: __uint2float_ru(u /*unsigned int*/); // __UINT2FLOAT_RU-NEXT: Is migrated to: -// __UINT2FLOAT_RU-NEXT: sycl::vec{u}.convert()[0]; +// __UINT2FLOAT_RU-NEXT: sycl::vec{u}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__uint2float_rz | FileCheck %s -check-prefix=__UINT2FLOAT_RZ // __UINT2FLOAT_RZ: CUDA API: // __UINT2FLOAT_RZ-NEXT: __uint2float_rz(u /*unsigned int*/); // __UINT2FLOAT_RZ-NEXT: Is migrated to: -// __UINT2FLOAT_RZ-NEXT: sycl::vec{u}.convert()[0]; +// __UINT2FLOAT_RZ-NEXT: sycl::vec{u}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__uint_as_float | FileCheck %s -check-prefix=__UINT_AS_FLOAT // __UINT_AS_FLOAT: CUDA API: @@ -382,46 +382,46 @@ // __ULL2DOUBLE_RD: CUDA API: // __ULL2DOUBLE_RD-NEXT: __ull2double_rd(ull /*unsigned long long int*/); // __ULL2DOUBLE_RD-NEXT: Is migrated to: -// __ULL2DOUBLE_RD-NEXT: sycl::vec{ull}.convert()[0]; +// __ULL2DOUBLE_RD-NEXT: sycl::vec{ull}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__ull2double_rn | FileCheck %s -check-prefix=__ULL2DOUBLE_RN // __ULL2DOUBLE_RN: CUDA API: // __ULL2DOUBLE_RN-NEXT: __ull2double_rn(ull /*unsigned long long int*/); // __ULL2DOUBLE_RN-NEXT: Is migrated to: -// __ULL2DOUBLE_RN-NEXT: sycl::vec{ull}.convert()[0]; +// __ULL2DOUBLE_RN-NEXT: sycl::vec{ull}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__ull2double_ru | FileCheck %s -check-prefix=__ULL2DOUBLE_RU // __ULL2DOUBLE_RU: CUDA API: // __ULL2DOUBLE_RU-NEXT: __ull2double_ru(ull /*unsigned long long int*/); // __ULL2DOUBLE_RU-NEXT: Is migrated to: -// __ULL2DOUBLE_RU-NEXT: sycl::vec{ull}.convert()[0]; +// __ULL2DOUBLE_RU-NEXT: sycl::vec{ull}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__ull2double_rz | FileCheck %s -check-prefix=__ULL2DOUBLE_RZ // __ULL2DOUBLE_RZ: CUDA API: // __ULL2DOUBLE_RZ-NEXT: __ull2double_rz(ull /*unsigned long long int*/); // __ULL2DOUBLE_RZ-NEXT: Is migrated to: -// __ULL2DOUBLE_RZ-NEXT: sycl::vec{ull}.convert()[0]; +// __ULL2DOUBLE_RZ-NEXT: sycl::vec{ull}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__ull2float_rd | FileCheck %s -check-prefix=__ULL2FLOAT_RD // __ULL2FLOAT_RD: CUDA API: // __ULL2FLOAT_RD-NEXT: __ull2float_rd(ull /*unsigned long long int*/); // __ULL2FLOAT_RD-NEXT: Is migrated to: -// __ULL2FLOAT_RD-NEXT: sycl::vec{ull}.convert()[0]; +// __ULL2FLOAT_RD-NEXT: sycl::vec{ull}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__ull2float_rn | FileCheck %s -check-prefix=__ULL2FLOAT_RN // __ULL2FLOAT_RN: CUDA API: // __ULL2FLOAT_RN-NEXT: __ull2float_rn(ull /*unsigned long long int*/); // __ULL2FLOAT_RN-NEXT: Is migrated to: -// __ULL2FLOAT_RN-NEXT: sycl::vec{ull}.convert()[0]; +// __ULL2FLOAT_RN-NEXT: sycl::vec{ull}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__ull2float_ru | FileCheck %s -check-prefix=__ULL2FLOAT_RU // __ULL2FLOAT_RU: CUDA API: // __ULL2FLOAT_RU-NEXT: __ull2float_ru(ull /*unsigned long long int*/); // __ULL2FLOAT_RU-NEXT: Is migrated to: -// __ULL2FLOAT_RU-NEXT: sycl::vec{ull}.convert()[0]; +// __ULL2FLOAT_RU-NEXT: sycl::vec{ull}.template convert()[0]; // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__ull2float_rz | FileCheck %s -check-prefix=__ULL2FLOAT_RZ // __ULL2FLOAT_RZ: CUDA API: // __ULL2FLOAT_RZ-NEXT: __ull2float_rz(ull /*unsigned long long int*/); // __ULL2FLOAT_RZ-NEXT: Is migrated to: -// __ULL2FLOAT_RZ-NEXT: sycl::vec{ull}.convert()[0]; +// __ULL2FLOAT_RZ-NEXT: sycl::vec{ull}.template convert()[0]; diff --git a/clang/test/dpct/query_api_mapping/NoLib/test.cu b/clang/test/dpct/query_api_mapping/NoLib/test.cu index a6171d9fb5ef..4df53a6da16b 100644 --- a/clang/test/dpct/query_api_mapping/NoLib/test.cu +++ b/clang/test/dpct/query_api_mapping/NoLib/test.cu @@ -89,13 +89,13 @@ // CUCOMPLEXDOUBLETOFLOAT: CUDA API: // CUCOMPLEXDOUBLETOFLOAT-NEXT: cuComplexDoubleToFloat(c /*cuDoubleComplex*/); // CUCOMPLEXDOUBLETOFLOAT-NEXT: Is migrated to: -// CUCOMPLEXDOUBLETOFLOAT-NEXT: c.convert(); +// CUCOMPLEXDOUBLETOFLOAT-NEXT: c.template convert(); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuComplexFloatToDouble | FileCheck %s -check-prefix=CUCOMPLEXFLOATTODOUBLE // CUCOMPLEXFLOATTODOUBLE: CUDA API: // CUCOMPLEXFLOATTODOUBLE-NEXT: cuComplexFloatToDouble(c /*cuFloatComplex*/); // CUCOMPLEXFLOATTODOUBLE-NEXT: Is migrated to: -// CUCOMPLEXFLOATTODOUBLE-NEXT: c.convert(); +// CUCOMPLEXFLOATTODOUBLE-NEXT: c.template convert(); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuConj | FileCheck %s -check-prefix=CUCONJ // CUCONJ: CUDA API: