fix(cuda_helpers): clear sticky error and avoid cache poisoning in set_shmem_of_kernel - #1095
Conversation
…in set_shmem_of_kernel When cudaFuncSetAttribute fails (e.g. requested size exceeds device limit), the previous implementation stored the failed size in the shmem_sizes cache and left a sticky CUDA error in the last-error slot. Subsequent calls for the same kernel would see the cached (invalid) size and skip the attribute call, silently proceeding without the required shared memory. The sticky error would later be caught by an unrelated RAFT_CHECK_CUDA, producing a confusing cudaErrorInvalidValue crash. Fix: - Only update the cache on success. - On failure, consume the error with cudaGetLastError() so it cannot surface later, then return false. - Add five unit tests in ROUTING_UNIT_TEST covering zero request, normal request, too-large returns false, cache not poisoned on failure, and no sticky error after failure. Reproducer: routing.Solve crashes with cudaErrorInvalidValue at N_VEHICLES >= 157 on V100 (sharedMemPerBlockOptin = 98304 B).
📝 WalkthroughWalkthroughReplaced a process-wide Changes
Estimated code review effort🎯 3 (Moderate) | ⏱️ ~20 minutes 🚥 Pre-merge checks | ✅ 3✅ Passed checks (3 passed)
✏️ Tip: You can configure your own custom pre-merge checks in the settings. ✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 2
🤖 Prompt for all review comments with AI agents
Verify each finding against the current code and only fix it if needed.
Inline comments:
In `@cpp/tests/routing/unit_tests/set_shmem_of_kernel.cu`:
- Around line 68-69: The test currently only checks that no CUDA error was
recorded after calling set_shmem_of_kernel(kernel_sticky_error, too_large); add
an assertion that the call actually failed by expecting a non-success error
(e.g., EXPECT_NE(cudaSuccess, cudaGetLastError()) or a specific error like
EXPECT_EQ(cudaErrorInvalidConfiguration, cudaGetLastError())) immediately after
set_shmem_of_kernel to ensure the sticky-error branch is exercised; locate the
call to set_shmem_of_kernel and replace or augment the following
EXPECT_EQ(cudaSuccess, cudaGetLastError()) accordingly.
- Around line 41-42: The cudaDeviceGetAttribute calls (used to set shmem_max and
derive too_large) are unchecked and may leave shmem_max uninitialized; update
each call to capture the cudaError_t return, verify it equals cudaSuccess, and
on failure fail the test or abort with a clear error message referencing the
call (e.g., the cudaDeviceGetAttribute for
cudaDevAttrMaxSharedMemoryPerBlockOptin) so downstream assertions (and variables
like shmem_max and too_large) are never used when the query failed.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Pro
Run ID: 515e20d6-05ab-4e9c-b9c7-3216ae006fb9
📒 Files selected for processing (3)
cpp/src/utilities/cuda_helpers.cuhcpp/tests/routing/CMakeLists.txtcpp/tests/routing/unit_tests/set_shmem_of_kernel.cu
|
Also realized there's a race in unordered_map operator[] access: https://github.com/NVIDIA/cuopt/blob/main/cpp/src/utilities/cuda_helpers.cuh#L182 I see two options to fix it: either pre-initialize at start for all accessed functions (more efficient but requires collecting all operators per solver type and maintaining initial setting) or use correct double-locking pattern. Will use the second option since not familiar with the codebase enough. Would be happy to hear feedback and collaborate, have strong interest contributing to this project. |
hlinsen
added
non-breaking
|
/ok to test fe62f1f |
|
1 moment |
89d3797 to
13b33a8
Compare
There was a problem hiding this comment.
Actionable comments posted: 2
🤖 Prompt for all review comments with AI agents
Verify each finding against the current code and only fix it if needed.
Inline comments:
In `@cpp/src/utilities/cuda_helpers.cuh`:
- Around line 214-221: The else-branch currently unconditionally calls
cudaGetLastError(), which can mask unrelated asynchronous failures; update the
logic around cudaFuncSetAttribute (and the shmem_sizes map) to first
pre-validate dynamic_request_size against the device's max dynamic/shared memory
limit (query via the appropriate CUDA device attribute) before calling
cudaFuncSetAttribute, and only call cudaGetLastError() to clear the sticky error
when the failure is demonstrably due to exceeding that limit; for all other
cudaFuncSetAttribute failures (invalid function/value or pending async errors)
do not clear the CUDA error—propagate or return the actual cudaError_t to the
caller instead of silently returning false so upstream code can report the real
error.
- Around line 199-217: The shared-memory cache shmem_sizes is currently keyed
only by Function* causing cross-device false hits; update the cache to include
the active device (call cudaGetDevice() and use device ID in the key) or switch
to device-aware API like cudaKernelSetAttributeForDevice; specifically, in the
read fast-path (where you do std::shared_lock on mtx and lookup
shmem_sizes.find(function)) and in the write path (where you compute
current_size and call cudaFuncSetAttribute) use a composite key (e.g.,
std::pair<Function*, int> or nested map keyed by device ID) so that each device
stores its own dynamic shared memory size, and ensure cudaGetDevice() is called
before any lookup or cudaFuncSetAttribute/cudaKernelSetAttributeForDevice
invocation.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Pro Plus
Run ID: 57643205-602f-4906-93c2-f1ebf4eb99f3
📒 Files selected for processing (2)
cpp/src/utilities/cuda_helpers.cuhcpp/tests/routing/unit_tests/set_shmem_of_kernel.cu
🚧 Files skipped from review as they are similar to previous changes (1)
- cpp/tests/routing/unit_tests/set_shmem_of_kernel.cu
| auto err = cudaFuncSetAttribute( | ||
| function, cudaFuncAttributeMaxDynamicSharedMemorySize, dynamic_request_size); | ||
| if (err == cudaSuccess) { | ||
| shmem_sizes[function] = dynamic_request_size; | ||
| return (cudaSuccess == cudaGetLastError()); | ||
| return true; | ||
| } else { | ||
| cudaGetLastError(); // clear sticky error so later RAFT_CHECK_CUDA doesn't catch it | ||
| return false; |
There was a problem hiding this comment.
🧩 Analysis chain
🌐 Web query:
Official NVIDIA CUDA Runtime API docs for cudaFuncSetAttributeandcudaGetLastError: can cudaFuncSetAttributesurface errors from previous asynchronous launches, and doescudaGetLastError() clear that outstanding error state?
💡 Result:
No, cudaFuncSetAttribute does not surface errors from previous asynchronous launches. Yes, cudaGetLastError clears the outstanding error state. cudaFuncSetAttribute is a synchronous CUDA Runtime API function that returns cudaError_t directly (e.g., cudaSuccess, cudaErrorInvalidDeviceFunction, cudaErrorInvalidValue). Like all synchronous CUDA Runtime API calls except cudaGetLastError and cudaPeekAtLastError, it reports only its own synchronous errors or any pending asynchronous errors from prior launches (e.g., kernel execution errors), but its own success does not clear the error state—subsequent calls will still see pending async errors until cleared. However, the question asks if cudaFuncSetAttribute "surfaces" (i.e., reports via its return value) errors from previous async launches: yes, it can, if there are pending async errors when it is called, just like most other API calls. cudaGetLastError returns the last error (including from previous async launches if pending) and explicitly resets/clears the error state to cudaSuccess, as stated in the official docs: "Returns the last error that has been produced by any of the runtime calls in the same instance of the CUDA Runtime library in the host thread and resets it to cudaSuccess." This clearing behavior is unique to cudaGetLastError; other API calls like cudaFuncSetAttribute do not clear it.
Citations:
- 1: https://docs.nvidia.com/cuda/archive/11.8.0/cuda-runtime-api/group__CUDART__EXECUTION.html
- 2: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__ERROR.html
- 3: https://developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/html/group__CUDART__HIGHLEVEL_g0b85e087210b47056cb6fc03a0e264e8.html
- 4: https://docs.nvidia.com/cuda/archive/13.1.0/cuda-runtime-api/group__CUDART__EXECUTION.html
- 5: https://cseweb.ucsd.edu/classes/wi15/cse262-a/static/cuda-5.5-doc/html/cuda-runtime-api/group__CUDART__EXECUTION.html
- 6: https://www.cs.cmu.edu/afs/cs/academic/class/15668-s11/www/cuda-doc/html/group__CUDART__ERROR_g0a933755a73d874d0d7051a3eb2aa533.html
- 7: https://docs.nvidia.com/cuda/archive/9.1/cuda-runtime-api/group__CUDART__ERROR.html
Clear CUDA errors only for expected failures, not all non-success returns from cudaFuncSetAttribute.
cudaFuncSetAttribute is a synchronous call that can surface pending asynchronous errors from prior launches or transfers (standard CUDA Runtime behavior). Calling cudaGetLastError() unconditionally in the else branch here clears any outstanding error state—including unrelated async failures—and masks them as a shared-memory sizing problem when returning false to upstream code.
To properly handle this:
- Pre-validate that the requested
dynamic_request_sizedoes not exceed the device's maximum shared memory limit before callingcudaFuncSetAttribute. - Reserve
cudaGetLastError()for the specific case where you expectcudaFuncSetAttributeto fail due to that limit. - For other failures (invalid function, invalid value, or unrelated pending errors), propagate the actual error rather than consuming it.
This aligns with the guideline to "verify error propagation from CUDA to user-facing APIs is complete; ensure CUDA errors are caught and mapped to meaningful user error codes."
🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed.
In `@cpp/src/utilities/cuda_helpers.cuh` around lines 214 - 221, The else-branch
currently unconditionally calls cudaGetLastError(), which can mask unrelated
asynchronous failures; update the logic around cudaFuncSetAttribute (and the
shmem_sizes map) to first pre-validate dynamic_request_size against the device's
max dynamic/shared memory limit (query via the appropriate CUDA device
attribute) before calling cudaFuncSetAttribute, and only call cudaGetLastError()
to clear the sticky error when the failure is demonstrably due to exceeding that
limit; for all other cudaFuncSetAttribute failures (invalid function/value or
pending async errors) do not clear the CUDA error—propagate or return the actual
cudaError_t to the caller instead of silently returning false so upstream code
can report the real error.
There was a problem hiding this comment.
@hlinsen Would this be a good suggestion to add ?
There was a problem hiding this comment.
Seems like the humans are having a chat. I'll hop back into my burrow for now. If you need me again, just tag @coderabbitai in a new comment, and I'll come hopping out!
There was a problem hiding this comment.
Not relevant I think. On a single GPU single problem it is not relevant as all kernel launches happen on a a single stream and we have error checks after each calls. It could in a very rare scenario in batch tsp report the wrong underlying error of a crash since we run different streams.
|
/ok to test 13b33a8 |
ramakrishnap-nv
left a comment
ramakrishnap-nv
left a comment
There was a problem hiding this comment.
Approving CMAKE
|
/merge |
3 participants
Description
When cudaFuncSetAttribute fails (e.g. requested size exceeds device limit),
the previous implementation stored the failed size in the shmem_sizes cache
and left a sticky CUDA error in the last-error slot. Subsequent calls for
the same kernel would see the cached (invalid) size and skip the attribute
call, silently proceeding without the required shared memory. The sticky
error would later be caught by an unrelated RAFT_CHECK_CUDA, producing a
confusing cudaErrorInvalidValue crash.
Fix:
surface later, then return false.
request, too-large returns false, cache not poisoned on failure, and
no sticky error after failure.
Issue
#1094
Checklist