This is an archive of the discontinued LLVM Phabricator instance.

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

Authored by tianshilei1992 on Mar 18 2021, 2:56 PM.

Details

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.

Diff Detail

Event Timeline

tianshilei1992 created this revision.Mar 18 2021, 2:56 PM
tianshilei1992 requested review of this revision.Mar 18 2021, 2:56 PM
Herald added a project: Restricted Project. · View Herald Transcript
tra added a subscriber: tra.Mar 18 2021, 3:09 PM

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.

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.

tianshilei1992 edited the summary of this revision. (Show Details)Mar 18 2021, 3:15 PM

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

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

tra added a comment.Mar 19 2021, 10:17 AM

@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.

@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 added a comment.EditedMar 22 2021, 10:36 AM

@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.

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...

kkwli0 accepted this revision.Apr 13 2021, 10:19 AM

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

This revision is now accepted and ready to land.Apr 13 2021, 10:19 AM
This revision was landed with ongoing or failed builds.Apr 13 2021, 10:22 AM
This revision was automatically updated to reflect the committed changes.