Skip to content

Commit 4c57bc6

Browse files
zhiweij1zhiweij1
committed
[SYCLomatic] Always add "template" before sycl::vec::convert
Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>
1 parent c03712d commit 4c57bc6

12 files changed

Lines changed: 350 additions & 382 deletions

‎clang/lib/DPCT/RulesLang/APINamesComplex.inc‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -95,5 +95,7 @@ BINARY_OP_FACTORY_ENTRY("cuCfmaf", BinaryOperatorKind::BO_Add,
9595
makeCallArgCreatorWithCall(1)),
9696
makeCallArgCreatorWithCall(2))
9797

98-
MEMBER_CALL_FACTORY_ENTRY("cuComplexDoubleToFloat", ARG(0), false, "convert<float>")
99-
MEMBER_CALL_FACTORY_ENTRY("cuComplexFloatToDouble", ARG(0), false, "convert<double>")
98+
MEMBER_CALL_FACTORY_ENTRY("cuComplexDoubleToFloat", ARG(0), false,
99+
"template convert<float>")
100+
MEMBER_CALL_FACTORY_ENTRY("cuComplexFloatToDouble", ARG(0), false,
101+
"template convert<double>")

‎clang/lib/DPCT/RulesLang/Math/CallExprRewriterMath.cpp‎

Lines changed: 17 additions & 68 deletions
Original file line numberDiff line numberDiff line change
@@ -288,75 +288,24 @@ std::optional<std::string> MathTypeCastRewriter::rewrite() {
288288
const StringRef &FuncName = SourceCalleeName;
289289
std::string ReplStr;
290290
llvm::raw_string_ostream OS(ReplStr);
291-
292291
auto MigratedArg0 = getMigratedArgWithExtraParens(0);
293-
if (FuncName == "__float22half2_rn") {
294-
OS << MigratedArg0
295-
<< ".convert<" + MapNames::getClNamespace() + "half, " +
296-
MapNames::getClNamespace() + "rounding_mode::rte>()";
297-
} else if (FuncName == "__float2half2_rn") {
298-
OS << MapNames::getClNamespace() + "float2{" << MigratedArg0 << ","
299-
<< MigratedArg0
300-
<< "}.convert<" + MapNames::getClNamespace() + "half, " +
301-
MapNames::getClNamespace() + "rounding_mode::rte>()";
302-
} else if (FuncName == "__floats2half2_rn") {
303-
auto MigratedArg1 = getMigratedArg(1);
304-
OS << MapNames::getClNamespace() + "float2{" << MigratedArg0 << ","
305-
<< MigratedArg1
306-
<< "}.convert<" + MapNames::getClNamespace() + "half, " +
307-
MapNames::getClNamespace() + "rounding_mode::rte>()";
308-
} else if (FuncName == "__half22float2") {
309-
OS << MigratedArg0
310-
<< ".convert<float, " + MapNames::getClNamespace() +
311-
"rounding_mode::automatic>()";
312-
} else if (FuncName == "__half2half2") {
313-
OS << MapNames::getClNamespace() + "half2{" << MigratedArg0 << ","
314-
<< MigratedArg0 << "}";
315-
} else if (FuncName == "__halves2half2") {
316-
auto MigratedArg1 = getMigratedArg(1);
317-
OS << MapNames::getClNamespace() + "half2{" << MigratedArg0 << ","
318-
<< MigratedArg1 << "}";
319-
} else if (FuncName == "__high2half") {
320-
OS << MigratedArg0 << "[0]";
321-
} else if (FuncName == "__high2half2") {
322-
OS << MapNames::getClNamespace() + "half2{" << MigratedArg0 << "[0], "
323-
<< MigratedArg0 << "[0]}";
324-
} else if (FuncName == "__highs2half2") {
325-
auto MigratedArg1 = getMigratedArgWithExtraParens(1);
326-
OS << MapNames::getClNamespace() + "half2{" << MigratedArg0 << "[0], "
327-
<< MigratedArg1 << "[0]}";
328-
} else if (FuncName == "__low2half") {
329-
OS << MigratedArg0 << "[1]";
330-
} else if (FuncName == "__low2half2") {
331-
OS << MapNames::getClNamespace() + "half2{" << MigratedArg0 << "[1], "
332-
<< MigratedArg0 << "[1]}";
333-
} else if (FuncName == "__lowhigh2highlow") {
334-
OS << MapNames::getClNamespace() + "half2{" << MigratedArg0 << "[1], "
335-
<< MigratedArg0 << "[0]}";
336-
} else if (FuncName == "__lows2half2") {
337-
auto MigratedArg1 = getMigratedArgWithExtraParens(1);
338-
OS << MapNames::getClNamespace() + "half2{" << MigratedArg0 << "[1], "
339-
<< MigratedArg1 << "[1]}";
340-
} else {
341-
//__half2short_rd and __half2float
342-
static SSMap TypeMap{{"ll", "long long"},
343-
{"ull", "unsigned long long"},
344-
{"ushort", "unsigned short"},
345-
{"uint", "unsigned int"},
346-
{"half", MapNames::getClNamespace() + "half"}};
347-
std::string RoundingMode;
348-
if (FuncName[FuncName.size() - 3] == '_')
349-
RoundingMode = FuncName.substr(FuncName.size() - 2).str();
350-
auto FN = FuncName.substr(2, FuncName.find('_', 2) - 2).str();
351-
auto Types = split(FN, '2');
352-
assert(Types.size() == 2);
353-
MapNames::replaceName(TypeMap, Types[0]);
354-
MapNames::replaceName(TypeMap, Types[1]);
355-
OS << MapNames::getClNamespace() + "vec<" << Types[0] << ", 1>{"
356-
<< MigratedArg0 << "}.convert<" << Types[1]
357-
<< ", " + MapNames::getClNamespace() + "rounding_mode::"
358-
<< RoundingModeMap[RoundingMode] << ">()[0]";
359-
}
292+
static SSMap TypeMap{{"ll", "long long"},
293+
{"ull", "unsigned long long"},
294+
{"ushort", "unsigned short"},
295+
{"uint", "unsigned int"},
296+
{"half", MapNames::getClNamespace() + "half"}};
297+
std::string RoundingMode;
298+
if (FuncName[FuncName.size() - 3] == '_')
299+
RoundingMode = FuncName.substr(FuncName.size() - 2).str();
300+
auto FN = FuncName.substr(2, FuncName.find('_', 2) - 2).str();
301+
auto Types = split(FN, '2');
302+
assert(Types.size() == 2);
303+
MapNames::replaceName(TypeMap, Types[0]);
304+
MapNames::replaceName(TypeMap, Types[1]);
305+
OS << MapNames::getClNamespace() + "vec<" << Types[0] << ", 1>{"
306+
<< MigratedArg0 << "}.template convert<" << Types[1]
307+
<< ", " + MapNames::getClNamespace() + "rounding_mode::"
308+
<< RoundingModeMap[RoundingMode] << ">()[0]";
360309
OS.flush();
361310
return ReplStr;
362311
}

‎clang/lib/DPCT/RulesLang/Math/RewriterBfloat16PrecisionConversionAndDataMovement.cpp‎

Lines changed: 32 additions & 32 deletions
Original file line numberDiff line numberDiff line change
@@ -67,7 +67,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
6767
MEMBER_CALL(
6868
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
6969
false,
70-
"convert<int, " + MapNames::getClNamespace() +
70+
"template convert<int, " + MapNames::getClNamespace() +
7171
"rounding_mode::rtn>"),
7272
LITERAL("0")))
7373
// __bfloat162int_rn
@@ -84,7 +84,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
8484
MEMBER_CALL(
8585
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
8686
false,
87-
"convert<int, " + MapNames::getClNamespace() +
87+
"template convert<int, " + MapNames::getClNamespace() +
8888
"rounding_mode::rte>"),
8989
LITERAL("0")))
9090
// __bfloat162int_ru
@@ -101,7 +101,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
101101
MEMBER_CALL(
102102
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
103103
false,
104-
"convert<int, " + MapNames::getClNamespace() +
104+
"template convert<int, " + MapNames::getClNamespace() +
105105
"rounding_mode::rtp>"),
106106
LITERAL("0")))
107107
// __bfloat162int_rz
@@ -118,7 +118,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
118118
MEMBER_CALL(
119119
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
120120
false,
121-
"convert<int, " + MapNames::getClNamespace() +
121+
"template convert<int, " + MapNames::getClNamespace() +
122122
"rounding_mode::rtz>"),
123123
LITERAL("0")))
124124
// __bfloat162ll_rd
@@ -135,7 +135,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
135135
MEMBER_CALL(
136136
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
137137
false,
138-
"convert<long long, " + MapNames::getClNamespace() +
138+
"template convert<long long, " + MapNames::getClNamespace() +
139139
"rounding_mode::rtn>"),
140140
LITERAL("0")))
141141
// __bfloat162ll_rn
@@ -152,7 +152,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
152152
MEMBER_CALL(
153153
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
154154
false,
155-
"convert<long long, " + MapNames::getClNamespace() +
155+
"template convert<long long, " + MapNames::getClNamespace() +
156156
"rounding_mode::rte>"),
157157
LITERAL("0")))
158158
// __bfloat162ll_ru
@@ -169,7 +169,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
169169
MEMBER_CALL(
170170
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
171171
false,
172-
"convert<long long, " + MapNames::getClNamespace() +
172+
"template convert<long long, " + MapNames::getClNamespace() +
173173
"rounding_mode::rtp>"),
174174
LITERAL("0")))
175175
// __bfloat162ll_rz
@@ -186,7 +186,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
186186
MEMBER_CALL(
187187
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
188188
false,
189-
"convert<long long, " + MapNames::getClNamespace() +
189+
"template convert<long long, " + MapNames::getClNamespace() +
190190
"rounding_mode::rtz>"),
191191
LITERAL("0")))
192192
// __bfloat162short_rd
@@ -203,7 +203,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
203203
MEMBER_CALL(
204204
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
205205
false,
206-
"convert<short, " + MapNames::getClNamespace() +
206+
"template convert<short, " + MapNames::getClNamespace() +
207207
"rounding_mode::rtn>"),
208208
LITERAL("0")))
209209
// __bfloat162short_rn
@@ -220,7 +220,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
220220
MEMBER_CALL(
221221
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
222222
false,
223-
"convert<short, " + MapNames::getClNamespace() +
223+
"template convert<short, " + MapNames::getClNamespace() +
224224
"rounding_mode::rte>"),
225225
LITERAL("0")))
226226
// __bfloat162short_ru
@@ -237,7 +237,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
237237
MEMBER_CALL(
238238
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
239239
false,
240-
"convert<short, " + MapNames::getClNamespace() +
240+
"template convert<short, " + MapNames::getClNamespace() +
241241
"rounding_mode::rtp>"),
242242
LITERAL("0")))
243243
// __bfloat162short_rz
@@ -254,7 +254,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
254254
MEMBER_CALL(
255255
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
256256
false,
257-
"convert<short, " + MapNames::getClNamespace() +
257+
"template convert<short, " + MapNames::getClNamespace() +
258258
"rounding_mode::rtz>"),
259259
LITERAL("0")))
260260
// __bfloat162uint_rd
@@ -271,7 +271,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
271271
MEMBER_CALL(
272272
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
273273
false,
274-
"convert<unsigned, " + MapNames::getClNamespace() +
274+
"template convert<unsigned, " + MapNames::getClNamespace() +
275275
"rounding_mode::rtn>"),
276276
LITERAL("0")))
277277
// __bfloat162uint_rn
@@ -288,7 +288,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
288288
MEMBER_CALL(
289289
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
290290
false,
291-
"convert<unsigned, " + MapNames::getClNamespace() +
291+
"template convert<unsigned, " + MapNames::getClNamespace() +
292292
"rounding_mode::rte>"),
293293
LITERAL("0")))
294294
// __bfloat162uint_ru
@@ -305,7 +305,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
305305
MEMBER_CALL(
306306
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
307307
false,
308-
"convert<unsigned, " + MapNames::getClNamespace() +
308+
"template convert<unsigned, " + MapNames::getClNamespace() +
309309
"rounding_mode::rtp>"),
310310
LITERAL("0")))
311311
// __bfloat162uint_rz
@@ -322,7 +322,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
322322
MEMBER_CALL(
323323
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
324324
false,
325-
"convert<unsigned, " + MapNames::getClNamespace() +
325+
"template convert<unsigned, " + MapNames::getClNamespace() +
326326
"rounding_mode::rtz>"),
327327
LITERAL("0")))
328328
// __bfloat162ull_rd
@@ -339,8 +339,8 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
339339
MEMBER_CALL(
340340
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
341341
false,
342-
"convert<unsigned long long, " + MapNames::getClNamespace() +
343-
"rounding_mode::rtn>"),
342+
"template convert<unsigned long long, " +
343+
MapNames::getClNamespace() + "rounding_mode::rtn>"),
344344
LITERAL("0")))
345345
// __bfloat162ull_rn
346346
CONDITIONAL_FACTORY_ENTRY(
@@ -356,8 +356,8 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
356356
MEMBER_CALL(
357357
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
358358
false,
359-
"convert<unsigned long long, " + MapNames::getClNamespace() +
360-
"rounding_mode::rte>"),
359+
"template convert<unsigned long long, " +
360+
MapNames::getClNamespace() + "rounding_mode::rte>"),
361361
LITERAL("0")))
362362
// __bfloat162ull_ru
363363
CONDITIONAL_FACTORY_ENTRY(
@@ -373,8 +373,8 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
373373
MEMBER_CALL(
374374
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
375375
false,
376-
"convert<unsigned long long, " + MapNames::getClNamespace() +
377-
"rounding_mode::rtp>"),
376+
"template convert<unsigned long long, " +
377+
MapNames::getClNamespace() + "rounding_mode::rtp>"),
378378
LITERAL("0")))
379379
// __bfloat162ull_rz
380380
CONDITIONAL_FACTORY_ENTRY(
@@ -390,8 +390,8 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
390390
MEMBER_CALL(
391391
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
392392
false,
393-
"convert<unsigned long long, " + MapNames::getClNamespace() +
394-
"rounding_mode::rtz>"),
393+
"template convert<unsigned long long, " +
394+
MapNames::getClNamespace() + "rounding_mode::rtz>"),
395395
LITERAL("0")))
396396
// __bfloat162ushort_rd
397397
CONDITIONAL_FACTORY_ENTRY(
@@ -408,8 +408,8 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
408408
MEMBER_CALL(
409409
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
410410
false,
411-
"convert<unsigned short, " + MapNames::getClNamespace() +
412-
"rounding_mode::rtn>"),
411+
"template convert<unsigned short, " +
412+
MapNames::getClNamespace() + "rounding_mode::rtn>"),
413413
LITERAL("0")))
414414
// __bfloat162ushort_rn
415415
CONDITIONAL_FACTORY_ENTRY(
@@ -426,8 +426,8 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
426426
MEMBER_CALL(
427427
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
428428
false,
429-
"convert<unsigned short, " + MapNames::getClNamespace() +
430-
"rounding_mode::rte>"),
429+
"template convert<unsigned short, " +
430+
MapNames::getClNamespace() + "rounding_mode::rte>"),
431431
LITERAL("0")))
432432
// __bfloat162ushort_ru
433433
CONDITIONAL_FACTORY_ENTRY(
@@ -444,8 +444,8 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
444444
MEMBER_CALL(
445445
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
446446
false,
447-
"convert<unsigned short, " + MapNames::getClNamespace() +
448-
"rounding_mode::rtp>"),
447+
"template convert<unsigned short, " +
448+
MapNames::getClNamespace() + "rounding_mode::rtp>"),
449449
LITERAL("0")))
450450
// __bfloat162ushort_rz
451451
CONDITIONAL_FACTORY_ENTRY(
@@ -462,8 +462,8 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
462462
MEMBER_CALL(
463463
CALL(MapNames::getClNamespace() + "vec<float, 1>", ARG(0)),
464464
false,
465-
"convert<unsigned short, " + MapNames::getClNamespace() +
466-
"rounding_mode::rtz>"),
465+
"template convert<unsigned short, " +
466+
MapNames::getClNamespace() + "rounding_mode::rtz>"),
467467
LITERAL("0")))
468468
// __bfloat16_as_short
469469
CONDITIONAL_FACTORY_ENTRY(

0 commit comments

Comments
 (0)