🧠 Demystifying NVIDIA’s SM–Warp–Thread Execution Model

A Developer-Friendly Deep Dive into How Modern GPUs Really Run Code (Using RTX 3090)

Understanding GPU execution from first principles with real hardware numbers

If you’ve ever wondered “How do GPUs actually execute thousands of parallel threads?” or “What does an SM or warp really mean?”, this post will finally give you crystal-clear understanding.

Instead of abstract CUDA theory, we’ll use a real GPU — the NVIDIA RTX 3090 — to explain the math, architecture, and performance characteristics that matter for actual development.

By the end of this post, you’ll understand exactly how your CUDA kernels transform into hardware execution, why warp-level thinking matters, and how to reason about GPU performance.

🟩 Part 1: The Fundamental Paradigm Shift

CPUs vs GPUs: Different Design Philosophy

CPUs are optimized for sequential logic:

• Few powerful cores (8–64 typically)
• Deep pipelines
• Large caches
• Branch prediction
• Out-of-order execution
• Low latency per operation

GPUs are optimized for massive parallel execution:

• Thousands of simpler cores
• Shallow pipelines
• Smaller caches per core
• Minimal branch prediction
• In-order execution within warps
• High throughput via parallelism

The NVIDIA RTX 3090 exemplifies this philosophy:

RTX 3090 Specifications:
├── 82 SMs physically (80 enabled)
├── 10,240 CUDA cores total
├── 122,880 maximum resident threads
└── Hardware designed to execute 32 threads at a time (warp granularity)

To understand how a kernel or shader runs on a GPU, we must understand three hierarchical layers:

1. SM (Streaming Multiprocessor) — The parallel execution unit
2. Warp — 32 threads that execute in lockstep
3. Thread — Your code’s individual execution instance

This hierarchy is the foundation of CUDA and all GPU programming models.

🟦 Part 2: The SM — GPU’s Parallel Supercomputer

What Is a Streaming Multiprocessor?

The SM is the GPU’s equivalent of a CPU core — but designed for massive parallelism rather than sequential performance.

On the RTX 3090 (Ampere GA102), each SM contains:

SM Architecture (Ampere):
├── 128 CUDA cores (FP32 units)
├── 4 warp schedulers
├── 8 dispatch units
├── 4 Tensor Cores (for AI/ML)
├── 64 KB register file
├── 128 KB L1 cache / shared memory
├── 4 texture units
└── Load/Store units

Think of each SM as a tiny supercomputer capable of managing 48 concurrent warps and 1,536 resident threads.

RTX 3090: Full GPU Composition

Total CUDA Cores Calculation:
80 SMs × 128 CUDA cores per SM = 10,240 CUDA cores 

But here’s the critical insight most developers miss:

Not all CUDA cores execute every cycle — and that’s by design.

The GPU prioritizes latency hiding and warp-level parallelism over per-cycle utilization. We’ll see exactly why this matters.

🟨 Part 3: The Warp — Fundamental Execution Unit

The Concept That Changes Everything

This is where most developers’ mental model breaks down.

A warp is a group of 32 threads that execute in lockstep.

NVIDIA GPUs (including RTX 3090) always execute at warp granularity — not per individual thread.

How Warps Work

When you launch a kernel:

kernel<<<numBlocks, threadsPerBlock>>>(); 

Here’s what happens:

1. You launch threads (e.g., 1024 threads per block)
2. GPU groups them into warps (1024 / 32 = 32 warps)
3. SM schedules and executes warps (not individual threads)

Critical Warp Properties

SIMT Execution Example

__global__ void simpleKernel(float* data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    data[idx] = data[idx] * 2.0f;  // All 32 threads execute this together
}

In this kernel:

• Threads 0–31 form Warp 0
• Threads 32–63 form Warp 1
• Each warp executes data[idx] * 2.0f for all 32 threads simultaneously

🟥 Part 4: Maximum Occupancy and Resident Resources

SM Capacity Limits (RTX 3090)

Each SM can simultaneously hold:

Per-SM Limits:
├── Maximum threads: 1,536
├── Maximum warps: 48
├── Maximum blocks: 16 (depending on resources)
└── Register file: 65,536 registers

Mathematical relationship:

1,536 threads ÷ 32 threads per warp = 48 warps per SM

GPU-Wide Capacity

Across the entire RTX 3090:

Total Resident Threads:
80 SMs × 1,536 threads = 122,880 threads

Total Resident Warps:
80 SMs × 48 warps = 3,840 warps

Why This Matters: Latency Hiding

This massive thread capacity enables GPUs’ secret weapon:

Zero-overhead context switching between warps

When Warp A stalls (waiting for memory):

1. SM instantly switches to Warp B (already resident)
2. No OS involvement
3. No cache flushing
4. Sub-nanosecond switch time

This is how GPUs hide memory latency that would cripple CPU performance.

🧩 Part 5: From Kernel Launch to Hardware Execution

Example: Real-World Kernel Launch

// Launch kernel with 80 blocks, 1024 threads per block
myKernel<<<80, 1024>>>(data);

What happens behind the scenes:

Step 1: Thread-to-Warp Mapping

Threads per block: 1,024
Warps per block: 1,024 ÷ 32 = 32 warps

Step 2: Block-to-SM Assignment

Number of blocks: 80
Number of SMs: 80
Assignment: 1 block per SM (perfect match!)

Each SM receives:

• 1 block
• 32 warps
• 1,024 threads

Step 3: Occupancy Calculation

Active warps per SM: 32
Maximum warps per SM: 48
Occupancy: 32 ÷ 48 = 66.7%

Is 66.7% occupancy bad?

No! It’s actually excellent for most workloads. Higher occupancy only helps if your kernel is latency-bound. Compute-bound kernels often perform best at 50–75% occupancy.

🟦 Part 6: Warp Scheduling and Execution Architecture

SM Scheduling Hardware (RTX 3090 Ampere)

Each SM contains:

Scheduling Components:
├── 4 warp schedulers
├── 8 dispatch units
└── Can issue instructions to 2 warps per cycle

Execution Per Cycle

Key insight:

Each cycle, the SM can execute 2 warps × 32 threads = 64 threads

Calculation:

Active threads per cycle per SM: 64 threads
CUDA cores used per cycle per SM: 64 cores
Total CUDA cores per SM: 128 cores
Utilization per cycle: 50%

Why Only 64 of 128 Cores Active?

This confuses many developers. Here’s why it’s intentional:

Ampere’s dual-issue architecture allows:

• 2 warps issued per cycle (not 4)
• Each warp uses 32 cores
• Total: 64 cores active per cycle

The remaining 64 cores are utilized via:

1. Mixed FP32+INT32 pipelines — Simultaneous float and integer ops
2. Tensor cores — Dedicated AI acceleration
3. Dual-issue opportunities — Independent instruction pairs
4. High occupancy — More warps = better utilization

This design prioritizes:

• Latency hiding through high warp count
• Instruction-level parallelism
• Mixed workload efficiency
• Power efficiency

🔥 Part 7: GPU-Wide Execution Mathematics

Active Cores Per Cycle (Full GPU)

Calculation:
80 SMs × 64 active cores per cycle = 5,120 cores per cycle

Total Hardware Cores:
10,240 CUDA cores

Theoretical Utilization:
5,120 ÷ 10,240 = 50% per cycle

Understanding “50% Utilization”

This is not inefficiency — it’s intelligent design:

GPUs achieve teraflops of throughput precisely because they prioritize parallelism over per-cycle efficiency.

🌀 Part 8: Performance Pitfalls and Optimization

Pitfall 1: Warp Divergence

The Problem:

__global__ void divergentKernel(int* data) {
    int idx = threadIdx.x;
    if (idx % 2 == 0) {
        data[idx] = computeEven(idx);  // Half warp executes
    } else {
        data[idx] = computeOdd(idx);   // Other half executes
    }
}

What happens:

1. Warp must execute computeEven() for threads 0,2,4...30
2. Then execute computeOdd() for threads 1,3,5...31
3. Both paths run serially — 2× the time

Performance impact:

Without divergence: 1 instruction cycle
With divergence: 2 instruction cycles
Performance loss: 50%

Solutions:

// Option 1: Reorganize work to avoid divergence
__global__ void betterKernel(int* data) {
    int idx = threadIdx.x;
    // Process even threads first, odd threads separately
}

// Option 2: Use warp-level primitives
__global__ void warpAwareKernel(int* data) {
    int idx = threadIdx.x;
    int lane = idx % 32;
    bool predicate = (lane % 2 == 0);
    // Use __ballot_sync(), __shfl_sync() for warp-aware code
}

Pitfall 2: Low Occupancy

Causes of low occupancy:

1. Excessive register usage

// This kernel uses 128 registers per thread
__global__ void registerHeavy() {
    double temp[30];  // 60 registers
    // Complex computation...
}

// Max threads per SM:
// 65,536 registers ÷ 128 registers per thread = 512 threads
// vs potential 1,536 threads = 33% occupancy

1. Excessive shared memory

__global__ void sharedMemHeavy() {
    __shared__ float buffer[16384];  // 64 KB
    // Only 1 block per SM (vs potential 16 blocks)
}

Optimization strategies:

// Use compiler flags to control resources
// -maxrregcount=32 limits register usage
// Launch with smaller blocks and more occupancy

__global__ __launch_bounds__(512, 2)
void optimizedKernel() {
    // __launch_bounds__ hints:
    // - 512 threads per block
    // - 2 blocks per SM minimum
}

Pitfall 3: Uncoalesced Memory Access

// Bad: Strided access pattern
__global__ void stridedAccess(float* data) {
    int idx = threadIdx.x;
    float value = data[idx * 32];  // Each thread in warp accesses different cache line
}

// Good: Coalesced access pattern
__global__ void coalescedAccess(float* data) {
    int idx = threadIdx.x;
    float value = data[idx];  // Consecutive threads access consecutive memory
}

Performance difference:

• Coalesced: 1 memory transaction per warp
• Uncoalesced: Up to 32 memory transactions per warp
• Potential: 32× speedup

🟩 Part 9: Quick Reference Tables

RTX 3090 Architecture Summary

Per-SM Resources

Execution Metrics

Occupancy Impact

🧠 Part 10: Mental Model Summary

The Complete Execution Picture

Your CUDA Kernel
       ↓
[Grid of Blocks]
       ↓
  Block → SM Assignment (1 block per SM typically)
       ↓
  Threads → Warp Grouping (32 threads per warp)
       ↓
  Warp Scheduling (48 warps resident per SM)
       ↓
  Warp Execution (2 warps active per cycle)
       ↓
  CUDA Core Execution (64 cores active per SM per cycle)
       ↓
  Results written back to memory

Key Principles for GPU Performance

1. Think in warps, not threads

• Optimize for 32-thread groups
• Avoid divergence within warps
• Coalesce memory access patterns

2. Maximize occupancy (when beneficial)

• Keep SM busy with resident warps
• Balance registers vs threads
• But don’t sacrifice per-thread performance

3. Hide latency through parallelism

• Launch enough work to keep all SMs busy
• Use enough warps to hide memory latency
• Overlap computation with memory access

4. Respect the memory hierarchy

• Coalesced global memory access
• Leverage shared memory for reuse
• Use registers for temporary values
• Understand cache behavior

One-Sentence Summary

NVIDIA GPUs like the RTX 3090 execute your code by grouping threads into 32-thread warps, scheduling up to 48 warps per SM, and executing 2 warps per cycle across 80 SMs — enabling over 122,000 threads to be resident and over 5,000 CUDA cores to be active every cycle, with zero-overhead context switching for latency hiding.

🎯 Practical Takeaways for Developers

When writing CUDA kernels:

✅ DO:

• Launch multiples of 32 threads per block
• Aim for 128–512 threads per block (4–16 warps)
• Use __launch_bounds__ to guide the compiler
• Profile with nsys and ncu (NVIDIA Nsight)
• Test different block sizes for your workload

❌ DON’T:

• Assume more occupancy is always better
• Ignore warp divergence in conditionals
• Use excessive registers or shared memory without testing
• Forget to coalesce memory accesses
• Optimize blindly without profiling

Tools for Understanding Your Kernels

# Theoretical occupancy calculator
nsight compute --occupancy-calculator

# Actual kernel profiling
ncu --set full ./myApp

# Timeline visualization
nsys profile --trace=cuda,nvtx ./myApp

📚 Further Reading

NVIDIA Official Documentation:

• CUDA C++ Programming Guide
• Ampere Architecture Whitepaper
• CUDA Best Practices Guide

Advanced Topics:

• Cooperative Groups for flexible warp programming
• Tensor Cores for AI/ML acceleration
• CUDA Graphs for kernel launch optimization
• Stream-ordered memory allocation

🚀 Conclusion

Understanding the SM-Warp-Thread hierarchy transforms you from someone who writes CUDA code to someone who architects GPU solutions.

The RTX 3090’s 10,240 cores aren’t just numbers — they represent:

• 80 independent execution engines (SMs)
• 3,840 schedulable units (warps)
• 122,880 execution contexts (threads)
• A carefully balanced architecture for throughput computing

Master these concepts, and you’ll write faster kernels, debug performance issues effectively, and make intelligent decisions about GPU architecture selection.

Now go forth and parallelize everything. 🎮⚡

Found this helpful? Share it with your team and follow for more deep dives into GPU architecture, CUDA optimization, and high-performance computing.