[Clang][OpenMP] Add partial support for Static Device Libraries

[saiislam]

Summary

An archive containing device code object files can be passed to
clang command line for linking. For each given offload target
it creates a device specific archives which is either passed to llvm-link
if the target is amdgpu, or to clang-nvlink-wrapper if the target is
nvptx. -L/-l flags are used to specify these fat archives on the command
line. E.g.

clang++ -fopenmp -fopenmp-targets=nvptx64 main.cpp -L. -lmylib

It currently doesn't support linking an archive directly, like:

clang++ -fopenmp -fopenmp-targets=nvptx64 main.cpp libmylib.a

Linking with x86 offload also does not work.

[jdoerfert]

Not a thorough review but comments to address.

clang/lib/Driver/ToolChains/CommonArgs.cpp

1663

Coding convention.

1675

This will accumulate more and more entries in AOBFileNames, shouldn't you clear it at the beginning?

1685

1811

1819

Lots of places you create new std::strings for no reason, use StringRef or const std::string & to avoid that (especially in range loops).

Please no conditional like a != ".." && a != ".." ..... Use a compile time set or string switch.

[grokos]

clang/lib/Driver/ToolChains/CommonArgs.cpp

1690

"a" --> ".a" (add a dot)

1797–1799

I'm with @jdoerfert here, you can use a set of library names which are known to not have device-specific SDLs and check whether that set contains SDL_Name. Also, SDL_Names can be a set of unique entries, this way even if you try to add the same library twice it won't be added. This quadratic-complexity loop looks ugly...

[jdoerfert]

Tests?

[saiislam]

1. Simplified the nested loop to look for repeated SDLs.
2. Minimized creation of std::string in favour of StringRef wherever possible.
3. Added test cases

[saiislam]

Couple of wrong files got added in the last commit.

[saiislam]

clang/lib/Driver/ToolChains/Clang.cpp

7745–7758

It will be replaced with a method invocation to get GPUArch in an upcoming patch to support multi-architecture compilation.

clang/lib/Driver/ToolChains/CommonArgs.cpp

1690

"a" is second argument of GetTemporayPath and doesn't the prefix "."

[saiislam]

Added instructions to generate a fat archive and reduced the size of attached libFatArchive.a used for testing.

[jdoerfert]

x64 debian > Clang.Driver::fat_archive.cpp Failed

@ABataev @grokos any comments?

[ye-luo]

Do I must use llvm-ar/ranlib or system ar/ranlib is OK?

1. existing use case breaks

Use https://github.com/ye-luo/openmp-target/blob/master/tests/math/modf.cpp
$ clang++ -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target=nvptx64 -march=sm_80 modf.cpp # still OK
two steps
$ clang++ -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target=nvptx64 -march=sm_80 modf.cpp -c
$ clang++ -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target=nvptx64 -march=sm_80 modf.o
clang-14: warning: Unknown CUDA version. version.txt: 11.0.228. Assuming the latest supported version 10.1 [-Wunknown-cuda-version]
nvlink fatal : Could not open input file '/tmp/modf-0bf89b.cubin'
clang-14: error: nvlink command failed with exit code 1 (use -v to see invocation)

2. could you make my test case working?

https://github.com/ye-luo/openmp-target/tree/master/tests/link_static_fat_bin
both compile-amd.sh and compile.sh doesn't work for me.
linking was successful but no device code is in the executable and the run fails.

Does your test actually test a run?

[saiislam]

In D105191#2910961, @ye-luo wrote:

Do I must use llvm-ar/ranlib or system ar/ranlib is OK?

1. existing use case breaks

Use https://github.com/ye-luo/openmp-target/blob/master/tests/math/modf.cpp
$ clang++ -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target=nvptx64 -march=sm_80 modf.cpp # still OK
two steps
$ clang++ -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target=nvptx64 -march=sm_80 modf.cpp -c
$ clang++ -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target=nvptx64 -march=sm_80 modf.o
clang-14: warning: Unknown CUDA version. version.txt: 11.0.228. Assuming the latest supported version 10.1 [-Wunknown-cuda-version]
nvlink fatal : Could not open input file '/tmp/modf-0bf89b.cubin'
clang-14: error: nvlink command failed with exit code 1 (use -v to see invocation)

2. could you make my test case working?

https://github.com/ye-luo/openmp-target/tree/master/tests/link_static_fat_bin
both compile-amd.sh and compile.sh doesn't work for me.
linking was successful but no device code is in the executable and the run fails.

Does your test actually test a run?

1. System ar/ranlib is OK.
2. modf.cpp test case is breaking due to incompatibility between bundle entry ID formats which got introduced by my last patch D93525. I think D106809 should fix this, but I haven't tested it on this case yet.
3. link_static_fat_bin is working fine. The missing component was inability of nvlink to take archive of cubin files as input. It requires a wrapper over nvlink (D108291). Please review.

[tra]

clang/lib/Driver/ToolChains/CommonArgs.cpp

1599–1618

This could probably be collapsed into a couple of nested loops:

for (auto base: {"/libbc-", "/lib"}) {
  std::string ext = base == "/lib" ? ".bc" : ".a";
  for( auto suffix : {Lib + "-" + Arch + "-" + Target, Lib + "-" + Arch, Lib } }) {
    SDLs.append({"/libdevice" + base + suffix + ext,  base + suffix + ext })
  }
}

[ye-luo]

clang/lib/Driver/ToolChains/CommonArgs.h

62

Differentiate the names of all the three AddStaticDeviceLibs functions and add documentation. Be sure to do document every function added in this patch.

[ye-luo]

@saiislam since clang-nvlink-wrapper has landed, could you update this patch?

[saiislam]

In D105191#2985819, @ye-luo wrote:

@saiislam since clang-nvlink-wrapper has landed, could you update this patch?

@ye-luo , In the last multi-company OpenMP-dev meeting it was suggested that D106809 should be reviewed separately because it is logically a separate patch. But, D106809 is required for this patch so that it doesn't break our downstream pipeline.
Could you please take a look at D106809?

[saiislam]

Added documentation and other fixes suggested by reviewers.

[ye-luo]

@saiislam do my test cases work on your side? I tried this patch and still got linking failure.

[saiislam]

In D105191#2996201, @ye-luo wrote:

@saiislam do my test cases work on your side? I tried this patch and still got linking failure.

Hey @ye-luo,
I am able to successfully compile and run all (36) tests in https://github.com/ye-luo/openmp-target (branch master, commit id: dd0544f).
Here are the steps I followed (all seems to be straightforward):

1. Get updated trunk
2. Apply this patch
3. Build and install llvm-project from step 2.
4. Set CC and CXX variables to newly installed location from step 3.
5. Get openmp-target repo
6. cmake, make, make all, make test

Output of make test

/work/saiyedul/openmp-target/build$ make test
Running tests...
Test project /work/saiyedul/openmp-target/build

Start  1: cxx.complex_reduction_cpu

1/36 Test #1: cxx.complex_reduction_cpu .......................... Passed 0.01 sec

Start  2: cxx.complex_reduction

2/36 Test #2: cxx.complex_reduction .............................. Passed 0.01 sec

Start  3: cxx.complex

3/36 Test #3: cxx.complex ........................................ Passed 0.00 sec

Start  4: f.complex

4/36 Test #4: f.complex .......................................... Passed 0.00 sec

Start  5: cxx.global_static

5/36 Test #5: cxx.global_static .................................. Passed 0.00 sec

Start  6: cxx.constexpr

6/36 Test #6: cxx.constexpr ...................................... Passed 0.00 sec

Start  7: cxx.link_static_fat_bin

7/36 Test #7: cxx.link_static_fat_bin ............................ Passed 0.00 sec

Start  8: cxx.linker_outlined_function_collision

8/36 Test #8: cxx.linker_outlined_function_collision ............. Passed 0.00 sec

Start  9: cxx.math_FP_ZERO

9/36 Test #9: cxx.math_FP_ZERO ................................... Passed 0.00 sec

Start 10: cxx.math_header_only

10/36 Test #10: cxx.math_header_only ............................... Passed 0.00 sec

Start 11: cxx.math_modf

11/36 Test #11: cxx.math_modf ...................................... Passed 0.00 sec

Start 12: cxx.math_sqrt_simd

12/36 Test #12: cxx.math_sqrt_simd ................................. Passed 0.00 sec

Start 13: cxx.math_sin_cos

13/36 Test #13: cxx.math_sin_cos ................................... Passed 0.00 sec

Start 14: cxx.math_sin_simd

14/36 Test #14: cxx.math_sin_simd .................................. Passed 0.00 sec

Start 15: cxx.math_sincos

15/36 Test #15: cxx.math_sincos .................................... Passed 0.00 sec

Start 16: cxx.math_sincos_simd

16/36 Test #16: cxx.math_sincos_simd ............................... Passed 0.00 sec

Start 17: cxx.math_sincos_simd_template

17/36 Test #17: cxx.math_sincos_simd_template ...................... Passed 0.00 sec

Start 18: cxx.target_teams_private__distribute

18/36 Test #18: cxx.target_teams_private__distribute ............... Passed 0.00 sec

Start 19: cxx.target__teams_private__distribute

19/36 Test #19: cxx.targetteams_privatedistribute .............. Passed 0.00 sec

Start 20: cxx.target_teams_distribute_private

20/36 Test #20: cxx.target_teams_distribute_private ................ Passed 0.00 sec

Start 21: cxx.target_teams__distribute_private

21/36 Test #21: cxx.target_teams__distribute_private ............... Passed 0.00 sec

Start 22: cxx.target__teams_distribute_private

22/36 Test #22: cxx.target__teams_distribute_private ............... Passed 0.00 sec

Start 23: cxx.target__teams__distribute_private

23/36 Test #23: cxx.targetteamsdistribute_private .............. Passed 0.00 sec

Start 24: cxx.target_teams_distribute_parallel_for_private

24/36 Test #24: cxx.target_teams_distribute_parallel_for_private ... Passed 0.00 sec

Start 25: f.target_teams_distribute_parallel_for_private

25/36 Test #25: f.target_teams_distribute_parallel_for_private ..... Passed 0.00 sec

Start 26: f.target_teams_distribute_private

26/36 Test #26: f.target_teams_distribute_private .................. Passed 0.00 sec

Start 27: cxx.target_nowait_task

27/36 Test #27: cxx.target_nowait_task ............................. Passed 0.00 sec

Start 28: cxx.taskloop_offload_nowait

28/36 Test #28: cxx.taskloop_offload_nowait ........................ Passed 0.00 sec

Start 29: cxx.taskloop

29/36 Test #29: cxx.taskloop ....................................... Passed 0.00 sec

Start 30: cxx.omp-task-bug

30/36 Test #30: cxx.omp-task-bug ................................... Passed 0.00 sec

Start 31: cxx.host_bug_libomp

31/36 Test #31: cxx.host_bug_libomp ................................ Passed 0.00 sec

Start 32: f.use_device_ptr_target

32/36 Test #32: f.use_device_ptr_target ............................ Passed 0.00 sec

Start 33: f.allocatable_array_device

33/36 Test #33: f.allocatable_array_device ......................... Passed 0.00 sec

Start 34: f.allocatable_array_device_ptr

34/36 Test #34: f.allocatable_array_device_ptr ..................... Passed 0.00 sec

Start 35: f.allocatable_array_device_isptr

35/36 Test #35: f.allocatable_array_device_isptr ................... Passed 0.00 sec

Start 36: f.allocatable_array_device_ptr_isptr

36/36 Test #36: f.allocatable_array_device_ptr_isptr ............... Passed 0.00 sec

100% tests passed, 0 tests failed out of 36

Label Time Summary:
cxx = 0.08 sec*proc (28 tests)
fortran = 0.02 sec*proc (8 tests)

Total Test time (real) = 0.12 sec

[ye-luo]

@saiislam did you turn on offload? https://github.com/ye-luo/openmp-target/wiki/OpenMP-offload-compilers#llvm-clang
On NVIDIA, it fails at CMake step. On AMD, make step stops because of unrelated issue.

Please make the exact reproducer 1 working. Right now I got

$ clang++ -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target=nvptx64 -march=sm_80 modf.o
nvlink fatal   : Could not open input file '/tmp/modf-88b730.cubin'
/soft/llvm/main-patched/bin/clang-nvlink-wrapper: error: 'nvlink' failed
clang-14: error: nvlink command failed with exit code 1 (use -v to see invocation)

[saiislam]

Fix for file types in fat archive.

[saiislam]

In D105191#2997272, @ye-luo wrote:

@saiislam did you turn on offload? https://github.com/ye-luo/openmp-target/wiki/OpenMP-offload-compilers#llvm-clang
On NVIDIA, it fails at CMake step. On AMD, make step stops because of unrelated issue.

Please make the exact reproducer 1 working. Right now I got

$ clang++ -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target=nvptx64 -march=sm_80 modf.o
nvlink fatal   : Could not open input file '/tmp/modf-88b730.cubin'
/soft/llvm/main-patched/bin/clang-nvlink-wrapper: error: 'nvlink' failed
clang-14: error: nvlink command failed with exit code 1 (use -v to see invocation)

@ye-luo
Can you please take this latest update and verify if it is working for you?
It is working for me.

Following modifications were also required for the compile.sh and compile-amd.sh.
Assuming name of your fat archive is libmylib.a and it is present in the current directory, the command to use it will be as follows:

clang++ -fopenmp -fopenmp-targets=nvptx64 main.cpp -L. -lmylib

or

clang++ -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 main.cpp -L. -lmylib

So, compile.sh should be:

clang++ -fopenmp -fopenmp-targets=nvptx64 -c -Xopenmp-target=nvptx64 -march=sm_50 classA.cpp
rm -f libmylib.a
ar qc libmylib.a classA.o
ranlib libmylib.a
clang++ -fopenmp -fopenmp-targets=nvptx64 main.cpp -L. -lmylib
./a.out

And compile-amd.sh should be:

clang++ -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 -c classA.cpp
rm -f libmylib.a
llvm-ar qc libmylib.a classA.o
llvm-ranlib libmylib.a
clang++ -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 main.cpp -L. -lmylib
./a.out

[ye-luo]

the modf test still doesn't work. The issue was from unbundle.
case 1 works.

clang++ -fopenmp -fopenmp-targets=nvptx64 modf.cpp -c
clang++ -fopenmp -fopenmp-targets=nvptx64 modf.o

case 2

clang++ -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target=nvptx64 -march=sm_80 modf.cpp -c
clang++ -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target=nvptx64 -march=sm_80 modf.o

doesn't work. The issue was a failure in unbundling

"/soft/llvm/main-patched/bin/clang-offload-bundler" -type=o -targets=host-x86_64-unknown-linux-gnu,openmp-nvptx64-sm_80 -inputs=modf.o -outputs=/tmp/modf-99be57.o,/tmp/modf-88a9eb.cubin -unbundle -allow-missing-bundles

ends up an empty cubin file. If I remove -sm_80 on the unbundle line
It seems that we need to have sm_80 when the object file is created.

When clang is compiled CLANG_OPENMP_NVPTX_DEFAULT_ARCH can be used to set the default sm_XX so it doesn't needs to be provided at compile time. If not set default is sm_35.
So in general sm_XX is always known when the compiler is used for compiling.

[saiislam]

In D105191#2998372, @ye-luo wrote:

the modf test still doesn't work. The issue was from unbundle.
case 1 works.

clang++ -fopenmp -fopenmp-targets=nvptx64 modf.cpp -c
clang++ -fopenmp -fopenmp-targets=nvptx64 modf.o

case 2

clang++ -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target=nvptx64 -march=sm_80 modf.cpp -c
clang++ -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target=nvptx64 -march=sm_80 modf.o

doesn't work. The issue was a failure in unbundling

"/soft/llvm/main-patched/bin/clang-offload-bundler" -type=o -targets=host-x86_64-unknown-linux-gnu,openmp-nvptx64-sm_80 -inputs=modf.o -outputs=/tmp/modf-99be57.o,/tmp/modf-88a9eb.cubin -unbundle -allow-missing-bundles

ends up an empty cubin file. If I remove -sm_80 on the unbundle line
It seems that we need to have sm_80 when the object file is created.

When clang is compiled CLANG_OPENMP_NVPTX_DEFAULT_ARCH can be used to set the default sm_XX so it doesn't needs to be provided at compile time. If not set default is sm_35.
So in general sm_XX is always known when the compiler is used for compiling.

Both cases 1 and 2 are working fine in my setup using these commands. I have sm_50, though.
If you think unbundler is the the issue, can you please tell what all code objects are inside modf.o bundle in both cases with following command:

clang-offload-bundler -type=o --inputs=modf.o --list

In my case it prints the following for case 1:

openmp-nvptx64
host-x86_64-unknown-linux-gnu

for case 2:

openmp-nvptx64-sm_50
host-x86_64-unknown-linux-gnu

[ye-luo]

case 1

yeluo@epyc-server:~/opt/openmp-target/tests/math$ clang++ -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target=nvptx64 -march=sm_80 modf.cpp -c
yeluo@epyc-server:~/opt/openmp-target/tests/math$ clang-offload-bundler -type=o --inputs=modf.o --list
openmp-nvptx64
host-x86_64-unknown-linux-gnu

case2

yeluo@epyc-server:~/opt/openmp-target/tests/math$ clang++ -fopenmp -fopenmp-targets=nvptx64 modf.cpp -c
yeluo@epyc-server:~/opt/openmp-target/tests/math$ clang-offload-bundler -type=o --inputs=modf.o --list
openmp-nvptx64
host-x86_64-unknown-linux-gnu

also try nvptx64-nvidia-cuda

yeluo@epyc-server:~/opt/openmp-target/tests/math$ clang++ -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda modf.cpp -c
warning: linking module '/soft/llvm/main-20210910/lib/libomptarget-nvptx-sm_80.bc': Linking two modules of different target triples: '/soft/llvm/main-20210910/lib/libomptarget-nvptx-sm_80.bc' is 'nvptx64' whereas 'modf.cpp' is 'nvptx64-nvidia-cuda'
 [-Wlinker-warnings]
1 warning generated.
yeluo@epyc-server:~/opt/openmp-target/tests/math$ clang-offload-bundler -type=o --inputs=modf.o --list
openmp-nvptx64-nvidia-cuda
host-x86_64-unknown-linux-gnu

None of them have sm_80 added.

Here is my clang build recipe

cmake -DCMAKE_C_COMPILER=gcc -DCMAKE_CXX_COMPILER=g++ \
    -DCMAKE_BUILD_TYPE=Release \
    -DCMAKE_INSTALL_PREFIX=$INSTALL_FOLDER \
    -DLLVM_ENABLE_BACKTRACES=ON \
    -DLLVM_ENABLE_WERROR=OFF \
    -DBUILD_SHARED_LIBS=OFF \
    -DLLVM_ENABLE_RTTI=ON \
    -DLLVM_TARGETS_TO_BUILD="X86;AMDGPU;NVPTX" \
    -DLLVM_ENABLE_ASSERTIONS=ON \
    -DLLVM_ENABLE_PROJECTS="clang;lld" \
    -DLLVM_ENABLE_RUNTIMES="libcxxabi;libcxx;openmp" \
    -DLIBOMPTARGET_NVPTX_COMPUTE_CAPABILITIES="80,61" \
    -DCLANG_OPENMP_NVPTX_DEFAULT_ARCH=sm_80 \
    -DLIBOMPTARGET_NVPTX_MAX_SM=38 \
    -DLIBOMPTARGET_ENABLE_DEBUG=ON \
    ../llvm-project/llvm

[saiislam]

Rebase and a minor fix.

[ye-luo]

1. modf works now.

2. if I modify the complile.sh

clang++ -fopenmp -fopenmp-targets=nvptx64 -c classA.cpp
rm -f libmylib.a
ar qc libmylib.a classA.o
ranlib libmylib.a
clang++ -fopenmp -fopenmp-targets=nvptx64 main.cpp -L. -lmylib
./a.out

doesn't work. I think the solution is adding sm_XX to the module name regardless of user command line.

3, directly linking static archive doesn't work.

clang++ -fopenmp -fopenmp-targets=nvptx64 main.cpp libmylib.a

CMake generates this style of link line. So this really needs to work.

only the following case works right now.

clang++ -fopenmp -fopenmp-targets=nvptx64 main.cpp -L. -lmylib

[saiislam]

In D105191#3001056, @ye-luo wrote:

1. modf works now.

2. if I modify the complile.sh

clang++ -fopenmp -fopenmp-targets=nvptx64 -c classA.cpp
rm -f libmylib.a
ar qc libmylib.a classA.o
ranlib libmylib.a
clang++ -fopenmp -fopenmp-targets=nvptx64 main.cpp -L. -lmylib
./a.out

doesn't work. I think the solution is adding sm_XX to the module name regardless of user command line.

3, directly linking static archive doesn't work.

clang++ -fopenmp -fopenmp-targets=nvptx64 main.cpp libmylib.a

CMake generates this style of link line. So this really needs to work.

only the following case works right now.

clang++ -fopenmp -fopenmp-targets=nvptx64 main.cpp -L. -lmylib

Thanks for confirming that modf and this patch works with -l/-L options.

The option of adding sm_XX in Bundle Entry ID when user hasn't used -march flag, comes under command line simplification. I have a bunch of upcoming patches which will significantly simplify OpenMP command line for GPU offloading. But, don't you think this feature is different than supporting static device libraries and should be dealt separately?

Our plan was always to support SDLs using -l and -L options, as described in Greg Rodgers's presentation about static device libraries in last year's LLVM-CTH Workshop. It was also discussed in multi company meetings. We can explore supporting direct linking of SDLs as described by you, but it seems to me that it is out of scope of this patch. What do you think?

PS: We (mulit-company OpenMP-dev meetings) have been tracking this feature for a while and we would very much like it to be picked for llvm-13.

[JonChesterfield]

In D105191#3001508, @saiislam wrote:

PS: We (mulit-company OpenMP-dev meetings) have been tracking this feature for a while and we would very much like it to be picked for llvm-13.

I think llvm-13 is on -final now so I think we've missed the train for big feature work. That's not so bad, 14 branches in ~6 months.

[ye-luo]

The option of adding sm_XX in Bundle Entry ID when user hasn't used -march flag, comes under command line simplification. I have a bunch of upcoming patches which will significantly simplify OpenMP command line for GPU offloading. But, don't you think this feature is different than supporting static device libraries and should be dealt separately?

Simplifying command line options is a different topic and we have not agreed upon how to do the simplification. In general, that is unrelated to this broken test case. But it can probably be fixed later with a separate patch.

Our plan was always to support SDLs using -l and -L options, as described in Greg Rodgers's presentation about static device libraries in last year's LLVM-CTH Workshop. It was also discussed in multi company meetings. We can explore supporting direct linking of SDLs as described by you, but it seems to me that it is out of scope of this patch. What do you think?

To me, supporting both -L/-l and libXXX.a are both required for static libraries. It is not reasonable to request users to use one way not the other. AOMP had no issue working in both ways.

PS: We (mulit-company OpenMP-dev meetings) have been tracking this feature for a while and we would very much like it to be picked for llvm-13.

Although we wanted to deliver this feature to users, I don't think we should tell users that linking static archive works with the current quality. So in my view, It is not suitable to backport it o 13.

This patch is good. It makes part of the static linking feature work but we can not claim "Static Device Libraries" is fully working. So please update the tile and also in the description saying what is working and what is not. Also need a test case for nvptx64.
In addition, I'd like to see commitment of addressing both above failed test cases before accepting this patch to help incremental development.

by the way, I found offload to x86 with static linking doesn't work with -L. -lmylib
https://github.com/ye-luo/openmp-target/blob/master/tests/linking/link_static_fat_bin/compile-x86.sh
I think this can also be addressed separately. FYI, AOMP happily makes it working.

[ye-luo]

This patch doesn't seem to break anything on my side.
@saiislam could you

1. address all the in-source review comments
2. update the title to [Clang][OpenMP] Add partial support for Static Device Libraries
3. update the patch description about what works and what doesn't.

I will accept this patch once these minor revisions are added.

[saiislam]

In D105191#3013955, @ye-luo wrote:

This patch doesn't seem to break anything on my side.
@saiislam could you

1. address all the in-source review comments
2. update the title to [Clang][OpenMP] Add partial support for Static Device Libraries
3. update the patch description about what works and what doesn't.

I will accept this patch once these minor revisions are added.

Sure. I will do these changes. Thanks.

[pdhaliwal]

clang/test/Driver/fat_archive.cpp

9 ↗

(On Diff #372514)

Here, LIBRARY_PATH is specifying path to build directory of openmp runtime which might not be available when openmp is not built. Why not use %S/Inputs directory and put the required files into that?

[saiislam]

Added nvptx test cases, simplified amdgpu test case, modified commit message.

[saiislam]

fixed the typos in the test cases

[saiislam]

fixed nvptx test

[saiislam]

Ping.

[ye-luo]

clang/lib/Driver/ToolChains/CommonArgs.h

62

@saiislam Fix this?

[saiislam]

clang/lib/Driver/ToolChains/CommonArgs.h

62

I have added documentation along with function definition in CommonArgs.cpp. Should I move it here? I thought keep documentation and code at the same place will improve readability.

[ye-luo]

clang/lib/Driver/ToolChains/CommonArgs.h

62

The docs are good. I was looking for "Differentiate the names of all the three AddStaticDeviceLibs functions"

[ye-luo]

clang/lib/Driver/ToolChains/CommonArgs.h

62

ping @saiislam

[saiislam]

Function name refactoring for AddStaticDeviceLibs.

[ye-luo]

LGTM.

[saiislam]

Fixed windows build bot errors after landing this patch with 06404d5488ea505b00f711393973db3ae32d01e9

[thakis]

Looks like this breaks tests on Mac: http://45.33.8.238/mac/36543/step_7.txt

Please take a look and revert for now if it takes a while to fix.

[saiislam]

In D105191#3047596, @thakis wrote:

Looks like this breaks tests on Mac: http://45.33.8.238/mac/36543/step_7.txt

Please take a look and revert for now if it takes a while to fix.

Thanks for reporting. Taking a look.

[thakis]

It's been broken for a while now. Should we revert this for now?

[thakis]

Looks like the test is failing on Windows too: http://45.33.8.238/win/46512/step_7.txt

[thakis]

Here's a failure on an official LLVM buildbot: https://lab.llvm.org/buildbot/#/builders/123/builds/6536 https://lab.llvm.org/buildbot/#/builders/123 has been red for over 8 hours now.

[JonChesterfield]

Blocking windows + osx for a couple of hours is bad, let's revert it. Will do so myself when I get to a desktop

[saiislam]

In D105191#3047903, @JonChesterfield wrote:

Blocking windows + osx for a couple of hours is bad, let's revert it. Will do so myself when I get to a desktop

Here is the fix: D111311
Please review it.

[thakis]

Fix didn't help, still broken at least on Windows: http://45.33.8.238/win/46515/step_7.txt

[saiislam]

In D105191#3047954, @thakis wrote:

Fix didn't help, still broken at least on Windows: http://45.33.8.238/win/46515/step_7.txt

Ohk. Reverting this patch in a few minutes.

[thakis]

Looks like the test looks for "ld" but Windows prints "\link.exe" -- maybe you need to pass an explicit triple to the compiler? Or relax the check line.

[MyDeveloperDay]

clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp

120

This file now fails clang-format

[saiislam]

clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp

120

Thanks for reporting. Fixed it in D111545. Please have a look.

clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp

if (HasLibm) {
       SmallVector<std::string, 12> BCLibs =
           AMDGPUOpenMPTC.getCommonDeviceLibNames(Args, SubArchName.str());
       llvm::for_each(BCLibs, [&](StringRef BCFile) {
         CmdArgs.push_back(Args.MakeArgString(BCFile));
       });
     }
   }
 
+  AddStaticDeviceLibsLinking(C, *this, JA, Inputs, Args, CmdArgs, "amdgcn",
+                      SubArchName,
+                      /* bitcode SDL?*/ true,
+                      /* PostClang Link? */ false);
   // Add an intermediate output file.
   CmdArgs.push_back("-o");
   const char *OutputFileName =
       getOutputFileName(C, OutputFilePrefix, "-linked", "bc");
   CmdArgs.push_back(OutputFileName);
   const char *Exec =
       Args.MakeArgString(getToolChain().GetProgramPath("llvm-link"));
   C.addCommand(std::make_unique<Command>(

clang/lib/Driver/ToolChains/Clang.cpp

if (const auto *OA = dyn_cast<OffloadAction>(CurDep)) {
         assert(CurTC == nullptr && "Expected one dependence!");
         CurKind = A->getOffloadingDeviceKind();
         CurTC = TC;
       });
     }
     Triples += Action::GetOffloadKindName(CurKind);
     Triples += '-';
     Triples += CurTC->getTriple().normalize();
-    if ((CurKind == Action::OFK_HIP || CurKind == Action::OFK_OpenMP ||
-         CurKind == Action::OFK_Cuda) &&
+    if ((CurKind == Action::OFK_HIP || CurKind == Action::OFK_Cuda) &&
         CurDep->getOffloadingArch()) {
       Triples += '-';
       Triples += CurDep->getOffloadingArch();
     }
+
+    // TODO: Replace parsing of -march flag. Can be done by storing GPUArch
+    //       with each toolchain.
+    StringRef GPUArchName;
+    if (CurKind == Action::OFK_OpenMP) {
+      // Extract GPUArch from -march argument in TC argument list.
+      for (unsigned ArgIndex = 0; ArgIndex < TCArgs.size(); ArgIndex++) {
+        auto ArchStr = StringRef(TCArgs.getArgString(ArgIndex));
+        auto Arch = ArchStr.startswith_insensitive("-march=");
+        if (Arch) {
+          GPUArchName = ArchStr.substr(7);
+          Triples += "-";
+          break;
+        }
+      }
+      Triples += GPUArchName.str();
+    }
   }
   CmdArgs.push_back(TCArgs.MakeArgString(Triples));
 
   // Get bundled file command.
   CmdArgs.push_back(
       TCArgs.MakeArgString(Twine("-outputs=") + Output.getFilename()));
 
   // Get unbundled files command.
for (unsigned I = 0; I < DepInfo.size(); ++I) {
     if (I)
       Triples += ',';
 
     auto &Dep = DepInfo[I];
     Triples += Action::GetOffloadKindName(Dep.DependentOffloadKind);
     Triples += '-';
     Triples += Dep.DependentToolChain->getTriple().normalize();
     if ((Dep.DependentOffloadKind == Action::OFK_HIP ||
-         Dep.DependentOffloadKind == Action::OFK_OpenMP ||
          Dep.DependentOffloadKind == Action::OFK_Cuda) &&
         !Dep.DependentBoundArch.empty()) {
       Triples += '-';
       Triples += Dep.DependentBoundArch;
     }
+    // TODO: Replace parsing of -march flag. Can be done by storing GPUArch
+    //       with each toolchain.
+    StringRef GPUArchName;
+    if (Dep.DependentOffloadKind == Action::OFK_OpenMP) {
+      // Extract GPUArch from -march argument in TC argument list.
+      for (uint ArgIndex = 0; ArgIndex < TCArgs.size(); ArgIndex++) {
+        StringRef ArchStr = StringRef(TCArgs.getArgString(ArgIndex));
+        auto Arch = ArchStr.startswith_insensitive("-march=");
+        if (Arch) {
+          GPUArchName = ArchStr.substr(7);
+          Triples += "-";
+          break;
+        }
+      }
+      Triples += GPUArchName.str();
+    }
   }
 
   CmdArgs.push_back(TCArgs.MakeArgString(Triples));
 
   // Get bundled file command.
   CmdArgs.push_back(
       TCArgs.MakeArgString(Twine("-inputs=") + Input.getFilename()));
 

clang/lib/Driver/ToolChains/CommonArgs.h

 
 void linkXRayRuntimeDeps(const ToolChain &TC,
                          llvm::opt::ArgStringList &CmdArgs);
 
 void AddRunTimeLibs(const ToolChain &TC, const Driver &D,
                     llvm::opt::ArgStringList &CmdArgs,
                     const llvm::opt::ArgList &Args);
 
+void AddStaticDeviceLibsLinking(Compilation &C, const Tool &T,
+                                const JobAction &JA,
+                                const InputInfoList &Inputs,
+                                const llvm::opt::ArgList &DriverArgs,
+                                llvm::opt::ArgStringList &CmdArgs,
+                                StringRef Arch, StringRef Target,
+                                bool isBitCodeSDL, bool postClangLink);
+void AddStaticDeviceLibsPostLinking(const Driver &D,
+                                    const llvm::opt::ArgList &DriverArgs,
+                                    llvm::opt::ArgStringList &CmdArgs,
+                                    StringRef Arch, StringRef Target,
+                                    bool isBitCodeSDL, bool postClangLink);
+void AddStaticDeviceLibs(Compilation *C, const Tool *T, const JobAction *JA,
+                         const InputInfoList *Inputs, const Driver &D,
+                         const llvm::opt::ArgList &DriverArgs,
+                         llvm::opt::ArgStringList &CmdArgs, StringRef Arch,
+                         StringRef Target, bool isBitCodeSDL,
+                         bool postClangLink);
+
+bool SDLSearch(const Driver &D, const llvm::opt::ArgList &DriverArgs,
+               llvm::opt::ArgStringList &CmdArgs,
+               SmallVector<std::string, 8> LibraryPaths, std::string Lib,
+               StringRef Arch, StringRef Target, bool isBitCodeSDL,
+               bool postClangLink);
+
+bool GetSDLFromOffloadArchive(Compilation &C, const Driver &D, const Tool &T,
+                              const JobAction &JA, const InputInfoList &Inputs,
+                              const llvm::opt::ArgList &DriverArgs,
+                              llvm::opt::ArgStringList &CC1Args,
+                              SmallVector<std::string, 8> LibraryPaths,
+                              StringRef Lib, StringRef Arch, StringRef Target,
+                              bool isBitCodeSDL, bool postClangLink);
+
 const char *SplitDebugName(const JobAction &JA, const llvm::opt::ArgList &Args,
                            const InputInfo &Input, const InputInfo &Output);
 
 void SplitDebugInfo(const ToolChain &TC, Compilation &C, const Tool &T,
                     const JobAction &JA, const llvm::opt::ArgList &Args,
                     const InputInfo &Output, const char *OutFile);
 
 void addLTOOptions(const ToolChain &ToolChain, const llvm::opt::ArgList &Args,

clang/lib/Driver/ToolChains/CommonArgs.cpp

 #include "clang/Driver/InputInfo.h"
 #include "clang/Driver/Job.h"
 #include "clang/Driver/Options.h"
 #include "clang/Driver/SanitizerArgs.h"
 #include "clang/Driver/ToolChain.h"
 #include "clang/Driver/Util.h"
 #include "clang/Driver/XRayArgs.h"
 #include "llvm/ADT/STLExtras.h"
+#include "llvm/ADT/SmallSet.h"
 #include "llvm/ADT/SmallString.h"
 #include "llvm/ADT/StringExtras.h"
 #include "llvm/ADT/StringSwitch.h"
 #include "llvm/ADT/Twine.h"
 #include "llvm/Config/llvm-config.h"
 #include "llvm/Option/Arg.h"
 #include "llvm/Option/ArgList.h"
 #include "llvm/Option/Option.h"
if (Value.getAsInteger(10, PrefixSize)) {
       D.Diag(diag::err_drv_invalid_argument_to_option)
           << Value << A->getOption().getName();
     } else {
       addArg("-x86-pad-max-prefix-size=" + Twine(PrefixSize));
     }
   }
 }
 
+/// SDLSearch: Search for Static Device Library
+/// The search for SDL bitcode files is consistent with how static host
+/// libraries are discovered. That is, the -l option triggers a search for
+/// files in a set of directories called the LINKPATH. The host library search
+/// procedure looks for a specific filename in the LINKPATH.  The filename for
+/// a host library is lib<libname>.a or lib<libname>.so. For SDLs, there is an
+/// ordered-set of filenames that are searched. We call this ordered-set of
+/// filenames as SEARCH-ORDER. Since an SDL can either be device-type specific,
+/// architecture specific, or generic across all architectures, a naming
+/// convention and search order is used where the file name embeds the
+/// architecture name <arch-name> (nvptx or amdgcn) and the GPU device type
+/// <device-name> such as sm_30 and gfx906. <device-name> is absent in case of
+/// device-independent SDLs. To reduce congestion in host library directories,
+/// the search first looks for files in the “libdevice” subdirectory. SDLs that
+/// are bc files begin with the prefix “lib”.
+///
+/// Machine-code SDLs can also be managed as an archive (*.a file). The
+/// convention has been to use the prefix “lib”. To avoid confusion with host
+/// archive libraries, we use prefix "libbc-" for the bitcode SDL archives.
+///
+bool tools::SDLSearch(const Driver &D, const llvm::opt::ArgList &DriverArgs,
+                      llvm::opt::ArgStringList &CC1Args,
+                      SmallVector<std::string, 8> LibraryPaths, std::string Lib,
+                      StringRef Arch, StringRef Target, bool isBitCodeSDL,
+                      bool postClangLink) {
+  SmallVector<std::string, 12> SDLs;
+
+  std::string LibDeviceLoc = "/libdevice";
+  std::string LibBcPrefix = "/libbc-";
+  std::string LibPrefix = "/lib";
+
+  if (isBitCodeSDL) {
+    // SEARCH-ORDER for Bitcode SDLs:
+    //       libdevice/libbc-<libname>-<arch-name>-<device-type>.a
+    //       libbc-<libname>-<arch-name>-<device-type>.a
+    //       libdevice/libbc-<libname>-<arch-name>.a
+    //       libbc-<libname>-<arch-name>.a
+    //       libdevice/libbc-<libname>.a
+    //       libbc-<libname>.a
+    //       libdevice/lib<libname>-<arch-name>-<device-type>.bc
+    //       lib<libname>-<arch-name>-<device-type>.bc
+    //       libdevice/lib<libname>-<arch-name>.bc
+    //       lib<libname>-<arch-name>.bc
+    //       libdevice/lib<libname>.bc
+    //       lib<libname>.bc
+
+    for (StringRef Base : {LibBcPrefix, LibPrefix}) {
+      const auto *Ext = Base.contains(LibBcPrefix) ? ".a" : ".bc";
+
+      for (auto Suffix : {Twine(Lib + "-" + Arch + "-" + Target).str(),
+                          Twine(Lib + "-" + Arch).str(), Twine(Lib).str()}) {
+        SDLs.push_back(Twine(LibDeviceLoc + Base + Suffix + Ext).str());
+        SDLs.push_back(Twine(Base + Suffix + Ext).str());
+      }
+    }
+  } else {
+    // SEARCH-ORDER for Machine-code SDLs:
+    //    libdevice/lib<libname>-<arch-name>-<device-type>.a
+    //    lib<libname>-<arch-name>-<device-type>.a
+    //    libdevice/lib<libname>-<arch-name>.a
+    //    lib<libname>-<arch-name>.a
+
+    const auto *Ext = ".a";
+
+    for (auto Suffix : {Twine(Lib + "-" + Arch + "-" + Target).str(),
+                        Twine(Lib + "-" + Arch).str()}) {
+      SDLs.push_back(Twine(LibDeviceLoc + LibPrefix + Suffix + Ext).str());
+      SDLs.push_back(Twine(LibPrefix + Suffix + Ext).str());
+    }
+  }
+
+  // The CUDA toolchain does not use a global device llvm-link before the LLVM
+  // backend generates ptx. So currently, the use of bitcode SDL for nvptx is
+  // only possible with post-clang-cc1 linking. Clang cc1 has a feature that
+  // will link libraries after clang compilation while the LLVM IR is still in
+  // memory. This utilizes a clang cc1 option called “-mlink-builtin-bitcode”.
+  // This is a clang -cc1 option that is generated by the clang driver. The
+  // option value must a full path to an existing file.
+  bool FoundSDL = false;
+  for (auto LPath : LibraryPaths) {
+    for (auto SDL : SDLs) {
+      auto FullName = Twine(LPath + SDL).str();
+      if (llvm::sys::fs::exists(FullName)) {
+        if (postClangLink)
+          CC1Args.push_back("-mlink-builtin-bitcode");
+        CC1Args.push_back(DriverArgs.MakeArgString(FullName));
+        FoundSDL = true;
+        break;
+      }
+    }
+    if (FoundSDL)
+      break;
+  }
+  return FoundSDL;
+}
+
+/// Search if a user provided archive file lib<libname>.a exists in any of
+/// the library paths. If so, add a new command to clang-offload-bundler to
+/// unbundle this archive and create a temporary device specific archive. Name
+/// of this SDL is passed to the llvm-link (for amdgcn) or to the
+/// clang-nvlink-wrapper (for nvptx) commands by the driver.
+bool tools::GetSDLFromOffloadArchive(
+    Compilation &C, const Driver &D, const Tool &T, const JobAction &JA,
+    const InputInfoList &Inputs, const llvm::opt::ArgList &DriverArgs,
+    llvm::opt::ArgStringList &CC1Args, SmallVector<std::string, 8> LibraryPaths,
+    StringRef Lib, StringRef Arch, StringRef Target, bool isBitCodeSDL,
+    bool postClangLink) {
+
+  // We don't support bitcode archive bundles for nvptx
+  if (isBitCodeSDL && Arch.contains("nvptx"))
+    return false;
+
+  bool FoundAOB = false;
+  SmallVector<std::string, 2> AOBFileNames;
+  std::string ArchiveOfBundles;
+  for (auto LPath : LibraryPaths) {
+    ArchiveOfBundles.clear();
+
+    AOBFileNames.push_back(Twine(LPath + "/libdevice/lib" + Lib + ".a").str());
+    AOBFileNames.push_back(Twine(LPath + "/lib" + Lib + ".a").str());
+
+    for (auto AOB : AOBFileNames) {
+      if (llvm::sys::fs::exists(AOB)) {
+        ArchiveOfBundles = AOB;
+        FoundAOB = true;
+        break;
+      }
+    }
+
+    if (!FoundAOB)
+      continue;
+
+    StringRef Prefix = isBitCodeSDL ? "libbc-" : "lib";
+    std::string OutputLib = D.GetTemporaryPath(
+        Twine(Prefix + Lib + "-" + Arch + "-" + Target).str(), "a");
+
+    C.addTempFile(C.getArgs().MakeArgString(OutputLib.c_str()));
+
+    ArgStringList CmdArgs;
+    SmallString<128> DeviceTriple;
+    DeviceTriple += Action::GetOffloadKindName(JA.getOffloadingDeviceKind());
+    DeviceTriple += '-';
+    std::string NormalizedTriple = T.getToolChain().getTriple().normalize();
+    DeviceTriple += NormalizedTriple;
+    if (!Target.empty()) {
+      DeviceTriple += '-';
+      DeviceTriple += Target;
+    }
+
+    std::string UnbundleArg("-unbundle");
+    std::string TypeArg("-type=a");
+    std::string InputArg("-inputs=" + ArchiveOfBundles);
+    std::string OffloadArg("-targets=" + std::string(DeviceTriple));
+    std::string OutputArg("-outputs=" + OutputLib);
+
+    const char *UBProgram = DriverArgs.MakeArgString(
+        T.getToolChain().GetProgramPath("clang-offload-bundler"));
+
+    ArgStringList UBArgs;
+    UBArgs.push_back(C.getArgs().MakeArgString(UnbundleArg.c_str()));
+    UBArgs.push_back(C.getArgs().MakeArgString(TypeArg.c_str()));
+    UBArgs.push_back(C.getArgs().MakeArgString(InputArg.c_str()));
+    UBArgs.push_back(C.getArgs().MakeArgString(OffloadArg.c_str()));
+    UBArgs.push_back(C.getArgs().MakeArgString(OutputArg.c_str()));
+
+    // Add this flag to not exit from clang-offload-bundler if no compatible
+    // code object is found in heterogenous archive library.
+    std::string AdditionalArgs("-allow-missing-bundles");
+    UBArgs.push_back(C.getArgs().MakeArgString(AdditionalArgs.c_str()));
+
+    C.addCommand(std::make_unique<Command>(
+        JA, T, ResponseFileSupport::AtFileCurCP(), UBProgram, UBArgs, Inputs,
+        InputInfo(&JA, C.getArgs().MakeArgString(OutputLib.c_str()))));
+    if (postClangLink)
+      CC1Args.push_back("-mlink-builtin-bitcode");
+
+    CC1Args.push_back(DriverArgs.MakeArgString(OutputLib));
+    break;
+  }
+
+  return FoundAOB;
+}
+
+// Wrapper function used by driver for adding SDLs during link phase.
+void tools::AddStaticDeviceLibsLinking(Compilation &C, const Tool &T,
+                                const JobAction &JA,
+                                const InputInfoList &Inputs,
+                                const llvm::opt::ArgList &DriverArgs,
+                                llvm::opt::ArgStringList &CC1Args,
+                                StringRef Arch, StringRef Target,
+                                bool isBitCodeSDL, bool postClangLink) {
+  AddStaticDeviceLibs(&C, &T, &JA, &Inputs, C.getDriver(), DriverArgs, CC1Args,
+                      Arch, Target, isBitCodeSDL, postClangLink);
+}
+
+// Wrapper function used for post clang linking of bitcode SDLS for nvptx by
+// the CUDA toolchain.
+void tools::AddStaticDeviceLibsPostLinking(const Driver &D,
+                                const llvm::opt::ArgList &DriverArgs,
+                                llvm::opt::ArgStringList &CC1Args,
+                                StringRef Arch, StringRef Target,
+                                bool isBitCodeSDL, bool postClangLink) {
+  AddStaticDeviceLibs(nullptr, nullptr, nullptr, nullptr, D, DriverArgs,
+                      CC1Args, Arch, Target, isBitCodeSDL, postClangLink);
+}
+
+// User defined Static Device Libraries(SDLs) can be passed to clang for
+// offloading GPU compilers. Like static host libraries, the use of a SDL is
+// specified with the -l command line option. The primary difference between
+// host and SDLs is the filenames for SDLs (refer SEARCH-ORDER for Bitcode SDLs
+// and SEARCH-ORDER for Machine-code SDLs for the naming convention).
+// SDLs are of following types:
+//
+// * Bitcode SDLs: They can either be a *.bc file or an archive of *.bc files.
+//           For NVPTX, these libraries are post-clang linked following each
+//           compilation. For AMDGPU, these libraries are linked one time
+//           during the application link phase.
+//
+// * Machine-code SDLs: They are archive files. For NVPTX, the archive members
+//           contain cubin for Nvidia GPUs and are linked one time during the
+//           link phase by the CUDA SDK linker called nvlink. For AMDGPU, the
+//           process for machine code SDLs is still in development. But they
+//           will be linked by the LLVM tool lld.
+//
+// * Bundled objects that contain both host and device codes: Bundled objects
+//           may also contain library code compiled from source. For NVPTX, the
+//           bundle contains cubin. For AMDGPU, the bundle contains bitcode.
+//
+// For Bitcode and Machine-code SDLs, current compiler toolchains hardcode the
+// inclusion of specific SDLs such as math libraries and the OpenMP device
+// library libomptarget.
+void tools::AddStaticDeviceLibs(Compilation *C, const Tool *T,
+                                const JobAction *JA,
+                                const InputInfoList *Inputs, const Driver &D,
+                                const llvm::opt::ArgList &DriverArgs,
+                                llvm::opt::ArgStringList &CC1Args,
+                                StringRef Arch, StringRef Target,
+                                bool isBitCodeSDL, bool postClangLink) {
+
+  SmallVector<std::string, 8> LibraryPaths;
+  // Add search directories from LIBRARY_PATH env variable
+  llvm::Optional<std::string> LibPath =
+      llvm::sys::Process::GetEnv("LIBRARY_PATH");
+  if (LibPath) {
+    SmallVector<StringRef, 8> Frags;
+    const char EnvPathSeparatorStr[] = {llvm::sys::EnvPathSeparator, '\0'};
+    llvm::SplitString(*LibPath, Frags, EnvPathSeparatorStr);
+    for (StringRef Path : Frags)
+      LibraryPaths.emplace_back(Path.trim());
+  }
+
+  // Add directories from user-specified -L options
+  for (std::string Search_Dir : DriverArgs.getAllArgValues(options::OPT_L))
+    LibraryPaths.emplace_back(Search_Dir);
+
+  // Add path to lib-debug folders
+  SmallString<256> DefaultLibPath = llvm::sys::path::parent_path(D.Dir);
+  llvm::sys::path::append(DefaultLibPath, Twine("lib") + CLANG_LIBDIR_SUFFIX);
+  LibraryPaths.emplace_back(DefaultLibPath.c_str());
+
+  // Build list of Static Device Libraries SDLs specified by -l option
+  llvm::SmallSet<std::string, 16> SDLNames;
+  static const StringRef HostOnlyArchives[] = {
+      "omp", "cudart", "m", "gcc", "gcc_s", "pthread", "hip_hcc"};
+  for (auto SDLName : DriverArgs.getAllArgValues(options::OPT_l)) {
+    if (!HostOnlyArchives->contains(SDLName)) {
+      SDLNames.insert(SDLName);
+    }
+  }
+
+  // The search stops as soon as an SDL file is found. The driver then provides
+  // the full filename of the SDL to the llvm-link or clang-nvlink-wrapper
+  // command. If no SDL is found after searching each LINKPATH with
+  // SEARCH-ORDER, it is possible that an archive file lib<libname>.a exists
+  // and may contain bundled object files.
+  for (auto SDLName : SDLNames) {
+    // This is the only call to SDLSearch
+    if (!SDLSearch(D, DriverArgs, CC1Args, LibraryPaths, SDLName, Arch, Target,
+                   isBitCodeSDL, postClangLink)) {
+      GetSDLFromOffloadArchive(*C, D, *T, *JA, *Inputs, DriverArgs, CC1Args,
+                               LibraryPaths, SDLName, Arch, Target,
+                               isBitCodeSDL, postClangLink);
+    }
+  }
+}
+
 static llvm::opt::Arg *
 getAMDGPUCodeObjectArgument(const Driver &D, const llvm::opt::ArgList &Args) {
   // The last of -mcode-object-v3, -mno-code-object-v3 and
   // -mcode-object-version=<version> wins.
   return Args.getLastArg(options::OPT_mcode_object_v3_legacy,
                          options::OPT_mno_code_object_v3_legacy,
                          options::OPT_mcode_object_version_EQ);
 }

clang/lib/Driver/ToolChains/Cuda.cpp

if (!II.isFilename())
       continue;
 
     const char *CubinF = C.addTempFile(
         C.getArgs().MakeArgString(getToolChain().getInputFilename(II)));
 
     CmdArgs.push_back(CubinF);
   }
 
+  AddStaticDeviceLibsLinking(C, *this, JA, Inputs, Args, CmdArgs, "nvptx", GPUArch,
+                      false, false);
+
   const char *Exec =
-      Args.MakeArgString(getToolChain().GetProgramPath("nvlink"));
+      Args.MakeArgString(getToolChain().GetProgramPath("clang-nvlink-wrapper"));
   C.addCommand(std::make_unique<Command>(
       JA, *this,
       ResponseFileSupport{ResponseFileSupport::RF_Full, llvm::sys::WEM_UTF8,
                           "--options-file"},
       Exec, CmdArgs, Inputs, Output));
 }
 
 /// CUDA toolchain.  Our assembler is ptxas, and our "linker" is fatbinary,
if (DeviceOffloadingKind == Action::OFK_OpenMP) {
     if (DriverArgs.hasFlag(options::OPT_fopenmp_target_new_runtime,
                            options::OPT_fno_openmp_target_new_runtime, false))
       BitcodeSuffix = "new-nvptx-" + GpuArch.str();
     else
       BitcodeSuffix = "nvptx-" + GpuArch.str();
 
     addOpenMPDeviceRTL(getDriver(), DriverArgs, CC1Args, BitcodeSuffix,
                        getTriple());
+    AddStaticDeviceLibsPostLinking(getDriver(), DriverArgs, CC1Args, "nvptx", GpuArch,
+                        /* bitcode SDL?*/ true, /* PostClang Link? */ true);
   }
 }
 
 llvm::DenormalMode CudaToolChain::getDefaultDenormalModeForType(
     const llvm::opt::ArgList &DriverArgs, const JobAction &JA,
     const llvm::fltSemantics *FPType) const {
   if (JA.getOffloadingDeviceKind() == Action::OFK_Cuda) {
     if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&

clang/test/Driver/fat_archive_amdgpu.cpp

+// REQUIRES: clang-driver
+// REQUIRES: x86-registered-target
+// REQUIRES: amdgpu-registered-target
+
+// See the steps to create a fat archive are given at the end of the file.
+
+// Given a FatArchive, clang-offload-bundler should be called to create a
+// device specific archive, which should be passed to llvm-link.
+// RUN: %clang -O2 -### -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 %s -L%S/Inputs/openmp_static_device_link -lFatArchive 2>&1 | FileCheck %s
+// CHECK: clang{{.*}}"-cc1"{{.*}}"-triple" "amdgcn-amd-amdhsa"{{.*}}"-emit-llvm-bc"{{.*}}"-target-cpu" "[[GPU:gfx[0-9]+]]"{{.*}}"-o" "[[HOSTBC:.*.bc]]" "-x" "c++"{{.*}}.cpp
+// CHECK: clang-offload-bundler" "-unbundle" "-type=a" "-inputs={{.*}}/Inputs/openmp_static_device_link/libFatArchive.a" "-targets=openmp-amdgcn-amd-amdhsa-[[GPU]]" "-outputs=[[DEVICESPECIFICARCHIVE:.*.a]]" "-allow-missing-bundles"
+// CHECK: llvm-link{{.*}}"[[HOSTBC]]" "[[DEVICESPECIFICARCHIVE]]" "-o" "{{.*}}-[[GPU]]-linked-{{.*}}.bc"
+// CHECK: ld"{{.*}}" "-L{{.*}}/Inputs/openmp_static_device_link" "{{.*}} "-lFatArchive" "{{.*}}" "-lomp{{.*}}-lomptarget"
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+#define N 10
+
+#pragma omp declare target
+// Functions defined in Fat Archive.
+extern "C" void func_present(float *, float *, unsigned);
+
+#ifdef MISSING
+// Function not defined in the fat archive.
+extern "C" void func_missing(float *, float *, unsigned);
+#endif
+
+#pragma omp end declare target
+
+int main() {
+  float in[N], out[N], sum = 0;
+  unsigned i;
+
+#pragma omp parallel for
+  for (i = 0; i < N; ++i) {
+    in[i] = i;
+  }
+
+  func_present(in, out, N); // Returns out[i] = a[i] * 0
+
+#ifdef MISSING
+  func_missing(in, out, N); // Should throw an error here
+#endif
+
+#pragma omp parallel for reduction(+ \
+                                   : sum)
+  for (i = 0; i < N; ++i)
+    sum += out[i];
+
+  if (!sum)
+    return 0;
+  return sum;
+}
+
+#endif
+
+/***********************************************
+   Steps to create Fat Archive (libFatArchive.a)
+************************************************
+***************** File: func_1.c ***************
+void func_present(float* in, float* out, unsigned n){
+  unsigned i;
+  #pragma omp target teams distribute parallel for map(to: in[0:n]) map(from: out[0:n])
+  for(i=0; i<n; ++i){
+    out[i] = in[i] * 0;
+  }
+}
+*************************************************
+1. Compile source file(s) to generate object file(s)
+    clang -O2 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 -c func_1.c -o func_1_gfx906.o
+    clang -O2 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx908 -c func_1.c -o func_1_gfx908.o
+    clang -O2 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 -c func_2.c -o func_2_gfx906.o
+    clang -O2 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx908 -c func_2.c -o func_2_gfx908.o
+    clang -O2 -fopenmp -fopenmp-targets=nvptx64 -c func_1.c -o func_1_nvptx.o
+    clang -O2 -fopenmp -fopenmp-targets=nvptx64 -c func_2.c -o func_2_nvptx.o
+
+2. Create a fat archive by combining all the object file(s)
+    llvm-ar cr libFatArchive.a func_1_gfx906.o func_1_gfx908.o func_2_gfx906.o func_2_gfx908.o func_1_nvptx.o func_2_nvptx.o
+************************************************/

clang/test/Driver/fat_archive_nvptx.cpp

+// REQUIRES: clang-driver
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// See the steps to create a fat archive are given at the end of the file.
+
+// Given a FatArchive, clang-offload-bundler should be called to create a
+// device specific archive, which should be passed to clang-nvlink-wrapper.
+// RUN: %clang -O2 -### -fopenmp -fopenmp-targets=nvptx64 %s -L%S/Inputs/openmp_static_device_link -lFatArchive 2>&1 | FileCheck %s
+// CHECK: clang{{.*}}"-cc1"{{.*}}"-triple" "nvptx64"{{.*}}"-target-cpu" "[[GPU:sm_[0-9]+]]"{{.*}}"-o" "[[HOSTBC:.*.s]]" "-x" "c++"{{.*}}.cpp
+// CHECK: clang-offload-bundler" "-unbundle" "-type=a" "-inputs={{.*}}/Inputs/openmp_static_device_link/libFatArchive.a" "-targets=openmp-nvptx64-[[GPU]]" "-outputs=[[DEVICESPECIFICARCHIVE:.*.a]]" "-allow-missing-bundles"
+// CHECK: clang-nvlink-wrapper{{.*}}"-o" "{{.*}}.out" "-arch" "[[GPU]]" "{{.*}}[[DEVICESPECIFICARCHIVE]]"
+// CHECK: ld"{{.*}}" "-L{{.*}}/Inputs/openmp_static_device_link" "{{.*}} "-lFatArchive" "{{.*}}" "-lomp{{.*}}-lomptarget"
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+#define N 10
+
+#pragma omp declare target
+// Functions defined in Fat Archive.
+extern "C" void func_present(float *, float *, unsigned);
+
+#ifdef MISSING
+// Function not defined in the fat archive.
+extern "C" void func_missing(float *, float *, unsigned);
+#endif
+
+#pragma omp end declare target
+
+int main() {
+  float in[N], out[N], sum = 0;
+  unsigned i;
+
+#pragma omp parallel for
+  for (i = 0; i < N; ++i) {
+    in[i] = i;
+  }
+
+  func_present(in, out, N); // Returns out[i] = a[i] * 0
+
+#ifdef MISSING
+  func_missing(in, out, N); // Should throw an error here
+#endif
+
+#pragma omp parallel for reduction(+ \
+                                   : sum)
+  for (i = 0; i < N; ++i)
+    sum += out[i];
+
+  if (!sum)
+    return 0;
+  return sum;
+}
+
+#endif
+
+/***********************************************
+   Steps to create Fat Archive (libFatArchive.a)
+************************************************
+***************** File: func_1.c ***************
+void func_present(float* in, float* out, unsigned n){
+  unsigned i;
+  #pragma omp target teams distribute parallel for map(to: in[0:n]) map(from: out[0:n])
+  for(i=0; i<n; ++i){
+    out[i] = in[i] * 0;
+  }
+}
+*************************************************
+1. Compile source file(s) to generate object file(s)
+    clang -O2 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 -c func_1.c -o func_1_gfx906.o
+    clang -O2 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx908 -c func_1.c -o func_1_gfx908.o
+    clang -O2 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 -c func_2.c -o func_2_gfx906.o
+    clang -O2 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx908 -c func_2.c -o func_2_gfx908.o
+    clang -O2 -fopenmp -fopenmp-targets=nvptx64 -c func_1.c -o func_1_nvptx.o
+    clang -O2 -fopenmp -fopenmp-targets=nvptx64 -c func_2.c -o func_2_nvptx.o
+
+2. Create a fat archive by combining all the object file(s)
+    llvm-ar cr libFatArchive.a func_1_gfx906.o func_1_gfx908.o func_2_gfx906.o func_2_gfx908.o func_1_nvptx.o func_2_nvptx.o
+************************************************/

clang/tools/clang-offload-bundler/ClangOffloadBundler.cpp

return OffloadKind == Target.OffloadKind &&
            Triple.isCompatibleWith(Target.Triple) && GPUArch == Target.GPUArch;
   }
 
   std::string str() {
     return Twine(OffloadKind + "-" + Triple.str() + "-" + GPUArch).str();
   }
 };
 
+static StringRef getDeviceFileExtension(StringRef Device) {
+  if (Device.contains("gfx"))
+    return ".bc";
+  if (Device.contains("sm_"))
+    return ".cubin";
+
+  WithColor::warning() << "Could not determine extension for archive"
+                          "members, using \".o\"\n";
+  return ".o";
+}
+
+static std::string getDeviceLibraryFileName(StringRef BundleFileName,
+                                            StringRef Device) {
+  StringRef LibName = sys::path::stem(BundleFileName);
+  StringRef Extension = getDeviceFileExtension(Device);
+
+  std::string Result;
+  Result += LibName;
+  Result += Extension;
+  return Result;
+}
+
 /// Generic file handler interface.
 class FileHandler {
 public:
   struct BundleInfo {
     StringRef BundleID;
   };
 
   FileHandler() {}
while (!CodeObject.empty()) {
                 FileHandler.get()->ReadBundle(DataStream, *CodeObjectBuffer))
           return Err;
 
         for (auto &CompatibleTarget : CompatibleTargets) {
           SmallString<128> BundledObjectFileName;
           BundledObjectFileName.assign(BundledObjectFile);
           auto OutputBundleName =
               Twine(llvm::sys::path::stem(BundledObjectFileName) + "-" +
-                    CodeObject)
+                    CodeObject +
+                    getDeviceLibraryFileName(BundledObjectFileName,
+                                             CodeObjectInfo.GPUArch))
                   .str();
           // Replace ':' in optional target feature list with '_' to ensure
           // cross-platform validity.
           std::replace(OutputBundleName.begin(), OutputBundleName.end(), ':',
                        '_');
 
           std::unique_ptr<MemoryBuffer> MemBuf = MemoryBuffer::getMemBufferCopy(
               DataStream.str(), OutputBundleName);