[SYCLomatic] Add query-api-mapping for 43 cooperative_groups APIs

clang/examples/DPCT/Runtime/cooperative_groups$$bit_and.cu

@@ -0,0 +1,14 @@
+#include "cooperative_groups.h"
+#include <cooperative_groups/reduce.h>
+
+__global__ void test() {
+  int *sdata;
+  cooperative_groups::thread_block cta =
+      cooperative_groups::this_thread_block();
+  const unsigned int tid = cta.thread_rank();
+  cooperative_groups::thread_block_tile<32> tile32 = cooperative_groups::tiled_partition<32>(cta);
+  int *idata;
+  // Start
+  cooperative_groups::reduce(tile32 /*thread_block_tile<32>*/, sdata[tid]/*data*/, cooperative_groups::bit_and<int>()/*cg::bit_and<T>*/);
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$bit_or.cu

@@ -0,0 +1,14 @@
+#include "cooperative_groups.h"
+#include <cooperative_groups/reduce.h>
+
+__global__ void test() {
+  int *sdata;
+  cooperative_groups::thread_block cta =
+      cooperative_groups::this_thread_block();
+  const unsigned int tid = cta.thread_rank();
+  cooperative_groups::thread_block_tile<32> tile32 = cooperative_groups::tiled_partition<32>(cta);
+  int *idata;
+  // Start
+  cooperative_groups::reduce(tile32 /*thread_block_tile<32>*/, sdata[tid]/*data*/, cooperative_groups::bit_or<int>()/*cg::bit_or<T>*/);
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$bit_xor.cu

@@ -0,0 +1,14 @@
+#include "cooperative_groups.h"
+#include <cooperative_groups/reduce.h>
+
+__global__ void test() {
+  int *sdata;
+  cooperative_groups::thread_block cta =
+      cooperative_groups::this_thread_block();
+  const unsigned int tid = cta.thread_rank();
+  cooperative_groups::thread_block_tile<32> tile32 = cooperative_groups::tiled_partition<32>(cta);
+  int *idata;
+  // Start
+  cooperative_groups::reduce(tile32 /*thread_block_tile<32>*/, sdata[tid]/*data*/, cooperative_groups::bit_xor<int>()/*cg::bit_xor<T>*/);
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$coalesced_group$$shfl.cu

@@ -0,0 +1,10 @@
+// Option: --use-experimental-features=non-uniform-groups
+#include <cooperative_groups.h>
+
+__global__ void test() {
+
+  // Start
+  cooperative_groups::coalesced_group active = cooperative_groups::coalesced_threads();
+  active.shfl(0, 0);
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$coalesced_group$$size.cu

@@ -0,0 +1,10 @@
+// Option: --use-experimental-features=non-uniform-groups
+#include <cooperative_groups.h>
+
+__global__ void test() {
+
+  // Start
+  cooperative_groups::coalesced_group active = cooperative_groups::coalesced_threads();
+  active.size();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$coalesced_group$$sync.cu

@@ -0,0 +1,9 @@
+// Option: --use-experimental-features=non-uniform-groups
+#include <cooperative_groups.h>
+
+__global__ void test() {
+  // Start
+  cooperative_groups::coalesced_group active = cooperative_groups::coalesced_threads();
+  active.sync();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$coalesced_group$$thread_rank.cu

@@ -0,0 +1,11 @@
+// Option: --use-experimental-features=non-uniform-groups
+#include <cooperative_groups.h>
+
+
+__global__ void test() {
+  // Start
+  cooperative_groups::coalesced_group active =
+      cooperative_groups::coalesced_threads();
+  active.thread_rank();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$coalesced_threads.cu

@@ -0,0 +1,9 @@
+// Option: --use-experimental-features=non-uniform-groups
+#include <cooperative_groups.h>
+#include <cooperative_groups/reduce.h>
+
+__global__ void test() {
+  // Start
+   cooperative_groups::coalesced_group active = cooperative_groups::coalesced_threads();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$exclusive_scan.cu

@@ -0,0 +1,18 @@
+#include <cooperative_groups.h>
+#include <cooperative_groups/scan.h>
+
+__global__ void test() {
+  double *sdata;
+  cooperative_groups::thread_block cta =
+      cooperative_groups::this_thread_block();
+  const unsigned int tid = cta.thread_rank();
+  cooperative_groups::thread_block_tile<32> tile32 =
+      cooperative_groups::tiled_partition<32>(cta);
+  // Start
+  cooperative_groups::exclusive_scan(
+      tile32 /* type group */, sdata[tid] /* type value */,
+      cooperative_groups::plus<double>() /* type operator */);
+  cooperative_groups::exclusive_scan(tile32 /* type group */,
+                                     sdata[tid] /* type value */);
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$greater.cu

@@ -0,0 +1,14 @@
+#include "cooperative_groups.h"
+#include <cooperative_groups/reduce.h>
+
+__global__ void test() {
+  double *sdata;
+  cooperative_groups::thread_block cta =
+      cooperative_groups::this_thread_block();
+  const unsigned int tid = cta.thread_rank();
+  cooperative_groups::thread_block_tile<32> tile32 = cooperative_groups::tiled_partition<32>(cta);
+  int *idata;
+  // Start
+  cooperative_groups::reduce(tile32 /*thread_block_tile<32>*/, sdata[tid]/*data*/, cooperative_groups::greater<double>()/*cg::greater<T>*/);
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$grid_group$$block_rank.cu

@@ -0,0 +1,11 @@
+// Option: --use-experimental-features=nd_range_barrier,root-group
+#include <cooperative_groups.h>
+#include <cooperative_groups/scan.h>
+
+__global__ void test() {
+
+  // Start
+  cooperative_groups::grid_group grid = cooperative_groups::this_grid();
+  grid.block_rank();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$grid_group$$num_blocks.cu

@@ -0,0 +1,11 @@
+// Option: --use-experimental-features=nd_range_barrier,root-group
+#include <cooperative_groups.h>
+#include <cooperative_groups/scan.h>
+
+__global__ void test() {
+
+  // Start
+  cooperative_groups::grid_group grid = cooperative_groups::this_grid();
+  grid.num_blocks();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$grid_group$$num_threads.cu

@@ -0,0 +1,11 @@
+// Option: --use-experimental-features=nd_range_barrier,root-group
+#include <cooperative_groups.h>
+#include <cooperative_groups/scan.h>
+
+__global__ void test() {
+
+  // Start
+  cooperative_groups::grid_group grid = cooperative_groups::this_grid();
+  grid.num_threads();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$grid_group$$size.cu

@@ -0,0 +1,11 @@
+// Option: --use-experimental-features=nd_range_barrier,root-group
+#include <cooperative_groups.h>
+#include <cooperative_groups/scan.h>
+
+__global__ void test() {
+
+  // Start
+  cooperative_groups::grid_group grid = cooperative_groups::this_grid();
+  grid.size();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$grid_group$$sync.cu

@@ -0,0 +1,11 @@
+// Option: --use-experimental-features=nd_range_barrier,root-group
+#include <cooperative_groups.h>
+#include <cooperative_groups/scan.h>
+
+__global__ void test() {
+
+  // Start
+  cooperative_groups::grid_group grid = cooperative_groups::this_grid();
+  grid.sync();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$grid_group$$thread_rank.cu

@@ -0,0 +1,9 @@
+#include <cooperative_groups.h>
+
+__global__ void test() {
+
+  // Start
+  cooperative_groups::grid_group grid = cooperative_groups::this_grid();
+  grid.thread_rank() /* grid_group::thread_rank */;
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$inclusive_scan.cu

@@ -0,0 +1,21 @@
+#include <cooperative_groups.h>
+#include <cooperative_groups/scan.h>
+
+__global__ void test() {
+  double *sdata;
+  cooperative_groups::thread_block cta =
+      cooperative_groups::this_thread_block();
+  const unsigned int tid = cta.thread_rank();
+  cooperative_groups::thread_block_tile<32> tile32 =
+      cooperative_groups::tiled_partition<32>(cta);
+
+  // Start
+  cooperative_groups::inclusive_scan(
+      tile32 /* type group */, sdata[tid] /* type value */,
+      cooperative_groups::plus<double>() /* type operator */);
+
+  cooperative_groups::inclusive_scan(tile32 /* type group */,
+                                     sdata[tid] /* type value */);
+
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$less.cu

@@ -0,0 +1,14 @@
+#include "cooperative_groups.h"
+#include <cooperative_groups/reduce.h>
+
+__global__ void test() {
+  double *sdata;
+  cooperative_groups::thread_block cta =
+      cooperative_groups::this_thread_block();
+  const unsigned int tid = cta.thread_rank();
+  cooperative_groups::thread_block_tile<32> tile32 = cooperative_groups::tiled_partition<32>(cta);
+  int *idata;
+  // Start
+  cooperative_groups::reduce(tile32 /*thread_block_tile<32>*/, sdata[tid]/*data*/, cooperative_groups::less<double>()/*cg::less<T>*/);
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$plus.cu

@@ -0,0 +1,14 @@
+#include "cooperative_groups.h"
+#include <cooperative_groups/reduce.h>
+
+__global__ void test() {
+  double *sdata;
+  cooperative_groups::thread_block cta =
+      cooperative_groups::this_thread_block();
+  const unsigned int tid = cta.thread_rank();
+  cooperative_groups::thread_block_tile<32> tile32 = cooperative_groups::tiled_partition<32>(cta);
+  int *idata;
+  // Start
+  cooperative_groups::reduce(tile32 /*thread_block_tile<32>*/, sdata[tid]/*data*/, cooperative_groups::plus<double>()/*cg::plus<T>*/);
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$sync.cu

@@ -0,0 +1,16 @@
+// Option: --use-experimental-features=logical-group 
+#include "cooperative_groups.h"
+#include <cooperative_groups/reduce.h>
+
+__global__ void test() {
+
+  // Start
+  cooperative_groups::thread_block tb = cooperative_groups::this_thread_block();
+  cooperative_groups::thread_block_tile<32> tbt32 =
+      cooperative_groups::tiled_partition<32>(tb);
+  tb.sync();
+  tbt32.sync();
+  cooperative_groups::sync(tb);
+  cooperative_groups::sync(tbt32);
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$this_grid.cu

@@ -0,0 +1,9 @@
+// Option: --use-experimental-features=root-group
+#include "cooperative_groups.h"
+#include <cooperative_groups/reduce.h>
+
+__global__ void test() {
+  // Start
+cooperative_groups::grid_group grid = cooperative_groups::this_grid();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$this_thread.cu

@@ -0,0 +1,9 @@
+// Option: --use-experimental-features=logical-group
+#include "cooperative_groups.h"
+#include <cooperative_groups/reduce.h>
+
+__global__ void test() {
+  // Start
+  auto thread = cooperative_groups::this_thread();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$thread_block$$group_index.cu

@@ -0,0 +1,10 @@
+#include <cooperative_groups.h>
+#include <cooperative_groups/scan.h>
+
+__global__ void test() {
+
+  // Start
+  cooperative_groups::thread_block tb = cooperative_groups::this_thread_block();
+  tb.thread_index();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$thread_block$$num_threads.cu

@@ -0,0 +1,11 @@
+#include <cooperative_groups.h>
+#include <cooperative_groups/scan.h>
+
+__global__ void test() {
+
+  // Start
+  cooperative_groups::thread_block tb = cooperative_groups::this_thread_block();
+
+  tb.num_threads();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$thread_block$$size.cu

@@ -0,0 +1,9 @@
+#include <cooperative_groups.h>
+#include <cooperative_groups/scan.h>
+
+__global__ void test() {
+  // Start
+  cooperative_groups::thread_block tb = cooperative_groups::this_thread_block();
+  tb.size();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$thread_block$$sync.cu

@@ -0,0 +1,10 @@
+#include <cooperative_groups.h>
+#include <cooperative_groups/scan.h>
+
+__global__ void test() {
+
+  // Start
+  cooperative_groups::thread_block tb = cooperative_groups::this_thread_block();
+  tb.sync();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$thread_block$$thread_index.cu

@@ -0,0 +1,11 @@
+#include <cooperative_groups.h>
+#include <cooperative_groups/scan.h>
+
+__global__ void test() {
+
+  // Start
+    cooperative_groups::thread_block tb =
+      cooperative_groups::this_thread_block();
+    tb.sync();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$thread_block$$thread_rank.cu

@@ -0,0 +1,10 @@
+#include <cooperative_groups.h>
+#include <cooperative_groups/scan.h>
+
+__global__ void test() {
+
+  // Start
+  cooperative_groups::thread_block tb = cooperative_groups::this_thread_block();
+  tb.thread_rank();
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$thread_block_tile$$meta_group_rank.cu

@@ -0,0 +1,12 @@
+#include <cooperative_groups.h>
+#include <cooperative_groups/scan.h>
+
+__global__ void test() {
+
+  cooperative_groups::thread_block block = cooperative_groups::this_thread_block();
+  cooperative_groups::thread_block_tile<32> ctile32 = cooperative_groups::tiled_partition<32>(block);
+
+  // Start
+  ctile32.meta_group_rank();// thread_block_tile<tile size>::meta_group_rank
+  // End
+}

clang/examples/DPCT/Runtime/cooperative_groups$$thread_block_tile$$meta_group_size.cu

@@ -0,0 +1,14 @@
+#include <cooperative_groups.h>
+#include <cooperative_groups/scan.h>
+
+__global__ void test() {
+
+  cooperative_groups::thread_block block =
+      cooperative_groups::this_thread_block();
+
+  // Start
+  cooperative_groups::thread_block_tile<32> ctile32 =
+      cooperative_groups::tiled_partition<32>(block);
+  ctile32.meta_group_size();
+  // End
+}