如何打破memory bandwidth limited

加速计算 CUDA
41

还有一个问题，刚刚想起来，我的矩阵实际上用int类型就可以了，但是cublasDgemm()函数没有int类型的参数，那我是不是要把之前的矩阵改为int类型呢？改为int类型对计算速度有影响吗？

42

浮点型比较快， 建议不要换。

43

我把之前的kernel修改如下，是不是比较符合合并访问了呢？

__global__ void KernelV(float *dev_TranXMatrix,double *dev_V)         //GPU端运行，初始化主效方差
{
	 int k=blockIdx.x*512+threadIdx.x;
	 int i=blockIdx.y;
	
	if((i<XROW)&&(k<XROW))
	{
		for(int j=0;j<XCOL;j++)
		{
			dev_V[i*XROW+k]=dev_V[i*XROW+k]+dev_TranXMatrix[j*XROW+i]*dev_TranXMatrix[j*XROW+k]*1e-30;	
		}
	}
}

还有一段kernel，我不太明白，反正是用for循环，像dev_VI[i+k*XROW]这种是不是就无所谓了啊？

__global__ void KernelVOne(double *dev_V,double dev_VK,int *dev_XMatrix,int j,double s0)
{
	 int k=blockIdx.x*512+threadIdx.x;
	 int i=blockIdx.y;
	
	if((i<XROW)&&(k<XROW))
		{
			dev_V[i*XROW+k]=dev_V[i*XROW+k]+dev_XMatrix[i*XCOL+j]*dev_XMatrix[k*XCOL+j]*(dev_VK-s0);
	    }
}
44

简单地，你只要考虑一个warp内相邻的threads访存是不是连续的，或者在一小段连续的范围内。

dev_V[i*XROW+k]=dev_V[i*XROW+k]+dev_TranXMatrix[j*XROW+i]dev_TranXMatrix[jXROW+k]*1e-30;

左：合并
右：1）合并；2）广播，而且可能受到内部循环j的影响，使得局部性变差；3）同一步循环内是连续的，但是和2）一样可能受到j的影响。

dev_V[i*XROW+k]=dev_V[i*XROW+k]+dev_XMatrix[i*XCOL+j]dev_XMatrix[kXCOL+j]*(dev_VK-s0);
左：合并
右：1）合并；2）广播；3）完全不合并。

祝您编码愉快~

45

右三为什么会完全不合并？我特意定义了一个原矩阵的转置矩阵dev_TranXMatrix[jXROW+k]，原来的矩阵是dev_XMatrix[j+kXCOL]，我知道之前这个是不合并的，我认为是k是一列一列的读，两个k之间总是差一行的距离，但是现在您说完全不合并我就不理解了，k不是连续的行元素吗？解决办法呢？用shared memory吗？

46

int k=blockIdx.x*512+threadIdx.x

dev_XMatrix[k*XCOL+j]

考虑一个warp里面的相邻两个threads，不妨假定是第一个block里面的前两个threads，也就是blockIdx.x==0，threadIdx.x==0,1
如此，对于这两个线程，k分别为0,1。
此时，需要访问dev_XMatrix[k*XCOL+j]的元素分别为dev_XMatrix[j]和
dev_XMatrix[XCOL+j]

考虑到XCOL一般是一个比较大的值，这会使得一个warp内部每个线程申请的都是相距非常远的数据，每次读取只能为一个线程服务，实际由于读取粒度的问题，会浪费大量的带宽。
即使XCOL是一个比较小的值，也会带来若干损失。

至于解决方法，您可以考虑自行修改访存操作的实现看看。

祝您编码愉快。

47

可能我们说的不是一个kernel,dev_XMatrix[kXCOL+j]这个不合并，我知道，我以为你说的是dev_TranXMatrix[jXROW+K]不合并呢

48

我是按照您43#提供的代码说的，无论是44#还是46#均是如此。
如果您隐藏了自己的最新科研成果，那我也无法看到，无能为力~

祝您编码愉快~

49

我那么烂的水平，怎能称科研成果！今天在电梯里又遇到老板了，他又问我加速状况了，最近确实在努力弄，但是进展缓慢，昨天就一个矩阵转置的库倒腾了一个晚上，本来打算用cublas，但是结果一直不对，后来，索性用cula的culaSgeTranspose（）函数，结果，发现把数据从.txt中读取，和直接分配空间然后赋值调用的row和col大小是不一样的，弄了半天才把这个问题搞定~

50

我在48#的意思是，你写好了合并访问的版本，却拿一个不合并的代码贴出来…

当然也无妨的，您改好了就行。

祝您编码愉快~

51

嗯，有位专家跟我讲让我用Texture memory来处理广播的问题，我正在看一些资料，遇到新的问题，再向各位请教，O(∩_∩)O谢谢~

52

好的，欢迎您继续深入地研究此问题，祝您编码愉快~