Skip to content

Commit 9b15e50

Browse files
authored
[SYCLomatic] Add query-api-mapping for 43 cooperative_groups APIs (#2901)
Signed-off-by: Chen, Sheng S <sheng.s.chen@intel.com>
1 parent a9679ba commit 9b15e50

File tree

47 files changed

+976
-2
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

47 files changed

+976
-2
lines changed
Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,14 @@
1+
#include "cooperative_groups.h"
2+
#include <cooperative_groups/reduce.h>
3+
4+
__global__ void test() {
5+
int *sdata;
6+
cooperative_groups::thread_block cta =
7+
cooperative_groups::this_thread_block();
8+
const unsigned int tid = cta.thread_rank();
9+
cooperative_groups::thread_block_tile<32> tile32 = cooperative_groups::tiled_partition<32>(cta);
10+
int *idata;
11+
// Start
12+
cooperative_groups::reduce(tile32 /*thread_block_tile<32>*/, sdata[tid]/*data*/, cooperative_groups::bit_and<int>()/*cg::bit_and<T>*/);
13+
// End
14+
}
Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,14 @@
1+
#include "cooperative_groups.h"
2+
#include <cooperative_groups/reduce.h>
3+
4+
__global__ void test() {
5+
int *sdata;
6+
cooperative_groups::thread_block cta =
7+
cooperative_groups::this_thread_block();
8+
const unsigned int tid = cta.thread_rank();
9+
cooperative_groups::thread_block_tile<32> tile32 = cooperative_groups::tiled_partition<32>(cta);
10+
int *idata;
11+
// Start
12+
cooperative_groups::reduce(tile32 /*thread_block_tile<32>*/, sdata[tid]/*data*/, cooperative_groups::bit_or<int>()/*cg::bit_or<T>*/);
13+
// End
14+
}
Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,14 @@
1+
#include "cooperative_groups.h"
2+
#include <cooperative_groups/reduce.h>
3+
4+
__global__ void test() {
5+
int *sdata;
6+
cooperative_groups::thread_block cta =
7+
cooperative_groups::this_thread_block();
8+
const unsigned int tid = cta.thread_rank();
9+
cooperative_groups::thread_block_tile<32> tile32 = cooperative_groups::tiled_partition<32>(cta);
10+
int *idata;
11+
// Start
12+
cooperative_groups::reduce(tile32 /*thread_block_tile<32>*/, sdata[tid]/*data*/, cooperative_groups::bit_xor<int>()/*cg::bit_xor<T>*/);
13+
// End
14+
}
Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,10 @@
1+
// Option: --use-experimental-features=non-uniform-groups
2+
#include <cooperative_groups.h>
3+
4+
__global__ void test() {
5+
6+
// Start
7+
cooperative_groups::coalesced_group active = cooperative_groups::coalesced_threads();
8+
active.shfl(0, 0);
9+
// End
10+
}
Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,10 @@
1+
// Option: --use-experimental-features=non-uniform-groups
2+
#include <cooperative_groups.h>
3+
4+
__global__ void test() {
5+
6+
// Start
7+
cooperative_groups::coalesced_group active = cooperative_groups::coalesced_threads();
8+
active.size();
9+
// End
10+
}
Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,9 @@
1+
// Option: --use-experimental-features=non-uniform-groups
2+
#include <cooperative_groups.h>
3+
4+
__global__ void test() {
5+
// Start
6+
cooperative_groups::coalesced_group active = cooperative_groups::coalesced_threads();
7+
active.sync();
8+
// End
9+
}
Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,11 @@
1+
// Option: --use-experimental-features=non-uniform-groups
2+
#include <cooperative_groups.h>
3+
4+
5+
__global__ void test() {
6+
// Start
7+
cooperative_groups::coalesced_group active =
8+
cooperative_groups::coalesced_threads();
9+
active.thread_rank();
10+
// End
11+
}
Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,9 @@
1+
// Option: --use-experimental-features=non-uniform-groups
2+
#include <cooperative_groups.h>
3+
#include <cooperative_groups/reduce.h>
4+
5+
__global__ void test() {
6+
// Start
7+
cooperative_groups::coalesced_group active = cooperative_groups::coalesced_threads();
8+
// End
9+
}
Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,18 @@
1+
#include <cooperative_groups.h>
2+
#include <cooperative_groups/scan.h>
3+
4+
__global__ void test() {
5+
double *sdata;
6+
cooperative_groups::thread_block cta =
7+
cooperative_groups::this_thread_block();
8+
const unsigned int tid = cta.thread_rank();
9+
cooperative_groups::thread_block_tile<32> tile32 =
10+
cooperative_groups::tiled_partition<32>(cta);
11+
// Start
12+
cooperative_groups::exclusive_scan(
13+
tile32 /* type group */, sdata[tid] /* type value */,
14+
cooperative_groups::plus<double>() /* type operator */);
15+
cooperative_groups::exclusive_scan(tile32 /* type group */,
16+
sdata[tid] /* type value */);
17+
// End
18+
}
Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,14 @@
1+
#include "cooperative_groups.h"
2+
#include <cooperative_groups/reduce.h>
3+
4+
__global__ void test() {
5+
double *sdata;
6+
cooperative_groups::thread_block cta =
7+
cooperative_groups::this_thread_block();
8+
const unsigned int tid = cta.thread_rank();
9+
cooperative_groups::thread_block_tile<32> tile32 = cooperative_groups::tiled_partition<32>(cta);
10+
int *idata;
11+
// Start
12+
cooperative_groups::reduce(tile32 /*thread_block_tile<32>*/, sdata[tid]/*data*/, cooperative_groups::greater<double>()/*cg::greater<T>*/);
13+
// End
14+
}

0 commit comments

Comments
 (0)