Grid group (cooperative launch)

this_grid().sync() is a GRID-WIDE barrier. It only means anything when every block is resident, which is what a cooperative launch guarantees — so use LAUNCH_COOP / LAUNCH_DYN_COOP (see below) for kernels that call it. Under an ordinary LAUNCH the sim runs one block at a time and can only degrade it to the block barrier (which is exact for a single-block grid).

Defined in code/simulator.h.

This page documents 35 symbols: 2 classes, 3 structs, 30 functions.

Classes

Signature	Description	Availability	Location	Example
class coalesced_group : public thread_group	Coalesced group. A real coalesced_group is the set of *currently-converged* lanes of a warp; tracking that needs per-branch divergence state, which the one-fiber-at-a-time scheduler does not keep. On …	sim only	simulator.h:1270
class grid_group	this_grid().sync() is a GRID-WIDE barrier.	sim only	simulator.h:1233

Structs

Signature	Description	Availability	Location	Example
struct greater		sim only	simulator.h:1252	View
struct less		sim only	simulator.h:1251
struct plus	cg::reduce (CUDA 11+) and its common ops.	sim only	simulator.h:1250	View

Functions

Signature	Description	Availability	Location	Example
int all(int pred) const		sim only	simulator.h:1282	View
int any(int pred) const		sim only	simulator.h:1281	View
unsigned ballot(int pred) const		sim only	simulator.h:1283	View
inline coalesced_group binary_partition(const G&, bool )		sim only	simulator.h:1287	View
unsigned block_rank() const		sim only	simulator.h:1238
inline coalesced_group coalesced_threads()		sim only	simulator.h:1285	View
dim3 group_dim() const		sim only	simulator.h:1244
bool is_valid() const		sim only	simulator.h:1236
inline coalesced_group labeled_partition(const G&, L )		sim only	simulator.h:1288	View
unsigned meta_group_rank() const		sim only	simulator.h:1276	View
unsigned meta_group_size() const		sim only	simulator.h:1277	View
unsigned num_blocks() const		sim only	simulator.h:1237
unsigned num_threads() const		sim only	simulator.h:1241
unsigned num_threads() const		sim only	simulator.h:1275
inline T reduce(const thread_block_tile<N>& tile, T val, Op op)		sim only	simulator.h:1253	View
inline T reduce(const coalesced_group&, T val, Op)	reduce over a singleton coalesced group returns the thread’s own value.	sim only	simulator.h:1290	View
inline T reduce(const coalesced_group&, T val, plus<T>)		sim only	simulator.h:1291	View
T shfl (T v, unsigned ) const		sim only	simulator.h:1278	View
T shfl_down(T v, unsigned ) const		sim only	simulator.h:1280	View
T shfl_up (T v, unsigned ) const		sim only	simulator.h:1279
unsigned size() const		sim only	simulator.h:1242	View
unsigned size() const		sim only	simulator.h:1274	View
void sync() const		sim only	simulator.h:1235	View
inline void sync(const grid_group& g)		sim only	simulator.h:1247	View
inline void sync(const thread_block_tile<N>& t)	Free sync() overloads (real CG provides sync(g) for every group type).	sim only	simulator.h:1259	View
void sync() const		sim only	simulator.h:1272	View
inline void sync(const coalesced_group& g)		sim only	simulator.h:1286	View
inline grid_group this_grid()		sim only	simulator.h:1246	View
unsigned thread_rank() const		sim only	simulator.h:1243	View
unsigned thread_rank() const		sim only	simulator.h:1273	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.

greater (test_web_cs_warpaggr_multi.cu:85)

        const int bucket = srcArr[i];
        if (bucket < numOfBuckets) indicesBuckets[atomicAggIncMulti(bucket, bucketCounters)] = i;
    }
}

// Warp-aggregated atomic max in a multi-bucket layout (verbatim).
__device__ void atomicAggMaxMulti(const int bucket, int* counter, const int valueForMax) {
    cg::coalesced_group active = cg::coalesced_threads();
    auto labeledGroup = cg::labeled_partition(active, bucket);

    const int maxValueInGroup = cg::reduce(labeledGroup, valueForMax, cg::greater<int>());
    if (labeledGroup.thread_rank() == 0) atomicMax(&counter[bucket], maxValueInGroup);
}

__global__ void calculateMaxInEachBuckets(const int* srcArr, const int* valueInBuckets,
                                          int* bucketsMax, const int srcSize,
                                          const int numOfBuckets) {
    cg::grid_group grid = cg::this_grid();

    for (int i = (int)grid.thread_rank(); i < srcSize; i += (int)grid.size()) {
        const int bucket = srcArr[i];
        if (bucket < numOfBuckets) atomicAggMaxMulti(bucket, bucketsMax, valueInBuckets[i]);

plus (test_coop_groups.cu:73)

    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.
__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) {

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;
}

binary_partition (test_web_cs_binarypart.cu:63)

__global__ void oddEvenCountAndSumCG(int *inputArr, int *numOfOdds, int *sumOfOddAndEvens, unsigned int size)
{
    cg::thread_block          cta    = cg::this_thread_block();
    cg::thread_block_tile<32> tile32 = cg::tiled_partition<32>(cta);

    // [SIM] was: cg::grid_group grid = cg::this_grid();
    //            for (int i = grid.thread_rank(); i < size; i += grid.size())
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < (int)size;
         i += gridDim.x * blockDim.x) {
        int  elem    = inputArr[i];
        auto subTile = cg::binary_partition(tile32, elem & 1);
        if (elem & 1) // Odd numbers group
        {
            int oddGroupSum = cg::reduce(subTile, elem, cg::plus<int>());

            if (subTile.thread_rank() == 0) {
                // Add number of odds present in this group of Odds.
                atomicAdd(numOfOdds, subTile.size());

                // Add local reduction of odds present in this group of Odds.
                atomicAdd(&sumOfOddAndEvens[0], oddGroupSum);
            }

coalesced_threads (test_web_cs_warpaggr.cu:54)

            std::exit(1);                                                      \
        }                                                                      \
    } while (0)

const int NUM_ELEMS             = 1 << 16;   // [SIM] shrunk from 10000000
const int NUM_THREADS_PER_BLOCK = 256;       // [SIM] source used 512

// verbatim device code — warp-aggregated atomic increment
__device__ int atomicAggInc(int *counter)
{
    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

labeled_partition (test_web_cs_warpaggr_multi.cu:61)

        cudaError_t _e = (call);                                                 \
        if (_e != cudaSuccess) {                                                 \
            std::printf("CUDA error %s at line %d\n", cudaGetErrorString(_e), __LINE__); \
            g_ok = false;                                                        \
        }                                                                        \
    } while (0)

// Warp-aggregated atomic multi-bucket increment (verbatim).
__device__ int atomicAggIncMulti(const int bucket, int* counter) {
    cg::coalesced_group active = cg::coalesced_threads();
    auto labeledGroup = cg::labeled_partition(active, bucket);

    int res = 0;
    if (labeledGroup.thread_rank() == 0) res = atomicAdd(&counter[bucket], labeledGroup.size());
    res = labeledGroup.shfl(res, 0);
    return res + labeledGroup.thread_rank();
}

// Places individual value indices into their corresponding buckets (verbatim).
__global__ void mapToBuckets(const int* srcArr, int* indicesBuckets, int* bucketCounters,
                             const int srcSize, const int numOfBuckets) {
    cg::grid_group grid = cg::this_grid();

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);

reduce (test_coop_groups.cu:73)

    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.
__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) {

reduce (test_coop_groups.cu:73)

    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.
__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) {

reduce (test_coop_groups.cu:73)

    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.
__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) {

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;

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_grid (test_web_cs_awbarrier.cu:101)

                *result = beta;
        }
    }
}
#endif

__global__ void normVecByDotProductAWBarrier(float* vecA, float* vecB,
                                             double* partialResults, int size) {
#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 700
    cg::thread_block cta  = cg::this_thread_block();
    cg::grid_group   grid = cg::this_grid();

    cg::thread_block_tile<32> tile32 = cg::tiled_partition<32>(cta);

    __shared__ cuda::barrier<cuda::thread_scope_block> barrier;

    if (threadIdx.x == 0) {
        init(&barrier, blockDim.x);
    }

    cg::sync(cta);

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];