[SYCLomatic] Add __nv_half and __nv_half2 migration. (#2714)

Signed-off-by: Chen, Sheng S <sheng.s.chen@intel.com>

Author: ShengchenJ

clang/lib/DPCT/RuleInfra/MapNames.cpp

@@ -327,6 +327,8 @@ void MapNames::setExplicitNamespaceMap(
       {"__half2", std::make_shared<TypeNameRule>(getClNamespace() + "half2")},
       {"half", std::make_shared<TypeNameRule>(getClNamespace() + "half")},
       {"half2", std::make_shared<TypeNameRule>(getClNamespace() + "half2")},
+      {"__nv_half2", std::make_shared<TypeNameRule>(getClNamespace() + "half2")},
+      {"__nv_half", std::make_shared<TypeNameRule>(getClNamespace() + "half")},
       {"cudaEvent_t",
        std::make_shared<TypeNameRule>(getDpctNamespace() + "event_ptr",
                                       HelperFeatureEnum::device_ext)},

clang/lib/DPCT/RulesLang/MapNamesLang.cpp

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

clang/lib/DPCT/RulesLang/MapNamesLang.h

@@ -29,7 +29,8 @@ 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", "__half2_raw"
+      "__nv_bfloat162", "nv_bfloat162", "__half_raw", "__half2_raw",           \
+      "__nv_half2", "__nv_half"
 
 /// Record mapping between names
 class MapNamesLang {

clang/lib/DPCT/RulesLang/RulesLang.cpp

@@ -346,7 +346,8 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
               "cublasLtMatmulHeuristicResult_t", "CUjit_target",
               "cublasLtMatrixTransformDesc_t", "cudaGraphicsMapFlags",
               "cudaGraphicsRegisterFlags", "cudaExternalMemoryHandleType",
-              "CUstreamCallback", "cudaHostFn_t"))))))
+              "CUstreamCallback", "cudaHostFn_t", "__nv_half2",
+              "__nv_half"))))))
           .bind("cudaTypeDef"),
       this);
 
@@ -1189,10 +1190,6 @@ void VectorTypeNamespaceRule::registerMatcher(MatchFinder &MF) {
                     .bind("vectorTypeTL"),
                 this);
 
-  MF.addMatcher(
-      cxxRecordDecl(isDirectlyDerivedFrom(hasAnyName(SUPPORTEDVECTORTYPENAMES)))
-          .bind("inheritanceType"),
-      this);
 
   auto Vec3Types = [&]() {
     return hasAnyName("char3", "uchar3", "short3", "ushort3", "int3", "uint3",
@@ -5034,8 +5031,9 @@ void DeviceFunctionDeclRule::registerMatcher(ast_matchers::MatchFinder &MF) {
                 this);
 
   MF.addMatcher(typeLoc(hasAncestor(DeviceFunctionMatcher),
-                        loc(qualType(hasDeclaration(namedDecl(hasAnyName(
-                            "__half", "half", "__half2", "half2"))))))
+                        loc(qualType(hasDeclaration(namedDecl(
+                            hasAnyName("__half", "half", "__half2", "half2",
+                                       "__nv_half2", "__nv_half"))))))
                     .bind("fp16"),
                 this);
 

clang/test/dpct/types008.cu

@@ -18,3 +18,23 @@
 void foo(){
   thrust::reverse_iterator<int *> d_tmp(0);
 }
+
+void foo1() {
+  //CHECK:sycl::half nv_h;
+  //CHECK-NEXT:int a = sizeof(sycl::half);
+  //CHECK-NEXT:a = sizeof(nv_h);
+  //CHECK-NEXT:a = sizeof nv_h;
+  __nv_half nv_h;
+  int a = sizeof(__nv_half);
+  a = sizeof(nv_h);
+  a = sizeof nv_h;
+
+  //CHECK:sycl::half2 nv_h2;
+  //CHECK-NEXT:a = sizeof(sycl::half2);
+  //CHECK-NEXT:a = sizeof(nv_h2);
+  //CHECK-NEXT:a = sizeof nv_h2;
+  __nv_half2 nv_h2;
+  a = sizeof(__nv_half2);
+  a = sizeof(nv_h2);
+  a = sizeof nv_h2;
+}