system
1
我的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.xblockDim.x+threadIdx.x;
int j=blockIdx.yblockDim.y+threadIdx.y;
int idx=jcol+i;
if(i<256&&j<256)
{
c[idx] = a[idx] + b[idx];
}
}
如果我想一个线程算两个数据,或者4个数据,该如何修改代码?
system
2
楼主您好,
这个逻辑上应该是很清晰的。例如您可以直接计算相邻的两个点:
(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;
您觉得呢?
system
3
谢谢玫瑰版主
其实我这个程序 本来是复数运算的,我为了简化说明改成了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下标该如何设置呢?
system
4
楼主以后请不要简化代码,任何给出的范例均在您当前给出的代码的情况下适用的。
如果您给出的当前代码是某种代码的简化,则范例不一定适用于它。
以及,此问题,如果要一般的用,建议您还可以这样:
元素规模是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
这个一般通用点。
请注意,这样依然是连续访问的。
感谢来访。
system
5
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呢?
system
6
楼主您好,请问题建议您新开主题。
(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的第一章,阅读完几页再回来看我的回复,会加深理解。
(手册在论坛的资源下载区有下载)
感谢您的来访。