cuda - 如何将内核代码保存在主 .cpp 之外的单独 .cu 文件中？

Question

如何将 cuda 内核代码和项目中的其他 cpp 代码分开？我想将所有内核定义收集在一个文件中，因为其他 cpp 文件在需要时调用它们。我试图在 kernel.cu 中编写所有内核并通过包含 kernel.cu 文件来调用内核，但它在编译时出现以下错误。

/usr/bin/ld: error: ./vector_summation.o: multiple definition of 

'perform_summation_method1(int*, int)'
/usr/bin/ld: ./kernels.o: previous definition here
/usr/bin/ld: error: ./vector_summation.o: multiple definition of '__device_stub__Z25perform_summation_method1Pii(int*, int)'
/usr/bin/ld: ./kernels.o: previous definition here
/usr/bin/ld: error: ./vector_summation.o: multiple definition of '__device_stub__Z25perform_summation_method2PiS_i(int*, int*, int)'
/usr/bin/ld: ./kernels.o: previous definition here
/usr/bin/ld: error: ./vector_summation.o: multiple definition of 'perform_summation_method2(int*, int*, int)'
/usr/bin/ld: ./kernels.o: previous definition here
/usr/bin/ld: error: ./vector_summation.o: multiple definition of '__device_stub__Z25perform_summation_method3PiS_i(int*, int*, int)'
/usr/bin/ld: ./kernels.o: previous definition here
/usr/bin/ld: error: ./vector_summation.o: multiple definition of 'perform_summation_method3(int*, int*, int)'
/usr/bin/ld: ./kernels.o: previous definition here

Answer 1

您的操作方式与使用普通 cpp 文件/模块的方式基本相同。在 c++ 中，当您想从另一个文件访问函数时，通常不会在另一个文件中包含一个 .cpp 文件。您包括通常仅包含函数原型(prototype)的 header 。

这是一个例子:

测试.h:

void my_cuda_func();

主.cpp:

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

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

测试.cu:

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


__global__ void my_kernel(){
  printf("Hello!\n");
}

void my_cuda_func(){
  my_kernel<<<1,1>>>();
  cudaDeviceSynchronize();
}

使用以下命令构建:

g++ -c main.cpp
nvcc -arch=sm_20 -c test.cu
g++  -o test main.o test.o -L/usr/local/cuda/lib64 -lcudart

当然还有其他方法。如果要链接到 C 而不是 C++，则需要考虑到这一点。如果您想直接从其他模块调用内核而不是使用包装函数，那么您需要通过 nvcc 而不是 g++ 传递所有模块(并且它们都应该是 .cu 文件)。此外，如果您想拥有多个带有 GPU 设备代码的文件(例如内核定义)，那么您需要熟悉使用 device code linker。 .

为了完整起见，这里对上面的示例进行了重新处理，以显示如果您希望将所有内核定义都放在一个文件中，但能够直接从另一个模块调用内核时该怎么做:

测试.h:

__global__ void my_kernel();

main.cu:

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

int main(){
  my_kernel<<<1,1>>>();
  cudaDeviceSynchronize();
  return 0;
}

测试.cu:

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


__global__ void my_kernel(){
  printf("Hello!\n");
}

构建:

nvcc -arch=sm_20 -c main.cu
nvcc -arch=sm_20 -c test.cu
nvcc -arch=sm_20 -o test main.o test.o