From a39e929a44f99de5cdd59c316ee9774950f1e7be Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Fri, 27 Jun 2025 10:11:58 +0800 Subject: [PATCH 1/5] [SYCLomatic] Remove 2 experimental test cooperative_groups APIs in INC file. Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/SrcAPI/APINames.inc | 2 -- 1 file changed, 2 deletions(-) diff --git a/clang/lib/DPCT/SrcAPI/APINames.inc b/clang/lib/DPCT/SrcAPI/APINames.inc index 8f17aaa8d9a2..cd0b7cdb2cf9 100644 --- a/clang/lib/DPCT/SrcAPI/APINames.inc +++ b/clang/lib/DPCT/SrcAPI/APINames.inc @@ -2140,9 +2140,7 @@ ENTRY(cooperative_groups::this_thread, cooperative_groups::__v1::this_thread, tr ENTRY(cooperative_groups::this_grid, cooperative_groups::__v1::this_grid, true, NO_FLAG, P4, "Successful") ENTRY(cooperative_groups::this_multi_grid, cooperative_groups::__v1::this_multi_grid, false, NO_FLAG, P4, "comment") ENTRY(cooperative_groups::this_thread_block, cooperative_groups::__v1::this_thread_block, true, NO_FLAG, P4, "Successful") -ENTRY(cooperative_groups::experimental::this_thread_block, cooperative_groups::__v1::experimental::this_thread_block, true, NO_FLAG, P4, "Successful") ENTRY(cooperative_groups::tiled_partition, cooperative_groups::__v1::tiled_partition, true, NO_FLAG, P4, "If size equals to 32, tool will migrate it, else warning is emitted.") -ENTRY(cooperative_groups::experimental::tiled_partition, cooperative_groups::__v1::experimental::tiled_partition, true, NO_FLAG, P4, "If size equals to 32, tool will migrate it, else warning is emitted.") ENTRY(cooperative_groups::labeled_partition, cooperative_groups::__v1::labeled_partition, false, NO_FLAG, P4, "comment") ENTRY(cooperative_groups::binary_partition, cooperative_groups::__v1::binary_partition, false, NO_FLAG, P4, "comment") ENTRY(cooperative_groups::sync, cooperative_groups::__v1::sync, true, NO_FLAG, P4, "comment") From a3f381ac6ea94306e525d9fb6536ef96903ec970 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Mon, 7 Jul 2025 09:53:47 +0800 Subject: [PATCH 2/5] up Signed-off-by: Chen, Sheng S --- ...ve_groups$$experimental$$this_thread_block.cu | 16 ++++++++++++++++ ...tive_groups$$experimental$$tiled_partition.cu | 12 ++++++++++++ 2 files changed, 28 insertions(+) create mode 100644 clang/examples/DPCT/Runtime/cooperative_groups$$experimental$$this_thread_block.cu create mode 100644 clang/examples/DPCT/Runtime/cooperative_groups$$experimental$$tiled_partition.cu 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..bfe97a44c895 --- /dev/null +++ b/clang/examples/DPCT/Runtime/cooperative_groups$$experimental$$this_thread_block.cu @@ -0,0 +1,16 @@ + +#define _CG_ABI_EXPERIMENTAL + +#include +#include "cooperative_groups.h" + +void test() { + cooperative_groups::experimental::block_tile_memory<1, 1> mem; + + // Start + cooperative_groups::thread_block tb = + cooperative_groups::experimental::this_thread_block(); + // 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..0564fbb73c7c --- /dev/null +++ b/clang/examples/DPCT/Runtime/cooperative_groups$$experimental$$tiled_partition.cu @@ -0,0 +1,12 @@ + +#define _CG_ABI_EXPERIMENTAL +#include +#include "cooperative_groups.h" + +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 +} From e6e67c25a43357aebf7a3dc771c92e245ee3f345 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Mon, 7 Jul 2025 14:07:45 +0800 Subject: [PATCH 3/5] [SYCLomatic] Add 2 cooperative_groups experimental query API mappings. Signed-off-by: Chen, Sheng S --- ...groups$$experimental$$this_thread_block.cu | 21 +++++++------------ ...e_groups$$experimental$$tiled_partition.cu | 3 +-- clang/lib/DPCT/DPCT.cpp | 4 ++-- clang/lib/DPCT/SrcAPI/APINames.inc | 2 ++ 4 files changed, 13 insertions(+), 17 deletions(-) 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 index bfe97a44c895..18263f81ab19 100644 --- a/clang/examples/DPCT/Runtime/cooperative_groups$$experimental$$this_thread_block.cu +++ b/clang/examples/DPCT/Runtime/cooperative_groups$$experimental$$this_thread_block.cu @@ -1,16 +1,11 @@ - #define _CG_ABI_EXPERIMENTAL -#include -#include "cooperative_groups.h" - -void test() { - cooperative_groups::experimental::block_tile_memory<1, 1> mem; - - // Start - cooperative_groups::thread_block tb = - cooperative_groups::experimental::this_thread_block(); - // End - - +#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 index 0564fbb73c7c..8c34894333ad 100644 --- a/clang/examples/DPCT/Runtime/cooperative_groups$$experimental$$tiled_partition.cu +++ b/clang/examples/DPCT/Runtime/cooperative_groups$$experimental$$tiled_partition.cu @@ -3,9 +3,8 @@ #include #include "cooperative_groups.h" -void test(cooperative_groups::thread_block tb) { +__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..7671503820c5 100644 --- a/clang/lib/DPCT/DPCT.cpp +++ b/clang/lib/DPCT/DPCT.cpp @@ -1446,9 +1446,9 @@ int runDPCT(int argc, const char **argv) { std::string Err = getDpctTermStr(); StringRef ErrStr = Err; // Avoid the "Visual Studio version" error on windows platform. - if (ErrStr.find("error:") == ErrStr.rfind("error:") && + if (ErrStr.find("error:") == ErrStr.rfind("error:") && (ErrStr.contains("no function template matches function template specialization 'this_multi_grid'") || ErrStr.contains( - "error -- unsupported Microsoft Visual Studio version")) { + "error -- unsupported Microsoft Visual Studio version"))) { break; } if (ErrStr.contains("use of undeclared identifier")) { diff --git a/clang/lib/DPCT/SrcAPI/APINames.inc b/clang/lib/DPCT/SrcAPI/APINames.inc index cd0b7cdb2cf9..8f17aaa8d9a2 100644 --- a/clang/lib/DPCT/SrcAPI/APINames.inc +++ b/clang/lib/DPCT/SrcAPI/APINames.inc @@ -2140,7 +2140,9 @@ ENTRY(cooperative_groups::this_thread, cooperative_groups::__v1::this_thread, tr ENTRY(cooperative_groups::this_grid, cooperative_groups::__v1::this_grid, true, NO_FLAG, P4, "Successful") ENTRY(cooperative_groups::this_multi_grid, cooperative_groups::__v1::this_multi_grid, false, NO_FLAG, P4, "comment") ENTRY(cooperative_groups::this_thread_block, cooperative_groups::__v1::this_thread_block, true, NO_FLAG, P4, "Successful") +ENTRY(cooperative_groups::experimental::this_thread_block, cooperative_groups::__v1::experimental::this_thread_block, true, NO_FLAG, P4, "Successful") ENTRY(cooperative_groups::tiled_partition, cooperative_groups::__v1::tiled_partition, true, NO_FLAG, P4, "If size equals to 32, tool will migrate it, else warning is emitted.") +ENTRY(cooperative_groups::experimental::tiled_partition, cooperative_groups::__v1::experimental::tiled_partition, true, NO_FLAG, P4, "If size equals to 32, tool will migrate it, else warning is emitted.") ENTRY(cooperative_groups::labeled_partition, cooperative_groups::__v1::labeled_partition, false, NO_FLAG, P4, "comment") ENTRY(cooperative_groups::binary_partition, cooperative_groups::__v1::binary_partition, false, NO_FLAG, P4, "comment") ENTRY(cooperative_groups::sync, cooperative_groups::__v1::sync, true, NO_FLAG, P4, "comment") From 425fc61dabec9e8211113a48196538ce26a04e3c Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Mon, 7 Jul 2025 14:10:24 +0800 Subject: [PATCH 4/5] up Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/DPCT.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/clang/lib/DPCT/DPCT.cpp b/clang/lib/DPCT/DPCT.cpp index 7671503820c5..71c7e236c073 100644 --- a/clang/lib/DPCT/DPCT.cpp +++ b/clang/lib/DPCT/DPCT.cpp @@ -1446,9 +1446,11 @@ int runDPCT(int argc, const char **argv) { std::string Err = getDpctTermStr(); StringRef ErrStr = Err; // Avoid the "Visual Studio version" error on windows platform. - if (ErrStr.find("error:") == ErrStr.rfind("error:") && (ErrStr.contains("no function template matches function template specialization 'this_multi_grid'") || - ErrStr.contains( - "error -- unsupported Microsoft Visual Studio version"))) { + if (ErrStr.find("error:") == ErrStr.rfind("error:") && + (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")) { From 705a3c0894598d3690675068846da5004c693c85 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Mon, 7 Jul 2025 15:36:42 +0800 Subject: [PATCH 5/5] up Signed-off-by: Chen, Sheng S --- .../test_cooperative_groups_experimental.cu | 27 +++++++++++++++++++ clang/test/dpct/query_api_mapping/test_all.cu | 2 ++ 2 files changed, 29 insertions(+) create mode 100644 clang/test/dpct/query_api_mapping/Runtime/test_cooperative_groups_experimental.cu 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