[Mlir-commits] [mlir] [MLIR][NVGPU] Add `mbarrier.get` Op (PR #133221)
Durgadoss R
llvmlistbot at llvm.org
Thu Mar 27 04:40:29 PDT 2025
================
@@ -322,6 +322,25 @@ def NVGPU_MBarrierCreateOp : NVGPU_Op<"mbarrier.create", []> {
}];
}
+def NVGPU_MBarrierGetOp : NVGPU_Op<"mbarrier.get", []> {
+ let summary = "Return a pointer to an `nvgpu.mbarrier`.";
+ let description = [{
+ The `nvgpu.mbarrier.get` operation retrieves a pointer to a specific
+ `mbarrier` object from a group of barriers created by the `nvgpu.mbarrier.create` operation.
+
+ Example:
+ ```mlir
+ %mbars = nvgpu.mbarrier.create -> !nvgpu.mbarrier.group<memorySpace = #gpu.address_space<workgroup>, num_barriers = 10>
+ %mbar_pointer = nvgpu.mbarrier.get %mbars[%c2] : !nvgpu.mbarrier.group<memorySpace = #gpu.address_space<workgroup>>
----------------
durga4github wrote:
I was wondering whether we should have a check to verify if the given `idx` is lesser than `num_barriers`.
But it seems the `idx` can be a runtime value also. So, I think we are good.
https://github.com/llvm/llvm-project/pull/133221
More information about the Mlir-commits
mailing list