Fusion Warm-Up: Remove the Scratch Round Trip

Lesson

You already know the first expression from SAXPY: alpha * x[i] + y[i]. The starter stores that result in scratch[i], finishes the first grid-stride loop, and then starts a second loop that reads scratch[i] and multiplies it by beta.

The scratch array costs memory in two different ways. The kernel writes and reads the whole array, which adds traffic. The host pipeline also allocates N floats for it, which raises the run's peak device-memory footprint.

First, put both expressions in one loop in solution.cu. A local variable can hold the intermediate in a register. Then open pipeline.cu and remove the N-sized cudaMalloc/cudaFree pair. The trusted main still owns x, y, out, the fixed test sizes, and correctness.

After the run, use the kernel's SASS memory traffic to verify that two array sweeps disappeared. Then open Session Details > Memory and compare Peak in flight. Those are related improvements, but they are not the same measurement. The next stage applies this fusion pattern to SwiGLU inside a Transformer MLP.

This is a guided warm-up: the same-kernel two-loop scratch pattern makes traffic visible. Real framework pipelines commonly launch separate operators. Read issued sectors and allocation footprint separately.

What you will learn

• Count full-array reads and writes in a chain of elementwise operations.
• Keep a thread-private intermediate in a local variable instead of global scratch memory.
• Remove an intermediate allocation once no kernel needs it.
• Distinguish global-memory traffic from allocated device-memory footprint.
• Use SASS sectors and the Memory tab to verify both changes.

Starter code

    int start = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;

    // Pass 1: familiar SAXPY, but the intermediate is written to global memory.
    for (int i = start; i < N; i += stride) {
        scratch[i] = alpha * x[i] + y[i];
    }

    // Pass 2: every scratch value is read straight back.
    for (int i = start; i < N; i += stride) {
        out[i] = beta * scratch[i];
    }