CUDA SDK VolumeRender 分析 (1)

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

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

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

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

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

实现部分

  1: // volumeRender_kernel.cu
  2: __global__ void
  3: d_render(uint *d_output, uint imageW, uint imageH,
  4:          float density, float brightness,
  5:          float transferOffset, float transferScale)
  6: {
  7: // implemention
  8: }

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

  1: // volumeRender_kernel.cu
  2: extern "C"
  3: void render_kernel(dim3 gridSize, dim3 blockSize, uint *d_output, uint imageW, uint imageH, 
  4:            float density, float brightness, float transferOffset, float transferScale)
  5: {
  6:   d_render<<<gridSize, blockSize>>>( d_output, imageW, imageH, density, 
  7:                     brightness, transferOffset, transferScale);
  8: }

调用端使用声明绕开包含

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

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