Understand GPUs

One thread per element

Take the loop for (i = 0; i < N; i++) c[i] = a[i] + b[i];. Assume the input and output arrays do not overlap. The addition for c[5] then needs only a[5] and b[5]; it does not wait for c[4]. All N additions can therefore happen in any order. A CPU could split the loop across cores and vectorize each chunk (Lesson 19). In CUDA, NVIDIA's GPU programming model, a function launched on the GPU is called a kernel. This kernel describes the work for one element. Each logical thread executes the function for its assigned index:

The __global__ qualifier marks the function as a kernel. The launch shown in the comments chooses a number of blocks and a number of threads per block; these two counts appear between the triple angle brackets. The arrays are assumed to be allocated in memory accessible to the GPU. Every thread receives its position within the launch: blockIdx.x names its block, blockDim.x is the number of threads in a block, and threadIdx.x is its position within that block. All indices start at 0. Multiplying the block number by its size, then adding the position within it, gives a unique array index. Because the launch rounds up to whole blocks, the guard if (i < n) keeps surplus threads from accessing beyond the arrays.

For N = 70 and 64 threads per block, block 0 handles indices 0–63 and block 1 starts at 64. Thread 5 of block 1 computes 1 × 64 + 5 = 69 and does the addition. Thread 6 computes 70 and fails the guard. The formula assigns work; it does not say where on the chip either block will run.

In SIMT (single instruction, multiple threads), the program describes individual threads, each with its own index, private values and control flow. The hardware issues instructions for groups of threads, usually using SIMD-style execution lanes. The thread model determines which values and branches belong to each thread; the lanes perform the operations for the participating threads.

How threads are grouped and placed on the GPU

In CUDA, a block is a group of threads that can cooperate through shared storage and synchronization. All the blocks in one kernel launch form a grid. NVIDIA hardware groups the threads within each block into warps of up to 32 threads, then assigns whole blocks to streaming multiprocessors (SMs). A thread's position within its warp is its lane, numbered from 0 to 31. The diagram shows how a grid's blocks use the finite space on those processing units.

The table gives corresponding names across programming systems. Hardware limits still depend on the device: a CUDA warp has 32 lane positions, while AMD wavefront widths and compute-unit organization vary by architecture. During ordinary execution, a CUDA block runs on one SM. Several blocks can share that SM if their state fits; blocks awaiting space remain unscheduled.

Blocks must work regardless of their scheduling order. If block 0 waits for block 1 to signal it, block 1 might still be waiting for an SM slot. If all the slots are occupied by waiting blocks, the program cannot make progress: it has deadlocked. Separate kernels ordered in the same execution queue can provide a handoff across the whole grid (Lesson 24). The number of blocks that fit at once determines when this scheduling problem can occur.

Try it: map elements to blocks, warps and lanes

Follow one value through the machine

For an index within the array, thread i loads a[i] and b[i], adds them in a private register, and stores c[i]. Its neighbours in the warp run the same instructions on their own values: register r1 in lane 5 and r1 in lane 6 are different storage. A warp instruction describes the group's operation, but the hardware may carry it out over several cycles or on specialized units. Warp width specifies the grouping of threads; the instruction and hardware implementation determine how many physical ALUs execute it and how long it takes.

CUDA calls the CPU the host and the GPU the device. The CPU enqueues the launch by adding it to a queue of GPU work, and it can continue before the GPU finishes. To use the result on the CPU, the program must wait for completion and, when necessary, copy c back to host memory. Within an SM, a scheduler can issue an instruction from another resident warp while one waits for data. This can keep the execution units working when the other warp has a ready instruction. It does not remove dependencies within a thread: following one linked-list pointer to discover the next still requires sequential loads. Lesson 21 examines what limits the number of resident warps and what makes their instructions ready to issue.

Map an element to its block, thread, warp and lane, then distinguish those coordinates from the SM that runs the block and the arithmetic units that execute its instructions.

Resident, eligible, issued

Lesson 20 showed that an SM can hold several warps. A scheduler can issue an instruction from one warp while another waits, provided the first warp has ready work. Three terms describe the steps from keeping a warp on the SM to executing its next instruction:

• A resident warp belongs to a block admitted to the SM. Its execution state and allocated registers remain on the SM while it waits.
• An eligible warp can issue its next instruction now: the instruction is available, its operands are ready, the required execution pipeline can accept it, and no synchronization is holding it back.
• An issued warp is one whose instruction the scheduler selects and sends to an execution pipeline. Other eligible warps may have to wait for a later issue slot.

A warp can be resident for hundreds of cycles without being eligible, for example while its next instruction waits for a load result. Occupancy is the number of resident warps divided by the SM's maximum resident-warp count. It measures capacity in use. Even at full occupancy, a scheduler can have no instruction to issue if all its warps are waiting.

Try it: what fits on a teaching SM

Our invented SM has room for at most 8 warps, 256 threads, 4 blocks, 8,192 registers (32-bit) and 16 KiB of shared memory. Suppose a block has T threads, each using R registers, and the block reserves S KiB of shared memory. At T = 64, that block occupies two warp slots, 64 × R registers and S KiB. We can divide each SM capacity by the corresponding block cost; the smallest whole-block count decides how many blocks fit:

This is the potential occupancy the resource limits allow. The achieved occupancy measured on a real run can be lower, for example when the grid has too few blocks to fill every SM, or near the end as blocks finish and nothing replaces them.

With T = 64, R = 32 and S = 4, four blocks fit: 8 warps, 100%. Double R to 64 and the register file only holds two blocks: 50%. Keep R = 32 but double S to 8 KiB and shared memory also allows only two blocks: 50%. Both configurations have 50% occupancy. Allowing more blocks to fit would require reducing register use in the first case or shared-memory use in the second.

Try it: hide latency by picking another ready warp

Start with just one warp executing LD → ADD → STORE. If its load issues in cycle 1 and takes four cycles, the ADD cannot use the value until cycle 5. Issue slots 2–4 sit empty. Other ready warps could fill those slots without shortening that first load.

The explorer uses one issue slot per cycle: an opportunity for the scheduler to send one instruction to an execution unit. A load issued in cycle c supplies its value at the start of cycle c + L, and an ADD supplies its result the next cycle. Among eligible warps, the scheduler chooses the lowest numbered. The model omits bandwidth and instruction-fetch delays. A warp counts as done when its STORE issues, without modelling when that store completes or becomes visible.

At L = 4, four warps can issue their loads in cycles 1–4. Warp 0's ADD issues in cycle 5 and its STORE in cycle 6; the scheduler then selects the other ready warps. All 12 issue slots are used. Counting each three-instruction sequence as one unit of work, four warps issue four units in 12 cycles (1/3 unit per cycle), compared with one warp's one unit in 6 cycles (1/6 per cycle). Throughput doubles by this measure, even though warp 3 takes longer to reach its STORE than it would alone. Every load still has the same four-cycle latency.

Try L = 16: more resident warps are needed to cover the longer wait. Real GPUs have multiple issue paths, other instruction latencies, and finite memory queues and bandwidth. Independent instructions within one warp can also fill gaps. Once a memory channel is saturated, however, issuing from more warps cannot make it transfer more bytes.

Try it: follow the active lanes through a branch

Issuing a warp instruction does not mean all 32 lanes participate. When lanes choose different branch paths, the warp must execute instructions from both paths. Each instruction has an active mask, whose set bits identify the participating lanes; lanes on the other path do no work for that instruction. As in Lesson 19, a compiler may instead predicate a short branch, issuing instructions with a per-lane condition that suppresses their effects where they are not needed.

Use the 8-lane model below to see the cost; CUDA warps have 32 lanes. Each lane starts with its lane number + 1. The true path does +10, ×2 (2 instructions); the false path does −1, ×3, +5 (3 instructions). A path with no active lanes is skipped.

If every lane in a warp chooses the same side, that warp runs only one path; other warps can choose differently. The warp avoids issuing instructions with lanes disabled by the other path, although the branch and the chosen path still take time. Grouping data so neighbouring lanes make the same choice can reduce this cost. Any speedup must also cover the cost of rearranging the data and any increase in memory traffic.

For example, a block of exactly one warp (32 threads, all active) exchanges values through shared memory:

__shared__ int slot[32];
int lane = threadIdx.x % 32;
slot[lane] = my_value;                // each lane writes its own slot
__syncwarp();                         // default full mask: all 32 lanes must reach it
int left = slot[(lane + 31) % 32];    // now safe to read a neighbour's slot

Without __syncwarp(), lane 5 could read slot[4] before lane 4 writes it. The call waits for its participating lanes and orders their shared-memory accesses, so each read here follows its neighbour's write. All 32 lanes in this example are active and reach the default full-mask call.

Every non-exited lane named in the mask must call __syncwarp with the same mask. Only the named lanes in that warp participate; __syncthreads coordinates the whole block. With independent thread scheduling, the hardware still groups participating threads from a warp to issue a common instruction. Instructions from different branch paths therefore issue for separate groups of active threads.

When profiling a kernel, check how many warps are resident, how many are eligible, how often the scheduler issues an instruction, and how many lanes participate in each issued instruction.

Name the storage and who can see it

The array-add thread keeps its sum in a register and reads its inputs from a global array. Other kernels also need a place for threads in one block to exchange data. A memory space defines which threads can access its data; its physical storage determines how those accesses are served. The diagram shows both, including the caches between an SM and device memory.

CUDA local describes thread-private visibility; it does not promise on-chip storage. A cache miss can require the same device-memory access as a global-memory load. Likewise, a global load that hits in a cache can be served without reading DRAM. AMD calls its corresponding shared scratchpad Local Data Share (LDS).

Try it: count sectors for one load instruction

Imagine 32 active lanes, each loading one 4-byte word: 128 useful bytes in all. If their words are contiguous and start on a 32-byte boundary, eight words fill each 32-byte sector, so four sectors cover the load. If the addresses spread out, the memory system must fetch more sectors to deliver the same 128 useful bytes.

In the explorer, lane l requests byte address base + 4 × stride × l. Each word is aligned, sectors are aligned 32-byte units, and all requested sectors initially lie outside the cache being modelled. Count distinct sectors touched by one load instruction.

That count can guide data layout. Suppose each record is two 4-byte floats, {x, y}, and a warp loads only x from 32 records. In an aligned array of structures, neighbouring x values are 8 bytes apart: 128 useful bytes occupy 8 sectors, or 256 sector bytes. With a separate, aligned x[] array—a structure of arrays layout—the x values are contiguous and occupy 4 sectors, or 128 sector bytes. If the kernel also needs y, compare both loads before choosing a layout.

This 32-byte sector model reflects recent NVIDIA GPUs. The table counts the sectors needed to serve one instruction's requests, assuming none of those sectors is already in the cache being modelled. Neighbouring warps may reuse the fetched sectors, and caches or request merging can reduce traffic farther down the memory hierarchy. The sector count therefore does not specify how many commands reach DRAM.

Try it: shared-memory banks

Shared memory is faster than a trip to device memory, but it has limited service capacity too. It is divided into banks. In our 8-bank, 8-lane model, bank w mod 8 holds 4-byte word w; each bank serves one distinct word per round. For example, word 0 and word 8 both map to bank 0. If different lanes ask for those two words together, that bank needs two rounds: a bank conflict. If they ask for the same word, the bank can broadcast it in one round.

Use the bank diagram to compare a column of an 8×8 tile (word stride 8) with the same tile padded to a row length of 9. The padding changes the addresses of successive values in that column, and therefore the banks that serve them. To evaluate another layout, compute the addresses an instruction actually requests and map each to its bank.

Code and storage choices change traffic at different places. Spilled register values add private-memory loads and stores. Poor coalescing fetches extra sector bytes; a cache hit may prevent some of those bytes from reaching DRAM. Shared-memory staging adds instructions but can avoid repeated global loads.

In the vector-add example, each input is used for one output, so staging it in shared memory would not avoid later global loads. Each result needs two 4-byte inputs and one 4-byte output, or 12 useful bytes per element. If memory bandwidth limits this kernel, adding arithmetic units cannot increase the rate at which those bytes arrive. Matrix multiplication, in the next lesson, uses the same inputs for multiple outputs and can benefit from staging them.

For an access pattern, compute the useful bytes, the sector bytes and the bank-conflict rounds. State the cache assumptions and bank-mapping rules used in each calculation.

Why matrix multiplication rewards reuse

In C = A × B, output C[r, c] is the dot product of row r of A and column c of B: multiply matching entries, then add the products. In symbols, C[r, c] = Σ_k A[r, k] × B[k, c], where Σ means sum over k. For example, C[0, 0] and C[0, 1] both use A[0, 0], then A[0, 1], and so on. Outputs in one row share A's row; outputs in one column share B's column. A block can load each needed input entry once and share it among the threads computing those outputs.

A tile is a small region of outputs computed together. One block owns a T × T tile of C. For T successive values of k, the block loads a T × T piece of A and a T × T piece of B into shared memory, where its threads can reuse them. It then replaces those pieces for the next T values of k. Each thread keeps its own output's running sum in a private register throughout.

Try it: a complete 2 × 2 tile, step by step

A is 2 × 4 and B is 4 × 2, so C is 2 × 2. Four threads in one block own its four outputs: thread (r, c) accumulates C[r, c]. Summing the products over k is a reduction: it combines several values into one. Here each output sums four products, so the reduction length is 4. The shared tiles have width 2, and the block processes k = 0, 1 in the first round, then k = 2, 3 in the second.

In each round, every thread loads one A entry and one B entry into shared memory. Four threads load eight words in total, then each thread uses some entries its neighbours loaded. Follow the loads, the two barriers, and one private sum in the stepper.

After the first tile (k = 0, 1) the partial sums are [[7, 10], [23, 34]]. Thread (0, 0), for instance, computes 1 × 1 + 2 × 3 = 7. The second tile (k = 2, 3) adds 3 × 5 + 4 × 7 = 43 to it, giving 50. The final result is [[50, 60], [114, 140]]. C's dimensions specify four outputs; the reduction length specifies the four products summed for each output.

Thread (0, 1) loads A[0, 1] = 2 into As[0][1] and B[0, 1] = 2 into Bs[0][1] during the first load phase. Computing its own C[0, 1] also requires values loaded by neighbours: As[0][0] = 1 and Bs[1][1] = 4. Its first partial sum is 1 × 2 + 2 × 4 = 10. A thread's load assignment supplies entries for several outputs, while computing one output uses entries supplied by several threads.

The two barriers protect different phases

__syncthreads() is a block barrier: threads that have not exited must reach the same barrier call before any can proceed. Shared-memory writes made before the barrier are visible to the participating threads after it. The tiled loop uses one barrier after loading to make the entries ready for all readers, and another after computing to let all readers finish before any thread replaces an entry:

Here is the kernel for square n × n row-major matrices. Row-major storage places the entries of each row next to one another, so entry (row, col) is at array index row × n + col. Assume C does not overlap either input array. Launch ⌈n / T⌉ × ⌈n / T⌉ blocks of T × T threads, with T small enough to satisfy the device's block and resource limits:

At an edge tile, some source positions lie outside A or B. The corresponding threads put zero into shared memory instead of reading an invalid address; zero adds nothing to a sum. A thread whose own output lies outside C must still load its assigned tile entries, because another thread may need them. That is why the load-and-compute loop runs for every thread and only individual global accesses and the final store are guarded.

Every thread in this kernel executes the same two barriers on each round. If some threads that have not exited branch around __syncthreads(), the synchronization is invalid and may hang or produce incorrect results. A thread that has exited no longer participates in the barrier, but returning before its assigned loads can leave shared entries missing for other threads. Keeping edge threads in the loop preserves both the cooperative loads and the common barrier sequence. Production libraries often give each thread several outputs, use other tile shapes, and overlap loading with computing.

Which threads form a warp in a 2D block?

The tiled kernel uses 2D blocks, but warps are still formed from a linear numbering, with x varying fastest:

The compute phase has a different shared-memory pattern from the load phase. At one k, threads computing outputs in the same row all read As[y][k], so that one word can be broadcast to them. Neighbouring output columns read neighbouring Bs[k][x] words, which can be served by different banks. The bank count depends on the instruction's actual addresses; loading and computing with the same tile can produce different counts.

What a tile costs in resources

With one thread per output, a T × T tile uses T² threads per block. Its two T × T float input tiles take 2T² × 4 = 8T² bytes of shared memory before any double buffering. At T = 16, that is 256 threads and 2 KiB; at T = 32, 1,024 threads and 8 KiB.

In this mapping, the thread count is often the tighter constraint: 1,024 threads is a common per-block maximum. Production kernels commonly give each thread several outputs instead. Check the limits of the actual device. Lesson 24's “tile 64” is a traffic model, not a proposed 64 × 64-thread block.

Count the reuse

For one T × T output tile, let the reduction length K be a multiple of T. Each of the K/T rounds loads two T × T input tiles: (K/T) × 2T² = 2TK input words. At the end, the block writes T² outputs. Each output performs K multiplications and K additions. Counting each as one floating-point operation, or FLOP, gives 2T²K FLOPs. Dividing this work by the requested bytes gives the arithmetic intensity for these requests. With 4-byte words:

For our example (T = 2, K = 4), the tiled version requests 16 input words and writes 4 outputs: 80 bytes for 32 FLOPs. Fetching both operands separately for every multiply-add would request 32 input words and write the same 4 outputs: 144 bytes for 32 FLOPs. A wider tile can raise arithmetic work per requested byte by reusing each loaded value more often.

These byte counts come from the program's load and store requests. Actual DRAM traffic can differ: a cache might serve repeated loads in the untiled version, while tiling adds instructions and barriers. The ratio 144 / 80 therefore does not predict a 1.8× speedup.

A double-buffered version keeps two sets of A and B input tiles, loading the next set while computing on the current one. Threads must wait for each buffer's loads to complete before reading it, and for its readers to finish before overwriting it. An asynchronous copy starts a transfer that may finish later; its completion must be established before using the data.

To estimate how these saved requests affect running time, compare arithmetic work with bytes moved across a particular memory boundary, such as DRAM. Lesson 24 makes that comparison and explains which quantities must be measured.

For every shared tile entry and every private sum, say who writes it, who reads it, and how long it must live. Explain each barrier with a concrete producer/consumer race.

From diagrams to a performance hypothesis

Lesson 23 cut the number of requested bytes, but a smaller byte count alone does not tell us how much faster the kernel can run. We need two rates: how quickly the GPU can perform the arithmetic, and how quickly the chosen memory boundary can supply the data.

Call the kernel's floating-point work W, counted in FLOPs as in Lesson 23. Call the bytes it moves across a named memory boundary, such as between the GPU's caches and DRAM, Q. Its arithmetic intensity is I = W / Q FLOPs per byte. Let P be the compute ceiling for this kind of arithmetic and β the bandwidth at that boundary. The work needs at least W / P seconds and the data movement at least Q / β seconds. Even if they overlap perfectly, the kernel cannot finish before the slower requirement:

The resulting limit is called the Roofline model because its performance graph has a sloping bandwidth ceiling and a flat compute ceiling. It gives a best-case bound on execution time. Dependencies, issue limits, synchronization or too little parallel work can make the actual run slower. Use ceilings that match the instructions and precision: a matrix-unit peak is not the compute ceiling for scalar addition. At intensities below the ridge point, memory movement sets the lower performance ceiling; above it, computation does. These are the model's memory-bound and compute-bound regions. A measured point below the sloped roof does not, by itself, prove that DRAM bandwidth is saturated.

Try it: a Roofline you can check by hand

Use a made-up device with P = 10 TFLOP/s for our arithmetic and β = 200 GB/s. These are decimal units: 1 TFLOP/s is 10¹² floating-point operations per second, 1 GFLOP/s is 10⁹, and 1 GB is 10⁹ bytes. Start with vector addition. In this explorer, every counted byte is assumed to cross the DRAM boundary exactly once, so requested bytes stand in for DRAM bytes; real caches can change that. The matrix presets use 1024 × 1024 inputs, 2n³ FLOPs and requested bytes 8n³/T + 4n², derived below. “Tile 64” describes the reuse used to count traffic, not a 64 × 64-thread launch.

Here is the matrix byte formula. If T divides n, there are (n/T)² output tiles. For one output tile, the reduction covers n positions and requests 2Tn input words: T A words and T B words at each position. Multiply by the number of output tiles to get 2n³/T input words, or 8n³/T bytes for 4-byte floats. Writing n² outputs adds 4n² bytes. This counts requested bytes with no reuse between output tiles; a cache may keep some of them from crossing DRAM.

For vector addition, one million float additions perform W = 10⁶ FLOPs and request Q = 12 × 10⁶ bytes, so I = 1/12 FLOP per byte. At the stated ceilings, computation takes at least 0.1 μs, while moving the bytes takes at least 60 μs. The longer memory requirement sets the time bound, giving a performance ceiling of about 16.7 GFLOP/s. Doubling P cannot shorten that 60 μs bound; doubling β could halve it in this model.

The ridge point, where the two ceilings meet, lies at P / β = 10 × 10¹² / (200 × 10⁹) = 50 FLOPs per byte. The vector add's 1/12 is far to its left. Reusing loaded values, as in tiled matrix multiplication, raises intensity by reducing Q for the same work W. Adding arithmetic merely to raise the FLOP count increases W and does not make the original task faster.

Add H = 10 μs of fixed overhead that does not overlap with the kernel. The modelled vector add takes 10 + 60 = 70 μs, an effective 10⁶ FLOPs ÷ 70 μs ≈ 14.3 GFLOP/s including H. A hypothetical version that halves its traffic would take 10 + 30 = 40 μs, a 1.75× speedup rather than 2× because the same 10 μs remains. These effective rates and the kernel-only ceiling describe different intervals. A measurement must specify which interval it covers.

Decide what "time" means

A kernel launch is normally asynchronous with respect to the CPU: the call submits work to a GPU stream and can return before that work finishes. Submission itself takes time, and the CUDA runtime may stall the host when resources are unavailable. Timing just the launch call therefore measures host-side submission, not how long the GPU takes to execute the kernel. Choose the interval that answers your performance question:

cudaEvent_t start, stop;
cudaEventCreate(&start);  cudaEventCreate(&stop);   // timing enabled by default
cudaEventRecord(start, stream);
add_arrays<<<blocks, threads, 0, stream>>>(a, b, c, n);
cudaEventRecord(stop, stream);
cudaEventSynchronize(stop);                        // wait until the GPU passes 'stop'
float ms = 0;
cudaEventElapsedTime(&ms, start, stop);            // device time between the events
// check the return codes and cudaGetLastError() in real code

For device time, put the start event, kernel and stop event in the same stream, in that order. An event marks the GPU's progress through that stream; waiting for the stop event establishes that the preceding work has completed. Read the elapsed time only after that wait. Other work on the GPU can still stretch the interval, so record what else was running. For a pipeline that overlaps copies and computation, measure steady state after it has filled and before it drains. Divide the completed work by elapsed time to report throughput, or invert that ratio to report time per item.

Repeating the kernel R times and dividing the total can reduce timer-resolution error. It can also change the workload: later runs on the same data may hit in cache if the working set fits, and a slowly enqueueing CPU may leave gaps between launches. Report whether the data was warm or cold and inspect the timeline when the numbers surprise you.

Order dependent phases with streams and events

An algorithm may need every block to finish producing intermediate results before any block starts consuming them. Within an ordinary kernel, a block cannot assume that the other blocks are running; spinning until an unscheduled block signals it can deadlock. A kernel boundary supplies the ordering. With the ordinary launches shown below, work queued in one stream runs in order. A consumer kernel launched after a producer in that stream starts after the producer completes and sees its global-memory results. For the cross-stream example, assume s1 and s2 are independent streams:

produce<<<grid, block, 0, s>>>(buf);   // phase 1 writes buf
consume<<<grid, block, 0, s>>>(buf);   // same stream: starts after produce completes

// Across independent streams, add an explicit dependency:
produce<<<grid, block, 0, s1>>>(buf);
cudaEventRecord(done, s1);
cudaStreamWaitEvent(s2, done, 0);          // s2 waits for the event, the host doesn't
consume<<<grid, block, 0, s2>>>(buf);

In independent streams, launching the consumer without the event supplies no producer-to-consumer dependency: the consumer may start before the producer finishes. Recording the event after produce and waiting for it in s2 creates that dependency. The CPU can continue submitting work while s2 waits. CUDA's legacy default stream can impose implicit ordering with other streams; the explicit event makes the required dependency clear without relying on those rules.

Beyond ordinary blocks. CUDA also supports constrained forms of cooperation across blocks. A cooperative launch, on a supported device, limits the grid to the number of blocks that can reside concurrently for that kernel and launch configuration, enabling grid-wide synchronization. Thread-block clusters, on supporting hardware, group blocks that are scheduled together and can access each other's shared memory (distributed shared memory) with cluster-level synchronization. Both features have launch rules and size limits; neither guarantees simultaneous execution for arbitrary blocks. This unit assumes a discrete GPU with its own memory, reached by copies; Unit 6 covers integrated and coherent arrangements.

Measure a change against a stated hypothesis

1. Check correctness first, against a reference, including sizes that aren't multiples of the block size.
2. Record the setup: GPU model, compiler, launch configuration, precision, input shape, whether inputs start on the device.
3. Warm up to complete one-time initialization and bring the device to its intended operating state. Collect repeated samples and report a representative value and its spread. Keep cold-cache and warm-cache results apart; warm-up can also leave data in caches.
4. Profile to test a specific hypothesis. A profiler collects counters and timelines that describe execution. If an access pattern appears to waste sectors, inspect requests and bytes at the relevant cache level. If operands appear to hold up issue, inspect eligible warps and stall reasons. For low residency, inspect register and shared-memory allocation. For launch overhead, inspect the CPU/GPU timeline. Profiling may replay kernels, alter cache state and perturb execution, so collect baseline timings separately.
5. Change one thing, such as layout, tile shape, block size or reuse, and predict which metric should move before you measure. A speedup with an unexpected metric change is still evidence, but of a different explanation.

Before explaining a speedup, state the work, the byte boundary, the applicable ceiling, the measurement interval and the uncertainty.

Launch N = 70 float additions in two blocks of 64 threads, with every array starting on a 32-byte boundary. Predict the thread indices and active lanes, then count sectors for each load and store instruction.

To analyze a kernel, identify its independent work and the resources needed to keep it resident. Then trace which instructions are ready, which lanes participate, which addresses are accessed together, and where values are reused. Choose a measurement interval that tests the resulting explanation. The CS149 Lecture 2 companion develops these ideas through programming examples. Unit 6 adds matrix hardware and number formats, then follows data through whole systems of chips.

Based on the Unit 5 outline and worked answers in chatgpt-reference/. The diagrams and explorers use the teaching assumptions listed below.

Model assumptions and further reading

The teaching SM's capacities (4 blocks, 256 threads, 8 warps, 8,192 registers, 16 KiB), the 8-lane groups in the divergence explorer, the 8 shared-memory banks, and the 10 TFLOP/s, 200 GB/s device are invented round numbers chosen to make the arithmetic visible. Real limits differ by GPU generation; look them up for your device. Primary sources:

• NVIDIA CUDA Programming Guide: SIMT execution and independent thread scheduling, __syncthreads() and other built-ins, and streams and events.
• NVIDIA CUDA C++ Best Practices Guide: coalesced access, shared memory and occupancy.
• NVIDIA Nsight Compute Profiling Guide: the Roofline charts and the memory boundaries a profiler actually measures.
• AMD GPU architecture documentation (ROCm): compute units and wavefronts.