Program with device code in multiple translation units fails on CUDA

[sergey-semenov]

Describe the bug
A simple program with device code in multiple translation units fails in runtime with CUDA_ERROR_INVALID_IMAGE as of #3735

To Reproduce
h.hpp:

#include <CL/sycl.hpp>

void submit_kernelB();

b.cpp

#include "h.hpp"

class KernelNameB;

void submit_kernelB() {
  sycl::queue q;
  q.submit([&](sycl::handler &cgh) { cgh.single_task<KernelNameB>([]() {}); });
}

main.cpp:

#include "h.hpp"
#include <CL/sycl.hpp>

class KernelNameA;
void submit_kernelA() {
  sycl::queue q;
  q.submit([&](sycl::handler &cgh) { cgh.single_task<KernelNameA>([]() {}); });
}

int main() { submit_kernelA(); }

clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda-sycldevice main.cpp b.cpp
./a.out

This reproducer fails with CUDA_ERROR_INVALID_IMAGE, note that compiling this results in 2 device images as of #3735, but in only one with it reverted. The error disappears once the number of device images in the application is reduced to 1 (either by moving submit_kernelB to the same translation unit as submit_kernelA, by using -fsycl-device-code-split=off or by reverting #3735).

Environment:

• OS: Linux
• Target device and vendor: CUDA, Titan RTX.
• DPC++ version: e9d308e
• Dependencies version: CUDA 10.1

[bader]

@sergey-semenov, I think this issue is already fixed by 351af24. Could you check with newer version of the compiler?

[bader]

#4088 and #4079 seems to be about the same problem.

[sergey-semenov]

Ah, I specified the guilty commit rather than the version of the compiler I reproduced this on. This problem is still reproducible on e9d308e, so it seems 351af24 didn't address this.

[bader]

Okay, thanks for the update.
Maybe #4107 will help this case.

[Michoumichmich]

@sergey-semenov can you try that? Reverting only the driver is enough to make cuda work again (temporarily)

[sergey-semenov]

@Michoumichmich Reverting the driver part of f7ce532 didn't help, the reproducer is still failing as before.

[bader]

A simple program with device code in multiple translation units fails in runtime with CUDA_ERROR_INVALID_IMAGE as of #3735

@steffenlarsen, could you take a look, please? It looks like #3735 introduced a significant functional regression.

[steffenlarsen]

#4107 introduces a better solution for the driver changes in #3735. @sergey-semenov would you be able to check if that solves this issue?

[sergey-semenov]

This is indeed resolved by #4107