[SYCLomatic] Fix the issue that no viable conversion from 'sycl::ext::oneapi::experimental::device_global' to pointer (#2432)

Signed-off-by: intwanghao <[email protected]>

Author: intwanghao

clang/lib/DPCT/ASTTraversal.cpp

@@ -8651,7 +8651,17 @@ void MemVarRefMigrationRule::runRule(const MatchFinder::MatchResult &Result) {
           Info->setInitForDeviceGlobal(InitStr);
         }
       }
-      if (!Info->getType()->isArray()) {
+      auto VarType = Info->getType();
+      if (VarType->isArray()) {
+        if (const auto *const ICE =
+                dyn_cast_or_null<ImplicitCastExpr>(Parent)) {
+          if (ICE->getCastKind() == CK_ArrayToPointerDecay) {
+            if (!dyn_cast_or_null<ArraySubscriptExpr>(getParentStmt(ICE))) {
+              emplaceTransformation(new InsertAfterStmt(MemVarRef, ".get()"));
+            }
+          }
+        }
+      } else {
         emplaceTransformation(new InsertAfterStmt(MemVarRef, ".get()"));
       }
       return;

clang/lib/DPCT/AnalysisInfo.cpp

@@ -3282,15 +3282,10 @@ std::string MemVarInfo::getMemoryType(const std::string &MemoryType,
                        VarType->getDimension(), ">");
   } else if (isUseDeviceGlobal()) {
     std::string Dims;
-    std::string Specifier;
     for (auto &D : VarType->getRange()) {
       Dims = Dims + "[" + D.getSize() + "]";
     }
-    if (isConstant()) {
-      Specifier = "const ";
-    }
-    return buildString(MemoryType, "<", Specifier, VarType->getBaseName(), Dims,
-                       ">");
+    return buildString(MemoryType, "<", VarType->getBaseName(), Dims, ">");
   } else {
     return buildString(MemoryType, VarType->getBaseNameWithoutQualifiers());
   }

clang/test/dpct/device_global/device_global.cu

@@ -23,13 +23,13 @@ struct B{
 // CHECK: /*
 // CHECK: DPCT1127:{{[0-9]+}}: The constant compile-time initialization for device_global is supported when compiling with C++20. You may need to adjust the compile commands.
 // CHECK: */
-// CHECK: static sycl::ext::oneapi::experimental::device_global<const float> var_c{2.f};
+// CHECK: static sycl::ext::oneapi::experimental::device_global<float> var_c{2.f};
 // CHECK: /*
 // CHECK: DPCT1127:{{[0-9]+}}: The constant compile-time initialization for device_global is supported when compiling with C++20. You may need to adjust the compile commands.
 // CHECK: */
-// CHECK: static sycl::ext::oneapi::experimental::device_global<const float> var_d{1.f};
+// CHECK: static sycl::ext::oneapi::experimental::device_global<float> var_d{1.f};
 // CHECK: static sycl::ext::oneapi::experimental::device_global<A> var_e;
-// CHECK: static sycl::ext::oneapi::experimental::device_global<const B> var_f;
+// CHECK: static sycl::ext::oneapi::experimental::device_global<B> var_f;
 // CHECK: static dpct::global_memory<int, 0> var_g;
 __device__ int var_a;
 __device__ int var_b = 0;
@@ -46,12 +46,12 @@ __device__ int var_g;
 // CHECK: /*
 // CHECK: DPCT1127:{{[0-9]+}}: The constant compile-time initialization for device_global is supported when compiling with C++20. You may need to adjust the compile commands.
 // CHECK: */
-// CHECK: static sycl::ext::oneapi::experimental::device_global<const int[10]> arr_c{3, 2, 1};
+// CHECK: static sycl::ext::oneapi::experimental::device_global<int[10]> arr_c{3, 2, 1};
 // CHECK: /*
 // CHECK: DPCT1127:{{[0-9]+}}: The constant compile-time initialization for device_global is supported when compiling with C++20. You may need to adjust the compile commands.
 // CHECK: */
-// CHECK: static sycl::ext::oneapi::experimental::device_global<const int[10]> arr_d{2};
-// CHECK: static sycl::ext::oneapi::experimental::device_global<const A[10]> arr_e;
+// CHECK: static sycl::ext::oneapi::experimental::device_global<int[10]> arr_d{2};
+// CHECK: static sycl::ext::oneapi::experimental::device_global<A[10]> arr_e;
 // CHECK: static sycl::ext::oneapi::experimental::device_global<B[10]> arr_f;
 __device__ float arr_a[10];
 __device__ float arr_b[10] = {1, 2, 3};
@@ -62,6 +62,7 @@ __device__ B arr_f[10];
 
 
 // CHECK: int device_func() {
+// CHECK:   float *p = arra_a.get();
 // CHECK:   arr_a[0] = 1;
 // CHECK:   return arr_a[0] + arr_b[0] + arr_c[0] + arr_d[0] + arr_e[0].data + arr_f[0].data;
 // CHECK: }
@@ -72,6 +73,7 @@ __device__ B arr_f[10];
 // CHECK:   *ptr = var_a.get() + var_b.get() + var_c.get() + var_d.get() + var_e.get().data + var_f.get().data + device_func();
 // CHECK: }
 __device__ int device_func() {
+  float *p = arra_a;
   arr_a[0] = 1;
   return arr_a[0] + arr_b[0] + arr_c[0] + arr_d[0] + arr_e[0].data + arr_f[0].data;
 }