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

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...