ProtonGPUOps¶
proton_gpu.circular_store (triton::proton::gpu::CircularStoreOp)¶
将值存入循环缓冲区
语法
operation ::= `proton_gpu.circular_store` (`start` $isStart^):(`end`)? $segment `,` $counter attr-dict `:`
qualified(type($segment)) `,` type($counter)
将一个指标 counter 存入由内部内存 segment 支持的循环缓冲区。自动更新。如果 segment 缓冲区已满，较旧的指标计数器将被丢弃。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Read on ::mlir::SideEffects::DefaultResource, MemoryEffects::Write on ::mlir::SideEffects::DefaultResource}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
isStart | ::mlir::UnitAttr | 单元属性 |
scopeId | ::mlir::IntegerAttr | 32位无符号整数属性 |
操作数:¶
操作数 |
描述 |
|---|---|
|
内部缓冲区中的一个段 |
|
32 位无符号整数或 64 位无符号整数 |
proton_gpu.finalize (triton::proton::gpu::FinalizeOp)¶
完成内核内分析器
语法
operation ::= `proton_gpu.finalize` $segment `,` $scratchPtr attr-dict `:` qualified(type($segment)) `,` qualified(type($scratchPtr))
将元数据和分析结果写回全局内存。segment 是包含分析数据的内部分析缓冲区的段。scratchPtr 是分析暂存缓冲区的地址。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Read on ::mlir::triton::GlobalMemory}、MemoryEffects::Effect{MemoryEffects::Read on ::mlir::triton::gpu::SharedMemory}、MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
操作数：¶
操作数 |
描述 |
|---|---|
|
内部缓冲区中的一个段 |
|
ptr |
proton_gpu.global_scratch_alloc (triton::proton::gpu::GlobalScratchAllocOp)¶
分配一个全局内存分析缓冲区
语法
operation ::= `proton_gpu.global_scratch_alloc` attr-dict `:` qualified(type($result))
此操作在全局内存中分配一个分析缓冲区。每个被分析的 Triton 程序都将持有一个指向该缓冲区的指针，直到分析完成。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Allocate on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
nbytes | ::mlir::IntegerAttr | 32位无符号整数属性 |
alignment | ::mlir::IntegerAttr | 32位无符号整数属性 |
结果:¶
结果 |
描述 |
|---|---|
|
ptr |
proton_gpu.init_ctx (triton::proton::gpu::InitCtxOp)¶
初始化内核内分析器的 Warp 级上下文
语法
operation ::= `proton_gpu.init_ctx` $scratchPtr attr-dict `:` qualified(type($scratchPtr))
为 scratchPtr（分析暂存缓冲区的基地址）中的所有 Warp 初始化内核内分析器的 Warp 级上下文。它不能在 ttg.warp_specialize 内部调用。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
操作数：¶
操作数 |
描述 |
|---|---|
|
ptr |
proton_gpu.initialize (triton::proton::gpu::InitializeOp)¶
初始化内核内分析器
语法
operation ::= `proton_gpu.initialize` $scratchPtr attr-dict `:` qualified(type($scratchPtr))
通过将辅助元数据填充到头部来初始化内核内分析器。scratchPtr 是存储头部的分析暂存缓冲区的基地址。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
操作数:¶
操作数 |
描述 |
|---|---|
|
ptr |
proton_gpu.read_counter (triton::proton::gpu::ReadCounterOp)¶
将 GPU 指标计数器读入标量寄存器
语法
operation ::= `proton_gpu.read_counter` attr-dict `:` type($counter)
将一个 GPU 指标计数器读入标量寄存器。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Read on ::mlir::SideEffects::DefaultResource, MemoryEffects::Write on ::mlir::SideEffects::DefaultResource}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
metric | ::mlir::triton::proton::MetricTypeAttr | 要分析的指标类型{{% markdown %}} 用于指示要分析的指标的属性。支持以下指标： - CYCLE: 周期计数指标。{{% /markdown %}} |
结果:¶
结果 |
描述 |
|---|---|
|
32 位无符号整数或 64 位无符号整数 |
proton_gpu.restore_ctx (triton::proton::gpu::RestoreCtxOp)¶
恢复当前的 Warp 级上下文
语法
operation ::= `proton_gpu.restore_ctx` $segment `,` $scratchPtr attr-dict `:` qualified(type($segment)) `,` qualified(type($scratchPtr))
从 scratchPtr（分析暂存缓冲区的基地址）恢复 $segment 中的当前 Warp 上下文。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Read on ::mlir::triton::GlobalMemory}、MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
操作数：¶
操作数 |
描述 |
|---|---|
|
内部缓冲区中的一个段 |
|
ptr |
proton_gpu.save_ctx (triton::proton::gpu::SaveCtxOp)¶
保存当前的 Warp 级上下文
语法
operation ::= `proton_gpu.save_ctx` $segment `,` $scratchPtr attr-dict `:` qualified(type($segment)) `,` qualified(type($scratchPtr))
将 $segment 中的当前 Warp 上下文保存到 scratchPtr（分析暂存缓冲区的基地址）。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
操作数:¶
操作数 |
描述 |
|---|---|
|
内部缓冲区中的一个段 |
|
ptr |
proton_gpu.segment_alloc (triton::proton::gpu::SegmentAllocOp)¶
获取内部缓冲区段的基偏移量
语法
operation ::= `proton_gpu.segment_alloc` $buffer attr-dict `:` qualified(type($buffer)) `->` type($segment)
内部缓冲区为每个分析“单元”划分成段。此操作获取内部缓冲区中内存段的位置。
特性: AlwaysSpeculatableImplTrait
接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
效果: MemoryEffects::Effect{}
操作数:¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
结果：¶
结果 |
描述 |
|---|---|
|
内部缓冲区中的一个段 |