矩阵乘法问题

[wts_net]

在深入浅出的例子中矩阵乘法的中的“进一步改良”中的块跟线程是怎么划分的，看得很糊涂，不知道谁还有更详细的解释没有？如果明白的也可以给我解释一下，kernel部分的过程完全不明白！
即：Kernel 程式的部份，則改成：

__global__ static void matMultCUDA(const float* a, size_t lda,
const float* b, size_t ldb, float* c, size_t ldc, int n)
{
__shared__ float matA[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float matB[BLOCK_SIZE][BLOCK_SIZE];
const int tidc = threadIdx.x;
const int tidr = threadIdx.y;
const int bidc = blockIdx.x * BLOCK_SIZE;
const int bidr = blockIdx.y * BLOCK_SIZE;
int i, j;
float results = 0;
float comp = 0;
for(j = 0; j < n; j += BLOCK_SIZE) {
if(tidr + bidr < n && tidc + j < n) {
matA[tidr][tidc] = a[(tidr + bidr) * lda + tidc + j];
}
else {
matA[tidr][tidc] = 0;
}
if(tidr + j < n && tidc + bidc < n) {
matB[tidr][tidc] = b[(tidr + j) * ldb + tidc + bidc];
}
else {
matB[tidr][tidc] = 0;
}
__syncthreads();
for(i = 0; i < BLOCK_SIZE; i++) {
float t;
comp -= matA[tidr][i] * matB[i][tidc];
t = results - comp;
comp = (t - results) + comp;
results = t;
}
__syncthreads();
}
if(tidr + bidr < n && tidc + bidc < n) {
c[(tidr + bidr) * ldc + tidc + bidc] = results;
}
}

[无心人_过过小日子]

[Quote=引用楼主 wts_net 的帖子:]
__syncthreads();
for(i = 0; i < BLOCK_SIZE; i++) {
float t;
comp -= matA[tidr][i] * matB[i][tidc];
t = results - comp;
comp = (t - results) + comp;
results = t;
}
__syncthreads();
[/Quote]

这一段是Kahan's Summation Formula算法,提高累加精度的.等价于:
result += matA[tidr][i] * matB[i][tidc];

[无心人_过过小日子]

就对应关系来说,最外面的j循环可以不用考虑,只是让每个线程多做一些工作,如果没有这个j循环,让block数多n/BLOCK_SIZE倍也是一样的.

这样,对应关系清楚些了吧?A和B及C都分割成BLOCK_SIZExBLOCK_SIZE个块，然后，用一个block计算C=AxB的一个分块。

[wts_net]

[Quote=引用 1 楼 l7331014 的回复:]
引用楼主 wts_net 的帖子:
__syncthreads();
for(i = 0; i < BLOCK_SIZE; i++) {
float t;
comp -= matA[tidr][i] * matB[i][tidc];
t = results - comp;
comp = (t - results) + comp;
results = t;
}
__syncthreads();


这一段是Kahan's Summation Formula算法,提高累加精度的.等价于:
…
[/Quote]

if(tidr + j < n && tidc + bidc < n) {
matB[tidr][tidc] = b[(tidr + j) * ldb + tidc + bidc];
}
这里b[]的下标不明白是怎么回事？能给解释一下吗？

[无心人_过过小日子]

1)mxn的二维数组a[x][y]的一维下标为x*n+y
2)j和bidc是分块的开始地址,tidr和tidc是相对与分块点开始的偏移.
3)是把b中分块(以<j,bidc>开始的)中的<tidr,tidc>点赋给matB的[tidr][tidc].