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

[BlackMark29A]

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:

#include <cassert>
#include <numeric>

#include <sycl/sycl.hpp>

int main() {
    constexpr auto REDUCTION_SIZE = 1024;

    auto queue = sycl::queue();

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

    sycl::buffer<int> maxBuf{ 1 };

    queue.submit([&](sycl::handler& cgh) {
        auto inputValues = valuesBuf.get_access<sycl::access_mode::read>(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: ec34869
• Dependencies version: CUDA 11

[JackAKirk]

Interesting, the program runs fine on two different cuda devices I've tried, but I don't have access to an rtx 2070. Could you check that when you run the program nothing else is executing on the gpu: if you do $ nvidia-smi it should show no processes before you execute the program. Just for completeness: how did you compile the program?

[BlackMark29A]

Okay, this is very interesting.

So, first of all, there were no other applications running that could influence the result, I already made sure of that.
I now also tested the code again on another system with a similar software stack but with an RTX 3090 card, and it also crashes there.

As for the compiler invocation, that's where it gets interesting. I was using cmake, and invoking the compiler with the same flags (/path/to/dpcpp/clang++ -g -fsycl -fsycl-targets=nvptx64-nvidia-cuda -std=c++17 -o main main.cpp) also results in the exception being thrown. However, I wanted to get rid of some flags that aren't necessary and found that -g actually causes the problem. Without specifying -g everything works as expected.

[JackAKirk]

Okay, this is very interesting.

So, first of all, there were no other applications running that could influence the result, I already made sure of that. I now also tested the code again on another system with a similar software stack but with an RTX 3090 card, and it also crashes there.

As for the compiler invocation, that's where it gets interesting. I was using cmake, and invoking the compiler with the same flags (/path/to/dpcpp/clang++ -g -fsycl -fsycl-targets=nvptx64-nvidia-cuda -std=c++17 -o main main.cpp) also results in the exception being thrown. However, I wanted to get rid of some flags that aren't necessary and found that -g actually causes the problem. Without specifying -g everything works as expected.

Ah OK that could make sense. I think this could be because -g compiles at -O0. We had a few issues before with -O0. We will investigate this further. Thanks

[zjin-lcf]

I can reproduce the error ("-g") on a RTX-series GPU while the memory usage is busy for other programs. However, "-O3" will not cause the error.

[BlackMark29A]

I just tested all combinations of debug symbol and optimization flags. This is the result:

	no -O	-O0	-O1	-O2	-O3
no -g	✔️	✔️	✔️	✔️	✔️
-g0	✔️	✔️	✔️	✔️	✔️
-g1	❌	✔️	✔️	✔️	✔️
-g2	❌	✔️	✔️	✔️	✔️
-g3	❌	✔️	✔️	✔️	✔️
-g	❌	✔️	✔️	✔️	✔️

Where ✔️ represents that the program works as expected and ❌ throws the out of resources exception.

[zjin-lcf]

The compiler may optimize the kernel compiled with "-g" while a user may compile the kernel with the option

"
The target architecture may also be specified for the CUDA backend, with -Xsycl-target-backend --cuda-gpu-arch=. Specifying the architecture is necessary if an application aims to use newer hardware features, such as native atomic operations or tensor core operations. Moreover, it is possible to pass specific options to CUDA ptxas (such as --maxrregcount= for limiting the register usage or --verbose for printing generation statistics) using the -Xcuda-ptxas flag.

clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda
simple-sycl-app.cpp -o simple-sycl-app-cuda.exe
-Xcuda-ptxas --maxrregcount=128 -Xcuda-ptxas --verbose
-Xsycl-target-backend --cuda-gpu-arch=sm_80
"

[ldrumm]

I just tested all combinations of debug symbol and optimization flags. This is the result:

I've recreated your exact matrix above on an A100. I'm investigating this now.

[ldrumm]

I compared two compilations:
-g (CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES) and no -g (executes fine), and note that the assemblies themselves are functionally identical. The only difference in the generated programs are the label, location, and debug section directives that are obviously required for encoding debug info.

I've attached the two compilation results (github wouldn't let met upload .s so I renamed them as .s.txt). I'm going to dig deeper, but my working theory is that this feels like a CUDA driver issue rather than a SYCL issue, since the executed instructions are identical.

There's still a distinct possibility that we're doing something wrong, so I'll see if I can find the root cause; watch this space

g.s
no-g.s

[tom91136]

I've observed something similar with or without -g.

In a reduction with complex kernel body we're seeing ptxas (enabled with -Xcuda-ptxas --verbose) reporting values as high as 66 registers.
This in itself is not a problem as long as we don't blow the per thread-block register budget (see https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#features-and-technical-specifications-technical-specifications-per-compute-capability, anything after sm_70 is 64k.)

Now, in a reduction, the range appears to be passed directly as the threads-per-block for a cuLaunchKernel:

llvm/sycl/plugins/cuda/pi_cuda.cpp

Line 3147 in 2953d25

retError = PI_CHECK_ERROR(cuLaunchKernel(

So for a register count of 66 and a sycl::range<1>(1024), we get 66*1024=67584 which is larger than 64k.
For normal parallel_for calls, the thread-per-block size in the CUDA PI appears to be reduced appropriately by the grid count so we never blow the budget.

I think we just need to guard the launch by checking for register count (cuFuncGetAttribute(&_reg, CU_FUNC_ATTRIBUTE_NUM_REGS, cuFunc)) and then making sure we increase the grid size such that threads-per-block*reg doesn't exceed the limit.

We probably also need to check the logic around cuOccupancyMaxPotentialBlockSize , as used in:

llvm/sycl/plugins/cuda/pi_cuda.cpp

Line 281 in 2953d25

void guessLocalWorkSize(_pi_device *device, size_t *threadsPerBlock,

[ldrumm]

@tom91136 We've seen something like this internally with a kernel from blender, but I don't think there's a public ticket for it. @GeorgeWeb might have some advice.

[tom91136]

Should I open a separate ticket for this? (Or maybe we can file it in our CoE bug tracker)

[ldrumm]

@BlackMark29A

It seems that without -g, the compiler driver always passes -O3 to ptxas (which isn't really an assembler, but actually an optimizing compiler of the PTX intermediate language. It's responsible for register allocation et el).

When we're generating debug info, we can't actually pass that flag:

$ clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda main.cpp -g -Xcuda-ptxas -O3 
ptxas fatal   : Optimized debugging not supported

So this issue is not as a result of debug info at all, but appears to be a general resource allocation problem. Thus, I think it's equivalent to @tom91136's described issue, and one that @GeorgeWeb has already seen.

We'll look at ways to gracefully handle this.