[SYCLomatic] Support migration for type __half2_raw (#2713)

Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>

Author: zhiweij1

clang/lib/DPCT/RuleInfra/MapNames.cpp

@@ -813,6 +813,8 @@ void MapNames::setExplicitNamespaceMap(
       {"cusparseSpGEMMAlg_t", std::make_shared<TypeNameRule>("int")},
       {"cusparseSpSVAlg_t", std::make_shared<TypeNameRule>("int")},
       {"__half_raw", std::make_shared<TypeNameRule>("uint16_t")},
+      {"__half2_raw",
+       std::make_shared<TypeNameRule>(MapNames::getClNamespace() + "ushort2")},
       {"cudaFuncAttributes",
        std::make_shared<TypeNameRule>(MapNames::getDpctNamespace() +
                                       "kernel_function_info")},

clang/lib/DPCT/RulesLang/MapNamesLang.cpp

@@ -91,7 +91,8 @@ const std::map<std::string, int> MapNamesLang::VectorTypeMigratedTypeSizeMap{
     {"ulonglong4", 32}, {"float1", 4},      {"float2", 8},
     {"float3", 16},     {"float4", 16},     {"double1", 8},
     {"double2", 16},    {"double3", 32},    {"double4", 32},
-    {"__half", 2},      {"__half2", 4},     {"__half_raw", 2}};
+    {"__half", 2},      {"__half2", 4},     {"__half_raw", 2},
+    {"__half2_raw", 4}};
 
 const std::map<clang::dpct::KernelArgType, int>
     MapNamesLang::KernelArgTypeSizeMap{

clang/lib/DPCT/RulesLang/MapNamesLang.h

@@ -29,7 +29,7 @@ const std::string StringLiteralUnsupported{"UNSUPPORTED"};
       "ulonglong1", "longlong2", "ulonglong2", "longlong3", "ulonglong3",      \
       "longlong4", "ulonglong4", "double1", "double2", "double3", "double4",   \
       "__half", "__half2", "half", "half2", "__nv_bfloat16", "nv_bfloat16",    \
-      "__nv_bfloat162", "nv_bfloat162", "__half_raw"
+      "__nv_bfloat162", "nv_bfloat162", "__half_raw", "__half2_raw"
 #define VECTORTYPE2MARRAYNAMES "__nv_bfloat162", "nv_bfloat162"
 
 /// Record mapping between names

clang/lib/DPCT/RulesLang/RulesLang.cpp

@@ -1218,12 +1218,12 @@ void VectorTypeNamespaceRule::registerMatcher(MatchFinder &MF) {
                                   "longlong1", "ulonglong1", "double1", "__half_raw")))
                     .bind("inherit"),
                 this);
-  // Matcher for __half_raw implicitly convert to half.
+  // Matcher for __half_raw/__half2_raw implicitly convert to half/half2.
   MF.addMatcher(
       declRefExpr(allOf(unless(hasParent(memberExpr())),
                         unless(hasParent(unaryOperator(hasOperatorName("&")))),
-                        to(varDecl(hasType(qualType(hasDeclaration(
-                                       namedDecl(hasAnyName("__half_raw"))))))),
+                        to(varDecl(hasType(qualType(hasDeclaration(namedDecl(
+                            hasAnyName("__half_raw", "__half2_raw"))))))),
                         hasParent(implicitCastExpr())))
           .bind("halfRawExpr"),
       this);
@@ -1405,7 +1405,7 @@ void VectorTypeNamespaceRule::runRule(const MatchFinder::MatchResult &Result) {
         UETT, Diagnostics::SIZEOF_WARNING, true, argTypeName,
         "Check that the allocated memory size in the migrated code is correct");
   }
-  // Runrule for __half_raw implicitly convert to half.
+  // Run rule for __half_raw/__half2_raw implicitly convert to half/half2.
   if (auto DRE = getNodeAsType<DeclRefExpr>(Result, "halfRawExpr")) {
     if (const auto *RT =
             DRE->getType().getCanonicalType()->getAs<RecordType>()) {
@@ -1414,13 +1414,17 @@ void VectorTypeNamespaceRule::runRule(const MatchFinder::MatchResult &Result) {
     }
     ExprAnalysis EA;
     std::string Replacement;
-    llvm::raw_string_ostream OS(Replacement);
-    OS << MapNames::getClNamespace() + "bit_cast<" +
-              MapNames::getClNamespace() + "half>(";
     EA.analyze(DRE);
-    OS << EA.getReplacedString();
-    OS << ")";
-    OS.flush();
+    if (DRE->getType().getCanonicalType().getAsString() == "__half2_raw") {
+      llvm::raw_string_ostream OS(Replacement);
+      OS << EA.getReplacedString() << ".as<" << MapNames::getClNamespace()
+         << "half2>()";
+    } else {
+      llvm::raw_string_ostream OS(Replacement);
+      OS << MapNames::getClNamespace() << "bit_cast<"
+         << MapNames::getClNamespace() << "half>(" << EA.getReplacedString()
+         << ")";
+    }
     emplaceTransformation(new ReplaceStmt(DRE, Replacement));
     return;
   }

clang/test/dpct/half_raw.cu

@@ -6,7 +6,8 @@
 
 #include <cuda_fp16.h>
 #include <cuda_runtime.h>
-int main() {
+
+void foo1() {
   // CHECK: uint16_t one_h{0x3C00};
   __half_raw one_h{0x3C00};
   // CHECK: uint16_t zero_h{0};
@@ -40,3 +41,54 @@ int main() {
   // CHECK: uint16_t *ptr1 = &one_h;
   __half_raw *ptr1 = &one_h;
 }
+
+void foo2() {
+  // CHECK: sycl::ushort2 one_h{0x3C00};
+  __half2_raw one_h{0x3C00};
+  // CHECK: sycl::ushort2 zero_h{0};
+  __half2_raw zero_h{0};
+  // CHECK: sycl::ushort2 *ptr = new sycl::ushort2{0};
+  __half2_raw *ptr = new __half2_raw{0};
+  // clang-format off
+  // CHECK: ptr->x() = 0x3C00;
+  ptr->x = 0x3C00;
+  // CHECK: ptr ->x() = 0x3C00;
+  ptr ->x = 0x3C00;
+  // CHECK: ptr-> x() = 0x3C00;
+  ptr-> x = 0x3C00;
+  // CHECK: ptr -> x() = 0x3C00;
+  ptr -> x = 0x3C00;
+  // CHECK: zero_h.x() = 0x3C00;
+  zero_h.x = 0x3C00;
+  // CHECK: zero_h .x() = 0x3C00;
+  zero_h .x = 0x3C00;
+  // CHECK: zero_h. x() = 0x3C00;
+  zero_h. x = 0x3C00;
+  // CHECK: zero_h . x() = 0x3C00;
+  zero_h . x = 0x3C00;
+  // CHECK: ptr->y() = 0x3C00;
+  ptr->y = 0x3C00;
+  // CHECK: ptr ->y() = 0x3C00;
+  ptr ->y = 0x3C00;
+  // CHECK: ptr-> y() = 0x3C00;
+  ptr-> y = 0x3C00;
+  // CHECK: ptr -> y() = 0x3C00;
+  ptr -> y = 0x3C00;
+  // CHECK: zero_h.y() = 0x3C00;
+  zero_h.y = 0x3C00;
+  // CHECK: zero_h .y() = 0x3C00;
+  zero_h .y = 0x3C00;
+  // CHECK: zero_h. y() = 0x3C00;
+  zero_h. y = 0x3C00;
+  // CHECK: zero_h . y() = 0x3C00;
+  zero_h . y = 0x3C00;
+  // clang-format on
+  // CHECK: sycl::half2 alpha = one_h.as<sycl::half2>();
+  half2 alpha = one_h;
+  // CHECK: alpha = one_h.as<sycl::half2>();
+  alpha = one_h;
+  // CHECK: uint16_t as = zero_h.x();
+  uint16_t as = zero_h.x;
+  // CHECK: sycl::ushort2 *ptr1 = &one_h;
+  __half2_raw *ptr1 = &one_h;
+}