SAXPY: Make Adjacent Lanes Access Adjacent Data
The kernel computes SAXPY correctly, but its thread-to-element mapping makes neighboring lanes touch elements far apart in memory. Keep the same arithmetic and fix only the indexing so each warp accesses contiguous values. GPUFlight's SASS coalescing efficiency is the judge, making this a direct lesson in how a source-level index expression becomes memory transactions.
Sign in to run this lab. GPU availability and credits are shown in the app.
Languages: CUDA C++, Python, Rust
Lesson
A warp executes 32 lanes together. Global memory is most efficient when those lanes request neighboring addresses, because the requests collapse into the minimum number of memory sectors. The starter deliberately transposes the usual thread mapping: lane number is multiplied by the grid size, so adjacent lanes are separated by thousands of floats. The result is correct, but each warp fragments a simple SAXPY load and store into many transactions. Use the standard global thread index so lane i handles the element next to lane i-1. The arithmetic does not change; only the mapping changes. GPUFlight's coalescing-efficiency signal compares the ideal sector count with the sectors actually issued, so it exposes the problem even when caches make wall-clock timing noisy.
The small diagram illustrates the mapping of adjacent lanes. Count sectors across a full 32-lane warp for the actual metric; do not scale the pictured lane count into measured totals.
What you will learn
- Relate a warp's lane-to-address mapping to global-memory transaction count.
- Recognize that correct output does not imply efficient memory access.
- Use the standard global thread index to make adjacent lanes process adjacent elements.
- Read SASS coalescing efficiency as a direct signal instead of relying only on runtime.
Starter code
// Correct but badly strided: adjacent lanes are gridDim.x elements apart.
int i = threadIdx.x * gridDim.x + blockIdx.x;
if (i < N) {
y[i] = alpha * x[i] + y[i];
}