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:¶
| Attribute | MLIR Type | Description |
|---|---|---|
isStart | ::mlir::BoolAttr | bool attribute |
regionId | ::mlir::IntegerAttr | 32-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 %}} |