diff --git a/clang/examples/DPCT/Runtime/cooperative_groups$$experimental$$this_thread_block.cu b/clang/examples/DPCT/Runtime/cooperative_groups$$experimental$$this_thread_block.cu new file mode 100644 index 000000000000..18263f81ab19 --- /dev/null +++ b/clang/examples/DPCT/Runtime/cooperative_groups$$experimental$$this_thread_block.cu @@ -0,0 +1,11 @@ +#define _CG_ABI_EXPERIMENTAL + +#include +__device__ +void _Copy() { + __shared__ cooperative_groups::experimental::block_tile_memory<8> shared; +// Start + cooperative_groups::thread_block thb = cooperative_groups::experimental::this_thread_block(shared/*cooperative_groups::experimental::block_tile_memory*/); +// End +} + diff --git a/clang/examples/DPCT/Runtime/cooperative_groups$$experimental$$tiled_partition.cu b/clang/examples/DPCT/Runtime/cooperative_groups$$experimental$$tiled_partition.cu new file mode 100644 index 000000000000..8c34894333ad --- /dev/null +++ b/clang/examples/DPCT/Runtime/cooperative_groups$$experimental$$tiled_partition.cu @@ -0,0 +1,11 @@ + +#define _CG_ABI_EXPERIMENTAL +#include +#include "cooperative_groups.h" + +__device__ void test(cooperative_groups::thread_block tb) { + cooperative_groups::experimental::block_tile_memory<1, 1> mem; + // Start + cooperative_groups::thread_block_tile<32> tbt32 = cooperative_groups::experimental::tiled_partition<32>(tb/*cooperative_groups::thread_block*/); + // End +} diff --git a/clang/lib/DPCT/DPCT.cpp b/clang/lib/DPCT/DPCT.cpp index 2f06c278c1ea..71c7e236c073 100644 --- a/clang/lib/DPCT/DPCT.cpp +++ b/clang/lib/DPCT/DPCT.cpp @@ -1447,8 +1447,10 @@ int runDPCT(int argc, const char **argv) { StringRef ErrStr = Err; // Avoid the "Visual Studio version" error on windows platform. if (ErrStr.find("error:") == ErrStr.rfind("error:") && - ErrStr.contains( - "error -- unsupported Microsoft Visual Studio version")) { + (ErrStr.contains("no function template matches function " + "template specialization 'this_multi_grid'") || + ErrStr.contains( + "error -- unsupported Microsoft Visual Studio version"))) { break; } if (ErrStr.contains("use of undeclared identifier")) { diff --git a/clang/test/dpct/query_api_mapping/Runtime/test_cooperative_groups_experimental.cu b/clang/test/dpct/query_api_mapping/Runtime/test_cooperative_groups_experimental.cu new file mode 100644 index 000000000000..1d205aa4f8c4 --- /dev/null +++ b/clang/test/dpct/query_api_mapping/Runtime/test_cooperative_groups_experimental.cu @@ -0,0 +1,27 @@ +// UNSUPPORTED: system-windows +// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2, v11.0, v12.1, v12.2, v12.3, v12.4, v12.5, v12.6, v12.7, v12.8, v12.9 +// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2, cuda-11.0, cuda-12.1, cuda-12.2, cuda-12.3, cuda-12.4, cuda-12.5, cuda-12.6, cuda-12.7, cuda-12.8, cuda-12.9 + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cooperative_groups::exclusive_scan | FileCheck %s -check-prefix=CG_EXCLUSIVE_SCAN +// CG_EXCLUSIVE_SCAN: CUDA API: +// CG_EXCLUSIVE_SCAN-NEXT: cooperative_groups::exclusive_scan( +// CG_EXCLUSIVE_SCAN-NEXT: tile32 /* type group */, sdata[tid] /* type value */, +// CG_EXCLUSIVE_SCAN-NEXT: cooperative_groups::plus() /* type operator */); +// CG_EXCLUSIVE_SCAN-NEXT: cooperative_groups::exclusive_scan(tile32 /* type group */, +// CG_EXCLUSIVE_SCAN-NEXT sdata[tid] /* type value */); +// CG_EXCLUSIVE_SCAN: Is migrated to: +// CG_EXCLUSIVE_SCAN-NEXT: sycl::exclusive_scan_over_group(sycl::ext::oneapi::this_work_item::get_sub_group(), sdata[tid], sycl::plus()); +// CG_EXCLUSIVE_SCAN-NEXT: sycl::exclusive_scan_over_group(sycl::ext::oneapi::this_work_item::get_sub_group(), sdata[tid], sycl::plus<>()); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cooperative_groups::experimental::this_thread_block +// CG_EXP_THIS_THREAD_BLOCK: CUDA API: +// CG_EXP_THIS_THREAD_BLOCK-NEXT: cooperative_groups::thread_block thb = cooperative_groups::experimental::this_thread_block(shared/*cooperative_groups::experimental::block_tile_memory*/); +// CG_EXP_THIS_THREAD_BLOCK-NEXT: Is migrated to: +// CG_EXP_THIS_THREAD_BLOCK-NEXT: sycl::group<3> thb = sycl::ext::oneapi::this_work_item::get_work_group<3>(); + + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cooperative_groups::experimental::tiled_partition +// CG_EXP_THIS_TILED_PARTITION: CUDA API: +// CG_EXP_THIS_TILED_PARTITION-NEXT: cooperative_groups::thread_block_tile<32> tbt32 = cooperative_groups::experimental::tiled_partition<32>(tb/*cooperative_groups::thread_block*/); +// CG_EXP_THIS_TILED_PARTITION-NEXT: Is migrated to: +// CG_EXP_THIS_TILED_PARTITION-NEXT: sycl::sub_group tbt32 = sycl::ext::oneapi::this_work_item::get_sub_group(); diff --git a/clang/test/dpct/query_api_mapping/test_all.cu b/clang/test/dpct/query_api_mapping/test_all.cu index afa9c8e49a9d..4550891e2a32 100644 --- a/clang/test/dpct/query_api_mapping/test_all.cu +++ b/clang/test/dpct/query_api_mapping/test_all.cu @@ -629,6 +629,8 @@ // CHECK-NEXT: cooperative_groups::coalesced_group::thread_rank // CHECK-NEXT: cooperative_groups::coalesced_threads // CHECK-NEXT: cooperative_groups::exclusive_scan +// CHECK-NEXT: cooperative_groups::experimental::this_thread_block +// CHECK-NEXT: cooperative_groups::experimental::tiled_partition // CHECK-NEXT: cooperative_groups::greater // CHECK-NEXT: cooperative_groups::grid_group::block_rank // CHECK-NEXT: cooperative_groups::grid_group::num_blocks