Sample 1 - LeetGPU — reduction, from shared memory to warp shuffles

Classic warm-up problem: sum an array of N = 2^26 floats. The interesting part is not correctness but how close you can get to the memory-bandwidth roofline.

Iterations

Version	Technique	Time (ms)	Effective BW
v1	Shared memory, interleaved addressing	3.91	137 GB/s
v2	Sequential addressing (no bank conflicts)	2.44	220 GB/s
v3	Grid-stride loop, 4 elements/thread	1.32	407 GB/s
v4	Warp shuffle + float4 vectorised loads	0.86	624 GB/s

The final inner reduction

__inline__ __device__ float warpReduce(float v) {
  for (int off = warpSize / 2; off > 0; off >>= 1)
    v += __shfl_down_sync(0xffffffff, v, off);
  return v;
}

Two observations worth keeping:

1. Vectorised loads carried most of v4's gain, not the shuffles. Moving from float to float4 cut the number of memory instructions by 4× and let the loads issue back to back.
2. Beyond v3 the kernel is fully bandwidth bound — the arithmetic is free, so any further tuning has to reduce bytes moved, not instructions executed.

The same shuffle-based pattern feeds directly into the thesis work: the best/second-best scan in the auction bidding kernel is structurally the same reduction with a different operator.