Question

深入浅出cuda里第二个程序将矩阵乘法编程小块相乘的不太懂，希望解释下BLOCK_NUM和kernel里的循环什么意思

clock_t matmultCUDA(const float* a, int lda, const float* b, int ldb, float* c, int ldc, int n)
{
float *ac, *bc, *cc;
clock_t start, end;

start = clock();
size_t pitch_a, pitch_b, pitch_c;
cudaMallocPitch((void**) &ac, &pitch_a, sizeof(float) * n, n);
cudaMallocPitch((void**) &bc, &pitch_b, sizeof(float) * n, n);
cudaMallocPitch((void**) &cc, &pitch_c, sizeof(float) * n, n);//自动以最佳的倍数来配置内存//

cudaMemcpy2D(ac, pitch_a, a, sizeof(float) * lda,sizeof(float) * n, n, cudaMemcpyHostToDevice);
cudaMemcpy2D(bc, pitch_b, b, sizeof(float) * ldb,sizeof(float) * n, n, cudaMemcpyHostToDevice);

int bx = (n + BLOCK_SIZE - 1) / BLOCK_SIZE;
dim3 blocks(bx, bx);
dim3 threads(BLOCK_SIZE, BLOCK_SIZE);
matMultCUDA<<<blocks, threads>>>(ac, pitch_a / sizeof(float),bc, pitch_b / sizeof(float), cc, pitch_c / sizeof(float), n);

cudaMemcpy2D(c, sizeof(float) * ldc, cc, pitch_c,sizeof(float) * n, n, cudaMemcpyDeviceToHost);

cudaFree(ac);
cudaFree(bc);
cudaFree(cc);

end = clock();

return end - start;
}
__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;
}
}

Answer 1

BLOCK_NUM是那个变量？

追问

给出的定义为16

追答

那是BLOCK_SIZE，同学。
因为矩阵被分块了，所有要有循环，这个得在纸上画下才看得出来。

追问

哦，提问的时候打错了，想问BLOCK_SIZE具体是做什么用的，因为直接给了个定义16，看不出来具体代表什么意思，我知道这个事矩阵分块，可能由于编程基础有点生疏，有些地方看不懂
for(j = 0; j < n; j += BLOCK_SIZE) {
if(tidr + bidr < n && tidc + j < n) {
matA[tidr][tidc] = a[(tidr + bidr) * lda + tidc + j];
}
这里面的tidr+bidr tidc+j还有下面的这些if里的变量啥意思啊

追答

要仔细阅读编程指南，BLOCK_SIZE是block每维的大小. if的作用保证a[(tidr + bidr) * lda + tidc + j]的下标不越界.
强烈推荐你看编程指南里面的那个矩阵相乘.