fix(hip): query the device LDS limit instead of assuming 64 KB - #284
Open
demandal25 wants to merge 2 commits into
Open
fix(hip): query the device LDS limit instead of assuming 64 KB#284demandal25 wants to merge 2 commits into
demandal25 wants to merge 2 commits into
Conversation
Single-decode rejected any configuration needing more than 65536 bytes of shared memory, with the message "exceeds CDNA3 limit of 64KB". That number is correct for CDNA3 and wrong for CDNA4: measured on an MI350X, hipDeviceProp_t reports sharedMemPerBlock = 163840 (160 KB) maxSharedMemoryPerMultiProcessor = 163840 (160 KB) so the hard-coded guard rejects decode configurations gfx950 can actually run, at 2.5x the assumed ceiling. Query getMaxSharedMemPerBlock(dev_id) instead. The helper already exists in gpu_iface/gpu_runtime_compat.hpp and is already exercised by three call sites in prefill.cuh, so this adds no new machinery -- decode simply was not using it. The error message now quotes the queried limit rather than a fixed "64KB", so a future failure names the real ceiling of the device it happened on. gpuGetDevice had to be hoisted: the existing call sits inside the partition-kv `else` branch, which is reached only when seq_len > 256 && tmp != nullptr, and the check runs before that. It is now done once above the check and reused below, so the query costs nothing extra. CDNA3 impact: none. gfx942 reports sharedMemPerBlock = 65536, so the comparison is byte-identical to the previous literal and no configuration changes status. Only the message wording differs. To be confirmed on CDNA3 hardware; the targeted check is scripted. Verified on MI350X (gfx950), ROCm 7.2.0, torch 2.9.1: the full batch-decode suite (1872 tests, which exercises the single-decode dispatch this guard sits in) passes with no failures. Co-Authored-By: Claude <noreply@anthropic.com>
There was a problem hiding this comment.
Pull request overview
Updates the ROCm single-decode shared-memory guard to use the device-reported per-block shared memory limit (LDS) instead of a hard-coded 64 KB assumption, enabling valid gfx950 (CDNA4) configurations that require >64 KB dynamic shared memory.
Changes:
- Replace the
smem_size > 65536Ucheck withsmem_size > getMaxSharedMemPerBlock(dev_id). - Improve the error message to report the queried device limit in bytes.
- Hoist
gpuGetDevice(&dev_id)so the device ID can be reused later in the partition-KV path.
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
getMaxSharedMemPerBlock() copies out a full device-properties struct on every call, and the previous commit put it on the single-decode dispatch path -- including the seq_len <= 256 fast path, which previously issued no device query at all. Memoize per device id, matching getMultiProcessorCount() directly above. Measured on MI350X / ROCm 7.2: the query costs ~6.4 us on first call and ~0.17 us thereafter, against ~1.0 us for the hipLaunchKernel it precedes; cached it drops to ~0.002 us. The three prefill.cuh call sites benefit too. thread_local, so concurrent callers never race. 0 means "not cached" -- a valid limit is always > 0, so a device reporting 0 is simply not memoized rather than poisoning the cache.
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.
Summary
Single-decode rejected any configuration needing more than 65536 bytes of shared memory, with the message "exceeds CDNA3 limit of 64KB". That number is correct for CDNA3 and wrong for CDNA4 — measured on an MI350X:
So the hard-coded guard rejects decode configurations gfx950 can actually run, at 2.5x the assumed ceiling.
What changed
include/flashinfer/attention/generic/decode.cuh— the guard compares againstgetMaxSharedMemPerBlock(dev_id)instead of the literal65536U, and the error message quotes the queried limit rather than a fixed "64KB", so a future failure names the real ceiling of the device it happened on.include/gpu_iface/gpu_runtime_compat.hpp— memoizegetMaxSharedMemPerBlockper device id, so moving the guard onto a queried limit does not add a device query to every decode launch.Architecture / design notes
The helper already exists in
include/gpu_iface/gpu_runtime_compat.hppand is already exercised by three call sites inprefill.cuh, so this adds no new machinery — decode simply was not using it.gpuGetDevicehad to be hoisted. The existing call sits inside the partition-kvelsebranch, reached only whenseq_len > 256 && tmp != nullptr, while the check runs before that. It is now done once above the check and reused below.Why the helper is now cached
Querying the limit put
getMaxSharedMemPerBlockon everySingleDecodeWithKVCacheDispatchedcall, including theseq_len <= 256 || tmp == nullptrfast path that previously issued no device query at all — and the helper copies out a wholehipDeviceProp_trather than reading a single attribute. Measured on MI350X / ROCm 7.2:hipGetDeviceProperties, first callhipGetDeviceProperties, steady statehipFuncSetAttribute(already on this path)hipLaunchKernel(already on this path)getMaxSharedMemPerBlock, memoizedROCm memoizes device properties internally, so the steady-state cost was never dramatic — roughly 16% of the kernel launch the same dispatch already pays. Caching removes it anyway, along with the 6.4 us first-call hit, and the three
prefill.cuhcall sites benefit for free.The cache mirrors
getMultiProcessorCountdirectly above it:thread_localso concurrent callers never race, and0treated as "not cached" so a device reporting 0 is simply not memoized rather than poisoning the entry. Device ids outside[0, 64)fall back to querying every time rather than indexing out of bounds.getMaxSharedMemPerMultiprocessoris deliberately left uncached — identical shape, but it is only reached frompod.cuh/batch_pod.cuh, off this PR's path.CDNA3 impact
None. gfx942 reports
sharedMemPerBlock = 65536, so the comparison is byte-identical to the previous literal and no configuration changes status — only the message wording differs.Test plan
On gfx950 (MI350X), ROCm 7.2.0, torch 2.9.1:
sharedMemPerBlock = 163840.hipGetDevicePropertiesquery: identical (163840), stable over 1000 calls.ctr-rack31-mi300x-2, commit3708813f, pinned detached worktree): 27476 passed, 3585 skipped, 0 failed. This covers the second commit's memoization ofgetMaxSharedMemPerBlock, which matters because that helper is also called from three sites inprefill.cuh— so the change sits on the prefill dispatch path, not only decode.pre-commit run -aOn gfx942 (MI300X,
ctr-rack31-mi300x-2, same toolchain — torch 2.9.1+rocm7.2.0, HIP 7.2.26015, amd-aiter 0.1.10):sharedMemPerBlock = 65536, so the comparison is byte-identical to the previous literal and no configuration changes status on CDNA3.Side by side, same container stack, architecture the only variable:
sharedMemPerBlockwarpSize