[BUG]: `cub::DeviceScan::ExclusiveScan` and `InclusiveScan` produce incorrect results with custom binary operator

[shrshi]

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

Runtime Error

Component

CUB

Describe the bug

Environment

• GPU: NVIDIA RTX PRO 6000 Blackwell Server Edition (Compute Capability 12.0)
• Driver: 580.95.05
• CUDA: 13.1 (V13.1.115)
• CCCL Version: 3.4.0 (CCCL_VERSION 3004000)
• OS: Ubuntu 24.04.4 LTS

Description

cub::DeviceScan::ExclusiveScan produces incorrect results when using a custom binary operator that propagates values conditionally.

Additional Notes

• cub::DeviceScan::InclusiveScan exhibits the same bug at the same positions

How to Reproduce

Minimal Reproducible Example

#include <cub/cub.cuh>
#include <cuda_runtime.h>
#include <numeric>
#include <vector>
#include <cstdio>

// Custom operator that propagates the last "write" value through "read" values
template <typename StackSymbolT>
struct PropagateLastWrite {
  __host__ __device__ StackSymbolT operator()(StackSymbolT const& lhs,
                                               StackSymbolT const& rhs) const
  {
    bool is_rhs_read  = (rhs == read_symbol);
    bool is_lhs_write = (lhs != read_symbol);
    return (is_rhs_read && is_lhs_write) ? lhs : rhs;
  }
  StackSymbolT read_symbol;
};

int main()
{
  constexpr size_t num_elements     = 8160;
  constexpr char read_symbol        = 'x';
  constexpr char empty_stack_symbol = '_';

  // Create input - mostly read_symbol with scattered write values
  std::vector<char> h_input(num_elements, read_symbol);

  // Key write values that trigger the bug
  h_input[8073] = '[';
  h_input[8074] = '{';
  h_input[8075] = '[';
  h_input[8076] = '{';  // Last '{' before positions 8077-8147 (all 'x')
  h_input[8148] = '_';
  h_input[8151] = '{';
  h_input[8152] = '_';
  h_input[8154] = '[';
  h_input[8155] = '_';
  h_input[8157] = '[';
  h_input[8159] = '_';

  // Earlier scattered values
  h_input[8020] = '_';
  h_input[8023] = '{';
  h_input[8057] = '[';
  h_input[8060] = '{';
  h_input[8061] = '[';
  h_input[8068] = '{';

  // Compute CPU reference
  std::vector<char> h_expected(num_elements);
  std::exclusive_scan(h_input.begin(),
                      h_input.end(),
                      h_expected.begin(),
                      empty_stack_symbol,
                      PropagateLastWrite<char>{read_symbol});

  // Allocate device memory
  char* d_input;
  char* d_output;
  cudaMalloc(&d_input, num_elements);
  cudaMalloc(&d_output, num_elements);
  cudaMemcpy(d_input, h_input.data(), num_elements, cudaMemcpyHostToDevice);

  // Get temp storage size
  size_t temp_storage_bytes = 0;
  PropagateLastWrite<char> op{read_symbol};

  cub::DeviceScan::ExclusiveScan(nullptr, temp_storage_bytes,
                                  d_input, d_output, op,
                                  empty_stack_symbol, num_elements);

  char* d_temp;
  cudaMalloc(&d_temp, temp_storage_bytes);

  // Run GPU scan
  cub::DeviceScan::ExclusiveScan(d_temp, temp_storage_bytes,
                                  d_input, d_output, op,
                                  empty_stack_symbol, num_elements);

  // Copy back and compare
  std::vector<char> h_output(num_elements);
  cudaMemcpy(h_output.data(), d_output, num_elements, cudaMemcpyDeviceToHost);

  int mismatches = 0;
  for (size_t i = 0; i < num_elements; i++) {
    if (h_output[i] != h_expected[i]) {
      mismatches++;
    }
  }
  printf("Total mismatches: %d\n", mismatches);

  cudaFree(d_input);
  cudaFree(d_output);
  cudaFree(d_temp);

  return mismatches > 0 ? 1 : 0;
}

Actual Output on Blackwell

Mismatch at 8127: GPU='[' (0x5B), CPU='{' (0x7B)
Mismatch at 8128: GPU='[' (0x5B), CPU='{' (0x7B)
...
Mismatch at 8148: GPU='[' (0x5B), CPU='{' (0x7B)
Total mismatches: 22

Analysis

• Position 8076 contains '{' (the last write before a gap)
• Positions 8077-8147 contain 'x' (read_symbol)
• The exclusive scan should propagate '{' through all read positions
• Instead, GPU produces '[' (from position 8075) starting at position 8127

Expected behavior

Expected Output

All positions should match (0 mismatches).

[github-actions]

Hi @shrshi!

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.

[jrhemstad]

@bernhardmgruber I suspect this is the same issue as #8134 ?

[davidwendt]

For what it's worth, compute-sanitizer memcheck does not report any errors.
The position of the first error is at 8127 which is 64 from 8191 (0x1FFF) if that helps.

[davidwendt]

But compute-sanitizer --tool synccheck does report an error

========= COMPUTE-SANITIZER
========= Barrier error detected. Divergent thread(s) in block.
=========     at void cub::_V_300400_SM_1200::detail::scan::DeviceScanKernel<cub::_V_300400_SM_1200::detail::scan::policy_selector_from_types<char, char, char, unsigned long long, PropagateLastWrite<char>>, char *, char *, cub::_V_300400_SM_1200::ScanTileState<char, (bool)1>, PropagateLastWrite<char>, cub::_V_300400_SM_1200::detail::InputValue<char, char *>, unsigned long long, char, (bool)0, char>(T2, T3, cub::_V_300400_SM_1200::detail::scan::tile_state_kernel_arg_t<T4, T8>, int, T5, T6, T7, int)+0x6040
=========     by thread (192,0,0) in block (0,0,0)
=========     Saved host backtrace up to driver entry point at kernel launch time
=========         Host Frame: cudaLaunchKernelExC [0x8b73a] in exclscan
=========         Host Frame: cudaLaunchKernelEx<char*, char*, cub::_V_300400_SM_1200::detail::scan::tile_state_kernel_arg_t<cub::_V_300400_SM_1200::ScanTileState<char, true>, char>, int, PropagateLastWrite<char>, cub::_V_300400_SM_1200::detail::InputValue<char, char*>, unsigned long long, int, char* const&, char* const&, cub::_V_300400_SM_1200::detail::scan::tile_state_kernel_arg_t<cub::_V_300400_SM_1200::ScanTileState<char, true>, char> const&, int const&, PropagateLastWrite<char> const&, cub::_V_300400_SM_1200::detail::InputValue<char, char*> const&, unsigned long long const&, int const&>(cudaLaunchConfig_st const*, void (*)(char*, char*, cub::_V_300400_SM_1200::detail::scan::tile_state_kernel_arg_t<cub::_V_300400_SM_1200::ScanTileState<char, true>, char>, int, PropagateLastWrite<char>, cub::_V_300400_SM_1200::detail::InputValue<char, char*>, unsigned long long, int), char* const&, char* const&, cub::_V_300400_SM_1200::detail::scan::tile_state_kernel_arg_t<cub::_V_300400_SM_1200::ScanTileState<char, true>, char> const&, int const&, PropagateLastWrite<char> const&, cub::_V_300400_SM_1200::detail::InputValue<char, char*> const&, unsigned long long const&, int const&)::{lambda(char*, char*, cub::_V_300400_SM_1200::detail::scan::tile_state_kernel_arg_t<cub::_V_300400_SM_1200::ScanTileState<char, true>, char>, int, PropagateLastWrite<char>, cub::_V_300400_SM_1200::detail::InputValue<char, char*>, unsigned long long, int)#1}::operator()(char*, char*, cub::_V_300400_SM_1200::detail::scan::tile_state_kernel_arg_t<cub::_V_300400_SM_1200::ScanTileState<char, true>, char>, int, PropagateLastWrite<char>, cub::_V_300400_SM_1200::detail::InputValue<char, char*>, unsigned long long, int) const [0x19095] in exclscan
=========         Host Frame: cudaError cudaLaunchKernelEx<char*, char*, cub::_V_300400_SM_1200::detail::scan::tile_state_kernel_arg_t<cub::_V_300400_SM_1200::ScanTileState<char, true>, char>, int, PropagateLastWrite<char>, cub::_V_300400_SM_1200::detail::InputValue<char, char*>, unsigned long long, int, char* const&, char* const&, cub::_V_300400_SM_1200::detail::scan::tile_state_kernel_arg_t<cub::_V_300400_SM_1200::ScanTileState<char, true>, char> const&, int const&, PropagateLastWrite<char> const&, cub::_V_300400_SM_1200::detail::InputValue<char, char*> const&, unsigned long long const&, int const&>(cudaLaunchConfig_st const*, void (*)(char*, char*, cub::_V_300400_SM_1200::detail::scan::tile_state_kernel_arg_t<cub::_V_300400_SM_1200::ScanTileState<char, true>, char>, int, PropagateLastWrite<char>, cub::_V_300400_SM_1200::detail::InputValue<char, char*>, unsigned long long, int), char* const&, char* const&, cub::_V_300400_SM_1200::detail::scan::tile_state_kernel_arg_t<cub::_V_300400_SM_1200::ScanTileState<char, true>, char> const&, int const&, PropagateLastWrite<char> const&, cub::_V_300400_SM_1200::detail::InputValue<char, char*> const&, unsigned long long const&, int const&) [0x191c8] in exclscan
=========         Host Frame: cudaError thrust::_V_300400_SM_1200::cuda_cub::detail::triple_chevron::doit_host<void (*)(char*, char*, cub::_V_300400_SM_1200::detail::scan::tile_state_kernel_arg_t<cub::_V_300400_SM_1200::ScanTileState<char, true>, char>, int, PropagateLastWrite<char>, cub::_V_300400_SM_1200::detail::InputValue<char, char*>, unsigned long long, int), char*, char*, cub::_V_300400_SM_1200::detail::scan::tile_state_kernel_arg_t<cub::_V_300400_SM_1200::ScanTileState<char, true>, char>, int, PropagateLastWrite<char>, cub::_V_300400_SM_1200::detail::InputValue<char, char*>, unsigned long long, int>(void (*)(char*, char*, cub::_V_300400_SM_1200::detail::scan::tile_state_kernel_arg_t<cub::_V_300400_SM_1200::ScanTileState<char, true>, char>, int, PropagateLastWrite<char>, cub::_V_300400_SM_1200::detail::InputValue<char, char*>, unsigned long long, int), char* const&, char* const&, cub::_V_300400_SM_1200::detail::scan::tile_state_kernel_arg_t<cub::_V_300400_SM_1200::ScanTileState<char, true>, char> const&, int const&, PropagateLastWrite<char> const&, cub::_V_300400_SM_1200::detail::InputValue<char, char*> const&, unsigned long long const&, int const&) const [0x187be] in exclscan
=========         Host Frame: auto cub::_V_300400_SM_1200::detail::scan::dispatch<(cub::_V_300400_SM_1200::ForceInclusive)1, char*, char*, PropagateLastWrite<char>, cub::_V_300400_SM_1200::detail::InputValue<char, char*>, unsigned long long, char, cub::_V_300400_SM_1200::detail::scan::policy_selector_from_types<char, char, char, unsigned long long, PropagateLastWrite<char> >, cub::_V_300400_SM_1200::detail::scan::DeviceScanKernelSource<cub::_V_300400_SM_1200::detail::scan::policy_selector_from_types<char, char, char, unsigned long long, PropagateLastWrite<char> >, char*, char*, PropagateLastWrite<char>, cub::_V_300400_SM_1200::detail::InputValue<char, char*>, unsigned long long, char, (cub::_V_300400_SM_1200::ForceInclusive)1>, cub::_V_300400_SM_1200::detail::TripleChevronFactory>(void*, unsigned long&, char*, char*, PropagateLastWrite<char>, cub::_V_300400_SM_1200::detail::InputValue<char, char*>, unsigned long long, CUstream_st*, cub::_V_300400_SM_1200::detail::scan::policy_selector_from_types<char, char, char, unsigned long long, PropagateLastWrite<char> >, cub::_V_300400_SM_1200::detail::scan::DeviceScanKernelSource<cub::_V_300400_SM_1200::detail::scan::policy_selector_from_types<char, char, char, unsigned long long, PropagateLastWrite<char> >, char*, char*, PropagateLastWrite<char>, cub::_V_300400_SM_1200::detail::InputValue<char, char*>, unsigned long long, char, (cub::_V_300400_SM_1200::ForceInclusive)1>, cub::_V_300400_SM_1200::detail::TripleChevronFactory)::{lambda(auto:1)#1}::operator()<cuda::std::__4::integral_constant<cub::_V_300400_SM_1200::detail::scan::scan_policy, cub::_V_300400_SM_1200::detail::scan::scan_policy{128, 30, (cub::_V_300400_SM_1200::BlockLoadAlgorithm)4, (cub::_V_300400_SM_1200::CacheLoadModifier)0, (cub::_V_300400_SM_1200::BlockStoreAlgorithm)4, (cub::_V_300400_SM_1200::BlockScanAlgorithm)2, cub::_V_300400_SM_1200::detail::delay_constructor_policy{(cub::_V_300400_SM_1200::detail::delay_constructor_kind)1, 350u, 450u}, cub::_V_300400_SM_1200::detail::scan::scan_warpspeed_policy{true, 4, 4, 1, 1, 1, 4, 352, 63, 8064}}> >(cuda::std::__4::integral_constant<cub::_V_300400_SM_1200::detail::scan::scan_policy, cub::_V_300400_SM_1200::detail::scan::scan_policy{128, 30, (cub::_V_300400_SM_1200::BlockLoadAlgorithm)4, (cub::_V_300400_SM_1200::CacheLoadModifier)0, (cub::_V_300400_SM_1200::BlockStoreAlgorithm)4, (cub::_V_300400_SM_1200::BlockScanAlgorithm)2, cub::_V_300400_SM_1200::detail::delay_constructor_policy{(cub::_V_300400_SM_1200::detail::delay_constructor_kind)1, 350u, 450u}, cub::_V_300400_SM_1200::detail::scan::scan_warpspeed_policy{true, 4, 4, 1, 1, 1, 4, 352, 63, 8064}}>) const [0x17215] in exclscan
=========         Host Frame: cudaError cub::_V_300400_SM_1200::DeviceScan::ExclusiveScan<char*, char*, PropagateLastWrite<char>, char, unsigned long>(void*, unsigned long&, char*, char*, PropagateLastWrite<char>, char, unsigned long, CUstream_st*) [0x148b7] in exclscan
=========         Host Frame: main [0xd335] in exclscan
...

compute-sanitizer --tool racecheck does not report any errors.