[Mlir-commits] [mlir] [mlir][nvvm] Fix the PTX lowering of wgmma.mma_async (PR #76150)
Adam Paszke
llvmlistbot at llvm.org
Thu Dec 21 05:08:04 PST 2023
https://github.com/apaszke updated https://github.com/llvm/llvm-project/pull/76150
>From 0aff6aee640217b2cf4b771dde63dac8493d4fdf Mon Sep 17 00:00:00 2001
From: Adam Paszke <apaszke at google.com>
Date: Thu, 21 Dec 2023 13:05:49 +0000
Subject: [PATCH] [mlir][nvvm] Fix the PTX lowering of wgmma.mma_async
The default layout of A and B matrices is row- and column-major
respectively, meaning that the transpose flags have opposite meanings
between those two operands.
---
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 4f5d71e10f68c1..a4de89d928e1be 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -1003,7 +1003,7 @@ void NVVM::WgmmaMmaAsyncOp::getAsmValues(
{makeConstantI32(rewriter, static_cast<int>(getLayoutA())),
mlir::NVVM::PTXRegisterMod::Read});
asmValues.push_back(
- {makeConstantI32(rewriter, static_cast<int>(getLayoutB())),
+ {makeConstantI32(rewriter, 1 - static_cast<int>(getLayoutB())),
mlir::NVVM::PTXRegisterMod::Read});
}
}
More information about the Mlir-commits
mailing list