rocblas自身的实现为架构通用的tensile 机器码，很少会顾及isa架构相关的优化如内存读入，寄存器分配，block size大小等。在对rocblas_sgemm_strided_batch 和自己写的naive版本的batch进行profiling和extractkernel后，着重发现和了解了几个重难点。首先是块与线程部分。第一最高线程速率，硬件生成Threads的速率将直接影响最终程序的效率， 例如GPU显存的读写速度，测试发现gfx906获得64 threads/Cycles的极限性能；第二是1D形状的线程速率曲线。测试得到仅仅当 BlockDim = 256， 512, 1024时， 线程产生速度达到峰值。也即是如果能够将原来4 threads的工作合并到一个thread，每个线程处理的事务随之提高到4倍，例如读写操作，将极大地提高理论极限。在测试了2D 和 3D 之后得出以 256倍的BlockDim性能最佳。其次，在Instruction cahceline的部分，结果表明增加越多的线程能更有效的增加SIMD的效率。在对dump后的VGPR寄存器分析后，发现：HIPCC大多数时候不使用s_load_dword，C ++ Statemetn不能将内联汇编的输出用作操作数，C ++的操作数必须来自C ++，64位地址或数据在内联汇编中很难调用，内联汇编是一个很难用C ++变量，if-esle，for-loop等控制的程序。HIPCC分配SGPR / VGPR的情况不太好GCN限制：s_load_DWORDx4，x8，X16 SGPR必须以x4地址开头内联汇编很难具有正确的VGPR / SGPR设置。由此引出了我们的解决方案：每个 WorkGroup 的 Macro-Tile和Micro-Tile 的分配问题，也即是VLP SGEMM 大小为128的WorkGroup的Macro使用 M=64，N=128，256 的Macro则为 M=64，N=256，每个线程为的Micro tile 大小为M=64， N=1，即每个线程运算Matrix A= 64xK, Matrix B = Kx64, 结果在 Matrix-C 64 x1。对于64个线程，每个M的Matrix-C地址是连续的。每个Wave 的 Matrix A 的Basic Offset 被设定为 N/64 *64 *lda;而B为M/64 *64 lda，取数据的指令为s_load_dwordx8 s[32:39], s[12:13], s18。GCN架构总共有96个可用的SGPR 这个算法使用s32到95，只有64个读取A，而88的并行读入设计使得效率提升。对B来说，每个线程使用微区块大小M = 64，N = 1。每个线程需要8个VGPR来加载1个N的8xK数据。该算法使用global_load_dwordx4来获得最佳的缓存行命中率。下一条存储器读取指令读取同一高速缓存行的下4个DWORD。关于VGPR分配，每个线程需要V [2：3]作为矩阵B的每个线程偏移量。矩阵B的双缓冲区加载需要16倍VGPR。这样总共83个剩下的VGPR负责每个SIMD 3个Waves 得到了很好的性能表现。还有，这种先由先分配变量至寄存器再反编译到机器码（相当于inline 静态库）的方式使得完全没有调度器带来的barrier 和LDS（LDS访存慢于L1 和VGPR）。最终，完成这些操作能使gdx906最高达到77%的性能释放。

doc

编译运行

make
./lib/sgemm_strided_batch_final -m 512 -n 512 -k 256 --batch_count 10

生成

make compile_co

Creativity

深入研读gpu，也即是gcn架构的体系结构相关知识。用汇编和反编译代码的方式优化。主要为优化代码，尤其是gpu代码提供一种优化思路，即先编译分配好VGPR的inline函数和其他一些工具到机器码，再反编译到.co文件，被需要的cpp文件当作外部库来使用，可以极大地利用体系机构地优势从而加速sgemm。

Applications

SgemmBatchedStrided 的应用领域非常多。但是我认为最能体现本答卷价值的是CNN的Convolution，即用体系架构的优化代码方式优化现有CNN代码，用profile 和dump工具对现有的CNN Convolution 汇编分析，比如可以看的点主要有一/二级缓存命中延时、缓存行长度，接下来就可以用简单的汇编代码inlin 再反汇编进行优化。