如何將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的int main()事件，並從那裏運行。在我的.cpp代碼中的特定點處，我希望能夠引用CUDA位。所以我的想法（並糾正我，如果這裏有一個更標準的約定）是我將把CUDA內核代碼放入它們的.cu文件中，然後有一個支持.cu文件，它將負責與設備通話並呼叫內核函數和什麼不是。

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代碼。

Answer 2

簡單的解決方案是關閉建立MyKernel.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; 
} 

Answer 3

如果您看看CUDA SDK代碼示例，它們具有extern C定義了從.cu文件編譯的引用函數。通過這種方式，.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（其中主（）所在的）剛下了一個定義：

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文件中。

Answer 4

獲得分離其實很簡單，請查看this answer瞭解如何設置它。然後，您只需將.cpp文件中的主機代碼和設備代碼放入.cu文件中，構建規則就會告訴Visual Studio如何將它們鏈接到最終的可執行文件中。

代碼中的直接問題是您定義了__global__ TestDevice函數兩次，一次是當您#include MyKernel.cu和一次獨立編譯MyKernel.cu。

您還需要將一個包裝文件放入一個.cu文件中 - 目前您正在從主函數調用TestDevice<<<>>>，但是當您將其移入.cpp文件時，它將與cl.exe一起編譯，不理解<<<>>>語法。因此，您只需在.cpp文件中調用TestDeviceWrapper(griddim, blockdim, params)並在.cu文件中提供此功能。

如果您想要一個示例，SDK中的SobolQRNG示例實現了很好的分離，儘管它仍然使用cutil，我總是建議避免cutil。