[BUG]: cub::DeviceReduce::Min (Max too?) returns wrong value for inputs with INF

[leofang]

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

INF should be returned but it is not, see cupy/cupy#8180.

It seems to be the same issue as ArgMin/ArgMax (NVIDIA/cub#642).

How to Reproduce

#include <iostream>
#include <cub/cub.cuh>

constexpr int N = 10;
constexpr float INF = std::numeric_limits<float>::infinity();

int main() {
  float* in;
  float* d_in;
  float* out;
  float* d_out;

  in = new float[N];
  out = new float[1];
  std::fill(in, in + N, INF);

  size_t temp_storage_bytes = 1024;
  void* d_temp_storage;
  cudaMalloc((void**)&d_temp_storage, temp_storage_bytes);

  cudaMalloc((void**)&d_in, sizeof(float) * N);
  cudaMalloc((void**)&d_out, sizeof(float) * 1);

  cudaMemcpy(d_in, in, sizeof(float) * N, cudaMemcpyHostToDevice);
  cub::DeviceReduce::Min(d_temp_storage, temp_storage_bytes, d_in, d_out, N);
  cudaMemcpy(out, d_out, sizeof(float), cudaMemcpyDeviceToHost);

  std::cout << *out << std::endl;  // Expected: inf
}

Expected behavior

return inf

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response

[leofang]

Hmmm I just found someone actually asked the same question. @gevtushenko does your comment back then (NVIDIA/cub#662 (comment)) still hold? If so, we can probably close this issue without any action on the CCCL side.

[leofang]

(Thinking about it a bit more, I am not sure if this is a regression in CUB. Let me do some checks first...)

[leofang]

Checked all the way back to CCCL/CUB v2.0.1, doesn't seem to be a regression on the CUB side, more like an overlook in CuPy's test suite.

[leofang]

Hmmm I just found someone actually asked the same question. @gevtushenko does your comment back then (NVIDIA/cub#662 (comment)) still hold? If so, we can probably close this issue without any action on the CCCL side.

Confirmed with Georgii offline that this still holds, closing.