Fea/6482 enable grid constant for all non-mutable CUB kernel parameters by toxicteddy00077 · Pull Request #6642 · NVIDIA/cccl

toxicteddy00077 · 2025-11-17T17:00:31Z

Description

closes #6482

Just added _CCCL_GRID_CONSTANT const to the immutable params for the necessary kernels

copy-pr-bot · 2025-11-17T17:00:35Z

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

…ea/6482-enable-grid-constant Merge upstream

bernhardmgruber · 2025-11-17T17:20:29Z

/ok to test bf617f5

bernhardmgruber · 2025-11-17T17:21:30Z

Thank you for this contribution! I think we should also do a benchmark and a SASS check somewhere, just because I am curious about the impact of this change. @toxicteddy00077 you can leave this to us.

cub/cub/device/dispatch/kernels/kernel_for_each.cuh

fbusato

thanks @toxicteddy00077. It is great to see this feature in CUB.

Initial feedbacl: This work could be extended to more parameters, e.g. cuda::std::array and pointers. Also, we need to double-check the implications for custom operators

fbusato · 2025-11-18T00:06:31Z

cub/cub/device/dispatch/kernels/kernel_histogram.cuh

  CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceHistogramSweepKernel(
-    SampleIteratorT d_samples,
+    _CCCL_GRID_CONSTANT const SampleIteratorT d_samples,
    ::cuda::std::array<int, NumActiveChannels> num_output_bins_wrapper,


why _CCCL_GRID_CONSTANT is skipped on cuda::std::array parameters?

I tried it earlier, but when i tested the benches I ran into errors regarding AgentHistogram since the .data() function in it is expected to return type CounterT** but with _CCCL_GRID_CONSTANT const seems to return something else. Please let me know if I've missed something

I didn't check the code. It is a bit unexpected that we modify internal cuda::std::array pointers. Maybe @bernhardmgruber knows more about this point.

Excluding pointers, other ::cuda::std::array parameters should work.

I assume this had historical reasons, because we could not change AgentHistogram to take pointers by const before CCCL 3.0, where we moved them to a detail namespace.

Please add _CCCL_GRID_CONSTANT const to the ::cuda::std::array for histogram and propagatge const throughout the agent.

I have added _CCCL_GRID_CONSTANT const to the ::cuda::std::array, just had to make the necessary AgentHistogram members const(I marked them as read-only) and also make some kernel_histogram.cuh methods in Transform which are utilized as const

fbusato · 2025-11-18T00:09:07Z

cub/cub/device/dispatch/kernels/kernel_merge_sort.cuh

    ValueT* items_pong,
-    CompareOpT compare_op,
+    _CCCL_GRID_CONSTANT const CompareOpT compare_op,
    OffsetT* merge_partitions,


is it a write memory location? or _CCCL_GRID_CONSTANT was not applied to pointers

From what i understand the _CCCL_GRID_CONSTANT const was applied to the outer pointers, not the inner pointers which actually point to the write memory.

yes, this is the reason I asked. It doesn't matter if the pointer is for read or write

cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh

fbusato · 2025-11-18T17:55:46Z

cub/cub/device/dispatch/kernels/kernel_histogram.cuh

  CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceHistogramSweepKernel(
-    SampleIteratorT d_samples,
+    _CCCL_GRID_CONSTANT const SampleIteratorT d_samples,
    ::cuda::std::array<int, NumActiveChannels> num_output_bins_wrapper,


I didn't check the code. It is a bit unexpected that we modify internal cuda::std::array pointers. Maybe @bernhardmgruber knows more about this point.

Excluding pointers, other ::cuda::std::array parameters should work.

fbusato · 2025-11-18T17:57:16Z

cub/cub/device/dispatch/kernels/kernel_merge_sort.cuh

    KeyT* tmp_keys_out,
    ValueT* tmp_items_out,
-    CompareOpT compare_op,
+    _CCCL_GRID_CONSTANT const CompareOpT compare_op,


side effects here are extremely rare, but I would remove _CCCL_GRID_CONSTANT from compare_op

cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh

Fixes

Fixes More Fixes

toxicteddy00077 · 2025-11-23T07:50:50Z

Apologies for the delay. I will be making changes and adding _CCCL_GRID_CONSTANT const to ::cuda::std::array for histogram and make sure it propogates const

fbusato · 2025-11-24T17:37:54Z

cub/cub/device/dispatch/kernels/kernel_histogram.cuh

@@ -445,17 +445,17 @@ template <typename ChainedPolicyT,
          typename OffsetT>
 __launch_bounds__(int(ChainedPolicyT::ActivePolicy::AgentHistogramPolicyT::BLOCK_THREADS))
  CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceHistogramSweepKernel(


the previous kernel DeviceHistogramInitKernel misses _CCCL_GRID_CONSTANT

made corrections

cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh

fbusato · 2025-11-24T17:42:45Z

cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh

    AtomicOffsetT* d_ctrs,
    OffsetT* d_bins_out,
-    const OffsetT* d_bins_in,
+    _CCCL_GRID_CONSTANT const OffsetT* const d_bins_in,


DeviceRadixSortExclusiveSumKernel (below) is missing

made corrections

cub/cub/device/dispatch/kernels/kernel_reduce.cuh

cub/cub/device/dispatch/kernels/kernel_scan.cuh

cub/cub/device/dispatch/kernels/kernel_merge_sort.cuh

cub/cub/device/dispatch/kernels/kernel_transform.cuh

toxicteddy00077 · 2026-02-22T14:01:26Z

done @bernhardmgruber

bernhardmgruber · 2026-02-22T14:27:24Z

/ok to test 53dbe0b

bernhardmgruber · 2026-02-25T09:32:42Z

Some CI jobs are getting stuck compiling now. I have seen this previously with the warpspeed scan, where adding _CCCL_GRID_CONSTANT to one specific kernel parameter caused nvcc < 12.8 to get stuck. See:

cccl/cub/cub/device/dispatch/kernels/kernel_scan.cuh

Lines 187 to 190 in 94bd6e4

    
           // nvcc 12.0 gets stuck compiling some TUs like `cub.bench.scan.exclusive.sum.base`, so only enable for newer versions 
        
           #if _CCCL_CUDACC_AT_LEAST(12, 8) 
        
               _CCCL_GRID_CONSTANT 
        
           #endif // _CCCL_CUDACC_AT_LEAST(12, 8)

@fbusato what do you want to do? Use _CCCL_GRID_CONSTANT only for nvcc >= 12.8 or start finding out which line(s) of this PR cause an issue?

fbusato · 2026-02-25T16:52:22Z

thanks, @bernhardmgruber. This is very helpful to address the problem. Yes, I strongly suggest to enable it only for nvcc >= 12.8. I will track the potential bug offline and chat with the compiler team.

toxicteddy00077 · 2026-03-11T13:40:54Z

@fbusato any update on this? thanks

fbusato · 2026-03-11T16:18:37Z

@toxicteddy00077 sorry again, this not a known issue. Please go ahead and enable _CCCL_GRID_CONSTANT only for nvcc >= 12.8

fbusato · 2026-03-11T18:00:52Z

/ok to test fb8b5cb

toxicteddy00077 · 2026-03-12T05:11:13Z

@fbusato there's still one test failing, what could be the reason? Anything i can fix?

fbusato · 2026-03-12T16:29:19Z

this is unrelated to your changes. It has been fixed quite recently. Let me try to rebase the PR.

fbusato · 2026-03-12T16:30:02Z

/ok to test 4a7004e

github-actions · 2026-03-12T19:30:47Z

😬 CI Workflow Results

🟥 Finished in 2h 58m: Pass: 99%/298 | Total: 12d 04h | Max: 2h 57m | Hits: 69%/372635

See results here.

fbusato · 2026-03-14T00:01:19Z

Python nvcc GCC / SQ / [CTK13.1 GCC13 py3.13] Test cuda.compute(amd64, L4)

AssertionError: LDL instruction found in SASS

@NaderAlAwar could you please check?

NaderAlAwar · 2026-03-16T20:27:51Z

@fbusato Looked into this. The LDL/STL instructions appear even with NVCC, this is not specific to cuda.compute. I ran some benchmarks and the performance after the changes in this PR are 0-2% better, so this is not an issue. I will open a PR to disable the checks on the python side. I think we should disable all the LDL/STL checks until I spend some more time looking into #7978.

NaderAlAwar · 2026-03-16T23:32:59Z

Disabled sass checks in #8053. Once it lands, you should merge from main to fix CI

toxicteddy00077 added 3 commits November 15, 2025 18:39

FEA: Added const _CCCL_GRID_CONSTANT to non-mutable kernel parameters

b7f0800

FEA: Added _CCCL_GRID_CONSTANT const to more kernels

6f71b09

FEA: Added const _CCCL_GRID_CONSTANT to non-mutable kernel parameters

ffd18c1

toxicteddy00077 requested a review from a team as a code owner November 17, 2025 17:00

toxicteddy00077 requested a review from davebayer November 17, 2025 17:00

github-project-automation bot added this to CCCL Nov 17, 2025

github-project-automation bot moved this to Todo in CCCL Nov 17, 2025

cccl-authenticator-app bot moved this from Todo to In Review in CCCL Nov 17, 2025

Merge branch 'main' of https://github.com/toxicteddy00077/cccl into f…

bf617f5

…ea/6482-enable-grid-constant Merge upstream

toxicteddy00077 force-pushed the fea/6482-enable-grid-constant branch from f690f5e to bf617f5 Compare November 17, 2025 17:16

This comment has been minimized.

Sign in to view

fbusato reviewed Nov 18, 2025

View reviewed changes

cub/cub/device/dispatch/kernels/kernel_for_each.cuh Outdated Show resolved Hide resolved

fbusato requested changes Nov 18, 2025

View reviewed changes

github-project-automation bot moved this from In Review to In Progress in CCCL Nov 18, 2025

Removed _CCCL_GRID_CONSTANT const from all operators

413d9bb

toxicteddy00077 requested a review from fbusato November 18, 2025 11:40

fbusato requested changes Nov 18, 2025

View reviewed changes

toxicteddy00077 and others added 6 commits November 20, 2025 11:54

Merge branch 'main' into fea/6482-enable-grid-constant

f8ac0ed

Merge branch 'NVIDIA:main' into fea/6482-enable-grid-constant

5d51587

FEA: Added const _CCCL_GRID_CONSTANT to non-mutable kernel parameters

5b62360

Fixes

Removed _CCCL_GRID_CONSTANT const from all operators

3240916

Fixes

Fixes

5a88802

Fixes More Fixes

More Fixes

bd65528

fbusato requested changes Nov 24, 2025

View reviewed changes

Minor Fixes

0e50ba3

toxicteddy00077 force-pushed the fea/6482-enable-grid-constant branch from 774eb51 to 0e50ba3 Compare November 26, 2025 15:04

Merge branch 'main' into fea/6482-enable-grid-constant

53dbe0b

toxicteddy00077 force-pushed the fea/6482-enable-grid-constant branch from bfbcdf4 to 53dbe0b Compare February 22, 2026 13:59

This comment has been minimized.

Sign in to view

Merge branch 'main' into fea/6482-enable-grid-constant

0835e6d

toxicteddy00077 force-pushed the fea/6482-enable-grid-constant branch from 9ffb226 to e384d92 Compare March 11, 2026 11:18

Merge branch 'main' into fea/6482-enable-grid-constant

fb8b5cb

toxicteddy00077 force-pushed the fea/6482-enable-grid-constant branch from e384d92 to fb8b5cb Compare March 11, 2026 16:43

toxicteddy00077 requested a review from a team as a code owner March 11, 2026 16:43

toxicteddy00077 requested a review from fbusato March 11, 2026 17:32

This comment has been minimized.

Sign in to view

Merge branch 'main' into fea/6482-enable-grid-constant

4a7004e

fbusato approved these changes Mar 12, 2026

View reviewed changes

NaderAlAwar mentioned this pull request Mar 16, 2026

[cuda.compute]: Disable sass checks for float16 merge sort #8053

Open

2 tasks

Conversation

toxicteddy00077 commented Nov 17, 2025

Description

Uh oh!

copy-pr-bot bot commented Nov 17, 2025

Uh oh!

bernhardmgruber commented Nov 17, 2025

Uh oh!

bernhardmgruber commented Nov 17, 2025

Uh oh!

This comment has been minimized.

Uh oh!

fbusato left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

fbusato Nov 18, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

fbusato Nov 18, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

toxicteddy00077 commented Nov 23, 2025

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

toxicteddy00077 commented Feb 22, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

bernhardmgruber commented Feb 22, 2026

Uh oh!

This comment has been minimized.

This comment has been minimized.

This comment has been minimized.

This comment has been minimized.

bernhardmgruber commented Feb 25, 2026

Uh oh!

fbusato commented Feb 25, 2026

Uh oh!

toxicteddy00077 commented Mar 11, 2026

Uh oh!

fbusato commented Mar 11, 2026

Uh oh!

fbusato commented Mar 11, 2026

Uh oh!

This comment has been minimized.

toxicteddy00077 commented Mar 12, 2026

Uh oh!

fbusato Nov 18, 2025 •

edited

Loading

fbusato Nov 18, 2025 •

edited

Loading

toxicteddy00077 commented Feb 22, 2026 •

edited

Loading