[Mlir-commits] [mlir] [MLIR][NVVM] Add `inline_ptx` op (PR #139923)
Durgadoss R
llvmlistbot at llvm.org
Thu May 15 06:18:40 PDT 2025
================
@@ -236,6 +236,76 @@ foreach index = !range(0, 32) in {
def NVVM_EnvReg # index # Op : NVVM_SpecialRegisterOp<"read.ptx.sreg.envreg" # index>;
}
+//===----------------------------------------------------------------------===//
+// Inline PTX op definition
+//===----------------------------------------------------------------------===//
+
+def NVVM_InlinePtxOp : NVVM_Op<"inline_ptx",
+ [DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>,
+ AttrSizedOperandSegments]>
+{
+ let summary = "Inline PTX Op";
+ let description = [{This op allows using PTX directly within the NVVM
+ dialect, while greatly simplifying llvm.inline_asm generation. It
+ automatically handles register size selection and sets the correct
+ read/write access for each operand. The operation leverages the
+ `BasicPtxBuilderInterface` to abstract away low-level details of
+ PTX assembly formatting.
+
+ The `predicate` attribute is used to specify a predicate for the
+ PTX instruction.
+
+ Example 1: Read-only Parameters
+ ```mlir
+ nvvm.inline_ptx "mbarrier.init.b64 [$0], $1;" (%barrier_gen, %count) : !llvm.ptr, i32
+
+ // Lowers to:
+ llvm.inline_asm has_side_effects asm_dialect = att
+ "mbarrier.init.b64 [$0], $1;", "l,r" %arg0, %arg2 : (!llvm.ptr, i32) -> ()
----------------
durga4github wrote:
> > In all these given examples, the mlir looks clean.
>
> What do you mean by clean?
I meant the mlir version exactly matches the unit-test but the lowered version still refers to (!llvm.ptr, i32) instead of "l, r"
https://github.com/llvm/llvm-project/pull/139923
More information about the Mlir-commits
mailing list