Optimizing-SGEMM-on-NVIDIA-Turing-GPUs icon indicating copy to clipboard operation
Optimizing-SGEMM-on-NVIDIA-Turing-GPUs copied to clipboard

Published 20 hours ago • verified publisher icon yzhaiustc

kernel3

Open liuqi123123 opened this issue 3 years ago • 7 comments

您好，我就用中文提问了呀。在kernel3中，你把blocksize从(32,32)改为(1024),这种做法的优点你说有3点好处:1.storing threadIdx.x before re-using it massively 2. in order to reduce living registers 3. benefit the compiler optimization 这几点我都不太懂是啥意思。在书中和网上都找不到对应的解释，能麻烦您能说的详细一些吗? 如果还能给出参考资料那也是最好不过的！

liuqi123123 avatar Dec 12 '21 08:12 liuqi123123

谢谢你对这个开源项目的兴趣. 这里说的是有问题的。CUDA11之前编译器对TID之类的特殊寄存器的编译我记得有些不完善，所以我写的时候就加上了这一步。 我刚刚检查了下 kernel2 和 kernel3的汇编，没看出有啥区别, 所以这一步并不对性能有帮助. 我会抽空更新这个代码，更新完之前不会close这个issue. 感谢你提出宝贵的建议 :)

yzhaiustc avatar Dec 12 '21 18:12 yzhaiustc

嗯嗯好，谢谢你的项目，找工作还得靠它呢！

liuqi123123 avatar Dec 13 '21 01:12 liuqi123123

作者你好，请问cuda程序在运行时，你是用的什么工具测得它的计算速度(FLOPS/s)的呀？

liuqi123123 avatar Dec 16 '21 13:12 liuqi123123

用手算. performance (GFLOPS) = 2 * m * n * k * 1e-9 / elapsed_time 这里运行时间的单位是 秒 (second).

yzhaiustc avatar Dec 16 '21 20:12 yzhaiustc

谢谢指导。准确来说，1GB = 2^30B，那表达式是不是这样：performance (GFLOPS) = 2 * m * n * k /（2^30 * elapsed_time)

liuqi123123 avatar Dec 17 '21 01:12 liuqi123123

不是的. 性能上 1G = 10^9 存储上 1G = 2^30.

yzhaiustc avatar Dec 17 '21 02:12 yzhaiustc

您好，我又来问问题了。在gemm优化中，一个线程最终计算了矩阵C中8乘8个元素，这样做使得计算密度更大了。那如果继续增大处理元素的个数，如16乘16,32乘32，这样做带来的弊端将会是什么呢？谢谢您的解答！

liuqi123123 avatar Dec 21 '21 05:12 liuqi123123