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（作用域 ID）	::mlir::IntegerAttr	32位无符号整数属性

操作数：

操作数	描述
segment（段）	内部缓冲区中的一个段
counter（计数器）	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}

操作数:

操作数	描述
segment（段）	内部缓冲区中的一个段
scratchPtr（暂存指针）	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}

操作数：

操作数	描述
scratchPtr（暂存指针）	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}

操作数:

操作数	描述
scratchPtr（暂存指针）	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 %}}

结果：

结果	描述
counter（计数器）	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}

操作数：

操作数	描述
segment（段）	内部缓冲区中的一个段
scratchPtr（暂存指针）	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))

将当前的 Warp 上下文从 $segment 保存到 scratchPtr（分析暂存缓冲区的基地址）中。

接口：MemoryEffectOpInterface (MemoryEffectOpInterface)

效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

操作数：

操作数	描述
segment（段）	内部缓冲区中的一个段
scratchPtr（暂存指针）	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{}

操作数：

操作数	描述
buffer（缓冲区）	Triton IR 类型系统中的内存描述符类型（::mlir::triton::gpu::MemDescType）或指针

结果：

结果	描述
segment（段）	内部缓冲区中的一个段