CUDA Cooperative Groups Guide

// Traditional block synchronization
__global__ void kernel_traditional(float* data, int n) {
    __shared__ float sdata[256];
    int tid = threadIdx.x;

    // Phase 1
    sdata[tid] = data[blockIdx.x * blockDim.x + tid];
    __syncthreads();  // Traditional sync

    // Phase 2
    if (tid < 128) sdata[tid] += sdata[tid + 128];
    __syncthreads();

    // More phases...
}

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

__global__ void kernel_coop_groups(float* data, int n) {
    // Get thread block group
    cg::thread_block block = cg::this_thread_block();

    // Create warp-sized tile
    cg::thread_block_tile<32> warp = cg::tiled_partition<32>(block);

    // Warp shuffle using cooperative groups API
    int lane = warp.thread_rank();
    float val = data[blockIdx.x * blockDim.x + threadIdx.x];

    // Warp-level reduction
    for (int offset = 16; offset > 0; offset /= 2) {
        val += warp.shfl_down(val, offset);
    }

    // Block sync
    block.sync();

    // Coalesced group for active threads only
    if (threadIdx.x < n) {
        cg::coalesced_group active = cg::coalesced_threads();
        // Only active threads participate
        float sum = cg::reduce(active, val, cg::plus<float>());
    }
}

// Grid-wide synchronization
__global__ void grid_sync_kernel(float* data, int n) {
    cg::grid_group grid = cg::this_grid();

    // Phase 1: All blocks process their portion
    int idx = grid.thread_rank();
    if (idx < n) data[idx] *= 2.0f;

    // Grid-wide barrier - ALL blocks wait here
    grid.sync();

    // Phase 2: Now all data is updated
    // Can safely read neighbors from other blocks
}

// Launch with cooperative kernel API
void launch_grid_sync() {
    int numBlocksPerSm;
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocksPerSm,
        grid_sync_kernel, 256, 0);

    void* args[] = { &d_data, &n };
    cudaLaunchCooperativeKernel((void*)grid_sync_kernel,
        numBlocksPerSm * num_sms, 256, args);
}