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

[guraypp]

Summary

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

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

def WGMMAScaleOut : I32EnumAttr<"WGMMAScaleOut", "WGMMA input predicate",
   [WGMMAScaleOutZero, WGMMAScaleOutOne]> {
   let genSpecializedAttr = 0;
   let cppNamespace = "::mlir::NVVM";
 }
 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>>,]> 
 {
   let results = (outs LLVM_AnyAggregate:$results);
   let arguments = (ins 
     LLVM_AnyAggregate:$inouts,
     I64:$descriptorA, 

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

Type dstType = LLVM::LLVMStructType::getLiteral(
         getContext(), SmallVector<Type>(getNum(), i32));
     if (getType() != dstType)
       return emitOpError("expected destination type is a structure of ")
              << getNum() << " elements of type i32";
   }
   return success();
 }
 
-LogicalResult NVVM::WgmmaMmaSyncOp::verify() {
+LogicalResult NVVM::WgmmaMmaAsyncOp::verify() {
   Value outValue = getResults();
   auto stype = dyn_cast<LLVM::LLVMStructType>(outValue.getType());
   if (!stype)
     return emitOpError() << "expected results to be struct";
   Type outputType = stype.getBody().front();
   int outputSize = stype.getBody().size();
   for (Type t : stype.getBody()) {
     if (t != outputType)
return emitOpError()
            << " `satfinite` can be only used with s32 accumulator, however "
               "the current accumulator is "
            << outputType;
   }
 
   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 ||
                getTypeA() == mlir::NVVM::MMATypes::bf16;
 
   Value outValue = getResults() ? getResults() : getInouts();
   auto stype = dyn_cast<LLVM::LLVMStructType>(outValue.getType());
   Type outputType = stype.getBody().front();
ss << ", $" << (expectedOutputRegisters + 5) << ",  $"
        << (expectedOutputRegisters + 6);
   }
   ss << ";\n"
      << "}\n";
   ss.flush();
   return ptx;
 }
 
-void NVVM::WgmmaMmaSyncOp::getAsmValues(
+void NVVM::WgmmaMmaAsyncOp::getAsmValues(
     RewriterBase &rewriter,
     llvm::SmallVectorImpl<std::pair<mlir::Value, mlir::NVVM::PTXRegisterMod>>
         &asmValues) {
   Value outValue = getResults() ? getResults() : getInouts();
   auto stype = dyn_cast<LLVM::LLVMStructType>(outValue.getType());
   Type outputType = stype.getBody().front();
   bool isF16 = getTypeA() == mlir::NVVM::MMATypes::f16 ||
                getTypeA() == mlir::NVVM::MMATypes::bf16;