[SYCLomatic] Add 2 cooperative_groups experimental query API mappings. (#2917)

Signed-off-by: Chen, Sheng S <sheng.s.chen@intel.com>

Author: zhimingwang36

clang/examples/DPCT/Runtime/cooperative_groups$$experimental$$this_thread_block.cu

@@ -0,0 +1,11 @@
+#define _CG_ABI_EXPERIMENTAL
+
+#include <cooperative_groups.h>
+__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
+}
+

clang/examples/DPCT/Runtime/cooperative_groups$$experimental$$tiled_partition.cu

@@ -0,0 +1,11 @@
+
+#define _CG_ABI_EXPERIMENTAL
+#include <cuda.h>
+#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
+}

clang/lib/DPCT/DPCT.cpp

@@ -1433,8 +1433,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")) {

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<double>() /* 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<double>());
+// 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();

clang/test/dpct/query_api_mapping/test_all.cu

@@ -630,6 +630,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