[BUG]: `cub.test.block_scan.alg_1.mode_0` is broken for `uchar3`
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
### Is this a duplicate?
- [x] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)
### Type of Bug
Runtime Error
### Component
CUB
### Describe the bug
The test `cub.test.block_scan.alg_1.mode_0` fails for `uchar3`. However, this is not detected on `main` because of a bug in `c2h/generators.h` that causes only zero data to be generated. If this is fixed, the CUB test fails with mismatching results between host (reference) and device paths.
### How to Reproduce
Apply the following diff:
```c++
diff --git a/c2h/include/c2h/test_util_vec.h b/c2h/include/c2h/test_util_vec.h
index f663c49e3..94c93ce25 100644
--- a/c2h/include/c2h/test_util_vec.h
+++ b/c2h/include/c2h/test_util_vec.h
@@ -353,7 +353,25 @@ C2H_VEC_OVERLOAD(double, double)
return retval; \
} \
}; \
- CUB_NAMESPACE_END
+ CUB_NAMESPACE_END \
+ \
+ namespace std \
+ { \
+ template <> \
+ class numeric_limits \
+ { \
+ public: \
+ static constexpr bool is_specialized = true; \
+ static __host__ __device__ T max() \
+ { \
+ return ::cub::NumericTraits::Max(); \
+ } \
+ static __host__ __device__ T lowest() \
+ { \
+ return ::cub::NumericTraits::Lowest(); \
+ } \
+ }; \
+ }
/**
* Vector4 overloads
```
Which makes `c2h::gen(...)` correctly generate test data for `uchar3`. Then build and run `cub.test.block_scan.alg_1.mode_0`. It fails with:
```
-------------------------------------------------------------------------------
"Block scan works with vec types"(2) - types_380 - 1
-------------------------------------------------------------------------------
/home/bgruber/dev/cccl/cub/test/catch2_test_block_scan.cu:380
...............................................................................
/home/bgruber/dev/cccl/cub/test/catch2_test_block_scan.cu:401: FAILED:
REQUIRE( h_out == d_out )
with expansion:
{ (228,141,83), (255,221,58), (100,135,180), (70,69,167), (194,79,12), (221,
48,237), (218,254,217), (247,199,193), (247,169,186), (87,50,26), (165,81,16)
...
```
I also generated a sequence of `(0, 0, 0), (1, 1, 1), ...` as input and can confirm that the host/reference result is correct. The device result is wrong starting at the 672th element, but only the first of the 3 values in `uchar3`. The second and third value are correct.
### Expected behavior
The test should pass
### Reproduction link
_No response_
### Operating System
Ubuntu 24.10
### nvidia-smi output
```
Mon Feb 17 14:48:53 2025
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 570.86.15 Driver Version: 570.86.15 CUDA Version: 12.8 |
|-----------------------------------------+------------------------+----------------------+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+========================+======================|
| 0 NVIDIA RTX A6000 Off | 00000000:41:00.0 On | Off |
| 30% 51C P8 34W / 300W | 1143MiB / 49140MiB | 0% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+
+-----------------------------------------------------------------------------------------+
| Processes: |
| GPU GI CI PID Type Process name GPU Memory |
| ID ID Usage |
|=========================================================================================|
| 0 N/A N/A 5317 G /usr/lib/xorg/Xorg 510MiB |
| 0 N/A N/A 5685 G /usr/bin/gnome-shell 299MiB |
| 0 N/A N/A 7303 G ...ersion=20250214-130114.277000 94MiB |
| 0 N/A N/A 55879 G /usr/bin/gnome-text-editor 15MiB |
+-----------------------------------------------------------------------------------------+
```
### NVCC version
```
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Wed_Jan_15_19:20:09_PST_2025
Cuda compilation tools, release 12.8, V12.8.61
Build cuda_12.8.r12.8/compiler.35404655_0
```
Contributor guide
Assessment
This issue has not been assessed yet.