[SYCLomatic] Fix migration bug of dim3 with --enable-ctad. (oneapi-sr…

…c#2149) Signed-off-by: Tang, Jiajun [email protected]
OuadiElfarouki · Jul 16, 2024 · b7f0338 · b7f0338
1 parent 6019a73
commit b7f0338
Show file tree

Hide file tree

Showing 2 changed files with 15 additions and 5 deletions.
diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp
@@ -5411,9 +5411,7 @@ void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) {
     Printer.line("[=](" + MapNames::getClNamespace() + "nd_item<1> ",
                  getItemName(), ")", ExecutionConfig.SubGroupSize, " {");
   } else {
-    DpctGlobalInfo::printCtadClass(Printer.indent(),
-                                   MapNames::getClNamespace() + "nd_range", 3)
-        << "(";
+    Printer << MapNames::getClNamespace() + "nd_range<3>(";
     if (ExecutionConfig.GroupSize == CanIgnoreRangeStr3D) {
       Printer << ExecutionConfig.LocalSize;
     } else if (ExecutionConfig.LocalSize == CanIgnoreRangeStr3D) {

diff --git a/clang/test/dpct/ctad.cu b/clang/test/dpct/ctad.cu
@@ -92,7 +92,7 @@ int main() {
   // CHECK-NEXT:       sycl::local_accessor<int, 1> k_acc_ct1(sycl::range(32), cgh);
   // CHECK-EMPTY:
   // CHECK-NEXT:       cgh.parallel_for<dpct_kernel_name<class kernel_{{[a-f0-9]+}}>>(
-  // CHECK-NEXT:         sycl::nd_range(sycl::range(1, 1, 1), sycl::range(1, 1, 1)),
+  // CHECK-NEXT:         sycl::nd_range<3>(sycl::range(1, 1, 1), sycl::range(1, 1, 1)),
   // CHECK-NEXT:         [=](sycl::nd_item<3> item_ct1) {
   // CHECK-NEXT:           kernel(1, k_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get());
   // CHECK-NEXT:         });
@@ -103,12 +103,24 @@ int main() {
   // CHECK-NEXT:       sycl::local_accessor<int, 1> k_acc_ct1(sycl::range(32), cgh);
   // CHECK-EMPTY:
   // CHECK-NEXT:       cgh.parallel_for<dpct_kernel_name<class kernel_{{[a-f0-9]+}}>>(
-  // CHECK-NEXT:         sycl::nd_range(sycl::range(1, 1, NUM) * sycl::range(1, 1, NUM), sycl::range(1, 1, NUM)),
+  // CHECK-NEXT:         sycl::nd_range<3>(sycl::range(1, 1, NUM) * sycl::range(1, 1, NUM), sycl::range(1, 1, NUM)),
   // CHECK-NEXT:         [=](sycl::nd_item<3> item_ct1) {
   // CHECK-NEXT:           kernel(1, k_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get());
   // CHECK-NEXT:         });
   // CHECK-NEXT:     });
   kernel<<<NUM, NUM>>>(1);
+
+  // CHECK:   q_ct1.submit(
+  // CHECK-NEXT:     [&](sycl::handler &cgh) {
+  // CHECK-NEXT:       sycl::local_accessor<int, 1> k_acc_ct1(sycl::range(32), cgh);
+  // CHECK-EMPTY:
+  // CHECK-NEXT:       cgh.parallel_for<dpct_kernel_name<class kernel_{{[a-f0-9]+}}>>(
+  // CHECK-NEXT:         sycl::nd_range<3>(deflt * deflt, deflt),
+  // CHECK-NEXT:         [=](sycl::nd_item<3> item_ct1) {
+  // CHECK-NEXT:           kernel(1, k_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get());
+  // CHECK-NEXT:         });
+  // CHECK-NEXT:     });
+  kernel<<<deflt, deflt>>>(1);
 }