[CUDA] __global__ kernels pre-declared in header results into nv-link error in separable compilation (gpu-rdc)

[geomois]

When enabling compilation with relocatable device code flag (--gpu-rdc), with cuda code with header and implementation files, I ran into a linking error:

nvlink error: Function A() not declared __global__ in all source files 
nvlink fatal: merge_elf failed

The error can be reproduced with the following set-up:

// bug.cuh
__global__ void foo();

// bug.cu
__global__ void foo() {}

// main.cu
#include "bug.cuh"

int main() { foo<<<1, 1, 1>>>(); }

>clang++ bug.cu main.cu --offload-arch=sm_89 -fgpu-rdc --offload-new-driver -lcudart

@jhuber6's insight:
Solution is to have .extern .entry instead of .extern .func

>clang++ -x cuda main.cu --offload-arch=sm_89 --offload-device-only -o - -S -fgpu-rdc

.version 8.5
.target sm_89
.address_size 64

.extern .func _Z3foov
()
;
.global .align 8 .u64 __clang_gpu_used_external[1] = {_Z3foov};

Workaround, enable -foffload-lto.

[Artem-B]

Right now it's not clear where the things went wrong.

It would help to include the PTX produced in the broken case. Maybe, relevant IR, too. Including compiler output with -v would also be useful.
In short, re-do the compilation with --save-temps=obj -v and post the file with intermediate compiler outputs and whatever is printed on the console.

[jhuber6]

Right now it's not clear where the things went wrong.

It would help to include the PTX produced in the broken case. Maybe, relevant IR, too. Including compiler output with -v would also be useful. In short, re-do the compilation with --save-temps=obj -v and post the file with intermediate compiler outputs and whatever is printed on the console.

Pretty sure the solution is just to emit kernel metadata on function declarations. If you check the output PTX we have this, which as you can see does not declare this as a kernel. Thus when nvlink tries to link it, it gets confused.

> clang++ -x cuda main.cu --offload-arch=sm_89 --offload-device-only -o - -S -fgpu-rdc

.version 8.5
.target sm_89
.address_size 64

.extern .func _Z3foov
()
;
.global .align 8 .u64 __clang_gpu_used_external[1] = {_Z3foov};```

[Artem-B]

Pretty sure the solution...

It sounds plausible. I'm just pointing out that the bug, as filed, could use additional info info to support that conclusion.

Right now I can't get clang to generate any IR for a kernel declaration at all, so it's not clear where the .extern .func _Z3foov is coming from.
https://godbolt.org/z/hnqqaWGGc

[jhuber6]

Pretty sure the solution...

It sounds plausible. I'm just pointing out that the bug, as filed, could use additional info info to support that conclusion.

Right now I can't get clang to generate any IR for a kernel declaration at all, so it's not clear where the .extern .func _Z3foov is coming from. https://godbolt.org/z/hnqqaWGGc

You need a call https://godbolt.org/z/vEcz5ohb3.

[Artem-B]

I think in this case the problem is actually the fact that we attempt to put the kernel reference into __clang_gpu_used_external even though there's no definition of that kernel here. The reference should be created in the same TU whenre the kernel itself lives, IMO.

Emitting .func for the kernel is a problem, but it's a secondary failure here. We would probably need to fix that if/when we implement GPU-side kernel launches, but we currently do not have them.

Clang currently does not allow taking kernel referencesfor any other purpose, so I can't think of a way for the user to create such a reference explicitly:
https://godbolt.org/z/vaYn4b3zT

[jhuber6]

I think in this case the problem is actually the fact that we attempt to put the kernel reference into __clang_gpu_used_external even though there's no definition of that kernel here. The reference should be created in the same TU whenre the kernel itself lives, IMO.

Emitting .func for the kernel is a problem, but it's a secondary failure here. We would probably need to fix that if/when we implement GPU-side kernel launches, but we currently do not have them.

Clang currently does not allow taking kernel referencesfor any other purpose, so I can't think of a way for the user to create such a reference explicitly: https://godbolt.org/z/vaYn4b3zT

Fixing the lack of the attribute is trivial, just stop returning early in

llvm-project/clang/lib/CodeGen/Targets/NVPTX.cpp

Line 236 in 5f881b7

if (GV->isDeclaration())

. That's probably something that should be fixed anyway.

[Artem-B]

Fixing the lack of the attribute is trivial, just stop returning early in

I'm fine doing that along with not emitting __clang_gpu_used_external entries that are not actually used by the TU.