← BACK TO LESSONPARALLEL HORIZONSREAD // 01

LESSON READING // 01

MEMORY HIERARCHY

HISTORICAL ANCHOR2010 · FERMI MEMORY TURNING POINTREAD TIME · ABOUT 60 MIN

Once enough threads exist, stop adding more. Shape each warp's memory requests and stage reused data in the right memory space.

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 // 01
FOUNDATIONALPROFILINGNSIGHT COMPUTE
THE OFFICIAL PATH EMPHASIZES

Move beyond running code: use industry tools to locate memory bottlenecks and let evidence drive optimization.

HOW THIS LESSON CONNECTS

The lesson visualizes coalescing, reuse, and shared memory; the lab requires profiler evidence for transactions, throughput, and bank behavior.

SUGGESTED PREREQUISITES
  • Lesson 00 or a basic kernel
  • Warp and global-index concepts
  • Reference-output comparison
PRACTICE TOOLS
  • Nsight Compute
  • CUDA C++
  • CUDA C++ Best Practices Guide
01 · 05 MIN

ORIENT + OBJECTIVES

REMEMBER THIS FIRST

GPU performance depends not only on how much you compute, but on the shape in which data reaches the compute units.

BY THE END, YOU SHOULD BE ABLE TO

  1. 01

    Inspect a warp's address set and predict excess global-memory transactions.

  2. 02

    Choose registers, shared memory, or global memory from reuse scope.

  3. 03

    Explain halos, cooperative loading, and synchronization in neighborhood work.

  4. 04

    Recognize bank conflicts, unsafe barriers, and over-caching.

02 · 08 MIN

BUILD THE MENTAL MODEL

  1. 01
    ADDRESS

    A warp issues 32 addresses together; hardware coalesces them into as few memory transactions as possible.

  2. 02
    REUSE

    When a block reuses neighboring values, its threads can cooperatively stage them in shared memory.

  3. 03
    SYNC

    Readers can safely consume the complete tile only after all participating threads finish loading it.

03 · 15 MIN

CONCEPT DEEP DIVES

01

COALESCED ACCESS

Adjacent threads reading adjacent, suitably aligned data commonly reduces the global transactions needed to serve a warp.

02

SCOPE CHOOSES THE SPACE

Registers belong to a thread, shared memory to a block, and global memory is visible across the grid and host.

03

SHARED MEMORY IS MANAGED

It is fast but finite, can suffer bank conflicts, and requires explicit loading, layout, and synchronization.

PART 01

COUNT TRANSACTIONS, NOT JUST ADDRESSES

When a warp executes a global load, it presents a set of addresses together. Hardware services the required memory segments, so the useful question is not merely '32 values were read' but 'how many bytes moved to deliver them?' Contiguous addresses commonly improve useful-byte efficiency.

Exact coalescing rules vary with compute capability, data width, and cache path. The durable method is to inspect address distribution, alignment, and useful bytes, then confirm with a profiler such as Nsight Compute.

PAUSE AND REASONWhich usually wastes more transfer: contiguous floats or floats 4 KB apart?

REFERENCE ANSWERThe 4 KB stride usually touches many more segments and transfers more unused data; contiguous addresses coalesce more readily.

PART 02

TILES TURN REPEATED MOVEMENT INTO COOPERATION

Neighborhood, convolution, and matrix algorithms often make adjacent threads reread the same input. A shared-memory tile lets a block load a region cooperatively, then reuse it on chip. Halos provide the extra edge values needed by boundary threads in the tile.

Tile size trades reuse against shared-memory capacity, block size, and halo overhead. Tiny tiles spend more on halos; large tiles can reduce resident blocks. Compare saved global traffic with added barriers, indexing, and resource cost.

PAUSE AND REASONHow can 66 input values serve 64 outputs in a three-point stencil?

REFERENCE ANSWERThe 64 central values plus two halos cover every neighbor. Central values are reused rather than loaded three times per thread.

PART 03

THE WHOLE PARTICIPATING GROUP MUST REACH A BARRIER

__syncthreads() is a block barrier. If only some threads enter a branch containing the barrier, others may never arrive. A common safe pattern conditionally performs load duties but reaches the barrier unconditionally.

Shared memory is divided into banks. Threads mapping different addresses to the same bank can serialize; padding a 2D tile is a common remedy. A barrier ensures visibility but does not repair bank conflicts or bad indexing.

PAUSE AND REASONWhy is if (threadIdx.x < 16) __syncthreads(); unsafe?

REFERENCE ANSWEROnly part of the block reaches the barrier, producing undefined behavior and possible deadlock.

04 · 12 MIN

PUT IT BACK INTO CODE

PROGRAM MODELA block cooperatively loads and reuses a tile
01__shared__ float tile[BLOCK_SIZE + 2];02int i = blockIdx.x * blockDim.x + threadIdx.x;03tile[threadIdx.x + 1] = input[i];04// boundary threads also load left/right halos05__syncthreads();06output[i] = tile[t] + tile[t + 1] + tile[t + 2];

TRACE THE CODE

  1. 01
    ALLOCRESERVE CENTRAL VALUES + HALOS

    BLOCK_SIZE + 2 leaves one location on each side; 2D kernels commonly add borders in both dimensions.

  2. 02
    LOADCOALESCE THE CENTRAL LOAD

    Adjacent threads map to adjacent central values; boundary threads take a small number of halo duties.

  3. 03
    BARRIERFINISH COOPERATION BEFORE USE

    Every thread reaches __syncthreads() after completing its assigned loads.

  4. 04
    REUSECOMBINE OUTPUTS ON CHIP

    Each global input can serve several neighboring outputs, with global boundaries still guarded.

REF · 01

KNOWLEDGE ATLAS

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

CORE VOCABULARY

Memory Transaction
A memory segment moved to satisfy a warp request; requested bytes and transferred bytes are not identical.
Coalescing
Arranging warp addresses into as few aligned segments as practical to raise useful-byte efficiency.
Shared Memory
Program-managed on-chip storage shared within a block for cooperative loading and reuse.
Halo
Extra input around a tile that lets edge threads compute neighborhood outputs.
Bank Conflict
Threads request different addresses in one bank, forcing a shared-memory request to split or serialize.
EXAMPLE

WORKED EXAMPLE

  1. 01
    DIRECT LOADS

    64 × 3 = 192 scalar requests, with neighboring threads asking for many of the same values.

  2. 02
    COOPERATIVE TILE

    Load 64 body values plus one halo value on each side: 66 values total.

  3. 03
    ESTIMATE REUSE

    192 ÷ 66 ≈ 2.9, so the source-level request count falls to roughly one third.

RESULT

The ideal count falls from 192 requests to 66 cooperative loads, before synchronization and boundary cost.

WHY IT MATTERS

This estimates reuse, not transactions or speed. Profile actual traffic, cache behavior, and barrier stalls.

DIAGNOSTIC PLAYBOOK

01SYMPTOM

Low bandwidth with many load requests

INSPECT FIRST
Warp continuity, alignment, and whether adjacent threads touch adjacent data
EVIDENCE
Nsight Compute transactions, throughput, and warp-address shape
02SYMPTOM

Shared memory made it slower

INSPECT FIRST
Reuse, barriers, halo ratio, capacity, and occupancy
EVIDENCE
Before/after global bytes, barrier stalls, and active warps
03SYMPTOM

Unexpected latency in a 2D tile

INSPECT FIRST
Whether row stride maps many lanes to one bank
EVIDENCE
Bank-conflict metrics and a controlled +1-padding comparison

HARDWARE → ECOSYSTEM

  1. 01GLOBAL-MEMORY DISCIPLINE

    Early GPGPU demanded careful address organization.

    Coalescing became a foundational optimization.
  2. 02CACHE + SHARED MEMORY

    Cache hierarchy, ECC, and on-chip reuse grew stronger.

    Reasoning expanded to reuse, reliability, and resources.
  3. 03ASYNCHRONOUS MOVEMENT

    Newer architectures strengthened asynchronous tile movement.

    Software pipelines load the next tile while computing the current one.
05 · 08 MIN

HARDWARE + ECOSYSTEM COORDINATE

2010 · FERMI

Fermi strengthened caches, ECC, and the memory system required for general-purpose computing. Data-center GPU programming increasingly became a question of data movement and reliability, not just arithmetic-unit count.

06 · 12 MIN

PRACTICE + REVIEW

A warp reads input[i * 16]. What should you inspect first?

  1. A. Whether addresses spread across many memory transactions
  2. B. Whether the block count always equals 16
  3. C. Whether to add more CPU threads
REVEAL ANSWER
A

The stride spreads addresses within the warp. Inspect the request shape before reordering data or introducing shared memory.

02REASON

A warp reads 32 consecutive floats starting one float off alignment. Is it always as efficient as aligned access?

HINT

Consider whether an extra segment is touched.

REFERENCE ANSWER

Not always. Misalignment can add a transaction, although cache reuse may soften the cost. Check alignment and measure.

03CALCULATE

64 threads each read left, center, and right. How many scalar reads versus tile values?

HINT

Ignore caches for the conceptual count.

REFERENCE ANSWER

192 scalar reads versus 66 tile inputs: 64 centers and two halos. Hardware transaction count still depends on access shape.

04DEBUG

Only halo-loading threads execute __syncthreads(). How do you fix it?

HINT

Separate load duty from barrier participation.

REFERENCE ANSWER

Keep conditional halo loads in branches, then move __syncthreads() after the branches so the full block reaches it.

05DESIGN

Should a value used once by one thread go into shared memory?

HINT

Look at scope and reuse.

REFERENCE ANSWER

Usually no. A short-lived thread-private value belongs in a register; shared memory is for block-wide sharing or meaningful reuse.

LAB · 35–45 MIN

OPTIONAL HANDS-ON LAB

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

LAB GOAL

MEASURE CONTIGUOUS, STRIDED, AND SHARED-MEMORY PATHS

Profile three correct variants and connect address shape to global-memory efficiency and shared-memory behavior.

PROCEDURE

  1. 01

    Implement input[i] and input[(i * stride) % n] with fixed size and warmup.

  2. 02

    Collect global load/store, memory-throughput, and warp metrics.

  3. 03

    Add a three-point shared-memory tile with halos and a safe barrier.

  4. 04

    Compare direct and tiled code; explain unchanged results rather than forcing a speedup.

EVIDENCE OF COMPLETION

Provide correctness, a metric table, and one causal explanation linking address/reuse behavior to profiler evidence.

STRETCH CHALLENGE

Compare a 2D transpose tile with and without +1 padding for bank conflicts.

OFFICIAL LEARNING ENTRY POINTSCUDA C++ Best Practices GuideNsight Compute Documentation
REF

GO DEEPER

01
CUDA C++ Best Practices · Device MemoryCoalescing, shared memory, and bank conflicts
OFFICIAL SOURCE · OPENS IN A NEW TAB02
CUDA Programming Guide · Memory PerformanceGlobal transactions from a warp's requests
OFFICIAL SOURCE · OPENS IN A NEW TAB