矩阵相乘

加速计算 CUDA
1

#include <cutil_inline.h>
#define BLOCK_SIZE 16
#define DATA_SIZE 20
global void Kernel(float Ad, float Bd, floatCd)
{
int bx=blockIdx.x;
int by=blockIdx.y;
int tx=threadIdx.x;
int ty=threadIdx.y;
if(bxBLOCK_SIZE+tx<DATA_SIZE && byBLOCK_SIZE+ty<DATA_SIZE)//只有符合要求的线程才能执行
{
int aBegin=DATA_SIZEBLOCK_SIZEby;
int aEnd=aBegin+DATA_SIZE-1;
int bBegin=BLOCK_SIZEbx;
float sub=0;
for(int a=aBegin,b=bBegin;a<=aEnd;a+=BLOCK_SIZE,b+=BLOCK_SIZEDATA_SIZE)
{
shared float A[BLOCK_SIZE][BLOCK_SIZE];
shared float B[BLOCK_SIZE][BLOCK_SIZE];
A[tx][ty]=Ad[a+DATA_SIZEty+tx];
B[tx][ty]=Bd[b+DATA_SIZE*ty+tx];
__syncthreads();
for(int k=0;k<BLOCK_SIZE;k++)
{
if(a-aBegin+k<DATA_SIZE)
sub+=A[k][ty]B[tx][k];
}
__syncthreads();
}
int c=DATA_SIZEBLOCK_SIZEby+BLOCK_SIZEbx;
Cd[c+DATA_SIZEty+tx]=sub;
}
}
高性能运算之CUDA中的例子，改造了下
主要是解决矩阵的维数不是块中线程个数的整数倍的情况
但还是有问题
在不足块中线程个数的哪个块结果是错的
我的矩阵里面都是1，两个矩阵都是2020
这是一行的结果，，最后四个是错的
20.000000 20.000000 20.000000 20.000000 20.000000 20.000000 20.000000 20.000000 20.000000 20.000000 20.000000 20.000000 20.000000 20.000000 20.000000 20.000000 -0.001327 -0.001327 -0.001327 -0.001327

2

问题解决了，谁想知道联系我:)

3

我想问题是__syncthreads()不能放在if里面，猜对否？

4

#include <cutil_inline.h>
#define BLOCK_SIZE 16
#define DATA_SIZE 512
global void Kernel(float Ad, float Bd, floatCd)
{
int bx=blockIdx.x;
int by=blockIdx.y;
int tx=threadIdx.x;
int ty=threadIdx.y;
int aBegin=DATA_SIZEBLOCK_SIZEby;
int aEnd=aBegin+DATA_SIZE-1;
int bBegin=BLOCK_SIZEbx;
float sub=0;
for(int a=aBegin,b=bBegin;a<=aEnd;a+=BLOCK_SIZE,b+=BLOCK_SIZEDATA_SIZE)
{
shared float A[BLOCK_SIZE][BLOCK_SIZE];
shared float B[BLOCK_SIZE][BLOCK_SIZE];
if(a+tx<aBegin+DATA_SIZE && byBLOCK_SIZE+ty<DATA_SIZE) //对A,B的限制是不同的
{
A[tx][ty]=Ad[a+DATA_SIZE*ty+tx];
}
else
{
A[tx][ty]=0;

	}
	if( b+ty*DATA_SIZE<bBegin+DATA_SIZE*DATA_SIZE && bx*BLOCK_SIZE+tx<DATA_SIZE )
	{
		B[tx][ty]=Bd[b+DATA_SIZE*ty+tx];
	}
	else
	{
		B[tx][ty]=0;

	}
	__syncthreads();

	for(int k=0;k<BLOCK_SIZE;k++)
	{
		sub+=A[k][ty]*B[tx][k];
	}
	__syncthreads();

}
if(bx*BLOCK_SIZE+tx<DATA_SIZE && by*BLOCK_SIZE+ty<DATA_SIZE)
{
	int c=DATA_SIZE*BLOCK_SIZE*by+BLOCK_SIZE*bx;
	Cd[c+DATA_SIZE*ty+tx]=sub;  

}

}

5

能不能修改成kernel函数和block和thread的数目无关的算法呢？

6

应该不行吧，必须要考虑线程越界的问题啊。
如果数组正好是分块数的整数倍那就简单了
但是必须得考虑一般的情况

7

必须要考虑线程越界的问题