Attribute | MLIR Type | Description |
atomic_rmw_op | ::mlir::triton::RMWOpAttr | allowed 32-bit signless integer cases: 1, 2, 3, 4, 5, 6, 7, 8, 9, 10{{% markdown %}}Enum cases:
* and (`AND`)
* or (`OR`)
* xor (`XOR`)
* add (`ADD`)
* fadd (`FADD`)
* max (`MAX`)
* min (`MIN`)
* umax (`UMAX`)
* umin (`UMIN`)
* exch (`XCHG`){{% /markdown %}} |
sem | ::mlir::triton::MemSemanticAttr | allowed 32-bit signless integer cases: 1, 2, 3, 4{{% markdown %}}Enum cases:
* relaxed (`RELAXED`)
* acquire (`ACQUIRE`)
* release (`RELEASE`)
* acq_rel (`ACQUIRE_RELEASE`){{% /markdown %}} |
scope | ::mlir::triton::MemSyncScopeAttr | allowed 32-bit signless integer cases: 1, 2, 3{{% markdown %}}Enum cases:
* gpu (`GPU`)
* cta (`CTA`)
* sys (`SYSTEM`){{% /markdown %}} |
#### Operands:
| Operand | Description |
| :-----: | ----------- |
| `ptr` | ptr or ranked tensor of ptr values
| `val` | floating-point or ranked tensor of floating-point values or integer or ranked tensor of integer values or ptr or ranked tensor of ptr values or ptr
| `mask` | 1-bit signless integer or ranked tensor of 1-bit signless integer values
#### Results:
| Result | Description |
| :----: | ----------- |
| `result` | floating-point or ranked tensor of floating-point values or integer or ranked tensor of integer values or ptr or ranked tensor of ptr values or ptr
### `tt.bitcast` (triton::BitcastOp)
_Cast between types of the same bitwidth_
Syntax:
```
operation ::= `tt.bitcast` $src attr-dict `:` type($src) `->` type($result)
```
Traits: `AlwaysSpeculatableImplTrait`, `Elementwise`, `SameOperandsAndResultEncoding`, `SameOperandsAndResultShape`, `TensorSizeTrait`, `VerifyTensorLayoutsTrait`
Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`
Effects: `MemoryEffects::Effect{}`
#### Operands:
| Operand | Description |
| :-----: | ----------- |
| `src` | floating-point or ranked tensor of floating-point values or integer or ranked tensor of integer values or ptr or ranked tensor of ptr values or ptr
#### Results:
| Result | Description |
| :----: | ----------- |
| `result` | floating-point or ranked tensor of floating-point values or integer or ranked tensor of integer values or ptr or ranked tensor of ptr values or ptr
### `tt.broadcast` (triton::BroadcastOp)
_Broadcast a tensor_
Syntax:
```
operation ::= `tt.broadcast` $src attr-dict `:` type($src) `->` type($result)
```
For a given tensor, broadcast changes one or more dimensions with size 1
to a new size, e.g. tensor<1x32x1xf32> -> tensor<2x32x4xf32>. You cannot
change the size of a non-1 dimension.
Traits: `AlwaysSpeculatableImplTrait`, `SameOperandsAndResultElementType`, `SameOperandsAndResultEncoding`, `TensorSizeTrait`, `VerifyTensorLayoutsTrait`
Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`
Effects: `MemoryEffects::Effect{}`
#### Operands:
| Operand | Description |
| :-----: | ----------- |
| `src` | ranked tensor of floating-point or integer or ptr values
#### Results:
| Result | Description |
| :----: | ----------- |
| `result` | ranked tensor of floating-point or integer or ptr values
### `tt.cat` (triton::CatOp)
_Concatenate 2 tensors_
Syntax:
```
operation ::= `tt.cat` $lhs `,` $rhs attr-dict `:` type($lhs) `->` type($result)
```
Traits: `SameOperandsAndResultElementType`, `SameTypeOperands`, `TensorSizeTrait`, `VerifyTensorLayoutsTrait`
Interfaces: `NoMemoryEffect (MemoryEffectOpInterface)`
Effects: `MemoryEffects::Effect{}`
#### Operands:
| Operand | Description |
| :-----: | ----------- |
| `lhs` | ranked tensor of floating-point or integer or ptr values
| `rhs` | ranked tensor of floating-point or integer or ptr values
#### Results:
| Result | Description |
| :----: | ----------- |
| `result` | ranked tensor of floating-point or integer or ptr values
### `tt.clampf` (triton::ClampFOp)
_Clamp operation for floating point types_
Syntax:
```
operation ::= `tt.clampf` $x `,` $min `,` $max `,` `propagateNan` `=` $propagateNan attr-dict `:` type($result)
```
Clamp operation for floating point types.
The operation takes three arguments: x, min, and max. It returns a tensor of the same shape as x with its values clamped to the range [min, max].
Traits: `AlwaysSpeculatableImplTrait`, `Elementwise`, `SameOperandsAndResultType`, `TensorSizeTrait`, `VerifyTensorLayoutsTrait`
Interfaces: `ConditionallySpeculatable`, `InferTypeOpInterface`, `NoMemoryEffect (MemoryEffectOpInterface)`
Effects: `MemoryEffects::Effect{}`
#### Attributes:
Attribute | MLIR Type | Description |
boundaryCheck | ::mlir::DenseI32ArrayAttr | i32 dense array attribute |
padding | ::mlir::triton::PaddingOptionAttr | allowed 32-bit signless integer cases: 1, 2{{% markdown %}}Enum cases:
* zero (`PAD_ZERO`)
* nan (`PAD_NAN`){{% /markdown %}} |
cache | ::mlir::triton::CacheModifierAttr | allowed 32-bit signless integer cases: 1, 2, 3, 4, 5, 6{{% markdown %}}Enum cases:
* none (`NONE`)
* ca (`CA`)
* cg (`CG`)
* wb (`WB`)
* cs (`CS`)
* wt (`WT`){{% /markdown %}} |
evict | ::mlir::triton::EvictionPolicyAttr | allowed 32-bit signless integer cases: 1, 2, 3{{% markdown %}}Enum cases:
* evict_normal (`NORMAL`)
* evict_first (`EVICT_FIRST`)
* evict_last (`EVICT_LAST`){{% /markdown %}} |
isVolatile | ::mlir::BoolAttr | bool attribute |
#### Operands:
| Operand | Description |
| :-----: | ----------- |
| `ptr` | ptr or ranked tensor of ptr values or ptr
| `mask` | 1-bit signless integer or ranked tensor of 1-bit signless integer values
| `other` | floating-point or ranked tensor of floating-point values or integer or ranked tensor of integer values or ptr or ranked tensor of ptr values or ptr
#### Results:
| Result | Description |
| :----: | ----------- |
| `result` | floating-point or ranked tensor of floating-point values or integer or ranked tensor of integer values or ptr or ranked tensor of ptr values or ptr
### `tt.make_range` (triton::MakeRangeOp)
_Make range_
Syntax:
```
operation ::= `tt.make_range` attr-dict `:` type($result)
```
Returns an 1D int32 tensor.
Values span from $start to $end (exclusive), with step = 1
Traits: `AlwaysSpeculatableImplTrait`, `TensorSizeTrait`, `VerifyTensorLayoutsTrait`
Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`
Effects: `MemoryEffects::Effect{}`
#### Attributes: