CUDA编程架之构解析篇(0)

CUDA编程架之构解析篇(0) [已结贴,结贴人:Cyrosly]

Cyrosly
该用户很懒，没有填写自己的个人昵称
等级:

发表于：2008-05-01 21:13:53 楼主

这个斑竹没有做好啊,近期一直被公司项目所累,所以来得少,以后再接再厉.仓促中写的,请大家指教.转自我的blog:http://blog.csdn.net/Cyrosly/

在开始写这篇文章前的数天里,我曾经为写个公司的个人工作描述文档而大伤脑筋,项目组的负责人在让我们每个人写个自己目前所做项目的描述和时间安排等,而我们开始却认为不过是写简单的文字描述罢了.结果第一次我们全部都要求重写.而我的原因是认为是废话的都不许要写,结果写完后仅寥寥几行,由于过于简单所以重写.而在他的"指导"下完成的最终的内容在我看来却仍然几乎全部是废话外加一些看似增加了文档内容的充实却又简单的过于笼统且不能从中得到任何可用信息的示意图.而他的理由是公司检查文档的那些人都是非技术人员,她们只是看你的文档的书写规范和格式.所以我现在真的不知道如何着手写这篇文章,以前都只是网友提相关的问题,我根据他们的需要解答或写些什么.而现在却不知道该如何写才是CUDA技术学习人员所需要的.所以我决定:避繁就简,量少而精.(前面这些文字也是废话^-^)

执行kernel时并非任何时候都需要所有的MPS和其相关资源.例如对于下面如此简单的kernel,仅一个block就可以了,也就是仅需要一个MSP既足够了:

#define IMUL(X,Y) __mul24(X,Y)
#define IMAD(X,Y,Z) __mul24(X,Y)+Z
extern "C"
__global__ void add_kernel(float* o,float* i0,float* i1)
{
unsigned int gtidx=IMAD(blockIdx.x,blockDim.x,threadIdx.x);
unsigned int gtidy=IMAD(blockIdx.y,blockDim.y,threadIdx.y);
unsigned int otid =IMAD(IMUL(gridDim.x,blockDim.x),gtidy,gtidx);
o[otid]=i0[otid]+i1[otid];
}

而在应用程序中如何设置:

cuFuncSetBlockShape(func,16,16,1);
...
cuLaunch(func);
or
add_kernel < < <dim3(1,1,1),dim3(16,16,1)>>>(o,i0,i1);

注意(重点):另外,如果thread的总数量不超过768(当然,对于cuda CC <=1.1,至少目为至,因为将来的cuda CC1.2或2.0情况也可能不变),则各个MSP可以同时处理8个block,也即对于被分配的block group的同一个MSP,如果希望8个SP同时处理8个block,则每个block的thread数量不应大于96.而是否这样处理取决于你在程序中的设置.为了简单起见我们假设当前thread总数即为768,那么如下的代码将决定MSP[0]的8个SP各处理一个block,切至少理论上8个SP应该是并行执行的(以上面的kernel为例):

...

cuFuncSetBlockShape(func,4,24,1);//It have 96 threads per block

cuFuncSetSharedSize(func,28);

unsigned int offset=0;

cuParamSeti(func,oDevPtr,offset );offset+=sizeof(CUdeviceptr);

cuParamSeti(func,iDevPtr0,offset);offset+=sizeof(CUdeviceptr);

cuParamSeti(func,iDevPtr1,offset);offset+=sizeof(CUdeviceptr);

cuParamSetSize(func,offset);

cuLaunchGrid(func,8,1);//gridDim be set to 8x1,it means there are 8 blocks

cuCtxSyncthronize();

...

还应该注意kernel代码中的blockDim与gridDim的使用,如果使用不当,即使正确执行切是希望的效果,但内部MSP的执行过程却未必如你所愿

示意图是用用window系统自带的画图程序画的(不容易啊),所以虽然很多地方感觉表达不当,但也懒得改了^-^

问题点数：20 回复次数：9 修改删除举报引用回复