CUDA Thread Indexing

CUDA Thread Indexing — Complete Tutorial
(with expanded practical code examples + timing)

1. What is Thread Indexing?

In CUDA, thousands or millions of threads can execute the same kernel.
The problem is:

How does each thread know which piece of data it should process?

The answer is thread indexing.

For example, suppose:

textCopy
1A = [10, 20, 30, 40, 50]

We want:

textCopy
1Thread 0 → A[0]
2Thread 1 → A[1]
3Thread 2 → A[2]
4Thread 3 → A[3]
5Thread 4 → A[4]

CUDA provides built-in variables:

textCopy
1threadIdx
2blockIdx
3blockDim
4gridDim

2. CUDA Execution Hierarchy

textCopy
1GPU
2 │
3 └── Grid
4      │
5      ├── Block 0
6      │   ├── Thread 0
7      │   ├── Thread 1
8      │   ├── Thread 2
9      │   └── ...
10      │
11      ├── Block 1
12      │   ├── Thread 0
13      │   ├── Thread 1
14      │   └── ...
15      │
16      └── Block 2

textCopy
1grid → blocks → threads

3. threadIdx

threadIdx identifies the thread inside its block.

For 1-D:

cppCopy
1threadIdx.x

Practice – print threadIdx only:

cppCopy
1#include <stdio.h>
2#include <cuda_runtime.h>
3
4__global__ void print_threadIdx()
5{
6    printf("Block %d | threadIdx.x = %d\n", blockIdx.x, threadIdx.x);
7}
8
9int main()
10{
11    // 2 blocks × 4 threads
12    print_threadIdx<<<2, 4>>>();
13    cudaDeviceSynchronize();
14    return 0;
15}

Expected conceptual output:

textCopy
1Block 0 | threadIdx.x = 0
2Block 0 | threadIdx.x = 1
3Block 0 | threadIdx.x = 2
4Block 0 | threadIdx.x = 3
5Block 1 | threadIdx.x = 0   ← starts again from 0
6Block 1 | threadIdx.x = 1
7...

threadIdx.x alone is not globally unique.

4. blockIdx

blockIdx identifies the block inside the grid.

Practice:

cppCopy
1__global__ void print_blockIdx()
2{
3    printf("blockIdx.x = %d\n", blockIdx.x);
4}
5
6int main()
7{
8    print_blockIdx<<<4, 1>>>();
9    cudaDeviceSynchronize();
10    return 0;
11}

5. blockDim

blockDim tells you the size of the block.

cppCopy
1__global__ void print_blockDim()
2{
3    printf("blockDim.x = %d\n", blockDim.x);
4}
5
6int main()
7{
8    print_blockDim<<<2, 256>>>();   // every thread sees 256
9    cudaDeviceSynchronize();
10    return 0;
11}

6. gridDim

gridDim tells you the number of blocks.

cppCopy
1__global__ void print_gridDim()
2{
3    printf("Block %d of %d\n", blockIdx.x, gridDim.x);
4}
5
6int main()
7{
8    print_gridDim<<<4, 32>>>();
9    cudaDeviceSynchronize();
10    return 0;
11}

7. The Four Important Variables

8–10. The Most Important CUDA Formula

cppCopy
1int i = blockIdx.x * blockDim.x + threadIdx.x;

Complete demonstration with verification:

cppCopy
1#include <stdio.h>
2#include <cuda_runtime.h>
3
4__global__ void print_global_id()
5{
6    int i = blockIdx.x * blockDim.x + threadIdx.x;
7    printf("Block %d, Thread %d → global i = %d\n",
8           blockIdx.x, threadIdx.x, i);
9}
10
11int main()
12{
13    // 3 blocks × 4 threads = 12 threads
14    print_global_id<<<3, 4>>>();
15    cudaDeviceSynchronize();
16    return 0;
17}

Expected mapping:

textCopy
1Block 0 → i = 0 1 2 3
2Block 1 → i = 4 5 6 7
3Block 2 → i = 8 9 10 11

11–12. Vector Addition (Full Timed Example)

cppCopy
1#include <stdio.h>
2#include <cuda_runtime.h>
3
4__global__ void vector_add(const float* A, const float* B, float* C, int N)
5{
6    int i = blockIdx.x * blockDim.x + threadIdx.x;
7    if (i < N)
8        C[i] = A[i] + B[i];
9}
10
11int main()
12{
13    const int N = 1 << 20;          // 1M elements
14    size_t bytes = N * sizeof(float);
15
16    float *h_A = new float[N];
17    float *h_B = new float[N];
18    float *h_C = new float[N];
19
20    for (int i = 0; i < N; i++) {
21        h_A[i] = 1.0f;
22        h_B[i] = 2.0f;
23    }
24
25    float *d_A, *d_B, *d_C;
26    cudaMalloc(&d_A, bytes);
27    cudaMalloc(&d_B, bytes);
28    cudaMalloc(&d_C, bytes);
29
30    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
31    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);
32
33    int threads = 256;
34    int blocks  = (N + threads - 1) / threads;
35
36    // ---- Timing ----
37    cudaEvent_t start, stop;
38    cudaEventCreate(&start);
39    cudaEventCreate(&stop);
40
41    // Warm-up
42    vector_add<<<blocks, threads>>>(d_A, d_B, d_C, N);
43    cudaDeviceSynchronize();
44
45    // Timed run
46    cudaEventRecord(start);
47    vector_add<<<blocks, threads>>>(d_A, d_B, d_C, N);
48    cudaEventRecord(stop);
49    cudaEventSynchronize(stop);
50
51    float ms = 0.0f;
52    cudaEventElapsedTime(&ms, start, stop);
53    printf("Vector-add time: %.3f ms\n", ms);
54
55    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);
56    printf("C[0] = %.1f (expect 3.0)\n", h_C[0]);
57
58    // Cleanup
59    cudaEventDestroy(start);
60    cudaEventDestroy(stop);
61    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
62    delete[] h_A; delete[] h_B; delete[] h_C;
63    return 0;
64}

Why if (i < N)?
With N = 1000 and threads = 256 you launch 4 blocks → 1024 threads.
Threads 1000–1023 must not access the arrays.

13. gridDim in Practice

Already shown above. Use it when a kernel needs to know the total number of blocks (rare for simple element-wise kernels, useful for reductions, multi-pass algorithms, etc.).

14–18. 2D Indexing (Full Timed Matrix Example)

cppCopy
1#include <stdio.h>
2#include <cuda_runtime.h>
3
4__global__ void matrix_scale(float* matrix, int width, int height, float factor)
5{
6    int col = blockIdx.x * blockDim.x + threadIdx.x;
7    int row = blockIdx.y * blockDim.y + threadIdx.y;
8
9    if (row < height && col < width) {
10        int idx = row * width + col;   // row-major linear index
11        matrix[idx] *= factor;
12    }
13}
14
15int main()
16{
17    const int WIDTH  = 1024;
18    const int HEIGHT = 1024;
19    size_t bytes = WIDTH * HEIGHT * sizeof(float);
20
21    float *h_mat = new float[WIDTH * HEIGHT];
22    for (int i = 0; i < WIDTH * HEIGHT; i++)
23        h_mat[i] = 1.0f;
24
25    float *d_mat;
26    cudaMalloc(&d_mat, bytes);
27    cudaMemcpy(d_mat, h_mat, bytes, cudaMemcpyHostToDevice);
28
29    dim3 threads(16, 16);                     // 256 threads per block
30    dim3 blocks( (WIDTH  + 15) / 16,
31                 (HEIGHT + 15) / 16 );
32
33    cudaEvent_t start, stop;
34    cudaEventCreate(&start);
35    cudaEventCreate(&stop);
36
37    // Warm-up
38    matrix_scale<<<blocks, threads>>>(d_mat, WIDTH, HEIGHT, 2.0f);
39    cudaDeviceSynchronize();
40
41    // Timed
42    cudaEventRecord(start);
43    matrix_scale<<<blocks, threads>>>(d_mat, WIDTH, HEIGHT, 2.0f);
44    cudaEventRecord(stop);
45    cudaEventSynchronize(stop);
46
47    float ms = 0.0f;
48    cudaEventElapsedTime(&ms, start, stop);
49    printf("2-D matrix scale time: %.3f ms\n", ms);
50
51    cudaMemcpy(h_mat, d_mat, bytes, cudaMemcpyDeviceToHost);
52    printf("h_mat[0] = %.1f (expect 4.0 after two multiplies)\n", h_mat[0]);
53
54    cudaEventDestroy(start);
55    cudaEventDestroy(stop);
56    cudaFree(d_mat);
57    delete[] h_mat;
58    return 0;
59}

Linear index formula reminder:

cppCopy
1index = row * width + col;

19. 3D Indexing

cppCopy
1int x = blockIdx.x * blockDim.x + threadIdx.x;
2int y = blockIdx.y * blockDim.y + threadIdx.y;
3int z = blockIdx.z * blockDim.z + threadIdx.z;

Useful for volumetric data / 3-D tensors.
Launch with dim3 of rank 3.

20. 1D vs 2D vs 3D

21. Thread Indexing and Warps

cppCopy
1kernel<<<1, 256>>>();

→ 8 warps:

textCopy
1Warp 0 → threads 0–31
2Warp 1 → threads 32–63
3...
4Warp 7 → threads 224–255

Good consecutive indexing → good coalescing.

22. Thread Indexing and Memory Access

Consecutive global IDs inside a warp produce contiguous addresses → coalesced loads/stores.

23–24. Common Beginner Mistakes

Wrong:

cppCopy
1int i = threadIdx.x;          // resets every block → collisions

Correct:

cppCopy
1int i = blockIdx.x * blockDim.x + threadIdx.x;

Do not confuse blockIdx with threadIdx.

25. Complete Mental Model

textCopy
1                 GRID
2                  │
3             gridDim
4                  │
5        ┌─────────┴─────────┐
6     Block 0             Block 1
7        │                   │
8    blockIdx.x          blockIdx.x
9        │                   │
10    blockDim.x          blockDim.x
11        │                   │
12   threadIdx.x         threadIdx.x

Global position:

cppCopy
1int i = blockIdx.x * blockDim.x + threadIdx.x;

26. CUDA Indexing Cheat Sheet

1D

cppCopy
1int i = blockIdx.x * blockDim.x + threadIdx.x;

2D

cppCopy
1int row = blockIdx.y * blockDim.y + threadIdx.y;
2int col = blockIdx.x * blockDim.x + threadIdx.x;

3D

cppCopy
1int x = blockIdx.x * blockDim.x + threadIdx.x;
2int y = blockIdx.y * blockDim.y + threadIdx.y;
3int z = blockIdx.z * blockDim.z + threadIdx.z;

27. Connection to AI Kernel Programming

textCopy
1Tensor
2   ↓
3Index mapping (this tutorial)
4   ↓
5Threads → Warps → Blocks → SMs

This mapping is the bridge between mathematics / tensors and real GPU execution.

28. What You Should Practice

Implement these yourself (use the timed skeletons above):

1. Print threadIdx.x
2. Print blockIdx.x
3. Print blockDim.x / gridDim.x
4. Generate and print global thread ID
5. Vector addition (timed)
6. Vector multiplication / scaling
7. ReLU / SiLU / GELU element-wise kernels
8. 2-D matrix scaling (timed)
9. 2-D matrix addition
10. 2-D matrix transpose (naive)
11. Simple 3-D tensor element-wise operation

Key milestone:
When you look at

cppCopy
1int i = blockIdx.x * blockDim.x + threadIdx.x;

you should immediately understand:

“This converts a thread’s local position inside its block into its unique global position across the entire 1-D grid.”

Once that is automatic, you are ready for the next major topics: memory coalescing → shared memory → __syncthreads() → reductions → tiled matrix multiplication — the real foundations of high-performance AI kernels.