如何将内核代码保存在主 .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 设备代码的文件（例如内核定义），那么您需要熟悉使用 device code linker。

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

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

Answer 2

您可以创建 cuda 头文件 *.cuh 并将其作为标准头文件包含在内。我不会只将内核放在一个单独的文件中，而是将它们与一些初始化函数放在一起，然后只将这些函数放在头文件中，因为您通常不只是从外部代码调用内核，而是调用一些负责处理的函数内存等。我通常制作这样的标题：

#ifndef __CUDAHEADER_CUH__
#define __CUDAHEADER_CUH__

/** Initialize cuda stuff */
void cudaInit(Data * host_data);

/** Cleanup, frees resources used by the device. */
void cudaFinalize();

#endif

然后是包含内核、设备方法和处理 cuda 内容的这些方法的文件：

#include "cudaHeader.cuh"

//some global variables like:
Data * device_data;

//some kernels and device functions:
__global__ void someKernel(data * device_data) {
    ...
}

void cudaInit(Data * host_data) {
    some cudaMalloc()
    some cudaMemcpy()
    someKernel<<< gridRes, blockRes >>>(device_data);
}


void cudaFinalize() {
    cudaFree(device_data);
}

但是还有更多方法可以处理您的代码...

Answer 3

第一个例子：我认为你必须像这样构建（最后一行）：

g++ -L/usr/local/cuda/lib64 -o test main.o test.o -lcudart

即将库放在链接的最后。 （无论如何我都需要）