[libcu++] Move `<cuda/fpmeow>` headers to `<cuda/experimental/fpmeow>` by davebayer · Pull Request #11218 · NVIDIA/cccl

davebayer · 2026-09-04T15:54:04Z

Since these are purely experimental headers, I think it would be better to move the to cuda/experimental directory. The name of the headers is not final and changing it in the future might be surprising for some users.

coderabbitai · 2026-09-04T16:07:02Z

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: eb9c6649-d58a-4cc3-8d41-adf1492b95ba

📥 Commits

Reviewing files that changed from the base of the PR and between 13462e2 and d11e45b.

📒 Files selected for processing (62)

libcudacxx/include/cuda/__fp/fpmp.h
libcudacxx/include/cuda/__fp/fpmp_impl.h
libcudacxx/include/cuda/__fp/fpmp_math.h
libcudacxx/include/cuda/__fp/fptool_custom.h
libcudacxx/include/cuda/__fp/fptool_stat.h
libcudacxx/include/cuda/__fp/fptool_stat_math.h
libcudacxx/include/cuda/experimental/fpemu
libcudacxx/include/cuda/experimental/fpmp
libcudacxx/include/cuda/experimental/fpmp_math
libcudacxx/include/cuda/experimental/fptool
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/accuracy.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/api.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/api_unpacked.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/bit_cast.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/charbool.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/cmp.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/conv.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/copyable.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/core.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/core_unpacked.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/cuda_std.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/div.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/exp.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/float.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/float64.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/from_int.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/mixed.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/pi.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/sqrt.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/template.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/to_int.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/unary.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/volatile.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/accuracy.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/accuracy_conv.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/accuracy_mixed.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/api.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/atomic_dd.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/atomic_ff.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/black_scholes.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/cast.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/charbool.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/copyable.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/cross_prec.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/cuda_std.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/fpemu_interop.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/implicit_conv.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/int.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/limits.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/lut.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/math.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/pi.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/qualified_calls.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/reduce.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/shfl.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/unary.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/volatile.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fptool/custom.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fptool/fpemu_interop.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fptool/qualified_calls.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fptool/stat.pass.cpp
libcudacxx/test/libcudacxx/cuda/fp/units/fptool/unary.pass.cpp

Included review availability: Your plan provides up to 12 included reviews per hour; 10 remain after this review.

📝 Summary

Summary by CodeRabbit

New Features
- Added the experimental cuda/experimental/fpemu header for floating-point emulation functionality.
- Standardized public include paths for FP emulation, FPMP, FPMP math, and FP tools under cuda/experimental.
Documentation
- Updated usage examples, quick-start guidance, and API references to show the current experimental header paths.
Tests
- Updated floating-point emulation, FPMP, and FP tool coverage to use the experimental headers while preserving existing test behavior.

Walkthrough

The change adds cuda/experimental/fpemu and updates fpemu, fpmp, and fptool documentation and tests to use experimental include paths. Test logic and public declarations remain unchanged.

Changes

Experimental floating-point headers

Layer / File(s)	Summary
Header wrapper and documentation paths `libcudacxx/include/cuda/experimental/`, `libcudacxx/include/cuda/__fp/`	Adds the `cuda/experimental/fpemu` wrapper and updates floating-point header documentation paths.
FP emulation test include migration `libcudacxx/test/libcudacxx/cuda/fp/units/fpemu/*`	Replaces legacy `cuda/fpemu` includes with `cuda/experimental/fpemu`.
FPMP test include migration `libcudacxx/test/libcudacxx/cuda/fp/units/fpmp/*`	Replaces legacy fpmp and fpmp math includes with experimental paths.
FP tool test include migration `libcudacxx/test/libcudacxx/cuda/fp/units/fptool/*`	Replaces legacy fptool and fpemu includes with experimental paths.

Merge Risk: ⚪ Minimal · up to d11e4

This change relocates experimental floating-point header usage to cuda/experimental paths and updates related documentation and tests. No remaining merge-readiness risk is identified.

_{Comment @coderabbitai help to get the list of available commands.}

github-actions · 2026-09-04T16:19:12Z

🔬 CUB benchmark SASS comparison

⚠️ The SASS changed for 4 of 84 CUB benchmark target(s). A benchmark run may be necessary

How to request a benchmark run

Request a CUB benchmark run for this PR:

1. Replace the `benchmarks:` block of ci/bench.yaml with exactly this:

benchmarks:
  filters:
    cub:
      - '^cub\.bench\.segmented_topk\.variable\.indexed\.base$'
      - '^cub\.bench\.segmented_topk\.variable\.indexed\.cluster\.base$'
      - '^cub\.bench\.segmented_topk\.variable\.keys\.base$'
      - '^cub\.bench\.segmented_topk\.variable\.keys\.cluster\.base$'
  gpus:
    - "h100"   # pick the GPUs that this change can affect

2. Commit with `[bench-only]` at the end of the commit summary, so that
   the unrelated CI jobs are skipped. Then push.

ci/bench.yaml must match ci/bench.template.yaml before the PR can merge.
Reset it once the measurement is done.

Run	Value
Baseline	`13462e2a7cabda1cf36b739da765a83f8f06fb61`
Tested	`HEAD`
Architectures	`75-real;80-real;90-real;100-real;110-real;120-real;120-virtual`

Targets with a SASS change

Target	Architectures with a SASS change
`cub.bench.segmented_topk.variable.indexed.base`	`sm_120`, `sm_110`, `sm_75`, `sm_90`, `sm_100`, `sm_80`
`cub.bench.segmented_topk.variable.indexed.cluster.base`	`sm_120`, `sm_110`, `sm_75`, `sm_90`, `sm_100`, `sm_80`
`cub.bench.segmented_topk.variable.keys.base`	`sm_120`, `sm_110`, `sm_75`, `sm_90`, `sm_100`, `sm_80`
`cub.bench.segmented_topk.variable.keys.cluster.base`	`sm_120`, `sm_110`, `sm_75`, `sm_90`, `sm_100`, `sm_80`

‼️ Summary of Differences ‼️

Showing 4/4 summaries.

cub.bench.segmented_topk.variable.indexed.base - sm_120

Showing 40/137 diff lines, 30 changes. - ⬇️ Full diff

--- base/cub.bench.segmented_topk.variable.indexed.base.sm_120
+++ test/cub.bench.segmented_topk.variable.indexed.base.sm_120
@@ -137845,7 +137845,7 @@
 STG.E.ENL2.256 desc[UR10][R28.64+0x2000], R20, R24 ;
 EXIT ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)1024, (int)1024>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)1024, (int)1024>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 LDC R1, c[0x0][0x37c] ;
 S2R R7, SR_CTAID.X ;
 LDC.U8 R0, c[0x0][0x38c] ;
@@ -148337,7 +148337,7 @@
 HFMA2 R3, -RZ, RZ, 0, 0 ;
 RET.REL.NODEC R2 0x0 ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)1024, (int)2048>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)1024, (int)2048>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 LDC R1, c[0x0][0x37c] ;
 S2R R7, SR_CTAID.X ;
 LDC.U8 R0, c[0x0][0x38c] ;
@@ -158829,7 +158829,7 @@
 HFMA2 R3, -RZ, RZ, 0, 0 ;
 RET.REL.NODEC R2 0x0 ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)1024, (int)512>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)1024, (int)512>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 LDC R1, c[0x0][0x37c] ;
 S2R R7, SR_CTAID.X ;
 LDC.U8 R0, c[0x0][0x38c] ;
@@ -169321,7 +169321,7 @@
 HFMA2 R3, -RZ, RZ, 0, 0 ;
 RET.REL.NODEC R2 0x0 ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)2048, (int)1024>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)2048, (int)1024>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 LDC R1, c[0x0][0x37c] ;
 S2R R7, SR_CTAID.X ;
 LDC.U8 R0, c[0x0][0x38c] ;
@@ -179813,7 +179813,7 @@
 HFMA2 R3, -RZ, RZ, 0, 0 ;

cub.bench.segmented_topk.variable.indexed.cluster.base - sm_120

Showing 40/137 diff lines, 30 changes. - ⬇️ Full diff

--- base/cub.bench.segmented_topk.variable.indexed.cluster.base.sm_120
+++ test/cub.bench.segmented_topk.variable.indexed.cluster.base.sm_120
@@ -221336,7 +221336,7 @@
 STG.E.ENL2.256 desc[UR10][R28.64+0x2000], R20, R24 ;
 EXIT ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)1024, (int)1024>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)1024, (int)1024>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 LDC R1, c[0x0][0x37c] ;
 S2R R7, SR_CTAID.X ;
 LDC.U8 R0, c[0x0][0x38c] ;
@@ -231828,7 +231828,7 @@
 HFMA2 R3, -RZ, RZ, 0, 0 ;
 RET.REL.NODEC R2 0x0 ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)1024, (int)2048>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)1024, (int)2048>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 LDC R1, c[0x0][0x37c] ;
 S2R R7, SR_CTAID.X ;
 LDC.U8 R0, c[0x0][0x38c] ;
@@ -242320,7 +242320,7 @@
 HFMA2 R3, -RZ, RZ, 0, 0 ;
 RET.REL.NODEC R2 0x0 ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)1024, (int)512>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)1024, (int)512>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 LDC R1, c[0x0][0x37c] ;
 S2R R7, SR_CTAID.X ;
 LDC.U8 R0, c[0x0][0x38c] ;
@@ -252812,7 +252812,7 @@
 HFMA2 R3, -RZ, RZ, 0, 0 ;
 RET.REL.NODEC R2 0x0 ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)2048, (int)1024>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)2048, (int)1024>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 LDC R1, c[0x0][0x37c] ;
 S2R R7, SR_CTAID.X ;
 LDC.U8 R0, c[0x0][0x38c] ;
@@ -263304,7 +263304,7 @@
 HFMA2 R3, -RZ, RZ, 0, 0 ;

cub.bench.segmented_topk.variable.keys.base - sm_120

Showing 40/137 diff lines, 30 changes. - ⬇️ Full diff

--- base/cub.bench.segmented_topk.variable.keys.base.sm_120
+++ test/cub.bench.segmented_topk.variable.keys.base.sm_120
@@ -135370,7 +135370,7 @@
 STG.E.ENL2.256 desc[UR10][R28.64+0x2000], R20, R24 ;
 EXIT ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)1024, (int)1024>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)1024, (int)1024>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 LDC R1, c[0x0][0x37c] ;
 S2R R7, SR_CTAID.X ;
 LDC.U8 R0, c[0x0][0x38c] ;
@@ -145862,7 +145862,7 @@
 HFMA2 R3, -RZ, RZ, 0, 0 ;
 RET.REL.NODEC R2 0x0 ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)1024, (int)2048>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)1024, (int)2048>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 LDC R1, c[0x0][0x37c] ;
 S2R R7, SR_CTAID.X ;
 LDC.U8 R0, c[0x0][0x38c] ;
@@ -156354,7 +156354,7 @@
 HFMA2 R3, -RZ, RZ, 0, 0 ;
 RET.REL.NODEC R2 0x0 ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)1024, (int)512>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)1024, (int)512>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 LDC R1, c[0x0][0x37c] ;
 S2R R7, SR_CTAID.X ;
 LDC.U8 R0, c[0x0][0x38c] ;
@@ -166846,7 +166846,7 @@
 HFMA2 R3, -RZ, RZ, 0, 0 ;
 RET.REL.NODEC R2 0x0 ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)2048, (int)1024>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)2048, (int)1024>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 LDC R1, c[0x0][0x37c] ;
 S2R R7, SR_CTAID.X ;
 LDC.U8 R0, c[0x0][0x38c] ;
@@ -177338,7 +177338,7 @@
 HFMA2 R3, -RZ, RZ, 0, 0 ;

cub.bench.segmented_topk.variable.keys.cluster.base - sm_120

Showing 40/137 diff lines, 30 changes. - ⬇️ Full diff

--- base/cub.bench.segmented_topk.variable.keys.cluster.base.sm_120
+++ test/cub.bench.segmented_topk.variable.keys.cluster.base.sm_120
@@ -135370,7 +135370,7 @@
 STG.E.ENL2.256 desc[UR10][R28.64+0x2000], R20, R24 ;
 EXIT ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)1024, (int)1024>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)1024, (int)1024>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 LDC R1, c[0x0][0x37c] ;
 S2R R7, SR_CTAID.X ;
 LDC.U8 R0, c[0x0][0x38c] ;
@@ -145862,7 +145862,7 @@
 HFMA2 R3, -RZ, RZ, 0, 0 ;
 RET.REL.NODEC R2 0x0 ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)1024, (int)2048>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)1024, (int)2048>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 LDC R1, c[0x0][0x37c] ;
 S2R R7, SR_CTAID.X ;
 LDC.U8 R0, c[0x0][0x38c] ;
@@ -156354,7 +156354,7 @@
 HFMA2 R3, -RZ, RZ, 0, 0 ;
 RET.REL.NODEC R2 0x0 ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)1024, (int)512>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)1024, (int)512>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 LDC R1, c[0x0][0x37c] ;
 S2R R7, SR_CTAID.X ;
 LDC.U8 R0, c[0x0][0x38c] ;
@@ -166846,7 +166846,7 @@
 HFMA2 R3, -RZ, RZ, 0, 0 ;
 RET.REL.NODEC R2 0x0 ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)2048, (int)1024>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)2048, (int)1024>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 LDC R1, c[0x0][0x37c] ;
 S2R R7, SR_CTAID.X ;
 LDC.U8 R0, c[0x0][0x38c] ;
@@ -177338,7 +177338,7 @@
 HFMA2 R3, -RZ, RZ, 0, 0 ;

iburyl · 2026-09-04T16:32:15Z

This feature is already being experimented with by several parties, please, do not move it.

copy-pr-bot · 2026-09-04T16:46:32Z

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

[libcu++] Move <cuda/fpmeow> headers to <cuda/experimental/fpmeow>

d11e45b

davebayer requested a review from akolesov-nvidia September 4, 2026 15:54

davebayer requested a review from a team as a code owner September 4, 2026 15:54

davebayer requested a review from pciolkosz September 4, 2026 15:54

github-project-automation Bot added this to CCCL Sep 4, 2026

github-project-automation Bot moved this to Todo in CCCL Sep 4, 2026

cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Sep 4, 2026

davebayer marked this pull request as draft September 4, 2026 16:46

cccl-authenticator-app Bot moved this from In Review to In Progress in CCCL Sep 4, 2026

Navigation Menu

Uh oh!

[libcu++] Move `<cuda/fpmeow>` headers to `<cuda/experimental/fpmeow>` - #11218

[libcu++] Move `<cuda/fpmeow>` headers to `<cuda/experimental/fpmeow>`#11218
davebayer wants to merge 1 commit into
NVIDIA:mainfrom
davebayer:move_fp_to_experimental

davebayer commented Sep 4, 2026

Uh oh!

coderabbitai Bot commented Sep 4, 2026

Summary by CodeRabbit

Walkthrough

Changes

Uh oh!

github-actions Bot commented Sep 4, 2026

Uh oh!

iburyl commented Sep 4, 2026

Uh oh!

copy-pr-bot Bot commented Sep 4, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

davebayer commented Sep 4, 2026

Uh oh!

coderabbitai Bot commented Sep 4, 2026

Summary by CodeRabbit

Walkthrough

Changes

Uh oh!

github-actions Bot commented Sep 4, 2026

🔬 CUB benchmark SASS comparison

‼️ Summary of Differences ‼️

Uh oh!

iburyl commented Sep 4, 2026

Uh oh!

copy-pr-bot Bot commented Sep 4, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants