Undefined symbol when trying to link with shared library built from CUDA objects

Question

I'm experimenting with building a simple application from a couple of .cu source files and a very simple C++ main that calls a function from one of the .cu files. I'm making a shared library (.so file) from the compiled .cu files. I'm finding that everything builds without trouble, but when I try to run the application, I get a linker undefined symbol error, with the mangled name of the .cu function I'm calling from main(). If I build a static library instead, my application runs just fine. Here's the makefile I've set up:

.PHONY: clean
NVCCFLAGS = -std=c++11 --compiler-options '-fPIC'
CXXFLAGS = -std=c++11
HLIB = libhello.a
SHLIB = libhello.so
CUDA_OBJECTS = bridge.o add.o

all: driver

%.o :: %.cu
    nvcc -o $@ $(NVCCFLAGS) -c -I. $<

%.o :: %.cpp
    c++ $(CXXFLAGS) -o $@ -c -I. $<

$(HLIB): $(CUDA_OBJECTS)
    ar rcs $@ $^

$(SHLIB): $(CUDA_OBJECTS)
    nvcc $(NVCCFLAGS) --shared  -o $@ $^

#driver : driver.o $(HLIB)
#   c++ -std=c++11 -fPIC -o $@ driver.o -L. -lhello -L/usr/local/cuda-10.1/targets/x86_64-linux/lib -lcudart

driver : driver.o $(SHLIB)
    c++ -std=c++11 -fPIC -o $@ driver.o -L. -lhello

clean:
    -rm -f driver *.o *.so *.a

Here are the various source files that the makefile takes as fodder. add.cu:

__global__ void add(int n, int* a, int* b, int* c) {
    int index = threadIdx.x;
    int stride = blockDim.x;

    for (int ii = index; ii < n; ii += stride) {
        c[ii] = a[ii] + b[ii];
    }
}

add.h:

extern __global__ void add(int n, int* a, int* b, int* c);

bridge.cu:

#include <iostream>
#include "add.h"

void bridge() {
    int N = 1 << 16;
    int blockSize = 256;
    int numBlocks = (N + blockSize - 1)/blockSize;

    int* a;
    int* b;
    int* c;

    cudaMallocManaged(&a, N*sizeof(int));
    cudaMallocManaged(&b, N*sizeof(int));
    cudaMallocManaged(&c, N*sizeof(int));

    for (int ii = 0; ii < N; ii++) {
        a[ii] = ii;
        b[ii] = 2*ii;
    }

    add<<<numBlocks, blockSize>>>(N, a, b, c);

    cudaDeviceSynchronize();

    for (int ii = 0; ii < N; ii++) {
        std::cout << a[ii] << " + " << b[ii] << " = " << c[ii] << std::endl;
    }

    cudaFree(a);
    cudaFree(b);
    cudaFree(c);
}

bridge.h:

extern void bridge();

driver.cpp:

#include "bridge.h"

int main() {
    bridge();
    return 0;
}

I'm very new to cuda, so I expect that's where I'm doing something wrong. I've played a bit with using extern "C" declarations, but that just seems to move the "undefined symbol" error from run time to build time.

I'm familiar with various ways that one can end up with an undefined symbol, and I've mentioned various experiments I've already performed (static linking, extern "C" declarations) that make me think that this problem isn't addressed by the proposed duplicate question.

My unresolved symbol is _Z6bridgev

It looks to me as though the linker should be able resolve the symbol. If I can nm on driver.o, I see:

0000000000000000 T main
                 U _Z6bridgev

And if I run nm on libhello.so, I see:

0000000000006e56 T _Z6bridgev

Answer 1

When Robert Crovella was able to get my example to work on his machine, while I wasn't able to get his example to work on mine, I started realizing that my problem had nothing to do with cuda or nvcc. It was the fact that with a shared library, the loader has to resolve symbols at runtime, and my shared library wasn't in a "well-known location". I built a simple test case just now, purely with c++ sources, and repeated my failure. Once I copied libhello.so to /usr/local/lib, I was able to run driver successfully. So, I'm OK with closing my original question, if that's the will of the people.