James的成长之路

CUDA 生态与框架：从 cuBLAS 到多 GPU 训练 Q1：cuBLAS、cuDNN、NCCL 分别是什么？各自用途？ CUDA 生态全景 NVIDIA 构建了一个以 CUDA Runtime 为基础的完整软件栈： 123456789应用层（PyTorch / TensorFlow / JAX） ↓框架加速库（cuDNN / cuBLAS / cuSPARSE / cuFFT） ↓通信库（NCCL / NVSHMEM） ↓CUDA Runtime & Driver ↓GPU 硬件（SM / HBM / NVLink / NVSwitch） cuBLAS：GPU 上的线性代数库 cuBLAS（CUDA Basic Linear Algebra Subroutines）是 NVIDIA 官方实现的 BLAS 标准接口，专为 GPU 加速的密集矩阵运算而设计。 核心能力： Level 操作类型 典型函数 说明 BLAS-1 向量操作 cublasSaxpy, cublasSnrm2 向量加法、范数 ...

CUDA 工具与调试：从 Profiling 到 CUDA Graph Q1：如何 Profile 一个 CUDA 程序？常用工具有哪些？ 为什么要 Profile？ 在 CUDA 程序优化中，"猜测瓶颈"几乎总是错的。先 Profile、后优化是正确顺序。Profile 能告诉你： 每个 Kernel 各耗时多少 内存带宽利用率是否饱和 SM 的 Occupancy（占用率）是否合理 CPU 与 GPU 之间是否在流水线重叠上出了问题 L1/L2 缓存命中率如何 主要工具一览 工具 用途 状态 Nsight Systems 系统级时间线（CPU+GPU 全局视角） 现役，推荐 Nsight Compute Kernel 级性能分析（指标深挖） 现役，推荐 nvprof 命令行 Profiler（旧版） 已废弃（CUDA 11.x 后不再维护） Visual Profiler nvprof 的 GUI 版 已废弃 cuda-memcheck 内存错误检测 已被 compute-sanitizer 取代 compute-sanitizer ...

CUDA 并行算法：从 Reduce 到矩阵转置 Q1：如何在 GPU 上实现 Reduce（规约）操作？有哪些优化方法？ 什么是 Reduce？ Reduce（规约）是并行计算中最基础的操作之一：将 N 个元素通过某种二元运算（求和、求最大值、求最小值等）合并为一个标量结果。CPU 上的串行实现是 O(N)；GPU 上的并行实现可以达到 O(log N) 步，但需要精心设计。 朴素实现与问题 最直观的想法是让每个线程将相邻元素两两相加，迭代 log₂(blockDim.x) 轮： 123456789101112131415161718__global__ void reduce_naive(float* input, float* output, int n) { extern __shared__ float sdata[]; int tid = threadIdx.x; int gid = blockIdx.x * blockDim.x + threadIdx.x; sdata[tid] = (gid < n) ? input[gid] : ...

CUDA 性能优化 Q&A Q1：优化一个 CUDA Kernel 可以从哪些角度入手？ 性能优化没有万能药，但有清晰的分析框架。一个 kernel 的性能瓶颈通常来自以下三个方向之一，找到瓶颈再针对性优化，效果远好于盲目调参。 优化层次全景 1234567891011┌─────────────────────────────────────────────────┐│ 算法层 选择计算复杂度更低的算法（最高回报） │├─────────────────────────────────────────────────┤│ 内存层 减少全局内存访问、提升合并、利用共享内存 │├─────────────────────────────────────────────────┤│ 并行层 提高 Occupancy、消除 Warp Divergence │├─────────────────────────────────────────────────┤│ 指令层 Loop Unroll、向量化加载、减少低效指令 │ ...

CUDA 执行模型与调度 Q&A Q1：什么是 Warp Divergence（线程束分歧）？如何减少其影响？ 什么是 Warp Divergence GPU 以 Warp（32 个线程）为单位同时发射指令。当 Warp 内的线程由于条件分支（if/else、switch、while）而走向不同的代码路径时，就发生了 Warp Divergence（线程束分歧）。 发生分歧时，GPU 必须串行执行每条分支路径——先执行走某条路径的线程（其余线程被屏蔽），再执行走另一条路径的线程。所有路径都执行完毕后，Warp 才重新汇合（reconverge）继续往下执行。 极端情况： 若 Warp 内 32 个线程走了 32 条完全不同的路径，性能退化至理论峰值的 1/32。 分歧仅在 Warp 内部发生，不同 Warp 之间完全独立。 两个 Warp 各自走不同分支是完全并行的，没有任何性能损失。 CC 7.0 前后的差异 在 CC 7.0（Volta）之前，Warp 内所有 32 个线程共享一个程序计数器，必须严格锁步（lockstep）执行。 CC 7.0 起引入了独立线程调 ...

CUDA 内存模型 Q&A Q1：CUDA 中有哪几种内存类型？各自的特点、带宽、延迟是什么？ GPU 的存储层次和 CPU 类似，都是"越快越小、越慢越大"，但 GPU 拥有更丰富的片上可编程内存，这是高性能 CUDA 编程的核心所在。 CUDA 设备上的内存类型汇总如下： 内存类型 物理位置 访问范围 生命周期 速度 大小 寄存器（Register） SM 片上 单线程私有 Kernel 执行期间 最快 极小（每线程约 255 个 32 位寄存器） 共享内存（Shared Memory） SM 片上 线程块内所有线程 Kernel 执行期间 极快（低延迟） 每 SM 最大约 100~228 KB（因架构而异） L1 Cache SM 片上（与共享内存共用物理空间） 自动缓存全局内存访问 自动管理 极快 与共享内存共享配额 L2 Cache 设备片上，全局共享 所有 SM 共享 自动管理 较快 数 MB 至数十 MB 全局内存（Global Memory） 设备 DRAM（显存） Grid 内所有线程 应用程序生命周期 高带宽但高延迟 数 ...

CUDA 基础概念 Q&A Q1：CPU 与 GPU 的核心区别是什么？各自适合哪类任务？ CPU 和 GPU 的设计哲学从根本上就不同——CPU 追求的是把单条线程跑得尽可能快，GPU 追求的是同时把海量线程都跑起来。 CPU 擅长执行一条接一条的串行操作，可以并行运行的线程数通常只有几十个。GPU 则相反，它愿意牺牲单线程的执行速度，换取同时运行数万乃至数百万个线程的能力。这种设计差异直接体现在芯片的晶体管分配上：CPU 把大量晶体管用在数据缓存和流控制逻辑上，以减少单线程的延迟；GPU 则把更多晶体管堆在计算单元上，优先保证整体吞吐量。 正因如此，GPU 在相近的价格和功耗下，能提供远高于 CPU 的指令吞吐量和内存带宽。 核心对比： 维度 CPU GPU 设计目标 单线程串行执行速度 大规模并行吞吐量 并发线程数 数十个 数万至数百万 晶体管用途 缓存 + 流控制 计算单元（ALU） 单线程延迟 极低 较高 内存带宽 相对较低 极高 各自适合的任务： CPU 擅长：逻辑复杂、分支密集、线程间依赖强的串行任务。比如操作系统调度 ...

环境：华为刀片服务器，操作系统 CentOS Linux 7.9 (Core) x86_64，内网无外网访问，只能通过 U 盘等方式离线安装软件。 这次的目标很简单，但过程不算短： 在 完全离线的 CentOS 7.9 上安装 ipmitool； 用 ipmitool 查出华为刀片的 BMC（管理口）IP； 把原本 DHCP 获取的管理口 IP 固定下来（保持不变，但改为静态配置）； 顺带解决 “跳板机 ping 不到管理口”、“ipmitool timeout” 等一系列小坑。 下面按时间线把过程梳理一下。 一、确认系统版本和架构：别一上来就装错包 在离线环境下安装 RPM 包，第一步必须搞清楚系统版本和架构，否则 100% 会装错。 在服务器上执行： 12cat /etc/centos-releaseuname -m 得到结果： 12CentOS Linux release 7.9.2009 (Core)x86_64 这两个信息非常关键： 说明这是 CentOS 7 系列，而不是 Stream 9； 架构是 x86_64，所以必须找 ...x86_64.rpm 的包。 ...

概述 本文深入分析 Mini-SGLang 的模型加载架构，对比传统方式（如 HuggingFace Transformers）与 Mini-SGLang 的设计差异，理解结构与权重分离的设计哲学及其带来的优势。 核心发现：Mini-SGLang 不是通过解析配置文件动态构建模型，而是先手动定义网络结构，再加载权重。这种设计使得框架具有极高的可定制性和优化空间。 1. 支持的模型 1.1 当前支持 Mini-SGLang 目前支持 2 个模型系列： 1. Llama 系列 Llama-2 (7B, 13B, 70B) Llama-3 (8B, 70B) Llama-3.1 (8B, 70B, 405B) Llama-3.2 其他 Llama 架构变体（Mistral、Yi 等） 2. Qwen3 系列 Qwen3-0.6B Qwen3-14B Qwen3-32B 1.2 模型判断逻辑 1234567891011# models/__init__.pydef create_model(model_path: str, model_config: ModelConfig) ...

概述 阶段八学习 Mini-SGLang 的高级特性，主要是自定义 CUDA Kernels。这些 Kernels 通过 C++/CUDA 实现，提供比 PyTorch 更高的性能。 核心问题：为什么需要自定义 CUDA Kernels？ 答案： 性能优化：PyTorch 的通用实现有额外开销 内存效率：直接操作 GPU 内存，避免中间拷贝 并行优化：针对特定场景优化并行策略 内存对齐：利用 GPU 缓存行对齐提高带宽 1. Radix 树操作：kernel/radix.py 1.1 核心思想 问题：在 Radix Tree 中，需要频繁比较两个 token 序列是否相等。Python 的比较很慢。 解决方案：用 C++ 实现快速比较，通过 TVM FFI 调用。 1.2 关键代码 1234567@lru_cache(maxsize=None)def _load_radix_module() -> Module: return load_aot("radix", cpp_files=["radix.cpp"])def fast_compare_key(x: t ...

学习文件：distributed/info.py, distributed/impl.py, layers/linear.py, message/backend.py 1. 为什么需要分布式推理？ 挑战1：模型太大 大模型参数量巨大（例如 Llama-70B 有 700 亿参数） 单个 GPU 显存不够（例如 A100 只有 80GB） 需要多 GPU 并行 挑战2：推理速度慢 单 GPU 计算速度有限 需要并行计算提高吞吐量 解决方案：Tensor Parallelism（TP） 核心思想： 将模型的张量（权重、激活）切分到多个 GPU 每个 GPU 计算一部分 通过通信合并结果 类比： 就像一个大任务分给多个工人 每个工人负责一部分 最后汇总结果 2. DistributedInfo：分布式信息 2.1 核心概念 12345678910@dataclass(frozen=True)class DistributedInfo: rank: int # 当前进程的编号（0, 1, 2, ...） size: int # 总进程数（例如 4 ...

学习文件：kvcache/base.py, kvcache/mha_pool.py, kvcache/naive_manager.py, kvcache/radix_manager.py 1. 为什么需要 KV Cache 管理？ 在阶段一我们学习了 KV Cache 的基本概念：缓存已计算的 Key 和 Value，避免重复计算。 但在实际系统中，KV Cache 管理面临更多挑战： 挑战1：内存有限 GPU 内存有限（例如 24GB） 每个请求的 KV Cache 可能很大（长文本、长对话） 需要驱逐策略来释放空间 挑战2：前缀复用 多个请求可能有相同的前缀（例如 System Prompt） 重复存储浪费内存 需要前缀共享机制 挑战3：并发安全 多个请求同时访问缓存 驱逐时不能删除正在使用的缓存 需要锁机制 2. KV Cache 系统架构 Mini-SGLang 的 KV Cache 系统分为两层： 123456789┌─────────────────────────────────────────┐│ BaseCacheManage ...

学习文件：attention/base.py (66 行), attention/fa.py (187 行), attention/fi.py (279 行) 1. Attention Backend 的核心职责 Attention Backend 是整个推理系统的计算核心，负责执行 Attention 计算： 1.1 职责定位 输入：Q (Query)、K (Key)、V (Value) 三个矩阵 输出：Attention 输出（加权求和后的结果） 核心：实现高效的 Attention 计算 1.2 为什么需要抽象基类？ Mini-SGLang 支持多种 Attention 实现： FlashAttention：适合 Prefill 阶段（计算密集） FlashInfer：适合 Decode 阶段（访存密集） 但 Engine 只需要调用 attn_backend.forward()，不关心具体实现。 123456789# engine.pyclass Engine: def __init__(self, config): self.attn_ba ...

学习文件：engine/engine.py (209 行), engine/graph.py (145 行), engine/sample.py (76 行), models/base.py (15 行), models/llama.py (86 行) 1. Engine 的核心职责 Engine 是整个推理系统的数据平面，负责实际执行 GPU 计算： 1.1 职责定位 Scheduler：控制平面，决定"做什么"（调度哪些请求） Engine：数据平面，执行"怎么做"（前向传播 + 采样） 1.2 核心功能 模型加载和权重管理 KV Cache 内存分配 前向传播执行 Token 采样 CUDA Graph 优化 跨 Stream 异步同步 2. Engine 初始化流程 2.1 初始化的 6 个步骤 12345678910111213141516171819202122232425262728def __init__(self, config: EngineConfig): # 1. 设置设备和通信 self.device = torch.devi ...

学习目标 深入理解 Mini-SGLang 的调度系统,包括: Scheduler 主调度器的核心逻辑 PrefillManager 的 Chunked Prefill 实现 DecodeManager 的请求管理 重叠调度 (Overlap Scheduling) 的原理 1. 调度系统整体架构 1.1 核心组件 调度系统由 5 个核心组件组成: 123456Scheduler (主调度器) ├─→ PrefillManager (Prefill 阶段调度) ├─→ DecodeManager (Decode 阶段调度) ├─→ CacheManager (缓存管理) ├─→ TableManager (页表管理) └─→ Engine (GPU 计算引擎) 1.2 数据流概览 123456789101112131415用户请求 (UserMsg) ↓PrefillManager.pending_list (待处理队列) ↓Scheduler._schedule_next_batch() (调度批次) ↓Scheduler._prepare_batch( ...

本文深入分析 Mini-SGLang 的完整推理流程，探讨多进程架构、ZMQ 消息传递、流式反分词和重叠调度等核心技术。 环境：Mini-SGLang v0.1.0 | Python 3.10+ | PyTorch 2.0+ 源码位置：python/minisgl/server/, python/minisgl/scheduler/, python/minisgl/tokenizer/ 1. 背景：推理系统的架构挑战 LLM 推理系统需要平衡三个关键目标： 低延迟：用户体验要求快速响应（首 token 延迟 < 100ms） 高吞吐：服务端需要同时处理大量请求（> 1000 QPS） 资源利用：GPU 利用率需要保持在 80% 以上 Mini-SGLang 通过多进程架构 + 异步消息传递实现了这些目标。 2. 系统架构：3 进程模型 2.1 架构全景 1234567┌─────────────┐ ZMQ ┌─────────────┐ ZMQ ┌─────────────┐│ Frontend │ ────→ │ Tokenizer ...