Skip to content

Navigation Menu

Sign in
Appearance settings

Search code, repositories, users, issues, pull requests...

Provide feedback

We read every piece of feedback, and take your input very seriously.

Saved searches

Use saved searches to filter your results more quickly
Sign up
Appearance settings

Dandelion111/CUDA-Learn-Notes

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 📖HGEMM/SGEMM🔥🔥 which has been fully optimized, check 📖HGEMM/SGEMM 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(sky blue🔵) can achieve 95%~99% of its(orange🟠) performance. Please check hgemm benchmark 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(Shfl)
✔️ ✔️ ✔️ ✔️
Row Major(NN) Col Major(TN) SGEMM TF32 SMEM Swizzle(Permuted)
✔️ ✔️ ✔️ ...

📖 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 ⭐️⭐️⭐️
✔️ 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_1_fwd_f32 f32 f32 link ⭐️⭐️⭐️
✔️ flash_attn_2_fwd_f16_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👆🏻)

📖 类型-标题 📖 作者
[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

💡说明: 大佬们写的文章实在是太棒了,学到了很多东西。欢迎大家提PR推荐更多优秀的文章!

©️License (©️back👆🏻)

GNU General Public License v3.0

🎉Contribute (©️back👆🏻)

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

📖 References

References (©️back👆🏻)

About

📚Modern CUDA Learn Notes with PyTorch: Tensor/CUDA Cores, 📖150+ CUDA Kernels with PyTorch bindings, 📖HGEMM/SGEMM (95%~99% cuBLAS performance), 📖100+ LLM/CUDA Blogs.

Resources

License

Stars

Watchers

Forks

Releases

No releases published

Packages

Contributors

Languages

  • Cuda 85.1%
  • Python 14.7%
  • C++ 0.2%

AltStyle によって変換されたページ (->オリジナル) /