根据几年的CUDA开发经验，简单的介绍下CUDA程序的大概开发步骤，按照先修改CPU串行程序后移植到GPU平台的原理，把需要在GPU上做的工作尽量先在CPU平台上修改，降低了程序的开发难度，同时有利用bug的调试。通过实现一种快速、有效地CUDA并行程序开发的方法，提高CUDA并行程序开发效率，降低CUDA并行程序开发周期和难度。
（1） CPU串行程序分析
对于CPU串行程序，首先需要测试串行程序中的热点函数，以及分析热点函数的并行性：
a) 热点测试
根据时间的测试结果确定热点函数，作为后面移植的重点代码模块。
b) 并行性分析
找出热点代码后，需要分析热点部分的算法、数据特点，根据算法和数据的特点分析其是否可以并行，是否适合细粒度并行。
c) 确定CUDA内核使用的数组
根据对串行程序的分析，确定哪些模块需要移植到GPU平台上运行，对于需要运行上GPU平台上的代码内的数据进行分析，确定哪些数组需要在CUDA内核中使用，分析CUDA计算时，这些数组传递的方向是CPUtoGPU还是GPUtoCPU，以及每次传递时的数据大小等信息，然后设计这些数组的定义方式和大小。
（2） 仿CUDA格式的CPU串行程序
CUDA程序相对CPU程序比较复杂，当出现bug时，调试的难度也要比CPU程序大很多，为了降低CUDA程序开发难度和周期，可以把一些GPU上的移植工作提前在CPU平台上实现，具体涉及下面几个方面：
a) 修改成可并行算法
对于CPU串行程序，有些代码理论上可以并行，但经过CPU版本的优化之后导致代码不能直接并行化，这时需要根据并行算法的要求修改原程序，改成可并行的模式；有些模块理论上是可以并行的，但串行算法无法直接并行化，需要重新设计并行算法。
b) 数组修改
CPU串行程序中使用的数组定义的形式有可能无法在CUDA内核中直接使用，这时需要对数组的定义进行修改，如C语言程序，结构体中的指针需要改变成单独的指针/数组，才能进行CPU与GPU之间的数据传递。另外，考虑到全局存储器合并访问的问题，有时还需要对数组的访问方向进行修改，从而也需要改变数组的定义形式（如做行列变换）。总之，根据CUDA对数组使用和CPU串行程序之间的区别，提前把数组修改，方便程序的调试。
根据前面几条的修改方式，对原CPU程序修改成一个仿CUDA格式的CPU串行程序，为后面的移植工作做大量的准备，有利于后面CUDA程序的移植。
（3） GPU数组设计、并行模型设计
GPU数组设计：设计GPU数组大小、类型、维度等信息；
设计CPU与GPU之间的数组通信方式；
并行模型设计：block和grid的设计满足算法的数据特点。
（4） CUDA并行程序基本版本
根据对原程序数组的分析，把CPU串行程序移植到GPU平台，根据热点模块的算法和代码实现CUDA内核代码。
a) 设计调用语句
Kernel<<<grid, block, …>>>(…);
b) 设计CUDA内核
根据算法的并行性分析，设计内核，划分每个线程的计算任务，利用同步语句满足内核程序的逻辑正确性。
（5） CUDA并行程序优化版本
根据步骤4实现的基本版本的CUDA并行程序，利用CUDA的优化技术进一步提高并行程序的性能，主要优化包括2个方面：通信优化和内核优化。
a) CPU与GPU通信优化
GPU计算需要CPU与GPU之间进行数据的传递，合理的利用通信优化技术有利用提高GPU并行程序的性能，如流技术隐藏CPU与GPU的通信。
b) CUDA内核优化
CUDA内核的优化对其性能更为重要，主要涉及存储器访问优化和指令流优化，存储器访问优化包括：全局存储器合并访问，利用共享存储器、常量存储器、纹理存储器替换全局存储器的访问，提高访问速度；指令流优化是指利用高效的指令代替低效的指令，如CUDA中的快速函数。
以上方法只是大概介绍，难免有介绍的不全的地方，希望对大家的开发有点帮助。
原文地址：http://blog.csdn.net/zhang0311/article/details/39479913
