[libcu++] Add CUDA logical endpoint abstractions - #11219
Open
pciolkosz wants to merge 6 commits into
Open
Conversation
Contributor
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (1)
💤 Files with no reviewable changes (1)
Included review availability: Your plan provides up to 12 included reviews per hour; 10 remain after this review. 📝 SummarySummary by CodeRabbit
WalkthroughThe change adds CUDA 13.3 logical endpoint driver APIs, unicast and multicast C++ wrappers, creation-only ownership semantics, launch transformations, and lifecycle and fabric-operation tests. ChangesLogical Endpoint API
Merge Risk: 🟡 Moderate · up to This adds CUDA logical endpoint APIs and tests, but older-driver use may terminate rather than return an error, readiness can be reported after timeout, and NVRTC test configurations may fail to compile. These issues should be resolved before merge. Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Actionable comments posted: 6
🧹 Nitpick comments (3)
libcudacxx/test/support/logical_endpoint_test_helper.h (1)
244-244: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion: This loop repeats
wait_for_mbarrier_completionwith a hard-coded1000000instead oflogical_endpoint_test::wait_iterations. The same duplication is at lines 302 and 526. Reuse the helper so the timeout bound stays in one place.- for (int iteration = 0; iteration < 1000000; ++iteration) - { - if (cuda::ptx::mbarrier_try_wait_parity(cuda::ptx::sem_acquire, cuda::ptx::scope_cta, &barrier, 0)) - { - *status = logical_endpoint_test::status_success; - return; - } - } - - *status = logical_endpoint_test::status_timeout; + *status = logical_endpoint_test::wait_for_mbarrier_completion(&barrier) + ? logical_endpoint_test::status_success + : logical_endpoint_test::status_timeout;libcudacxx/test/libcudacxx/cuda/logical_endpoint/logical_endpoint.pass.cpp (1)
343-344: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: include
<cuda.h>directly.The test uses driver symbols
CUmemGenericAllocationHandle(Line 174),::CUlogicalEndpointId(Line 247), and theCU_LOGICAL_ENDPOINT_IPC_HANDLE_TYPE_*macros here. It gets them transitively from<cuda/logical_endpoint>. Include the driver header directly so the test does not break when the public header changes its includes.As per coding guidelines: "Include all headers needed by the symbols being used; do not rely on transitive includes."
Source: Coding guidelines
libcudacxx/include/cuda/__logical_endpoint/common.h (1)
500-500: 🩺 Stability & Availability | 🔵 Trivial | ⚡ Quick winsuggestion: Use host
::std::chrono::steady_clockfor both samples and convert the elapsed value to::cuda::std::chrono::nanoseconds. Here,cuda::std::chrono::high_resolution_clockaliases non-steadysystem_clockbecause_LIBCUDACXX_HAS_MONOTONIC_CLOCK()is always0; wall-clock adjustments can cause early or extended timeouts. Do not usecuda::std::chrono::steady_clock, which is unavailable in this configuration.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 74f706fb-25c0-40e6-8347-5ea2e96cc5e7
📒 Files selected for processing (9)
libcudacxx/include/cuda/__driver/driver_api.hlibcudacxx/include/cuda/__logical_endpoint/common.hlibcudacxx/include/cuda/__logical_endpoint/multicast.hlibcudacxx/include/cuda/__logical_endpoint/unicast.hlibcudacxx/include/cuda/logical_endpointlibcudacxx/test/libcudacxx/cuda/ccclrt/logical_endpoint/logical_endpoint.culibcudacxx/test/libcudacxx/cuda/logical_endpoint/logical_endpoint.pass.cpplibcudacxx/test/libcudacxx/cuda/logical_endpoint/logical_endpoint_multiprocess.pass.cpplibcudacxx/test/support/logical_endpoint_test_helper.h
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.
Comment on lines
+1198
to
+1204
| [[nodiscard]] _CCCL_HOST_API inline ::cudaError_t | ||
| __logicalEndpointIdReleaseNoThrow(::CUlogicalEndpointId __base_le_id, ::cuuint32_t __count) noexcept | ||
| { | ||
| static auto __driver_fn = | ||
| _CCCLRT_GET_DRIVER_FUNCTION_VERSIONED(cuLogicalEndpointIdRelease, cuLogicalEndpointIdRelease, 13, 3); | ||
| return static_cast<::cudaError_t>(__driver_fn(__base_le_id, __count)); | ||
| } |
Contributor
There was a problem hiding this comment.
🩺 Stability & Availability | 🟠 Major | 🏗️ Heavy lift
🔎 Supported by static analysis
🏁 Script executed:
#!/bin/bash
# Inspect the versioned driver-function macros and the exception-escape annotations.
fd -t f 'driver_api.h' -x rg -n -C4 '_CCCLRT_GET_DRIVER_FUNCTION_VERSIONED|_CCCLRT_GET_DRIVER_FUNCTION\b' {} \; | head -60
# Confirm every existing NoThrow wrapper that resolves a versioned entry point.
rg -n -B2 -A6 'NoThrow\(' --type=cpp -g '**/__driver/driver_api.h' | rg -n 'NoThrow|VERSIONED|NOLINT' | head -60
# Check whether the repository offers a non-throwing entry-point lookup helper.
rg -n '__get_driver_entry_point[a-zA-Z_]*' --type=cpp | head -30Repository: NVIDIA/cccl
Length of output: 2738
🏁 Script executed:
#!/bin/bash
set -e
# Inspect the bound lookup implementation, its exception behavior, and the affected wrappers.
rg -n -C12 '__get_driver_entry_point|__driver_fn_lookup|__logicalEndpoint(IdRelease|Destroy|BindAddr|Import)NoThrow|cuLogicalEndpoint' libcudacxx/include/cuda/__driver libcudacxx/include/cuda --glob '*.h' | head -240
# Inspect the relevant tests and all versioned NoThrow uses in the driver API.
rg -n -C8 'logicalEndpoint|GET_DRIVER_FUNCTION_VERSIONED|NOLINT\(bugprone-exception-escape\)' libcudacxx tests --glob '*.{h,hpp,cpp,cu}' | head -260Repository: NVIDIA/cccl
Length of output: 44909
🤖 get_repo_knowledge executed:
get_repo_knowledge NVIDIA/cccl /tmp/coderabbit-repo-knowledge/nvidia-cccl-2f773195/learnings /tmp/coderabbit-repo-knowledge/nvidia-cccl-2f773195/conventions
Length of output: 47796
🏁 Script executed:
#!/bin/bash
set -e
# Read the exact declarations and definitions without truncating the causal chain.
rg -l '__get_driver_entry_point' libcudacxx/include --glob '*.h' | while read -r f; do
echo "=== $f ==="
rg -n -C18 '__get_driver_entry_point' "$f"
done
echo "=== affected wrapper ranges ==="
sed -n '1170,1250p;1270,1320p' libcudacxx/include/cuda/__driver/driver_api.h
echo "=== relevant tests ==="
rg -n -C10 'logicalEndpoint|LogicalEndpoint' libcudacxx tests --glob '*.{h,hpp,cpp,cu}' | head -260Repository: NVIDIA/cccl
Length of output: 37128
important: Guard versioned entry-point lookup in the four noexcept wrappers. When __get_driver_entry_point_impl reports CU_GET_PROC_ADDRESS_VERSION_NOT_SUFFICIENT, it throws cuda_error. The static lookup in __logicalEndpointIdReleaseNoThrow, __logicalEndpointDestroyNoThrow, __logicalEndpointBindAddrNoThrow, and __logicalEndpointImportNoThrow can therefore escape their noexcept functions and call std::terminate. Catch the lookup failure and return cudaErrorNotSupported.
Comment on lines
+25
to
+32
| #include <cstdio> | ||
| #include <stdexcept> | ||
|
|
||
| #include <cuda_runtime_api.h> | ||
|
|
||
| #include "test_macros.h" | ||
|
|
||
| #if _CCCL_CTK_AT_LEAST(13, 3) && !TEST_COMPILER(NVRTC) |
Contributor
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟠 Major | ⚡ Quick win
🔎 Supported by static analysis
🏁 Script executed:
#!/bin/bash
# Verify how sibling libcudacxx tests gate NVRTC and confirm the lit feature name.
fd -t f 'lit.cfg|*.cfg.in|*.py' libcudacxx/test --exec rg -n 'nvrtc' {} \; | head -50
rg -n 'UNSUPPORTED: nvrtc' libcudacxx/test/libcudacxx/cuda -g '*.pass.cpp' | head -20Repository: NVIDIA/cccl
Length of output: 2538
🏁 Script executed:
#!/bin/bash
set -eu
printf '%s\n' '--- target ---'
cat -n libcudacxx/test/libcudacxx/cuda/logical_endpoint/logical_endpoint.pass.cpp | sed -n '1,55p'
printf '%s\n' '--- sibling ---'
cat -n libcudacxx/test/libcudacxx/cuda/logical_endpoint/logical_endpoint_multiprocess.pass.cpp | sed -n '1,35p'
printf '%s\n' '--- nearby NVRTC-gated test ---'
cat -n libcudacxx/test/libcudacxx/cuda/complex/std_interop/complex.conv.pass.cpp | sed -n '1,25p'
printf '%s\n' '--- lit configuration references ---'
rg -n -g '*.py' -g '*lit.cfg*' -g '*.cfg.in' 'nvrtc|UNSUPPORTED' libcudacxx/test | head -80Repository: NVIDIA/cccl
Length of output: 6825
important: Add // UNSUPPORTED: nvrtc before the includes.
The #if at line 32 does not guard any include. Lit exposes nvrtc as a feature, and sibling tests use this directive. Without it, NVRTC can process the host-only headers before the guard and fail.
🧰 Tools
🪛 Cppcheck (2.21.0)
[error] 32-32: failed to evaluate #if condition, undefined function-like macro invocation
(syntaxError)
Source: Coding guidelines
Contributor
There was a problem hiding this comment.
Actionable comments posted: 1
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 2a7cfb2d-b3e3-45bf-ae1a-75b2c6e65a40
📒 Files selected for processing (1)
libcudacxx/test/support/logical_endpoint_test_helper.h
Included review availability: Your plan provides up to 12 included reviews per hour; 10 remain after this review.
| { | ||
| template <class Config> | ||
| TEST_DEVICE_FUNC void operator()( | ||
| Config config, |
Contributor
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win
🔎 Supported by static analysis
🤖 get_repo_knowledge executed:
get_repo_knowledge NVIDIA/cccl /tmp/coderabbit-repo-knowledge/nvidia-cccl-2f773195/conventions /tmp/coderabbit-repo-knowledge/nvidia-cccl-2f773195/learnings
Length of output: 47856
🏁 Script executed:
#!/bin/bash
set -eu
file="libcudacxx/test/support/logical_endpoint_test_helper.h"
printf '%s\n' '--- changed hunks ---'
git diff -- "$file"
printf '%s\n' '--- target declarations and nearby bodies ---'
sed -n '180,245p;250,305p;315,365p;375,420p;435,480p;505,550p' "$file"
printf '%s\n' '--- Config definition and uses in this file ---'
rg -n -C 4 '\b(Config|config)\b' "$file"
printf '%s\n' '--- applicable local guidance ---'
for f in references/libcudacxx.md .agent/skills/cccl-style/SKILL.md .agent/skills/cccl-test/SKILL.md; do
if [ -f "$f" ]; then
printf '\n--- %s ---\n' "$f"
sed -n '1,240p' "$f"
fi
doneRepository: NVIDIA/cccl
Length of output: 17166
🏁 Script executed:
#!/bin/bash
set -eu
file="libcudacxx/test/support/logical_endpoint_test_helper.h"
printf '%s\n' '--- declarations with exact line numbers ---'
nl -ba "$file" | sed -n '205,225p;266,286p;329,349p;384,404p;444,464p;514,534p'
printf '%s\n' '--- all Config parameter declarations ---'
rg -n -C 12 'Config config' "$file"Repository: NVIDIA/cccl
Length of output: 5968
🏁 Script executed:
#!/bin/bash
set -eu
printf '%s\n' '--- path-specific test references ---'
find . -path '*/references/libcudacxx.md' -print
printf '%s\n' '--- Config-related API declarations and definitions ---'
rg -n -C 6 'gpu_thread|block\.rank|struct .*config|class .*config|template.*Config' libcudacxx/include libcudacxx/test/support 2>/dev/null | head -240
printf '%s\n' '--- nearby helper call sites ---'
rg -n -C 5 'fabric_(try_put|ring_put)|logical_endpoint.*kernel|smoke_kernel' libcudacxx/test | head -260Repository: NVIDIA/cccl
Length of output: 43943
🏁 Script executed:
#!/bin/bash
set -eu
file="libcudacxx/test/support/logical_endpoint_test_helper.h"
printf '%s\n' '--- exact bodies for all six operators ---'
sed -n '212,270p' "$file"
sed -n '273,333p' "$file"
sed -n '336,388p' "$file"
sed -n '391,449p' "$file"
sed -n '451,519p' "$file"
sed -n '521,580p' "$file"Repository: NVIDIA/cccl
Length of output: 14592
🏁 Script executed:
#!/bin/bash
set -eu
for f in .agent/skills/cccl-test/references/libcudacxx.md .agent/skills/cccl-style/references/libcudacxx.md; do
printf '\n--- %s ---\n' "$f"
sed -n '1,260p' "$f"
done
printf '%s\n' '--- resolved hierarchy API definitions ---'
rg -n -C 8 'struct (gpu_thread|block|grid)|gpu_thread.*count|block.*rank|constexpr.*count|constexpr.*rank' libcudacxx/include/cuda | head -260Repository: NVIDIA/cccl
Length of output: 29839
🏁 Script executed:
#!/bin/bash
set -eu
printf '%s\n' '--- hierarchy declarations ---'
rg -n -C 10 'count\(.*hierarchy|rank\(.*hierarchy|struct (gpu_thread|block|grid)|class (gpu_thread|block|grid)' libcudacxx/include/cuda/__hierarchy libcudacxx/include/cuda | head -320Repository: NVIDIA/cccl
Length of output: 1769
🏁 Script executed:
#!/bin/bash
set -eu
printf '%s\n' '--- hierarchy files ---'
find libcudacxx/include/cuda/__hierarchy -maxdepth 3 -type f -print
printf '%s\n' '--- rank/count definitions ---'
rg -n -C 8 'rank\s*\(|count\s*\(' libcudacxx/include/cuda/__hierarchy
printf '%s\n' '--- nearby test helper parameter style ---'
rg -n -C 2 'operator\(\)\(|const Config config|Config config' libcudacxx/test/support | head -240Repository: NVIDIA/cccl
Length of output: 50367
important: Qualify all six read-only Config parameters as const. The operators only pass config to cuda::gpu_thread.count and cuda::block.rank; apply const Config config at lines 216, 277, 340, 395, 455, and 525.
📍 Affects 1 file
libcudacxx/test/support/logical_endpoint_test_helper.h#L216-L216(this comment)libcudacxx/test/support/logical_endpoint_test_helper.h#L277-L277libcudacxx/test/support/logical_endpoint_test_helper.h#L340-L340libcudacxx/test/support/logical_endpoint_test_helper.h#L395-L395libcudacxx/test/support/logical_endpoint_test_helper.h#L455-L455libcudacxx/test/support/logical_endpoint_test_helper.h#L525-L525
Sources: Coding guidelines, Path instructions
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
Contributor
There was a problem hiding this comment.
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (1)
libcudacxx/include/cuda/__logical_endpoint/common.h (1)
506-508: 🩺 Stability & Availability | 🟡 Minor | ⚡ Quick winimportant: Check the deadline before returning success.
The loop calls
__is_ready()at Line 506 before it checks__elapsed. If readiness changes after__timeoutexpires, this path returnstrueafter the deadline. Check the deadline before accepting readiness, or prevent a post-deadline readiness result from returningtrue.
🧹 Nitpick comments (1)
libcudacxx/include/cuda/__logical_endpoint/common.h (1)
393-393: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Add
_CCCL_HOST_APIto the default constructor.This constructor is in a
libcudacxx/includeheader. The repository rule requires every function to use an appropriate_CCCL_*_APIannotation. Mark this host-only constructor consistently with the other host operations.Source: Coding guidelines
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 2a47aa6f-a5d4-4961-afb6-70f2b85eb2ae
📒 Files selected for processing (3)
libcudacxx/include/cuda/__driver/driver_api.hlibcudacxx/include/cuda/__logical_endpoint/common.hlibcudacxx/test/libcudacxx/cuda/logical_endpoint/logical_endpoint.pass.cpp
🚧 Files skipped from review as they are similar to previous changes (1)
- libcudacxx/include/cuda/__driver/driver_api.h
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
Contributor
🔬 CUB benchmark SASS comparisonHow to request a benchmark run
Targets with a SASS change
|
Contributor
🥳 CI Workflow Results🟩 Finished in 3h 31m: Pass: 100%/195 | Total: 4d 03h | Max: 2h 06m | Hits: 59%/892060See results here. AI failure analysis1. cuVS generated-kernel compilation: distributed sccache failure · 1 jobExplanation: The cuVS build delegated compilation to distributed sccache, and the remote compilation failure became fatal because local fallback was explicitly disabled. The job log contains no compiler diagnostic, while thousands of objects compiled successfully, so the evidence indicates an infrastructure failure rather than a demonstrated error in the PR's logical-endpoint changes. Evidence: Copy this prompt into a coding agentJobs: |
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
1 participant
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.
This PR adds types that manage logical endpoint lifetime, import/export, binding memory and passing them to kernels.
Future work will provide device-side interfaces that use them on top of
cuda::ptx::, but for now testing is done using the PTX layer.In a high level overview:
We have process-local logical endpoint ids that you can you use create the endpoints. They come in a range and the endpoint owning types refcount the range. You can get an export handle that you can share with other processes or nodes and they will import them to a new process-local id.
The endpoint creation takes a separate spec object that you can use to check the size limit and if the endpoint it supported in the first place.
Once you established the endpoint you can bind memory to it at a specific offset and everyone who imported the endpoint can access it through at that offset.
The endpoint types are strongly typed, because each endpoint type supports a separate set of instructions and for a given instruction encoding there is only one endpoint type that is legal to use for it.
The owning endpoint type also has a launch transform to cast it to a _ref.
Finally, you can also just use
logical_endpoint_idas an arithmetic type to address multiple types of endpoints in the same range from a kernel and only convert it to an endpoint ref at the fabric operation callThe import/export APIs will come in a separate future PR