address comments

Author: grypp

mlir/include/mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.h

@@ -26,11 +26,11 @@ namespace NVVM {
 enum class PTXRegisterMod {
   /// Read register with no modifier
   Read = 0,
-  /// Read register with '=' modifier
+  /// Write register with '=' modifier
   Write = 2,
-  /// Read register with '+' modifier.
-  /// Note that, this is not natively supported by LLVM, but it is possible to
-  /// set read and write for the same operand.
+  /// ReadWrite register with '+' modifier.
+  /// Note that, this is not natively supported by LLVM, the Interface does
+  /// mapping
   ReadWrite = 1,
 };
 
@@ -69,15 +69,17 @@ class PtxBuilder {
   std::string registerConstraints;
   // Modifiers
   SmallVector<PTXRegisterMod> registerModifiers;
+  // Has return value as write-only or read-write
   bool hasResult = false;
-  bool needsManualMapping = false;
+  // Indicates if the Op will handle the register mapping manually.
+  bool needsManualRegisterMapping = false;
 
 public:
   /// Single constructor that only initializes members.
   PtxBuilder(Operation *op, PatternRewriter &rewriter,
-             bool needsManualMapping = false)
+             bool needsManualRegisterMapping = false)
       : interfaceOp(op), rewriter(rewriter),
-        needsManualMapping(needsManualMapping) {}
+        needsManualRegisterMapping(needsManualRegisterMapping) {}
 
   /// Add an operand with the read/write input type.
   void insertValue(Value v, PTXRegisterMod itype = PTXRegisterMod::Read);

mlir/include/mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td

@@ -125,7 +125,7 @@ def BasicPtxBuilderOpInterface : OpInterface<"BasicPtxBuilderInterface"> {
              1) Adds results 
              2) Adds operands 
              3) Adds attributes
-             Returns true if it does the mapping manually
+             Returns true if the OP is going to do register mapping itself
           }],
          /*retType=*/"bool",
          /*methodName=*/"getAsmValues",
@@ -151,7 +151,7 @@ def BasicPtxBuilderOpInterface : OpInterface<"BasicPtxBuilderInterface"> {
              asmValues.push_back({val, mlir::NVVM::PTXRegisterMod::Read});
              }
            }
-           return false; // No needs manual mapping
+           return false; // No manual mapping needed
          }]
        >
   ];

mlir/lib/Dialect/LLVMIR/IR/BasicPtxBuilderInterface.cpp

@@ -129,9 +129,10 @@ void PtxBuilder::insertValue(Value v, PTXRegisterMod itype) {
 
 /// Check if the operation needs to pack and unpack results.
 static bool
-needsPackUnpack(BasicPtxBuilderInterface interfaceOp, bool needsManualMapping,
+needsPackUnpack(BasicPtxBuilderInterface interfaceOp,
+                bool needsManualRegisterMapping,
                 SmallVectorImpl<PTXRegisterMod> &registerModifiers) {
-  if (needsManualMapping)
+  if (needsManualRegisterMapping)
     return false;
   const unsigned writeOnly = interfaceOp->getNumResults();
   const unsigned readWrite =
@@ -145,13 +146,15 @@ needsPackUnpack(BasicPtxBuilderInterface interfaceOp, bool needsManualMapping,
 /// If the operation has multiple results, it packs them into a struct
 /// type. Otherwise, it returns the original result types.
 static SmallVector<Type>
-packResultTypes(BasicPtxBuilderInterface interfaceOp, bool needsManualMapping,
+packResultTypes(BasicPtxBuilderInterface interfaceOp,
+                bool needsManualRegisterMapping,
                 SmallVectorImpl<PTXRegisterMod> &registerModifiers,
                 SmallVectorImpl<Value> &ptxOperands) {
   MLIRContext *ctx = interfaceOp->getContext();
   TypeRange resultRange = interfaceOp->getResultTypes();
 
-  if (!needsPackUnpack(interfaceOp, needsManualMapping, registerModifiers)) {
+  if (!needsPackUnpack(interfaceOp, needsManualRegisterMapping,
+                       registerModifiers)) {
     // Single value path:
     if (interfaceOp->getResults().size() == 1)
       return SmallVector<Type>{resultRange.front()};
@@ -333,7 +336,7 @@ LLVM::InlineAsmOp PtxBuilder::build() {
                                                   LLVM::AsmDialect::AD_ATT);
 
   SmallVector<Type> resultTypes = packResultTypes(
-      interfaceOp, needsManualMapping, registerModifiers, ptxOperands);
+      interfaceOp, needsManualRegisterMapping, registerModifiers, ptxOperands);
 
   // Remove the last comma from the constraints string.
   if (!registerConstraints.empty() &&
@@ -342,7 +345,7 @@ LLVM::InlineAsmOp PtxBuilder::build() {
   registerConstraints = canonicalizeRegisterConstraints(registerConstraints);
 
   std::string ptxInstruction = interfaceOp.getPtx();
-  if (!needsManualMapping)
+  if (!needsManualRegisterMapping)
     ptxInstruction = rewriteAsmPlaceholders(ptxInstruction);
 
   // Add the predicate to the asm string.
@@ -379,13 +382,14 @@ void PtxBuilder::buildAndReplaceOp() {
     return;
   }
 
-  if (needsManualMapping) {
+  if (needsManualRegisterMapping) {
     rewriter.replaceOp(interfaceOp, inlineAsmOp->getResults());
     return;
   }
 
   // Case 1: Simple path, return single scalar
-  if (!needsPackUnpack(interfaceOp, needsManualMapping, registerModifiers)) {
+  if (!needsPackUnpack(interfaceOp, needsManualRegisterMapping,
+                       registerModifiers)) {
     if (inlineAsmOp->getNumResults() > 0) {
       rewriter.replaceOp(interfaceOp, inlineAsmOp->getResults());
     } else {

mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp

@@ -1884,7 +1884,7 @@ bool NVVM::InlinePtxOp::getAsmValues(
     asmValues.push_back({arg, mlir::NVVM::PTXRegisterMod::Read});
   if (getPredicate())
     asmValues.push_back({getPredicate(), mlir::NVVM::PTXRegisterMod::Read});
-  return false; // Needs manual mapping
+  return false; // No manual mapping needed
 }
 
 //===----------------------------------------------------------------------===//