[BUG]: CUB BLOCK_REDUCE_WARP_REDUCTIONS_NONDETERMINISTIC specialization doesn't support custom operators #11647

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

Issue found during the investigation in

• Fix [QA Agent] Cooperative reduction - Block #11646

BLOCK_REDUCE_WARP_REDUCTIONS_NONDETERMINISTIC fails to compile for a user-defined type. Also, it relies on T{} for the identity.

How to Reproduce

minimal reproducer:

#include <cub/block/block_reduce.cuh>

struct custom_t
{
  int value;
  __host__ __device__ custom_t operator+(custom_t other) const
  {
    return {value + other.value};
  }
};

__global__ void kernel(custom_t* data)
{
  using block_reduce_t = cub::BlockReduce<custom_t, 64, cub::BLOCK_REDUCE_WARP_REDUCTIONS_NONDETERMINISTIC>;
  __shared__ block_reduce_t::TempStorage storage;
  const custom_t aggregate = block_reduce_t(storage).Sum(data[threadIdx.x]);
  if (threadIdx.x == 0)
  {
    data[0] = aggregate;
  }
}

Expected behavior

no compile error and the correct result of the reduction

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response