CUDA学习笔记

来源：互联网发布：增值税发票抵扣软件编辑：程序博客网时间：2024/05/17 02:19

CUDA学习笔记

使用3.2.16版的SDK
生成64位的程序前，要在VS2005/VS2008上安装64位的编译器，默认是不安装的。
shrUtils32.lib/shrUtils32D.lib出现错误时（一般是链接错误，特别是找不到某些函数的定义），可以把找到他们的工程(一般在NVIDIA Corporation\NVIDIA GPU Computing SDK 3.2\shared目录下)，重新生成一下就可以了，原因不是很清楚，猜测是生成这个LIB时打不打SP1会有所不同。
3.2版的SDK与2.x版的SDK差别还是比较大的，很多东西都不再一样，那个cuda.rules好像也不用了。
错误信息为“模块计算机类型“X86”与目标计算机类型“x64”冲突”——解决方法：项目属性里“链接器”-“高级”-“目标计算机”与“CUDA Runtime API”-“host”里的“Target machine platform"选项要对应上。
cuda并不是一种完全硬件透明的语言，程序员需要根据硬件先点将任务进行合理的分解，在编程时对数据传输和存储器访问进行优化。
CUDA_SAFE_CALL(cudaMalloc(size))在显存golbal memory上分配大小为size字节的线性空间。注意：cudaMalloc()必须与cudaFree()结合使用。
extern关键字在设备端表示动态分配，而在主机端函数中则表示外部变量。如果要静态分配一块shared memory，那么在__shared__之前不要加extern，还必须在[]中写上要分配的字节数。
CUDA主机端代码可以支持完整的C++语法，而设备端代码则不能完全支持。
kernel函数完整的参数配置形式是<<<Dg,Db,Ns,S>>>，Dg用于定义整个grid的维度和尺寸，为dim3类型，但实际上只有前两维可以不为1.前两维的最大为65535，也就是说一个grid中最多也只能有65535*65535个block；参数Db为dim3类型，用于定义每个block的维度和尺寸。Db.x,Db.y最大值为512，Db.z最大值为64，三个维度之积小于768或1024；参数Ns是一个可选参数，用于设置每个block除了静态分配的shared memory以外，最多能够动态分配的shared memory大小，单位为Byte。
如果希望在VS里用Visual AssistX来帮助写CUDA的.cu文件，可以参考保存在wiz里的网页部分，或者http://www.wholetomato.com/forum/topic.asp?TOPIC_ID=5481
在CUDA的核函数里，常常会用到"if(tid<N)"的语句，主要为了避免读取或者写入不应该操作的内存，造成不可预料的错误。
CUDA核函数里，很重要的一步是计算tid(或者叫offset)，即该线程在总数组（矩阵）里的位置，计算的不正确产生的错误很不好查。
规约计算的本质是并行计算，总的计算次数并没有减少，但如果每一次规约都是并行的话，时间复杂度为log₂N
规约计算的最后可能只有很少的数据，可以直接放到CPU上来算，因为那样会浪费GPU的计算资源。
线程同步__syncthreads()函数相当好用，但也要慎用，必须保证线程块里的每一个线程都会执行，否则会一直等待下去。
初看矩阵相乘的例子，先自己设想了一个思路：每一个block计算一个C矩阵的值，也就是说这个block里的thread负责矩阵里两个一维向量的内积；看了一下书中的例子，都是先进行分块的，比如分割成16*16的小块，每一个小块用一个block进行计算；想了一会儿才明白：使用shared memory的话，每一小块只读进行一次，比如A的前16行，B的前16列，都只读一次的话就可以计算得到16*16个C的数；而自己原来的想法里每一个block只计算得C里的一个值，每一个值都要对A的一行与B的一列进行一次读取；分块的算法相当于我的初始想法的1/16的读取（不是1/16*16）。当然这个块并不是分的越大越好，因为还有shared memory的大小限制，如果进行权衡，也是很重要的。
blockIdx,BlockDim等变量定义在device_launch_parameters.h里
有时候，主机端定义的grid或者block的尺寸与核函数里的使用不想符的话，得不到正确的结果，但这种错误是很难调试出来的。
一定要注意x是横向的而y是纵向的。

http://www.cnblogs.com/yanglp/archive/2011/04/24/2026123.html