1. What Is a Warp?

A warp is a group of 32 CUDA threads that the NVIDIA GPU schedules and executes together.

The basic hierarchy is:

textCopy
1Grid
2 ↓
3Block
4 ↓
5Warp
6 ↓
7Threads

For example:

textCopy
1Block
2│
3├── Warp 0
4│   ├── Thread 0
5│   ├── Thread 1
6│   ├── ...
7│   └── Thread 31
8│
9├── Warp 1
10│   ├── Thread 32
11│   ├── Thread 33
12│   ├── ...
13│   └── Thread 63
14│
15└── ...

The key fact:

textCopy
11 warp = 32 threads

2. Why Does CUDA Use Warps?

Your CUDA code appears to operate at the thread level:

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

But the NVIDIA GPU hardware schedules threads in groups called warps.

Think of it as:

textCopy
1CUDA programming model:
2
3Thread
4  ↓
5Block
6  ↓
7Grid

while execution has an important additional grouping:

textCopy
1Hardware execution:
2
3Thread
4  ↓
5Warp (32 threads)
6  ↓
7SM

So:

textCopy
1Grid
2 ↓
3Blocks
4 ↓
5Warps
6 ↓
7Threads
8 ↓
9SM execution resources

3. Example: 32 Threads

Suppose:

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

You have:

textCopy
11 block
232 threads

Therefore:

textCopy
11 block
2└── 1 warp
3    ├── Thread 0
4    ├── Thread 1
5    ├── ...
6    └── Thread 31

4. Example: 64 Threads

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

You have:

textCopy
164 threads

Therefore:

textCopy
1Block
2├── Warp 0 → Threads 0–31
3└── Warp 1 → Threads 32–63

So:

textCopy
164 threads
2÷
332 threads/warp
4=
52 warps

5. Example: 256 Threads

A very common block size is:

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

Then:

textCopy
1256 / 32
2= 8 warps

Therefore:

textCopy
1Block
2│
3├── Warp 0 → 0–31
4├── Warp 1 → 32–63
5├── Warp 2 → 64–95
6├── Warp 3 → 96–127
7├── Warp 4 → 128–159
8├── Warp 5 → 160–191
9├── Warp 6 → 192–223
10└── Warp 7 → 224–255

This is why block sizes such as:

textCopy
132
264
3128
4256
5512

are common: they are multiples of the warp size.

6. What Is SIMT?

SIMT stands for:

Single Instruction, Multiple Threads

This is one of the most important concepts in NVIDIA GPU programming.

Suppose we have:

cppCopy
1C[i] = A[i] + B[i];

A warp contains 32 threads.

Conceptually:

textCopy
1Instruction:
2     ADD
3      │
4      ├── Thread 0 → A[0] + B[0]
5      ├── Thread 1 → A[1] + B[1]
6      ├── Thread 2 → A[2] + B[2]
7      ├── ...
8      └── Thread 31 → A[31] + B[31]

The operation is the same:

textCopy
1+

but each thread operates on different data.

That's the basic idea of:

textCopy
1Single Instruction
2        ↓
3Multiple Threads
4        ↓
5Different Data

7. SIMT vs CPU

A simplified CPU view might look like:

textCopy
1CPU
2│
3└── Core
4     └── instruction

A GPU uses massive thread parallelism:

textCopy
1GPU
2│
3├── Warp
4│   ├── Thread
5│   ├── Thread
6│   ├── ...
7│   └── Thread
8│
9├── Warp
10│   └── ...
11│
12└── ...

The GPU is designed to efficiently process large amounts of parallel data.

That's why workloads such as:

textCopy
1Matrix multiplication
2Vector operations
3Convolutions
4Attention
5Elementwise activations

are excellent GPU workloads.

8. Warp ID

CUDA doesn't provide a simple built-in variable called warpIdx.

But you can calculate the warp ID from the thread's global or block-local index.

For a block:

cppCopy
1int warp_id =
2    threadIdx.x / 32;

For example:

textCopy
1threadIdx.x = 0
2→ warp 0
3
4threadIdx.x = 31
5→ warp 0
6
7threadIdx.x = 32
8→ warp 1
9
10threadIdx.x = 63
11→ warp 1

9. Lane ID

Inside a warp, every thread has a lane.

Lane numbers are:

textCopy
10 → 31

You can calculate:

cppCopy
1int lane_id =
2    threadIdx.x % 32;

For example:

textCopy
1Thread 0  → Lane 0
2Thread 1  → Lane 1
3...
4Thread 31 → Lane 31

Then:

textCopy
1Thread 32 → Warp 1, Lane 0
2Thread 33 → Warp 1, Lane 1

Therefore:

textCopy
1Thread ID
2    ↓
3┌─────────────┐
4│             │
5Warp ID     Lane ID

For a simple 1D block:

cppCopy
1int warp_id =
2    threadIdx.x / 32;
3
4int lane_id =
5    threadIdx.x % 32;

10. Complete Warp/Lane Example

cppCopy
1#include <stdio.h>
2#include <cuda_runtime.h>
3
4__global__ void show_warp_info()
5{
6    int thread_id = threadIdx.x;
7
8    int warp_id =
9        thread_id / 32;
10
11    int lane_id =
12        thread_id % 32;
13
14    printf(
15        "Thread %d | Warp %d | Lane %d\n",
16        thread_id,
17        warp_id,
18        lane_id
19    );
20}
21
22int main()
23{
24    show_warp_info<<<1, 64>>>();
25
26    cudaError_t err =
27        cudaGetLastError();
28
29    if (err != cudaSuccess)
30    {
31        printf(
32            "Launch error: %s\n",
33            cudaGetErrorString(err)
34        );
35
36        return 1;
37    }
38
39    err = cudaDeviceSynchronize();
40
41    if (err != cudaSuccess)
42    {
43        printf(
44            "Execution error: %s\n",
45            cudaGetErrorString(err)
46        );
47
48        return 1;
49    }
50
51    return 0;
52}

Compile:

bashCopy
1nvcc warp_info.cu -o warp_info

Run:

bashCopy
1./warp_info

You will see entries conceptually like:

textCopy
1Thread 0  | Warp 0 | Lane 0
2Thread 1  | Warp 0 | Lane 1
3...
4Thread 31 | Warp 0 | Lane 31
5
6Thread 32 | Warp 1 | Lane 0
7Thread 33 | Warp 1 | Lane 1
8...
9Thread 63 | Warp 1 | Lane 31

The print order itself is not guaranteed.

11. Warp Scheduling

Now let's connect warps to the SM.

Suppose your GPU has an SM:

textCopy
1SM
2│
3├── Warp 0
4├── Warp 1
5├── Warp 2
6├── Warp 3
7└── ...

The SM has limited hardware resources.

It doesn't execute every thread from your entire grid simultaneously.

Instead, the GPU scheduler selects eligible warps and issues their instructions as resources allow.

Conceptually:

textCopy
1Many blocks
2     ↓
3Many warps
4     ↓
5SM scheduler
6     ↓
7Eligible warp
8     ↓
9Instruction execution

This is one reason GPUs can hide memory latency.

12. Why Have Many Warps?

Suppose Warp 0 is waiting for data from memory.

Instead of allowing the SM to sit idle:

textCopy
1Warp 0
2  ↓
3waiting
4  ↓
5GPU idle ❌

the scheduler can issue work from another eligible warp:

textCopy
1Warp 0 → waiting
2
3Warp 1 → execute
4Warp 2 → execute
5Warp 3 → execute

Conceptually:

textCopy
1Memory latency
2      ↓
3Scheduler
4      ↓
5Other eligible warps
6      ↓
7More useful work

This is called latency hiding.

It is a fundamental GPU performance concept.

13. Warp Divergence

Now we reach one of the most important optimization topics.

Consider:

cppCopy
1if (condition)
2{
3    A;
4}
5else
6{
7    B;
8}

What happens if all 32 threads in a warp agree?

textCopy
1Thread 0 → condition = true
2Thread 1 → condition = true
3...
4Thread 31 → condition = true

Excellent.

The warp can execute the true path without splitting.

But suppose:

textCopy
1Thread 0 → true
2Thread 1 → false
3Thread 2 → true
4Thread 3 → false
5...

Now the warp has divergent execution paths.

This is called:

Warp divergence

14. Simple Divergence Example

cppCopy
1__global__ void divergence_example()
2{
3    int i = threadIdx.x;
4
5    if (i % 2 == 0)
6    {
7        // Even threads
8        printf("Even\n");
9    }
10    else
11    {
12        // Odd threads
13        printf("Odd\n");
14    }
15}

Inside one warp:

textCopy
1Lane 0  → even
2Lane 1  → odd
3Lane 2  → even
4Lane 3  → odd
5...

The threads don't agree on the branch.

That creates divergence.

15. Why Divergence Can Hurt Performance

A simplified mental model is:

textCopy
1Warp
2│
3├── Threads taking TRUE path
4└── Threads taking FALSE path

The hardware must handle both paths while inactive lanes are masked appropriately.

Conceptually:

textCopy
1TRUE path
2████████████████
3some lanes active
4
5FALSE path
6████████████████
7other lanes active

The exact hardware behavior is more nuanced than "the warp executes each branch completely sequentially," especially on newer NVIDIA architectures, but divergent control flow can reduce useful work per instruction issue.

The key idea:

A warp is most efficient when its threads follow the same control flow.

16. Good vs Bad Branching

Good

cppCopy
1if (i < N)
2{
3    output[i] = input[i];
4}

For most full warps inside a sufficiently large array, all 32 threads commonly satisfy the condition.

Near the end, some lanes may become inactive, but only a boundary warp is affected.

Potentially bad

cppCopy
1if (i % 2 == 0)
2{
3    ...
4}
5else
6{
7    ...
8}

Now threads in the same warp may frequently take different paths.

17. Branch Efficiency

Branch efficiency is a performance metric that helps identify divergent branching.

Conceptually:

textCopy
1Branch efficiency
2=
3threads following the intended/common branch
4/
5total participating threads

For example, if all 32 lanes follow the same path:

textCopy
132 / 32
2= 100%

If only half follow a path:

textCopy
116 / 32
2= 50%

In real profiling, NVIDIA tools use hardware performance metrics with more precise definitions, so treat the simple ratio as an intuition rather than the exact profiler formula.

18. Complete Divergence Demonstration

cppCopy
1#include <stdio.h>
2#include <cuda_runtime.h>
3
4__global__ void divergence_kernel()
5{
6    int lane =
7        threadIdx.x % 32;
8
9    if (lane < 16)
10    {
11        printf(
12            "Lane %d -> TRUE\n",
13            lane
14        );
15    }
16    else
17    {
18        printf(
19            "Lane %d -> FALSE\n",
20            lane
21        );
22    }
23}
24
25int main()
26{
27    divergence_kernel<<<1, 32>>>();
28
29    cudaError_t err =
30        cudaGetLastError();
31
32    if (err != cudaSuccess)
33    {
34        printf(
35            "Launch error: %s\n",
36            cudaGetErrorString(err)
37        );
38
39        return 1;
40    }
41
42    err = cudaDeviceSynchronize();
43
44    if (err != cudaSuccess)
45    {
46        printf(
47            "Execution error: %s\n",
48            cudaGetErrorString(err)
49        );
50
51        return 1;
52    }
53
54    return 0;
55}

Here:

textCopy
1Lane 0–15
2    ↓
3TRUE
4
5Lane 16–31
6    ↓
7FALSE

One warp has two different control-flow groups.

This is a simple experiment to understand divergence.

19. A Better Way to Think About Divergence

Don't memorize:

"if statements are bad."

That's wrong.

This is perfectly normal:

cppCopy
1if (i < N)

The important question is:

Do threads within the same warp tend to make the same branch decision?

For example:

textCopy
1Warp 0:
2
3Thread 0 → true
4Thread 1 → true
5...
6Thread 31 → true

Good.

But:

textCopy
1Warp 0:
2
3Thread 0 → true
4Thread 1 → false
5Thread 2 → true
6Thread 3 → false
7...

Potentially inefficient.

20. Divergence in AI Kernels

This matters in AI because many kernels process large tensors.

For example:

textCopy
1Activation
2Attention
3Normalization
4Quantization
5Masking
6Top-k
7Mixture-of-Experts routing

Some operations contain conditions.

For example:

cppCopy
1if (mask[i])
2{
3    output[i] = value;
4}

If neighboring threads have similar mask values:

textCopy
1Warp
2████████████████████████████████
3all active

execution can be efficient.

If mask values are highly irregular:

textCopy
1█ █ ██ █ █ █ ██ █ █ ██ █

the warp may experience more divergent control flow.

21. Warp Divergence vs Data Parallelism

You want:

textCopy
1Warp
2│
3├── Thread 0 → same operation
4├── Thread 1 → same operation
5├── Thread 2 → same operation
6│
7└── Thread 31 → same operation

with different data:

textCopy
1Data 0
2Data 1
3Data 2
4...
5Data 31

This matches the GPU's strengths.

22. Warp Size and Block Size

A useful rule:

textCopy
1warp size = 32

Therefore:

textCopy
132 threads   → 1 warp
264 threads   → 2 warps
3128 threads  → 4 warps
4256 threads  → 8 warps
5512 threads  → 16 warps

If a block contains:

textCopy
1100 threads

then:

textCopy
1100 / 32 = 3.125

The hardware needs:

textCopy
14 warps

The final warp has inactive lanes.

Conceptually:

textCopy
1Warp 0 → 32 active
2Warp 1 → 32 active
3Warp 2 → 32 active
4Warp 3 → 4 active

This is called partial warp utilization.

23. Why Multiples of 32 Are Often Useful

Suppose:

textCopy
1Block = 256 threads

Then:

textCopy
1256 / 32 = 8 warps

Nice.

But:

textCopy
1Block = 250 threads

gives:

textCopy
1250 / 32 = 7.8125

So the block requires:

textCopy
18 warps

with some unused lanes in the last warp.

Therefore, multiples of 32 are often a good starting point.

But don't assume:

"256 is always faster than every other block size."

Actual performance depends on:

textCopy
1register usage
2shared memory
3occupancy
4memory access
5instruction mix
6architecture
7kernel workload

24. Warp-Level Mental Model

Imagine:

textCopy
1             SM
2              │
3       ┌──────┴──────┐
4       ↓             ↓
5    Warp 0         Warp 1
6       │             │
7  ┌────┴────┐   ┌────┴────┐
8  ↓         ↓   ↓         ↓
9 T0        T1  T32       T33
10  ...       ... ...       ...
11  ↓         ↓   ↓         ↓
12 T31           T63

Each warp contains:

textCopy
132 lanes

Each lane corresponds to one thread's position within that warp.

25. Important CUDA Vocabulary

You should now know these terms:

26. Complete Practical Example

Let's combine everything:

cppCopy
1#include <stdio.h>
2#include <cuda_runtime.h>
3
4__global__ void analyze_threads(int N)
5{
6    int global_id =
7        blockIdx.x * blockDim.x
8        + threadIdx.x;
9
10    int warp_id =
11        threadIdx.x / 32;
12
13    int lane_id =
14        threadIdx.x % 32;
15
16    if (global_id < N)
17    {
18        printf(
19            "Block=%d | "
20            "Thread=%d | "
21            "Global=%d | "
22            "Warp=%d | "
23            "Lane=%d\n",
24            blockIdx.x,
25            threadIdx.x,
26            global_id,
27            warp_id,
28            lane_id
29        );
30    }
31}
32
33int main()
34{
35    int N = 64;
36
37    int threads = 32;
38    int blocks = 2;
39
40    analyze_threads<<<blocks, threads>>>(N);
41
42    cudaError_t err =
43        cudaGetLastError();
44
45    if (err != cudaSuccess)
46    {
47        printf(
48            "Launch error: %s\n",
49            cudaGetErrorString(err)
50        );
51
52        return 1;
53    }
54
55    err = cudaDeviceSynchronize();
56
57    if (err != cudaSuccess)
58    {
59        printf(
60            "Execution error: %s\n",
61            cudaGetErrorString(err)
62        );
63
64        return 1;
65    }
66
67    return 0;
68}

Here:

textCopy
12 blocks
2×
332 threads
4=
564 threads

Each block contains exactly:

textCopy
11 warp

So:

textCopy
1Grid
2│
3├── Block 0
4│   └── Warp 0
5│       ├── Lane 0
6│       ├── Lane 1
7│       └── ...
8│
9└── Block 1
10    └── Warp 0
11        ├── Lane 0
12        ├── Lane 1
13        └── ...

Notice that warp_id is calculated within each block here. Therefore, both blocks have a local warp_id = 0.

27. What You Should Understand Before Moving On

You should now be able to explain this:

cppCopy
1int global_id =
2    blockIdx.x * blockDim.x
3    + threadIdx.x;
4
5int warp_id =
6    threadIdx.x / 32;
7
8int lane_id =
9    threadIdx.x % 32;

in plain English:

textCopy
1global_id
2    ↓
3Which data element does this thread process?
4
5warp_id
6    ↓
7Which warp inside this block?
8
9lane_id
10    ↓
11Which position does this thread occupy
12inside its warp?

And:

textCopy
1Thread
2   ↓
3Warp (32 threads)
4   ↓
5Block
6   ↓
7Grid
8   ↓
9GPU

28. Connection to Kernel Optimization

This topic is the foundation for the next optimization layers:

textCopy
1Warps & SIMT
2      ↓
3Warp divergence
4      ↓
5Memory coalescing
6      ↓
7Shared memory
8      ↓
9Warp-level primitives
10      ↓
11Reductions
12      ↓
13Occupancy
14      ↓
15Kernel optimization

For your AI kernel programming, the particularly important chain is:

textCopy
1Tensor
2 ↓
3Threads
4 ↓
5Warp
6 ↓
7Memory access
8 ↓
9Coalescing
10 ↓
11Shared memory
12 ↓
13Synchronization
14 ↓
15Compute utilization
16 ↓
17Performance

Confidence test

Before moving to the next topic, try answering these without looking:

1. Why does a warp contain 32 threads?
2. What is SIMT?
3. What is a lane?
4. How do you calculate a lane ID?
5. What is warp divergence?
6. Why can divergence reduce performance?
7. How many warps are in a 256-thread block?
8. How many warps are required for a 100-thread block?
9. Why can the GPU execute another warp while one warp waits?
10. Why are block sizes that are multiples of 32 often a good starting point?

If you can answer these and run the programs above, your Thread → Block → Grid → Warp → SIMT foundation is strong enough to move to Phase 2.6: CUDA Synchronization (__syncthreads()), race conditions, and thread cooperation.