================
@@ -695,6 +705,37 @@ def NVVM_MBarrierInvalOp :
NVVM_VoidIntrinsicOp<"mbarrier.inval">,
let assemblyFormat = "$addr attr-dict `:` type(operands)";
}
+def NVVM_MBarrierCheckLayoutOp :
+ NVVM_SingleResultIntrinsicOp<"mbarrier.check_layout",
+ [NVVMRequiresSM<90>]> {
+ let summary = "MBarrier Check-Layout Operation";
+ let description = [{
+ The `nvvm.mbarrier.check_layout` operation tests whether the *mbarrier
+ object* at `addr` was initialized with the layout named by `layout`.
+
+ - `res`: An `i1` that is `true` when the *mbarrier object* has the queried
+ layout and `false` otherwise.
+
+ The operation takes the following operand and attribute:
+ - `addr`: A pointer to the memory location of the *mbarrier object*. The
+ `addr` must be a pointer to generic or shared::cta memory. When it is
+ generic, the underlying address must be within the shared::cta memory
+ space; otherwise the behavior is undefined.
+ - `layout`: The mbarrier layout version to test for. Only `0` and `1` are
+ valid values.
+
+ [For more information, see PTX
ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-check-layout)
+ }];
+
+ let results = (outs I1:$res);
+ let arguments = (ins
+ AnyTypeOf<[LLVM_PointerGeneric, LLVM_PointerShared]>:$addr,
+ ConfinedAttr<I32Attr, [IntMinValue<0>, IntMaxValue<1>]>:$layout);
----------------
grypp wrote:
```suggestion
AnyTypeOf<[LLVM_PointerGeneric, LLVM_PointerShared]>:$addr,
DefaultValuedAttr<ConfinedAttr<I32Attr,
[IntMinValue<0>, IntMaxValue<1>]>, "0">:$layout);
```
https://github.com/llvm/llvm-project/pull/217252
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
