Efficient AI

Tail Bounds(White Board) & GPU Programming Model

Aditya Desai · August 12th 2026

Many thanks to all the resources listed in References on the main website. Almost all figures are taken from these resources; some were generated with Gemini. Some images are taken from Google Images and still need proper citation (in progress).

Example Workload: $C = A + B$

Problem. Given two matrices $A, B \in \mathbb{R}^{m \times n}$, compute $C \in \mathbb{R}^{m \times n}$ by element-wise addition and store the result:

$$C_{ij} = A_{ij} + B_{ij} \quad \text{for all } i \in [m],\; j \in [n]$$

Example Workload: $C = A + B$

Traditional C (row-major, nested loops):

void mat_add(const float *A, const float *B, float *C,
             int m, int n) {
  for (int i = 0; i < m; i++) {
    for (int j = 0; j < n; j++) {
      int idx = i * n + j;
      C[idx] = A[idx] + B[idx];
    }
  }
}

• One thread of control; $mn$ additions in sequence
• Entry $(i,j)$ lives at offset $i\cdot n + j$

Example Workload: $C = A + B$

• Massive parallelization opportunity: $mn$ additions can be done in parallel.

OpenMP Code for $C = A + B$

Shared-memory parallelism via OpenMP pragmas:

#include <omp.h>

void mat_add(const float *A, const float *B, float *C,
             int m, int n) {
  #pragma omp parallel for collapse(2)
  for (int i = 0; i < m; i++) {
    for (int j = 0; j < n; j++) {
      int idx = i * n + j;
      C[idx] = A[idx] + B[idx];
    }
  }
}

• #pragma omp parallel for — split loop iterations across threads
• collapse(2) — flatten the nested loops into one $mn$-iteration space

CPU vs. GPU architecture choice

GPU Code for $C = A + B$

CUDA C: one thread per entry

__global__ void mat_add(const float *A, const float *B,
                        float *C, int m, int n) {
  int i = blockIdx.y * blockDim.y + threadIdx.y;
  int j = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < m && j < n) {
    int idx = i * n + j;
    C[idx] = A[idx] + B[idx];
  }
}

// launch in main:
main() {
  ....
  dim3 block(16, 16);
  dim3 grid((n + 15) / 16, (m + 15) / 16);
  mat_add<<<grid, block>>>(A_d, B_d, C_d, m, n);
  ....
}

• Defines how threads are mapped to hardware.
• And thus how they are allowed to interact.

GPU systems

Streaming Multiprocessors

Block / Thread mapping to hardware

1. All threads in a block are executed on a single SM.
2. These threads can sync among each other and use shared memory for information sharing.
3. Threads across blocks cannot sync.

Warp: unit of computing and scheduling

Warp: Wrap divergence

Memory: A note

Tensor Cores

Deep learning needs matrix multiplication

Most deep-learning compute reduces to GEMMs (\(C = AB\)). One output entry needs \(K\) multiply-adds: \(C_{ij} = \sum_k A_{ik} B_{kj}\).

CUDA cores vs. Tensor Cores

Think: scalar ALU vs. a tiny matrix-multiplication engine.

What does a Tensor Core actually compute?

Conceptually, an MMA computes \( D = A B + C \)

Tile shapes and supported datatypes depend on GPU architecture and instruction.

Tensor Core evolution

Why lower precision?

From GEMM to Tensor Cores: tiling

Large GEMMs are decomposed into tiles; warps cooperate on tiles that map naturally to MMA instructions.

Warp-level matrix operations

A Tensor Core instruction is coordinated across threads in a warp.

• Fragment layout is architecture/instruction dependent.
• Higher-level APIs hide much of this complexity.

Tensor Cores and sparsity

Tensor Cores in Transformers

This is why Tensor Core throughput is central to modern LLM training and inference.

References

• Mitzenmacher and Upfal. “Probability and Computing.” Chapters 3–4.
• NVIDIA. “CUDA Programming Guide.”