Kernel launch#

blocks run one at a time in random order; a block’s threads run as fibers synchronised at __syncthreads(), visited in a fresh random order each pass.

grid/block are dim3, so full 3-D launches work; because dim3 has an unsigned-defaulting constructor, a plain LAUNCH(k, blocks, threads, …) with int arguments still compiles unchanged (1-D as the .x extent). Linear block/thread ids are decoded to uint3 with x fastest, then y, then z — the same ordering CUDA uses for lanes within a block.

Defined in code/simulator.h.

This page documents 12 symbols: 4 macros, 1 constants, 7 functions.

Macros#

Signature

Description

Availability

Location

Example

LAUNCH(kernel, grid, block, ...)

Expands to ::sim_launch(kernel, (grid), (block), __VA_ARGS__)

sim + nvcc

simulator.h:1695

View

LAUNCH_DYN(kernel, grid, block, shmem, ...)

Expands to ::sim_launch_shm(kernel, (grid), (block), (shmem), __VA_ARGS__)

sim + nvcc

simulator.h:1696

View

LAUNCH_DYN_STREAM(kernel, grid, block, shmem, stream, ...)

Expands to ::sim_launch_stream(kernel, (grid), (block), (shmem), (stream), __VA_ARGS__)

sim + nvcc

simulator.h:1701

LAUNCH_STREAM(kernel, grid, block, stream, ...)

Stream-aware launch: deferred into the stream’s queue (see drainStreams).

sim + nvcc

simulator.h:1699

View

Constants#

Signature

Description

Availability

Location

Example

const unsigned ub = static_cast<unsigned>(b)

sim only

simulator.h:1521

Functions#

Signature

Description

Availability

Location

Example

void sim_launch(F kernel, dim3 grid, dim3 block, Args... args)

Public entry points: LAUNCH forwards with no dynamic shared memory; LAUNCH_DYN passes the runtime shared-memory size (the <<<g,b,shmem>>> arg).

sim only

simulator.h:1667

void sim_launch_impl(F kernel, dim3 grid, dim3 block, size_t shmemBytes, Args... args)

grid/block are dim3, so full 3-D launches work; because dim3 has an unsigned-defaulting constructor, a plain LAUNCH(k, blocks, threads, ...) with int arguments still compiles unchanged (1-D as the …

sim only

simulator.h:1417

void sim_launch_shm(F kernel, dim3 grid, dim3 block, size_t shmemBytes, Args... args)

sim only

simulator.h:1672

void sim_launch_stream(F kernel, dim3 grid, dim3 block, size_t shmemBytes, simStream_st* stream, Args... args)

Stream-aware launch: on the NULL stream it drains and runs inline; on a real stream it is DEFERRED into that stream’s queue, like the real runtime.

sim only

simulator.h:1679

std::vector<uint3> tid3(static_cast<size_t>(numThreads))

Precompute each linear thread id’s 3-D index (x fastest).

sim only

simulator.h:1480

std::vector<size_t> torder(workers.size())

sim only

simulator.h:1491

std::vector<Worker> workers(static_cast<size_t>(numThreads))

One reusable fiber per thread, created once and reused for every block.

sim only

simulator.h:1473

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 (test_atomics.cu:120)
110}
111
112// Runs the full battery once under a given schedule seed.
113static bool runAll(unsigned seed) {
114#ifndef __CUDACC__
115    simSetSeed(seed);
116#endif
117    bool ok = true;
118
119    // Integer / unsigned-long-long / float / double counters -> exactly N.
120    { int* d = devFrom<int>(0); LAUNCH(kCounterInt, BLOCKS, BLOCK, d);
121      ok &= report("atomicAdd int counter", devRead(d) == N, seed); }
122    { unsigned long long* d = devFrom<unsigned long long>(0ULL);
123      LAUNCH(kCounterUll, BLOCKS, BLOCK, d);
124      ok &= report("atomicAdd ull counter", devRead(d) == (unsigned long long)N, seed); }
125    { float* d = devFrom<float>(0.0f); LAUNCH(kCounterFloat, BLOCKS, BLOCK, d);
126      ok &= report("atomicAdd float counter", devRead(d) == (float)N, seed); }
127    { double* d = devFrom<double>(0.0); LAUNCH(kCounterDouble, BLOCKS, BLOCK, d);
128      ok &= report("atomicAdd double counter", devRead(d) == (double)N, seed); }
129
130    // Histogram vs. serial reference.
131    {
LAUNCH_DYN (test_shared.cu:113)
103        LAUNCH(kStaticReduce, BLK, BLOCK, dIn, dOut);
104        std::vector<float> got(BLK); CUDA_CHECK(cudaMemcpy(got.data(), dOut, BLK * sizeof(float), cudaMemcpyDeviceToHost));
105        ok &= report("static __shared__ reduction", got == refBlk, seed);
106        CUDA_CHECK(cudaFree(dIn)); CUDA_CHECK(cudaFree(dOut));
107    }
108
109    // Dynamic shared reduction (size = BLOCK floats).
110    {
111        float* dIn = upload(in);
112        float* dOut = nullptr; CUDA_CHECK(cudaMalloc(&dOut, BLK * sizeof(float)));
113        LAUNCH_DYN(kDynReduce, BLK, BLOCK, BLOCK * sizeof(float), dIn, dOut);
114        std::vector<float> got(BLK); CUDA_CHECK(cudaMemcpy(got.data(), dOut, BLK * sizeof(float), cudaMemcpyDeviceToHost));
115        ok &= report("dynamic shared reduction", got == refBlk, seed);
116        CUDA_CHECK(cudaFree(dIn)); CUDA_CHECK(cudaFree(dOut));
117    }
118
119    // Dynamic shared with a DIFFERENT launch size (half the block) — the buffer
120    // must be re-sized; only the first half of each chunk is reduced.
121    {
122        const int hb = BLOCK / 2;
123        const int hblocks = N / hb;
124        std::vector<float> refH(hblocks, 0.0f);
LAUNCH_STREAM (test_web_bf_stream_host_kernel.cu:83)
73    CHECK(cudaMalloc(&d_x, M));
74    CHECK(cudaMalloc(&d_y, M));
75    CHECK(cudaMalloc(&d_z, M));
76    CHECK(cudaMemcpy(d_x, h_x.data(), M, cudaMemcpyHostToDevice));
77    CHECK(cudaMemcpy(d_y, h_y.data(), M, cudaMemcpyHostToDevice));
78
79    cudaStream_t stream;
80    CHECK(cudaStreamCreate(&stream));
81
82    // launch the kernel on the stream, then do host work "concurrently"
83    LAUNCH_STREAM(gpu_sum, grid_size, block_size, stream, d_x, d_y, d_z);
84    cpu_sum(h_x.data(), h_y.data(), h_z.data(), N);   // overlaps the async kernel
85    CHECK(cudaStreamSynchronize(stream));
86
87    std::vector<real> gpu_z(N);
88    CHECK(cudaMemcpy(gpu_z.data(), d_z, M, cudaMemcpyDeviceToHost));
89
90    // Self-check: device result and host result both equal x + y.
91    bool ok = true;
92    for (int n = 0; n < N; ++n) {
93        real ref = h_x[n] + h_y[n];
94        if (std::fabs(gpu_z[n] - ref) > 1e-4 || std::fabs(h_z[n] - ref) > 1e-4) { ok = false; break; }