[NVGPU] Add debug in nvgpu (nfc)

mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp

 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/Dialect/LLVMIR/LLVMTypes.h"
 #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
 #include "mlir/Dialect/NVGPU/IR/NVGPUDialect.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/IR/TypeUtilities.h"
 #include "mlir/Pass/Pass.h"
+#include "llvm/Support/Debug.h"
 #include "llvm/Support/raw_ostream.h"
 
+#define DEBUG_TYPE "nvgpu-to-nvvm"
+#define DBGS() (llvm::dbgs() << '[' << DEBUG_TYPE << "] ")
+#define DBGSE() (llvm::dbgs())
+
 namespace mlir {
 #define GEN_PASS_DEF_CONVERTNVGPUTONVVMPASS
 #include "mlir/Conversion/Passes.h.inc"
 } // namespace mlir
 
 using namespace mlir;
 
 /// GPU has 32 bit registers, this function truncates values when larger width
auto shiftRight = [&](Value value, unsigned shift) -> Value {
       return rewriter.create<LLVM::LShrOp>(loc, ti64, value, makeConst(shift));
     };
     auto insertBit = [&](Value desc, Value val, int startBit) {
       return rewriter.create<LLVM::OrOp>(loc, ti64, desc,
                                          shiftLeft(val, startBit));
     };
 
     int ex4LSB = 4;
-    Value strideDim = makeConst((layout << 3) >> ex4LSB);
     int64_t sizeN = op.getTensorMap().getType().getTensor().getDimSize(0);
-    Value leadDim = makeConst((sizeN * layout) >> ex4LSB);
+    uint64_t strideDimVal = (layout << 3) >> ex4LSB;
+    uint64_t leadDimVal = (sizeN * layout) >> ex4LSB;
+    uint64_t offsetVal = 0;
+
+    Value strideDim = makeConst(strideDimVal);
+    Value leadDim = makeConst(leadDimVal);
+
     Value baseAddr = getStridedElementPtr(
         op->getLoc(), cast<MemRefType>(op.getTensor().getType()),
         adaptor.getTensor(), {}, rewriter);
     Value basePtr = rewriter.create<LLVM::PtrToIntOp>(loc, ti64, baseAddr);
     // Just use 14 bits for base address
     Value basePtr14bit = shiftRight(shiftLeft(basePtr, 46), 50);
 
     int startSwizzleBit = 62, startOffsetBit = 49, startStrideBit = 32,
         startLeadBit = 16, startBaseAddrBit = 0;
     Value dsc = makeConst(0);
     // // [62,64)  swizzle type
     dsc = insertBit(dsc, makeConst(swizzle), startSwizzleBit);
     // // [49,52)  base_offset
-    dsc = insertBit(dsc, makeConst(0), startOffsetBit);
+    dsc = insertBit(dsc, makeConst(offsetVal), startOffsetBit);
     // // [32,46)  stride
     dsc = insertBit(dsc, strideDim, startStrideBit);
     // // [16,30)  leading dimension
     dsc = insertBit(dsc, leadDim, startLeadBit);
     // // [0,14)   start_address
     dsc = insertBit(dsc, basePtr14bit, startBaseAddrBit);
 
+    LLVM_DEBUG(DBGS() << "Generating wgmma.descriptor: "
+                      << "leading_off:" << leadDimVal << "\t"
+                      << "stride_off :" << strideDimVal << "\t"
+                      << "base_offset:" << offsetVal << "\t"
+                      << "layout_type:" << swizzle << " ("
+                      << nvgpu::stringifyTensorMapSwizzleKind(swizzleKind)
+                      << ")\n start_addr :  " << baseAddr << "\n");
+
     rewriter.replaceOp(op, dsc);
     return success();
   }
 };
 
 static Value makeI64Const(RewriterBase &rewriter, Operation *op,
                           int32_t index) {
   return rewriter.create<LLVM::ConstantOp>(op->getLoc(),