抱歉，您的浏览器无法访问本站

Cutlass通俗理解核心目的 核心目标: 消除矩阵乘法中对A, B矩阵的重复读取 核心: 内积变外积 -> 每个数据只会读第一次, 避免对A/B矩阵的重复读取 最内层外积, 可复用的地方也用外积, 体现就是MMA的warp level op 内积的问题: 如果每个thread负责C矩阵的一个元素C[m, n], 并且用下面这种方式循环(内积), 可以发现A, B矩阵会被重复...

shfl: warp-level primitives一个warp有32个thread, warp内的线程称为通道(lanes), lane id的计算方法是threadid % 32, warp id的计算方法是threadid / 32。 线程束洗牌: warp-level原语 可以直接获取warp内的线程的寄存器值，直接使用寄存器交换 每次调用都会同步warp内的线程, sync m...

用cutlass cute实现flash attentionflash attention自顶向下(虽然我学cutlass是自底向上学的但是感觉快速上手应该自顶向下学)。因为有了cutlass cute用户就可以方便的实现一些功能了, 即一些cuda编程的范式: cuda程序范式: global mem -> share mem -> reg -> compute blo...

CUDA编程小记这个教程的小记/速查 核函数(kernel)kernel的定义和启动形如，使用__global__修饰的函数就是kernel，由host端启动，在device端运行。如果使用__host__就是就是host端代码，cuda编译器不会编译到gpu上。使用__device__修饰就是设备端代码，只能在设备上运行 1234__global__ void cuda_kernel(.....

通过例子理解CUDA driver api和CUDA PTXcuda程序经过编译器编译后后添加很多对cuda driver api的调用, 这样用户就不用关心gpu module管理, context管理, kernel管理等的细节了。 不过我们就是想要知道细节, 所以这篇文章通过直接使用cuda driver api加载cuda ptx执行的方式体会其中的细节。 我们的目标是将下面这个cu...