Cooperative groups (common subset)

A shim for the widely-used cooperative_groups API so kernels written against it run on the simulator: group sync maps to the block/warp barriers and tiles map to warp-width shuffles. Covers thread_block, thread_block_tile<N> (N a power of two, 1..32), tiled_partition<N>, sync / thread_rank / size / num_threads / meta_group_*, tile shfl/shfl_up/down/xor + any/all/ballot, and cg::reduce with plus/less/greater. coalesced_group / coalesced_threads() / binary_partition / labeled_partition are modelled as per-thread singletons (see coalesced_group below). NOT modeled: grid / multi-grid groups (this_grid/grid.sync — cooperative launch), dynamic tiled_partition(g, n). A tile of N<32 syncs at warp granularity (correct when the whole block uses one tile size). Porting a real CG kernel: guard its own header include with #ifdef __CUDACC__ #include <cooperative_groups.h> // (+ …/reduce.h if it uses cg::reduce) #endif since the simulator supplies the namespace itself (see CLAUDE.md).

Defined in code/simulator.h.

This page documents 30 symbols: 3 classes, 1 constructors, 26 functions.

Classes

Signature	Description	Availability	Location	Example
class thread_block : public thread_group		sim only	simulator.h:1182
class thread_block_tile : public thread_group	A static-size sub-group of N consecutive lanes.	sim only	simulator.h:1193	View
class thread_group	A group is either the whole block (tile_ == 0) or a runtime-sized tile of tile_ consecutive lanes.	sim only	simulator.h:1170

Constructors

Signature	Description	Availability	Location	Example
thread_group() = default		sim only	simulator.h:1172

Functions

Signature	Description	Availability	Location	Example
int all(int pred) const		sim only	simulator.h:1218	View
int any(int pred) const		sim only	simulator.h:1217	View
unsigned ballot(int pred) const		sim only	simulator.h:1219	View
dim3 group_dim() const		sim only	simulator.h:1186
uint3 group_index() const		sim only	simulator.h:1184	View
unsigned meta_group_rank() const		sim only	simulator.h:1210	View
unsigned meta_group_size() const		sim only	simulator.h:1211	View
unsigned num_threads() const		sim only	simulator.h:1177
unsigned num_threads() const		sim only	simulator.h:1209
T shfl(T v, unsigned src) const		sim only	simulator.h:1213	View
T shfl_down(T v, unsigned d) const		sim only	simulator.h:1215	View
T shfl_up(T v, unsigned d) const		sim only	simulator.h:1214
T shfl_xor(T v, unsigned m) const		sim only	simulator.h:1216
unsigned size() const		sim only	simulator.h:1176	View
unsigned size() const		sim only	simulator.h:1208	View
void sync() const		sim only	simulator.h:1174	View
inline void sync(const thread_group& g)		sim only	simulator.h:1190	View
void sync() const		sim only	simulator.h:1206	View
inline thread_block this_thread_block()		sim only	simulator.h:1189	View
explicit thread_group(unsigned tile) : tile_(tile)		sim only	simulator.h:1173
uint3 thread_index() const		sim only	simulator.h:1185	View
unsigned thread_rank() const		sim only	simulator.h:1175	View
unsigned thread_rank() const		sim only	simulator.h:1207	View
static unsigned tile_mask()	The participating lanes are exactly this tile’s N lanes, so the tile rendezvous releases as soon as they arrive - other tiles of the same warp may be somewhere else entirely (independent thread …	sim only	simulator.h:1200
inline thread_block_tile<N> tiled_partition(const thread_group&)		sim only	simulator.h:1222	View
inline thread_group tiled_partition(const thread_group&, unsigned n)	Old-API dynamic partition: the width is a runtime value.	sim only	simulator.h:1225	View

Examples

Code from the repository’s example corpus (code\*.cu); the highlighted line is the call site. These blocks are what the View links above open.

thread_block_tile (test_web_bf_warp_cg.cu:133)

    {
        if (tid < offset)
        {
            s_y[tid] += s_y[tid + offset];
        }
        __syncthreads();
    }

    real y = s_y[tid];

    thread_block_tile<32> g = tiled_partition<32>(this_thread_block());
    for (int i = g.size() >> 1; i > 0; i >>= 1)
    {
        y += g.shfl_down(y, i);
    }

    if (tid == 0)
    {
        atomicAdd(d_y, y);
    }
}

all (test_web_warp_primitives.cu:124)

        if (v_down[t]      != e_down)    ok = false;
        if (v_xor[t]       != e_xor)     ok = false;
    }

    std::printf("threadIdx: ");   for (int t=0;t<n;++t) std::printf("%2d ", t);           std::printf("\n");
    std::printf("lane_id:   ");   for (int t=0;t<n;++t) std::printf("%2d ", lane[t]);     std::printf("\n");
    std::printf("shfl(l2):  ");   for (int t=0;t<n;++t) std::printf("%2d ", v_shfl[t]);   std::printf("\n");
    std::printf("shfl_up:   ");   for (int t=0;t<n;++t) std::printf("%2d ", v_up[t]);     std::printf("\n");
    std::printf("shfl_down: ");   for (int t=0;t<n;++t) std::printf("%2d ", v_down[t]);   std::printf("\n");
    std::printf("shfl_xor:  ");   for (int t=0;t<n;++t) std::printf("%2d ", v_xor[t]);    std::printf("\n");
    std::printf("ballot(tid>0)=%x  ballot(tid==0)=%x  all(FULL)=%d  any(FULL)=%d\n",
                ballot_pos[0], ballot_zero[0], all_full[0], any_full[0]);
    std::printf("warp primitives (%u lanes, width %u): %s\n",
                BLOCK_SIZE, WIDTH, ok ? "PASSED" : "FAILED");

    CHECK(cudaFree(d_lane)); CHECK(cudaFree(d_all)); CHECK(cudaFree(d_any));
    CHECK(cudaFree(d_shfl)); CHECK(cudaFree(d_up)); CHECK(cudaFree(d_down));
    CHECK(cudaFree(d_xor)); CHECK(cudaFree(d_bpos)); CHECK(cudaFree(d_bzero));
    return ok ? 0 : 1;
}

any (test_web_cs_reduction_full.cu:111)

    }

    // write result for this block to global mem
    if (cta.thread_rank() == 0)
        g_odata[blockIdx.x] = mySum;
}

/*
    This version is completely unrolled, unless warp shuffle is available, then
    shuffle is used within a loop. It uses a template parameter to achieve
    optimal code for any (power of 2) number of threads. (verbatim)
*/
template <class T, unsigned int blockSize> __global__ void reduce5(T *g_idata, T *g_odata, unsigned int n) {
    // Handle to thread block group
    cg::thread_block cta   = cg::this_thread_block();
    T               *sdata = SharedMemory<T>();

    // perform first level of reduction,
    // reading from global memory, writing to shared memory
    unsigned int tid = threadIdx.x;
    unsigned int i   = blockIdx.x * (blockSize * 2) + threadIdx.x;

ballot (test_web_warp_primitives.cu:124)

        if (v_down[t]      != e_down)    ok = false;
        if (v_xor[t]       != e_xor)     ok = false;
    }

    std::printf("threadIdx: ");   for (int t=0;t<n;++t) std::printf("%2d ", t);           std::printf("\n");
    std::printf("lane_id:   ");   for (int t=0;t<n;++t) std::printf("%2d ", lane[t]);     std::printf("\n");
    std::printf("shfl(l2):  ");   for (int t=0;t<n;++t) std::printf("%2d ", v_shfl[t]);   std::printf("\n");
    std::printf("shfl_up:   ");   for (int t=0;t<n;++t) std::printf("%2d ", v_up[t]);     std::printf("\n");
    std::printf("shfl_down: ");   for (int t=0;t<n;++t) std::printf("%2d ", v_down[t]);   std::printf("\n");
    std::printf("shfl_xor:  ");   for (int t=0;t<n;++t) std::printf("%2d ", v_xor[t]);    std::printf("\n");
    std::printf("ballot(tid>0)=%x  ballot(tid==0)=%x  all(FULL)=%d  any(FULL)=%d\n",
                ballot_pos[0], ballot_zero[0], all_full[0], any_full[0]);
    std::printf("warp primitives (%u lanes, width %u): %s\n",
                BLOCK_SIZE, WIDTH, ok ? "PASSED" : "FAILED");

    CHECK(cudaFree(d_lane)); CHECK(cudaFree(d_all)); CHECK(cudaFree(d_any));
    CHECK(cudaFree(d_shfl)); CHECK(cudaFree(d_up)); CHECK(cudaFree(d_down));
    CHECK(cudaFree(d_xor)); CHECK(cudaFree(d_bpos)); CHECK(cudaFree(d_bzero));
    return ok ? 0 : 1;
}

group_index (test_web_olcf_hw9_task1.cu:46)

        g.sync();
        if (lane < i) val += x[lane + i];
        g.sync();
    }
    return val;
}

__global__ void my_reduce_kernel(int* data, int* out) {
    __shared__ int sdata[nTPB];
    auto g1 = this_thread_block();
    size_t gindex = g1.group_index().x * nTPB + g1.thread_index().x;

    auto g2 = tiled_partition(g1, 32);      // 32-wide tile of the block
    auto g3 = tiled_partition(g2, 16);      // 16-wide tile of that tile
    auto g  = g3;

    int sdata_offset = (int)(g1.thread_index().x / g.size()) * (int)g.size();
    int val = reduce(g, sdata + sdata_offset, data[gindex]);

    // [SIM] publish each group's partial sum so the host can verify it.
    if (g.thread_rank() == 0) out[g1.thread_index().x / g.size()] = val;
}

meta_group_rank (test_coop_groups.cu:81)

    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    auto tile = cg::tiled_partition<32>(cg::this_thread_block());
    out[gid] = cg::reduce(tile, in[gid], cg::plus<int>());
}

// tile.thread_rank() and meta_group_rank() identify a lane and its warp.
__global__ void kTileInfo(int* rank, int* warp) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    auto tile = cg::tiled_partition<32>(cg::this_thread_block());
    rank[gid] = static_cast<int>(tile.thread_rank());
    warp[gid] = static_cast<int>(tile.meta_group_rank());
}

static int* upload(const std::vector<int>& h) {
    int* d = nullptr;
    CUDA_CHECK(cudaMalloc(&d, h.size() * sizeof(int)));
    CUDA_CHECK(cudaMemcpy(d, h.data(), h.size() * sizeof(int), cudaMemcpyHostToDevice));
    return d;
}
static bool report(const char* name, bool ok, unsigned seed) {
    std::printf("  %-26s seed=%-5u %s\n", name, seed, ok ? "PASSED" : "FAILED");
    return ok;

meta_group_size (test_web_cs_awbarrier.cu:78)

    if (tile32.thread_rank() == 0) {
        tmp[tile32.meta_group_rank()] = threadSum;
    }

    auto token = barrier.arrive();

    barrier.wait(std::move(token));

    // The warp 0 will perform last round of reduction
    if (tile32.meta_group_rank() == 0) {
        double beta = tile32.thread_rank() < tile32.meta_group_size()
                          ? tmp[tile32.thread_rank()]
                          : 0.0;

#pragma unroll
        for (int offset = tile32.size() / 2; offset > 0; offset /= 2) {
            beta += tile32.shfl_down(beta, offset);
        }

        if (tile32.thread_rank() == 0) {
            if (writeSquareRoot)
                *result = sqrt(beta);

shfl (test_web_cs_warpaggr.cu:63)

{
    cg::coalesced_group active = cg::coalesced_threads();

    // leader does the update
    int res = 0;
    if (active.thread_rank() == 0) {
        res = atomicAdd(counter, active.size());
    }

    // broadcast result
    res = active.shfl(res, 0);

    // each thread computes its own value
    return res + active.thread_rank();
}

// verbatim device code — filter positives, packing them via atomicAggInc
__global__ void filter_arr(int *dst, int *nres, const int *src, int n)
{
    int id = threadIdx.x + blockIdx.x * blockDim.x;

    for (int i = id; i < n; i += gridDim.x * blockDim.x) {

shfl_down (test_coop_groups.cu:65)

    int bsum = reduce_sum(cg::this_thread_block(), temp, in[gid]);
    if (threadIdx.x == 0) outBlk[blockIdx.x] = bsum;
}

// Warp-tile reduction via tile.shfl_down; rank 0 of each tile writes the sum.
__global__ void kTileShfl(const int* in, int* outWarp) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    auto block = cg::this_thread_block();
    auto tile  = cg::tiled_partition<32>(block);
    int v = in[gid];
    for (unsigned d = tile.size() / 2; d > 0; d >>= 1) v += tile.shfl_down(v, d);
    if (tile.thread_rank() == 0) outWarp[gid / 32u] = v;
}

// cg::reduce broadcasts the warp sum to every lane.
__global__ void kCgReduce(const int* in, int* out) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    auto tile = cg::tiled_partition<32>(cg::this_thread_block());
    out[gid] = cg::reduce(tile, in[gid], cg::plus<int>());
}

// tile.thread_rank() and meta_group_rank() identify a lane and its warp.

size (test_coop_groups.cu:43)

    } while (0)

constexpr int N     = 256;   // 8 warps
constexpr int BLOCK = 128;   // 4 warps/block
constexpr int WARPS = N / 32;

// The CoffeeBeforeArch reduce_sum: reduce a thread_group to a single value in
// thread 0, using shared scratch and group syncs.
__device__ int reduce_sum(cg::thread_group g, int* temp, int val) {
    int lane = static_cast<int>(g.thread_rank());
    for (int i = static_cast<int>(g.size()) / 2; i > 0; i /= 2) {
        temp[lane] = val;
        g.sync();
        if (lane < i) val += temp[lane + i];
        g.sync();
    }
    return val;  // full sum in rank 0
}

__global__ void kBlockReduce(const int* in, int* outBlk) {
    __shared__ int temp[BLOCK];
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;

size (test_coop_groups.cu:43)

    } while (0)

constexpr int N     = 256;   // 8 warps
constexpr int BLOCK = 128;   // 4 warps/block
constexpr int WARPS = N / 32;

// The CoffeeBeforeArch reduce_sum: reduce a thread_group to a single value in
// thread 0, using shared scratch and group syncs.
__device__ int reduce_sum(cg::thread_group g, int* temp, int val) {
    int lane = static_cast<int>(g.thread_rank());
    for (int i = static_cast<int>(g.size()) / 2; i > 0; i /= 2) {
        temp[lane] = val;
        g.sync();
        if (lane < i) val += temp[lane + i];
        g.sync();
    }
    return val;  // full sum in rank 0
}

__global__ void kBlockReduce(const int* in, int* outBlk) {
    __shared__ int temp[BLOCK];
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;

sync (test_coop_groups.cu:45)

constexpr int N     = 256;   // 8 warps
constexpr int BLOCK = 128;   // 4 warps/block
constexpr int WARPS = N / 32;

// The CoffeeBeforeArch reduce_sum: reduce a thread_group to a single value in
// thread 0, using shared scratch and group syncs.
__device__ int reduce_sum(cg::thread_group g, int* temp, int val) {
    int lane = static_cast<int>(g.thread_rank());
    for (int i = static_cast<int>(g.size()) / 2; i > 0; i /= 2) {
        temp[lane] = val;
        g.sync();
        if (lane < i) val += temp[lane + i];
        g.sync();
    }
    return val;  // full sum in rank 0
}

__global__ void kBlockReduce(const int* in, int* outBlk) {
    __shared__ int temp[BLOCK];
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    int bsum = reduce_sum(cg::this_thread_block(), temp, in[gid]);
    if (threadIdx.x == 0) outBlk[blockIdx.x] = bsum;

sync (test_coop_groups.cu:45)

constexpr int N     = 256;   // 8 warps
constexpr int BLOCK = 128;   // 4 warps/block
constexpr int WARPS = N / 32;

// The CoffeeBeforeArch reduce_sum: reduce a thread_group to a single value in
// thread 0, using shared scratch and group syncs.
__device__ int reduce_sum(cg::thread_group g, int* temp, int val) {
    int lane = static_cast<int>(g.thread_rank());
    for (int i = static_cast<int>(g.size()) / 2; i > 0; i /= 2) {
        temp[lane] = val;
        g.sync();
        if (lane < i) val += temp[lane + i];
        g.sync();
    }
    return val;  // full sum in rank 0
}

__global__ void kBlockReduce(const int* in, int* outBlk) {
    __shared__ int temp[BLOCK];
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    int bsum = reduce_sum(cg::this_thread_block(), temp, in[gid]);
    if (threadIdx.x == 0) outBlk[blockIdx.x] = bsum;

sync (test_coop_groups.cu:45)

constexpr int N     = 256;   // 8 warps
constexpr int BLOCK = 128;   // 4 warps/block
constexpr int WARPS = N / 32;

// The CoffeeBeforeArch reduce_sum: reduce a thread_group to a single value in
// thread 0, using shared scratch and group syncs.
__device__ int reduce_sum(cg::thread_group g, int* temp, int val) {
    int lane = static_cast<int>(g.thread_rank());
    for (int i = static_cast<int>(g.size()) / 2; i > 0; i /= 2) {
        temp[lane] = val;
        g.sync();
        if (lane < i) val += temp[lane + i];
        g.sync();
    }
    return val;  // full sum in rank 0
}

__global__ void kBlockReduce(const int* in, int* outBlk) {
    __shared__ int temp[BLOCK];
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    int bsum = reduce_sum(cg::this_thread_block(), temp, in[gid]);
    if (threadIdx.x == 0) outBlk[blockIdx.x] = bsum;

this_thread_block (test_coop_groups.cu:55)

        g.sync();
        if (lane < i) val += temp[lane + i];
        g.sync();
    }
    return val;  // full sum in rank 0
}

__global__ void kBlockReduce(const int* in, int* outBlk) {
    __shared__ int temp[BLOCK];
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    int bsum = reduce_sum(cg::this_thread_block(), temp, in[gid]);
    if (threadIdx.x == 0) outBlk[blockIdx.x] = bsum;
}

// Warp-tile reduction via tile.shfl_down; rank 0 of each tile writes the sum.
__global__ void kTileShfl(const int* in, int* outWarp) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    auto block = cg::this_thread_block();
    auto tile  = cg::tiled_partition<32>(block);
    int v = in[gid];
    for (unsigned d = tile.size() / 2; d > 0; d >>= 1) v += tile.shfl_down(v, d);
    if (tile.thread_rank() == 0) outWarp[gid / 32u] = v;

thread_index (test_web_olcf_hw9_task1.cu:46)

        g.sync();
        if (lane < i) val += x[lane + i];
        g.sync();
    }
    return val;
}

__global__ void my_reduce_kernel(int* data, int* out) {
    __shared__ int sdata[nTPB];
    auto g1 = this_thread_block();
    size_t gindex = g1.group_index().x * nTPB + g1.thread_index().x;

    auto g2 = tiled_partition(g1, 32);      // 32-wide tile of the block
    auto g3 = tiled_partition(g2, 16);      // 16-wide tile of that tile
    auto g  = g3;

    int sdata_offset = (int)(g1.thread_index().x / g.size()) * (int)g.size();
    int val = reduce(g, sdata + sdata_offset, data[gindex]);

    // [SIM] publish each group's partial sum so the host can verify it.
    if (g.thread_rank() == 0) out[g1.thread_index().x / g.size()] = val;
}

thread_rank (test_coop_groups.cu:42)

        }                                                                    \
    } while (0)

constexpr int N     = 256;   // 8 warps
constexpr int BLOCK = 128;   // 4 warps/block
constexpr int WARPS = N / 32;

// The CoffeeBeforeArch reduce_sum: reduce a thread_group to a single value in
// thread 0, using shared scratch and group syncs.
__device__ int reduce_sum(cg::thread_group g, int* temp, int val) {
    int lane = static_cast<int>(g.thread_rank());
    for (int i = static_cast<int>(g.size()) / 2; i > 0; i /= 2) {
        temp[lane] = val;
        g.sync();
        if (lane < i) val += temp[lane + i];
        g.sync();
    }
    return val;  // full sum in rank 0
}

__global__ void kBlockReduce(const int* in, int* outBlk) {
    __shared__ int temp[BLOCK];

thread_rank (test_coop_groups.cu:42)

        }                                                                    \
    } while (0)

constexpr int N     = 256;   // 8 warps
constexpr int BLOCK = 128;   // 4 warps/block
constexpr int WARPS = N / 32;

// The CoffeeBeforeArch reduce_sum: reduce a thread_group to a single value in
// thread 0, using shared scratch and group syncs.
__device__ int reduce_sum(cg::thread_group g, int* temp, int val) {
    int lane = static_cast<int>(g.thread_rank());
    for (int i = static_cast<int>(g.size()) / 2; i > 0; i /= 2) {
        temp[lane] = val;
        g.sync();
        if (lane < i) val += temp[lane + i];
        g.sync();
    }
    return val;  // full sum in rank 0
}

__global__ void kBlockReduce(const int* in, int* outBlk) {
    __shared__ int temp[BLOCK];

tiled_partition (test_coop_groups.cu:63)

    __shared__ int temp[BLOCK];
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    int bsum = reduce_sum(cg::this_thread_block(), temp, in[gid]);
    if (threadIdx.x == 0) outBlk[blockIdx.x] = bsum;
}

// Warp-tile reduction via tile.shfl_down; rank 0 of each tile writes the sum.
__global__ void kTileShfl(const int* in, int* outWarp) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    auto block = cg::this_thread_block();
    auto tile  = cg::tiled_partition<32>(block);
    int v = in[gid];
    for (unsigned d = tile.size() / 2; d > 0; d >>= 1) v += tile.shfl_down(v, d);
    if (tile.thread_rank() == 0) outWarp[gid / 32u] = v;
}

// cg::reduce broadcasts the warp sum to every lane.
__global__ void kCgReduce(const int* in, int* out) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    auto tile = cg::tiled_partition<32>(cg::this_thread_block());
    out[gid] = cg::reduce(tile, in[gid], cg::plus<int>());
}

tiled_partition (test_coop_groups.cu:63)

    __shared__ int temp[BLOCK];
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    int bsum = reduce_sum(cg::this_thread_block(), temp, in[gid]);
    if (threadIdx.x == 0) outBlk[blockIdx.x] = bsum;
}

// Warp-tile reduction via tile.shfl_down; rank 0 of each tile writes the sum.
__global__ void kTileShfl(const int* in, int* outWarp) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    auto block = cg::this_thread_block();
    auto tile  = cg::tiled_partition<32>(block);
    int v = in[gid];
    for (unsigned d = tile.size() / 2; d > 0; d >>= 1) v += tile.shfl_down(v, d);
    if (tile.thread_rank() == 0) outWarp[gid / 32u] = v;
}

// cg::reduce broadcasts the warp sum to every lane.
__global__ void kCgReduce(const int* in, int* out) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    auto tile = cg::tiled_partition<32>(cg::this_thread_block());
    out[gid] = cg::reduce(tile, in[gid], cg::plus<int>());
}