nvlink, перемещаемый код устройства и статические библиотеки устройств
Исследуя некоторые проблемы с перемещаемым кодом устройства, я наткнулся на то, что я не совсем понимаю.
Это пример использования того, что изображено на слайде 6. Я использовал ответ Роберта Кровеллы в качестве основы для кода воспроизведения. Идея состоит в том, что у нас есть некоторый перемещаемый код устройства, скомпилированный в статическую библиотеку (например, некоторую библиотеку math/toolbox), и мы хотим использовать некоторые функции этой предварительно скомпилированной библиотеки в другую библиотеку устройств нашей программы:
libutil.a ---> libtest.so ---> test_pgm
Допустим, что эта внешняя библиотека содержит следующую функцию:
__device__ int my_square (int a);
libutil.a например, был сгенерирован следующим образом (в другом проекте):
nvcc ${NVCC_FLAGS} -dc util.cu
nvcc ${NVCC_FLAGS} -dlink util.o -o util_dlink.o
nvcc ${NVCC_FLAGS} -lib util_dlink.o util.o -o libutil.a
Затем в нашем проекте, чтобы сгенерировать libtest.so:
nvcc ${NVCC_FLAGS} -dc test.cu
nvcc ${NVCC_FLAGS} -dlink test.o libutil.a -o test_dlink.o
g++ -shared -Wl,-soname,libtest.so -o libtest.so test.o test_dlink.o libutil.a -L${CUDA_LIBDIR} -lcudart
Но я получаю следующую ошибку при генерации test_dlink.o:
nvlink error : Undefined reference to '_Z9my_squarei' in 'test.o'
Линкер не находит наш манекен my_square(int) функция. Если мы вместо этого используем (при условии, что у нас был доступ к util.o):
nvcc ${NVCC_FLAGS} -dlink test.o util.o -o test_dlink.o
Компоновщик преуспевает, и все работает хорошо после.
Расследование дальше:
$ nm -C libutil.a
util_dlink.o:
U atexit
U __cudaRegisterFatBinary
0000000000000015 T __cudaRegisterLinkedBinary_39_tmpxft_0000106a_00000000_6_util_cpp1_ii_843d693d
...
util.o:
U __cudaInitModule
U __cudaRegisterLinkedBinary_39_tmpxft_0000106a_00000000_6_util_cpp1_ii_843d693d
...
0000000000000015 T my_square(int)
...
Символ есть в архиве util.o, но nvlink (вызывается nvcc Кажется, не найти его. Это почему? Согласно официальной документации:
Компоновщик устройства может читать статические форматы хост-библиотеки (.a в Linux и Mac,.lib в Windows).
Конечно, мы можем извлечь объектный файл и связать его с ним:
ar x libutil.a `ar t libutil.a | grep -v "dlink"`
nvcc ${NVCC_FLAGS} -dlink test.o util.o -o test_dlink.o
Но это не похоже на ожидаемое решение... Так чего мне здесь не хватает? Другая nvcc вариант, который решает это? Есть ли ошибка при генерации libutil.a и / или libtest.so?
Обратите внимание, что это было протестировано с CUDA 6.5 на Arch Linux.
РЕДАКТИРОВАТЬ: исправлен код репро с комментариями
Makefile
NVCC_FLAGS=-m64 -arch=sm_20 -Xcompiler '-fPIC'
CUDA_LIBDIR=${CUDA_HOME}/lib64
testmain : main.cpp libtest.so
g++ -c main.cpp
g++ -o testmain -L. -ldl -Wl,-rpath,. -ltest -L${CUDA_LIBDIR} -lcudart main.o
libutil.a : util.cu util.cuh
nvcc ${NVCC_FLAGS} -dc util.cu
# ---> FOLLOWING LINES THAT WERE WRONG <---
# nvcc ${NVCC_FLAGS} -dlink util.o -o util_dlink.o
# nvcc ${NVCC_FLAGS} -lib util.o util_dlink.o -o libutil.a
# INSTEAD:
nvcc ${NVCC_FLAGS} -lib util.o -o libutil.a
# Assuming util is an external library, so util.o is not available
rm util.o
libtest.so : test.cu test.h libutil.a util.cuh
nvcc ${NVCC_FLAGS} -dc test.cu
# Use NVCC for device linking + G++
nvcc -v ${NVCC_FLAGS} -dlink test.o libutil.a -o test_dlink.o
g++ -shared -o libtest.so test.o test_dlink.o libutil.a -L${CUDA_LIBDIR} -lcudart
# Or let NVCC generate the shared library
#nvcc -v ${NVCC_FLAGS} -shared -L. -lutil test.o -o libtest.so
clean :
rm -f testmain *.o *.a *.so
test.h
#ifndef TEST_H
# define TEST_H
int my_test_func();
#endif //! TEST_H
test.cu
#include <stdio.h>
#include "test.h"
#include "util.cuh"
#define DSIZE 1024
#define DVAL 10
#define SQVAL 3
#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_kernel(int *data){
int idx = threadIdx.x + (blockDim.x *blockIdx.x);
if (idx < DSIZE) data[idx] =+ DVAL + my_square (SQVAL);
}
int my_test_func()
{
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_kernel<<<((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 + SQVAL*SQVAL)
{
printf("Results check failed at offset %d, data was: %d, should be %d\n",
i, h_data[i], DVAL);
exit(1);
}
printf("Results check passed!\n");
return 0;
}
util.cuh
#ifndef UTIL_CUH
# define UTIL_CUH
__device__ int my_square (int a);
#endif //! UTIL_CUH
util.cu
#include "util.cuh"
__device__ int my_square (int a)
{
return a * a;
}
main.cpp
#include "test.h"
int main()
{
my_test_func();
return 0;
}
1 ответ
Я предлагаю привести в пример полный простой пример, как я это сделал ниже. Внешние ссылки на код не одобрены. Когда они устаревают, вопрос становится менее ценным.
Да, у вас есть ошибка при генерации libutil.a Создание статической библиотеки с открытой привязкой к устройству отличается от создания общей библиотеки без (по определению) открытой привязки к устройству. Обратите внимание на мое упоминание "Обертки без CUDA" в предыдущем вопросе, который вы указали. Пример в этом вопросе показал связывание устройства, потому что my_square находится в библиотеке, но используется кодом, внешним по отношению к библиотеке.
Посмотрите примеры компиляции кода перемещаемого устройства nvcc, и вы найдете тот, который генерирует статическую библиотеку, связанную с устройством. В создании статической библиотеки нет этапа привязки устройства. Этап привязки к устройству выполняется при окончательном создании исполняемого файла (или в этом случае при создании объекта, то есть "границы CUDA"). "Дополнительная" операция связывания устройства при создании статической библиотеки является проксимальной причиной ошибки, которую вы наблюдаете.
Вот полностью проработанный пример:
$ cat util.h
__device__ float my_square(float);
$ cat util.cu
__device__ float my_square(float val){ return val*val;}
$ cat test.h
float dbl_sq(float val);
$ cat test.cu
#include "util.h"
__global__ void my_dbl_sq(float *val){
*val = 2*my_square(*val);
}
float dbl_sq(float val){
float *d_val, h_val;
cudaMalloc(&d_val, sizeof(float));
h_val = val;
cudaMemcpy(d_val, &h_val, sizeof(float), cudaMemcpyHostToDevice);
my_dbl_sq<<<1,1>>>(d_val);
cudaMemcpy(&h_val, d_val, sizeof(float), cudaMemcpyDeviceToHost);
return h_val;
}
$ cat main.cpp
#include <stdio.h>
#include "test.h"
int main(){
printf("%f\n", dbl_sq(2.0f));
return 0;
}
$ nvcc -arch=sm_35 -Xcompiler -fPIC -dc util.cu
$ nvcc -arch=sm_35 -Xcompiler -fPIC -lib util.o -o libutil.a
$ nvcc -arch=sm_35 -Xcompiler -fPIC -dc test.cu
$ nvcc -arch=sm_35 -shared -Xcompiler -fPIC -L. -lutil test.o -o libtest.so
$ g++ -o main main.cpp libtest.so
$ cuda-memcheck ./main
========= CUDA-MEMCHECK
8.000000
========= ERROR SUMMARY: 0 errors
$
В этом примере привязка устройства происходит автоматически в nvcc вызов, который используется для создания библиотеки.so. В моем примере здесь я уже установил мой LD_LIBRARY_PATH переменная окружения, чтобы включить мой рабочий каталог. Протестировано с использованием CUDA 6.5 на CentOS 6.2 (обратите внимание, что во время создания исполняемого файла можно выполнять несколько операций связывания устройства, но эти операции связывания устройства должны находиться в отдельных доменах связи, то есть в точках входа кода пользователя или кода пользователя). не может быть разделен между доменами. Это не тот случай.)