Cooperative launch

A cooperative launch guarantees every block is resident at once, which is what makes this_grid().sync() a legal grid-wide barrier. The sim honours that by giving EVERY thread of EVERY block its own fiber up front, then running in “phases”: each phase visits the blocks in a random order and runs each one until all of its threads have reached the grid barrier (or exited); when the whole grid has arrived, every thread is released and the next phase starts.

CAVEAT (documented, not fixed): __shared__ is static in the sim, so a block’s static shared memory is one buffer reused by every block. Between two grid barriers that is still correct — a block runs to completion of the phase before the next one starts — but shared data does NOT survive a grid.sync(). Kernels that coordinate across a grid barrier do so through global memory, which is modelled exactly.

Defined in code/simulator.h.

This page documents 7 symbols: 2 macros, 5 functions.

Macros

Signature	Description	Availability	Location	Example
LAUNCH_COOP(kernel, grid, block, ...)	Portable spelling for a cooperative launch (the counterpart of LAUNCH): LAUNCH_COOP(kernel, grid, block, args…) becomes cudaLaunchCooperativeKernel under nvcc and the co-resident scheduler above …	sim + nvcc	simulator.h:1886	View
LAUNCH_DYN_COOP(kernel, grid, block, shmem, ...)	Expands to ::sim_launch_coop_shm(kernel, (grid), (block), (shmem), __VA_ARGS__)	sim + nvcc	simulator.h:1888	View

Functions

Signature	Description	Availability	Location	Example
void sim_launch_coop(F kernel, dim3 grid, dim3 block, Args... args)		sim only	simulator.h:1871
void sim_launch_coop_impl(F kernel, dim3 grid, dim3 block, size_t shmemBytes, Args... args)	CAVEAT (documented, not fixed): __shared__ is static in the sim, so a block’s static shared memory is one buffer reused by every block.	sim only	simulator.h:1718
void sim_launch_coop_shm(F kernel, dim3 grid, dim3 block, size_t shmemBytes, Args... args)		sim only	simulator.h:1876
std::vector<size_t> torder(static_cast<size_t>(numThreads))		sim only	simulator.h:1758
std::vector<Worker> workers(static_cast<size_t>(numBlocks) * numThreads)		sim only	simulator.h:1748

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.

LAUNCH_COOP (test_web_olcf_hw9_task2.cu:137)

    std::printf("cooperative grid: %d blocks of %d threads\n", nBlocks, nTPB);

    bool ok = true;

    // Test 1: nothing is removed.
    for (int i = 0; i < test_dsize; i++) h_data[i] = i;
    cudaMemcpy(d_idata, h_data.data(), tsize, cudaMemcpyHostToDevice);
    mytype   remove_val = -1;
    unsigned ds         = test_dsize;
    const mytype* c_idata = d_idata;
    LAUNCH_COOP(my_remove_if<mytype>, grid, block, c_idata, remove_val, d_odata, d_idxs, ds);
    cudaDeviceSynchronize();
    cudaMemcpy(h_data.data(), d_odata, tsize, cudaMemcpyDeviceToHost);
    for (int i = 0; i < test_dsize; i++)
        if (h_data[i] != i) {
            std::printf("mismatch 1 at %d, was: %d, should be: %d\n", i, h_data[i], i);
            ok = false; break;
        }

    // Test 2: remove the -1 entries.
    std::vector<mytype> h_in(test_dsize);
    int val = 0;

LAUNCH_DYN_COOP (test_web_cs_awbarrier.cu:191)

    int smemSize  = ((blockSize / 32) + 1) * (int)sizeof(double);
    check(cudaMalloc(reinterpret_cast<void**>(&d_partialResults), numBlocks * sizeof(double)),
          "cudaMalloc partialResults");
    check(cudaStreamSynchronize(stream), "cudaStreamSynchronize (H2D)");

    std::printf("Launching normVecByDotProductAWBarrier kernel with numBlocks = %d blockSize = %d\n",
                numBlocks, blockSize);

    dim3 dimGrid(numBlocks, 1, 1), dimBlock(blockSize, 1, 1);
    int  kSize = size;
    LAUNCH_DYN_COOP(normVecByDotProductAWBarrier, dimGrid, dimBlock, (size_t)smemSize,
                    d_vecA, d_vecB, d_partialResults, kSize);
    check(cudaDeviceSynchronize(), "cudaDeviceSynchronize");

    check(cudaMemcpy(vecA, d_vecA, sizeof(float) * size, cudaMemcpyDeviceToHost), "D2H A");

    const float expectedResult = (float)(baseVal / std::sqrt((double)size * baseVal * baseVal));
    unsigned int matches = 0;
    double maxErr = 0.0;
    for (int i = 0; i < size; i++) {
        maxErr = std::max(maxErr, std::fabs((double)vecA[i] - expectedResult));
        if ((vecA[i] - expectedResult) > 0.00001) {