FMA Contraction: Read the Arithmetic

Lesson

Run the starter and find its arithmetic in the PTX. alpha arrives with ld.param.f32 from saxpy_param_2, x[i] and y[i] with ld.global.f32, and between them sit mul.rn.f32 and add.rn.f32. The suffix is a rounding modifier: PTX arithmetic can be written with .rn (round to nearest even), .rz (toward zero), .rm (toward minus infinity) or .rp (toward plus infinity). An instruction that carries a modifier must deliver exactly that rounded result, so ptxas may not fuse it with its neighbours. A mul.f32 or add.f32 written with no modifier also rounds to nearest, but it gives the assembler permission to fuse. That difference is the whole lesson.

Fused multiply-add computes a * b + c with the product kept exact and a single rounding at the end. nvcc contracts a multiply followed by an add into an fma by default (the switch is -fmad, and it is on), and the PTX for alpha * x[i] + y[i] is one fma.rn.f32. The fused result can differ from the two-step result by one unit in the last place, and sometimes that matters: a difference of two nearly equal products, or code that must match a CPU bit for bit. That is what __fmul_rn and __fadd_rn are for. They are not faster multiplies; they are contraction fences, and their PTX is the mul.rn.f32 and add.rn.f32 you just read.

For this lab, remove the fence. Write the update as fmaf(alpha, x[i], y[i]), which makes the fusion explicit, or as alpha * x[i] + y[i], which the compiler contracts on its own. Run again and compare the two listings: one fma.rn.f32 where there were two instructions, and FFMA in the SASS where there were FMUL and FADD.

Do not expect the clock to move. SAXPY reads eight bytes and writes four per element, and whether that costs one arithmetic instruction or two makes no difference to a kernel that is waiting on memory. The Trace pass's GPU time for both versions is the same within noise, and the SassMetrics sector count is identical. The harness computes its reference with std::fmaf; both versions pass its tolerance because they differ by at most one unit in the last place. This lab is judged on correctness and on the arithmetic in your kernel's captured PTX, one fma.rn.f32 and no mul.rn.f32 or add.rn.f32; the thing to take away is that you can now read the arithmetic in any listing and say which results were rounded, how many times, and why.

What you will learn

• Read rounding modifiers on PTX arithmetic and explain what .rn commits the assembler to.
• Recognize fma.rn.f32 as a contracted multiply-add with one rounding, and mul.rn.f32 plus add.rn.f32 as two.
• Explain what __fmul_rn and __fadd_rn do to the generated code and when a kernel needs them.
• Predict when an instruction change will and will not change GPU time.

Starter code

    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        // Two separately rounded operations. The _rn intrinsics forbid fusing them.
        float product = __fmul_rn(alpha, x[i]);
        y[i] = __fadd_rn(product, y[i]);
    }