Skip to content

Commit b5ff576

Browse files
authored
[MLIR] Better document for basic PTX Builder Interface (NFC) (#67016)
1 parent 3301fd2 commit b5ff576

File tree

1 file changed

+25
-18
lines changed

1 file changed

+25
-18
lines changed

mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td

Lines changed: 25 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -102,41 +102,48 @@ def BasicPtxBuilderOpInterface : OpInterface<"BasicPtxBuilderInterface"> {
102102
let description = [{
103103
Interface to generate inline assembly with PTX for basic operations.
104104

105-
Interface is used in `convert-nvvm-to-llvm` pass that lowers Ops supports this interface to inline assembly Op. Interface has several methods and they are used for this lowering.
105+
Interface is used in `convert-nvvm-to-llvm` pass that lowers Ops supports
106+
this interface to inline assembly Op. Interface has several methods and
107+
they are used for this lowering.
106108

107109
`getPtx` method returns PTX code.
108110

109-
`hasSideEffect` is used to set whether the op has any side effect on the memory.
111+
`hasSideEffect` is used to set whether the op has any side effect on the
112+
memory.
110113

111-
`hasIntrinsic` returns whether the operation has intrinsic support in LLVM. This is useful for the Ops that don't have intrinsic support for each case.
114+
`hasIntrinsic` returns whether the operation has intrinsic support in LLVM.
115+
This is useful for the Ops that don't have intrinsic support for each case.
112116

113-
`getAsmValues` returns arguments to pass PTX code. The order of arguments is started from the results and they are used as write, followed by the operands and attributes.
117+
`getAsmValues` returns arguments to pass PTX code. The order of arguments
118+
is started from the results and they are used as write, followed by the
119+
operands and attributes.
114120

115121
Example:
116-
117122
If we have following Op definition that returns PTX code by `getPtx`.
118123

119124
```tablegen
120-
def NVVM_MBarrierArriveExpectTxOp : NVVM_Op<\"mbarrier.arrive.expect_tx\",
121-
[DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>]>,
122-
Results<(outs LLVM_Type:$res)>, Arguments<(ins LLVM_i64ptr_any:$addr, I32:$txcount)> {
123-
...
124-
let extraClassDefinition = [{
125-
std::string $cppClass::getPtx() { return std::string(\"mbarrier.arrive.expect_tx.b64 %0, [%1], %2;\"); }
126-
}\];
127-
}
125+
def NVVM_MyOp : NVVM_Op<"myop",
126+
[DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>]>,
127+
Results<(outs LLVM_Type:$res)>,
128+
Arguments<(ins LLVM_i64ptr_any:$op1, I32:$op2)> {
129+
...
130+
let extraClassDefinition = [{
131+
std::string $cppClass::getPtx() {
132+
return std::string("my.ptx.code %0, %1, %2;");
133+
}
134+
} ];
128135
```
129136

130137
The NVVM Op will look like below:
131138
```mlir
132-
%0 = nvvm.mbarrier.arrive.expect_tx %barrier, %txcount : !llvm.ptr, i32 -> i32
139+
%0 = my.ptx.code %1, %2 : !llvm.ptr, i32 -> i32
133140
```
134141

135-
The `convert-nvvm-to-llvm` Pass returns the PTX code below. The order of
136-
arguments are kept the same. The read/write modifiers are set based on the
137-
input and result types.
142+
The `convert-nvvm-to-llvm` Pass generates the PTX code below. The order of
143+
arguments are kept the same. The read and write modifiers are set based on
144+
the input and result types.
138145
```mlir
139-
%0 = llvm.inline_asm has_side_effects asm_dialect = att "mbarrier.arrive.expect_tx.b64 %0, [%1], %2;", "=r,l,r" %arg0, %arg1 : (!llvm.ptr, i32) -> i32
146+
%0 = llvm.inline_asm has_side_effects asm_dialect = att "my.ptx.code %0, %1, %2;", "=r,l,r" %arg0, %arg1 : (!llvm.ptr, i32) -> i32
140147
```
141148

142149
}];

0 commit comments

Comments
 (0)