【IT168 技术】开发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里面了。

        更多内容请点击：

        CUDA专区：http://cuda.it168.com/

        CUDA论坛：http://cudabbs.it168.com/