Reimplement elementwise vector addition with a large tile size so each program handles more work. For each index i in [0,N):
out[i]=x[i]+y[i]The kernel logic is identical to the standard vector-add. The differentiator is the block-size choice. With a 1024-wide block, a one-million-element add launches roughly 1000 programs; with a 4096-wide block the same job launches 244. Fewer programs amortize launch overhead, widen memory transactions when the compiler can vectorize the loads, and tend to improve achieved bandwidth on long arrays.
The trade-off is register pressure. At fp32 with 4096 elements per program, the per-program register footprint is 16KB - comfortably under the per-SM budget but worth tracking on hotter kernels. This problem isolates the launcher-level decision so the effect is measurable.
Input: x = [1, 2, 3, 4], y = [5, 6, 7, 8]
Output: [6, 8, 10, 12]
Input: x = [0.5, 1.5, 2.5], y = [1, 1, 1]
Output: [1.5, 2.5, 3.5]
Only the launcher constant changes from a standard vector add. The kernel body is byte-for-byte the same; the point is to measure the bandwidth effect of fewer, larger programs.
The tail mask matters more here than at a 1024-wide block: a single partial tile now covers up to 4095 lanes that must not touch memory.
No reduction, no shared memory. The only knob is the tile width, which the launcher fixes as a tl.constexpr on the kernel.
torch.float32 on cuda.atol=1e-3 (default elementwise tolerance, no fp32 accumulation drift here).Accepts: array
Accepts: array
Submit a correct solution to see your kernel performance.
Reimplement elementwise vector addition with a large tile size so each program handles more work. For each index i in [0,N):
out[i]=x[i]+y[i]The kernel logic is identical to the standard vector-add. The differentiator is the block-size choice. With a 1024-wide block, a one-million-element add launches roughly 1000 programs; with a 4096-wide block the same job launches 244. Fewer programs amortize launch overhead, widen memory transactions when the compiler can vectorize the loads, and tend to improve achieved bandwidth on long arrays.
The trade-off is register pressure. At fp32 with 4096 elements per program, the per-program register footprint is 16KB - comfortably under the per-SM budget but worth tracking on hotter kernels. This problem isolates the launcher-level decision so the effect is measurable.
Input: x = [1, 2, 3, 4], y = [5, 6, 7, 8]
Output: [6, 8, 10, 12]
Input: x = [0.5, 1.5, 2.5], y = [1, 1, 1]
Output: [1.5, 2.5, 3.5]
Only the launcher constant changes from a standard vector add. The kernel body is byte-for-byte the same; the point is to measure the bandwidth effect of fewer, larger programs.
The tail mask matters more here than at a 1024-wide block: a single partial tile now covers up to 4095 lanes that must not touch memory.
No reduction, no shared memory. The only knob is the tile width, which the launcher fixes as a tl.constexpr on the kernel.
torch.float32 on cuda.atol=1e-3 (default elementwise tolerance, no fp32 accumulation drift here).Accepts: array
Accepts: array
Submit a correct solution to see your kernel performance.