如何将CUDA代码分离到多个文件中

Question

我正在尝试将一个CUDA程序分成两个独立的.cu文件，以便更接近使用C++编写真实应用程序。一个简单的程序需要：

在主机和设备上分配内存。
将主机数组初始化为一系列数字。 将主机数组复制到设备数组 使用设备核函数找到数组中所有元素的平方 将设备数组复制回主机数组 打印结果

如果我把所有内容放在一个.cu文件中运行，它可以正常工作。但当我将其拆分为两个文件时，就会出现链接错误。就像我最近提出的所有问题一样，我知道这是一个小问题，但是它是什么呢？

KernelSupport.cu

#ifndef _KERNEL_SUPPORT_
#define _KERNEL_SUPPORT_

#include <iostream>
#include <MyKernel.cu>

int main( int argc, char** argv) 
{
    int* hostArray;
    int* deviceArray;
    const int arrayLength = 16;
    const unsigned int memSize = sizeof(int) * arrayLength;

    hostArray = (int*)malloc(memSize);
    cudaMalloc((void**) &deviceArray, memSize);

    std::cout << "Before device\n";
    for(int i=0;i<arrayLength;i++)
    {
        hostArray[i] = i+1;
        std::cout << hostArray[i] << "\n";
    }
    std::cout << "\n";

    cudaMemcpy(deviceArray, hostArray, memSize, cudaMemcpyHostToDevice);
    TestDevice <<< 4, 4 >>> (deviceArray);
    cudaMemcpy(hostArray, deviceArray, memSize, cudaMemcpyDeviceToHost);

    std::cout << "After device\n";
    for(int i=0;i<arrayLength;i++)
    {
        std::cout << hostArray[i] << "\n";
    }

    cudaFree(deviceArray);
    free(hostArray);

    std::cout << "Done\n";
}

#endif

MyKernel.cu

#ifndef _MY_KERNEL_
#define _MY_KERNEL_

__global__ void TestDevice(int *deviceArray)
{
    int idx = blockIdx.x*blockDim.x + threadIdx.x;
    deviceArray[idx] = deviceArray[idx]*deviceArray[idx];
}


#endif

构建日志：

1>------ Build started: Project: CUDASandbox, Configuration: Debug x64 ------
1>Compiling with CUDA Build Rule...
1>"C:\CUDA\bin64\nvcc.exe"    -arch sm_10 -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\bin"    -Xcompiler "/EHsc /W3 /nologo /O2 /Zi   /MT  "  -maxrregcount=32  --compile -o "x64\Debug\KernelSupport.cu.obj" "d:\Stuff\Programming\Visual Studio 2008\Projects\CUDASandbox\CUDASandbox\KernelSupport.cu" 
1>KernelSupport.cu
1>tmpxft_000016f4_00000000-3_KernelSupport.cudafe1.gpu
1>tmpxft_000016f4_00000000-8_KernelSupport.cudafe2.gpu
1>tmpxft_000016f4_00000000-3_KernelSupport.cudafe1.cpp
1>tmpxft_000016f4_00000000-12_KernelSupport.ii
1>Linking...
1>KernelSupport.cu.obj : error LNK2005: __device_stub__Z10TestDevicePi already defined in MyKernel.cu.obj
1>KernelSupport.cu.obj : error LNK2005: "void __cdecl TestDevice__entry(int *)" (?TestDevice__entry@@YAXPEAH@Z) already defined in MyKernel.cu.obj
1>D:\Stuff\Programming\Visual Studio 2008\Projects\CUDASandbox\x64\Debug\CUDASandbox.exe : fatal error LNK1169: one or more multiply defined symbols found
1>Build log was saved at "file://d:\Stuff\Programming\Visual Studio 2008\Projects\CUDASandbox\CUDASandbox\x64\Debug\BuildLog.htm"
1>CUDASandbox - 3 error(s), 0 warning(s)
========== Build: 0 succeeded, 1 failed, 0 up-to-date, 0 skipped ==========

我正在Windows 7 64位上运行Visual Studio 2008。

---

编辑：

我觉得我需要解释一下。 我想要的结果是创建一个普通的C++应用程序，例如Main.cpp，并从那里启动应用程序。在我的.cpp代码的某些点上，我想能够引用CUDA库。 因此，我的想法是将CUDA内核代码放入它们自己的.cu文件中，然后有一个支持性的.cu文件来处理与设备的通信以及调用内核函数等事项。

- Mr Bell

Answer 1

在kernelsupport.cu中包含了mykernel.cu，因此当您尝试进行链接时，编译器会发现mykernel.cu出现了两次。您需要创建一个定义TestDevice的头文件并包含它。

针对评论的回复:

像这样做应该就可以了。

// MyKernel.h
#ifndef mykernel_h
#define mykernel_h
__global__ void TestDevice(int* devicearray);
#endif

然后将包含文件更改为

//KernelSupport.cu
#ifndef _KERNEL_SUPPORT_
#define _KERNEL_SUPPORT_

#include <iostream>
#include <MyKernel.h>
// ...

关于你的编辑

只要你在C++代码中使用的头文件没有任何CUDA特定的内容（例如__kernel__，__global__等），你就可以将C++和CUDA代码链接起来。

- Scott Wales

---

您的 MyKernel.h 应该有 void TestDeviceWrapper(dim3 grid, dim3 block, int *devicearray)，因为当 KernelSupport.cu 变为 KernelSupport.cpp 时，cl.exe 将无法理解 global 语法。然后在 MyKernel.cu 中，TestDeviceWrapper() 只需调用TestDevice<<<>>>。 - Tom

听起来很合理，所给的代码假定它将被包含在一个CUDA文件中，正如问题中所述。 - Scott Wales

是的，但他还说：“我在这里寻找的最终结果是拥有一个正常的C++应用程序，例如带有int main()事件的Main.cpp，并从那里运行事物。”不过，这是在编辑问题时添加的。 - Tom

Answer 2

如果您查看CUDA SDK的代码示例，它们具有引用从.cu文件编译的函数的extern C定义。这样，.cu文件将由nvcc编译，仅在主程序中链接，而.cpp文件将正常编译。

例如，在marchingCubes_kernel.cu中有以下函数体：

extern "C" void
launch_classifyVoxel( dim3 grid, dim3 threads, uint* voxelVerts, uint *voxelOccupied, uchar *volume,
                      uint3 gridSize, uint3 gridSizeShift, uint3 gridSizeMask, uint numVoxels,
                      float3 voxelSize, float isoValue)
{
    // calculate number of vertices need per voxel
    classifyVoxel<<<grid, threads>>>(voxelVerts, voxelOccupied, volume, 
                                     gridSize, gridSizeShift, gridSizeMask, 
                                     numVoxels, voxelSize, isoValue);
    cutilCheckMsg("classifyVoxel failed");
}

在 marchingCubes.cpp 中（其中包含 main() 函数），只有一个定义：

extern "C" void
launch_classifyVoxel( dim3 grid, dim3 threads, uint* voxelVerts, uint *voxelOccupied, uchar *volume,
                      uint3 gridSize, uint3 gridSizeShift, uint3 gridSizeMask, uint numVoxels,
                      float3 voxelSize, float isoValue);

您也可以将这些内容放在 .h 文件中。

- tkerwin

---

在最近版本的CUDA工具包中，您不应该需要使用extern "C"。过去，由于nvcc将主机代码视为C语言，因此需要使用它，但现在默认情况下是C ++。删除extern "C"，它会使代码变得模糊不清！ - Tom

很好知道。他们应该更新SDK示例以反映这一点。但是，您仍然需要执行CUDA调用包装，我认为没有任何简单的方法可以解决这个问题。 - tkerwin

是的，SDK示例自创建以来就没有更新过，因此虽然较新的示例反映了最新的标准，但旧的示例有些过时。不过它们仍然说明了编码技术，只是风格有所不同。您是正确的，无法避免CUDA调用包装。不过这完全合理，三重尖括号语法（<<< >>>）是CUDA C的一部分，而不是C，因此您需要一个CUDA C编译器（即nvcc）来编译它。我认为这是为了获得运行时API的优雅而付出的小代价。 - Tom

Answer 3

获取分离其实很简单，请查看这个答案来设置它。然后，您只需将主机代码放在.cpp文件中，将设备代码放在.cu文件中，构建规则告诉Visual Studio如何将它们链接成最终可执行文件。

您的代码中存在的直接问题是，您定义了__global__ TestDevice函数两次，一次在您#include MyKernel.cu时，另一次在编译MyKernel.cu时。

您还需要在.cu文件中添加一个包装器 - 目前您正在从主函数中调用TestDevice<<<>>>，但当您将其移动到.cpp文件中时，它将与cl.exe一起编译，而cl.exe不理解<<<>>>语法。因此，在.cpp文件中，您将简单地调用TestDeviceWrapper(griddim, blockdim, params)，并在.cu文件中提供此函数。

如果您需要示例，则SDK中的SobolQRNG示例实现了良好的分离，尽管它仍使用cutil，但我始终建议避免使用cutil。

- Tom

Answer 4

简单的解决方案是关闭 MyKernel.cu 文件的构建。

属性 -> 常规 -> 从构建中排除

我认为更好的解决方案是将您的内核拆分为 cu 和 cuh 文件，并包含它，例如：

//kernel.cu
#include "kernel.cuh"
#include <cuda_runtime.h>

__global__ void increment_by_one_kernel(int* vals) {
  vals[threadIdx.x] += 1;
}

void increment_by_one(int* a) {
  int* a_d;

  cudaMalloc(&a_d, 1);
  cudaMemcpy(a_d, a, 1, cudaMemcpyHostToDevice);
  increment_by_one_kernel<<<1, 1>>>(a_d);
  cudaMemcpy(a, a_d, 1, cudaMemcpyDeviceToHost);

  cudaFree(a_d);
}

 

//kernel.cuh
#pragma once

void increment_by_one(int* a);

 

//main.cpp
#include "kernel.cuh"

int main() {
  int a[] = {1};

  increment_by_one(a);

  return 0;
}

- thebaldwin

---

这只能在您的主文件是 .cu 文件时工作。一旦将其放入 .cpp 文件中，它就不适用了。 - Tom

一旦您将所有CUDA/内核代码拆分为适当的cu/cuh文件，重命名或将主函数移动到cpp文件中就不应该有问题。请参考我的示例，我不清楚为什么它不合适。 - thebaldwin