Skip to content

Commit 8926314

Browse files
tangjj11tangjj11
authored
[SYCLomatic] Fix the migration of __low2float and __high2float. (#889)
 Signed-off-by: Tang, Jiajun [email protected]
1 parent e9efb05 commit 8926314

File tree

4 files changed

+30
-33
lines changed

4 files changed

+30
-33
lines changed

‎clang/lib/DPCT/APINamesMath.inc‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -388,7 +388,7 @@ ENTRY_TYPECAST("__half2ushort_rn")
388388
ENTRY_TYPECAST("__half2ushort_ru")
389389
ENTRY_TYPECAST("__half2ushort_rz")
390390
ENTRY_TYPECAST("__halves2half2")
391-
ENTRY_TYPECAST("__high2float")
391+
ENTRY_REWRITE("__high2float")
392392
ENTRY_TYPECAST("__high2half")
393393
ENTRY_TYPECAST("__high2half2")
394394
ENTRY_TYPECAST("__highs2half2")
@@ -406,7 +406,7 @@ ENTRY_TYPECAST("__ll2half_rd")
406406
ENTRY_TYPECAST("__ll2half_rn")
407407
ENTRY_TYPECAST("__ll2half_ru")
408408
ENTRY_TYPECAST("__ll2half_rz")
409-
ENTRY_TYPECAST("__low2float")
409+
ENTRY_REWRITE("__low2float")
410410
ENTRY_TYPECAST("__low2half")
411411
ENTRY_TYPECAST("__low2half2")
412412
ENTRY_TYPECAST("__lowhigh2highlow")

‎clang/lib/DPCT/APINamesMathRewrite.inc‎

Lines changed: 25 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -729,8 +729,8 @@ MATH_API_REWRITER_DEVICE(
729729
FEATURE_REQUEST_FACTORY(
730730
HelperFeatureEnum::Math_compare2,
731731
CALL_FACTORY_ENTRY(
732-
"__heq2", CALL(MapNames::getDpctNamespace() + "compare",
733-
ARG(0), ARG(1), LITERAL("std::equal_to<>()"))))))
732+
"__heq2", CALL(MapNames::getDpctNamespace() + "compare", ARG(0),
733+
ARG(1), LITERAL("std::equal_to<>()"))))))
734734

735735
MATH_API_REWRITER_DEVICE(
736736
"__hequ2",
@@ -761,10 +761,9 @@ MATH_API_REWRITER_DEVICE(
761761
EMPTY_FACTORY_ENTRY("__hge2"),
762762
FEATURE_REQUEST_FACTORY(
763763
HelperFeatureEnum::Math_compare2,
764-
CALL_FACTORY_ENTRY("__hge2",
765-
CALL(MapNames::getDpctNamespace() + "compare",
766-
ARG(0), ARG(1),
767-
LITERAL("std::greater_equal<>()"))))))
764+
CALL_FACTORY_ENTRY(
765+
"__hge2", CALL(MapNames::getDpctNamespace() + "compare", ARG(0),
766+
ARG(1), LITERAL("std::greater_equal<>()"))))))
768767

769768
MATH_API_REWRITER_DEVICE(
770769
"__hgeu2",
@@ -780,8 +779,8 @@ MATH_API_REWRITER_DEVICE(
780779
HelperFeatureEnum::Math_unordered_compare2,
781780
CALL_FACTORY_ENTRY(
782781
"__hgeu2",
783-
CALL(MapNames::getDpctNamespace() + "unordered_compare",
784-
ARG(0), ARG(1), LITERAL("std::greater_equal<>()"))))))
782+
CALL(MapNames::getDpctNamespace() + "unordered_compare", ARG(0),
783+
ARG(1), LITERAL("std::greater_equal<>()"))))))
785784

786785
MATH_API_REWRITER_DEVICE(
787786
"__hgt2",
@@ -796,8 +795,8 @@ MATH_API_REWRITER_DEVICE(
796795
FEATURE_REQUEST_FACTORY(
797796
HelperFeatureEnum::Math_compare2,
798797
CALL_FACTORY_ENTRY(
799-
"__hgt2", CALL(MapNames::getDpctNamespace() + "compare",
800-
ARG(0), ARG(1), LITERAL("std::greater<>()"))))))
798+
"__hgt2", CALL(MapNames::getDpctNamespace() + "compare", ARG(0),
799+
ARG(1), LITERAL("std::greater<>()"))))))
801800

802801
MATH_API_REWRITER_DEVICE(
803802
"__hgtu2",
@@ -845,10 +844,9 @@ MATH_API_REWRITER_DEVICE(
845844
EMPTY_FACTORY_ENTRY("__hle2"),
846845
FEATURE_REQUEST_FACTORY(
847846
HelperFeatureEnum::Math_compare2,
848-
CALL_FACTORY_ENTRY("__hle2",
849-
CALL(MapNames::getDpctNamespace() + "compare",
850-
ARG(0), ARG(1),
851-
LITERAL("std::less_equal<>()"))))))
847+
CALL_FACTORY_ENTRY(
848+
"__hle2", CALL(MapNames::getDpctNamespace() + "compare", ARG(0),
849+
ARG(1), LITERAL("std::less_equal<>()"))))))
852850

853851
MATH_API_REWRITER_DEVICE(
854852
"__hleu2",
@@ -864,8 +862,8 @@ MATH_API_REWRITER_DEVICE(
864862
HelperFeatureEnum::Math_unordered_compare2,
865863
CALL_FACTORY_ENTRY(
866864
"__hleu2",
867-
CALL(MapNames::getDpctNamespace() + "unordered_compare",
868-
ARG(0), ARG(1), LITERAL("std::less_equal<>()"))))))
865+
CALL(MapNames::getDpctNamespace() + "unordered_compare", ARG(0),
866+
ARG(1), LITERAL("std::less_equal<>()"))))))
869867

870868
MATH_API_REWRITER_DEVICE(
871869
"__hlt2",
@@ -880,8 +878,8 @@ MATH_API_REWRITER_DEVICE(
880878
FEATURE_REQUEST_FACTORY(
881879
HelperFeatureEnum::Math_compare2,
882880
CALL_FACTORY_ENTRY(
883-
"__hlt2", CALL(MapNames::getDpctNamespace() + "compare",
884-
ARG(0), ARG(1), LITERAL("std::less<>()"))))))
881+
"__hlt2", CALL(MapNames::getDpctNamespace() + "compare", ARG(0),
882+
ARG(1), LITERAL("std::less<>()"))))))
885883

886884
MATH_API_REWRITER_DEVICE(
887885
"__hltu2",
@@ -912,10 +910,9 @@ MATH_API_REWRITER_DEVICE(
912910
EMPTY_FACTORY_ENTRY("__hne2"),
913911
FEATURE_REQUEST_FACTORY(
914912
HelperFeatureEnum::Math_compare2,
915-
CALL_FACTORY_ENTRY("__hne2",
916-
CALL(MapNames::getDpctNamespace() + "compare",
917-
ARG(0), ARG(1),
918-
LITERAL("std::not_equal_to<>()"))))))
913+
CALL_FACTORY_ENTRY(
914+
"__hne2", CALL(MapNames::getDpctNamespace() + "compare", ARG(0),
915+
ARG(1), LITERAL("std::not_equal_to<>()"))))))
919916

920917
MATH_API_REWRITER_DEVICE(
921918
"__hneu2",
@@ -931,10 +928,12 @@ MATH_API_REWRITER_DEVICE(
931928
HelperFeatureEnum::Math_unordered_compare2,
932929
CALL_FACTORY_ENTRY(
933930
"__hneu2",
934-
CALL(MapNames::getDpctNamespace() + "unordered_compare",
935-
ARG(0), ARG(1), LITERAL("std::not_equal_to<>()"))))))
931+
CALL(MapNames::getDpctNamespace() + "unordered_compare", ARG(0),
932+
ARG(1), LITERAL("std::not_equal_to<>()"))))))
936933

937934
// Half Precision Conversion and Data Movement
935+
ARRAYSUBSCRIPT_EXPR_FACTORY_ENTRY("__high2float", ARG(0), LITERAL("1"))
936+
938937
MATH_API_REWRITER_DEVICE(
939938
"__ldca",
940939
MATH_API_DEVICE_NODES(
@@ -989,6 +988,8 @@ MATH_API_REWRITER_DEVICE(
989988
Diagnostics::MATH_EMULATION_EXPRESSION,
990989
std::string("__ldlu"), std::string("'*'"))))
991990

991+
ARRAYSUBSCRIPT_EXPR_FACTORY_ENTRY("__low2float", ARG(0), LITERAL("0"))
992+
992993
MATH_API_REWRITER_DEVICE(
993994
"__stcg",
994995
MATH_API_DEVICE_NODES(

‎clang/lib/DPCT/CallExprRewriterMath.cpp‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -438,8 +438,6 @@ std::optional<std::string> MathTypeCastRewriter::rewrite() {
438438
auto MigratedArg1 = getMigratedArg(1);
439439
OS << MapNames::getClNamespace() + "half2{" << MigratedArg0 << ","
440440
<< MigratedArg1 << "}";
441-
} else if (FuncName == "__high2float") {
442-
OS << MigratedArg0 << "[0]";
443441
} else if (FuncName == "__high2half") {
444442
OS << MigratedArg0 << "[0]";
445443
} else if (FuncName == "__high2half2") {
@@ -449,8 +447,6 @@ std::optional<std::string> MathTypeCastRewriter::rewrite() {
449447
auto MigratedArg1 = getMigratedArgWithExtraParens(1);
450448
OS << MapNames::getClNamespace() + "half2{" << MigratedArg0 << "[0], "
451449
<< MigratedArg1 << "[0]}";
452-
} else if (FuncName == "__low2float") {
453-
OS << MigratedArg0 << "[1]";
454450
} else if (FuncName == "__low2half") {
455451
OS << MigratedArg0 << "[1]";
456452
} else if (FuncName == "__low2half2") {

‎clang/test/dpct/cuda-math-intrinsics.cu‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1732,7 +1732,7 @@ __global__ void kernelFuncTypecasts() {
17321732
// CHECK: h2 = sycl::half2{h,h};
17331733
h2 = __halves2half2(h, h);
17341734

1735-
// CHECK: f = h2[0];
1735+
// CHECK: f = h2[1];
17361736
f = __high2float(h2);
17371737

17381738
// CHECK: h = h2[0];
@@ -1768,9 +1768,9 @@ __global__ void kernelFuncTypecasts() {
17681768
// CHECK: h = sycl::vec<long long, 1>{ll}.convert<sycl::half, sycl::rounding_mode::rtz>()[0];
17691769
h = __ll2half_rz(ll);
17701770

1771-
// CHECK: f = h2[1];
1771+
// CHECK: f = h2[0];
17721772
f = __low2float(h2);
1773-
// CHECK: f = (*(&h2))[1];
1773+
// CHECK: f = *(&h2)[0];
17741774
f = __low2float(*(&h2));
17751775

17761776
// CHECK: h = h2[1];

0 commit comments

Comments
 (0)