矩阵乘法问题

wts_net 2009-05-30 03:04:03
在深入浅出的例子中矩阵乘法的中的“进一步改良”中的块跟线程是怎么划分的,看得很糊涂,不知道谁还有更详细的解释没有?如果明白的也可以给我解释一下,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;
}
}

...全文
90 4 打赏 收藏 转发到动态 举报
AI 作业
写回复
用AI写文章
4 条回复
切换为时间正序
请发表友善的回复…
发表回复
  • 打赏
  • 举报
回复
1)mxn的二维数组a[x][y]的一维下标为x*n+y
2)j和bidc是分块的开始地址,tidr和tidc是相对与分块点开始的偏移.
3)是把b中分块(以<j,bidc>开始的)中的<tidr,tidc>点赋给matB的[tidr][tidc].
wts_net 2009-05-30
  • 打赏
  • 举报
回复
[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[]的下标不明白是怎么回事?能给解释一下吗?
  • 打赏
  • 举报
回复
就对应关系来说,最外面的j循环可以不用考虑,只是让每个线程多做一些工作,如果没有这个j循环,让block数多n/BLOCK_SIZE倍也是一样的.

这样,对应关系清楚些了吧?A和B及C都分割成BLOCK_SIZExBLOCK_SIZE个块,然后,用一个block计算C=AxB的一个分块。
  • 打赏
  • 举报
回复
[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];

589

社区成员

发帖
与我相关
我的任务
社区描述
CUDA™是一种由NVIDIA推出的通用并行计算架构,该架构使GPU能够解决复杂的计算问题。 它包含了CUDA指令集架构(ISA)以及GPU内部的并行计算引擎。
社区管理员
  • CUDA编程社区
加入社区
  • 近7日
  • 近30日
  • 至今
社区公告
暂无公告

试试用AI创作助手写篇文章吧