Skip to content

Commit 66c3d8c

Browse files
zhiweij1zhiweij1
authored
[SYCLomatic] Use .x()/.y()/... instead of [...] to migrate bfloat16 member access (#2712)
Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>
1 parent 4ea8029 commit 66c3d8c

24 files changed

Lines changed: 260 additions & 309 deletions

‎clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp‎

Lines changed: 1 addition & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -792,20 +792,7 @@ void ExprAnalysis::analyzeExpr(const MemberExpr *ME) {
792792
addReplacement(ME->getOperatorLoc(), ME->getEndLoc(), "");
793793
} else {
794794
std::string MemberName = ME->getMemberNameInfo().getAsString();
795-
const auto &MArrayIdx =
796-
MapNamesLang::MArrayMemberNamesMap.find(MemberName);
797-
if (MapNamesLang::VectorTypes2MArray.count(BaseType) &&
798-
MArrayIdx != MapNamesLang::MArrayMemberNamesMap.end()) {
799-
std::string RepStr = "";
800-
if (isImplicit) {
801-
RepStr = "(*this)";
802-
} else if (isPtr) {
803-
addReplacement(ME->getBeginLoc(), 0, "(*");
804-
RepStr = ")";
805-
}
806-
addReplacement(Begin, ME->getEndLoc(), RepStr + MArrayIdx->second);
807-
} else if (MapNames::replaceName(MapNamesLang::MemberNamesMap,
808-
MemberName)) {
795+
if (MapNames::replaceName(MapNamesLang::MemberNamesMap, MemberName)) {
809796
std::string RepStr = "";
810797
const auto *MD = DpctGlobalInfo::findAncestor<CXXMethodDecl>(ME);
811798
if (MD && MD->isVolatile()) {

‎clang/lib/DPCT/RuleInfra/MapNames.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -741,12 +741,12 @@ void MapNames::setExplicitNamespaceMap(
741741
{"__nv_bfloat16", std::make_shared<TypeNameRule>(
742742
getClNamespace() + "ext::oneapi::bfloat16")},
743743
{"__nv_bfloat162", std::make_shared<TypeNameRule>(
744-
getClNamespace() + "marray<" + getClNamespace() +
744+
getClNamespace() + "vec<" + getClNamespace() +
745745
"ext::oneapi::bfloat16, 2>")},
746746
{"nv_bfloat16", std::make_shared<TypeNameRule>(getClNamespace() +
747747
"ext::oneapi::bfloat16")},
748748
{"nv_bfloat162", std::make_shared<TypeNameRule>(
749-
getClNamespace() + "marray<" + getClNamespace() +
749+
getClNamespace() + "vec<" + getClNamespace() +
750750
"ext::oneapi::bfloat16, 2>")},
751751
{"libraryPropertyType_t",
752752
std::make_shared<TypeNameRule>(getLibraryHelperNamespace() +

‎clang/lib/DPCT/RulesLang/MapNamesLang.cpp‎

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -71,8 +71,6 @@ void MapNamesLang::setExplicitNamespaceMap(
7171
// Supported vector types
7272
const MapNamesLang::SetTy MapNamesLang::SupportedVectorTypes{
7373
SUPPORTEDVECTORTYPENAMES};
74-
const MapNamesLang::SetTy MapNamesLang::VectorTypes2MArray{
75-
VECTORTYPE2MARRAYNAMES};
7674

7775
const std::map<std::string, int> MapNamesLang::VectorTypeMigratedTypeSizeMap{
7876
{"char1", 1}, {"char2", 2}, {"char3", 4},

‎clang/lib/DPCT/RulesLang/MapNamesLang.h‎

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,6 @@ const std::string StringLiteralUnsupported{"UNSUPPORTED"};
3030
"longlong4", "ulonglong4", "double1", "double2", "double3", "double4", \
3131
"__half", "__half2", "half", "half2", "__nv_bfloat16", "nv_bfloat16", \
3232
"__nv_bfloat162", "nv_bfloat162", "__half_raw", "__half2_raw"
33-
#define VECTORTYPE2MARRAYNAMES "__nv_bfloat162", "nv_bfloat162"
3433

3534
/// Record mapping between names
3635
class MapNamesLang {
@@ -43,7 +42,6 @@ class MapNamesLang {
4342
using SetTy = std::set<std::string>;
4443

4544
static const SetTy SupportedVectorTypes;
46-
static const SetTy VectorTypes2MArray;
4745
static const std::map<std::string, int> VectorTypeMigratedTypeSizeMap;
4846
static const std::map<clang::dpct::KernelArgType, int> KernelArgTypeSizeMap;
4947
static int getArrayTypeSize(const int Dim);

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

Lines changed: 26 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -23,19 +23,19 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
2323
CALL(MapNames::getClNamespace() + "float2",
2424
CALL(MapNames::getClNamespace() +
2525
"ext::intel::math::bfloat162float",
26-
ARRAY_SUBSCRIPT(ARG(0), LITERAL("0"))),
26+
MEMBER_CALL(ARG(0), false, "x")),
2727
CALL(MapNames::getClNamespace() +
2828
"ext::intel::math::bfloat162float",
29-
ARRAY_SUBSCRIPT(ARG(0), LITERAL("1")))))),
29+
MEMBER_CALL(ARG(0), false, "y"))))),
3030
CALL_FACTORY_ENTRY("__bfloat1622float2",
3131
CALL(MapNames::getClNamespace() + "float2",
32-
ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")),
33-
ARRAY_SUBSCRIPT(ARG(0), LITERAL("1")))))
32+
MEMBER_CALL(ARG(0), false, "x"),
33+
MEMBER_CALL(ARG(0), false, "y"))))
3434
// __bfloat162bfloat162
3535
CONDITIONAL_FACTORY_ENTRY(
3636
math::UseBFloat16,
3737
CALL_FACTORY_ENTRY("__bfloat162bfloat162",
38-
CALL(MapNames::getClNamespace() + "marray<" +
38+
CALL(MapNames::getClNamespace() + "vec<" +
3939
MapNames::getClNamespace() +
4040
"ext::oneapi::bfloat16, 2>",
4141
ARG(0), ARG(0))),
@@ -512,11 +512,11 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
512512
CONDITIONAL_FACTORY_ENTRY(
513513
math::UseBFloat16,
514514
CALL_FACTORY_ENTRY("__float22bfloat162_rn",
515-
CALL(MapNames::getClNamespace() + "marray<" +
515+
CALL(MapNames::getClNamespace() + "vec<" +
516516
MapNames::getClNamespace() +
517517
"ext::oneapi::bfloat16, 2>",
518-
ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")),
519-
ARRAY_SUBSCRIPT(ARG(0), LITERAL("1")))),
518+
MEMBER_CALL(ARG(0), false, "x"),
519+
MEMBER_CALL(ARG(0), false, "y"))),
520520
UNSUPPORT_FACTORY_ENTRY("__float22bfloat162_rn",
521521
Diagnostics::API_NOT_MIGRATED,
522522
ARG("__float22bfloat162_rn")))
@@ -542,7 +542,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
542542
CONDITIONAL_FACTORY_ENTRY(
543543
math::UseBFloat16,
544544
CALL_FACTORY_ENTRY("__float2bfloat162_rn",
545-
CALL(MapNames::getClNamespace() + "marray<" +
545+
CALL(MapNames::getClNamespace() + "vec<" +
546546
MapNames::getClNamespace() +
547547
"ext::oneapi::bfloat16, 2>",
548548
ARG(0), ARG(0))),
@@ -601,7 +601,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
601601
CONDITIONAL_FACTORY_ENTRY(
602602
math::UseBFloat16,
603603
CALL_FACTORY_ENTRY("__floats2bfloat162_rn",
604-
CALL(MapNames::getClNamespace() + "marray<" +
604+
CALL(MapNames::getClNamespace() + "vec<" +
605605
MapNames::getClNamespace() +
606606
"ext::oneapi::bfloat16, 2>",
607607
ARG(0), ARG(1))),
@@ -612,7 +612,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
612612
CONDITIONAL_FACTORY_ENTRY(
613613
math::UseBFloat16,
614614
CALL_FACTORY_ENTRY("__halves2bfloat162",
615-
CALL(MapNames::getClNamespace() + "marray<" +
615+
CALL(MapNames::getClNamespace() + "vec<" +
616616
MapNames::getClNamespace() +
617617
"ext::oneapi::bfloat16, 2>",
618618
ARG(0), ARG(1))),
@@ -625,31 +625,31 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
625625
CALL_FACTORY_ENTRY(
626626
"__high2bfloat16",
627627
CALL(MapNames::getClNamespace() + "ext::oneapi::bfloat16",
628-
ARRAY_SUBSCRIPT(ARG(0), LITERAL("1")))),
628+
MEMBER_CALL(ARG(0), false, "y"))),
629629
UNSUPPORT_FACTORY_ENTRY("__high2bfloat16",
630630
Diagnostics::API_NOT_MIGRATED,
631631
ARG("__high2bfloat16")))
632632
// __high2bfloat162
633633
CONDITIONAL_FACTORY_ENTRY(
634634
math::UseBFloat16,
635635
CALL_FACTORY_ENTRY("__high2bfloat162",
636-
CALL(MapNames::getClNamespace() + "marray<" +
636+
CALL(MapNames::getClNamespace() + "vec<" +
637637
MapNames::getClNamespace() +
638638
"ext::oneapi::bfloat16, 2>",
639-
ARRAY_SUBSCRIPT(ARG(0), LITERAL("1")),
640-
ARRAY_SUBSCRIPT(ARG(0), LITERAL("1")))),
639+
MEMBER_CALL(ARG(0), false, "y"),
640+
MEMBER_CALL(ARG(0), false, "y"))),
641641
UNSUPPORT_FACTORY_ENTRY("__high2bfloat162",
642642
Diagnostics::API_NOT_MIGRATED,
643643
ARG("__high2bfloat162")))
644644
// __highs2bfloat162
645645
CONDITIONAL_FACTORY_ENTRY(
646646
math::UseBFloat16,
647647
CALL_FACTORY_ENTRY("__highs2bfloat162",
648-
CALL(MapNames::getClNamespace() + "marray<" +
648+
CALL(MapNames::getClNamespace() + "vec<" +
649649
MapNames::getClNamespace() +
650650
"ext::oneapi::bfloat16, 2>",
651-
ARRAY_SUBSCRIPT(ARG(0), LITERAL("1")),
652-
ARRAY_SUBSCRIPT(ARG(1), LITERAL("1")))),
651+
MEMBER_CALL(ARG(0), false, "y"),
652+
MEMBER_CALL(ARG(1), false, "y"))),
653653
UNSUPPORT_FACTORY_ENTRY("__highs2bfloat162",
654654
Diagnostics::API_NOT_MIGRATED,
655655
ARG("__highs2bfloat162")))
@@ -755,31 +755,31 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
755755
CALL_FACTORY_ENTRY(
756756
"__low2bfloat16",
757757
CALL(MapNames::getClNamespace() + "ext::oneapi::bfloat16",
758-
ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
758+
MEMBER_CALL(ARG(0), false, "x"))),
759759
UNSUPPORT_FACTORY_ENTRY("__low2bfloat16",
760760
Diagnostics::API_NOT_MIGRATED,
761761
ARG("__low2bfloat16")))
762762
// __low2bfloat162
763763
CONDITIONAL_FACTORY_ENTRY(
764764
math::UseBFloat16,
765765
CALL_FACTORY_ENTRY("__low2bfloat162",
766-
CALL(MapNames::getClNamespace() + "marray<" +
766+
CALL(MapNames::getClNamespace() + "vec<" +
767767
MapNames::getClNamespace() +
768768
"ext::oneapi::bfloat16, 2>",
769-
ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")),
770-
ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
769+
MEMBER_CALL(ARG(0), false, "x"),
770+
MEMBER_CALL(ARG(0), false, "x"))),
771771
UNSUPPORT_FACTORY_ENTRY("__low2bfloat162",
772772
Diagnostics::API_NOT_MIGRATED,
773773
ARG("__low2bfloat162")))
774774
// __lows2bfloat162
775775
CONDITIONAL_FACTORY_ENTRY(
776776
math::UseBFloat16,
777777
CALL_FACTORY_ENTRY("__lows2bfloat162",
778-
CALL(MapNames::getClNamespace() + "marray<" +
778+
CALL(MapNames::getClNamespace() + "vec<" +
779779
MapNames::getClNamespace() +
780780
"ext::oneapi::bfloat16, 2>",
781-
ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")),
782-
ARRAY_SUBSCRIPT(ARG(1), LITERAL("0")))),
781+
MEMBER_CALL(ARG(0), false, "x"),
782+
MEMBER_CALL(ARG(1), false, "x"))),
783783
UNSUPPORT_FACTORY_ENTRY("__lows2bfloat162",
784784
Diagnostics::API_NOT_MIGRATED,
785785
ARG("__lows2bfloat162")))
@@ -1005,7 +1005,7 @@ RewriterMap dpct::createBfloat16PrecisionConversionAndDataMovementRewriterMap()
10051005
Diagnostics::API_NOT_MIGRATED,
10061006
ARG("__ushort_as_bfloat16")))
10071007
// make_bfloat162
1008-
ENTRY_RENAMED("make_bfloat162", MapNames::getClNamespace() + "marray<" +
1008+
ENTRY_RENAMED("make_bfloat162", MapNames::getClNamespace() + "vec<" +
10091009
MapNames::getClNamespace() +
10101010
"ext::oneapi::bfloat16, 2>")};
10111011
}

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

Lines changed: 3 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -55,15 +55,13 @@ RewriterMap dpct::createHalf2ArithmeticFunctionsRewriterMap() {
5555
ARG(0))),
5656
CALL_FACTORY_ENTRY(
5757
"__habs2",
58-
CALL(MapNames::getClNamespace() + "marray<" +
58+
CALL(MapNames::getClNamespace() + "vec<" +
5959
MapNames::getClNamespace() +
6060
"ext::oneapi::bfloat16, 2>",
6161
CALL(MapNames::getClNamespace(false, true) + "fabs",
62-
CALL("float",
63-
ARRAY_SUBSCRIPT(ARG(0), LITERAL("0")))),
62+
CALL("float", MEMBER_CALL(ARG(0), false, "x"))),
6463
CALL(MapNames::getClNamespace(false, true) + "fabs",
65-
CALL("float",
66-
ARRAY_SUBSCRIPT(ARG(0), LITERAL("1"))))))))
64+
CALL("float", MEMBER_CALL(ARG(0), false, "y")))))))
6765
// __hadd2
6866
MATH_API_REWRITER_DEVICE(
6967
"__hadd2",

0 commit comments

Comments
 (0)