这个kernel可以这样理解吗?

global static void matMultCUDA(const float* a, size_t lda,
const float* b, size_t ldb, float* c, size_t ldc, int n)
{ extern shared float data;
const int tid = threadIdx.x;
const int row = blockIdx.x;
int i, j;
for(i = tid; i < n; i += blockDim.x) { //这里的blockDim.x是多少呢??是1吗?
data[i] = a[row * lda + i];
}
__syncthreads();
for(j = tid; j < n; j += blockDim.x) {
float t = 0;
float y = 0;
for(i = 0; i < n; i++) {
float r;
y -= data[i] * b[i * ldb + j];
r = t - y;
y = (r - t) + y;
t = r; }
c[row * ldc + j] = t;
}
}
程序是《深入浅出谈CUDA》中的第18页,计算2个矩阵相乘的。假设有n个block,n个thread。我的理解是block和thread均为一维的,在第一个for循环中,blockDim.x为1。
for(i = tid; i < n; i += blockDim.x) {
data[i] = a[row * lda + i];
}
比如第0个block中的thread分别0,1,2,3…n,第0个thread循环n次,将矩阵a中的第一行已经赋值给data了,第1个thread循环n-1次,将矩阵a中第一行中的a【0】【1】往后的值赋给了data,这样是不是有重复的读取呢?是这样理解吗?谢谢啊,交流一下[/font]

[ 本帖最后由 caoyuxin541 于 2010-5-3 16:16 编辑 ]