openmlsys/openmlsys-cuda

Tutorials for writing high-performance GPU operators in AI frameworks.

openmlsys-cuda's Issues

README Compilation 部分

这一行是不是得改一下
git submodule init && git submodule sync

TO:
git submodule init && git submodule update

Openmlsys 8.4 章节中 gemm_use_tile.cu 的代码问题

这一部分有关三个 Layout 的代码我一直没有看明白，在这篇知乎中看到一样的内容后，我发现这里的代码实现可能有 bug。
按照定义:

• LayoutTile 是每个 Block 有 (LayoutTile::m, LayoutTile::n) 个 float
• LayoutBlock 是每个 Block 有 （LayoutBlock::m, LayoutBlock::n) 个 thread
• LayoutThread 是每个 thread 中，每个 submatrix 有 (LayoutThread::m, LayoutThread::n) 个 float， 因为用的 float4，可以理解为 4*4。

那么此处 gemm_use_tile.cu 第10行和第11行 中对于m 和 n 的定义就有问题了，应该如下：

unsigned m= threadIdx.x* LayoutTile::m/LayoutBlock::m+ LayoutTile::m* blockIdx.x;
unsigned n= threadIdx.y* LayoutTile::n/LayoutBlock::n+ LayoutTile::n* blockIdx.y

同样的， gemm_use_tile.cu 第19行和第20行 中，iterationA 和 iterationB 应该分别指的是每个 thread 有多少个 (4,4) 的 subMatrix，这里应该是 2*2 = 4 个，那么 gemm_use_tile.cu 第21行和第22行 intervalA 和 intervalB 的定义就有问题了，按照后续代码，intervalA 和 intervalB 指的分别应该是每个 subMatrix 有多大，也就是 (LayoutThread::m, LayoutThread::n)

硬件加速器实践

在第7节给予了硬件加速器的若干介绍，在此基础上加入硬件加速器的实践，这个Repo将作为实践的代码库。

可能的一种形式是以一节出现，将7.3后半部分吸收到这一节，这一节将会成为7.4硬件加速器实践
本节将会以广义矩阵乘法为例，通过若干次的性能优化最终取得接近硬件加速器性能峰值的实现（部分场景超越或打平cuBlas），同时介绍若干性能优化的关键技术。

选择广义矩阵乘法的原因是在深度学习中全连接网络的重要组件就是广义矩阵乘法，事实上，卷积操作也往往是通过im2col等方法将其转化为广义矩阵乘法；此外实现一个高性能的广义矩阵乘法算子相比其他算子（如矩阵转置）对开发者的编程能力和底层硬件架构要求更高。

结构将以如下形式组织：
7.4.1 用最朴素的方法实现一个GPU核函数，相比CPU取得若干倍的加速
7.4.2 使用宽指令并提升计算密度(x5)
7.4.3 进一步提升计算密度(x1.3)
7.4.4 使用一级缓存对复用数据缓存(x10)
7.4.5 优化一级缓存的数据存储流程(x1.1)
7.4.6 优化流水线隐藏一级缓存读取延迟(x1.1)
7.4.7 优化流水线隐藏全局内存读取延迟(x1.1)

每个优化阶段将配合代码与性能分析工具解释策略及取得加速的原因。