cuda的_device_函数在另一个文件使用（分离编译和链接）

cuda的__device__函数在另一个文件使用（分开编译和链接）

类似c++，有这么一个使用场景。在一个文件定义一个函数，然后在另外一个文件使用。

刚开始发现总是编译出错。后来在http://developer.download.nvidia.com/compute/cuda/5_0/rel/docs/CUDA_Toolkit_Release_Notes_And_Errata.txt发现 这么一段文字

**  All __device__ functions can now be separately compiled and linked using NVCC. This allows creation of closed-source static libraries of __device__ functions and the ability for these libraries to call user-defined __device__ callback functions. The linker support is considered to be a BETA feature in this release. 

就是5.0之后，__device__函数是可以分开编译和链接的。

然后机智的我发现nsight这个IDE在创建cuda project的时候，有个选项可以勾选的，叫“separate compilation”。

这个选项在项目创建后，在项目属性的build->setting也可以修改的。

这里也提到分开编译链接。

http://stackoverflow.com/questions/13632180/static-library-with-multiple-h-and-cu-files-cant-resolve-functions

然后我就成功了。

函数头文件device.h

extern __device__ void helloworld();

其实就是声明一下函数

函数定义文件device.cu

#include <stdio.h>__device__  void helloworld(){	printf("Hello world!\n");}

函数调用文件test2.cu

#include <cuda.h>#include "device.h"__global__ void kernel(){	helloworld();}int main(void){	kernel<<<1,1>>>();	cudaDeviceSynchronize();return 0;}

编译

make all Building file: ../src/device.cuInvoking: NVCC Compiler/usr/local/cuda-5.5/bin/nvcc -G -g -O0 -gencode arch=compute_35,code=sm_35 -odir "src" -M -o "src/device.d" "../src/device.cu"/usr/local/cuda-5.5/bin/nvcc --device-c -G -O0 -g -gencode arch=compute_35,code=compute_35 -gencode arch=compute_35,code=sm_35  -x cu -o  "src/device.o" "../src/device.cu"Finished building: ../src/device.cu Building target: test2Invoking: NVCC Linker/usr/local/cuda-5.5/bin/nvcc --cudart static --relocatable-device-code=true -gencode arch=compute_35,code=compute_35 -gencode arch=compute_35,code=sm_35 -link -o  "test2"  ./src/device.o ./src/test2.o   -lcudadevrtFinished building target: test2

执行

Hello world!

其实也可以不用头文件，直接在调用文件加上extern __device__ void helloworld();一句就可以调用这个函数了。

同理，其实kernel函数也可以这样子。