CUDA处理非连续内存段的方法

最近在写GPUMeanShift，碰到一个非常头疼的问题，就是CPU与GPU之间的非连续内存拷贝。

 

按照以前的思路，对于一段连续CPU内存

float* h_data; // cpu data

float* d_data; // gpu data

cudaAlloc((void**)&d_data, size);

cudaMemcpy(d_data, h_data, size);

 

但是对于非连续的CPU内存 float** h_discontinue_data处理起来似乎就不如上面那面简单了。

 

我尝试了下列做法：

1. 将float** h_discontinue_data拷贝至float* h_continue_data，然后进行上述处理。

2. 将float** h_discontinue_data绑定到上传至三维纹理，然后注册到cuda端

3. 将float** h_discontinue_data依层拷贝至float** d_discontinue_data，然后写个kernel，将float** d_discontinue_data再拷贝到float* d_continue_data

 

最终结果

不管是时间复杂度和空间复杂度，2都是最优的。但是texture操作相对连续内存操作而言更不灵活，且不能对齐。

从实现速度上来看，1，是能解决问题的最快方案。但是空间复杂度较大。

从代码重用角度看，3是最好的。灵活程度也是最高的。

 

最终我的做法是：

将我的Volume数据分成若干层Image，然后分别得到这些Image的MemObject，最后用了一个ROI拷贝的Kernel完成所有工作。

事实上，我的代码都在写如何调用那些东西，大概只有20行的样子。

所有的可重用代码，都封装的CUDA Operation和CUDA Image里面了。