用户
 找回密码
 立即注册
quanzhang100 该用户已被删除
发表于 2013-10-29 10:47:20
132789
/****************************************************************
**    计算t1        t1 = D1.*conj(H1) + D2.*conj(H2);                           **
**    计算t2        t2 = abs(H1).^2+abs(H2).^2 + gama;                   **
*****************************************************************/
__global__ void   t1_t2(compx* H1_d,compx* H2_d,compx* D1_d,compx* D2_d,compx* t1_d,double *t2_d,int row,int col)
{
        int i=blockIdx.x*blockDim.x+threadIdx.x;
        int j=blockIdx.y*blockDim.y+threadIdx.y;
        int idx = i*col+j;
        if(i<row&&j<col)
        {

                t1_d[idx].real = (D1_d[idx].real * H1_d[idx].real + D1_d[idx].imag * H1_d[idx].imag)+
                        D2_d[idx].real * H2_d[idx].real + D2_d[idx].imag * H2_d[idx].imag;
                t1_d[idx].imag = (H1_d[idx].real * D1_d[idx].imag - D1_d[idx].real * H1_d[idx].imag)+
                        H2_d[idx].real * D2_d[idx].imag - D2_d[idx].real * H2_d[idx].imag;

                t2_d[idx] = (H1_d[idx].real * H1_d[idx].real + H1_d[idx].imag * H1_d[idx].imag) +
                        (H2_d[idx].real * H2_d[idx].real + H2_d[idx].imag * H2_d[idx].imag);
        }
}
上面是我的kernel,其中D1 H1 都重复使用,想把它们放的share memory 可是share memory只有48k,上面这个kernel有没有好的方法可以优化呢?

使用道具 举报 回复
发表于 2013-10-29 10:54:14
楼主您好:

如果想将大于最大shared memory容量的数组强制放入shared memory,
这个真心无法放入。也是强人所难。

以及,您线程间对D1[]和H1[]中的元素都无重复使用(您的下标是[idx]), 您可能需要使用shared memory的迫切性真心不大。

您最应该解决的是您的下标问题,您的计算:
int i=blockIdx.x*blockDim.x+threadIdx.x;
int j=blockIdx.y*blockDim.y+threadIdx.y;
int idx = i*col+j;
的多个线程的idx之间严重的不连续,造成了巨大的读取上的效率降低。

这才是您真的应该考虑的因素。
(而不是去考虑Shared memory)。

使用道具 举报 回复 支持 反对
发表于 2013-10-29 11:13:36
横扫千军 发表于 2013-10-29 10:54
楼主您好:

如果想将大于最大shared memory容量的数组强制放入shared memory,

版主我没看明白这里的idx为什么不是连续的,要怎么改才是连续的呢?谢谢哈
使用道具 举报 回复 支持 反对
发表于 2013-10-29 11:41:37
quanzhang100 发表于 2013-10-29 11:13
版主我没看明白这里的idx为什么不是连续的,要怎么改才是连续的呢?谢谢哈 ...

LZ您好:

您的写法是:

        int i=blockIdx.x*blockDim.x+threadIdx.x;
        int j=blockIdx.y*blockDim.y+threadIdx.y;
        int idx = i*col+j;
一般而言此时是二维的block和grid。

CUDA中线程编号是x优先变化的,考虑最简单的情况,您的第一个block,上式化简为
int i=threadIdx.x;
int j=threadIdx.y;
考虑block一行中相邻的线程,j是相同的,i是连续变化的。
此时您的
int idx = i*col+j;
一行中相邻的两个线程的idx相差col,并且一般而言col是一个较大的数。

这样您一个warp中的各个线程的访存位置就相差col个int元素,这将引起非常严重的global 非合并访问,将极大地影响您的程序效能。

如果您修正了这个问题,(直接只考虑合并访问的话,写成j*col+i就可以,但是这么写可能和您的算法安排就不对应了,请您通盘考虑),那么按照您的写法,无需再使用shared memory。因为您每个线程都只是访问自己对应的数组中的那个元素,这个值会自动使用该线程的寄存器保存,以及这里并没有线程之间需要“shared data”的情况,因此您直接用即可。

祝您编码顺利~
使用道具 举报 回复 支持 反对
发表于 2013-10-29 14:19:29
ice 发表于 2013-10-29 11:41
LZ您好:

您的写法是:

谢谢版主哈,我忘记线程编号是按照x优先变化了。
使用道具 举报 回复 支持 反对
发表于 2013-10-29 14:21:14
quanzhang100 发表于 2013-10-29 14:19
谢谢版主哈,我忘记线程编号是按照x优先变化了。

不客气的,欢迎您常来~

祝您编码顺利~
使用道具 举报 回复 支持 反对
发表于 2013-10-29 14:35:19
ice 发表于 2013-10-29 14:21
不客气的,欢迎您常来~

祝您编码顺利~

我代码中的col=row=256,我的block大小为(16,16),grid大小为(16,16)我现在用idx=j*col+i之后DRAM Utilization 为52.8%(96.81Gb/s)  请问一下还能否提高呢?
使用道具 举报 回复 支持 反对
发表于 2013-10-29 14:38:52
quanzhang100 发表于 2013-10-29 14:35
我代码中的col=row=256,我的block大小为(16,16),grid大小为(16,16)我现在用idx=j*col+i之后DRAM Ut ...

请记录当前的运行时间,

然后在您的if(i < row && j < col)里面的第一句加上:
compx a = D1_d[idx];
compx b = H1_d[idx];
然后将下文的所有D1_d[idx]改成a,将所有的H1_d[idx]改成b,
重新编译并运行。

再次继续这次的运行时间。

然后请您比较下时间变化。
(感谢ICE提供建议)
使用道具 举报 回复 支持 反对
发表于 2013-10-29 15:10:55
玫瑰幻想 发表于 2013-10-29 14:38
请记录当前的运行时间,

然后在您的if(i < row && j < col)里面的第一句加上:

确实提高了不少,现在的DRAM Utilization为67.8%(124.3GB),是不是这样修改后每个数据都是在register中,所以速度提高了呢?
使用道具 举报 回复 支持 反对
发表于 2013-10-29 15:12:09
quanzhang100 发表于 2013-10-29 15:10
确实提高了不少,现在的DRAM Utilization为67.8%(124.3GB),是不是这样修改后每个数据都是在register中 ...

嗯嗯。以及,建议您以时间为准,较好。

以及,您还可以考虑使用use_fast_math, 看看时间有无变化。

谢谢。
使用道具 举报 回复 支持 反对
发新帖
您需要登录后才可以回帖 登录 | 立即注册