Absolutely. Next is one of the most important CUDA topics for AI.
Phase 8 — NVIDIA Tensor Cores
Tensor Cores are specialized GPU hardware designed primarily for high-throughput matrix/tensor operations, which are everywhere in modern AI:
1Transformer 2 ↓ 3Linear layers 4 ↓ 5Matrix multiplication 6 ↓ 7GEMM 8 ↓ 9Tensor Cores
1. What Is a Tensor Core?
A Tensor Core is specialized hardware inside an NVIDIA GPU SM designed to perform matrix multiply-accumulate operations extremely efficiently.
Conceptually:
1CUDA Core 2 ↓ 3General arithmetic
while:
1Tensor Core 2 ↓ 3Matrix/tensor operations 4 ↓ 5AI workloads
For example:
1D = A × B + C
where A, B, C, and D are matrices/tiles.
This operation appears constantly in neural networks.
2. Why AI Uses Tensor Cores
Consider a Transformer linear layer:
1Y = XW
Suppose:
1X = [batch × hidden] 2W = [hidden × output]
This becomes matrix multiplication:
1 hidden 2 ┌───────┐ 3X │ │ 4batch │ │ 5 └───────┘ 6 7 × 8 9 ┌─────────┐ 10W │ │ 11hidden │ │ 12 │ output │ 13 └─────────┘ 14 15 ↓ 16 17 ┌─────────┐ 18Y │ │ 19batch │ output │ 20 └─────────┘
This is exactly the type of computation Tensor Cores accelerate.
3. CUDA Cores vs Tensor Cores
Think of an SM roughly like:
1SM 2│ 3├── CUDA Cores 4│ 5├── Tensor Cores 6│ 7├── Registers 8│ 9├── Shared Memory / L1 10│ 11└── Warp Scheduling
CUDA Cores
General-purpose arithmetic:
1addition 2multiplication 3integer operations 4floating-point operations
Tensor Cores
Specialized matrix operations:
1matrix multiply 2+ 3accumulate
For AI matrix-heavy workloads, Tensor Cores can provide dramatically higher throughput than relying only on ordinary CUDA-core arithmetic.
4. The Core Operation
The fundamental operation is:
1D = A × B + C
This is called:
Matrix Multiply-Accumulate (MMA)
Conceptually:
1A ──┐ 2 ├── Matrix Multiply ──┐ 3B ──┘ │ 4 ├── + C 5 ↓ 6 D
This operation is fundamental to:
1Linear layers 2Attention 3MLP 4Convolution 5GEMM
5. Why Matrix Multiplication Dominates AI
A Transformer contains many operations like:
1Q = XWq 2K = XWk 3V = XWv
Then:
1Attention = softmax(QKᵀ / √d)V
And:
1MLP = W2 activation(W1X)
The major matrix multiplications are ideal Tensor Core workloads.
6. Data Types Matter
Tensor Core performance depends heavily on datatype.
Important NVIDIA AI datatypes include:
1FP32 2TF32 3FP16 4BF16 5FP8 6INT8
Their characteristics differ in:
1precision 2range 3memory usage 4Tensor Core throughput
7. FP32
FP32:
132-bit floating point
Common for:
1training 2accumulation 3general computation
Example:
1float x = 1.0f;
FP32 has relatively high precision but uses more memory than lower-precision formats.
8. FP16
FP16:
116-bit floating point
Memory:
1FP32 → 4 bytes 2FP16 → 2 bytes
So:
1FP16 2 ↓ 3half the storage 4 ↓ 5less memory traffic 6 ↓ 7higher potential throughput
CUDA provides:
1__half
through CUDA's half-precision support.
9. BF16
BF16 is another 16-bit floating-point format.
Conceptually:
1FP32 2│ 3├── larger precision 4│ 5└── BF16 6 ↓ 7 fewer fraction bits 8 but large exponent range
BF16 is particularly important in modern deep-learning training.
You'll frequently encounter:
1BF16
in modern LLM training/inference systems.
10. FP16 vs BF16
A simplified comparison:
| Type | Bits | Typical AI Use |
|---|---|---|
| FP32 | 32 | High precision |
| FP16 | 16 | AI training/inference |
| BF16 | 16 | AI training/inference |
| FP8 | 8 | Modern high-performance AI |
| INT8 | 8 | Quantized inference |
The exact supported Tensor Core modes depend on the GPU architecture.
11. TF32
TF32 is important on NVIDIA Ampere-class and later architectures.
TF32 is designed to accelerate many FP32-style matrix workloads using Tensor Cores while maintaining a useful FP32-like exponent range.
Conceptually:
1FP32 input 2 ↓ 3Tensor Core matrix operation 4 ↓ 5TF32-style computation
This is why modern frameworks can accelerate certain FP32 matrix multiplications without requiring you to manually rewrite everything as FP16.
12. Tensor Core Architecture Depends on GPU Generation
This is very important from your previous question about hardware architecture.
Tensor Core capabilities vary by architecture:
1GPU Architecture 2 ↓ 3Tensor Core generation 4 ↓ 5Supported datatypes 6 ↓ 7Supported MMA shapes 8 ↓ 9Performance characteristics
For example, NVIDIA GPU generations have progressively added capabilities such as:
1FP16 2BF16 3TF32 4INT8 5FP8 6FP4
Support and performance depend on the specific GPU.
Therefore:
Never optimize a Tensor Core kernel without knowing the target GPU architecture.
13. Your RTX 4060
Your RTX 4060 Laptop GPU is based on NVIDIA's Ada Lovelace architecture.
That means you're dealing with:
1Ada Lovelace 2 ↓ 3SMs 4 ↓ 5CUDA Cores 6Tensor Cores 7 ↓ 8AI acceleration
For your own machine, this makes Tensor Core programming especially relevant.
14. Basic Tensor Core Programming Approaches
There are several levels.
Level 1 — Framework
PyTorch:
1C = A @ B
The framework/cuBLAS may choose highly optimized GPU implementations.
Level 2 — CUDA libraries
For example:
1cuBLAS 2cuBLASLt
These provide optimized matrix operations.
Level 3 — CUDA WMMA
CUDA provides:
1WMMA
which stands for:
Warp Matrix Multiply and Accumulate
Level 4 — Low-level MMA
More advanced code can use:
1MMA
and architecture-specific instructions.
15. WMMA
WMMA allows a warp to participate in matrix operations.
Conceptually:
1Warp 2 │ 3 ├── Thread 0 4 ├── Thread 1 5 ├── Thread 2 6 │ ... 7 └── Thread 31 8 ↓ 9 Matrix operation 10 ↓ 11 Tensor Core
This connects directly to what you learned earlier:
1Warp 2 ↓ 3SIMT 4 ↓ 5Tensor Core
16. Basic WMMA Structure
You may see code conceptually like:
1#include <mma.h> 2 3using namespace nvcuda; 4 5__global__ 6void tensor_kernel() 7{ 8 wmma::fragment< 9 wmma::matrix_a, 10 16, 11 16, 12 16, 13 half, 14 wmma::row_major 15 > a; 16 17 wmma::fragment< 18 wmma::matrix_b, 19 16, 20 16, 21 16, 22 half, 23 wmma::col_major 24 > b; 25 26 wmma::fragment< 27 wmma::accumulator, 28 16, 29 16, 30 16, 31 float 32 > c; 33 34 wmma::fill_fragment(c, 0.0f); 35 36 // Load matrix tiles 37 // Perform matrix multiply 38 // Store result 39}
The exact supported fragment configurations depend on CUDA version and GPU architecture.
17. The WMMA Mental Model
Think:
1Global Memory 2 ↓ 3Load matrix tile 4 ↓ 5WMMA fragment 6 ↓ 7Tensor Core operation 8 ↓ 9Accumulator 10 ↓ 11Store
More simply:
1A tile 2 × 3B tile 4 + 5C tile 6 ↓ 7D tile
18. Why Tiles?
Tensor Cores don't generally process an arbitrarily sized giant matrix in one magical operation.
Instead, matrix multiplication is broken into tiles.
For example:
1Large Matrix 2┌─────────────────────────┐ 3│ tile │ tile │ tile │... │ 4├──────┼──────┼──────┼─────┤ 5│ tile │ tile │ tile │... │ 6├──────┼──────┼──────┼─────┤ 7│ tile │ tile │ tile │... │ 8└─────────────────────────┘
Then:
1tile A × tile B 2 ↓ 3Tensor Core
This connects to the shared-memory tiling you learned earlier.
19. Tensor Core Optimization Pipeline
The complete flow becomes:
1Global Memory 2 ↓ 3Coalesced Loads 4 ↓ 5Shared Memory 6 ↓ 7Register / Fragment 8 ↓ 9Tensor Core 10 ↓ 11Accumulator 12 ↓ 13Global Memory
This is the architecture you should remember.
20. Tensor Cores + Shared Memory
Suppose:
1A 2B
are large matrices.
Instead of repeatedly accessing global memory:
1Global Memory 2 ↓ 3Tensor Core 4 ↓ 5Global Memory
a high-performance implementation may use:
1Global Memory 2 ↓ 3Shared Memory 4 ↓ 5Register fragments 6 ↓ 7Tensor Core
This reduces expensive memory traffic and improves data reuse.
21. Tensor Cores + Warp
You've learned:
132 threads 2 ↓ 3Warp
Tensor Core operations are tightly connected to warp-level execution.
Conceptually:
1Warp 2 ├── Thread 3 ├── Thread 4 ├── ... 5 └── Thread 6 ↓ 7Tensor Core MMA
This is why your earlier Warp/SIMT knowledge is directly useful here.
22. Tensor Cores + Occupancy
Now connect:
1Tensor Core 2 ↓ 3Registers 4 ↓ 5Shared Memory 6 ↓ 7Occupancy
A highly optimized Tensor Core kernel might use substantial:
1registers 2shared memory
which can reduce occupancy.
But again:
1100% occupancy
is not automatically the goal.
The real goal is:
1Maximum useful throughput
23. Tensor Core GEMM
GEMM means:
1General Matrix-Matrix Multiplication
Usually:
1C = αAB + βC
This is one of the most important operations in AI.
For example:
1Linear Layer: 2 3Y = XW + b
is essentially GEMM plus bias.
24. Why GEMM Is So Important for Transformers
Transformer:
1Input 2 ↓ 3Linear 4 ↓ 5Attention 6 ↓ 7Linear 8 ↓ 9MLP 10 ↓ 11Linear
Many of those operations contain matrix multiplications.
Therefore:
1GEMM 2 ↓ 3Tensor Cores 4 ↓ 5Transformer performance
This is why libraries such as cuBLAS/cuBLASLt are so important in AI systems.
25. Tensor Core Utilization
Having Tensor Cores doesn't mean your kernel automatically uses them efficiently.
Potential problems:
1wrong datatype 2wrong matrix dimensions 3poor memory access 4low arithmetic intensity 5bad tiling 6register pressure 7insufficient parallelism
So:
1Tensor Core available 2 ≠ 3Tensor Core fully utilized
26. Matrix Dimensions Matter
Suppose your matrix dimensions are poorly aligned with the efficient tile sizes of the target architecture.
You may get:
1padding 2boundary handling 3underutilized hardware
A well-designed GEMM generally considers:
1M 2N 3K 4tile sizes 5datatype 6alignment 7layout
This becomes important when writing custom kernels.
27. Tensor Core Performance Model
A simplified mental model:
1Performance 2 = 3 computation throughput 4 + 5 memory efficiency 6 + 7 parallelism 8 + 9 Tensor Core utilization
If memory is the bottleneck:
1Faster Tensor Core 2 ↓ 3doesn't help much
If Tensor Cores are underutilized:
1Better data movement 2 ↓ 3more Tensor Core work 4 ↓ 5higher throughput
28. Tensor Core Kernel Design
When writing an advanced AI kernel, think:
11. Choose datatype 2 ↓ 32. Choose tile dimensions 4 ↓ 53. Load data efficiently 6 ↓ 74. Use shared memory when beneficial 8 ↓ 95. Move data into registers/fragments 10 ↓ 116. Perform MMA 12 ↓ 137. Accumulate 14 ↓ 158. Store efficiently
29. Example: Conceptual Matrix Multiply
Suppose:
1A = M × K 2B = K × N 3C = M × N
We divide:
1A → tiles 2B → tiles
Then:
1for each tile: 2 3 load A tile 4 load B tile 5 6 Tensor Core: 7 C_tile += A_tile × B_tile
Conceptually:
1for (...) 2{ 3 load_A_tile(); 4 5 load_B_tile(); 6 7 mma( 8 A_tile, 9 B_tile, 10 C_tile 11 ); 12}
That is the fundamental idea behind high-performance GEMM.
30. Tensor Core vs CUDA Core Example
Imagine a huge matrix multiplication:
14096 × 4096
A naive implementation could perform scalar/vector operations across CUDA cores.
A Tensor Core implementation processes matrix tiles through specialized matrix hardware.
Conceptually:
1CUDA Core approach: 2 3element 4element 5element 6element 7...
versus:
1Tensor Core approach: 2 3matrix tile 4 × 5matrix tile 6 ↓ 7matrix tile
The latter matches the structure of neural-network computation much better.
31. PyTorch and Tensor Cores
You don't always need to write WMMA manually.
For example:
1import torch 2 3A = torch.randn( 4 4096, 5 4096, 6 device="cuda", 7 dtype=torch.float16 8) 9 10B = torch.randn( 11 4096, 12 4096, 13 device="cuda", 14 dtype=torch.float16 15) 16 17C = A @ B
PyTorch can dispatch to optimized GPU libraries/kernels that may use Tensor Cores when the hardware, datatype, dimensions, and configuration permit.
So the practical AI engineer often works at several layers:
1PyTorch 2 ↓ 3cuBLAS/cuBLASLt 4 ↓ 5CUDA kernels 6 ↓ 7Tensor Cores 8 ↓ 9GPU hardware
32. Why You Still Need CUDA Knowledge
You might ask:
"If PyTorch already uses Tensor Cores, why learn them?"
Because when performance isn't sufficient, you need to understand:
1Why is this operation slow?
Possible answer:
1Memory bottleneck
or:
1Tensor Core not being utilized
or:
1Bad matrix dimensions
or:
1Kernel launch overhead
or:
1Poor occupancy
or:
1Warp inefficiency
CUDA knowledge lets you investigate rather than blindly changing PyTorch code.
33. Tensor Cores and LLMs
For LLMs:
1Embedding 2 ↓ 3QKV projections 4 ↓ 5Attention 6 ↓ 7Output projection 8 ↓ 9MLP 10 ↓ 11Next layer
A huge percentage of the heavy computation involves matrix operations.
Therefore:
1LLM 2 ↓ 3GEMM 4 ↓ 5Tensor Cores 6 ↓ 7GPU throughput
This is why Tensor Core knowledge is highly valuable for:
1LLM inference 2LLM training 3quantization 4custom CUDA kernels 5FlashAttention 6Deep Learning Systems 7GPU performance engineering
34. The Full Stack You Now Understand
Your CUDA knowledge is becoming much more connected:
1 AI Workload 2 ↓ 3 Transformer 4 ↓ 5 Matrix Operations 6 ↓ 7 GEMM 8 ↓ 9 ┌────────┴────────┐ 10 ↓ ↓ 11 Memory Compute 12 ↓ ↓ 13 Coalescing Tensor Cores 14 ↓ ↓ 15 Shared Memory WMMA/MMA 16 ↓ ↓ 17 Registers Accumulation 18 └────────┬────────┘ 19 ↓ 20 Warp 21 ↓ 22 SM 23 ↓ 24 GPU
This is exactly the kind of architecture-level thinking you want for AI systems/kernel engineering.
35. What You Should Master
Don't try to memorize every WMMA API immediately.
Master these concepts first:
1✅ Tensor Core purpose 2✅ CUDA Core vs Tensor Core 3✅ GEMM 4✅ MMA 5✅ FP16 6✅ BF16 7✅ TF32 8✅ FP8/INT8 concepts 9✅ Warp-level matrix operations 10✅ WMMA 11✅ Matrix tiling 12✅ Shared memory + Tensor Cores 13✅ Register fragments 14✅ Accumulators 15✅ Tensor Core utilization 16✅ GPU architecture dependence
36. Your Current AI CUDA Roadmap
You now have:
1AI CUDA Kernel Engineering 2│ 3├── Warp / SIMT execution ✅ 4├── GPU memory hierarchy ✅ 5├── Memory coalescing ✅ 6├── Shared-memory optimization ✅ 7├── Occupancy ✅ 8├── Kernel optimization ✅ 9├── CUDA streams + async execution ✅ 10├── Tensor Cores ✅ 11│ 12├── CUDA Profiling / Nsight ← NEXT 13└── Multi-GPU / NCCL
Next: CUDA Profiling with Nsight
This is the critical practical step now.
You've learned:
1How GPU works 2 ↓ 3How kernels work 4 ↓ 5How memory works 6 ↓ 7How warps work 8 ↓ 9How to optimize kernels 10 ↓ 11How streams work 12 ↓ 13How Tensor Cores work
Now you need to learn:
How do I actually prove where the GPU is slow?
You'll learn:
1Nsight Systems 2 ↓ 3Timeline analysis 4 ↓ 5CPU ↔ GPU execution 6 7Nsight Compute 8 ↓ 9Kernel analysis 10 ↓ 11Warp efficiency 12Memory throughput 13Occupancy 14Tensor Core utilization 15Registers 16Shared memory
This will turn your CUDA knowledge from theoretical → performance-engineering level.