Cómo separar el código CUDA en varios archivos

Question

Estoy intentando separar un programa CUDA en dos archivos .cu separados en un esfuerzo por acercarme a escribir una aplicación real en C++. Tengo un pequeño programa simple que:

Asigna una memoria en el host y el dispositivo.
Inicializa la matriz de host en una serie de números. copias del conjunto anfitrión de una serie de dispositivos de encuentra la plaza de todos los elementos de la matriz usando un kernel dispositivo copias del array dispositivo de nuevo a la matriz huésped imprime los resultados

Esto funciona muy bien si pongo todo en un archivo .cu y ejecútelo. Cuando lo divido en dos archivos separados, empiezo a recibir errores de enlace. Al igual que todas mis preguntas recientes, sé que esto es algo pequeño, pero ¿qué es?

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 

Build Log:

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 ========== 

Me postulo Visual Studio 2008 en Windows 7 de 64 bits.

---

Editar:

creo que es necesario dar más detalles sobre esto un poco. El resultado final que estoy buscando aquí es tener una aplicación normal de C++ con algo así como Main.cpp con el evento int main() y ejecutar las cosas desde allí. En ciertos puntos de mi código .cpp, quiero poder hacer referencia a los bits de CUDA. Así que mi forma de pensar (y corregirme si hay una convención más estándar aquí) es que incluiré el código Kernel de CUDA en sus archivos .cu y luego tendré un archivo .cu de soporte que se encargará de hablar con el dispositivo y llamar funciones del núcleo y qué no.

Answer 1

Está incluyendo mykernel.cu en kernelsupport.cu, cuando intenta vincular el compilador ve mykernel.cu dos veces. Deberá crear un encabezado que defina TestDevice e incluirlo en su lugar.

nuevo comentario:

Algo como esto debería funcionar

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

y luego cambiar el archivo que incluye a

//KernelSupport.cu 
#ifndef _KERNEL_SUPPORT_ 
#define _KERNEL_SUPPORT_ 

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

re tu edición

Mientras el encabezado que el uso en el código de C++ no tiene ningún material específico de cuda (__kernel__, __global__, etc.) debería estar bien enlazando el código C++ y cuda.

Answer 2

La solución simple es desactivar la construcción de su archivo MyKernel.cu.

Propiedades -> General -> Quedan excluidos de la acumulación

La mejor solución imo es dividir el núcleo en un Cu y un archivo CUH, e incluyen que, por ejemplo:

//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

Si mira los ejemplos del código CUDA SDK, tienen C externo define las funciones de referencia compiladas a partir de archivos .cu. De esta forma, los archivos .cu son compilados por nvcc y solo se vinculan al programa principal mientras que los archivos .cpp se compilan normalmente.

Por ejemplo, en marchingCubes_kernel.cu tiene el cuerpo de la función:

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"); 
} 

Mientras que en marchingCubes.cpp (donde main() reside) solo tiene una definición:

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); 

Usted puede poner estos en un archivo .h también.

Answer 4

Obtener la separación es bastante simple. Consulte this answer para saber cómo configurarla. Luego, simplemente coloque su código de host en archivos .cpp y el código de su dispositivo en archivos .cu, las reglas de compilación le dicen a Visual Studio cómo vincularlos en el ejecutable final.

El problema inmediato en su código es que está definiendo la función __global__ TestDevice dos veces, una vez cuando #include MyKernel.cu y una cuando compila MyKernel.cu de forma independiente.

También deberá colocar un contenedor en un archivo .cu - en el momento en que llame al TestDevice<<<>>> desde su función principal, pero cuando lo mueva a un archivo .cpp se compilará con cl.exe, que no lo hace no entiendo la sintaxis <<<>>>. Por lo tanto, simplemente llame al TestDeviceWrapper(griddim, blockdim, params) en el archivo .cpp y proporcione esta función en su archivo .cu.

Si desea un ejemplo, la muestra SobolQRNG en el SDK logra una buena separación, aunque todavía utiliza cutil y yo siempre recomendaría evitar cutil.