[BUG]: ptxas error: Entry function with max regcount of 168 calls function with regcount of 254 on sm_120 (Blackwell)

[zkhatami]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Compile-time Error

Component

CUB

Describe the bug

For this test:

#include <execution>
#include <numeric>
#include <vector>

int32_t
main()
{
  std::vector<double> in(1000);
  std::vector<double> out(1000);
  auto orr = std::inclusive_scan(std::execution::par, in.begin(), in.end(), out.begin());
  return 0;
}

We’re seeing a ptxas register mismatch error when compiling for sm_120 (Blackwell) that didn’t occur on earlier architectures (e.g. sm_80).
This appears to be related to changes introduced in commit 94bd6e4
specifically the new get_device_scan_launch_bounds() function and its warpspeed path.

ptxas error   : Entry function '_ZN3cub22_V_300400_SM_120_NVHPC6detail4scan16DeviceScanKernelINS2_10policy_hubIdddmN4cuda3std3__44plusIvEEE10Policy1000EPdSC_NS0_13ScanTileStateIdLb1EEES9_NS0_8NullTypeEmdLb0ESF_EEvT0_T1_NS2_23tile_state_kernel_arg_tIT2_T6_EEiT3_T4_T5_i' with max regcount of 168 calls function '_ZN3cub22_V_300400_SM_120_NVHPC6detail4scan10kernelBodyINS2_10policy_hubIdddmN4cuda3std3__44plusIvEEE10Policy100015WarpspeedPolicyEdddS9_NS0_8NullTypeELb0EEEvNS1_9warpspeed5SquadENSE_16SpecialRegistersERKNS2_16scanKernelParamsIT0_T1_T2_EET3_T4_' with regcount of 254

A new call get_device_scan_launch_bounds makes 352 number of threads for sm_120 and thus 168 max regcount per thread. The warpspeed path leads to a different policy selection and higher thread count on sm_120, lowering the available registers per thread and triggering the error.
I found that a similar issue was previously discussed and addressed here: #902.

How to Reproduce

nvc++ -stdpar -fast --c++17  scan.cpp

Expected behavior

Test should compile fine after the fix.

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response

[github-actions]

Hi @zkhatami!

Thanks for submitting this issue - the CCCL team has been notified and we'll get back to you as soon as we can!
In the mean time, feel free to add any relevant information to this issue.

[miscco]

@bernhardmgruber I believe you tracked this down during the warpspeed development

[ahendriksen]

All functions need to be inlined to avoid this behavior.

If you compile with -opt-info inline, nvcc logs functions that are not inlined and the reason.

Looking at the error, either of these functions must be marked force_inline:

Entry function void cub::_V_300400_SM_120_NVHPC::detail::scan::DeviceScanKernel<cub::_V_300400_SM_120_NVHPC::detail::scan::policy_hub<double, double, double, unsigned long, cuda::std::__4::plus<void>>::Policy1000, double *, double *, cub::_V_300400_SM_120_NVHPC::ScanTileState<double, (bool)1>, cuda::std::__4::plus<void>, cub::_V_300400_SM_120_NVHPC::NullType, unsigned long, double, (bool)0, cub::_V_300400_SM_120_NVHPC::NullType>(T2, T3, cub::_V_300400_SM_120_NVHPC::detail::scan::tile_state_kernel_arg_t<T4, T8>, int, T5, T6, T7, int) with max regcount of 168 calls function void cub::_V_300400_SM_120_NVHPC::detail::scan::kernelBody<cub::_V_300400_SM_120_NVHPC::detail::scan::policy_hub<double, double, double, unsigned long, cuda::std::__4::plus<void>>::Policy1000::WarpspeedPolicy, double, double, double, cuda::std::__4::plus<void>, cub::_V_300400_SM_120_NVHPC::NullType, (bool)0>(cub::_V_300400_SM_120_NVHPC::detail::warpspeed::Squad, cub::_V_300400_SM_120_NVHPC::detail::warpspeed::SpecialRegisters, const cub::_V_300400_SM_120_NVHPC::detail::scan::scanKernelParams<T2, T3, T4> &, T5, T6) with regcount of 254

[dkolsen-pgi]

All functions need to be inlined to avoid this behavior.

This problem is happening with nvc++, not nvcc. With nvc++, RDC (a.k.a. separate compilation) is enabled by default. With RDC it is not possible to inline all functions.

From previous discussions, it was concluded (I think) that the __launch_bounds__ attribute could not be used reliably when RDC was enabled. The CUDA tools don't have a way to make that always work.

For same context, see

• Wrap launch bounds cub#570
• New indirection level for launch bounds #902
• NVBug 3720703
• NVBug 3384810

[bernhardmgruber]

Given the above context from @dkolsen-pgi, I don't see how we can fix this properly. We can disable the warpspeed kernel for NVHPC as a quick fix. The NVBugs above also suggest that a different tuning could lead to nvvm choose different register allocations which may "solve" the issue as well, but would be extremely brittle.

I am somewhat surprised this issue isn't surfacing more often with other kernels. I almost believe there could be some hidden mechanics that could allow us to massage nvvm to respect the launch bounds at the kernel. After all, almost every kernel in CUB has launch bounds and calls other device functions. And it has worked out everywhere else for years (with some hick-ups). So if anybody knows a good workaround, please reach out to me. We may also just try to inline the code ourselves in some places (by literally copying the function bodies into their call-sites).

[bernhardmgruber]

I have tried to reproduce this now in a git clone of commit 402dc0a but I get a different error:

$ /opt/nvidia/hpc_sdk/Linux_x86_64/2026/compilers/bin/nvc++ -stdpar -fast --c++17 -Icub -Ithrust -Ilibcudacxx/include  scan.cpp
DISubprogram contains null entry in `elements` field
!1834 = !DICompositeType(tag: DW_TAG_enumeration_type, name: "_ZN3cub25_V_300400_SM_86_120_NVHPC6detail9warpspeed6StagesE", file: !1802, size: 32, align: 32, elements: !1835)
DISubprogram contains null entry in `elements` field
!1836 = !DICompositeType(tag: DW_TAG_enumeration_type, name: "_ZN3cub25_V_300400_SM_86_120_NVHPC6detail9warpspeed5ElemsE", file: !1802, size: 32, align: 32, elements: !1835)
DISubprogram contains null entry in `elements` field
!2811 = !DICompositeType(tag: DW_TAG_enumeration_type, name: "_ZN4cuda3std4byteE", file: !1904, size: 8, align: 8, elements: !1835)
DISubprogram contains null entry in `elements` field
!3397 = !DICompositeType(tag: DW_TAG_enumeration_type, name: "_ZN4cuda3std4byteE", file: !3393, size: 8, align: 8, elements: !1835)
DISubprogram contains null entry in `elements` field
!3412 = !DICompositeType(tag: DW_TAG_enumeration_type, name: "_ZN4cuda3std4byteE", file: !3375, size: 8, align: 8, elements: !1835)
DISubprogram contains null entry in `elements` field
!3514 = !DICompositeType(tag: DW_TAG_enumeration_type, name: "_ZN3cub25_V_300400_SM_86_120_NVHPC6detail9warpspeed5ElemsE", file: !3511, size: 32, align: 32, elements: !1835)
DISubprogram contains null entry in `elements` field
!3521 = !DICompositeType(tag: DW_TAG_enumeration_type, name: "_ZN3cub25_V_300400_SM_86_120_NVHPC6detail9warpspeed6StagesE", file: !3511, size: 32, align: 32, elements: !1835)
DISubprogram contains null entry in `elements` field
!3529 = !DICompositeType(tag: DW_TAG_enumeration_type, name: "_ZN4cuda3std4byteE", file: !3489, size: 8, align: 8, elements: !1835)
DISubprogram contains null entry in `elements` field
!4449 = !DICompositeType(tag: DW_TAG_enumeration_type, name: "_ZN4cuda3std4byteE", file: !4442, size: 8, align: 8, elements: !1835)
nvlink error   : Undefined reference to '__nv_cluster_barrier_wait_impl' in '/tmp/nvc++ssMQ2RyOZ5O6.o'
nvlink error   : Undefined reference to '__nv_cluster_barrier_arrive_relaxed_impl' in '/tmp/nvc++ssMQ2RyOZ5O6.o'
pgacclnk: child process exit status 2: /opt/nvidia/hpc_sdk/Linux_x86_64/26.1/compilers/bin/tools/nvdd

Version:

$ /opt/nvidia/hpc_sdk/Linux_x86_64/2026/compilers/bin/nvc++ --version

nvc++ 26.1-0 64-bit target on x86-64 Linux -tp znver3 
NVIDIA Compilers and Tools
Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES.  All rights reserved.

Which compiler version did your use?

[ahendriksen]

If nvc++ doesn't support force inlining, then we should disable the warpspeed kernel on nvc++. I don't see how this can be fixed. For performance, the framework critically depends on inlining.

In this case, because the kernel (by design) uses such a large number of registers, we are also seeing compilation issues.

[bernhardmgruber]

I am currently investigating whether a reduction in reduce and scan warps can get us down to 256 threads per block, which can always launch.

[ahendriksen]

I would be surprised if that still has a performance advantage over the existing decoupled lookback implementation. Happy to be proven wrong :)