标题: 求复矩阵的每一个元素的模的平方
作者: LibAndLab 于 九月 07, 2018, 12:58:41 pm
假如有一个复矩阵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 << > >(d_trf, pitch);
标题:
作者: 屠戮人神 于 九月 10, 2018, 03:05:24 pm
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来的人都会弄错横向和纵向。
感谢来访。
标题:
作者: LibAndLab 于 九月 10, 2018, 07:22:35 pm
Lib and Lab:
代码可能存在诸多方面的问题:
在我将abs_Kernel << > >(d_trf, pitch);这个改为
abs_Kernel << > >(d_trf, pitch);之后结果就对了,不知道这是什么原因呢。还有就是我这里单独开一个核函数的原因是之前运行了cuFFT才得到的这个结果,所以只能分开写。我要那个样子访问其实是因为数据量起始会编程1024*1024*N,通俗的说我想把数据按照1024*1024这样的多个矩阵横着排在一起,所以我按照这个思维方式来做的,Pitch其实是我cudaMallocPitch得到的Pitch,如果我想要让这个核函数加速的话我应该怎么做呢,目前我测时间的话这个核函数需要大量的时间。
标题:
作者: LibAndLab 于 九月 10, 2018, 08:06:53 pm
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,这样来索引每一个数据
标题:
作者: LibAndLab 于 九月 11, 2018, 10:01:56 am
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 << > >(d_trf, pitch);会不对,而abs_Kernel << > >(d_trf, pitch);就计算对了,我的显卡每个Block最大线程数是1024
2)我的程序还有什么可优化空间吗?目前对1024*10240的复矩阵每一列做cuFFT的时间为0.02,之后求平方和耗时为0.01,我目前能想到的是把相乘换成求平方的API,不知道可以不可以
3)对于cuFFT,一次性做和分开循环做即每次做1024个数据的cuFFT循环做10240次,耗时时间差别大吗?
4)有没有一个标准,来比较我这个处理时间目前的程度是怎样的?根据我目前的显卡的型号怎么能估算出一个标准