找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 148|回复: 4

求复矩阵的每一个元素的模的平方

[复制链接]
发表于 2018-9-7 12:58:41 | 显示全部楼层 |阅读模式
GTC
假如有一个复矩阵H,大小为1024*2048,要求abs(H).^2,下面是我使用的代码,感觉没有问题但是就是计算结果不对,请各位指教N = 1024
circle = 2
tcol = 1024
size_t pitch;
cuComplex *d_trf;       
cudaMallocPitch((void **)&d_trf, &pitch, sizeof(cuComplex)*N*circle, tcol)
pitch = pitch / sizeof(cuComplex);
__global__ void abs_Kernel(cuComplex *abs_trf, size_t pitch)
{
        int tid = threadIdx.x;
        int bidx = blockIdx.x;
        int bidy = blockIdx.y;
        abs_trf[tid * pitch + bidx * gridDim.y + bidy].x = \
                abs_trf[tid * pitch + bidx * gridDim.y + bidy].x * \
                abs_trf[tid * pitch + bidx * gridDim.y + bidy].x + \
                abs_trf[tid * pitch + bidx * gridDim.y + bidy].y * \
                abs_trf[tid * pitch + bidx * gridDim.y + bidy].y;
}


abs_Kernel << <(circle, N, 1), (1024, 1, 1) >> >(d_trf, pitch);



回复

使用道具 举报

发表于 2018-9-10 15:05:24 | 显示全部楼层

回帖奖励 +1 金钱

Jetson TX2
Lib and Lab:
代码可能存在诸多方面的问题:

(1)正常人的访存写法是尽量连续合并的读写的。像是楼主这样的大跨步的跳跃式的读写(用threadIdx.x乘以一个大pitch),比较诡异。虽然这只是效率上的问题。但如果给出了错误的pitch值(例如你混淆了元素单位和字节单位),则可能出错或者kernel挂掉。因为并不知道这里Lib and Lab妹子你这样写的原因,所以请回头仔细检查一下。

(2)楼主你没有进行任何越界检查,虽然对于规整的形状(例如你这里的1024和2048元素,pitch往往总是等于行宽或者行元素数量),但可能会导致潜伏的BUG,并在形状变化的时候暴露。这里也建议进行检查。或者你可以直接使用cuda-memcheck 加上你的可执行文件的名字的方式, 进行快速检测。

(3)正常这种求平方计算不应当使用单独的kernel的,而应该集成在实际使用的kernel中(例如需要使用|A|^2的代码中,就地计算),单独开个kernel进行一次读取+一次写入来得到这个值,在目前你能找到的所有GPU卡中,都是不合算的。(这条对于只有1/32的double性能的卡, 在只使用double计算的情况下除外,但你不应该做这种卡上使用double的)。这条只是增强性建议。

请妹子仔细阅读,并进行实验和反馈。
PS: 大部分从matlab来的人都会弄错横向和纵向。

感谢来访。
回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-9-10 19:22:35 | 显示全部楼层
Tesla P100
屠戮人神 发表于 2018-9-10 15:05
Lib and Lab:
代码可能存在诸多方面的问题:

在我将abs_Kernel << <(circle, N, 1), (1024, 1, 1) >> >(d_trf, pitch);这个改为
abs_Kernel << <(circle, N, 1), 1024 >> >(d_trf, pitch);之后结果就对了,不知道这是什么原因呢。还有就是我这里单独开一个核函数的原因是之前运行了cuFFT才得到的这个结果,所以只能分开写。我要那个样子访问其实是因为数据量起始会编程1024*1024*N,通俗的说我想把数据按照1024*1024这样的多个矩阵横着排在一起,所以我按照这个思维方式来做的,Pitch其实是我cudaMallocPitch得到的Pitch,如果我想要让这个核函数加速的话我应该怎么做呢,目前我测时间的话这个核函数需要大量的时间。
回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-9-10 20:06:53 | 显示全部楼层
屠戮人神 发表于 2018-9-10 15:05
Lib and Lab:
代码可能存在诸多方面的问题:

我的问题其实是这样的,要对1024*1024*N个复数数据进行FFT,然后再求模的平方,我的思路是将这1024*1024*N个数据横着排列,即1024*1024作为一个矩阵,然后这N个矩阵横着排列,即一共有1024行,1024*N列,而我用cuFFTplanmany这个API对每一列(一列1024个数据)做FFT,做完之后就开始求每一个数据的平方和,这时我把行数1024对应Threadx,列数N中1024对应BlockIdx.x,N对应BlockIdx.y,这样来索引每一个数据
回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-9-11 10:01:56 | 显示全部楼层
屠戮人神 发表于 2018-9-10 15:05
Lib and Lab:
代码可能存在诸多方面的问题:

我又优化了一下程序,访问数据采用合并内存访问,之后平方采用访问shared memory方式,现在用时少了很多,但是我还不太清楚到底能到多少,就是这个标准该怎们计算,下面是我的程序
__global__ void abs_Kernel(cuComplex *abs_trf, size_t pitch)
{
        int tid = threadIdx.x;//映射为一行里面的1024 这样才可以提高内存访问效率
        int bidx = blockIdx.x;//映射为1024*circle里面的circle,即一行有circle个1024数据
        int bidy = blockIdx.y;//映射为列里面的1024即一列1024个数据,或者是一共有1024行
        __shared__ float trfx;
        __shared__ float trfy;
        trfx = abs_trf[bidy* pitch + bidx * gridDim.y + tid].x;
        trfy = abs_trf[bidy*pitch + bidx * gridDim.y + tid].y;
        abs_trf[bidy*pitch + bidx * gridDim.y + tid].x = trfx*trfx + trfy*trfy;
}
目前我的问题主要有四个:
1)为什么和函数调用的时候abs_Kernel << <(10,1024,1), (1024,1,1) >> >(d_trf, pitch);会不对,而abs_Kernel << <(10,1024,1),1024>> >(d_trf, pitch);就计算对了,我的显卡每个Block最大线程数是1024
2)我的程序还有什么可优化空间吗?目前对1024*10240的复矩阵每一列做cuFFT的时间为0.02,之后求平方和耗时为0.01,我目前能想到的是把相乘换成求平方的API,不知道可以不可以
3)对于cuFFT,一次性做和分开循环做即每次做1024个数据的cuFFT循环做10240次,耗时时间差别大吗?
4)有没有一个标准,来比较我这个处理时间目前的程度是怎样的?根据我目前的显卡的型号怎么能估算出一个标准
回复 支持 反对

使用道具 举报

您需要登录后才可以回帖 登录 | 立即注册

本版积分规则

关闭

站长推荐上一条 /1 下一条

快速回复 返回顶部 返回列表