← BACK TO LESSONPARALLEL HORIZONSREAD // 00

LESSON READING // 00

PROGRAMMABLE PARALLELISM

HISTORICAL ANCHOR2006–2007 · G80 / CUDAREAD TIME · ABOUT 60 MIN

Rewrite a serial loop as many threads, then see Grid, Block, and Thread as a language for describing parallel work—not as a product specification.

Write down answers as you go; the estimate includes code tracing and practice.

PLAY THE INTERACTIVE LESSON
COURSE PACINGABOUT 60 MINUTES
  1. 01ORIENT + OBJECTIVES05 MIN
  2. 02MENTAL MODEL08 MIN
  3. 03DEEP DIVES15 MIN
  4. 04CODE TRACE12 MIN
  5. 05HISTORY + ECOSYSTEM08 MIN
  6. 06PRACTICE + REVIEW12 MIN
ACCELERATED COMPUTING CAPABILITY MAPINDEPENDENT LEARNING COMPANION · NOT AN NVIDIA COURSE OR CERTIFICATION
PATH // 00
FOUNDATIONALMODERN CUDA C/C++CUDA PYTHON
THE OFFICIAL PATH EMPHASIZES

Build hands-on accelerated-computing fundamentals: identify parallel work, write kernels, configure the hierarchy, and validate correctness and performance.

HOW THIS LESSON CONNECTS

The game builds a visual Grid/Block/Thread model; the optional lab transfers that exact mapping into a real vector-add program.

SUGGESTED PREREQUISITES
  • Basic C/C++ or Python
  • Arrays, loops, and functions
  • Ability to run command-line examples
PRACTICE TOOLS
  • CUDA Toolkit / nvcc
  • CUDA Samples
  • Nsight Systems or CUDA Events
01 · 05 MIN

ORIENT + OBJECTIVES

REMEMBER THIS FIRST

Find independent work first; only then decide how many threads should express it.

BY THE END, YOU SHOULD BE ABLE TO

  1. 01

    Decide whether loop iterations are independent enough to map onto GPU threads.

  2. 02

    Compute a grid from data size and threads per block, including a safe boundary guard.

  3. 03

    Separate the programming semantics of grid, block, warp, and thread from physical scheduling.

  4. 04

    Name at least three reasons why more threads do not imply proportional speedup.

02 · 08 MIN

BUILD THE MENTAL MODEL

  1. 01
    WORK

    Every array element performs the same addition, with no dependency on another element.

  2. 02
    MAP

    One thread owns one element; blocks organize groups of threads, and a grid organizes all blocks.

  3. 03
    BOUND

    Launch sizes usually round up, so the kernel must guard its global index against the data boundary.

03 · 15 MIN

CONCEPT DEEP DIVES

01

GLOBAL THREAD INDEX

blockIdx.x × blockDim.x + threadIdx.x turns a block-local thread number into a unique index across the grid.

02

THE HIERARCHY HAS SEMANTICS

Threads in one block can synchronize and share on-chip memory. Blocks should normally be safe to schedule independently.

03

PARALLELISM IS NOT SPEEDUP

Thread count describes available work. Memory access, branches, resource use, and problem size still control observed performance.

PART 01

A LOGICAL HIERARCHY IS NOT AN SM FLOOR PLAN

Grid and block describe a kernel's logical work decomposition. A grid can contain far more blocks than the physical SM count; the runtime schedules blocks in waves as resources become available. Correct code cannot assume block 0 completes before block 1 or use an implicit cross-block barrier.

Threads inside a block execute in warps, commonly 32 threads in today's CUDA model. Warp behavior matters for branches, memory, and performance, while block scope defines shared memory and synchronization. Blocks are therefore more than arbitrary number ranges.

PAUSE AND REASONCan a grid with 10,000 blocks run correctly on a GPU with only dozens of SMs?

REFERENCE ANSWERYes. Blocks execute in waves. Physical concurrency changes elapsed time, not correctness, provided blocks do not depend on an undefined order.

PART 02

AN INDEX IS A DATA-OWNERSHIP RULE

The global index defines which thread owns which data. One-dimensional arrays often use blockIdx.x × blockDim.x + threadIdx.x. A 2D image computes x and y separately and often flattens them with y × width + x. A bad mapping causes duplicate writes, gaps, or out-of-range access.

Rounding a launch upward is normal because data sizes rarely divide evenly by block size. A boundary branch affects only a few threads in the final block and lets one kernel safely handle arbitrary n. Establish unique ownership, complete coverage, and safety before tuning.

PAUSE AND REASONWhat breaks if a 2D image kernel checks x < width but not y < height?

REFERENCE ANSWERThreads beyond the last valid row can still compute linear addresses and access out of bounds. Both dimensions need guards.

PART 03

BLOCK SIZE TRADES MULTIPLE RESOURCES

Threads per block determines warp count and combines with registers per thread and shared memory per block to limit how many blocks can reside on an SM. A larger block can reduce resident block count; a tiny block may expose too few warps to hide latency.

Occupancy is a diagnostic, not the only objective. Memory-bound, register-heavy, and compute-bound kernels may prefer different configurations. The game's 64-thread block satisfies one mapping exercise; it is not a universal optimum.

PAUSE AND REASONWhy might changing a block from 128 to 1,024 threads make a kernel slower?

REFERENCE ANSWERThe larger block may consume enough registers or shared memory to reduce resident blocks and warps, while increasing tail or branch waste.

04 · 12 MIN

PUT IT BACK INTO CODE

PROGRAM MODELOne thread per element, with a boundary guard
01__global__ void add(float* a, float* b, float* c, int n) {02  int i = blockIdx.x * blockDim.x + threadIdx.x;03  if (i < n) c[i] = a[i] + b[i];04}05int blocks = (n + threads - 1) / threads;06add<<<blocks, threads>>>(a, b, c, n);

TRACE THE CODE

  1. 01
    KERNEL__global__ DEFINES WHERE CODE RUNS

    The host launches the function and the device executes it; every thread runs the same kernel program.

  2. 02
    INDEXBUILD ONE UNIQUE i

    Block index, block width, and thread index combine into global ownership.

  3. 03
    GUARDONLY VALID THREADS WRITE

    Rounded-up launches create extra threads. The guard safely excludes them.

  4. 04
    LAUNCHCONFIGURATION FOLLOWS PROBLEM SIZE

    ceil(n / threads) covers all elements; choose block size by resource analysis and measurement after correctness.

REF · 00

KNOWLEDGE ATLAS

OPTIONAL REFERENCE · REVISIT AS NEEDED · OUTSIDE THE CORE 60 MINUTES

CORE VOCABULARY

Kernel
A host-launched function executed on the device by many threads; each runs the same program with a different index.
Grid
All blocks created by one launch—the logical problem space, not a diagram of physical GPU units.
Block
A schedulable group whose threads can share memory and use block-scoped synchronization.
Warp
The basic thread execution group on an SM, typically 32 threads; branches and addresses shape its efficiency.
Occupancy
Active warps relative to the maximum possible on an SM; useful for latency reasoning, but not a performance score.
EXAMPLE

WORKED EXAMPLE

  1. 01
    COUNT BLOCKS

    ceil(1,000 ÷ 256) = 4, so the grid needs four blocks.

  2. 02
    COUNT THREADS

    4 × 256 = 1,024 threads launch; the first 1,000 own real data.

  3. 03
    GUARD THE TAIL

    The final 24 threads own no element, so if (i < n) must stop them.

RESULT

Four blocks cover every element and leave 24 tail threads that exit safely.

WHY IT MATTERS

Rounding up provides coverage; the boundary guard provides correctness. Together they define basic thread ownership.

DIAGNOSTIC PLAYBOOK

01SYMPTOM

Intermittent wrong results or illegal address

INSPECT FIRST
Index ownership, rounded-up tail threads, array lengths, and error checks
EVIDENCE
Compute Sanitizer, the smallest failing size, and a CPU-reference diff
02SYMPTOM

GPU is slower than CPU

INSPECT FIRST
Problem size, cold start, movement, and launch share
EVIDENCE
Segmented timing and a warmed-up median—not one run
03SYMPTOM

A larger block is slower

INSPECT FIRST
Registers/thread, shared memory/block, active warps, and tail blocks
EVIDENCE
Compiler resource report, occupancy data, and a controlled block-size sweep

HARDWARE → ECOSYSTEM

  1. 01FIXED GRAPHICS PIPELINE

    Programs organized work around graphics stages.

    General computation had to masquerade as graphics.
  2. 02G80 + CUDA

    Unified shader hardware met the kernel/thread model.

    Developers could express data-parallel work directly.
  3. 03MODERN CUDA ECOSYSTEM

    Libraries, compilers, profilers, and frameworks share the execution model.

    Thread mapping became a common entry point to HPC and AI.
05 · 08 MIN

HARDWARE + ECOSYSTEM COORDINATE

2006–2007 · G80 AND CUDA

A unified shader architecture made many processing units available to more general programs; CUDA then exposed kernels, the thread hierarchy, memory, and synchronization. This was not the birth of parallel computing—it was a major shift in GPU programmability.

06 · 12 MIN

PRACTICE + REVIEW

With 1,000 elements and 256 threads per block, why does the kernel still need if (i < n)?

  1. A. To make threads execute faster
  2. B. Because 1,024 threads launch and the last 24 indices are out of range
  3. C. Because a warp contains 16 threads
REVEAL ANSWER
B

The block count rounds up to four, launching 1,024 threads. The guard lets extra threads exit safely.

02CALCULATE

n = 10,000 and threads per block = 256. How many blocks and extra threads launch?

HINT

Round 10,000 / 256 upward.

REFERENCE ANSWER

40 blocks launch 10,240 threads, so 240 extra threads exit through the boundary guard.

03DEBUG

The kernel writes c[i] with no if (i < n). When is the bug most likely to appear?

HINT

Use an n that is not divisible by block size.

REFERENCE ANSWER

The final block contains invalid indices and may read or write out of bounds. Guard every affected array access.

04MAP

For a 1920×1080 image and 16×16 blocks, what is the grid?

HINT

Round each dimension independently.

REFERENCE ANSWER

120×68. The y dimension covers 1,088 rows, so the kernel must check both x < 1920 and y < 1080.

05EXPLAIN

Why should block 1 not spin on a normal global variable until block 0 writes it?

HINT

Consider scheduling order and residency.

REFERENCE ANSWER

Blocks have no guaranteed order, so the scheme can race or deadlock. Split kernels, use supported cooperative launch, or redesign blocks to be independent.

LAB · 25–35 MIN

OPTIONAL HANDS-ON LAB

OUTSIDE THE CORE 60 MINUTES · REQUIRES A SUITABLE CUDA / GPU ENVIRONMENT

LAB GOAL

MOVE YOUR FIRST VECTOR ADD FROM CPU TO GPU

Match a CPU reference and explain launch configuration, the boundary guard, and host/device movement.

PROCEDURE

  1. 01

    Create deterministic 10,000-element input and a CPU correctness oracle.

  2. 02

    Write one-thread-per-element code with 256 threads per block and a rounded-up grid.

  3. 03

    Test n = 10,000, 10,001, and a sub-block input.

  4. 04

    Time allocation, H2D, kernel, and D2H separately from end-to-end time.

EVIDENCE OF COMPLETION

Submit the kernel, three correctness comparisons, and a table separating kernel and end-to-end time; do not treat one cold run as a conclusion.

STRETCH CHALLENGE

Map a 2D image with 16×16 blocks and explain both x and y guards.

OFFICIAL LEARNING ENTRY POINTSCUDA Programming GuideCUDA Samples
REF

GO DEEPER

01
CUDA Programming Guide · Thread HierarchyThreads, blocks, grids, and built-in indices
OFFICIAL SOURCE · OPENS IN A NEW TAB02
CUDA Programming Guide · Intro to CUDA C++Kernel launches through a vector-add example
OFFICIAL SOURCE · OPENS IN A NEW TAB