CUDA __device__ Funzione esterna non risolta

Question

Sto cercando di capire come disaccoppiare i codici CUDA __device__ in file di intestazione separati.

Ho tre file.

File: 1: int2.cuh

#ifndef INT2_H_ 
#define INT2_H_ 

#include "cuda.h" 
#include "cuda_runtime.h" 
#include "device_launch_parameters.h" 

__global__ void kernel(); 
__device__ int k2(int k); 

int launchKernel(int dim); 

#endif /* INT2_H_ */ 

File 2: int2.cu

#include "int2.cuh" 
#include "cstdio" 

__global__ void kernel() { 
    int tid = threadIdx.x; 
    printf("%d\n", k2(tid)); 
} 

__device__ int k2(int i) { 
    return i * i; 
} 

int launchKernel(int dim) { 
    kernel<<<1, dim>>>(); 
    cudaDeviceReset(); 
    return 0; 
} 

File 3: CUDASample.cu

include <stdio.h> 
#include <stdlib.h> 
#include "int2.cuh" 
#include "iostream" 

using namespace std; 

static const int WORK_SIZE = 256; 

__global__ void sampleCuda() { 
    int tid = threadIdx.x; 
// printf("%d\n", k2(tid)); //Can not call k2 
    printf("%d\n", tid * tid); 
} 

int main(void) { 

    int var; 
    var = launchKernel(16); 

    kernel<<<1, 16>>>(); 
    cudaDeviceReset(); 

    sampleCuda<<<1, 16>>>(); 
    cudaDeviceReset(); 

    return 0; 
} 

Il codice funziona nel file. Posso chiamare il kernel sampleCuda() (nello stesso file), chiamare la funzione C launchKernel() (in un altro file) e chiamare direttamente kernel() (in un altro file).

Il problema che sto affrontando sta chiamando la funzione __device__ dal kernel sampleCuda(). quindi mostra il seguente errore. Tuttavia, la stessa funzione è richiamabile in kernel().

10:58:11 **** Incremental Build of configuration Debug for project CUDASample **** 
make all 
Building file: ../src/CUDASample.cu 
Invoking: NVCC Compiler 
/Developer/NVIDIA/CUDA-6.5/bin/nvcc -G -g -O0 -gencode arch=compute_20,code=sm_20 -odir "src" -M -o "src/CUDASample.d" "../src/CUDASample.cu" 
/Developer/NVIDIA/CUDA-6.5/bin/nvcc -G -g -O0 --compile --relocatable-device-code=false -gencode arch=compute_20,code=compute_20 -gencode arch=compute_20,code=sm_20 -x cu -o "src/CUDASample.o" "../src/CUDASample.cu" 
../src/CUDASample.cu(18): warning: variable "var" was set but never used 

../src/CUDASample.cu(8): warning: variable "WORK_SIZE" was declared but never referenced 

../src/CUDASample.cu(18): warning: variable "var" was set but never used 

../src/CUDASample.cu(8): warning: variable "WORK_SIZE" was declared but never referenced 

ptxas fatal : Unresolved extern function '_Z2k2i' 
make: *** [src/CUDASample.o] Error 255 

10:58:14 Build Finished (took 2s.388ms) 

Answer 1

Il problema è che è stata definita una funzione di __device__ in unità di compilazione separata dal __global__ che lo chiama. È necessario abilitare esplicitamente la modalità del dispositivo rilocabile aggiungendo il flag -dc o spostare la definizione sulla stessa unità.

Da nvcc documentazione:

--device-c|-dc compilare ciascuno//// .c .cc cpp .cxx br ingresso in un file oggetto contenente codice del dispositivo rilocabile. È equivalente a --relocatable-device-code = true --compile.

Vedere Separate Compilation and Linking of CUDA C++ Device Code per ulteriori informazioni.