NVIDIA / NVIDIA/cccl

[BUG] inclusive_scan_by_key OOM on >= INT_MAX elements

Open
#766 12 comments 0 reactions 1 assignee Claimed by @griwes View on GitHub
thrust
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

Possibly related to the fix at https://github.com/NVIDIA/cub/issues/221, we're seeing an OOM in `inclusive_scan_by_key` when input size is close to (but not exactly) `INT_MAX`. The limit seems to depend on the accumulator type -- sometimes `(1 << 31) - 1` works, but other times it's `(1 << 31) - 2048` or `(1 << 31) - 4096`.

Tested in `nvidia/cuda:10.2-devel-ubuntu18.04` container with GCC 7.5.0/CUDA 10.2, but present in newer versions too.

```CMake
# CMakeLists.txt
cmake_minimum_required(VERSION 3.20)
project(test VERSION 1.0.0 LANGUAGES C CXX CUDA)
set(CMAKE_EXPORT_COMPILE_COMMANDS ON)

# https://github.com/cpm-cmake/CPM.cmake/blob/310efb9b17d0befe9ccd4f5bf2e39942869777fc/cmake/get_cpm.cmake
set(CPM_DOWNLOAD_VERSION 0.32.0)
if(CPM_SOURCE_CACHE)
# Expand relative path. This is important if the provided path contains a tilde (~)
get_filename_component(CPM_SOURCE_CACHE ${CPM_SOURCE_CACHE} ABSOLUTE)
set(CPM_DOWNLOAD_LOCATION "${CPM_SOURCE_CACHE}/cpm/CPM_${CPM_DOWNLOAD_VERSION}.cmake")
elseif(DEFINED ENV{CPM_SOURCE_CACHE})
set(CPM_DOWNLOAD_LOCATION "$ENV{CPM_SOURCE_CACHE}/cpm/CPM_${CPM_DOWNLOAD_VERSION}.cmake")
else()
set(CPM_DOWNLOAD_LOCATION "${CMAKE_BINARY_DIR}/cmake/CPM_${CPM_DOWNLOAD_VERSION}.cmake")
endif()
if(NOT (EXISTS ${CPM_DOWNLOAD_LOCATION}))
message(STATUS "Downloading CPM.cmake to ${CPM_DOWNLOAD_LOCATION}")
file(DOWNLOAD
https://github.com/cpm-cmake/CPM.cmake/releases/download/v${CPM_DOWNLOAD_VERSION}/CPM.cmake
${CPM_DOWNLOAD_LOCATION}
)
endif()
include(${CPM_DOWNLOAD_LOCATION})

# Get Thrust v1.12.0
CPMAddPackage(NAME Thrust
VERSION 1.12.0
GIT_REPOSITORY https://github.com/NVIDIA/thrust.git
GIT_TAG 1.12.0
GIT_SHALLOW TRUE
UPDATE_DISCONNECTED TRUE)

thrust_create_target(test::Thrust FROM_OPTIONS)

add_executable(test_inclusive_scan_limit test_inclusive_scan_limit.cu)

target_compile_features(test_inclusive_scan_limit
PRIVATE cxx_std_14 $)

target_link_libraries(test_inclusive_scan_limit test::Thrust)

execute_process(COMMAND cmake -E create_symlink ${CMAKE_CURRENT_BINARY_DIR}/compile_commands.json compile_commands.json)
```

```C++
// test_inclusive_scan_limit.cu
#include
#include
#include
#include
#include

int main(void) {
//
// limit is somewhere between these values, otherwise get this error:
// terminate called after throwing an instance of 'thrust::system::detail::bad_alloc'
// what(): std::bad_alloc: cudaErrorMemoryAllocation: out of memory
//
auto const scan_by_key_limit = (1uL << 31uL); // error
// auto const scan_by_key_limit = (1uL << 31uL) - 1uL; // no error
// auto const scan_by_key_limit = (1uL << 31uL) - 2048uL; // no error
// auto const scan_by_key_limit = (1uL << 31uL) - 4096uL; // no error

auto keys = thrust::make_constant_iterator(0);
auto vals = thrust::make_counting_iterator(uint64_t{0});

thrust::host_vector h_sum(1);
thrust::device_vector d_sums(scan_by_key_limit);

std::cout << "size: " << scan_by_key_limit << std::endl;

// OOM here
auto s_end = thrust::inclusive_scan_by_key(
thrust::device, keys, keys + scan_by_key_limit, vals, d_sums.begin());

thrust::copy(s_end - 1, s_end, h_sum.begin());

std::cout << "sums: " << h_sum[0] << std::endl;
}
```

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.