Listing P.1 — Vector Add
PTX target: sm_90C[i] = A[i] + B[i]; per-thread body only — addresses already computed, bounds guard omitted
Expected issue / dependency schematic
Position 1 issues ld.global.f32 %f1, [%rd1]; on the LSU pipeline. Position 2 issues ld.global.f32 %f2, [%rd2]; on the LSU pipeline. Position 3 issues add.rn.f32 %f3, %f1, %f2; on the FMA pipeline. Position 4 issues st.global.f32 [%rd3], %f3; on the LSU pipeline. Instruction 3 reads %f1 written by instruction 1 and %f2 written by instruction 2. Because issue within a warp is in-order, once instruction 3 becomes the warp's next instruction the warp is ineligible until those results have returned — a long-scoreboard stall. Instruction 4 reads %f3 written by instruction 3.
- Instruction 3 cannot issue until %f1, %f2 are ready
A fading line has no consumer in this excerpt — no register result is awaited
A dot marks where a result is required, not when it returns; two dots at the same position do not imply the loads return together.
Instruction details for ld.global.f32 %f1, [%rd1];
ld.global.f32 %f1, [%rd1];
Issue a load of A[i] to L1TEX. The LSU is not occupied for the full end-to-end latency; scattered or poorly coalesced lane addresses increase the sectors accessed and may require multiple serialized L1TEX wavefronts
Register Usage
Named PTX virtual register operands, private to each thread — one name, and each active lane has its own value. One name is not one hardware register: a 64-bit %rd occupies a pair of 32-bit GPRs, and ptxas assigns the physical SASS registers, reusing or eliminating these entirely.