From Pascal to Ampere: What Three Generations of NVIDIA GPU Architecture Mean for CUDA Programmers

1. The SM is the Resource You Actually Manage

A block size that achieves 100% occupancy on a Pascal P100 might throttle on an Ampere A100 due to register allocation rules. A shared memory tile size that saturates Turing's cache might leave Ampere's expanded 164 KB capacity half-empty. An algorithm that was hopelessly memory-bandwidth-bound on Pascal might suddenly become compute-bound on Ampere simply because HBM2e bandwidth tripled while Tensor Core throughput jumped by an order of magnitude.

None of that is an accident. It is the direct result of how NVIDIA has systematically redesigned the Streaming Multiprocessor (SM) — the fundamental execution engine of the GPU — across the Pascal, Turing, and Ampere architectures.

2. Pascal (GP100): The Baseline Architecture

Pascal is the cleanest reference architecture for classical CUDA programming. It represents the peak of the "pure SIMT" era before dedicated matrix-multiply hardware became the defining feature of high-end GPUs.

• SM Structure: A Pascal SM (GP100 in the P100) is partitioned into two 32-core processing blocks, each with its own warp scheduler, dispatch unit, and register file — 64 CUDA cores per SM total for GP100. Consumer GP102 parts use 128 CUDA cores per SM split into four 32-core partitions.
• CUDA Cores: Pascal cores execute one FP32 fused-multiply-add (FMA) per clock per core. GP100 additionally supports full-rate FP16 ($2\times$ FP32 throughput via packed half2 operations) and 1/2-rate FP64.
• Warp Scheduling: Each 32-core partition has its own warp scheduler issuing instructions to the core array. Latency hiding depends entirely on having enough resident warps per scheduler to cover arithmetic or memory latency pipelines.
• Registers & Shared Memory: 64K 32-bit registers per SM, up to 255 registers per thread, and 64KB-96KB of configurable shared memory/L1.
• Memory Bandwidth: GP100 pairs with HBM2, delivering $\approx 720\text{ GB/s}$ on the P100 — the first architecture where memory bandwidth stopped being the universal bottleneck for well-tiled GEMMs.
• Tensor Cores: None. Pascal has no Tensor Cores. Mixed precision on Pascal means manually packing FP16 storage with FP32 accumulation via intrinsics, not dedicated matrix hardware.

3. Turing (TU104 / Tesla T4): The Inference Shift

Turing is the pivotal transitional architecture. It introduced dedicated hardware matrix units to mainstream and datacenter inference tiers, fundamentally shifting what kind of compute dominates.

• SM Structure & Concurrent INT32: Turing splits the SM into four partitions (16 FP32 + 16 INT32 cores per partition = 64 FP32 + 64 INT32 per SM). The headline architectural leap is the addition of a dedicated INT32 datapath alongside FP32 — integer and floating-point instructions execute concurrently. Address generation, loop counters, and pointer arithmetic no longer steal cycles from your FP32 math.
• Tensor Cores (2nd Gen): Turing introduces 2nd-gen Tensor Cores (8 per SM), adding INT8 and INT4 support on top of FP16 matrix-multiply-accumulate (MMA). A single Tensor Core instruction executes an entire matrix MAC tile across a warp.
• Registers & Shared Memory: 64K registers per SM, with shared memory unified with L1 in a flexible split — up to 96KB combined per SM (e.g. 64KB shared / 32KB L1).
• L1 & L2 Cache: Doubled L2 cache size relative to Pascal at comparable tiers (up to 6MB L2), significantly accelerating irregular access patterns like embedding lookups and attention gathers.
• Memory Bandwidth: The T4 uses GDDR6 ($\approx 320\text{ GB/s}$), intentionally trading bandwidth for power efficiency (70W TDP) because Tensor Core compute throughput is the targeted bottleneck.

4. Ampere (GA100 / NVIDIA A100): The Memory & Sparsity Leap

Ampere is the modern standard for large-scale training and high-throughput inference. It restructures the SM around memory hierarchy acceleration and multi-precision versatility.

• SM Structure: GA100 divides the SM into four sub-partitions, each containing dedicated warp schedulers, register files, FP32/FP64 pipelines, and 3rd-gen Tensor Cores.
• Tensor Cores (3rd Gen): Adds native support for TF32 (19-bit format providing FP32 dynamic range at near-FP16 Tensor Core speed without code changes), BF16, full FP64 Tensor Cores, and 2:4 structured sparsity (~2× throughput for pruned weight matrices).
• Shared Memory & Asynchronous Copy (cp.async): Ampere nearly doubles the shared memory ceiling to 164KB per SM and introduces cp.async — an asynchronous copy instruction that transfers global memory directly into shared memory via hardware DMA without routing through registers.
• Massive 40MB L2 Cache: L2 cache balloons to 40MB on the A100 with residency controls (cudaFuncAttributePreferredSharedMemoryCarveout), enabling developers to pin frequently-reused structures (like attention KV caches) directly in L2.
• Memory Bandwidth: A100 delivers up to ~2.0 TB/s (HBM2e) — roughly $6\times$ Turing's T4 bandwidth, supercharging bandwidth-bound operations (Softmax, RMSNorm, GELU).

5. Memory Staging: Pre-Ampere vs. Ampere

The difference between how Pascal/Turing stage memory tiles versus Ampere's hardware DMA pipeline represents one of the most critical paradigm shifts in modern CUDA programming:

6. Side-by-Side Architectural Specifications

The definitive hardware parameters across the three architectures:

7. Six Rules of Thumb for CUDA Developers

// Ampere Asynchronous Tile Copy via cp.async
#if __CUDA_ARCH__ >= 800
__device__ inline void cp_async_16B(void* smem_ptr, const void* gmem_ptr) {
    unsigned int smem_addr = __cvta_generic_to_shared(smem_ptr);
    asm volatile("cp.async.ca.shared.global [%0], [%1], 16;\n" 
                 : : "r"(smem_addr), "l"(gmem_ptr));
}

__device__ inline void cp_async_commit_and_wait() {
    asm volatile("cp.async.commit_group;\n");
    asm volatile("cp.async.wait_group 0;\n");
}
#endif

8. Key Engineering Takeaways

Across all three generations, the trajectory is clear: NVIDIA grows on-chip resources per SM (shared memory, register flexibility, unified caches) faster than raw core counts, while shifting matrix-multiply arithmetic onto specialized Tensor Core units.

Read More Systems Research & Engineering