[Mlir-commits] [mlir] db11fd8 - [MLIR][NVVM] Change WgmmaMmaSyncOp to WgmmaMmaAsyncOp (NFC)

[Guray Ozen]

Author: Guray Ozen
Date: 2023-08-12T09:55:11+02:00
New Revision: db11fd8bf4ce1b1ce08db6011ea7c22d030eb473

URL: https://github.com/llvm/llvm-project/commit/db11fd8bf4ce1b1ce08db6011ea7c22d030eb473
DIFF: https://github.com/llvm/llvm-project/commit/db11fd8bf4ce1b1ce08db6011ea7c22d030eb473.diff

LOG: [MLIR][NVVM] Change WgmmaMmaSyncOp to WgmmaMmaAsyncOp (NFC)

WgmmaMmaSyncOp is asynchronous operation. There was a typo named op. This work fixes that.

Reviewed By: nicolasvasilache

Differential Revision: https://reviews.llvm.org/D157697

Added: 
    

Modified: 
    mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp

Removed: 
    


################################################################################
diff  --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 4845a226ec9c17..7ddf68e0a5b1a4 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -1497,7 +1497,7 @@ def WGMMAScaleOutAttr : EnumAttr<NVVM_Dialect, WGMMAScaleOut, "wgmma_scale_out">
   let assemblyFormat = "`<` $value `>`";
 }
 
-def NVVM_WgmmaMmaSyncOp : NVVM_Op<"wgmma.mma_async", 
+def NVVM_WgmmaMmaAsyncOp : NVVM_Op<"wgmma.mma_async", 
               [DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>,
                 PredOpTrait<"input struct and result struct must be the same type",
                   TCresIsSameAsOpBase<0, 0>>,]> 

diff  --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 25f837547b46f5..70c774233797bb 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -708,7 +708,7 @@ LogicalResult NVVM::LdMatrixOp::verify() {
   return success();
 }
 
-LogicalResult NVVM::WgmmaMmaSyncOp::verify() {
+LogicalResult NVVM::WgmmaMmaAsyncOp::verify() {
   Value outValue = getResults();
   auto stype = dyn_cast<LLVM::LLVMStructType>(outValue.getType());
   if (!stype)
@@ -889,7 +889,7 @@ LogicalResult NVVM::WgmmaMmaSyncOp::verify() {
   return success();
 }
 
-std::string NVVM::WgmmaMmaSyncOp::getPtx() {
+std::string NVVM::WgmmaMmaAsyncOp::getPtx() {
 
   int m = getShape().getM(), n = getShape().getN(), k = getShape().getK();
   bool isF16 = getTypeA() == mlir::NVVM::MMATypes::f16 ||
@@ -952,7 +952,7 @@ std::string NVVM::WgmmaMmaSyncOp::getPtx() {
   return ptx;
 }
 
-void NVVM::WgmmaMmaSyncOp::getAsmValues(
+void NVVM::WgmmaMmaAsyncOp::getAsmValues(
     RewriterBase &rewriter,
     llvm::SmallVectorImpl<std::pair<mlir::Value, mlir::NVVM::PTXRegisterMod>>
         &asmValues) {