Repository navigation
bend_dev on CUDA declares vote[3] with 0 dynamic shared memory, so SMs keep their L1 cache - #1602
Merged
Merged
Conversation
…s keep their L1 cache Today on main (since 5d6004a), gpu_kernel passes TG_HOLD * 8 (18,432 B) on both Metal and CUDA, while bend_dev on CUDA declares extern __shared__ u32 vote[] and only uses vote[0..2] (12 B). Reserving 18 KB of Apple's dedicated 32 KB threadgroup memory caps residency at one group per core on Metal without touching L1. On NVIDIA GPUs, shared memory and the L1 data cache share one 128 KB physical pool per SM, and default shared memory is 48 KB per block, so passing TG_HOLD * 8 to cuLaunchKernel neither caps residency at one block per SM (floor(48 / 18) = 2) nor leaves the L1 data cache intact: each SM carves 32 KB or 64 KB out of its 128 KB L1 cache for unused dynamic shared memory. This commit changes CUDA's bend_dev (#else) to declare __shared__ u32 vote[3] and changes CUDA's gpu_kernel (#elif BEND_CUDA) to pass 0 for sharedMemBytes in cuLaunchKernel. Under ncu on an RTX 3090 (CUDA 13.1), editdist's L1 hit rate rises from 6.32% to 84.14% and off-chip DRAM traffic drops from 315.55 GB to 128.65 GB, speeding up editdist 1.45x (1,101 -> 760 ms), terrain 1.35x (733 -> 544 ms), tree-matmul 1.20x (943 -> 787 ms), symreg 1.18x (528 -> 446 ms), queens 1.13x (1,258 -> 1,115 ms), bfs 1.11x (1,001 -> 906 ms), and the 21-workload bend-bench fast-gpu suite (8c09ff37) by 1.071x geomean with no regressions. TG_HOLD and Metal's setThreadgroupMemoryLength stay untouched. comp.ts 63495 -> 63493 ttok (+2 -2 lines) on d2aeefc (64092 -> 64090 on 234cb9b).
Lorenzobattistela
approved these changes
Oct 19, 2026
This was referenced Oct 19, 2026
Merged
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Today on main (since 5d6004a), gpu_kernel passes TG_HOLD * 8 (18,432 bytes) as the per-group shared-memory reservation on both Metal (setThreadgroupMemoryLength:TG_HOLD * 8 atIndex:0) and CUDA (cuLaunchKernel sharedMemBytes), while bend_dev on CUDA declares extern shared u32 vote[] and only uses three u32 words in vote (vote[0] as the round-robin cursor in monk_step, task_deal, and ring_pick; vote[1] for grew; vote[2] for has; 12 bytes total). On Apple GPUs, threadgroup memory is a dedicated 32 KB block separate from the L1 data cache, so reserving 18 KB on Metal caps residency at one group per core without touching L1. On NVIDIA GPUs, shared memory and the L1 data cache share one 128 KB physical pool per streaming multiprocessor (SM), and default shared memory is 48 KB per block, so passing TG_HOLD * 8 to cuLaunchKernel neither caps residency at one block per SM (floor(48 KB / 18 KB) = 2; and under l2 >> 16, CUBE_G is already at or below the SM count on RTX 3090 at 64 <= 82, RTX 4090 at 128 <= 128, and H100 at 128 <= 132) nor leaves the L1 data cache intact: each SM carves 32 KB (at 1 block/SM) or 64 KB (at 2 blocks/SM) out of its 128 KB L1 cache for unused dynamic shared memory.
NVIDIA Nsight Compute (ncu) on editdist (depth 12, RTX 3090, CUDA 13.1) measures main's 18,432-byte dynamic reservation (d2aeefc) against declaring vote[3] statically with 0 bytes of dynamic shared memory (higher L1 hit rate and lower off-chip memory traffic and kernel time are better):
This PR changes CUDA's bend_dev (#else) from extern shared u32 vote[] to shared u32 vote[3] and changes CUDA's gpu_kernel (#elif BEND_CUDA) to pass 0 instead of TG_HOLD * 8 for sharedMemBytes in cuLaunchKernel. TG_HOLD and Metal's setThreadgroupMemoryLength:TG_HOLD * 8 atIndex:0 stay untouched, as do host C and JS paths. comp.ts goes down by 2 ttok (+2 -2 lines): 63,495 -> 63,493 on d2aeefc, or 64,092 -> 64,090 on current main (234cb9b, where #1588 + #1591 pushed main over the 64,000 cap, so repo-gate is red from main until main drops below 64,000).
Measured with bend-bench
gpu fast(8c09ff37) on an RTX 3090 (CUDA 13.1, 4 GB heap, 3 measured runs per workload after a correctness check and warmup), d2aeefc against this PR; lower ms and speedup above 1x mean this PR is faster:Of the remaining 12 completed workloads, 8 are 1% to 3% faster (tree-radix 1.028x, bfs-shared 1.027x, gameoflife 1.026x, hashmap 1.019x, merkle 1.019x, lexer 1.017x, mandelbrot 1.013x, summation 1.011x) and 4 are within run-to-run noise (nbody 1.008x, pricing 0.998x, kmeans 0.997x, and Rodinia HotSpot: 6,162 -> 6,580 ms in the single-sweep run, then 6,351 -> 6,228 ms, -1.9%, p = 0.62 across 6 interleaved AB/BA pairs, dominated by ~2.2 s of host ASCII file I/O, which I plan to work on in the next version of bend-bench). Compact UTS hits the existing ERR_DEEP overflow on both.