ProtonOps

proton.record (triton::proton::RecordOp)

Record a GPU hardware event

Syntax:

operation ::= `proton.record` `(` operands `)` attr-dict

The operator records GPU events from performance counters. Currently only cycle counter is supported.

Example:

proton.record() {isStart = true, regionId = 4 : i32}
...
proton.record() {isStart = false, regionId = 4 : i32}
...
proton.record() {isStart = true, regionId = 1 : i32, granularity = 1 : i32}
...
proton.record() {isStart = false, regionId = 1 : i32, granularity = 1 : i32}

Interfaces: MemoryEffectOpInterface

Attributes:

AttributeMLIR TypeDescription
isStart::mlir::BoolAttrbool attribute
regionId::mlir::IntegerAttr32-bit signless integer attribute whose value is non-negative
metric::mlir::triton::proton::MetricAttr
allowed 32-bit signless integer cases: 0{{% markdown %}}Enum cases: * cycle (`CYCLE`){{% /markdown %}}
granularity::mlir::triton::proton::GranularityAttr
allowed 32-bit signless integer cases: 0, 1{{% markdown %}}Enum cases: * warpgroup (`WARPGROUP`) * warp (`WARP`){{% /markdown %}}