在任何计算能力的设备上的性能可以通过读取共享内存中的A片来提高,如使用共享内存来提高矩阵乘法中的全局内存负载效率。
使用共享内存来提高矩阵乘法中的全局内存负载效率
__global__ void coalescedMultiply(float *a, float* b, float *c,
int N)
{
__shared__ float aTile[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];
for (int i = 0; i < TILE_DIM; i++) {
sum += aTile[threadIdx.y][i] * b[i*N + col];
}
c[row*N + col] = sum;
}
在使用共享内存来提高矩阵乘法中的全局内存负载效率时,A的磁贴中的每个元素仅从全局内存读取一次,并以完全合并的方式(没有浪费的带宽)读取到共享内存。在for循环的每次迭代中,共享内存中的值将广播给warp中的所有线程。在将A的图块读入共享内存之后,不需要__syncthreads()
同步屏障调用,因为只有在warp内将数据写入共享内存的线程读取数据(注意:代替__syncthreads()
,__shared__
数组可能需要在计算能力2.0或更高的设备上被标记为易失性以确保正确性;请参阅NVIDIA Fermi兼容性指南)。该内核在NVIDIA Tesla K20X上的有效带宽为7.8GB / s。这说明当硬件L1高速缓存驱逐策略与应用程序的需求不匹配或L1高速缓存未用于从全局内存读取时,共享内存用作用户管理的高速缓存。
可以进一步改进如何使用共享存储器来提高矩阵乘法中的全局存储器负载效率处理矩阵B.在计算矩阵C的瓦片的每一行时,读取B的整个瓦片。重复读取B瓦片可以通过将其读入共享内存一次(通过读取共享内存中的附加数据来改进)来消除。
通过将附加数据读入共享内存来改进: