From 0e870f6682784036eaffe937bffc3d5f37f1c63c Mon Sep 17 00:00:00 2001 From: Ziran Zhang Date: Wed, 19 Mar 2025 10:30:53 +0800 Subject: [PATCH] [SYCLomatic] Enable APIs that SYCLcompat already supported Signed-off-by: Ziran Zhang --- .../Math/RewriterHalf2ArithmeticFunctions.cpp | 42 ++------ .../Math/RewriterHalf2ComparisonFunctions.cpp | 98 +++---------------- .../Math/RewriterHalfComparisonFunctions.cpp | 14 +-- 3 files changed, 22 insertions(+), 132 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/Math/RewriterHalf2ArithmeticFunctions.cpp b/clang/lib/DPCT/RulesLang/Math/RewriterHalf2ArithmeticFunctions.cpp index 107343d74c21..0f9a563e3430 100644 --- a/clang/lib/DPCT/RulesLang/Math/RewriterHalf2ArithmeticFunctions.cpp +++ b/clang/lib/DPCT/RulesLang/Math/RewriterHalf2ArithmeticFunctions.cpp @@ -119,17 +119,12 @@ RewriterMap dpct::createHalf2ArithmeticFunctionsRewriterMap() { MapNames::getClNamespace() + "half2>", BO(BinaryOperatorKind::BO_Add, ARG(0), ARG(1)), LITERAL("{0.f, 0.f}"), LITERAL("{1.f, 1.f}"))), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hadd2_sat", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hadd2_sat")), CALL_FACTORY_ENTRY( "__hadd2_sat", CALL(MapNames::getDpctNamespace() + "clamp", BO(BinaryOperatorKind::BO_Add, ARG(0), ARG(1)), LITERAL("{0.f, 0.f}"), - LITERAL("{1.f, 1.f}"))))))) + LITERAL("{1.f, 1.f}")))))) // __hcmadd MATH_API_REWRITER_DEVICE( "__hcmadd", @@ -144,17 +139,12 @@ RewriterMap dpct::createHalf2ArithmeticFunctionsRewriterMap() { "ext::intel::math::hcmadd", ARG(0), ARG(1), ARG(2))))), EMPTY_FACTORY_ENTRY("__hcmadd"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hcmadd", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hcmadd")), CALL_FACTORY_ENTRY("__hcmadd", CALL(MapNames::getDpctNamespace() + (DpctGlobalInfo::useSYCLCompat() ? "cmul_add" : "complex_mul_add"), - ARG(0), ARG(1), ARG(2)))))) + ARG(0), ARG(1), ARG(2))))) // __hfma2 MATH_API_REWRITER_DEVICE_OVERLOAD( CheckArgType(0, "__half2"), @@ -236,23 +226,13 @@ RewriterMap dpct::createHalf2ArithmeticFunctionsRewriterMap() { LITERAL("{0.f, 0.f}"), LITERAL("{1.f, 1.f}"))))), MATH_API_REWRITER_EXPERIMENTAL_BFLOAT16( "__hfma2_sat", - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hfma2_sat", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hfma2_sat")), CALL_FACTORY_ENTRY( "__hfma2_sat", CALL(MapNames::getDpctNamespace() + "clamp", CALL(MapNames::getClNamespace(false, true) + "ext::oneapi::experimental::fma", ARG(0), ARG(1), ARG(2)), - LITERAL("{0.f, 0.f}"), LITERAL("{1.f, 1.f}")))), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hfma2_sat", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hfma2_sat")), + LITERAL("{0.f, 0.f}"), LITERAL("{1.f, 1.f}"))), CALL_FACTORY_ENTRY( "__hfma2_sat", CALL(MapNames::getDpctNamespace() + "clamp", @@ -261,7 +241,7 @@ RewriterMap dpct::createHalf2ArithmeticFunctionsRewriterMap() { makeCallArgCreatorWithCall(0), makeCallArgCreatorWithCall(1)), makeCallArgCreatorWithCall(2)), - LITERAL("{0.f, 0.f}"), LITERAL("{1.f, 1.f}")))))) + LITERAL("{0.f, 0.f}"), LITERAL("{1.f, 1.f}"))))) // __hmul2 MATH_API_REWRITER_DEVICE( "__hmul2", @@ -317,17 +297,12 @@ RewriterMap dpct::createHalf2ArithmeticFunctionsRewriterMap() { MapNames::getClNamespace() + "half2>", BO(BinaryOperatorKind::BO_Mul, ARG(0), ARG(1)), LITERAL("{0.f, 0.f}"), LITERAL("{1.f, 1.f}"))), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hmul2_sat", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hmul2_sat")), CALL_FACTORY_ENTRY( "__hmul2_sat", CALL(MapNames::getDpctNamespace() + "clamp", BO(BinaryOperatorKind::BO_Mul, ARG(0), ARG(1)), LITERAL("{0.f, 0.f}"), - LITERAL("{1.f, 1.f}"))))))) + LITERAL("{1.f, 1.f}")))))) // __hneg2 MATH_API_REWRITER_DEVICE( "__hneg2", @@ -399,17 +374,12 @@ RewriterMap dpct::createHalf2ArithmeticFunctionsRewriterMap() { MapNames::getClNamespace() + "half2>", BO(BinaryOperatorKind::BO_Sub, ARG(0), ARG(1)), LITERAL("{0.f, 0.f}"), LITERAL("{1.f, 1.f}"))), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hsub2_sat", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hsub2_sat")), CALL_FACTORY_ENTRY( "__hsub2_sat", CALL(MapNames::getDpctNamespace() + "clamp", BO(BinaryOperatorKind::BO_Sub, ARG(0), ARG(1)), LITERAL("{0.f, 0.f}"), - LITERAL("{1.f, 1.f}"))))))) + LITERAL("{1.f, 1.f}")))))) // h2div MATH_API_REWRITER_DEVICE( "h2div", diff --git a/clang/lib/DPCT/RulesLang/Math/RewriterHalf2ComparisonFunctions.cpp b/clang/lib/DPCT/RulesLang/Math/RewriterHalf2ComparisonFunctions.cpp index 91f40274e45c..7c5d863d5c85 100644 --- a/clang/lib/DPCT/RulesLang/Math/RewriterHalf2ComparisonFunctions.cpp +++ b/clang/lib/DPCT/RulesLang/Math/RewriterHalf2ComparisonFunctions.cpp @@ -286,15 +286,10 @@ RewriterMap dpct::createHalf2ComparisonFunctionsRewriterMap() { EMPTY_FACTORY_ENTRY("__heq2_mask"), EMPTY_FACTORY_ENTRY("__heq2_mask"), EMPTY_FACTORY_ENTRY("__heq2_mask"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__heq2_mask", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__heq2_mask")), CALL_FACTORY_ENTRY( "__heq2_mask", CALL(MapNames::getDpctNamespace() + "compare_mask", - ARG(0), ARG(1), LITERAL("std::equal_to<>()")))))) + ARG(0), ARG(1), LITERAL("std::equal_to<>()"))))) // __hequ2 MATH_API_REWRITER_DEVICE( "__hequ2", @@ -322,16 +317,11 @@ RewriterMap dpct::createHalf2ComparisonFunctionsRewriterMap() { EMPTY_FACTORY_ENTRY("__hequ2_mask"), EMPTY_FACTORY_ENTRY("__hequ2_mask"), EMPTY_FACTORY_ENTRY("__hequ2_mask"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hequ2_mask", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hequ2_mask")), CALL_FACTORY_ENTRY("__hequ2_mask", CALL(MapNames::getDpctNamespace() + "unordered_compare_mask", ARG(0), ARG(1), - LITERAL("std::equal_to<>()")))))) + LITERAL("std::equal_to<>()"))))) // __hge2 MATH_API_REWRITER_DEVICE( "__hge2", @@ -359,16 +349,11 @@ RewriterMap dpct::createHalf2ComparisonFunctionsRewriterMap() { EMPTY_FACTORY_ENTRY("__hge2_mask"), EMPTY_FACTORY_ENTRY("__hge2_mask"), EMPTY_FACTORY_ENTRY("__hge2_mask"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hge2_mask", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hge2_mask")), CALL_FACTORY_ENTRY( "__hge2_mask", CALL(MapNames::getDpctNamespace() + "compare_mask", ARG(0), ARG(1), - LITERAL("std::greater_equal<>()")))))) + LITERAL("std::greater_equal<>()"))))) // __hgeu2 MATH_API_REWRITER_DEVICE( "__hgeu2", @@ -397,16 +382,11 @@ RewriterMap dpct::createHalf2ComparisonFunctionsRewriterMap() { EMPTY_FACTORY_ENTRY("__hgeu2_mask"), EMPTY_FACTORY_ENTRY("__hgeu2_mask"), EMPTY_FACTORY_ENTRY("__hgeu2_mask"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hgeu2_mask", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hgeu2_mask")), CALL_FACTORY_ENTRY("__hgeu2_mask", CALL(MapNames::getDpctNamespace() + "unordered_compare_mask", ARG(0), ARG(1), - LITERAL("std::greater_equal<>()")))))) + LITERAL("std::greater_equal<>()"))))) // __hgt2 MATH_API_REWRITER_DEVICE( "__hgt2", @@ -434,15 +414,10 @@ RewriterMap dpct::createHalf2ComparisonFunctionsRewriterMap() { EMPTY_FACTORY_ENTRY("__hgt2_mask"), EMPTY_FACTORY_ENTRY("__hgt2_mask"), EMPTY_FACTORY_ENTRY("__hgt2_mask"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hgt2_mask", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hgt2_mask")), CALL_FACTORY_ENTRY( "__hgt2_mask", CALL(MapNames::getDpctNamespace() + "compare_mask", - ARG(0), ARG(1), LITERAL("std::greater<>()")))))) + ARG(0), ARG(1), LITERAL("std::greater<>()"))))) // __hgtu2 MATH_API_REWRITER_DEVICE( "__hgtu2", @@ -470,16 +445,11 @@ RewriterMap dpct::createHalf2ComparisonFunctionsRewriterMap() { EMPTY_FACTORY_ENTRY("__hgtu2_mask"), EMPTY_FACTORY_ENTRY("__hgtu2_mask"), EMPTY_FACTORY_ENTRY("__hgtu2_mask"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hgtu2_mask", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hgtu2_mask")), CALL_FACTORY_ENTRY("__hgtu2_mask", CALL(MapNames::getDpctNamespace() + "unordered_compare_mask", ARG(0), ARG(1), - LITERAL("std::greater<>()")))))) + LITERAL("std::greater<>()"))))) // __hisnan2 MATH_API_REWRITER_DEVICE_OVERLOAD( CheckArgType(0, "__half2"), @@ -544,15 +514,10 @@ RewriterMap dpct::createHalf2ComparisonFunctionsRewriterMap() { EMPTY_FACTORY_ENTRY("__hle2_mask"), EMPTY_FACTORY_ENTRY("__hle2_mask"), EMPTY_FACTORY_ENTRY("__hle2_mask"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hle2_mask", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hle2_mask")), CALL_FACTORY_ENTRY( "__hle2_mask", CALL(MapNames::getDpctNamespace() + "compare_mask", - ARG(0), ARG(1), LITERAL("std::less_equal<>()")))))) + ARG(0), ARG(1), LITERAL("std::less_equal<>()"))))) // __hleu2 MATH_API_REWRITER_DEVICE( "__hleu2", @@ -580,16 +545,11 @@ RewriterMap dpct::createHalf2ComparisonFunctionsRewriterMap() { EMPTY_FACTORY_ENTRY("__hleu2_mask"), EMPTY_FACTORY_ENTRY("__hleu2_mask"), EMPTY_FACTORY_ENTRY("__hleu2_mask"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hleu2_mask", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hleu2_mask")), CALL_FACTORY_ENTRY("__hleu2_mask", CALL(MapNames::getDpctNamespace() + "unordered_compare_mask", ARG(0), ARG(1), - LITERAL("std::less_equal<>()")))))) + LITERAL("std::less_equal<>()"))))) // __hlt2 MATH_API_REWRITER_DEVICE( "__hlt2", @@ -617,15 +577,10 @@ RewriterMap dpct::createHalf2ComparisonFunctionsRewriterMap() { EMPTY_FACTORY_ENTRY("__hlt2_mask"), EMPTY_FACTORY_ENTRY("__hlt2_mask"), EMPTY_FACTORY_ENTRY("__hlt2_mask"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hlt2_mask", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hlt2_mask")), CALL_FACTORY_ENTRY( "__hlt2_mask", CALL(MapNames::getDpctNamespace() + "compare_mask", - ARG(0), ARG(1), LITERAL("std::less<>()")))))) + ARG(0), ARG(1), LITERAL("std::less<>()"))))) // __hltu2 MATH_API_REWRITER_DEVICE( "__hltu2", @@ -653,16 +608,11 @@ RewriterMap dpct::createHalf2ComparisonFunctionsRewriterMap() { EMPTY_FACTORY_ENTRY("__hltu2_mask"), EMPTY_FACTORY_ENTRY("__hltu2_mask"), EMPTY_FACTORY_ENTRY("__hltu2_mask"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hltu2_mask", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hltu2_mask")), CALL_FACTORY_ENTRY("__hltu2_mask", CALL(MapNames::getDpctNamespace() + "unordered_compare_mask", ARG(0), ARG(1), - LITERAL("std::less<>()")))))) + LITERAL("std::less<>()"))))) // __hmax2 MATH_API_REWRITER_DEVICE_OVERLOAD( CheckArgType(0, "__half2"), @@ -720,15 +670,10 @@ RewriterMap dpct::createHalf2ComparisonFunctionsRewriterMap() { "ext::intel::math::hmax2_nan", ARG(0), ARG(1))))), EMPTY_FACTORY_ENTRY("__hmax2_nan"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hmax2_nan", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hmax2_nan")), CALL_FACTORY_ENTRY( "__hmax2_nan", CALL(MapNames::getDpctNamespace() + "fmax_nan", ARG(0), - ARG(1)))))) + ARG(1))))) // __hmin2 MATH_API_REWRITER_DEVICE_OVERLOAD( CheckArgType(0, "__half2"), @@ -786,15 +731,10 @@ RewriterMap dpct::createHalf2ComparisonFunctionsRewriterMap() { "ext::intel::math::hmin2_nan", ARG(0), ARG(1))))), EMPTY_FACTORY_ENTRY("__hmin2_nan"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hmin2_nan", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hmin2_nan")), CALL_FACTORY_ENTRY( "__hmin2_nan", CALL(MapNames::getDpctNamespace() + "fmin_nan", ARG(0), - ARG(1)))))) + ARG(1))))) // __hne2 MATH_API_REWRITER_DEVICE( "__hne2", @@ -822,15 +762,10 @@ RewriterMap dpct::createHalf2ComparisonFunctionsRewriterMap() { EMPTY_FACTORY_ENTRY("__hne2_mask"), EMPTY_FACTORY_ENTRY("__hne2_mask"), EMPTY_FACTORY_ENTRY("__hne2_mask"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hne2_mask", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hne2_mask")), CALL_FACTORY_ENTRY( "__hne2_mask", CALL(MapNames::getDpctNamespace() + "compare_mask", - ARG(0), ARG(1), LITERAL("std::not_equal_to<>()")))))) + ARG(0), ARG(1), LITERAL("std::not_equal_to<>()"))))) // __hneu2 MATH_API_REWRITER_DEVICE( "__hneu2", @@ -858,15 +793,10 @@ RewriterMap dpct::createHalf2ComparisonFunctionsRewriterMap() { EMPTY_FACTORY_ENTRY("__hneu2_mask"), EMPTY_FACTORY_ENTRY("__hneu2_mask"), EMPTY_FACTORY_ENTRY("__hneu2_mask"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hneu2_mask", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hneu2_mask")), CALL_FACTORY_ENTRY( "__hneu2_mask", CALL(MapNames::getDpctNamespace() + "unordered_compare_mask", ARG(0), ARG(1), - LITERAL("std::not_equal_to<>()"))))))}; + LITERAL("std::not_equal_to<>()")))))}; } diff --git a/clang/lib/DPCT/RulesLang/Math/RewriterHalfComparisonFunctions.cpp b/clang/lib/DPCT/RulesLang/Math/RewriterHalfComparisonFunctions.cpp index a9831f1b27d6..16386d65246b 100644 --- a/clang/lib/DPCT/RulesLang/Math/RewriterHalfComparisonFunctions.cpp +++ b/clang/lib/DPCT/RulesLang/Math/RewriterHalfComparisonFunctions.cpp @@ -291,15 +291,10 @@ RewriterMap dpct::createHalfComparisonFunctionsRewriterMap() { "ext::intel::math::hmax_nan", ARG(0), ARG(1))))), EMPTY_FACTORY_ENTRY("__hmax_nan"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hmax_nan", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hmax_nan")), CALL_FACTORY_ENTRY( "__hmax_nan", CALL(MapNames::getDpctNamespace() + "fmax_nan", ARG(0), - ARG(1)))))) + ARG(1))))) // __hmin MATH_API_REWRITER_DEVICE_OVERLOAD( CheckArgType(0, "__half"), @@ -341,15 +336,10 @@ RewriterMap dpct::createHalfComparisonFunctionsRewriterMap() { "ext::intel::math::hmin_nan", ARG(0), ARG(1))))), EMPTY_FACTORY_ENTRY("__hmin_nan"), - CONDITIONAL_FACTORY_ENTRY( - UseSYCLCompat, - UNSUPPORT_FACTORY_ENTRY("__hmin_nan", - Diagnostics::UNSUPPORT_SYCLCOMPAT, - LITERAL("__hmin_nan")), CALL_FACTORY_ENTRY( "__hmin_nan", CALL(MapNames::getDpctNamespace() + "fmin_nan", ARG(0), - ARG(1)))))) + ARG(1))))) // __hne MATH_API_REWRITER_DEVICE( "__hne",