intel / intel/llvm

[SYCL][CUDA] Reductions with 1024+ elements crash with CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES

Open
#8,414 16 comments 1 reaction 2 assignees Claimed by @GeorgeWeb View on GitHub
bug compiler confirmed cuda runtime
Dominant language
LLVM
Stars
1.5k
Forks
854
Avg merge
3d 17h
Merged PRs (30d)
137

Description

**Describe the bug**

Reductions with the CUDA backend throw an `CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES` exception, even though there should be enough resources available for the reduction. On my setup any reduction with more than 640 `int` elements throws. I've tested this with `sum` and `max` as the reduction operator.

**To Reproduce**

The following code crashes with an `CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES` exception being thrown:

```cpp
#include
#include

#include

int main() {
constexpr auto REDUCTION_SIZE = 1024;

auto queue = sycl::queue();

sycl::buffer valuesBuf{ REDUCTION_SIZE };
{
sycl::host_accessor a{ valuesBuf };
std::iota(a.begin(), a.end(), 0);
}

sycl::buffer maxBuf{ 1 };

queue.submit([&](sycl::handler& cgh) {
auto inputValues = valuesBuf.get_access(cgh);
auto maxReduction = sycl::reduction(maxBuf, cgh, sycl::maximum<>());
cgh.parallel_for(valuesBuf.get_range(), maxReduction,
[=](sycl::id<1> idx, auto& max) { max.combine(inputValues[idx]); });
});

assert(maxBuf.get_host_access()[0] == REDUCTION_SIZE - 1);
}
```

Output:

```
PI CUDA ERROR:
Value: 701
Name: CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES
Description: too many resources requested for launch
Function: cuda_piEnqueueKernelLaunch
Source Location: /home/blackmark/dpcpp/llvm/sycl/plugins/cuda/pi_cuda.cpp:3055

terminate called after throwing an instance of 'sycl::_V1::runtime_error'
what(): Native API failed. Native API returns: -5 (PI_ERROR_OUT_OF_RESOURCES) -5 (PI_ERROR_OUT_OF_RESOURCES)
Aborted (core dumped)
```

Expected behavior would obviously be that reductions which should be possible given the hardware resources don't crash.

**Environment (please complete the following information):**

- OS: Ubuntu 22.04
- Target device and vendor: NVIDIA RTX 2070
- DPC++ version: ec348692e8d9272171a3204ecf7142c34607c9b3
- Dependencies version: CUDA 11

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.