NVIDIA / NVIDIA/cccl

[BUG]: `cub.test.block_scan.alg_1.mode_0` is broken for `uchar3`

Open
#3,835 2 comments 0 reactions 0 assignees View on GitHub
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.