diff --git a/clang/examples/DPCT/Math/mul24.cu b/clang/examples/DPCT/Math/mul24.cu new file mode 100644 index 000000000000..41a06cd9a88e --- /dev/null +++ b/clang/examples/DPCT/Math/mul24.cu @@ -0,0 +1,5 @@ +__global__ void test(int a, int b) { + // Start + mul24(a /*int*/, b /*int*/); + // End +} diff --git a/clang/examples/DPCT/Runtime/__barrier_sync.cu b/clang/examples/DPCT/Runtime/__barrier_sync.cu new file mode 100644 index 000000000000..1a57ea590bd4 --- /dev/null +++ b/clang/examples/DPCT/Runtime/__barrier_sync.cu @@ -0,0 +1,8 @@ +__global__ void test(unsigned id) { + goto label; + // Start + __barrier_sync(id /*unsigned*/); + // End +label: + int a; +} diff --git a/clang/test/dpct/query_api_mapping/Math/test-before12.cu b/clang/test/dpct/query_api_mapping/Math/test-before12.cu index ad9cf9c3989f..73421715a06c 100644 --- a/clang/test/dpct/query_api_mapping/Math/test-before12.cu +++ b/clang/test/dpct/query_api_mapping/Math/test-before12.cu @@ -6,3 +6,9 @@ // saturate-NEXT: saturate(f /*float*/); // saturate-NEXT: Is migrated to: // saturate-NEXT: dpct::clamp(f, 0.0f, 1.0f); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=mul24 | FileCheck %s -check-prefix=mul24 +// mul24: CUDA API: +// mul24-NEXT: mul24(a /*int*/, b /*int*/); +// mul24-NEXT: Is migrated to: +// mul24-NEXT: sycl::mul24(a, b); diff --git a/clang/test/dpct/query_api_mapping/Runtime/test-after9.cu b/clang/test/dpct/query_api_mapping/Runtime/test-after9.cu index 263ee5dc3c06..271f3d9b044c 100644 --- a/clang/test/dpct/query_api_mapping/Runtime/test-after9.cu +++ b/clang/test/dpct/query_api_mapping/Runtime/test-after9.cu @@ -17,3 +17,9 @@ // CUDALAUNCHCOOPERATIVEKERNEL-NEXT: sharedMem /*size_t*/, s /*cudaStream_t*/); // CUDALAUNCHCOOPERATIVEKERNEL-NEXT: Is migrated to: // CUDALAUNCHCOOPERATIVEKERNEL-NEXT: dpct::kernel_launcher::launch(f, gridDim, blockDim, args, sharedMem, s); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__barrier_sync | FileCheck %s -check-prefix=__barrier_sync +// __barrier_sync: CUDA API: +// __barrier_sync-NEXT: __barrier_sync(id /*unsigned*/); +// __barrier_sync-NEXT: Is migrated to: +// __barrier_sync-NEXT: sycl::ext::oneapi::this_work_item::get_nd_item<3>().barrier(); diff --git a/clang/test/dpct/query_api_mapping/test_all.cu b/clang/test/dpct/query_api_mapping/test_all.cu index afa9c8e49a9d..b34daf1b6edc 100644 --- a/clang/test/dpct/query_api_mapping/test_all.cu +++ b/clang/test/dpct/query_api_mapping/test_all.cu @@ -11,6 +11,7 @@ // CHECK-NEXT: __assertfail // CHECK-NEXT: __ballot // CHECK-NEXT: __ballot_sync +// CHECK-NEXT: __barrier_sync // CHECK-NEXT: __bfloat1622float2 // CHECK-NEXT: __bfloat162bfloat162 // CHECK-NEXT: __bfloat162float @@ -2293,6 +2294,7 @@ // CHECK-NEXT: min // CHECK-NEXT: modf // CHECK-NEXT: modff +// CHECK-NEXT: mul24 // CHECK-NEXT: nan // CHECK-NEXT: nanf // CHECK-NEXT: ncclAllReduce