未经优化的核函数:
__global__ void simpleMultiply(float* a,float* c,int M){
int row=blockIdx.y*blockDim.y+threadIdx.y;
int col=blockIdx.x*blockDim.x+threadIdx.x;
float sum=0.0f;
for(int i=0;i<TILE_DIM;i++){
sum+=a[row*TILE_DIM+i]*a[col*TILE_DIM+i];
}
c[row*M+col]=sum;
}
优化后的核函数:
__global__ void coalescedMultiply(float* a,float* c,int M){
__shared__ float aTile[TILE_DIM][TILE_DIM],
transposedTile[TILE_DIM][TILE_DIM];
int row=blockIdx.y*blockDim.y+threadIdx.y;
int col=blockIdx.x*blockDim.x+threadIdx.x;
float sum=0.0f;
aTile[threadIdx.y][threadIdx.x]=a[row*TILE_DIM+threadIdx.x];
transposedTile[threadIdx.x][threadIdx.y]=
a[(blockIdx.x*blockDim.x+threadIdx.y)*TILE_DIM+
threadIdx.x];
__syncthreads();
for(int i=0;i<TILE_DIM;i++){
sum+=aTile[threadIdx.y][i]*transposeTile[i][threadIdx.x];
}
c[row*M+col]=sum;}