1. What Is Occupancy?

Occupancy is the ratio between:

Active warps on an SM and the maximum number of warps that the SM can support.

Formula:

textCopy
1Occupancy =
2Active Warps
3────────────
4Maximum Warps

For example:

textCopy
1Maximum warps = 64
2Active warps  = 32

Then:

textCopy
1Occupancy = 32 / 64
2          = 0.5
3          = 50%

2. Why Does Occupancy Matter?

Remember the GPU execution model:

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

An SM executes many warps.

If one warp is waiting for memory:

textCopy
1Warp 0
2   ↓
3waiting for memory

the SM can potentially execute another ready warp:

textCopy
1Warp 1
2   ↓
3execute
4
5Warp 2
6   ↓
7execute
8
9Warp 3
10   ↓
11execute

This is one of the major reasons GPUs need many resident warps.

3. Occupancy Hides Latency

This is the key idea.

Suppose:

textCopy
1Warp 0
2   ↓
3global memory access
4   ↓
5waiting

Instead of leaving the SM idle:

textCopy
1SM
2 ↓
3nothing to execute

the scheduler can select another ready warp:

textCopy
1Warp 1 → execute
2Warp 2 → execute
3Warp 3 → execute

Therefore:

textCopy
1More resident warps
2        ↓
3More opportunities to hide latency
4        ↓
5Potentially better utilization

4. Occupancy Is Per SM

Occupancy isn't simply:

textCopy
1GPU = 80% occupancy

in the conceptual sense.

It is determined by how many warps/blocks can be resident on each SM under the resource limits.

For example:

textCopy
1GPU
2│
3├── SM0 → 32 active warps
4├── SM1 → 32 active warps
5├── SM2 → 32 active warps
6└── ...

Each SM has finite resources.

5. What Limits Occupancy?

The most important resources are:

textCopy
1Registers
2Shared Memory
3Threads
4Blocks
5Warps

Think:

textCopy
1             SM
2              │
3     ┌────────┼────────┐
4     ↓        ↓        ↓
5 Registers  Shared   Threads
6            Memory
7              │
8              ↓
9         Blocks/Warps
10              │
11              ↓
12          Occupancy

6. Register Limit

Every thread uses registers.

For example:

cppCopy
1__global__ void kernel()
2{
3    float a;
4    float b;
5    float c;
6    float d;
7}

The compiler may allocate registers for these variables and other intermediate values.

Suppose:

textCopy
1Registers per thread = 64

and:

textCopy
1Threads per block = 256

Then the approximate register requirement is:

textCopy
164 × 256
2=
316,384 registers

The SM has a finite register file.

Therefore, register-heavy kernels can limit how many blocks fit simultaneously.

7. Shared Memory Limit

Suppose:

cppCopy
1__global__ void kernel()
2{
3    __shared__
4    float tile[32][32];
5}

Memory:

textCopy
132 × 32 × 4 bytes
2=
34096 bytes
4=
54 KB

If each block needs:

textCopy
14 KB shared memory

and an SM can provide only a finite amount, the number of simultaneously resident blocks can be limited.

8. Thread Limit

Suppose an SM supports a maximum of:

textCopy
12048 threads

and your block has:

textCopy
1256 threads

The theoretical maximum from threads is:

textCopy
12048 / 256
2=
38 blocks

But this is only one constraint.

Registers or shared memory may reduce it further.

9. Block Limit

SMs also have a maximum number of resident blocks.

For example, suppose:

textCopy
1Maximum blocks per SM = 16

Even if registers and threads allow more blocks, you cannot exceed that architectural limit.

So:

textCopy
1Resident blocks
2=
3minimum of all resource constraints

10. The Core Formula

Conceptually:

textCopy
1Resident Blocks =
2min(
3    thread limit,
4    register limit,
5    shared-memory limit,
6    block limit,
7    other architectural limits
8)

Then:

textCopy
1Active Warps =
2Resident Blocks × Warps Per Block

And:

textCopy
1Occupancy =
2Active Warps
3────────────
4Maximum Warps Per SM

This is the most important mental model.

11. Example

Suppose:

textCopy
1Threads per block = 256

Since:

textCopy
1256 / 32 = 8

each block contains:

textCopy
18 warps

Suppose the SM can hold:

textCopy
164 warps

and resource limits allow:

textCopy
14 blocks

to be resident.

Then:

textCopy
1Active warps =
24 × 8
3=
432 warps

Occupancy:

textCopy
132 / 64
2=
350%

Therefore:

textCopy
1Occupancy = 50%

12. Another Example

Suppose:

textCopy
1Threads/block = 128

Then:

textCopy
1128 / 32 = 4 warps

If:

textCopy
18 blocks

can reside on the SM:

textCopy
1Active warps =
28 × 4
3=
432

With:

textCopy
1Maximum warps = 64

occupancy is:

textCopy
132 / 64
2=
350%

Notice something interesting:

textCopy
1256 threads/block
2× 4 blocks
3=
41024 threads
5
6128 threads/block
7× 8 blocks
8=
91024 threads

Both can have the same thread occupancy.

Block size alone does not determine occupancy.

13. Occupancy Does NOT Mean Performance

This is extremely important.

You should not assume:

textCopy
1100% occupancy
2=
3fastest kernel

That's false.

For example:

textCopy
1Kernel A
2Occupancy = 100%
3Performance = 80 TFLOPS
4
5Kernel B
6Occupancy = 50%
7Performance = 95 TFLOPS

Kernel B can be faster.

Why?

Because occupancy is only one factor.

14. Why Lower Occupancy Can Be Faster

Suppose increasing occupancy requires reducing register usage.

You might transform:

textCopy
164 registers/thread

into:

textCopy
132 registers/thread

This may allow more warps.

But perhaps the original kernel used registers efficiently and avoided expensive memory operations.

Reducing registers might cause:

textCopy
1register pressure
2        ↓
3spills
4        ↓
5local memory traffic
6        ↓
7slower kernel

So:

textCopy
1Higher occupancy
2      ≠
3Automatically better performance

15. Register Spilling

This is very important for AI kernels.

Suppose your kernel needs too many registers.

The compiler may spill some values into local memory.

Conceptually:

textCopy
1Registers
2   ↓
3not enough
4   ↓
5Local Memory
6   ↓
7additional memory traffic
8   ↓
9performance loss

So you should care about:

textCopy
1occupancy

and:

textCopy
1register pressure

together.

16. AI Kernel Example

Consider a simplified attention kernel:

textCopy
1Q
2K
3V
4 ↓
5Attention
6 ↓
7Output

The kernel may keep intermediate values in registers:

textCopy
1Q values
2K values
3partial results
4accumulators

This can require many registers per thread.

Therefore:

textCopy
1More registers/thread
2        ↓
3Fewer resident warps
4        ↓
5Lower occupancy

But those registers can also provide valuable data reuse.

So forcing maximum occupancy may be counterproductive.

17. Occupancy vs Register Reuse

Imagine:

textCopy
1Kernel A
264 registers/thread
350% occupancy

and:

textCopy
1Kernel B
232 registers/thread
3100% occupancy

But Kernel A has:

textCopy
1less memory traffic

while Kernel B has:

textCopy
1more loads/stores

Kernel A could easily win.

The real question is:

Does additional occupancy improve latency hiding enough to compensate for the additional resource usage?

18. Block Size

Common CUDA block sizes include:

textCopy
1128
2256
3512
41024

A common starting point for AI kernels is:

textCopy
1256 threads/block

because:

textCopy
1256 / 32
2=
38 warps

But this is not a universal best value.

You benchmark.

19. Why Warp Size Matters

You already learned:

textCopy
1Warp = 32 threads

Therefore:

textCopy
1128 threads
2=
34 warps
4
5256 threads
6=
78 warps
8
9512 threads
10=
1116 warps
12
131024 threads
14=
1532 warps

Try to think about block size in terms of warps.

For example:

textCopy
1256 threads

really means:

textCopy
18 warps

20. Occupancy and AI Workloads

Different AI kernels have different needs.

Elementwise kernel

Example:

textCopy
1Y = X + bias

Often:

textCopy
1low register usage
2high memory traffic

High occupancy can help hide memory latency.

GEMM

Matrix multiplication often has:

textCopy
1high arithmetic intensity

and may use:

textCopy
1shared memory
2register tiling
3Tensor Cores

Maximum occupancy isn't necessarily the main goal.

Attention

Attention kernels can have:

textCopy
1high register usage
2shared memory
3data reuse
4complex synchronization

Again:

textCopy
1100% occupancy

is not necessarily optimal.

21. Occupancy and Tensor Cores

This becomes especially important later.

Tensor Core kernels can use:

textCopy
1registers
2shared memory
3warp-level operations
4Tensor Core instructions

Resource consumption can be significant.

A Tensor Core kernel might have:

textCopy
1moderate occupancy

while still achieving excellent performance.

This is why later you'll combine:

textCopy
1Occupancy
2+
3Tensor Core utilization
4+
5Memory bandwidth
6+
7Instruction throughput

22. Occupancy Calculation Workflow

When analyzing a kernel:

Step 1

Find:

textCopy
1Threads per block

Step 2

Calculate:

textCopy
1Warps per block =
2ceil(threads / 32)

For a warp-aligned block:

textCopy
1256 / 32 = 8

Step 3

Find resource usage:

textCopy
1Registers/thread
2Shared memory/block

Step 4

Determine:

textCopy
1Maximum resident blocks

Step 5

Calculate:

textCopy
1Active warps

Step 6

Calculate:

textCopy
1Occupancy

23. CUDA Occupancy API

CUDA provides occupancy APIs.

For example:

cppCopy
1#include <cuda_runtime.h>
2#include <stdio.h>
3
4__global__ void my_kernel(float* data)
5{
6    int i =
7        blockIdx.x * blockDim.x
8        + threadIdx.x;
9
10    data[i] *= 2.0f;
11}
12
13int main()
14{
15    int blockSize;
16    int minGridSize;
17
18    cudaOccupancyMaxPotentialBlockSize(
19        &minGridSize,
20        &blockSize,
21        my_kernel,
22        0,
23        0
24    );
25
26    printf(
27        "Suggested block size: %d\n",
28        blockSize
29    );
30
31    printf(
32        "Minimum grid size: %d\n",
33        minGridSize
34    );
35
36    return 0;
37}

This lets CUDA estimate a block size that can provide good theoretical occupancy.

But remember:

The API optimizes occupancy, not necessarily end-to-end kernel performance.

You still benchmark.

24. NVIDIA Occupancy Calculator

NVIDIA provides occupancy analysis tools as part of the CUDA ecosystem.

They can help you examine:

textCopy
1Threads/block
2Registers/thread
3Shared memory/block
4Active blocks
5Active warps
6Occupancy

But for real kernel optimization, you'll eventually rely heavily on Nsight Compute.

We'll cover that later in your curriculum.

25. Important Command-Line Tools

For practical AI kernel optimization, you'll eventually use:

bashCopy
1ncu

for:

textCopy
1Nsight Compute

and:

bashCopy
1nsys

for:

textCopy
1Nsight Systems

Their roles are different.

Nsight Compute

Detailed kernel-level analysis:

textCopy
1occupancy
2memory
3warps
4instructions
5Tensor Cores
6SM utilization
7cache
8registers

Nsight Systems

System-level timeline:

textCopy
1CPU
2GPU
3CUDA kernels
4CUDA memcpy
5streams
6synchronization
7NCCL

We'll cover these later.

26. Occupancy vs Utilization

Do not confuse:

textCopy
1Occupancy

with:

textCopy
1GPU utilization

Occupancy

How many warps are resident relative to the maximum.

textCopy
1Resident warps
2───────────────
3Maximum warps

Utilization

How actively the GPU is doing useful work over time.

You can have:

textCopy
1100% occupancy

but poor utilization because warps are:

textCopy
1waiting
2diverging
3stalled on memory

27. A Useful Mental Model

Think of an SM as a restaurant.

textCopy
1SM
2│
3├── Tables = execution resources
4├── Customers = warps
5├── Kitchen = CUDA/Tensor cores
6├── Storage = registers/shared memory
7└── Waiters = warp schedulers

Occupancy means:

textCopy
1How many customers are available
2to keep the restaurant busy?

But:

textCopy
1100 customers

doesn't guarantee faster service if:

textCopy
1the kitchen is the bottleneck.

Likewise:

textCopy
1100% occupancy

doesn't guarantee maximum GPU performance.

28. The Most Important Rule

For AI kernel optimization:

Don't optimize occupancy in isolation.

Instead:

textCopy
1Measure
2  ↓
3Find bottleneck
4  ↓
5Change kernel
6  ↓
7Measure again

For example:

textCopy
1Kernel
2 ↓
350% occupancy
4 ↓
5Check profiler
6 ↓
7Memory-bound
8 ↓
9Improve memory access

rather than:

textCopy
150% occupancy
2 ↓
3force 100%

29. Your AI Kernel Optimization Model

You now have:

textCopy
1                 CUDA Kernel
2                      │
3        ┌─────────────┼─────────────┐
4        ↓             ↓             ↓
5      Warps        Memory        Compute
6        │             │             │
7      SIMT        Coalescing     CUDA Cores
8        │             │             │
9   Occupancy       Cache        Tensor Cores
10        │             │             │
11        └─────────────┼─────────────┘
12                      ↓
13                  Performance

This is the foundation for the next stage.

30. Practical Exercise

Create three versions of a kernel and compare them.

Version A

cppCopy
1__global__ void kernel(float* x)
2{
3    int i =
4        blockIdx.x * blockDim.x
5        + threadIdx.x;
6
7    float a = x[i];
8
9    x[i] = a * 2.0f;
10}

Version B

Add more intermediate values:

cppCopy
1__global__ void kernel(float* x)
2{
3    int i =
4        blockIdx.x * blockDim.x
5        + threadIdx.x;
6
7    float a = x[i];
8    float b = a * 2.0f;
9    float c = b + 1.0f;
10    float d = c * 3.0f;
11    float e = d + 2.0f;
12
13    x[i] = e;
14}

Version C

Build a more register-intensive computation.

Then profile them.

Look at:

textCopy
1Registers/thread
2Occupancy
3Memory throughput
4SM utilization
5Execution time

The goal isn't merely:

textCopy
1Which has highest occupancy?

Instead ask:

Which kernel is fastest, and why?

That's the mindset of a GPU kernel engineer.

31. What You Should Remember

The five most important points:

textCopy
11. Occupancy = active warps / maximum warps
2
32. Registers can limit occupancy
4
53. Shared memory can limit occupancy
6
74. Higher occupancy does NOT always mean
8   higher performance
9
105. Measure with profiling instead of guessing

32. Your AI-Focused CUDA Roadmap

You have now completed:

textCopy
1AI CUDA Kernel Engineering
2│
3├── Warp / SIMT execution              ✅
4├── GPU memory hierarchy               ✅
5├── Memory coalescing                  ✅
6├── Shared-memory optimization         ✅
7├── Occupancy                          ✅
8│
9├── Kernel optimization                ← NEXT
10├── CUDA streams + asynchronous exec
11├── Tensor Cores
12├── CUDA profiling / Nsight
13└── Multi-GPU / NCCL

Next: Kernel Optimization.

That topic will combine everything you've learned so far—warps, memory, coalescing, shared memory, registers, occupancy, synchronization, and arithmetic intensity—into a practical methodology for turning a slow CUDA kernel into a fast one.