Come separare il codice CUDA in più file

Question

Sto provando a separare un programma CUDA in due file separati .cu nel tentativo di avvicinarmi di più alla scrittura di un'app reale in C++. Ho un piccolo programma semplice che:

Assegna una memoria sull'host e sul dispositivo.
Inizializza l'array host su una serie di numeri. Copie l'array ospita una matrice dispositivo ritrovamenti quadrati di tutti gli elementi della matrice usando un kernel dispositivo Copie la matrice dispositivo torna alla matrice ospitante Stampa i risultati

Questo funziona se metto tutto in un file .cu ed eseguirlo. Quando lo divido in due file separati comincio a ricevere errori di collegamento. Come tutte le mie ultime domande, so che questo è qualcosa di piccolo, ma che cos'è?

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 

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

Sono in esecuzione Visual Studio 2008 su Windows 7 a 64 bit.

---

Edit:

penso di aver bisogno di elaborare su questo un po '. Il risultato finale che sto cercando qui è di avere una normale applicazione C++ con qualcosa come Main.cpp con l'evento int main() e le cose vanno da lì. A un certo punto nel mio codice .cpp voglio essere in grado di fare riferimento ai bit CUDA. Quindi il mio pensiero (e correggimi se c'è una convenzione più standard qui) è che inserirò il codice Kernel CUDA nei loro file .cu, e poi avrò un file .cu di supporto che si occuperà di parlare al dispositivo e di chiamare funzioni del kernel e cosa no.

Answer 1

Si sta includendo mykernel.cu in kernelsupport.cu, quando si tenta di collegare il compilatore vede mykernel.cu due volte. Dovrai creare un'intestazione che definisca TestDevice e includerlo.

re commento:

Qualcosa del genere dovrebbe funzionare

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

e quindi modificare il file incluso per

//KernelSupport.cu 
#ifndef _KERNEL_SUPPORT_ 
#define _KERNEL_SUPPORT_ 

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

re la tua modifica

Finché l'intestazione si l'uso nel codice C++ non ha alcun materiale specifico per cuda (__kernel__, __global__, ecc.) Si dovrebbe essere benissimo collegamento C++ e codice cuda.

Answer 2

La soluzione semplice consiste nel disattivare la creazione del file MyKernel.cu.

Proprietà -> Generale -> Sono esclusi dalla compilazione

La soluzione migliore è imo di dividere il kernel in un cu e un file cuh, e comprendono che, ad esempio:

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

Se si guardano gli esempi di codice dell'SDK CUDA, hanno extern C che definisce le funzioni di riferimento compilate da file .cu. In questo modo, i file .cu vengono compilati da nvcc e collegati solo al programma principale mentre i file .cpp vengono compilati normalmente.

Ad esempio, nel marchingCubes_kernel.cu ha il corpo della funzione:

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

Mentre in marchingCubes.cpp (dove risiede main()) appena ha una definizione:

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

Si può mettere anche questi in un file .h.

Answer 4

Ottenere la separazione è in realtà abbastanza semplice, si prega di controllare this answer per come configurarlo. Quindi inserisci semplicemente il codice host nei file .cpp e il codice del dispositivo nei file .cu, le regole di compilazione indicano a Visual Studio come collegarli insieme all'eseguibile finale.

Il problema immediato nel codice che si sta definendo la funzione __global__ TestDevice due volte, una volta quando si #include MyKernel.cu e una volta quando si compila il MyKernel.cu in modo indipendente.

Sarà necessario inserire un wrapper in un file .cu anche nel momento in cui si chiama TestDevice<<<>>> dalla propria funzione principale, ma quando lo si sposta in un file .cpp verrà compilato con cl.exe, che non Comprendere la sintassi <<<>>>. Pertanto chiameresti semplicemente TestDeviceWrapper(griddim, blockdim, params) nel file .cpp e fornisci questa funzione nel tuo file .cu.

Se si desidera un esempio, l'esempio SobolQRNG nell'SDK ottiene una buona separazione, sebbene utilizzi ancora cutil e consiglio sempre di evitare cutil.