CUDA SDK VolumeRender:跨编译单元调用

有关VolumeRender的介绍、和CUDA SDK中的VolumeRender解析在HERESY的一些文章中写的非常详细，这里我只想写写我对VolumeRender这个例子的一些理解。

　　曾经遇到过cuda函数在跨编译单元调用的问题，这个例子用到了一个很巧妙的解决方法。

　　首先描述下这个问题，当多个cu或cpp文件互相包含的时候cu文件中的实现会被nvcc生成在多个编译单元中，从而出现重定义链接错误。具体描述见这里。

　　CUDA SDK中避免此问题的方法是，添加一个extern层接口，需要被外部使用的功能封装为extern接口，接口函数和实现函数放在一个cu中，调用端通过extern的声明来找到接口，这样避免了实现出现在多个编译单元内(只有函数实现所在的编译单元)。

　　有点乱?下面给出VolumeRender中的用法：

　　实现部分：

// volumeRender_kernel.cu
__global__ void
d_render(uint *d_output, uint imageW, uint imageH,
 float density, float brightness,
 float transferOffset, float transferScale)
{
 // implemention
}

　　全局接口和实现在同一文件：

// volumeRender_kernel.cu
extern "C" void render_kernel(dim3 gridSize, dim3 blockSize, uint *d_output, uint imageW, uint imageH,
 float density, float brightness, float transferOffset, float transferScale)
{
 d_render<<>>( d_output, imageW, imageH, density,
 brightness, transferOffset, transferScale);
}

　　调用端使用声明绕开包含：

　　 1 : // volumeRender.cpp

　　 2 : extern " C " void render_kernel(dim3 gridSize, dim3 blockSize, uint * d_output, uint imageW, uint imageH,

　　 3 : float density, float brightness, float transferOffset, float transferScale);

　　如果你对这个问题很感兴趣，这里有一个完整的解决方案。

http://cuda.it168.com/a2011/1018/1260/000001260437.shtml