[Clang][OpenMP][NVPTX] Fixed failure in openmp-offload-gpu.c if the system has CUDA

[tianshilei1992]

Summary

https://lists.llvm.org/pipermail/openmp-dev/2021-March/003940.html reports
test failure in openmp-offload-gpu.c. The failure is, when using -S in the
clang driver, it still reports bitcode library doesn't exist. However, it is not
exposed in my local run and Phabiractor test. The reason it escaped from Phabricator
test is, the test machine doesn't have CUDA, so LibDeviceFile is empty. In this
case, the check of OPT_S will be hit, and we get "expected" result. However, if
the test machine has CUDA, LibDeviceFile will not be empty, then the check will
not be done, and it just proceeds, trying to add the bitcode library. The reason
it escaped from my local run is, I didn't build ALL targets, so this case was
marked UNSUPPORTED.

[tra]

My question is, if DeviceOffloadingKind == Action::OFK_Cuda, and we use -S, do we also want to skip as well?

I do not think so. Libdevice is needed to implement some libcalls that LLVM currently does not know how to handle.
We do need it even when we compile with -S. It may work without it in many cases, but it's still needed in general.

[tianshilei1992]

In D98902#2635930, @tra wrote:

My question is, if DeviceOffloadingKind == Action::OFK_Cuda, and we use -S, do we also want to skip as well?

I do not think so. Libdevice is needed to implement some libcalls that LLVM currently does not know how to handle.
We do need it even when we compile with -S. It may work without it in many cases, but it's still needed in general.

Thanks for the answer.

[kkwli0]

I tried the patch in our environment and it works. LG. Thanks.

[jdoerfert]

@tra, so you think we should not do this? The user will see a link error late I assume, might be better.

[tra]

In D98902#2636308, @jdoerfert wrote:

@tra, so you think we should not do this? The user will see a link error late I assume, might be better.

If I compile a __device__ float foo(float f) { return sin(f); } and it does compile to working GPU code, I If I compile the same code with -S, I would assume that produced PTX is still compileable with ptxas. After all, it was when the source was compiled with -c. That will no longer be the case if you disable linking with libdevice.

If the user wants to disable linking with the libdevice, there's already -nogpulib for that.

[tianshilei1992]

In D98902#2636308, @jdoerfert wrote:

@tra, so you think we should not do this? The user will see a link error late I assume, might be better.

I think @tra 's point is we should not do that for CUDA code. This change only affects OpenMP.

[tra]

In D98902#2640477, @tianshilei1992 wrote:

In D98902#2636308, @jdoerfert wrote:

@tra, so you think we should not do this? The user will see a link error late I assume, might be better.

I think @tra 's point is we should not do that for CUDA code. This change only affects OpenMP.

Correct. We do want to link with libdevice during CUDA compilation, even with -S. I don't have a strong opinion on what OpenMP does.

[Hahnfeld]

Can we get this fixed somehow? It's annoying that there is a test failure in Clang without building the OpenMP runtime, just because I have CUDA installed on my machine...

[tianshilei1992]

Ping...

[kkwli0]

No more comments from the community. I think it is okay to accept this revision. Thanks.

clang/lib/Driver/ToolChains/Cuda.cpp

if (DeviceOffloadingKind == Action::OFK_Cuda) {
     if (DriverArgs.hasFlag(options::OPT_fcuda_approx_transcendentals,
                            options::OPT_fno_cuda_approx_transcendentals, false))
       CC1Args.push_back("-fcuda-approx-transcendentals");
   }
 
   if (DriverArgs.hasArg(options::OPT_nogpulib))
     return;
 
-  std::string LibDeviceFile = CudaInstallation.getLibDeviceFile(GpuArch);
-
-  if (LibDeviceFile.empty()) {
-    if (DeviceOffloadingKind == Action::OFK_OpenMP &&
-        DriverArgs.hasArg(options::OPT_S))
-      return;
 
+  std::string LibDeviceFile = CudaInstallation.getLibDeviceFile(GpuArch);
+  if (LibDeviceFile.empty()) {
     getDriver().Diag(diag::err_drv_no_cuda_libdevice) << GpuArch;
     return;
   }
 
   CC1Args.push_back("-mlink-builtin-bitcode");
   CC1Args.push_back(DriverArgs.MakeArgString(LibDeviceFile));
 
   clang::CudaVersion CudaInstallationVersion = CudaInstallation.version();