Per-thread built-ins#

One instance each (single-threaded/cooperative, so the block scheduler simply rewrites these right before switching into a given thread’s fiber; no thread_local needed).

Defined in code/simulator.h.

This page documents 4 symbols: 4 constants.

Constants#

Signature

Description

Availability

Location

Example

inline dim3 blockDim{1, 1, 1}

sim only

simulator.h:422

View

inline uint3 blockIdx{0, 0, 0}

sim only

simulator.h:421

inline dim3 gridDim{1, 1, 1}

sim only

simulator.h:423

View

inline uint3 threadIdx{0, 0, 0}

One instance each (single-threaded/cooperative, so the block scheduler simply rewrites these right before switching into a given thread’s fiber; no thread_local needed).

sim only

simulator.h:420

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.

blockDim (test_web_pk_sgemm.cu:136)
126    deterministic_init(A);
127    deterministic_init(B);
128
129    float *d_A = nullptr, *d_B = nullptr, *d_C = nullptr;
130    cudaMalloc((void **)&d_A, M * K * sizeof(float));
131    cudaMalloc((void **)&d_B, K * N * sizeof(float));
132    cudaMalloc((void **)&d_C, M * N * sizeof(float));
133    cudaMemcpy(d_A, A.data(), M * K * sizeof(float), cudaMemcpyHostToDevice);
134    cudaMemcpy(d_B, B.data(), K * N * sizeof(float), cudaMemcpyHostToDevice);
135
136    dim3 blockDim(BLOCK_DIM, BLOCK_DIM);
137    dim3 gridDim((N + BLOCK_DIM - 1) / BLOCK_DIM, (M + BLOCK_DIM - 1) / BLOCK_DIM);
138
139    bool ok = true;
140    struct { const char *name; int which; } cases[2] = { {"naive", 1}, {"tiled/shared", 2} };
141    for (int ci = 0; ci < 2; ci++) {
142        // fresh C == 0 for each kernel (beta*C term depends on the initial value)
143        cudaMemcpy(d_C, C_init.data(), M * N * sizeof(float), cudaMemcpyHostToDevice);
144        if (cases[ci].which == 1)
145            LAUNCH(sgemm_kernel,    gridDim, blockDim, d_A, d_B, d_C, M, N, K, alpha, beta);   // [SIM]
146        else
147            LAUNCH(sgemm_kernel_v2, gridDim, blockDim, d_A, d_B, d_C, M, N, K, alpha, beta);   // [SIM]
gridDim (test_web_pk_sgemm.cu:137)
127    deterministic_init(B);
128
129    float *d_A = nullptr, *d_B = nullptr, *d_C = nullptr;
130    cudaMalloc((void **)&d_A, M * K * sizeof(float));
131    cudaMalloc((void **)&d_B, K * N * sizeof(float));
132    cudaMalloc((void **)&d_C, M * N * sizeof(float));
133    cudaMemcpy(d_A, A.data(), M * K * sizeof(float), cudaMemcpyHostToDevice);
134    cudaMemcpy(d_B, B.data(), K * N * sizeof(float), cudaMemcpyHostToDevice);
135
136    dim3 blockDim(BLOCK_DIM, BLOCK_DIM);
137    dim3 gridDim((N + BLOCK_DIM - 1) / BLOCK_DIM, (M + BLOCK_DIM - 1) / BLOCK_DIM);
138
139    bool ok = true;
140    struct { const char *name; int which; } cases[2] = { {"naive", 1}, {"tiled/shared", 2} };
141    for (int ci = 0; ci < 2; ci++) {
142        // fresh C == 0 for each kernel (beta*C term depends on the initial value)
143        cudaMemcpy(d_C, C_init.data(), M * N * sizeof(float), cudaMemcpyHostToDevice);
144        if (cases[ci].which == 1)
145            LAUNCH(sgemm_kernel,    gridDim, blockDim, d_A, d_B, d_C, M, N, K, alpha, beta);   // [SIM]
146        else
147            LAUNCH(sgemm_kernel_v2, gridDim, blockDim, d_A, d_B, d_C, M, N, K, alpha, beta);   // [SIM]
148        cudaMemcpy(C_gpu.data(), d_C, M * N * sizeof(float), cudaMemcpyDeviceToHost);
On this page