From 2a2b1d9adf5fe2ffd57f9a1666a2e871ac013e0e Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Thu, 28 Nov 2024 15:14:19 +0800 Subject: [PATCH] [SYCLomatic] Refine the "has_capability_or_fail" usage in the migrated code (#2515) Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/AnalysisInfo.cpp | 22 ++++++++-------------- clang/test/dpct/kernel_without_name-usm.cu | 2 +- clang/test/dpct/macro_test.cu | 3 +-- clang/test/dpct/mf-test.cu | 2 +- clang/test/dpct/sync_api.cu | 2 +- clang/test/dpct/template-kernel-call.cu | 2 +- clang/test/dpct/tm-complex-profiling.cu | 6 +++--- 7 files changed, 16 insertions(+), 23 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 717d345494c5..a34c12cd4b2c 100644 --- 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 index 26623cccb05c..eb20b7eedefa 100644 --- 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 index d03472313911..cbedd59037f3 100644 --- a/clang/test/dpct/macro_test.cu +++ b/clang/test/dpct/macro_test.cu @@ -1102,8 +1102,7 @@ template __global__ void foo31(); #define FOO31(DIMS) foo31<<<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 index b1da1260dbe6..68581ed28e42 100644 --- 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>( // 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 index 00807ea54c5d..1edafef39c59 100644 --- 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 index 326a046413aa..c16355d4e56c 100644 --- 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 index 43dd6f663168..37d4667b16c0 100644 --- 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>( // 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>( // 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>( // CHECK-NEXT: sycl::nd_range<3>(dimGrid * dimBlock, dimBlock),