链接通过clang编译.cu文件生成的.ll文件

我正在使用clang编译以下代码: clang++ -std=c++11 -emit-llvm -c -S $1 --cuda-gpu-arch=sm_30。这将生成vectoradd-cuda-nvptx64-nvidia-cuda-sm_30.llvectoradd.ll文件。运行某些LLVM分析的目标是传递给内核,这可能会对它进行检测。所以我想将分析后的IR链接到可执行文件,但是我不确定如何执行。当我尝试将.ll文件与llvm-link链接时,出现错误Linking globals named '_Z9vectoraddPiS_S_i': symbol multiply defined!。我不太确定如何实现此目标,因此不胜感激。

#define THREADS_PER_BLOCK 512

__global__ void vectoradd(int *A,int *B,int *C,int N) {
  int gi = threadIdx.x + blockIdx.x * blockDim.x;
  if ( gi < N) {
    C[gi] = A[gi] + B[gi];
  }
}

int main(int argc,char **argv) {
  int N = 10000,*d_A,*d_B,*d_C;

  /// allocate host memory
  std::vector<int> A(N);
  std::vector<int> B(N);
  std::vector<int> C(N);

  /// allocate device memory
  cudaMalloc((void **) &d_A,N * sizeof(int));
  cudaMalloc((void **) &d_B,N * sizeof(int));
  cudaMalloc((void **) &d_C,N * sizeof(int));

  /// populate host data
  for ( size_t i = 0; i < N; ++i) {
    A[i] = i; B[i] = i;
  }

  /// copy to device
  cudaMemcpy(d_A,&A[0],N * sizeof(int),cudaMemcpyHostToDevice);
  cudaMemcpy(d_B,&B[0],cudaMemcpyHostToDevice);

  dim3 block(THREADS_PER_BLOCK,1,1);
  dim3 grid((N + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK,1);

  vectoradd<<<grid,block>>>(d_A,d_B,d_C,N);
  cudaDeviceSynchronize();

  cudaMemcpy(&C[0],cudaMemcpyDeviceToHost);

  return 0;
}
hirxn45624 回答:链接通过clang编译.cu文件生成的.ll文件

Clang中的CUDA编译轨迹相当complicated(就像在NVIDIA工具链中一样),您尝试执行的操作无效。编译过程中每个分支的LLVM IR必须保持独立,直到可以直接链接的对象可用为止。因此,您需要手动执行许多中间步骤。

必须先将GPU的LLVM IR代码编译为PTX代码,然后将其组装为可以与主机目标文件链接的二进制有效负载。

因此,在您的示例中,您首先要做类似的事情:

clang++ -std=c++11 -emit-llvm -c -S test.cu --cuda-gpu-arch=sm_52

会发出两个llvm IR文件test-cuda-nvptx64-nvidia-cuda-sm_52.lltest.ll。然后需要将GPU代码编译为PTX(有关nvptx后端here的更多信息):

llc -mcpu=sm_52 test-cuda-nvptx64-nvidia-cuda-sm_52.ll -o test.ptx

现在,可以将PTX代码汇编成一个ELF文件,以后可以通过nvcc(或带有两个附加步骤的主机链接程序)以常规方式链接该文件:

ptxas --gpu-name=sm_52 test.ptx -o test.ptx.o
fatbinary --cuda -64 --create test.fatbin --image=profile=sm_52,file=test.ptx.o

对于主机代码,您可以执行以下操作

llc test.ll
clang -m64 -c test.s

从LLVM IR生成汇编程序输出,然后将其汇编到目标文件中。

现在使用包含CUDA的已编译代码的胖子文件和包含已编译的宿主代码的目标文件,您可以执行链接。我无法使用clang测试将主机对象文件与fatbinary链接,这是您需要自己解决的问题。研究CUDA编译调用期间clang的详细输出以及nvcc文档,以更好地了解设备代码构建系统的工作方式,将具有指导意义。

本文链接:https://www.f2er.com/3127478.html

大家都在问