tpoisonooo / how-to-optimize-gemm Goto Github PK

cuda代码中，MMult_cuda_7.cu中30行，b_ptr += 64 * k，应该是b_ptr += 64 * n，因为是方阵，所以结果对上了

https://github.com/tpoisonooo/how-to-optimize-gemm/blob/master/cuda/MMult_cuda_12.cu: 20,21
I'm a beginner of CUDA&&PTX, I want to know what does these two PTX use for?
"{.reg .pred p;\n"
"mov.b32 %0, 0;\n"
is it useless code?

Is that for performance reason or just comparison?

I have another question about MMult_cuda_12.cu
Honestly, I don't know how to overlap the share2register and computing process. Is it the asm(PTX) that make them run parallelly? The instructions are sequantially, so how could these two parts of code hide each other?
part1: loading shared-memory to panel
lds128(panelA[pp][0], panelA[pp][1], panelA[pp][2], panelA[pp][3],
aptr_base + ((subk + 1) % 8) * SMEM_LDA * sizeof(float));
lds128(panelA[pp][4], panelA[pp][5], panelA[pp][6], panelA[pp][7],
aptr_base + (((subk + 1) % 8) * SMEM_LDA + 64) * sizeof(float));
lds128(panelB[pp][0], panelB[pp][1], panelB[pp][2], panelB[pp][3],
bptr_base + ((subk + 1) % 8) * SMEM_LDB * sizeof(float));
lds128(panelB[pp][4], panelB[pp][5], panelB[pp][6], panelB[pp][7],
bptr_base + (((subk + 1) % 8) * SMEM_LDB + 64) * sizeof(float));

part2: computing the result of panel-data
#pragma unroll
for (int i = 0; i < 8; ++i) {
#pragma unroll
for (int j = 0; j < 8; ++j) {
sum[i][j] += panelA[subk % 2][i] * panelB[subk % 2][j];
}
}

Hi! I am learning your wonderful code but get puzzled by some assembly code. Excuse me that I am a new assembly learner. I searched for some documents and still find these two points strange:
v11, line 20: why we define a p but not used it?
line 29, why we use long type?

Thank you!!!

T4卡上运行时出现报错 CUDA error at test_MMult.cpp:110 code=999 "cudaEventSynchronize(stop)",想问一下是什么原因呢

gflops benchmark中的 #define OP_FLOATS (80), 这里的80是怎么计算的呢？

如kernel_v3中:
float *begin_a = a + by * BLOCK * k; //by->n
float *begin_b = b + bx * BLOCK; //bx->m

当A,B不为方阵时会出错,例如m=k=256,n=128.

A(i, j) = (2.0 * (float)drand48() - 1.0)
误差为
4.577637e-05

如果数字更大，误差更大

我现在跑的芯片型号是NVIDIA，ARMv8 Processor rev 0 (v8l)。
我看知乎文章里说测试浮点峰值时FMA指令的排布数量 = FMA的发射数 * FMA指令的延迟。我并没有查到上面这个芯片的手册。但是我看了A57的手册，里面是这样记录的：

FMA指令的延迟是10，吞吐量是2。我不太清楚这个吞吐是否代表着芯片可以同时发射两条FMA指令（是芯片发射吗），但是我分别放置了10条FMA指令（OP_FLOATS = 80）和20条FMA指令（OP_FLOATS = 160）都测试了，发现在10条的时候是16.095492 GFLOPS， 20条是 18.759214 GFLOPS。这是什么原因呢？
我的猜测有两个：
1.10条FMA指令确实不是测试这款芯片的浮点峰值所需要的指令数。
2.可能编译器自动开启了多线程？这个比较有可能，因为从4条指令到10条指令性能差不多翻倍，但是10-20只增加了一点。

I found the definition of MMult0.c and MMult1.c for multi-dimensional array storage, would you like to take a look and see whether both of them are correct?
For MMult0.c

For MMult1.c

As you see, the positions of i and j are inverted.

Hi! I think the a_sts_addr and a_base are the same? So maybe we can delete one?

Thank you!!!