#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
问题解决了,谁想知道联系我:)
我想问题是__syncthreads()不能放在if里面,猜对否?
#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;
}
}
能不能修改成kernel函数和block和thread的数目无关的算法呢?
应该不行吧,必须要考虑线程越界的问题啊。
如果数组正好是分块数的整数倍那就简单了
但是必须得考虑一般的情况
必须要考虑线程越界的问题