四、CUDA性能
CUDA中的block被划分成一个个的warp，在GeForce8800GTX上，一个warp有32个线程。若不够32个线程，则padding相应数目的线程。Warp中的线程ID是连续且递增的。对于二维组织的线程来说，先把threadIdx.y为0的线程按照threadIdx.x从小到大排，然后把threadIdx.y为1的线程按照threadIdx.x从小到大的顺序排列成warp。对于三维组织的线程来说，先排列 threadIdx.z为0的二维线程，再排列threadIdx.z为1的二维线程，以此类推。
任何时刻，硬件都只能一次选择执行一个warp。
下面2个图是执行元素总和的操作，不同的算法实现，其效率不同。第二种方法使得线程没有分支。
一个grid中的所有线程执行相同的内核函数，通过坐标进行区分。这些线程有两级的坐标，blockId和threadId，由CUDA runtime system指定。grimDim.x标识block在x维度上的数目，gridDim.y标识block在y维度上的数目。例如,
转自新浪博客： http://blog.sina.com.cn/s/blog_48b9e1f90100fm56.html
一、CUDA 编程模型
2009-10-21
CUDA的代码分成两部分，一部分在host(CPU)上运行，是普通的C代码；另一部分在device(GPU)上运行，是并行代码，称为kernel，由nvcc进行编译。
Kernel产生的所有线程成为Grid。在并行部分结束后，程序回到串行部分即到host上运行。
在CUDA中，host和device有不同的内存空间。所以在device上执行kernel时，程序员需要把host memory上的数据传送到分配的device memory上。在device执行完以后，需要把结果从device传送回host，并释放device memory。CUDA runtime system提供了API给程序员做这些事情。
