- 博客(76)
- 收藏
- 关注
RSS订阅原创 NVIDIA GPU PTX - mbarrier
因此，mbarrier支持Threads在mbarrier arrive之后，可以干一些其他的有效工作，再去wait mbarrier。mbarrier对象：一个内存中的不透明对象，充当一个有状态的同步计数器。mbarrier对象基于shared memory实现，个数是用户定义的，只受到shared memory大小的限制。
2025-12-04 17:16:09
466
原创 NVIDIA GPU PTX - Memory Consistency Model
在multi-threaded执行模型中，每个Thread所执行内存操作的side-effect，会以不完整且非一致的顺序对其他Thread可见。没有Memory Consistency Model的情况下，读取操作会返回对同一内存位置的写入操作所提交的值的集合中的任意一个。假设我们有两个Thread和两个初始值为 0 的共享变量 x 和 y。Memory Consistency Model约束了。
2025-11-25 15:38:12
649
原创 NVIDIA GPU - NVIDIA Nsight
结果会按照–trace里指定的内容（默认值是：cuda,nvtx,osrt,opengl）分表显示。有两种方式可以查看vector_add.ncu-rep，一种是GUI，一种是导出表格。可以看到profiling的结果被写入了：vector_add.nsys-rep。可以看到profiling的结果被写入了：vector_add.ncu-rep。包含GPU 内存事务的统计，分为2个表：(by Time/by Size)Branch Instructions Ratio：分支指令比例。
2025-11-20 16:11:43
882
原创 vLLM - GPUModelRunner
在execute_model中，会调用_calc_spec_decode_metadata计算投机推理的SpecDecodeMetadata。累积的调度token数：cu_num_scheduled_tokens = [4, 104, 107, 207, 209]在上一个Step每个Request已经通过draft model投机推理产生的token数：[3, 0, 2, 0, 1]累积的草稿token数cu_num_draft_tokens: [3, 3, 5, 5, 6]
2025-09-30 11:03:53
1091
原创 vLLM - Worker
Worker实现了WorkerBase的接口，是GPU Worker的具体实现，代码位于：D:\2025\vllm\vllm\vllm\v1\worker\gpu_worker.py。
2025-09-19 16:42:21
807
原创 vLLM - WorkerBase
vLLM定义了WorkerBase作为所有硬件平台Worker的基类，代码位于：D:\2025\vllm\vllm\vllm\v1\worker\worker_base.py。
2025-09-18 10:21:44
440
原创 vLLM - EngineCore
DPEngineCoreActor.run()调用了父类的方法：DPEngineCoreProc.run_busy_loop()提到使用Ray Backend时，会创建CoreEngineActorManager来管理EngineCore。基类，负责执行Inner Loop。
2025-09-17 14:31:18
714
原创 vLLM - EngineCoreClient
在SyncMPClient.__init__中，已经独立创建了一个线程，用于接收并解码output_socket的消息，并放入self.outputs_queue（queue.Queue类型）。所以SyncMPClient.get_output只需要调用self.outputs_queue.get()即可以实现同步IO：在未收到数据时，阻塞主线程。
2025-09-14 21:56:53
1607
原创 vLLM - LLMEngine
LLMEngine是vLLM的核心，它接收输入，通过语言模型生成输出，并将结果返回给客户端。离线批量推理的LLM，在线服务的AsyncLLMEngine都通过封装LLMEngine对外提供推理服务。LLMEngine有v0和v1两个版本的实现，目前主流是使用的v1版本，源代码位于：vllm/vllm/v1/engine/llm_engine.py。
2025-09-14 19:10:06
633
原创 vLLM - SchedulerInterface
SchedulerInterface定义了vLLM Scheduler的接口，源代码：vllm\v1\core\sched\interface.py。
2025-09-12 09:30:24
1215
原创 Graph Compilation - Lowering
本文重点展开讲解Inductor IR和FX Graph转换为Inductor IR的过程。
2025-09-09 14:34:11
699
原创 vLLM - Introduction
vLLM 是一个用于 LLM（大型语言模型）推理和部署的快且易用的库。最初由UC Berkeley的 Sky Computing Lab 开发，当前已发展为一个由社区驱动的项目，得到了学术界和工业界的贡献。硬件支持：支持 NVIDIA GPU、AMD CPU 和 GPU、Intel CPU 和 GPU、PowerPC CPU、TPU 和 AWS Neuron。量化支持：支持 GPTQ、AWQ、AutoRound、INT4、INT8 和 FP8 等。分布式推理：支持TP、PP、DP和EP。
2025-09-08 17:55:44
379
原创 MLIR - Pattern Rewriting
MLIR的一个独特优势是使用一种IR来表达多种抽象层次，一个MLIR Op可以是TensorFlow Op(高层抽象)，也可以是LLVM IR Instruction（底层抽象）。因此基于MLIR要解决的问题也是多样的，一个DAG-DAG的Rewriter框架可以解决很多问题。编译器的IR变换有很多种，Generic DAG Rewriter主要用于解决其中的一种：将一个Op组成的DAG，替换成另一种。在matchAndRewrite过程中，可能会由递归的产生：重写之后的结果又匹配上了Pattern。
2025-09-08 17:21:41
787
原创 Triton Linalg - PointerStrengthReductionPtrPass
在Triton程序中，有很多非常灵活的指针运算(如addptr，broadcast等)，这些运算比较亲和GPGPU这种有大量线程且通过SM调度去掩藏访存时延的微架构。对NPU这种线程数少且掩藏访存时延能力较弱的微架构，需要将Triton程序中的指针Op转换为等价的带偏移量的Op，实现性能优化，这就是PointerStrengthReductionPtrPass的功能。
2025-09-05 16:17:21
733
原创 MLIR - InlinerPass
在 MLIR中，用于内联函数调用。内联是指将函数调用替换为函数体的过程，可以减少函数调用的开销，并且有助于进一步的优化。
2025-09-04 19:45:45
729
原创 Triton Linalg - WrapFuncBodyWithSingleBlockPass
WrapFuncBodyWithSingleBlockPass用于将函数体中有多个Block的Function，用scf.execute_region封起来，让Function变为单Block的形式。示例：WrapFuncBodyWithSingleBlockPass简化了Function的控制流，解决了在scf.for循环中内联函数调用时由于多块（multi-block）问题导致的内联问题。假如inline优化需要把foo做内联，则原始的多Block的foo会导致生成的IR不正确或难以优化。
2025-09-03 11:36:11
913
原创 Triton Linalg编译
triton-linalg是一个将Triton Dialect转换为Linalg Dialect的项目。当前作为寒武纪Triton编译器的前端，几乎支持了所有Triton语言的特性。在Triton转换为Linalg的过程中，遵从了几个原则：尽可能使用结构化（Structured）的Operator，避免使用tt.reduce转换为linalg.reduce。tt.dot转换为linalg.matmul。arith和math中逐元素的Operator转换为对应的尽早识别Operator的语义。
2025-09-01 18:32:15
406
原创 MLIR - Linalg
linalg.generic可以不依赖程序上下文，完全通过Operands推导出迭代空间，即可以独立地生成所需的控制流。输入A是一个memref，layout是<?输入B是一个memref，layout是<?#attrs = {return上述代码可以lowering成如下形式：可以看到scf.for的控制流生成仅依赖linalg.generic自身。returnlinalg.generic显示定义了迭代器的类型：Parallel：并行。用于独立的维度，同时处理多个元素。
2025-08-30 17:55:41
1030
原创 Triton - Launcher
Triton编译在生成Device（如GPU）上的Kernel的同时，也会生成Host上的Kernel调用代码：Launcher。
2025-08-29 09:25:26
602
原创 TorchInductor - Autotune
Triton Kernel支持tl.constexpr类型的参数，这些参数会使用Triton的Autotune机制来寻优。
2025-08-26 15:23:35
894
原创 TorchInductor - Introduction
PyTorch 2.x通过TorchDynamo通过Python Bytecode的动态变换实现了图捕获功能，需要搭配一个Compiler Backend完成图编译。Pytorch尝试集成了多个后端，并使用一个轻量级的autotuner来选择最优的后端图编译结果。因此Pytorch需要原生的Compiler Backend：TorchInductor。
2025-08-21 19:37:23
827
原创 torch.compile
一个Deep Learning Compiler，为多种加速器生成高性能代码。对NVIDIA和AMD GPUs, 使用OpenAI Triton编译器作为Backend。规范化2000+ PyTorch Operators为250+ Primitive Operators, 极大降低了开发Pytorch后端的难度。基于Python Frame Evaluation Hook技术，实现安全的Pytorch的计算图捕获。AOT生成计算图的反向图。
2025-08-15 18:07:35
452
原创 AOT Autograd
AotAutograd（Ahead-of-Time Autograd）是一个用于在模型真正运行前跟踪模型的前向/反向图的模块，在Inductor中，通过aot_autograd调用了AOT Autograd的功能：生成反向传播图：在前向传播图的基础上，提前生成反向传播图。图编译：对生成的前向/反向图进行编译。
2025-08-15 18:05:01
405
原创 TorchDynamo - FX Graph
TorchDynamo在CPython的Frame Evaluation过程中增加了Hook，在Python字节码真正执行前，将Pytorch Operation捕获为FX Graph的形式，再动态修改字节码，实现JIT的编译和优化。TorchDynamo机制类似DynamoRIO的机制，所以取名为TorchDynamo。这个方案拥有如下优势：Python支持完备：可以轻松地fall back到原始的Bytecode执行（eager mode）。
2025-08-13 17:25:46
982
原创 TorchDynamo - Python PEP523
在Python PEP 523中提出了允许用户指定解释器的Frame Evaluation Function的提案，实现了允许外部C代码控制CPython的Frame Evaluation的机制。这个机制使JIT可以根据需要，来自由选择是编译Python字节码为机器码，还是Fall back到原始的字节码解释执行，这成为了。
2025-08-13 12:05:35
505
原创 TorchDynamo - API
PyTorch的torch.compile通过TorchDynamo捕获Python字节码并转换为FX图进行优化。支持两种使用方式：函数调用或装饰器形式。主要参数包括：backend（默认inductor）、fullgraph（是否强制全图捕获）、dynamic（动态形状）等。提供多种优化模式，如reduce-overhead（减少Python开销）和max-autotune（自动调优）。用户可自定义backend查看FX图结构，示例展示了如何捕获加法操作的计算图及输入输出。
2025-08-13 10:27:03
911
原创 Pytorch编译
查看官网说明，python版本可以在3.9和3.12之间任选：创建一个3.12的conda环境：源码下载CUDA版本号本地的CUDA版本号是12.4：查询官网：Previous PyTorch Versions可以看到2.6.0是一个可用的版本号。clone源代码：clone submodule:查看下载的submodule：查看分支：可以找到2个相关的分支：orig/release/2.6：2.6版本的初始分支。release/2.6：在2.6版本的初始分支基础上，合入了一
2025-08-12 15:17:37
500
原创 MLIR Defining Dialects
MLIR提供了一个强有力的声明式规范机制TableGen来维护领域特定的Record，基于Record可以自动生成C++的模板代码。基于TableGen，大大简化了Dialect的定义和维护工作。因此，使用声明式规范机制来定义新的Dialect是官方推荐的做法。
2025-08-08 11:55:03
219
原创 MLIR Language Reference
MLIR的IR形式类似传统的SSA形式，但是又导入了多面体循环优化的一些概念。这种混合IR的设计是为了方便表示，分析，变换高层数据流图，并生成特定后端的高性能代码。MLIR被设计为3种形式：textual form：方便人阅读的。in-memory form：方便程序变换和分析。compact serialized form：方便存储和传输。三种形式的语义内容是相同的，本文以textual form来介绍相关内容。
2025-08-07 11:38:11
880
原创 MLIR Data Layout
Data layout表达了一个特定类型的值在内存中的存储形式，使能了多样的线性内存寻址模式。Data layout的接口。
2025-08-06 14:36:06
299
原创 MLIR Bufferization
MLIR中的Bufferization是将tensor语义的ops转换为memref语义的ops。Bufferization的顶层目标是：使用尽可能少的内存。拷贝尽可能少的内存。这就意味着要尽可能地做Buffer复用，Bufferization就成了一个类似Register Allocation的算法问题。在实际的Use Case中，对Bufferization还有一些额外的需求，比如重计算和计算一次并拷贝的权衡等。
2025-08-05 19:58:08
1208
原创 MLIR Introduction
MLIR 项目是一个构建可重用和可扩展编译器基础设施创新项目，旨在解决软件碎片化问题，显著降低构建特定领域编译器的成本。
2025-08-04 20:18:52
410
原创 Triton源码分析 - 目录
本系列文章主要分析Triton-shared编译管线相关内容，后续视情况添加NVIDIA GPU编译管线相关内容。
2025-07-31 17:39:26
307
原创 MLIR TableGen
TableGen 是一种领域特定语言（DSL），TableGen 的设计目标是允许编写灵活的描述，并将记录的通用特性提取出来，从而减少重复代码并提高代码的可维护性。TableGen 的前端解析.td文件，这些文件包含了用 TableGen 语言编写的声明和定义。前端将这些声明和定义实例化，生成一个中间表示（IR），这个 IR 包含了所有定义的记录（records）和类（classes）。生成的中间表示（IR）会被传递给特定领域的后端进行处理。后端根据 IR 生成目标代码，通常是 C++ 代码。
2025-07-31 17:33:50
872
原创 Triton - Dialect
其中：%0：右边expression的结果值的名字（Value的name）tt：表示Dialect名称空间为tt（Triton）splat：operation的名字%1：operation的输入i32：%1的类型tensor<1024*i32>：operation的结果类型（即3%的类型）loc(%loc5)：对应源码的行号，调试信息。如下是一个pytorch cat算子的Triton DSL（inductor产生）
2025-07-28 00:00:50
1252
空空如也
空空如也
TA创建的收藏夹 TA关注的收藏夹
TA关注的人