gpt4 book ai didi

cuda - nvlink、可重定位设备代码和静态设备库

转载 作者:行者123 更新时间:2023-12-01 04:46:22 24 4
gpt4 key购买 nike

在调查可重定位设备代码的一些问题时,我偶然发现了一些我不太理解的东西。

这是 slide 6 上图片的用例.我用了 answer of Robert Crovella作为复制代码的基础。这个想法是我们将一些可重定位的设备代码编译成一个静态库(例如一些数学/工具箱库),我们想将该预编译库的一些函数用于我们程序的另一个设备库:

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 调用)似乎找不到它。这是为什么?根据official documentation :

The device linker has the ability to read the static host library formats (.a on Linux and Mac, .lib on 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 时是否有错误?

请注意,这是在 Arch Linux 上使用 CUDA 6.5 测试的。

编辑:修复了带有注释行的重现代码

生成文件

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

测试.h

#ifndef TEST_H
# define TEST_H

int my_test_func();

#endif //! TEST_H

测试.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;
}

工具.cuh

#ifndef UTIL_CUH
# define UTIL_CUH

__device__ int my_square (int a);

#endif //! UTIL_CUH

工具.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;
}

最佳答案

我建议在问题中放一个完整的简单示例,就像我在下面所做的那样。代码的外部链接不受欢迎。当它们变得陈旧时,问题就变得不那么有值(value)了。

是的,您在生成 libutil.a 时出错 创建具有公开设备链接的静态库与创建(根据定义)没有公开设备的共享库不同-链接。请注意我在您链接的上一个问题中提到的“无 CUDA 包装器”。这个问题中的示例暴露了设备链接,因为 my_square 在库中,但被库外部的代码使用。

查看 nvcc relocatable device code compiling examples你会发现一个生成设备可链接静态库的库。静态库创建 中没有设备链接步骤。设备链接步骤在最终可执行文件创建时完成(或者在本例中,在创建 so 时完成,即“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
$

在此示例中,设备链接自动发生在用于创建 .so 库的 nvcc 调用中。在我这里的示例中,我已经设置了我的 LD_LIBRARY_PATH 环境变量以包含我的工作目录。在 CentOS 6.2 上使用 CUDA 6.5 测试(请注意,在创建可执行文件期间可以执行多个设备链接操作,但这些设备链接操作必须在单独的链接域内,即用户代码或用户代码入口点不能在域之间共享。这里不是这种情况。)

关于cuda - nvlink、可重定位设备代码和静态设备库,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/26147981/

24 4 0
Copyright 2021 - 2024 cfsdn All Rights Reserved 蜀ICP备2022000587号
广告合作:1813099741@qq.com 6ren.com