[llvm-bugs] [Bug 49555] New: CUDA separable compilation support

[via llvm-bugs]

https://bugs.llvm.org/show_bug.cgi?id=49555

            Bug ID: 49555
           Summary: CUDA separable compilation support
           Product: clang
           Version: trunk
          Hardware: PC
                OS: Linux
            Status: NEW
          Severity: enhancement
          Priority: P
         Component: CUDA
          Assignee: unassignedclangbugs at nondot.org
          Reporter: carlosgalvezp at gmail.com
                CC: llvm-bugs at lists.llvm.org

Hi!

Since long time ago CUDA supports separable compilation:
https://developer.nvidia.com/blog/separate-compilation-linking-cuda-device-code/

Is this expected to be supported in Clang? The following example doesn't work:

// lib.h
__device__ int foo();

// lib.cu
#include "lib.h"
__device__ int foo()
{
    return 123;
}

// main.cu
#include "lib.h"

#include <cstdio>

__global__ kernel()
{
    printf("%d\n", foo());
}

int main()
{
   kernel<<<1,10>>>();
   cudaDeviceSynchronize();
}


Compiling with:

clang++ ... lib.cu main.cu

Gives a linker error: cannot find reference to extern function "foo".

If this is not supported out of the box, would the following procedure be
reasonable?

1) Compile object files with clang, adding -fgpu-rdc.
2) Device link them using "nvcc -dlink *.o -o device_link.o"
3) Link everything with clang: "clang *.o device_link.o"

Thanks!

-- 
You are receiving this mail because:
You are on the CC list for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-bugs/attachments/20210312/78dc579f/attachment.html>