CUDA __device__ Unresolved externe Funktion

Question

Ich versuche zu verstehen, wie CUDA __device__ Codes in separaten Header-Dateien zu entkoppeln.

Ich habe drei Dateien.

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

Datei 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; 
} 

Die Code-Datei. Ich kann den Kernel sampleCuda() aufrufen (in der gleichen Datei), die C-Funktion launchKernel() (in einer anderen Datei) aufrufen und kernel() direkt aufrufen (in einer anderen Datei).

Das Problem, mit dem ich konfrontiert bin, ruft die __device__ Funktion von sampleCuda() Kernel. dann zeigt es den folgenden Fehler an. Die gleiche Funktion ist jedoch in kernel() aufrufbar.

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

Das Problem ist, dass Sie eine __device__ Funktion in separater Übersetzungseinheit aus __global__ definiert, die es nennt. Sie müssen entweder verschiebbaren Gerätecode Modus explizit aktivieren, indem Sie -dc Flag hinzufügen oder Ihre Definition zu derselben Einheit verschieben.

Von nvcc Dokumentation:

--device-c|-dc Compile jede .c/.cc/CPP/.cxx/.cu Eingabedatei in eine Objektdatei, die relocatable Gerätecode enthält. Es entspricht --relocatable-device-code = True --compile.

Weitere Informationen finden Sie unter Separate Compilation and Linking of CUDA C++ Device Code.