Skip to content

Commit abc7b42

Browse files
committed
[mlir][GPU] Add builders to allow passing in integer upper_bounds
These are convendience builders to allow user code that knows (or optionally knows) the upper_bound on, say, a `gpu.thread_id` to specify that bound without needing to manually construct an IndexAttr.
1 parent 45f420e commit abc7b42

File tree

1 file changed

+13
-0
lines changed

1 file changed

+13
-0
lines changed

mlir/include/mlir/Dialect/GPU/IR/GPUOps.td

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -74,7 +74,20 @@ class GPU_IndexOp<string mnemonic, list<Trait> traits = []> :
7474
}]>,
7575
OpBuilder<(ins "::mlir::Type":$resultType, "::mlir::gpu::Dimension":$dimension), [{
7676
build($_builder, $_state, resultType, dimension, /*upperBound=*/nullptr);
77+
}]>,
78+
OpBuilder<(ins "::mlir::gpu::Dimension":$dimension, "std::optional<int64_t>":$upperBound), [{
79+
::mlir::IntegerAttr upperBoundAttr = nullptr;
80+
if (upperBound.has_value())
81+
upperBoundAttr = $_builder.getIndexAttr(*upperBound);
82+
build($_builder, $_state, dimension, upperBoundAttr);
83+
}]>,
84+
OpBuilder<(ins "::mlir::Type":$resultType, "::mlir::gpu::Dimension":$dimension, "std::optional<int64_t>":$upperBound), [{
85+
::mlir::IntegerAttr upperBoundAttr = nullptr;
86+
if (upperBound.has_value())
87+
upperBoundAttr = $_builder.getIndexAttr(*upperBound);
88+
build($_builder, $_state, resultType, dimension, upperBoundAttr);
7789
}]>
90+
7891
];
7992
}
8093

0 commit comments

Comments
 (0)