SwiGLU: Apply Fusion to a Transformer MLP

Lesson

SwiGLU is an activation used in many Transformer MLPs. Two matrix multiplications first produce arrays commonly called gate and up. You do not implement those matrix multiplications here; main.cu provides both arrays.

For one element, SiLU(g) = g / (1 + exp(-g)). SwiGLU multiplies that gated value by up: out[i] = SiLU(gate[i]) * up[i]. The starter already calculates the correct result.

The performance problem is the same pattern you saw in Stage 6. The first loop writes SiLU(gate[i]) to global scratch. The second loop reads it back and multiplies by up[i]. In solution.cu, combine those statements in one grid-stride loop and keep the SiLU result in a local variable.

Once the kernel no longer uses scratch, open pipeline.cu and remove the N-sized cudaMalloc/cudaFree pair. main.cu remains read-only and still owns the inputs, output, test sizes, and correctness checks.

After the run, use SASS memory traffic to verify that two full-array sweeps disappeared. Then open Session Details > Memory and compare Peak in flight. The sector count measures traffic; the Memory tab measures the allocation footprint.

The two scratch loops in one kernel are teaching scaffolding. A framework implementation may use separate operator launches. This lab covers the activation only, after the gate/up projections.

What you will learn

• Explain gate, up, SiLU, and the elementwise SwiGLU formula.
• Transfer the scratch-removal pattern from a familiar expression to Transformer code.
• Keep the SiLU intermediate in a local variable instead of global memory.
• Remove the intermediate allocation once the fused kernel no longer needs it.
• Verify both the traffic reduction and the lower allocation peak.

Starter code

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

    // Pass 1: compute SiLU and write the intermediate to global scratch.
    for (int i = start; i < N; i += stride) {
        float g = gate[i];
        scratch[i] = g / (1.0f + expf(-g));
    }

    // Pass 2: read that intermediate back and multiply by the up branch.
    for (int i = start; i < N; i += stride) {
        out[i] = scratch[i] * up[i];
    }