`reduce` fail due to integer overflow.

[trivialfis]

Platform

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2020 NVIDIA Corporation
Built on Wed_Jul_22_19:09:09_PDT_2020
Cuda compilation tools, release 11.0, V11.0.221
Build cuda_11.0_bu.TC445_37.28845127_0

Reproduce

#include <thrust/iterator/counting_iterator.h>
#include <thrust/reduce.h>
#include <limits>

void Sync() {
  assert(cudaDeviceSynchronize() == cudaSuccess);
  assert(cudaGetLastError() == cudaSuccess);
}

void TestReduceFails() {
  auto iter = thrust::make_counting_iterator<size_t>(0ul);
  size_t limit = std::numeric_limits<int32_t>::max();
  // after reduction step 2: cudaErrorInvalidConfiguration: invalid configuration argument
  auto ret = thrust::reduce(thrust::device, iter, iter + limit, size_t(0ul), thrust::maximum<size_t>{});
  Sync();
  assert(ret == limit - 1);
}

void TestReduceWorks() {
  auto iter = thrust::make_counting_iterator<size_t>(0ul);
  size_t limit = std::numeric_limits<int32_t>::max() / 2;
  auto ret = thrust::reduce(thrust::device, iter, iter + limit, size_t(0ul), thrust::maximum<size_t>{});
  Sync();
  assert(ret == limit - 1);
}

int main() {
  TestReduceFails();
  TestReduceWorks();
  return 0;
}

[alliepiper]

Probably related to NVIDIA/cccl#744.

[lilohuang]

@allisonvacanti @senior-zero

Confirmed the bug is still there with CUDA11.8 compiler with THRUST/CUB 1.17.2. It produces GPU hang issue with INT_MAX elements

#include <limits>
#include <iostream>
#include <thrust/reduce.h>
#include <thrust/execution_policy.h>
#include <thrust/iterator/constant_iterator.h>

int main() {
   
   auto begin = thrust::make_constant_iterator<size_t>(1);
   auto r = thrust::reduce(thrust::device,
      begin,
      begin + std::numeric_limits<int32_t>::max(),
      static_cast<size_t>(0));
   std::cout << r << std::endl;
   return 0;
}

However, the result is correct with INT_MAX + 1

#include <limits>
#include <iostream>
#include <thrust/reduce.h>
#include <thrust/execution_policy.h>
#include <thrust/iterator/constant_iterator.h>

int main() {
   
   auto begin = thrust::make_constant_iterator<size_t>(1);
   auto r = thrust::reduce(thrust::device,
      begin,
      begin + std::numeric_limits<int32_t>::max() +1 ,
      static_cast<size_t>(0));
   std::cout << r << std::endl;
   return 0;
}

This looks like NOT the same issue (exceeds 2^31) tracked by NVIDIA/cccl#744

[gevtushenko]

Hello @lilohuang! Recently, I've made changes to cub/thrust reduction related to offsets computation. On 2.X branch (that contains the changes), your code works fine for me, could you please verify?

[lilohuang]

Hi @senior-zero , confirmed the bug has been fixed on the latest thrust/cub branch. Looking forward to seeing a new release with the fix soon. It would be great if you can also backport the fix to thrust/cub 1.x. Thanks!