CUDA condiviso linking libreria: undefined reference to cudaRegisterLinkedBinary

Question

Obiettivo:

1. creare una libreria condivisa che contiene i miei kernel CUDA che ha un involucro CUDA-libero/intestazione.
2. creare un eseguibile test per la libreria condivisa.

Problema

1. libreria condivisa MYLIB.so sembra per compilare bene. (nessun problema).
2. errore nel collegamento:

./libMYLIB.so: undefined reference to __cudaRegisterLinkedBinary_39_tmpxft_000018cf_00000000_6_MYLIB_cpp1_ii_74c599a1

makefile semplificata:

libMYlib.so : MYLIB.o 
    g++ -shared -Wl,-soname,libMYLIB.so -o libMYLIB.so MYLIB.o -L/the/cuda/lib/dir -lcudart 


MYLIB.o : MYLIB.cu MYLIB.h 
    nvcc -m64 -arch=sm_20 -dc -Xcompiler '-fPIC' MYLIB.cu -o MYLIB.o -L/the/cuda/lib/dir -lcudart 


test : test.cpp libMYlib.so 
     g++ test.cpp -o test -L. -ldl -Wl,-rpath,. -lMYLIB -L/the/cuda/lib/dir -lcudart 

---

infatti

nm libMYLIB.so mostra che tutti CUDA API funct Gli ioni sono "simboli indefiniti":

  U __cudaRegisterFunction 
     U __cudaRegisterLinkedBinary_39_tmpxft_0000598c_00000000_6_CUPA_cpp1_ii_74c599a1 
     U cudaEventRecord 
     U cudaFree 
     U cudaGetDevice 
     U cudaGetDeviceProperties 
     U cudaGetErrorString 
     U cudaLaunch 
     U cudaMalloc 
     U cudaMemcpy 

Quindi CUDA in qualche modo non ha ottenuto collegato alla libreria condivisa MYLIB.so Che cosa mi manca?

---

CUDA non ha nemmeno ottenere collegato al file oggetto in qualche modo:

nm MYLIB.o

  U __cudaRegisterFunction 
     U __cudaRegisterLinkedBinary_39_tmpxft_0000598c_00000000_6_CUPA_cpp1_ii_74c599a1 
     U cudaEventRecord 
     U cudaFree 
     U cudaGetDevice 
     U cudaGetDeviceProperties 
     U cudaGetErrorString 
     U cudaLaunch 
     U cudaMalloc 
     U cudaMemcpy 

(come sopra)

Answer 1

Ecco una creazione di oggetti condivisi esempio linux lungo le linee hai indicato:

1. crea una libreria condivisa contenente i miei kernel CUDA che ha un wrapper/header CUDA-free.
2. creare un file eseguibile di test per la libreria condivisa.

Prima la libreria condivisa. I comandi di compilazione per questo sono i seguenti:

nvcc -arch=sm_20 -Xcompiler '-fPIC' -dc test1.cu test2.cu 
nvcc -arch=sm_20 -Xcompiler '-fPIC' -dlink test1.o test2.o -o link.o 
g++ -shared -o test.so test1.o test2.o link.o -L/usr/local/cuda/lib64 -lcudart 

Sembra si può mancare il secondo passo sopra nel makefile, ma non mi hanno analizzato se ci sono altri problemi con il tuo makefile.

Ora, per l'eseguibile di prova, i comandi di compilazione sono i seguenti:

g++ -c main.cpp 
g++ -o testmain main.o test.so 

per farlo funzionare, è sufficiente eseguire il file eseguibile testmain, ma essere sicuri biblioteca test.so è sulla vostra LD_LIBRARY_PATH.

Questi sono i file che ho usato per scopi di test:

test1.h:

int my_test_func1(); 

test1.cu:

#include <stdio.h> 
#include "test1.h" 

#define DSIZE 1024 
#define DVAL 10 
#define nTPB 256 

#define cudaCheckErrors(msg) \ 
    do { \ 
     cudaError_t __err = cudaGetLastError(); \ 
     if (__err != cudaSuccess) { \ 
      fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ 
       msg, cudaGetErrorString(__err), \ 
       __FILE__, __LINE__); \ 
      fprintf(stderr, "*** FAILED - ABORTING\n"); \ 
      exit(1); \ 
     } \ 
    } while (0) 

__global__ void my_kernel1(int *data){ 
    int idx = threadIdx.x + (blockDim.x *blockIdx.x); 
    if (idx < DSIZE) data[idx] =+ DVAL; 
} 

int my_test_func1(){ 

    int *d_data, *h_data; 
    h_data = (int *) malloc(DSIZE * sizeof(int)); 
    if (h_data == 0) {printf("malloc fail\n"); exit(1);} 
    cudaMalloc((void **)&d_data, DSIZE * sizeof(int)); 
    cudaCheckErrors("cudaMalloc fail"); 
    for (int i = 0; i < DSIZE; i++) h_data[i] = 0; 
    cudaMemcpy(d_data, h_data, DSIZE * sizeof(int), cudaMemcpyHostToDevice); 
    cudaCheckErrors("cudaMemcpy fail"); 
    my_kernel1<<<((DSIZE+nTPB-1)/nTPB), nTPB>>>(d_data); 
    cudaDeviceSynchronize(); 
    cudaCheckErrors("kernel"); 
    cudaMemcpy(h_data, d_data, DSIZE * sizeof(int), cudaMemcpyDeviceToHost); 
    cudaCheckErrors("cudaMemcpy 2"); 
    for (int i = 0; i < DSIZE; i++) 
    if (h_data[i] != DVAL) {printf("Results check failed at offset %d, data was: %d, should be %d\n", i, h_data[i], DVAL); exit(1);} 
    printf("Results check 1 passed!\n"); 
    return 0; 
} 

test2.h:

int my_test_func2(); 

test2.cu:

#include <stdio.h> 
#include "test2.h" 

#define DSIZE 1024 
#define DVAL 20 
#define nTPB 256 

#define cudaCheckErrors(msg) \ 
    do { \ 
     cudaError_t __err = cudaGetLastError(); \ 
     if (__err != cudaSuccess) { \ 
      fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ 
       msg, cudaGetErrorString(__err), \ 
       __FILE__, __LINE__); \ 
      fprintf(stderr, "*** FAILED - ABORTING\n"); \ 
      exit(1); \ 
     } \ 
    } while (0) 

__global__ void my_kernel2(int *data){ 
    int idx = threadIdx.x + (blockDim.x *blockIdx.x); 
    if (idx < DSIZE) data[idx] =+ DVAL; 
} 

int my_test_func2(){ 

    int *d_data, *h_data; 
    h_data = (int *) malloc(DSIZE * sizeof(int)); 
    if (h_data == 0) {printf("malloc fail\n"); exit(1);} 
    cudaMalloc((void **)&d_data, DSIZE * sizeof(int)); 
    cudaCheckErrors("cudaMalloc fail"); 
    for (int i = 0; i < DSIZE; i++) h_data[i] = 0; 
    cudaMemcpy(d_data, h_data, DSIZE * sizeof(int), cudaMemcpyHostToDevice); 
    cudaCheckErrors("cudaMemcpy fail"); 
    my_kernel2<<<((DSIZE+nTPB-1)/nTPB), nTPB>>>(d_data); 
    cudaDeviceSynchronize(); 
    cudaCheckErrors("kernel"); 
    cudaMemcpy(h_data, d_data, DSIZE * sizeof(int), cudaMemcpyDeviceToHost); 
    cudaCheckErrors("cudaMemcpy 2"); 
    for (int i = 0; i < DSIZE; i++) 
    if (h_data[i] != DVAL) {printf("Results check failed at offset %d, data was: %d, should be %d\n", i, h_data[i], DVAL); exit(1);} 
    printf("Results check 2 passed!\n"); 
    return 0; 
} 

main.cpp:

#include <stdio.h> 

#include "test1.h" 
#include "test2.h" 

int main(){ 

    my_test_func1(); 
    my_test_func2(); 
    return 0; 
} 

Quando compilo in base ai comandi impartiti, e correre ./testmain ottengo:

$ ./testmain 
Results check 1 passed! 
Results check 2 passed! 

Si noti che, se si preferisce, si può generare un libtest.so anziché test.so e quindi è possibile utilizzare una sequenza di build modificata per l'eseguibile di test:

g++ -c main.cpp 
g++ -o testmain main.o -L. -ltest 

Non credo che faccia alcuna differenza, ma potrebbe essere una sintassi più familiare.

Sono sicuro che c'è più di un modo per realizzare questo. Questo è solo un esempio. Si consiglia inoltre di rivedere la sezione pertinente di nvcc manual e di rivedere anche lo examples.

EDIT: ho provato questo sotto CUDA 5.5 RC, e il passo finale collegamento all'applicazione sono lamentati di non trovare il lib cudart (warning: libcudart.so.5.5., needed by ./libtest.so, not found). Tuttavia la seguente modifica relativamente semplice (esempio Makefile) dovrebbe funzionare sotto cuda 5.0 o cuda 5.5.

Makefile:

testmain : main.cpp libtest.so 
     g++ -c main.cpp 
     g++ -o testmain -L. -ldl -Wl,-rpath,. -ltest -L/usr/local/cuda/lib64 -lcudart main.o 

libtest.so : link.o 
     g++ -shared -Wl,-soname,libtest.so -o libtest.so test1.o test2.o link.o -L/usr/local/cuda/lib64 -lcudart 

link.o : test1.cu test2.cu test1.h test2.h 
     nvcc -m64 -arch=sm_20 -dc -Xcompiler '-fPIC' test1.cu test2.cu 
     nvcc -m64 -arch=sm_20 -Xcompiler '-fPIC' -dlink test1.o test2.o -o link.o 

clean : 
     rm -f testmain test1.o test2.o link.o libtest.so main.o 

Answer 2

Hai provato esplicitamente disabilitazione codice del dispositivo rilocabile? vale a dire -rdc=false? Ho ottenuto questo undefined reference to __cudaRegisterLinkedBinaryWhatever con -rdc=true e se ne sono andati quando l'ho rimosso. Anche se non sono abbastanza esperto per spiegare cosa sta succedendo esattamente.