如何将内核代码保存在除主.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文件.您包含通常只包含函数原型的标头.

这是一个例子:

test.h:

void my_cuda_func();

main.cpp中:

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

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

test.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设备代码的多个文件(例如内核定义),那么您需要熟悉使用设备代码链接器.

为了完整起见,上面的示例重新编写,以显示如果您希望所有内核定义都在一个文件中,但能够直接从另一个模块调用内核时该怎么做:

test.h:

__global__ void my_kernel();

main.cu:

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

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

test.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