Skip to content

[SYCLomatic] Enable APIs that SYCLcompat already supported #2780

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
wants to merge 1 commit into
base: SYCLomatic
Choose a base branch
from
Draft

[SYCLomatic] Enable APIs that SYCLcompat already supported #2780

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
42 changes: 6 additions & 36 deletions clang/lib/DPCT/RulesLang/Math/RewriterHalf2ArithmeticFunctions.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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",
Expand All @@ -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"),
Expand Down Expand Up @@ -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",
Expand All @@ -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",
Expand Down Expand Up @@ -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",
Expand Down Expand Up @@ -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",
Expand Down
98 changes: 14 additions & 84 deletions clang/lib/DPCT/RulesLang/Math/RewriterHalf2ComparisonFunctions.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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",
Expand Down Expand Up @@ -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",
Expand Down Expand Up @@ -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",
Expand Down Expand Up @@ -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",
Expand Down Expand Up @@ -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",
Expand Down Expand Up @@ -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"),
Expand Down Expand Up @@ -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",
Expand Down Expand Up @@ -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",
Expand Down Expand Up @@ -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",
Expand Down Expand Up @@ -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"),
Expand Down Expand Up @@ -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"),
Expand Down Expand Up @@ -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",
Expand Down Expand Up @@ -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",
Expand Down Expand Up @@ -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<>()")))))};
}
Loading
Loading