[libcu++] Use CCCL Runtime in CUB metric pre-runs pt 1#9802
Conversation
📝 WalkthroughSummary by CodeRabbit
WalkthroughCUDA benchmarks now allocate through device memory pools and CUDA buffers, synchronize benchmark streams directly, and construct stream-aware CUB environments. Shared helpers add stream access, pinned memory, default bounds, and device-buffer generators. ChangesCUDA benchmark memory and execution updates
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
🧹 Nitpick comments (4)
cub/benchmarks/bench/run_length_encode/encode.cu (1)
50-64: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion:
in_keys(lines 57-58) is only read via the const pointerd_in_keys; declare itconst autoper the const-correctness guideline (same note asby_key.cu).diff
- auto in_keys = + const auto in_keys = generate.uniform.key_segments(elements, min_segment_size, max_segment_size).device_buffer<T>(stream, device);Source: Coding guidelines
cub/benchmarks/bench/run_length_encode/non_trivial_runs.cu (1)
51-64: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion:
in_keys(lines 58-59) is read-only (onlyconst T* d_in_keysis ever extracted); declare itconst autoper the const-correctness guideline, same as the siblingencode.cu/by_key.cubenchmarks.diff
- auto in_keys = + const auto in_keys = generate.uniform.key_segments(elements, min_segment_size, max_segment_size).device_buffer<T>(stream, device);Source: Coding guidelines
cub/benchmarks/bench/select/unique_by_key.cu (1)
59-74: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion:
in_vals(line 64) andin_keys(lines 67-68) are only read through const pointers (d_in_vals,d_in_keys); declare themconst autoper the const-correctness guideline (same note asby_key.cu).diff
- auto in_vals = cuda::make_device_buffer<ValueT>(stream, device, elements, ValueT{}); + const auto in_vals = cuda::make_device_buffer<ValueT>(stream, device, elements, ValueT{}); auto out_vals = cuda::make_device_buffer<ValueT>(stream, device, elements, cuda::no_init); auto out_keys = cuda::make_device_buffer<KeyT>(stream, device, elements, cuda::no_init); - auto in_keys = + const auto in_keys = generate.uniform.key_segments(elements, min_segment_size, max_segment_size).device_buffer<KeyT>(stream, device);Source: Coding guidelines
cub/benchmarks/bench/reduce/by_key.cu (1)
49-64: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Make read-only buffers const
in_valsandin_keysare only read here, so declare bothconst autoto match the constdata()overloads and avoid accidental mutation.Source: Coding guidelines
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 9b15726a-0223-4e92-be80-d82ddfd1b442
📒 Files selected for processing (6)
cub/benchmarks/bench/reduce/by_key.cucub/benchmarks/bench/run_length_encode/encode.cucub/benchmarks/bench/run_length_encode/non_trivial_runs.cucub/benchmarks/bench/select/unique.cucub/benchmarks/bench/select/unique_by_key.cunvbench_helper/nvbench_helper/nvbench_helper.cuh
| // Run once to get num_runs for memory accounting | ||
| (void) cub::DeviceRunLengthEncode::Encode( | ||
| d_in_keys, d_out_keys, d_out_counts, d_num_runs_out, static_cast<OffsetT>(elements)); | ||
| cudaDeviceSynchronize(); | ||
| d_in_keys, | ||
| d_out_keys, | ||
| d_out_counts, | ||
| d_num_runs_out, | ||
| static_cast<OffsetT>(elements), | ||
| cub_bench_env(memory_resource, stream)); | ||
| stream.sync(); |
There was a problem hiding this comment.
🩺 Stability & Availability | 🟠 Major | ⚡ Quick win
important: this pre-run Encode call discards the return status via (void), while the equivalent one-time correctness call in every sibling benchmark file (by_key.cu, non_trivial_runs.cu, unique.cu, unique_by_key.cu) uses _CCCL_TRY_CUDA_API. A failed launch (e.g. temp-storage allocation failure through memory_resource) goes unnoticed, and num_runs_out[0] read right after may reflect an uninitialized/garbage value, silently corrupting the reported memory-access counters.
diff
- (void) cub::DeviceRunLengthEncode::Encode(
- d_in_keys,
- d_out_keys,
- d_out_counts,
- d_num_runs_out,
- static_cast<OffsetT>(elements),
- cub_bench_env(memory_resource, stream));
+ _CCCL_TRY_CUDA_API(
+ cub::DeviceRunLengthEncode::Encode,
+ "Encode failed",
+ d_in_keys,
+ d_out_keys,
+ d_out_counts,
+ d_num_runs_out,
+ static_cast<OffsetT>(elements),
+ cub_bench_env(memory_resource, stream));As per path instructions, "Focus on algorithm correctness, temporary-storage protocol, dispatch/policy selection, stream behavior, CUDA error handling, synchronization, memory access safety, performance regressions, and test coverage" for cub/**/*.
📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
| // Run once to get num_runs for memory accounting | |
| (void) cub::DeviceRunLengthEncode::Encode( | |
| d_in_keys, d_out_keys, d_out_counts, d_num_runs_out, static_cast<OffsetT>(elements)); | |
| cudaDeviceSynchronize(); | |
| d_in_keys, | |
| d_out_keys, | |
| d_out_counts, | |
| d_num_runs_out, | |
| static_cast<OffsetT>(elements), | |
| cub_bench_env(memory_resource, stream)); | |
| stream.sync(); | |
| // Run once to get num_runs for memory accounting | |
| _CCCL_TRY_CUDA_API( | |
| cub::DeviceRunLengthEncode::Encode, | |
| "Encode failed", | |
| d_in_keys, | |
| d_out_keys, | |
| d_out_counts, | |
| d_num_runs_out, | |
| static_cast<OffsetT>(elements), | |
| cub_bench_env(memory_resource, stream)); | |
| stream.sync(); |
Source: Path instructions
|
|
||
| #if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA | ||
| template <typename T> | ||
| auto device_buffer(cuda::stream_ref stream, cuda::device_ref device, T min, T max) | ||
| { | ||
| auto buffer = cuda::make_device_buffer<T>(stream, device, m_elements, cuda::no_init); | ||
| stream.sync(); | ||
|
|
||
| cuda::std::span<T> span(buffer.data(), buffer.size()); | ||
| gen_device(m_seed, span, m_entropy, min, max); | ||
|
|
||
| ++m_seed; | ||
| return buffer; | ||
| } | ||
| #endif // THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA |
There was a problem hiding this comment.
🩺 Stability & Availability | 🟠 Major | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash
# Inspect gen_device implementation for stream/sync semantics
fd -e cu -e cuh . nvbench_helper --exec rg -nP -C4 'gen_device' {} \;
rg -nP -C4 '\bgen_device\b' -g '*.cu' | head -60Repository: NVIDIA/cccl
Length of output: 2228
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Inspect the generator dispatch and the device executor implementation.
rg -n -C 6 'template <typename T>\s*void gen\(|gen\(executor::device|struct executor|class executor|namespace executor' nvbench_helper/nvbench_helper/nvbench_helper.cuh
rg -n -C 8 'executor::device|gen\(' nvbench_helper/nvbench_helper/nvbench_helper.cuh
# Show the relevant surrounding sections with line numbers.
sed -n '500,700p' nvbench_helper/nvbench_helper/nvbench_helper.cuh
sed -n '700,860p' nvbench_helper/nvbench_helper/nvbench_helper.cuhRepository: NVIDIA/cccl
Length of output: 149
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Find the executor::device implementation and how gen dispatches.
rg -n -C 6 'gen\(executor::device|executor::device|struct device|class device' nvbench_helper
rg -n -C 6 'template <typename.*>\s*void gen\(' nvbench_helper/nvbench_helper/nvbench_helper.cuh
# Read the relevant ranges around gen() and any executor definitions.
sed -n '540,660p' nvbench_helper/nvbench_helper/nvbench_helper.cuh
sed -n '660,780p' nvbench_helper/nvbench_helper/nvbench_helper.cuhRepository: NVIDIA/cccl
Length of output: 14002
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Read the device generator implementation and the executor dispatch around it.
sed -n '97,220p' nvbench_helper/nvbench_helper/nvbench_helper.cu
sed -n '260,320p' nvbench_helper/nvbench_helper/nvbench_helper.cuRepository: NVIDIA/cccl
Length of output: 5245
🌐 Web query:
CURAND curandGenerateUniformDouble asynchronous stream semantics curandSetStream documentation
💡 Result:
In the cuRAND library, functions such as curandGenerateUniformDouble operate asynchronously by launching kernels onto the specified CUDA stream [1][2]. When these functions are called, they initiate the random number generation process and return control to the host immediately, without waiting for the GPU to complete the operation [1]. The stream associated with these operations is controlled by the curandSetStream function [2]. Key details regarding this mechanism include: Stream Association: The curandSetStream(generator, stream) function configures the generator to use a specific cudaStream_t for all subsequent kernel launches [2]. If curandSetStream is not explicitly called for a given generator, the library defaults to using the null stream (stream 0) [2]. Asynchronous Semantics and Synchronization: Because these generation functions return asynchronously, the host-side code must ensure proper synchronization if subsequent operations depend on the generated results [1]. If other kernels or operations (e.g., memory copies) in a different stream need to consume the random numbers, the developer must use standard CUDA synchronization primitives, such as cudaDeviceSynchronize or stream/event management routines, to ensure the generation kernel has finished before the dependent operations begin [1]. In summary, curandGenerateUniformDouble and related generation functions queue work in the stream assigned via curandSetStream and provide asynchronous execution, requiring explicit CUDA synchronization for dependent tasks [1][2].
Citations:
- 1: https://docs.nvidia.com/cuda/curand/host-api-overview.html
- 2: https://docs.nvidia.com/cuda/archive/13.0.1/curand/group__HOST.html
important: gen_device can return before cuRAND finishes filling the buffer. curandGenerate* is async, and m_gen is never tied to stream, so the first CUB consumer on stream can race. Tie the generator to stream or sync after generation.
🥳 CI Workflow Results🟩 Finished in 2h 46m: Pass: 100%/263 | Total: 12d 11h | Max: 2h 45m | Hits: 14%/1141632See results here. |
2 participants
#2800 hitting cub benchmarks