用户
 找回密码
 立即注册
quanzhang100 该用户已被删除
发表于 2013-11-25 10:13:04
97817
我的kernel是的一个256x256的矩阵加法运算,我想找出占用率与ILP之间的平衡点
我现在线程结构是这样的
dim3 dimBlock(16,16);
dim3 dimGrid(16,16);
MatrixAdd<<<dimGrid,dimBlock>>>(a,b,c);
kernel如下:
__global__ void   MatrixAdd(double *a,double *b, double *cl)
{
        int i=blockIdx.x*blockDim.x+threadIdx.x;
        int j=blockIdx.y*blockDim.y+threadIdx.y;
        int idx=j*col+i;
        if(i<256&&j<256)
        {
                c[idx] = a[idx] + b[idx];
        }
}

如果我想一个线程算两个数据,或者4个数据,该如何修改代码?

使用道具 举报 回复
发表于 2013-11-25 12:48:41
楼主您好,

这个逻辑上应该是很清晰的。例如您可以直接计算相邻的两个点:
(1)保留您当前的block_shape不变。降低您的dimGrid到(8,16)
(2)将您的if改成:if (i < 128 && j < 256)
(3)将if里面改成:
double2 t0 = ((double2 *)a)[idx];
double2 t1 = ((double2 *)b)[idx];
double2 t2; t2.x = t0.x + t1.x; t2.y = t0.y + t1.y;
((double2 *)c)[idx] = t2;

您觉得呢?
使用道具 举报 回复 支持 反对
发表于 2013-11-25 14:43:28
玫瑰幻想 发表于 2013-11-25 12:48
楼主您好,

这个逻辑上应该是很清晰的。例如您可以直接计算相邻的两个点:

谢谢玫瑰版主
其实我这个程序 本来是复数运算的,我为了简化说明改成了double运算。
也就是说我如果想让每个thread做4次复数加法 每个thread怎么及要处理好合并内存访问又要处理好ILP
,kernel函数又怎么兼顾呢? 比如像你代码中说的那样要做多个double2类型的计算,
double2 t0 = ((double2 *)a)[idx];
double2 t1 = ((double2 *)b)[idx];
double2 t2; t2.x = t0.x + t1.x; t2.y = t0.y + t1.y;
((double2 *)c)[idx] = t2;

thread的idx下标该如何设置呢?
使用道具 举报 回复 支持 反对
发表于 2013-11-25 15:00:21
quanzhang100 发表于 2013-11-25 14:43
谢谢玫瑰版主
其实我这个程序 本来是复数运算的,我为了简化说明改成了double运算。
也就是说我如果想让 ...

楼主以后请不要简化代码,任何给出的范例均在您当前给出的代码的情况下适用的。

如果您给出的当前代码是某种代码的简化,则范例不一定适用于它。

以及,此问题,如果要一般的用,建议您还可以这样:
元素规模是N, 线程数量是M, (例如您可以选择M=N/4),
这样一个线程将分配到4个计算量。

然后每个线程里面:
(1)计算当前线程线性ID
(2)分别读取下标为[ID], [ID + M], [ID + 2 * M], [ID + 3 * M]的a和b
(3)分别计算4组a,b的和
(4)保存4组的结果到c

这个一般通用点。

请注意,这样依然是连续访问的。

感谢来访。
使用道具 举报 回复 支持 反对
发表于 2013-11-25 15:19:14
横扫千军 发表于 2013-11-25 15:00
楼主以后请不要简化代码,任何给出的范例均在您当前给出的代码的情况下适用的。

如果您给出的当前代码是 ...

dim3 dimBlock(16,16);
dim3 dimGrid(16,16);
MatrixAdd<<<dimGrid,dimBlock>>>(a,b,c);
我的kernel函数执行配置是用dim3配置的,我有如下几点疑惑
(1)配置线程执行结构时用一维的和二维的有什么区别?
(2)修改dimGrid 和修改dimBlock有什么区别,如果总的thread大小一定,缩小grid大小与缩小block大小对占用率有什么影响?
(3)每个block是分配到SM上执行,对于小的block是不是表示每个线程分配的资源比较多有利于ILP呢?
使用道具 举报 回复 支持 反对
发表于 2013-11-25 15:34:29
quanzhang100 发表于 2013-11-25 15:19
dim3 dimBlock(16,16);
dim3 dimGrid(16,16);
MatrixAdd(a,b,c);

楼主您好,请问题建议您新开主题。

(1)1维的,相当于第二和第三个维度都是1, 等于(..., 1, 1)这种。
同理,2维的,可以看成是(..., ..., 1)

(2)因为您限定了前提,总的threads数目一定,那么我们可以知道,如果block大了,自然grid里的总blocks数目会减少。反之,如果block过小,则总blocks数目会较大。
一般情况下,建议初学者不要使用过大的block(1024这种的),使用256之类的较好。初学者可能在较大的blocks上取得较低的achieved occupancy.
但一般情况下,如果代码写的好点,过大和过小的其实都没事。(对于您,我建议您总是使用256这样的,除非您的问题的确需要较大的block)

(3)一个线程占用的资源多少和block的大小无关的,因为属于线程的资源基本上主要是寄存器。你使用过小的block,对于block所拥有的资源(例如shared memory), 势必只能导致在较少数量的线程中共用,也是不利的。所以不要追求过小,不好的。

关于这些疑问,实际上因为您没有基本的对CUDA的认识造成的,建议您打开CUDA C Programming Guide的第一章,阅读完几页再回来看我的回复,会加深理解。
(手册在论坛的资源下载区有下载)

感谢您的来访。
使用道具 举报 回复 支持 反对
发表于 2013-11-25 15:41:00
玫瑰幻想 发表于 2013-11-25 15:34
楼主您好,请问题建议您新开主题。

(1)1维的,相当于第二和第三个维度都是1, 等于(..., 1, 1)这种。

好的 谢谢各位版主
使用道具 举报 回复 支持 反对
发表于 2013-11-25 15:47:09
quanzhang100 发表于 2013-11-25 15:41
好的 谢谢各位版主

您客气了,服务您是我们的荣幸。

期待您的下次莅临。
使用道具 举报 回复 支持 反对
发新帖
您需要登录后才可以回帖 登录 | 立即注册