• CUDA

CUDA

因为GPU不支持常规的Kernel递归，CPU上的很多递归算法只能换思路后进行改写，不能直接按原思路实现。而随着动态并行（Dynamic Parallelism）的引入，GPU现在能直接在Kernel中启动Kernel了。

https://zhuanlan.zhihu.com/p/674856090

CUDA动态并行详解（CDP2）

在“blocked”排列中，每个线程拥有一组连续的数据项；在“striped”排列中，所有线程拥有的数据项交错存储。

Local Thread Block：当前正在执行的线程所在的线程块。

Remote Thread Block：线程块簇中除本地线程块之外的其他线程块。

cuda_dlink: CUDA Dynamic Linking

早期的GPU硬件上只有一个execution engine，因此，不论是哪个进程、哪个线程发起的kernel launch，都在同一个队列里排队。

随着GPU的发展，GPU上面开始出现了多个execution engine。

一个stream就对应于一个执行队列（加一个执行单元），用户可以自行决定是否把两个kernel分开放在两个队列里。

https://zhuanlan.zhihu.com/p/699754357

一文读懂cuda stream与cuda event

原始地址：

https://github.com/ROCm/Tensile/wiki/Kernel-Parameters

Spill指的是当GPU的寄存器压力（Register Pressure）过高时，编译器无法将所有活跃变量分配到有限的物理寄存器中，被迫将部分数据临时存储到内存（通常是Local Memory或L1/L2缓存）中的现象。

pytorch CUDA RadixSort call stack：

MediumRadixSort
should_use_small_sort
sortKeyValueInplace
launch_stable_sort_kernel
segmented_sort_large_segments
radix_sort_pairs_impl
NO_ROCM(at_cuda_detail)::cub::DeviceRadixSort::SortPairs
cub::DeviceRadixSort::SortPairs
DeviceRadixSort::custom_radix_sort
DispatchRadixSort::Dispatch
DeviceRadixSortSingleTileKernel
triple_chevron
BlockRadixSort
BlockRadixSortT(temp_storage.sort).SortBlockedToStriped
RankKeys
DescendingBlockRadixRank
BlockRadixRank

float4的使用示例参见：

torch/aten/src/ATen/native/cuda/MemoryAccess.cuh

https://developer.nvidia.com/blog/even-easier-introduction-cuda/

An Even Easier Introduction to CUDA

http://ishare.iask.sina.com.cn/f/17211495.html

深入浅出谈CUDA技术

http://blog.csdn.net/xsc_c/article/category/2186063

某人的并行计算专栏

https://mp.weixin.qq.com/s/9D7uda3CV7volenhl-jchg

推荐几个不错的CUDA入门教程

https://mp.weixin.qq.com/s/bvNnzkOzGYYYewc3G9DOIw

GPU是如何优化运行机器学习算法的？

https://mp.weixin.qq.com/s/nAwxtOUi6HpIjVOREgEfaA

CUDA编程入门极简教程

https://mp.weixin.qq.com/s/-zdIWkuRZXhsLJmOZljOBw

《基于GPU-多核-集群等并行化编程》

https://mp.weixin.qq.com/s/bCb5VsH58JII886lpg9lvg

如何在CUDA中为Transformer编写一个PyTorch自定义层

https://mp.weixin.qq.com/s/OYSzol-vufiKPuU9YxtbuA

矩阵相乘在GPU上的终极优化：深度解析Maxas汇编器工作原理

https://zhuanlan.zhihu.com/p/358220419

PyTorch自定义CUDA算子教程与运行时间分析

https://zhuanlan.zhihu.com/p/358778742

详解PyTorch编译并调用自定义CUDA算子的三种方式

https://zhuanlan.zhihu.com/p/360441891

熬了几个通宵，我写了份CUDA新手入门代码

https://mp.weixin.qq.com/s/EZxO8IIBDJ4c7eQhUffc2w

怎样节省2/3的GPU？爱奇艺vGPU的探索与实践

https://mp.weixin.qq.com/s/3VjGpyXZSkJhy6sFPUsZzw

GPU虚拟化，算力隔离，和qGPU

https://zhuanlan.zhihu.com/p/383115932

大佬是怎么优雅实现矩阵乘法的？

https://zhuanlan.zhihu.com/p/410278370

CUDA矩阵乘法终极优化指南

https://www.zhihu.com/column/c_1437330196193640448

深入浅出GPU优化

https://www.zhihu.com/question/41060378

自己写的CUDA矩阵乘法能优化到多快？

https://zhuanlan.zhihu.com/p/559957579

简单谈谈CUDA的访存合并

https://zhuanlan.zhihu.com/p/565897763

GPGPU编程模型之CUDA

http://blog.csdn.net/augusdi/article/details/12833235

这是一篇转帖的CUDA教程，原帖比较分散，不好看。

https://zhuanlan.zhihu.com/p/544864997

cuda中threadIdx、blockIdx、blockDim和gridDim的使用

https://zhuanlan.zhihu.com/p/690717002

一文读懂cuda代码编译流程

https://zhuanlan.zhihu.com/p/690880124

并不太短的CUDA入门（The Not So Short Introduction to CUDA）

https://zhuanlan.zhihu.com/p/693690123

一文读懂nvidia-smi背后的nvml库

https://www.zhihu.com/question/445590537

问个CUDA并行上的小白问题，既然SM只能同时处理一个WARP，那是不是有的SP处于闲置？

https://zhuanlan.zhihu.com/p/2003147245525153066

现代CUDA编程指南

https://zhuanlan.zhihu.com/p/1898003094056470472

GPU GEMM优化之Stream-K