Skip to content

📚Modern CUDA Learn Notes with PyTorch: Tensor/CUDA Cores, 📖150+ CUDA Kernels, 📖HGEMM (achieve the performance of cuBLAS 🎉🎉).

License

Notifications You must be signed in to change notification settings

DefTruth/CUDA-Learn-Notes

Folders and files

NameName
Last commit message
Last commit date

Latest commit

 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 

Repository files navigation

📚 Modern CUDA Learn Notes with PyTorch for Beginners: It includes Tensor/CUDA Cores, TF32/F16/BF16/F8, 📖150+ CUDA Kernels🔥🔥 with PyTorch bindings, 📖30+ LLM/VLM🔥, 📖40+ CV/C++...🔥, 📖50+ CUDA/CuTe...🔥 Blogs and 📖toy-hgemm library🔥🔥 which can achieve the performance of cuBLAS, check 📖HGEMM Supported Matrix👇 for more details. Welcome to 🌟👆🏻star this repo to support me, many thanks ~ 🎉🎉

Currently, on NVIDIA L20, RTX 4090 and RTX 3090 Laptop, compared with cuBLAS's default Tensor Cores math algorithm CUBLAS_GEMM_DEFAULT_TENSOR_OP, the HGEMM (WMMA/MMA) implemented in this repo (blue🔵) can achieve 95%~99% of its (orange🟠) performance. Please check toy-hgemm library🔥🔥 for more details.

CUDA Cores Sliced K(Loop over K) Tile Block Tile Thread
✔️ ✔️ ✔️ ✔️
WMMA(m16n16k16) MMA(m16n8k16) Pack LDST(128 bits) SMEM Padding
✔️ ✔️ ✔️ ✔️
Copy Async Tile MMA(More Threads) Tile Warp(More Values) Multi Stages
✔️ ✔️ ✔️ ✔️
Reg Double Buffers Block Swizzle Warp Swizzle Collective Store(Warp Shfl)
✔️ ✔️ ✔️ ✔️
Row Major(NN) Col Major(TN) SGEMM TF32 SMEM Swizzle(CuTe)
✔️ ✔️ ✔️ ✔️

📖 150+ CUDA Kernels 🔥🔥 (面试常考题目) (©️back👆🏻)

Workflow: custom CUDA kernel impl -> PyTorch Python bindings -> Run tests. 👉TIPS: * = Tensor Cores(WMMA/MMA), otherwise, CUDA Cores; / = not supported; ✔️ = supported; = in my plan.

📖 cuda kernel 📖 elem dtype 📖 acc dtype 📖 docs 📖 level
✔️ nsys/ncu(timeline/ptx/sass) / / link ⭐️
✔️ elementwise_f32 f32 / link ⭐️
✔️ elementwise_f32x4 f32 / link ⭐️
✔️ elementwise_f16 f16 / link ⭐️
✔️ elementwise_f16x2 f16 / link ⭐️
✔️ elementwise_f16x8 f16 / link ⭐️
✔️ elementwise_f16x8_pack f16 / link ⭐️⭐️
✔️ histogram_i32 i32 / link ⭐️
✔️ histogram_i32x4 i32 / link ⭐️
✔️ sigmoid_f32 f32 / link ⭐️
✔️ sigmoid_f32x4 f32 / link ⭐️
✔️ sigmoid_f16 16 / link ⭐️
✔️ sigmoid_f16x2 f16 / link ⭐️
✔️ sigmoid_f16x8 f16 / link ⭐️
✔️ sigmoid_f16x8_pack f16 / link ⭐️⭐️
✔️ relu_f32 f32 / link ⭐️
✔️ relu_f32x4 f32 / link ⭐️
✔️ relu_f16 f16 / link ⭐️
✔️ relu_f16x2 f16 / link ⭐️
✔️ relu_f16x8 f16 / link ⭐️
✔️ relu_f16x8_pack f16 / link ⭐️⭐️
✔️ gelu_f32 f32 / link ⭐️
✔️ gelu_f32x4 f32 / link ⭐️
✔️ gelu_f16 f16 / link ⭐️
✔️ gelu_f16x2 f16 / link ⭐️
✔️ gelu_f16x8 f16 / link ⭐️
✔️ gelu_f16x8_pack f16 / link ⭐️⭐️
✔️ swish_f32 f32 / link ⭐️
✔️ swish_f32x4 f32 / link ⭐️
✔️ swish_f16 f16 / link ⭐️
✔️ swish_f16x2 f16 / link ⭐️
✔️ swish_f16x8 f16 / link ⭐️
✔️ swish_f16x8_pack f16 / link ⭐️⭐️
✔️ embedding_f32 f32 / link ⭐️
✔️ embedding_f32x4 f32 / link ⭐️
✔️ embedding_f32x4_pack f32 / link ⭐️
✔️ embedding_f16 f16 / link ⭐️
✔️ embedding_f16x2 f16 / link ⭐️
✔️ embedding_f16x8 f16 / link ⭐️
✔️ embedding_f16x8_pack f16 / link ⭐️⭐️
✔️ mat_trans_f32_col2row{2d} f32 / link ⭐️
✔️ mat_trans_f32_row2col{2d} f32 / link ⭐️
✔️ mat_trans_f32_diagonal2d f32 / link ⭐️⭐️
✔️ mat_trans_f32x4_col2row{2d} f32 / link ⭐️⭐️
✔️ mat_trans_f32x4_row2col{2d} f32 / link ⭐️⭐️
✔️ warp_reduce_[all] all all link ⭐️⭐️
✔️ reduce_f32_f32 f32 f32 link ⭐️⭐️
✔️ reduce_f32x4_f32 f32 f32 link ⭐️⭐️
✔️ reduce_f16_f16 f16 f16 link ⭐️⭐️
✔️ reduce_f16_f32 f16 f32 link ⭐️⭐️
✔️ reduce_f16x2_f16 f16 f16 link ⭐️⭐️
✔️ reduce_f16x2_f32 f16 f32 link ⭐️⭐️
✔️ reduce_f16x8_pack_f16 f16 f16 link ⭐️⭐️
✔️ reduce_f16x8_pack_f32 f16 f32 link ⭐️⭐️
✔️ reduce_bf16_bf16 bf16 bf16 link ⭐️⭐️
✔️ reduce_bf16_f32 bf16 f32 link ⭐️⭐️
✔️ reduce_bf16x2_bf16 bf16 bf16 link ⭐️⭐️
✔️ reduce_bf16x2_f32 bf16 f32 link ⭐️⭐️
✔️ reduce_bf16x8_pack_bf16 bf16 bf16 link ⭐️⭐️
✔️ reduce_bf16x8_pack_f32 bf16 f32 link ⭐️⭐️
✔️ reduce_fp8_e4m3_f16 fp8_e4m3 f16 link ⭐️⭐️
✔️ reduce_fp8_e5m2_f16 fp8_e5m2 f16 link ⭐️⭐️
✔️ reduce_fp8_e4m3x16_pack_f16 fp8_e4m3 f16 link ⭐️⭐️
✔️ reduce_fp8_e5m2x16_pack_f16 fp8_e5m2 f16 link ⭐️⭐️
✔️ reduce_i8_i32 i8 i32 link ⭐️⭐️
✔️ reduce_i8x16_pack_i32 i8 i32 link ⭐️⭐️
✔️ dot_product_f32 f32 f32 link ⭐️⭐️
✔️ dot_product_f32x4 f32 f32 link ⭐️⭐️
✔️ dot_product_f16_f32 f16 f32 link ⭐️⭐️
✔️ dot_product_f16x2_f32 f16 f32 link ⭐️⭐️
✔️ dot_product_f16x8_pack_f32 f16 f32 link ⭐️⭐️
✔️ softmax_f32(fence) f32 f32 link ⭐️⭐️
✔️ softmax_f32x4(fence) f32 f32 link ⭐️⭐️
✔️ softmax_f32 f32 f32 link ⭐️⭐️
✔️ softmax_f32x4 f32 f32 link ⭐️⭐️
✔️ safe_softmax_f32 f32 f32 link ⭐️⭐️
✔️ safe_softmax_f32x4 f32 f32 link ⭐️⭐️
✔️ safe_softmax_f16_f32 f16 f32 link ⭐️⭐️
✔️ safe_softmax_f16x2_f32 f16 f32 link ⭐️⭐️
✔️ safe_softmax_f16x8_pack_f32 f16 f32 link ⭐️⭐️
✔️ online_safe_softmax_f32 f32 f32 link ⭐️⭐️
✔️ online_safe_softmax_f32x4_pack f32 f32 link ⭐️⭐️
✔️ rope_f32 f32 f32 link ⭐️⭐️
✔️ rope_f32x4_pack f32 f32 link ⭐️⭐️
✔️ layer_norm_f32 f32 f32 link ⭐️⭐️
✔️ layer_norm_f32x4 f32 f32 link ⭐️⭐️
✔️ layer_norm_f16_f16 f16 f16 link ⭐️⭐️
✔️ layer_norm_f16x2_f16 f16 f16 link ⭐️⭐️
✔️ layer_norm_f16x8_f16 f16 f16 link ⭐️⭐️
✔️ layer_norm_f16x8_pack_f16 f16 f16 link ⭐️⭐️
✔️ layer_norm_f16x8_pack_f32 f16 f32 link ⭐️⭐️
✔️ layer_norm_f16_f32 f16 f32 link ⭐️⭐️
✔️ rms_norm_f32 f32 f32 link ⭐️⭐️
✔️ rms_norm_f32x4 f32 f32 link ⭐️⭐️
✔️ rms_norm_f16_f16 f16 f16 link ⭐️⭐️
✔️ rms_norm_f16x2_f16 f16 f16 link ⭐️⭐️
✔️ rms_norm_f16x8_f16 f16 f16 link ⭐️⭐️
✔️ rms_norm_f16x8_f32 f16 f32 link ⭐️⭐️
✔️ rms_norm_f16x8_pack_f16 f16 f16 link ⭐️⭐️
✔️ rms_norm_f16x8_pack_f32 f16 f32 link ⭐️⭐️
✔️ rms_norm_f16_f32 f16 f32 link ⭐️⭐️
✔️ sgemm_naive_f32 f32 f32 link ⭐️⭐️
✔️ sgemm_sliced_k_f32 f32 f32 link ⭐️⭐️⭐️
✔️ sgemm_t_8x8_sliced_k_f32x4 f32 f32 link ⭐️⭐️⭐️
✔️ sgemm_t_8x8_sliced_k...bcf f32 f32 link ⭐️⭐️⭐️
✔️ sgemm_t_8x8_sliced_k...dbuf f32 f32 link ⭐️⭐️⭐️
✔️ sgemm_t_8x8_sliced_k16...dbuf f32 f32 link ⭐️⭐️⭐️
✔️ sgemm_t_8x8_sliced_k16...async f32 f32 link ⭐️⭐️⭐️
✔️ sgemm_wmma_m16n16k8...stages* tf32 f32 link ⭐️⭐️⭐️
✔️ sgemm_wmma_m16n16k8...swizzle* tf32 f32 link ⭐️⭐️⭐️
✔️ hgemm_naive_f16 f16 f16 link ⭐️⭐️
✔️ hgemm_sliced_k_f16 f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_t_8x8_sliced_k_f16x4 f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_t_8x8_sliced_k_f16x4_pack f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_t_8x8_sliced_k_f16x8_pack f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_t_8x8_sliced_k...dbuf f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_t_8/16x8...k16/32...dbuf f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_t_8/16x8...k16/32...async f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_wmma_m16n16k16...naive* f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_wmma_m16n16k16...mma4x2* f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_wmma_m16n16k16...mma4x4* f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_wmma_m16n16k16...dbuf* f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_wmma_m32n8k16....dbuf* f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_wmma_m16n16k16...stages* f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_wmma_m16n16k16...swizzle* f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_mma_m16n8k16...naive* f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_mma_m16n8k16...mma2x4* f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_mma_m16n8k16...stages* f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_mma_m16n8k16...swizzle* f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_mma_stages{swizzle}...cute* f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_mma_cublas* f16 f16 link ⭐️⭐️
✔️ sgemv_k32_f32 f32 f32 link ⭐️⭐️⭐️
✔️ sgemv_k128_f32x4 f32 f32 link ⭐️⭐️⭐️
✔️ sgemv_k16_f32 f32 f32 link ⭐️⭐️⭐️
✔️ hgemv_k32_f16 f16 f16 link ⭐️⭐️⭐️
✔️ hgemv_k128_f16x4 f16 f16 link ⭐️⭐️⭐️
✔️ hgemv_k16_f16 f16 f16 link ⭐️⭐️⭐️
✔️ flash_attn_f32 f32 f32 link ⭐️⭐️⭐️
✔️ flash_attn_mma_m16n8k16* f16 f16 link ⭐️⭐️⭐️
✔️ nms_f32 f32 / link ⭐️⭐️
✔️ notes v1(deprecated) f32 f32 / ⭐️

📖 博客目录

📖 大模型|多模态|Diffusion|推理优化 (本人作者) (©️back👆🏻)

📖 类型-标题 📖 作者
[分布式训推][张量/序列并行]📖图解DeepSpeed-Ulysses&Megatron-LM TP/SP @DefTruth
[VLM推理优化][InternVL系列]📖InternLM2/.../InternVL1.5系列笔记: 核心点解析 @DefTruth
[LLM推理优化][TensorRT-LLM][5w字]📖TensorRT-LLM部署调优-指北 @DefTruth
[LLM推理优化][KV Cache优化]📖GQA/YOCO/CLA/MLKV: 层内和层间KV Cache共享 @DefTruth
[LLM推理优化][Prefill优化]📖图解vLLM Prefix Prefill Triton Kernel @DefTruth
[LLM推理优化][Prefill优化][万字]📖图解vLLM Automatic Prefix Caching: TTFT优化 @DefTruth
[LLM推理优化][Attention优化]📖图解:从Online-Softmax到FlashAttention V1/V2/V3 @DefTruth
[LLM推理优化][Decoding优化]📖原理&图解FlashDecoding/FlashDecoding++ @DefTruth
[VLM推理优化][LLaVA系列]📖CLIP/LLaVA/LLaVA1.5/VILA笔记: 核心点解析 @DefTruth
[LLM推理优化][Attention优化][万字]📖TensorRT MHA/Myelin vs FlashAttention-2 @DefTruth
[LLM推理优化][PTX汇编]📖CUDA 12 PTX汇编: PRMT指令详解-通用模式 @DefTruth
[LLM推理优化][PTX汇编]📖CUDA 12 PTX汇编: LOP3指令详解 @DefTruth
[LLM推理优化][CUDA][3w字]📖高频面试题汇总-大模型手撕CUDA @DefTruth
[LLM推理优化][Weight Only]📖WINT8/4-(00): 通俗易懂讲解-快速反量化算法 @DefTruth
[LLM推理优化][Weight Only]📖WINT8/4-(01): PRMT指令详解及FT源码解析 @DefTruth
[LLM推理优化][Weight Only]📖WINT8/4-(02): 快速反量化之INT8转BF16 @DefTruth
[LLM推理优化][Weight Only]📖WINT8/4-(03): LOP3指令详解及INT4转FP16/BF16 @DefTruth
[LLM推理优化][LLM Infra整理]📖100+篇: 大模型推理各方向新发展整理 @DefTruth
[LLM推理优化][LLM Infra整理]📖30+篇: LLM推理论文集-500页PDF @DefTruth
[LLM推理优化][LLM Infra整理]📖FlashDecoding++: 比FlashDecoding还要快! @DefTruth
[LLM推理优化][LLM Infra整理]📖TensorRT-LLM开源,TensorRT 9.1也来了 @DefTruth
[LLM推理优化][LLM Infra整理]📖20+篇: LLM推理论文集-300页PDF @DefTruth
[LLM推理优化][LLM Infra整理]📖PagedAttention论文新鲜出炉 @DefTruth

📖 CV推理部署|C++|算法|技术随笔 (本人作者) (©️back👆🏻)

📖 类型-标题 📖 作者
[推理部署][CV/NLP]📖FastDeploy三行代码搞定150+ CV、NLP模型部署 @DefTruth
[推理部署][CV]📖如何在lite.ai.toolkit(3.6k+ stars)中增加您的模型? @DefTruth
[推理部署][CV]📖美团 YOLOv6 ORT/MNN/TNN/NCNN C++推理部署 @DefTruth
[推理部署][ONNX]📖ONNX推理加速技术文档-杂记 @DefTruth
[推理部署][TensorFlow]📖Mac源码编译TensorFlow C++指北 @DefTruth
[推理部署][CV]📖1Mb!头部姿态估计: FSANet,一个小而美的模型(C++) @DefTruth
[推理部署][CV]📖opencv+ffmpeg编译打包全解指南 @DefTruth
[推理部署][CV]📖RobustVideoMatting视频抠图静态ONNX模型转换 @DefTruth
[推理部署][CV]📖190Kb!SSRNet年龄检测详细解读(含C++工程) @DefTruth
[推理部署][CV]📖MGMatting(CVPR2021)人像抠图C++应用记录 @DefTruth
[推理部署][CV]📖超准确人脸检测(带关键点)YOLO5Face C++工程详细记录 @DefTruth
[推理部署][ORT]📖解决: ONNXRuntime(Python) GPU 部署配置记录 @DefTruth
[推理部署][CV]📖记录SCRFD(CVPR2021)人脸检测C++工程化(含docker镜像) @DefTruth
[推理部署][NCNN]📖野路子:记录一个解决onnx转ncnn时op不支持的trick @DefTruth
[推理部署][CV]📖升级版轻量级NanoDet-Plus MNN/TNN/NCNN/ORT C++工程记录 @DefTruth
[推理部署][CV]📖超轻量级NanoDet MNN/TNN/NCNN/ORT C++工程记录 @DefTruth
[推理部署][CV]📖详细记录MGMatting之MNN、TNN和ORT C++移植 @DefTruth
[推理部署][CV]📖YOLOX NCNN/MNN/TNN/ONNXRuntime C++工程简记 @DefTruth
[推理部署][TNN]📖手动修改YoloX的tnnproto记录-TNN @DefTruth
[推理部署][ORT]📖全网最详细 ONNXRuntime C++/Java/Python 资料! @DefTruth
[推理部署][CV]📖RobustVideoMatting: C++工程化记录-实现篇 @DefTruth
[推理部署][CV]📖RobustVideoMatting: C++工程化记录-应用篇 @DefTruth
[推理部署][ORT]📖ONNXRuntime C++ CMake 工程分析及编译 @DefTruth
[推理部署][ORT]📖如何使用ORT C++ API处理NCHW和NHWC输入? @DefTruth
[推理部署][TNN]📖tnn-convert搭建简记-YOLOP转TNN @DefTruth
[推理部署][CV]📖YOLOP ONNXRuntime C++工程化记录 @DefTruth
[推理部署][NCNN]📖超有用NCNN参考资料整理 @DefTruth
[推理部署][MNN]📖超有用MNN参考资料整理 @DefTruth
[推理部署][TNN]📖超有用TNN参考资料整理 @DefTruth
[推理部署][ONNX]📖超有用ONNX参考资料整理 @DefTruth
[推理部署][ONNX]📖超有用ONNX模型结构参考资料整理 @DefTruth
[推理部署][OpenCV-DNN]📖超有用OpenCV-DNN参考资料整理 @DefTruth
[推理部署][Tensorflow]📖超有用Tensorflow C++工程化知识点 @DefTruth
[推理部署][模型转换]📖深度学习模型转换资料整理 @DefTruth
[技术随笔][C++][CMake]📖超有用CMake参考资料整理 @DefTruth
[技术随笔][C++][3W字]📖静态链接和静态库实践指北-原理篇 @DefTruth
[技术随笔][C++]📖Mac下C++内存检查指北(Valgrind VS Asan) @DefTruth
[技术随笔][CV]📖torchlm: 人脸关键点检测库 @DefTruth
[技术随笔][ML]📖《统计学习方法-李航: 笔记-从原理到实现-基于R》 @DefTruth
[技术随笔][Git]📖如何优雅地git clone和git submodule? @DefTruth
[技术随笔][3D]📖人脸重建3D参考资料整理 @DefTruth
[技术随笔][3D]📖BlendShapes参考资料整理 @DefTruth
[技术随笔][3D]📖从源码安装Pytorch3D详细记录及学习资料 @DefTruth
[技术随笔][ML]📖200页:《统计学习方法:李航》笔记 -从原理到实现 @DefTruth

📖 CUTLASS|CuTe|NCCL|CUDA|文章推荐 (其他作者) (©️back👆🏻)

💡说明: 本小节整理一些自己比较喜欢的文章。欢迎大家提PR推荐更多优秀的文章!

📖 类型-标题 📖 作者
[cute系列详解][入门]📖cutlass cute 101 @朱小霖
[cute系列详解][入门]📖CUTLASS 2.x & CUTLASS 3.x Intro 学习笔记 @BBuf
[cute系列详解][Layout]📖cute 之 Layout @reed
[cute系列详解][Layout]📖cute Layout 的代数和几何解释 @reed
[cute系列详解][Tensor]📖cute 之 Tensor @reed
[cute系列详解][MMA]📖cute 之 MMA抽象 @reed
[cute系列详解][Copy]📖cute 之 Copy抽象 @reed
[cute系列详解][Swizzle]📖cute 之 Swizzle @reed
[cute系列详解][Swizzle]📖cute Swizzle细谈 @进击的Killua
[cute系列详解][Swizzle]📖cutlass swizzle机制解析(一) @Titus
[cute系列详解][Swizzle]📖cutlass swizzle机制解析(二) @Titus
[cute系列详解][GEMM]📖cute 之 简单GEMM实现 @reed
[cute系列详解][GEMM]📖cute 之 GEMM流水线 @reed
[cute系列详解][GEMM]📖cute 之 高效GEMM实现 @reed
[cute系列详解][GEMM]📖GEMM流水线: single/multi-stage、pipeline @Titus
[cute系列详解][GEMM]📖GEMM细节分析(一): ldmatrix的选择 @Anonymous
[cute系列详解][GEMM]📖GEMM细节分析(二): TiledCopy与cp.async @Anonymous
[cute系列详解][GEMM]📖GEMM细节分析(三): Swizzle<B,M,S>参数取值 @Anonymous
[cute系列详解][实践]📖Hopper Mixed GEMM的CUTLASS实现笔记 @BBuf
[cute系列详解][实践]📖CUTLASS CuTe实战(一): 基础 @进击的Killua
[cute系列详解][实践]📖CUTLASS CuTe实战(二): 应用 @进击的Killua
[cute系列详解][实践]📖FlashAttention fp8实现(ada架构) @shengying.wei
[cute系列详解][实践]📖FlashAttention 笔记: tiny-flash-attention解读 @shengying.wei
[cute系列详解][实践]📖使用cutlass cute复现flash attention @66RING
[cutlass教程][入门]📖cutlass 基本认知 @JoeNomad
[cutlass教程][入门]📖cutlass 软件架构 @JoeNomad
[cutlass教程][入门]📖CUTLASS 基础介绍 @进击的Killua
[cutlass教程][入门]📖乱谈CUTLASS GTC2020 SLIDES @zzk again
[cutlass教程][深入]📖cutlass block swizzle 和 tile iterator @JoeNomad
[cutlass教程][深入]📖cutlass bank conflict free的smem layout @JoeNomad
[cutlass教程][深入]📖cutlass 多级流水线 @JoeNomad
[GPU指令集架构][精解]📖NVidia GPU指令集架构-前言 @reed
[GPU指令集架构][精解]📖NVidia GPU指令集架构-寄存器 @reed
[GPU指令集架构][精解]📖NVidia GPU指令集架构-Load和Cache @reed
[GPU指令集架构][精解]📖NVidia GPU指令集架构-浮点运算 @reed
[GPU指令集架构][精解]📖NVidia GPU指令集架构-整数运算 @reed
[GPU指令集架构][精解]📖NVidia GPU指令集架构-比特和逻辑操作 @reed
[CUDA优化][入门]📖CUDA(一):CUDA 编程基础 @紫气东来
[CUDA优化][入门]📖CUDA(二):GPU的内存体系及其优化指南 @紫气东来
[CUDA优化][实践]📖CUDA(三):通用矩阵乘法:从入门到熟练 @紫气东来
[CUDA优化][实践]📖ops(1):LayerNorm 算子的 CUDA 实现与优化 @紫气东来
[CUDA优化][实践]📖ops(2):SoftMax算子的 CUDA 实现 @紫气东来
[CUDA优化][实践]📖ops(3):Cross Entropy 的 CUDA 实现 @紫气东来
[CUDA优化][实践]📖ops(4):AdamW 优化器的 CUDA 实现 @紫气东来
[CUDA优化][实践]📖ops(5):激活函数与残差连接的 CUDA 实现 @紫气东来
[CUDA优化][实践]📖ops(6):embedding 层与 LM head 层的 CUDA 实现 @紫气东来
[CUDA优化][实践]📖ops(7):self-attention 的 CUDA 实现及优化 (上) @紫气东来
[CUDA优化][实践]📖ops(8):self-attention 的 CUDA 实现及优化 (下) @紫气东来
[CUDA优化][实践]📖CUDA(四):使用 CUDA 实现 Transformer 结构 @紫气东来
[CUDA优化][Copy]📖Async Copy及Memory Barrier指令的功能与实现 @Frank Wang
[CUDA优化][GEMV]📖深入浅出GPU优化系列:gemv优化 @有了琦琦的棍子
[Tensor Cores]📖Nvidia Tensor Core初探 @木子知
[Tensor Cores]📖Nvidia Tensor Core-WMMA API编程入门 @木子知
[Tensor Cores]📖Nvidia Tensor Core-MMA PTX编程入门 @木子知
[Tensor Cores]📖CUDA Ampere Tensor Core HGEMM 矩阵乘法优化 @nicholaswilde
[GPU通信架构][精解]📖NVIDIA GPGPU(四)- 通信架构 @Bruce

©️License (©️back👆🏻)

GNU General Public License v3.0

🎉Contribute (©️back👆🏻)

How to contribute? please check 🌤🌤CONTRIBUTE🎉🎉.

📖 References (©️back👆🏻)