Conversation
cabb2b2 to
15c236d
Compare
|
Tested locally on RTX 5070 / sm_120, CUDA 12.9, with both variants
Honest disclosure: small-N + equal-all (where Pass 2 is skipped because Pushed the swap as 15c236d. |
|
Hi @miscco — gentle ping when you have a moment, I think this is ready for another look. The int8_t → int suggestion is done (15c236d), with the before/after numbers in the table above. I also re-checked locally on a 5070 (CUDA 12.9): both tests pass, plus a 50M cross-check against std::lexicographical_compare across int/int64/uint8/float/double and the edge cases, and memcheck came back clean. Happy to make any changes — please let me know if anything's still blocking. Thanks! |
|
/ok to test 15c236d |
miscco
left a comment
miscco
left a comment
There was a problem hiding this comment.
Thanks a lot for the work. I believe we want to coalesce the two allocations we currently perform into a single one
…cation Address review feedback: - Replace the -1/0/+1 magic numbers with an `__lex_ordering` enum class (`__less`/`__equal`/`__greater`). - Coalesce the two temporary allocations into one. The frontend no longer composes the generic find_if + transform_reduce dispatches (each of which allocates its own scratch); instead a dedicated `__pstl_dispatch<__lexicographical_compare, __cuda>` queries both cub passes up front, allocates a single `__temporary_storage` sized to the larger pass, and reuses that scratch region for both FindIf and the 1-element TransformReduce read-back. Only the two 1-element result slots are extra. Also harden the common-length computation with `common_type_t` so iterators with differing difference_type compile.
|
@miscco thanks for the review! All three points are addressed in
Re-ran everything locally (RTX 5070, CUDA 12.9): tests pass and memcheck is clean. Let me know if anything else needs changing. |
🥳 CI Workflow Results🟩 Finished in 1h 43m: Pass: 100%/110 | Total: 2d 11h | Max: 1h 04m | Hits: 97%/317054See results here. |
miscco
left a comment
miscco
left a comment
There was a problem hiding this comment.
Sorry approved accidentally
Co-authored-by: Michael Schellenberger Costa <miscco@nvidia.com>
- Add a non-const `operator()` overload to `__lex_state_fn` for comparators with a non-const call operator. - Rename `__n1/__n2/__n` to `__count1/__count2/__count`. - Drop the read-back `TransformReduce`: a single-thread result kernel now resolves the answer in place on device from the find index, so only one value is copied back (no offset host round-trip, no second reduce launch, one sync instead of two). Also fixes the missing `is_nothrow_move_constructible` include.
|
No worries! Another look whenever you have time would be great. Thanks again! |
|
Looking at the implementation from a bird’s-eye view, I was thinking whether we could avoid calling an extra kernel if we would extend I think the current approach with using |
|
yeah I believe we could do better if we extended |
- Move the result kernel inside _CCCL_BEGIN_NAMESPACE_ARCH_DEPENDENT. - Rename the kernel's __k to __offset. - Simplify the kernel-arg casts to static_cast<void*> (drop the reinterpret_cast/const_cast pair; the launched locals are non-const). - Construct an __ensure_current_context guard so the stream's device is current for the device work. - Qualify ::cudaErrorMemoryAllocation and rethrow with a plain throw.
|
@bernhardmgruber @miscco agreed, extending FindIf is out of scope here. Happy to keep a tracking issue around to revisit avoiding the second load later. |
|
Sorry I totally dropped the ball on this. Ping me if I did not respond on monday |
No worries at all! Hope you have a great weekend :> |
|
@miscco When you have a chance, could you please take a look? Thanks! |
📝 SummarySummary by CodeRabbit
WalkthroughAdded the CUDA PSTL execution-policy overload and backend for ChangesParallel lexicographical comparison
Assessment against linked issues
Suggested reviewers: Priority: ➖ Normal Change: Feature Merge Risk: 🔵 Low · up to The new benchmark underreports logical reads by one compared element pair. Correct its accounting before using it for performance comparisons; no implementation correctness issue was established. 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: 1
ℹ️ Review info
⚙️ Run configuration
Configuration used: Repository: NVIDIA/cccl/.coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: bd6d63d9-e47a-4048-8d9b-35a9e44fe29d
📒 Files selected for processing (7)
libcudacxx/benchmarks/bench/lexicographical_compare/basic.culibcudacxx/include/cuda/std/__pstl/cuda/lexicographical_compare.hlibcudacxx/include/cuda/std/__pstl/dispatch.hlibcudacxx/include/cuda/std/__pstl/lexicographical_compare.hlibcudacxx/include/cuda/std/executionlibcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.lex.comparison/pstl_lexicographical_compare.culibcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.lex.comparison/pstl_lexicographical_compare_comp.cu
Included review availability: Your plan provides up to 12 included reviews per hour; 9 remain after this review.
| thrust::device_vector<T> dinput2(elements, thrust::no_init); | ||
| prepare_inputs(dinput1, dinput2, violation_point); | ||
|
|
||
| state.add_global_memory_reads<T>(2 * violation_point); |
There was a problem hiding this comment.
🚀 Performance & Scalability | 🟡 Minor | ⚡ Quick win
🔎 Supported by static analysis
🏁 Script executed:
sed -n '1,110p' libcudacxx/benchmarks/bench/lexicographical_compare/basic.cu
rg -n -C4 'add_global_memory_reads|violation_point|lexicographical_compare' libcudacxx/benchmarks libcudacxx/include/cuda/std/__pstl/cudaRepository: NVIDIA/cccl
Length of output: 50367
🏁 Script executed:
#!/bin/bash
set -eu
sed -n '145,235p' libcudacxx/include/cuda/std/__pstl/cuda/lexicographical_compare.h
printf '\\n-- CUB DeviceFind definitions and relevant FindIf code --\\n'
rg -n -l 'class DeviceFind|struct DeviceFind|DeviceFind::FindIf|FindIf\\(' --glob '*.{h,hpp,cu,cpp,cuh}' . | head -40
printf '\\n-- focused CUB references --\\n'
rg -n -C8 'DeviceFind|FindIf' libcudacxx/include third_party cub thrust --glob '*.{h,hpp,cu,cpp,cuh}' 2>/dev/null | head -240Repository: NVIDIA/cccl
Length of output: 20911
🏁 Script executed:
#!/bin/bash
set -eu
sed -n '68,170p' cub/cub/device/device_find.cuh
printf '\\n-- dispatch_find structure and kernels --\\n'
rg -n -C6 'dispatch|kernel|Scan|Flag|predicate|scan_op|num_items|FindIf' cub/cub/device/dispatch/dispatch_find.cuh | head -320
printf '\\n-- complete focused implementation slices --\\n'
sed -n '1,260p' cub/cub/device/dispatch/dispatch_find.cuhRepository: NVIDIA/cccl
Length of output: 18263
🏁 Script executed:
#!/bin/bash
set -eu
rg -n -C10 'class AgentFind|struct AgentFind|Process\\(|Load|num_items|tile_size|found_pos' cub/cub/agent/agent_find.cuh | head -360Repository: NVIDIA/cccl
Length of output: 287
🏁 Script executed:
#!/bin/bash
set -eu
rg -n -F 'Process(' cub/cub/agent/agent_find.cuh
rg -n -F 'num_items' cub/cub/agent/agent_find.cuh
wc -l cub/cub/agent/agent_find.cuhRepository: NVIDIA/cccl
Length of output: 587
🏁 Script executed:
#!/bin/bash
set -eu
sed -n '50,205p' cub/cub/agent/agent_find.cuhRepository: NVIDIA/cccl
Length of output: 5494
🏁 Script executed:
#!/bin/bash
set -eu
nl -ba libcudacxx/benchmarks/bench/lexicographical_compare/basic.cu | sed -n '15,82p'
printf '\\n-- lexicographical predicate --\\n'
nl -ba libcudacxx/include/cuda/std/__pstl/cuda/lexicographical_compare.h | sed -n '80,145p'Repository: NVIDIA/cccl
Length of output: 6166
important: Count the logical differing pair. violation_point is zero-based, so the logical comparison covers indices 0 through violation_point and accounts for 2 * (violation_point + 1) input elements. Update both add_global_memory_reads<T> calls and revise the comment at lines 22-23 to describe logical comparison work, not an upper bound on physical device reads. CUB DeviceFind::FindIf can read positions beyond the first differing pair because its parallel tiles process assigned elements concurrently.
4 participants
Resolves #7752. Implements the CUDA backend for
cuda::std::lexicographical_compare.Design
Following the suggestion on #7752: a
cuda::zip_transform_iteratorproduces anint8_tper pair (-1LESS,+1GREATER,0EQUIV),cub::DeviceFind::FindIffinds the first non-zero index with early termination, then a 1-element
cub::DeviceReduce::TransformReduce(withcuda::std::identity{}+plus<int8_t>{}over
int8_t{0}) reads the state at that position back to host. Final answer isstate < 0.Spec complexity bound (≤
2 * min(N1, N2)compcalls) is met: Pass 1 short-circuitsat
2kworst-case for early-termination indexk; Pass 2 adds 1; total ≤2 * min - 1.I prototyped four value-recovery variants locally (1-element FindIf, 1-element
TransformReduce, custom 1-thread kernel, two-FindIf baseline). TransformReduce-identity
won consistently by 5–30% in p50 because cub's 1-element reduce path has lower launch
overhead than 1-element FindIf in this regime. Happy to switch shape if you'd prefer
something different.
Performance (RTX 5070 / sm_120, CUDA 12.9)
Realistic workload — two independent random sequences (≈always diverges at k=0):
Worst case for the early-termination design (late divergence and equal-all at
scale): within 5–10% of the no-early-termination baseline.
Tests
pstl_lexicographical_compare.cu(defaultless) — 480 assertionspstl_lexicographical_compare_comp.cu(customCompare) — 192 assertionsBoth cover the
all_typesmatrix × 4 SECTIONs (default stream / provided stream /provided memory_resource / both), plus edge cases: empty vs empty, empty/non-empty
mixed, single-element ranges (equal/less/greater — exercises the
n=1path throughboth passes), prefix in either direction, divergence at
k=0/N/2/N-1,random_access_iteratorwrapper, andgreater<>/less<>duality(
lex_cmp(s2, s1, greater) == lex_cmp(s1, s2, less)).bench/lexicographical_compare/basic.cumirrorsis_heap_until/basic.cu's shapewith
2 * violation_pointmemory accounting andViolationAt ∈ {1.0, 0.5, 0.01}.