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

作者: 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)有没有一个标准，来比较我这个处理时间目前的程度是怎样的？根据我目前的显卡的型号怎么能估算出一个标准