[SYCLomatic] Refine the "has_capability_or_fail" usage in the migrate…

…d code (#2515) Signed-off-by: Jiang, Zhiwei <[email protected]>
oneapi-src · Nov 28, 2024 · 2a2b1d9 · 2a2b1d9
1 parent 2a562a7
commit 2a2b1d9
Show file tree

Hide file tree

Showing 7 changed files with 16 additions and 23 deletions.
diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp
@@ -6315,26 +6315,20 @@ void KernelCallExpr::addDevCapCheckStmt() {
   if (!AspectList.empty()) {
     std::string Str;
     llvm::raw_string_ostream OS(Str);
+    OS << MapNames::getDpctNamespace() << "has_capability_or_fail(";
     if (auto Iter = MapNames::CustomHelperFunctionMap.find(getQueueKind());
         Iter != MapNames::CustomHelperFunctionMap.end()) {
-      OS << MapNames::getDpctNamespace() << "has_capability_or_fail(";
-      OS << Iter->second << ".get_device(), ";
-      OS << "{" << AspectList.front();
-      for (size_t i = 1; i < AspectList.size(); ++i) {
-        OS << ", " << AspectList[i];
-      }
-      OS << "});";
+      OS << Iter->second << ".";
     } else {
       requestFeature(HelperFeatureEnum::device_ext);
-      OS << MapNames::getDpctNamespace() << "get_device(";
-      OS << MapNames::getDpctNamespace() << "get_device_id(";
       printStreamBase(OS);
-      OS << "get_device())).has_capability_or_fail({" << AspectList.front();
-      for (size_t i = 1; i < AspectList.size(); ++i) {
-        OS << ", " << AspectList[i];
-      }
-      OS << "});";
     }
+    OS << "get_device(), ";
+    OS << "{" << AspectList.front();
+    for (size_t i = 1; i < AspectList.size(); ++i) {
+      OS << ", " << AspectList[i];
+    }
+    OS << "});";
     OuterStmts.OthersList.emplace_back(OS.str());
   }
 }

diff --git a/clang/test/dpct/kernel_without_name-usm.cu b/clang/test/dpct/kernel_without_name-usm.cu
@@ -31,7 +31,7 @@ struct Mat {
   }
 
 void run_foo1(Mat mat) {
-  // CHECK: DISPATCH(mat.getType(), ([&] {   dpct::get_device(dpct::get_device_id(dpct::get_in_order_queue().get_device())).has_capability_or_fail({sycl::aspect::fp64});
+  // CHECK: DISPATCH(mat.getType(), ([&] {   dpct::has_capability_or_fail(dpct::get_in_order_queue().get_device(), {sycl::aspect::fp64});
   // CHECK-EMPTY:
   // CHECK-NEXT: dpct::get_in_order_queue().submit(
   // CHECK-NEXT:   [&](sycl::handler &cgh) {

diff --git a/clang/test/dpct/macro_test.cu b/clang/test/dpct/macro_test.cu
@@ -1102,8 +1102,7 @@ template<class T1, class T2, int N> __global__ void foo31();
 #define FOO31(DIMS) foo31<unsigned int, float, DIMS><<<1,1>>>();
 
 //CHECK: {
-//CHECK-NEXT:   dpct::get_device(dpct::get_device_id(q_ct1.get_device()))
-//CHECK-NEXT:       .has_capability_or_fail({sycl::aspect::fp64});
+//CHECK-NEXT:   dpct::has_capability_or_fail(q_ct1.get_device(), {sycl::aspect::fp64});
 //CHECK-EMPTY:
 //CHECK-NEXT:   q_ct1.submit([&](sycl::handler &cgh) {
 //CHECK-NEXT:     /*

diff --git a/clang/test/dpct/mf-test.cu b/clang/test/dpct/mf-test.cu
@@ -68,7 +68,7 @@ void test() {
   kernel_extern<<<1,1>>>();
 
   //      CHECK:   {
-  // CHECK-NEXT:     dpct::get_device(dpct::get_device_id(q_ct1.get_device())).has_capability_or_fail({sycl::aspect::fp64, sycl::aspect::fp16});
+  // CHECK-NEXT:     dpct::has_capability_or_fail(q_ct1.get_device(), {sycl::aspect::fp64, sycl::aspect::fp16});
   // CHECK-EMPTY:
   // CHECK-NEXT:     q_ct1.parallel_for<dpct_kernel_name<class test_fp_{{[a-f0-9]+}}>>(
   // CHECK-NEXT:       sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), 

diff --git a/clang/test/dpct/sync_api.cu b/clang/test/dpct/sync_api.cu
@@ -198,7 +198,7 @@ __global__ void foo_tile32() {
 int foo3() {
   // CHECK: {
   // CHECK-NEXT:  auto exp_props = sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::use_root_sync};
-  // CHECK-NEXT:  dpct::get_device(dpct::get_device_id(dpct::get_in_order_queue().get_device())).has_capability_or_fail({sycl::aspect::fp64});
+  // CHECK-NEXT:  dpct::has_capability_or_fail(dpct::get_in_order_queue().get_device(), {sycl::aspect::fp64});
   // CHECK-EMPTY:
   // CHECK-NEXT:  dpct::get_in_order_queue().parallel_for(
   // CHECK-NEXT:  sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),

diff --git a/clang/test/dpct/template-kernel-call.cu b/clang/test/dpct/template-kernel-call.cu
@@ -256,7 +256,7 @@ __global__ void convert_kernel(T b){
 // CHECK-NEXT:void convert(){
 // CHECK-NEXT:  T b;
 // CHECK-NEXT:  {
-// CHECK-NEXT:  dpct::get_device(dpct::get_device_id(dpct::get_out_of_order_queue().get_device())).has_capability_or_fail({sycl::aspect::fp64});
+// CHECK-NEXT:  dpct::has_capability_or_fail(dpct::get_out_of_order_queue().get_device(), {sycl::aspect::fp64});
 // CHECK-EMPTY:
 // CHECK-NEXT:  dpct::get_out_of_order_queue().submit(
 // CHECK-NEXT:    [&](sycl::handler &cgh) {

diff --git a/clang/test/dpct/tm-complex-profiling.cu b/clang/test/dpct/tm-complex-profiling.cu
@@ -356,7 +356,7 @@ void foo_test_4() {
   // CHECK-NEXT:  DPCT1049:{{[0-9]+}}: The work-group size passed to the SYCL kernel may exceed the limit. To get the device limit, query info::device::max_work_group_size. Adjust the work-group size if needed.
   // CHECK-NEXT:  */
   // CHECK-NEXT:  {
-  // CHECK-NEXT:    dpct::get_device(dpct::get_device_id(dpct::get_in_order_queue().get_device())).has_capability_or_fail({sycl::aspect::fp64});
+  // CHECK-NEXT:    dpct::has_capability_or_fail(dpct::get_in_order_queue().get_device(), {sycl::aspect::fp64});
   // CHECK-EMPTY:
   // CHECK-NEXT:    dpct::get_in_order_queue().parallel_for<dpct_kernel_name<class set_array_{{[a-z0-9]+}}>>(
   // CHECK-NEXT:                sycl::nd_range<3>(dimGrid * dimBlock, dimBlock),
@@ -377,7 +377,7 @@ void foo_test_4() {
     // CHECK-NEXT:    DPCT1049:{{[0-9]+}}: The work-group size passed to the SYCL kernel may exceed the limit. To get the device limit, query info::device::max_work_group_size. Adjust the work-group size if needed.
     // CHECK-NEXT:    */
     // CHECK-NEXT:    {
-    // CHECK-NEXT:        dpct::get_device(dpct::get_device_id(dpct::get_in_order_queue().get_device())).has_capability_or_fail({sycl::aspect::fp64});
+    // CHECK-NEXT:        dpct::has_capability_or_fail(dpct::get_in_order_queue().get_device(), {sycl::aspect::fp64});
     // CHECK-EMPTY:
     // CHECK-NEXT:        dpct::get_in_order_queue().parallel_for<dpct_kernel_name<class STREAM_Copy_{{[a-z0-9]+}}>>(
     // CHECK-NEXT:                    sycl::nd_range<3>(dimGrid * dimBlock, dimBlock),
@@ -399,7 +399,7 @@ void foo_test_4() {
     // CHECK-NEXT:    DPCT1049:{{[0-9]+}}: The work-group size passed to the SYCL kernel may exceed the limit. To get the device limit, query info::device::max_work_group_size. Adjust the work-group size if needed.
     // CHECK-NEXT:    */
     // CHECK-NEXT:    {
-    // CHECK-NEXT:        dpct::get_device(dpct::get_device_id(dpct::get_in_order_queue().get_device())).has_capability_or_fail({sycl::aspect::fp64});
+    // CHECK-NEXT:        dpct::has_capability_or_fail(dpct::get_in_order_queue().get_device(), {sycl::aspect::fp64});
     // CHECK-EMPTY:
     // CHECK-NEXT:        dpct::get_in_order_queue().parallel_for<dpct_kernel_name<class STREAM_Copy_Optimized_{{[a-z0-9]+}}>>(
     // CHECK-NEXT:                    sycl::nd_range<3>(dimGrid * dimBlock, dimBlock),