Inline PTX: Vector Loads by Hand

Lesson

An asm statement puts PTX text into the kernel. Its shape is asm volatile("instruction text" : outputs : inputs : clobbers). Operands are placeholders numbered from %0 in the order they are listed, outputs first. Each operand carries a constraint letter that says what kind of register binds it: "r" a 32-bit integer, "l" a 64-bit value such as an address, "f" a float, "d" a double, "h" a 16-bit value. An output is marked with = in front of the letter. volatile tells the compiler the statement must stay where it is and must not be removed, and a "memory" clobber tells it the statement touches memory it cannot see, so it must not keep a stale value in a register across it.

The instruction you are writing is ld.global.v4.f32 {%0, %1, %2, %3}, [%4]; with four "=f" outputs and one "l" input. The braces are part of PTX syntax for a vector operand, and the semicolon is part of the instruction. The store is st.global.v4.f32 [%0], {%1, %2, %3, %4}; with five inputs and a "memory" clobber. Keep the arithmetic in C++; the point of this lab is the memory instructions.

The address needs one more step. ld.global expects a global-space address, and a C++ pointer is a generic address. The compiler's own listing converts with cvta.to.global.u64 before every global access; from C++ you ask for the same conversion with __cvta_generic_to_global(pointer), which returns the global address as a size_t that binds to an "l" operand. On current GPUs the two addresses coincide and the conversion costs nothing, but the instruction is specified in terms of the global space, so write it that way.

Two rules from the previous lab still apply. A 16-byte load needs a 16-byte aligned address, which holds for cudaMalloc buffers at an offset that is a multiple of four floats. And N is not always a multiple of four: 2, 31, 255, 257 and 4097 are among the test sizes, so the last one to three elements need a scalar tail. Set the grid expression to ceilDiv(N, blockX * 4) once each thread handles four elements.

After the run, read your kernel's PTX. nvcc copies each asm statement into its PTX between a begin inline asm comment and an end inline asm comment, with %0 and the other placeholders replaced by the registers it chose, so the PTX pane shows exactly what reaches ptxas. The listing captured from a run is the same PTX with its comments removed, so there you find the statements by their opcodes. The SASS shows LDG.E.128 and STG.E.128, as it did when the compiler generated the vector load from float4. The sector count from the SassMetrics pass does not change, because width does not change the access pattern; the instruction count does. This lab is judged on correctness, on at least one inline asm statement in your code, and on your kernel's captured PTX, which must hold ld.global.v4.f32 and st.global.v4.f32.

If the compile fails, read the ptxas line in the log. A constraint letter that does not match the instruction's operand type is reported as an arguments mismatch for the instruction; a missing semicolon or a stray brace is a syntax error at the line of your statement. An "r" where an "l" belongs silently truncates an address to 32 bits, and the kernel faults instead of failing to assemble.

What you will learn

• Write an asm statement with outputs, inputs and a clobber, and bind each operand with the right constraint letter.
• Convert a generic pointer to the global-space address a ld.global or st.global expects.
• Issue 128-bit vector loads and a vector store by hand and keep the kernel correct for every N.
• Find your own instructions in the captured PTX and the corresponding LDG.E.128 and STG.E.128 in the SASS.

Starter code

    // Scalar SAXPY: the compiler emits ld.global.f32, fma.rn.f32 and st.global.f32.
    // Rewrite it so each thread handles four elements through your own
    // ld.global.v4.f32 and st.global.v4.f32 asm statements, with a scalar tail.
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        y[i] = fmaf(alpha, x[i], y[i]);
    }