[MLIR][NVVM] Run clang-tidy #135006
Merged
[MLIR][NVVM] Run clang-tidy #135006
+58
−51
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
There was a problem hiding this comment.
Copilot reviewed 1 out of 1 changed files in this pull request and generated no comments.
Comments suppressed due to low confidence (2)
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:60
- [nitpick] The function name has been changed to lower camel case for consistency. Ensure that the same naming convention is consistently applied throughout the project.
static LogicalResult cpAsyncBulkTensorCommonVerifier(size_t tensorDims,
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:1289
- [nitpick] The use of the macro _none to replace an empty token could be confusing; consider renaming it to a more descriptive identifier to avoid potential conflicts and improve clarity.
return GET_CVT_F2TF32_ID(rna, _none, _satfinite);
|
@llvm/pr-subscribers-mlir-llvm Author: Guray Ozen (grypp) ChangesFull diff: https://github.com/llvm/llvm-project/pull/135006.diff 1 Files Affected:
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 09bff6101edd3..cced84ded9592 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -33,6 +33,7 @@
#include "llvm/AsmParser/Parser.h"
#include "llvm/IR/Attributes.h"
#include "llvm/IR/Function.h"
+#include "llvm/IR/IntrinsicsNVPTX.h"
#include "llvm/IR/Type.h"
#include "llvm/Support/Casting.h"
#include "llvm/Support/FormatVariadic.h"
@@ -56,7 +57,7 @@ using namespace NVVM;
// CpAsyncBulkTensorGlobalToSharedClusterOp (TMA Load)
// CpAsyncBulkTensorPrefetchOp (TMA Prefetch)
// CpAsyncBulkTensorReduceOp (TMA Store-Reduce)
-static LogicalResult CpAsyncBulkTensorCommonVerifier(size_t tensorDims,
+static LogicalResult cpAsyncBulkTensorCommonVerifier(size_t tensorDims,
bool isIm2Col,
size_t numIm2ColOffsets,
Location loc) {
@@ -81,7 +82,7 @@ static LogicalResult CpAsyncBulkTensorCommonVerifier(size_t tensorDims,
LogicalResult CpAsyncBulkTensorGlobalToSharedClusterOp::verify() {
size_t numIm2ColOffsets = getIm2colOffsets().size();
bool isIm2Col = numIm2ColOffsets > 0;
- return CpAsyncBulkTensorCommonVerifier(getCoordinates().size(), isIm2Col,
+ return cpAsyncBulkTensorCommonVerifier(getCoordinates().size(), isIm2Col,
numIm2ColOffsets, getLoc());
}
@@ -105,13 +106,13 @@ LogicalResult CpAsyncOp::verify() {
LogicalResult CpAsyncBulkTensorPrefetchOp::verify() {
size_t numIm2ColOffsets = getIm2colOffsets().size();
bool isIm2Col = numIm2ColOffsets > 0;
- return CpAsyncBulkTensorCommonVerifier(getCoordinates().size(), isIm2Col,
+ return cpAsyncBulkTensorCommonVerifier(getCoordinates().size(), isIm2Col,
numIm2ColOffsets, getLoc());
}
LogicalResult CpAsyncBulkTensorReduceOp::verify() {
bool isIm2Col = (getMode() == TMAStoreMode::IM2COL);
- return CpAsyncBulkTensorCommonVerifier(getCoordinates().size(), isIm2Col, 0,
+ return cpAsyncBulkTensorCommonVerifier(getCoordinates().size(), isIm2Col, 0,
getLoc());
}
@@ -183,14 +184,14 @@ static bool isIntegerPtxType(MMATypes type) {
MMATypes MmaOp::accumPtxType() {
std::optional<mlir::NVVM::MMATypes> val = inferOperandMMAType(
- getODSOperands(2).getTypes().front(), /*isAccum=*/true);
+ getODSOperands(2).getTypes().front(), /*isAccumulator=*/true);
assert(val.has_value() && "accumulator PTX type should always be inferrable");
return val.value();
}
MMATypes MmaOp::resultPtxType() {
std::optional<mlir::NVVM::MMATypes> val =
- inferOperandMMAType(getResult().getType(), /*isAccum=*/true);
+ inferOperandMMAType(getResult().getType(), /*isAccumulator=*/true);
assert(val.has_value() && "result PTX type should always be inferrable");
return val.value();
}
@@ -224,7 +225,7 @@ void MmaOp::print(OpAsmPrinter &p) {
}
}
std::optional<MMATypes> inferredType =
- inferOperandMMAType(regTypes.back(), /*isAccum=*/fragIdx >= 2);
+ inferOperandMMAType(regTypes.back(), /*isAccumulator=*/fragIdx >= 2);
if (inferredType)
ignoreAttrNames.push_back(frag.ptxTypeAttr);
}
@@ -243,7 +244,8 @@ void MmaOp::print(OpAsmPrinter &p) {
p.printOptionalAttrDict(this->getOperation()->getAttrs(), ignoreAttrNames);
// Print the types of the operands and result.
- p << " : " << "(";
+ p << " : "
+ << "(";
llvm::interleaveComma(SmallVector<Type, 3>{frags[0].regs[0].getType(),
frags[1].regs[0].getType(),
frags[2].regs[0].getType()},
@@ -363,14 +365,14 @@ ParseResult MmaOp::parse(OpAsmParser &parser, OperationState &result) {
if (failed(parser.resolveOperands(frag.regs, frag.regTypes,
parser.getNameLoc(), result.operands)))
return failure();
- frag.elemtype =
- inferOperandMMAType(frag.regTypes[0], /*isAccum=*/iter.index() < 2);
+ frag.elemtype = inferOperandMMAType(frag.regTypes[0],
+ /*isAccumulator*/ iter.index() < 2);
}
Type resultType;
if (parser.parseArrow() || parser.parseType(resultType))
return failure();
- frags[3].elemtype = inferOperandMMAType(resultType, /*isAccum=*/true);
+ frags[3].elemtype = inferOperandMMAType(resultType, /*isAccumulator*/ true);
std::array<StringRef, 2> names{"multiplicandAPtxType",
"multiplicandBPtxType"};
@@ -992,7 +994,9 @@ std::string NVVM::WgmmaMmaAsyncOp::getPtx() {
ss << "},";
// Need to map read/write registers correctly.
regCnt = (regCnt * 2);
- ss << " $" << (regCnt) << "," << " $" << (regCnt + 1) << "," << " p";
+ ss << " $" << (regCnt) << ","
+ << " $" << (regCnt + 1) << ","
+ << " p";
if (getTypeD() != WGMMATypes::s32) {
ss << ", $" << (regCnt + 3) << ", $" << (regCnt + 4);
}
@@ -1118,9 +1122,9 @@ LogicalResult NVVM::Tcgen05CpOp::verify() {
LogicalResult NVVM::MatchSyncOp::verify() {
if (getKind() == NVVM::MatchSyncKind::all) {
- auto Type = llvm::dyn_cast<LLVM::LLVMStructType>(getType());
- if (!Type || Type.getBody().size() != 2 ||
- !Type.getBody()[0].isInteger(32) || !Type.getBody()[1].isInteger(1)) {
+ auto type = llvm::dyn_cast<LLVM::LLVMStructType>(getType());
+ if (!type || type.getBody().size() != 2 ||
+ !type.getBody()[0].isInteger(32) || !type.getBody()[1].isInteger(1)) {
return emitOpError("match.sync 'all' returns a two element struct with "
"first element as i32 and second element as i1");
}
@@ -1161,7 +1165,7 @@ CpAsyncOp::getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt,
llvm::Intrinsic::ID id;
auto cpAsyncOp = cast<NVVM::CpAsyncOp>(op);
- bool hasCpSize = cpAsyncOp.getCpSize() ? true : false;
+ bool hasCpSize = static_cast<bool>(cpAsyncOp.getCpSize());
switch (cpAsyncOp.getSize()) {
case 4:
id = GET_CP_ASYNC_ID(ca, 4, hasCpSize);
@@ -1219,7 +1223,7 @@ llvm::Intrinsic::ID CpAsyncBulkTensorPrefetchOp::getIntrinsicID(int tensorDims,
: CP_ASYNC_BULK_TENSOR_REDUCE_MODE(op, dim, tile)
#define GET_CP_ASYNC_BULK_TENSOR_ID(op, dims, is_im2col) \
- [&]() -> auto { \
+ [&]() -> auto{ \
switch (dims) { \
case 1: \
return CP_ASYNC_BULK_TENSOR_REDUCE_MODE(op, 1, tile); \
@@ -1234,7 +1238,8 @@ llvm::Intrinsic::ID CpAsyncBulkTensorPrefetchOp::getIntrinsicID(int tensorDims,
default: \
llvm_unreachable("Invalid TensorDim in CpAsyncBulkTensorReduceOp."); \
} \
- }()
+ } \
+ ()
llvm::Intrinsic::ID CpAsyncBulkTensorReduceOp::getIntrinsicID(
int tensorDims, NVVM::TMAReduxKind kind, bool isIm2Col) {
@@ -1260,6 +1265,8 @@ llvm::Intrinsic::ID CpAsyncBulkTensorReduceOp::getIntrinsicID(
llvm_unreachable("Invalid Reduction Op for CpAsyncBulkTensorReduceOp");
}
+#define _none
+
#define CVT_F2TF32_ID_IMPL(rnd, relu, sf) \
hasRelu ? llvm::Intrinsic::nvvm_f2tf32_##rnd##relu##sf \
: llvm::Intrinsic::nvvm_f2tf32_##rnd##sf
@@ -1279,7 +1286,7 @@ llvm::Intrinsic::ID CvtFloatToTF32Op::getIntrinsicID(NVVM::FPRoundingMode rnd,
case RndMode::RZ:
return GET_CVT_F2TF32_ID(rz, _relu, _satfinite);
case RndMode::RNA:
- return GET_CVT_F2TF32_ID(rna, , _satfinite);
+ return GET_CVT_F2TF32_ID(rna, _none, _satfinite);
default:
llvm_unreachable("Invalid RoundingMode for CvtFloatToTF32Op");
}
@@ -1290,9 +1297,9 @@ Tcgen05AllocOp::getIntrinsicIDAndArgs(Operation &op,
LLVM::ModuleTranslation &mt,
llvm::SmallVector<llvm::Value *> &args) {
auto curOp = cast<NVVM::Tcgen05AllocOp>(op);
- unsigned AS = llvm::cast<LLVM::LLVMPointerType>(curOp.getAddr().getType())
+ unsigned as = llvm::cast<LLVM::LLVMPointerType>(curOp.getAddr().getType())
.getAddressSpace();
- bool isShared = AS == NVVMMemorySpace::kSharedMemorySpace;
+ bool isShared = as == NVVMMemorySpace::kSharedMemorySpace;
bool is2CTAMode = curOp.getGroup() == Tcgen05GroupKind::CTA_2;
llvm::Intrinsic::ID id;
@@ -1339,14 +1346,15 @@ Tcgen05CommitOp::getIntrinsicIDAndArgs(Operation &op,
LLVM::ModuleTranslation &mt,
llvm::SmallVector<llvm::Value *> &args) {
auto curOp = cast<NVVM::Tcgen05CommitOp>(op);
- unsigned AS = llvm::cast<LLVM::LLVMPointerType>(curOp.getAddr().getType())
+ unsigned as = llvm::cast<LLVM::LLVMPointerType>(curOp.getAddr().getType())
.getAddressSpace();
- bool isShared = AS == NVVMMemorySpace::kSharedMemorySpace;
- bool hasMulticast = curOp.getMulticastMask() ? true : false;
+ bool isShared = as == NVVMMemorySpace::kSharedMemorySpace;
+ bool hasMulticast = static_cast<bool>(curOp.getMulticastMask());
bool is2CTAMode = curOp.getGroup() == Tcgen05GroupKind::CTA_2;
- auto id = is2CTAMode ? GET_TCGEN05_COMMIT_ID(cg2, isShared, hasMulticast)
- : GET_TCGEN05_COMMIT_ID(cg1, isShared, hasMulticast);
+ llvm::Intrinsic::ID id =
+ is2CTAMode ? GET_TCGEN05_COMMIT_ID(cg2, isShared, hasMulticast)
+ : GET_TCGEN05_COMMIT_ID(cg1, isShared, hasMulticast);
// Fill the Intrinsic Args
args.push_back(mt.lookupValue(curOp.getAddr()));
@@ -1364,13 +1372,14 @@ Tcgen05CommitOp::getIntrinsicIDAndArgs(Operation &op,
: TCGEN05_CP_IMPL(shape_mc, src_fmt, _cg1)
#define GET_TCGEN05_CP_ID(shape_mc, src_fmt, is_2cta) \
- [&]() -> auto { \
- if (src_fmt == Tcgen05CpSrcFormat::B6x16_P32) \
+ [&]() -> auto{ \
+ if ((src_fmt) == Tcgen05CpSrcFormat::B6x16_P32) \
return TCGEN05_CP_2CTA(shape_mc, _b6x16_p32, is_2cta); \
- if (src_fmt == Tcgen05CpSrcFormat::B4x16_P64) \
+ if ((src_fmt) == Tcgen05CpSrcFormat::B4x16_P64) \
return TCGEN05_CP_2CTA(shape_mc, _b4x16_p64, is_2cta); \
return TCGEN05_CP_2CTA(shape_mc, , is_2cta); \
- }()
+ } \
+ ()
llvm::Intrinsic::ID Tcgen05CpOp::getIntrinsicID(Operation &op) {
auto curOp = cast<NVVM::Tcgen05CpOp>(op);
@@ -1397,47 +1406,47 @@ llvm::Intrinsic::ID Tcgen05CpOp::getIntrinsicID(Operation &op) {
// Returns the valid vector length for a given shape and vector length, the
// function models the table mentioned in the tcgen05.{ld, st} Op description
-static unsigned isValidVectorLength(NVVM::Tcgen05LdStShape Shape,
- unsigned VecLen) {
- if (Shape == NVVM::Tcgen05LdStShape::SHAPE_16X128B)
- return VecLen >= 2;
- if (Shape == NVVM::Tcgen05LdStShape::SHAPE_16X256B)
- return VecLen >= 4;
+static unsigned isValidVectorLength(NVVM::Tcgen05LdStShape shape,
+ unsigned vecLen) {
+ if (shape == NVVM::Tcgen05LdStShape::SHAPE_16X128B)
+ return vecLen >= 2;
+ if (shape == NVVM::Tcgen05LdStShape::SHAPE_16X256B)
+ return vecLen >= 4;
return true;
}
LogicalResult Tcgen05LdOp::verify() {
- LogicalResult Result = success();
+ LogicalResult result = success();
if (getShape() == NVVM::Tcgen05LdStShape::SHAPE_16X32BX2 && !getOffset())
- Result = emitError("shape 16x32bx2 requires offset argument");
+ result = emitError("shape 16x32bx2 requires offset argument");
- auto ResTy = getRes().getType();
- unsigned ResLen = isa<VectorType>(ResTy)
- ? llvm::cast<VectorType>(ResTy).getNumElements()
+ auto resTy = getRes().getType();
+ unsigned resLen = isa<VectorType>(resTy)
+ ? llvm::cast<VectorType>(resTy).getNumElements()
: 1;
- if (!isValidVectorLength(getShape(), ResLen))
- Result = emitError(llvm::formatv("invalid result type length {0} for shape "
+ if (!isValidVectorLength(getShape(), resLen))
+ result = emitError(llvm::formatv("invalid result type length {0} for shape "
"{1} in tcgen05.ld Op",
- ResLen, stringifyEnum(getShape())));
+ resLen, stringifyEnum(getShape())));
- return Result;
+ return result;
}
LogicalResult Tcgen05StOp::verify() {
- LogicalResult Result = success();
+ LogicalResult result = success();
if (getShape() == NVVM::Tcgen05LdStShape::SHAPE_16X32BX2 && !getOffset())
- Result = emitError("shape 16x32bx2 requires offset argument");
+ result = emitError("shape 16x32bx2 requires offset argument");
- auto ValTy = getVal().getType();
- unsigned ValLen = isa<VectorType>(ValTy)
- ? llvm::cast<VectorType>(ValTy).getNumElements()
+ auto valTy = getVal().getType();
+ unsigned valLen = isa<VectorType>(valTy)
+ ? llvm::cast<VectorType>(valTy).getNumElements()
: 1;
- if (!isValidVectorLength(getShape(), ValLen))
- Result = emitError(llvm::formatv("invalid input length {0} for shape "
+ if (!isValidVectorLength(getShape(), valLen))
+ result = emitError(llvm::formatv("invalid input length {0} for shape "
"{1} in tcgen05.st Op",
- ValLen, stringifyEnum(getShape())));
+ valLen, stringifyEnum(getShape())));
- return Result;
+ return result;
}
/// Infer the result ranges for the NVVM SpecialRangeableRegisterOp that might
@@ -1557,7 +1566,7 @@ NVVMTargetAttr::verify(function_ref<InFlightDiagnostic()> emitError,
return failure();
}
if (files && !llvm::all_of(files, [](::mlir::Attribute attr) {
- return attr && mlir::isa<StringAttr>(attr);
+ return mlir::isa_and_nonnull<StringAttr>(attr);
})) {
emitError() << "All the elements in the `link` array must be strings.";
return failure();
|
|
@llvm/pr-subscribers-mlir Author: Guray Ozen (grypp) ChangesFull diff: https://github.com/llvm/llvm-project/pull/135006.diff 1 Files Affected:
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 09bff6101edd3..cced84ded9592 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -33,6 +33,7 @@
#include "llvm/AsmParser/Parser.h"
#include "llvm/IR/Attributes.h"
#include "llvm/IR/Function.h"
+#include "llvm/IR/IntrinsicsNVPTX.h"
#include "llvm/IR/Type.h"
#include "llvm/Support/Casting.h"
#include "llvm/Support/FormatVariadic.h"
@@ -56,7 +57,7 @@ using namespace NVVM;
// CpAsyncBulkTensorGlobalToSharedClusterOp (TMA Load)
// CpAsyncBulkTensorPrefetchOp (TMA Prefetch)
// CpAsyncBulkTensorReduceOp (TMA Store-Reduce)
-static LogicalResult CpAsyncBulkTensorCommonVerifier(size_t tensorDims,
+static LogicalResult cpAsyncBulkTensorCommonVerifier(size_t tensorDims,
bool isIm2Col,
size_t numIm2ColOffsets,
Location loc) {
@@ -81,7 +82,7 @@ static LogicalResult CpAsyncBulkTensorCommonVerifier(size_t tensorDims,
LogicalResult CpAsyncBulkTensorGlobalToSharedClusterOp::verify() {
size_t numIm2ColOffsets = getIm2colOffsets().size();
bool isIm2Col = numIm2ColOffsets > 0;
- return CpAsyncBulkTensorCommonVerifier(getCoordinates().size(), isIm2Col,
+ return cpAsyncBulkTensorCommonVerifier(getCoordinates().size(), isIm2Col,
numIm2ColOffsets, getLoc());
}
@@ -105,13 +106,13 @@ LogicalResult CpAsyncOp::verify() {
LogicalResult CpAsyncBulkTensorPrefetchOp::verify() {
size_t numIm2ColOffsets = getIm2colOffsets().size();
bool isIm2Col = numIm2ColOffsets > 0;
- return CpAsyncBulkTensorCommonVerifier(getCoordinates().size(), isIm2Col,
+ return cpAsyncBulkTensorCommonVerifier(getCoordinates().size(), isIm2Col,
numIm2ColOffsets, getLoc());
}
LogicalResult CpAsyncBulkTensorReduceOp::verify() {
bool isIm2Col = (getMode() == TMAStoreMode::IM2COL);
- return CpAsyncBulkTensorCommonVerifier(getCoordinates().size(), isIm2Col, 0,
+ return cpAsyncBulkTensorCommonVerifier(getCoordinates().size(), isIm2Col, 0,
getLoc());
}
@@ -183,14 +184,14 @@ static bool isIntegerPtxType(MMATypes type) {
MMATypes MmaOp::accumPtxType() {
std::optional<mlir::NVVM::MMATypes> val = inferOperandMMAType(
- getODSOperands(2).getTypes().front(), /*isAccum=*/true);
+ getODSOperands(2).getTypes().front(), /*isAccumulator=*/true);
assert(val.has_value() && "accumulator PTX type should always be inferrable");
return val.value();
}
MMATypes MmaOp::resultPtxType() {
std::optional<mlir::NVVM::MMATypes> val =
- inferOperandMMAType(getResult().getType(), /*isAccum=*/true);
+ inferOperandMMAType(getResult().getType(), /*isAccumulator=*/true);
assert(val.has_value() && "result PTX type should always be inferrable");
return val.value();
}
@@ -224,7 +225,7 @@ void MmaOp::print(OpAsmPrinter &p) {
}
}
std::optional<MMATypes> inferredType =
- inferOperandMMAType(regTypes.back(), /*isAccum=*/fragIdx >= 2);
+ inferOperandMMAType(regTypes.back(), /*isAccumulator=*/fragIdx >= 2);
if (inferredType)
ignoreAttrNames.push_back(frag.ptxTypeAttr);
}
@@ -243,7 +244,8 @@ void MmaOp::print(OpAsmPrinter &p) {
p.printOptionalAttrDict(this->getOperation()->getAttrs(), ignoreAttrNames);
// Print the types of the operands and result.
- p << " : " << "(";
+ p << " : "
+ << "(";
llvm::interleaveComma(SmallVector<Type, 3>{frags[0].regs[0].getType(),
frags[1].regs[0].getType(),
frags[2].regs[0].getType()},
@@ -363,14 +365,14 @@ ParseResult MmaOp::parse(OpAsmParser &parser, OperationState &result) {
if (failed(parser.resolveOperands(frag.regs, frag.regTypes,
parser.getNameLoc(), result.operands)))
return failure();
- frag.elemtype =
- inferOperandMMAType(frag.regTypes[0], /*isAccum=*/iter.index() < 2);
+ frag.elemtype = inferOperandMMAType(frag.regTypes[0],
+ /*isAccumulator*/ iter.index() < 2);
}
Type resultType;
if (parser.parseArrow() || parser.parseType(resultType))
return failure();
- frags[3].elemtype = inferOperandMMAType(resultType, /*isAccum=*/true);
+ frags[3].elemtype = inferOperandMMAType(resultType, /*isAccumulator*/ true);
std::array<StringRef, 2> names{"multiplicandAPtxType",
"multiplicandBPtxType"};
@@ -992,7 +994,9 @@ std::string NVVM::WgmmaMmaAsyncOp::getPtx() {
ss << "},";
// Need to map read/write registers correctly.
regCnt = (regCnt * 2);
- ss << " $" << (regCnt) << "," << " $" << (regCnt + 1) << "," << " p";
+ ss << " $" << (regCnt) << ","
+ << " $" << (regCnt + 1) << ","
+ << " p";
if (getTypeD() != WGMMATypes::s32) {
ss << ", $" << (regCnt + 3) << ", $" << (regCnt + 4);
}
@@ -1118,9 +1122,9 @@ LogicalResult NVVM::Tcgen05CpOp::verify() {
LogicalResult NVVM::MatchSyncOp::verify() {
if (getKind() == NVVM::MatchSyncKind::all) {
- auto Type = llvm::dyn_cast<LLVM::LLVMStructType>(getType());
- if (!Type || Type.getBody().size() != 2 ||
- !Type.getBody()[0].isInteger(32) || !Type.getBody()[1].isInteger(1)) {
+ auto type = llvm::dyn_cast<LLVM::LLVMStructType>(getType());
+ if (!type || type.getBody().size() != 2 ||
+ !type.getBody()[0].isInteger(32) || !type.getBody()[1].isInteger(1)) {
return emitOpError("match.sync 'all' returns a two element struct with "
"first element as i32 and second element as i1");
}
@@ -1161,7 +1165,7 @@ CpAsyncOp::getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt,
llvm::Intrinsic::ID id;
auto cpAsyncOp = cast<NVVM::CpAsyncOp>(op);
- bool hasCpSize = cpAsyncOp.getCpSize() ? true : false;
+ bool hasCpSize = static_cast<bool>(cpAsyncOp.getCpSize());
switch (cpAsyncOp.getSize()) {
case 4:
id = GET_CP_ASYNC_ID(ca, 4, hasCpSize);
@@ -1219,7 +1223,7 @@ llvm::Intrinsic::ID CpAsyncBulkTensorPrefetchOp::getIntrinsicID(int tensorDims,
: CP_ASYNC_BULK_TENSOR_REDUCE_MODE(op, dim, tile)
#define GET_CP_ASYNC_BULK_TENSOR_ID(op, dims, is_im2col) \
- [&]() -> auto { \
+ [&]() -> auto{ \
switch (dims) { \
case 1: \
return CP_ASYNC_BULK_TENSOR_REDUCE_MODE(op, 1, tile); \
@@ -1234,7 +1238,8 @@ llvm::Intrinsic::ID CpAsyncBulkTensorPrefetchOp::getIntrinsicID(int tensorDims,
default: \
llvm_unreachable("Invalid TensorDim in CpAsyncBulkTensorReduceOp."); \
} \
- }()
+ } \
+ ()
llvm::Intrinsic::ID CpAsyncBulkTensorReduceOp::getIntrinsicID(
int tensorDims, NVVM::TMAReduxKind kind, bool isIm2Col) {
@@ -1260,6 +1265,8 @@ llvm::Intrinsic::ID CpAsyncBulkTensorReduceOp::getIntrinsicID(
llvm_unreachable("Invalid Reduction Op for CpAsyncBulkTensorReduceOp");
}
+#define _none
+
#define CVT_F2TF32_ID_IMPL(rnd, relu, sf) \
hasRelu ? llvm::Intrinsic::nvvm_f2tf32_##rnd##relu##sf \
: llvm::Intrinsic::nvvm_f2tf32_##rnd##sf
@@ -1279,7 +1286,7 @@ llvm::Intrinsic::ID CvtFloatToTF32Op::getIntrinsicID(NVVM::FPRoundingMode rnd,
case RndMode::RZ:
return GET_CVT_F2TF32_ID(rz, _relu, _satfinite);
case RndMode::RNA:
- return GET_CVT_F2TF32_ID(rna, , _satfinite);
+ return GET_CVT_F2TF32_ID(rna, _none, _satfinite);
default:
llvm_unreachable("Invalid RoundingMode for CvtFloatToTF32Op");
}
@@ -1290,9 +1297,9 @@ Tcgen05AllocOp::getIntrinsicIDAndArgs(Operation &op,
LLVM::ModuleTranslation &mt,
llvm::SmallVector<llvm::Value *> &args) {
auto curOp = cast<NVVM::Tcgen05AllocOp>(op);
- unsigned AS = llvm::cast<LLVM::LLVMPointerType>(curOp.getAddr().getType())
+ unsigned as = llvm::cast<LLVM::LLVMPointerType>(curOp.getAddr().getType())
.getAddressSpace();
- bool isShared = AS == NVVMMemorySpace::kSharedMemorySpace;
+ bool isShared = as == NVVMMemorySpace::kSharedMemorySpace;
bool is2CTAMode = curOp.getGroup() == Tcgen05GroupKind::CTA_2;
llvm::Intrinsic::ID id;
@@ -1339,14 +1346,15 @@ Tcgen05CommitOp::getIntrinsicIDAndArgs(Operation &op,
LLVM::ModuleTranslation &mt,
llvm::SmallVector<llvm::Value *> &args) {
auto curOp = cast<NVVM::Tcgen05CommitOp>(op);
- unsigned AS = llvm::cast<LLVM::LLVMPointerType>(curOp.getAddr().getType())
+ unsigned as = llvm::cast<LLVM::LLVMPointerType>(curOp.getAddr().getType())
.getAddressSpace();
- bool isShared = AS == NVVMMemorySpace::kSharedMemorySpace;
- bool hasMulticast = curOp.getMulticastMask() ? true : false;
+ bool isShared = as == NVVMMemorySpace::kSharedMemorySpace;
+ bool hasMulticast = static_cast<bool>(curOp.getMulticastMask());
bool is2CTAMode = curOp.getGroup() == Tcgen05GroupKind::CTA_2;
- auto id = is2CTAMode ? GET_TCGEN05_COMMIT_ID(cg2, isShared, hasMulticast)
- : GET_TCGEN05_COMMIT_ID(cg1, isShared, hasMulticast);
+ llvm::Intrinsic::ID id =
+ is2CTAMode ? GET_TCGEN05_COMMIT_ID(cg2, isShared, hasMulticast)
+ : GET_TCGEN05_COMMIT_ID(cg1, isShared, hasMulticast);
// Fill the Intrinsic Args
args.push_back(mt.lookupValue(curOp.getAddr()));
@@ -1364,13 +1372,14 @@ Tcgen05CommitOp::getIntrinsicIDAndArgs(Operation &op,
: TCGEN05_CP_IMPL(shape_mc, src_fmt, _cg1)
#define GET_TCGEN05_CP_ID(shape_mc, src_fmt, is_2cta) \
- [&]() -> auto { \
- if (src_fmt == Tcgen05CpSrcFormat::B6x16_P32) \
+ [&]() -> auto{ \
+ if ((src_fmt) == Tcgen05CpSrcFormat::B6x16_P32) \
return TCGEN05_CP_2CTA(shape_mc, _b6x16_p32, is_2cta); \
- if (src_fmt == Tcgen05CpSrcFormat::B4x16_P64) \
+ if ((src_fmt) == Tcgen05CpSrcFormat::B4x16_P64) \
return TCGEN05_CP_2CTA(shape_mc, _b4x16_p64, is_2cta); \
return TCGEN05_CP_2CTA(shape_mc, , is_2cta); \
- }()
+ } \
+ ()
llvm::Intrinsic::ID Tcgen05CpOp::getIntrinsicID(Operation &op) {
auto curOp = cast<NVVM::Tcgen05CpOp>(op);
@@ -1397,47 +1406,47 @@ llvm::Intrinsic::ID Tcgen05CpOp::getIntrinsicID(Operation &op) {
// Returns the valid vector length for a given shape and vector length, the
// function models the table mentioned in the tcgen05.{ld, st} Op description
-static unsigned isValidVectorLength(NVVM::Tcgen05LdStShape Shape,
- unsigned VecLen) {
- if (Shape == NVVM::Tcgen05LdStShape::SHAPE_16X128B)
- return VecLen >= 2;
- if (Shape == NVVM::Tcgen05LdStShape::SHAPE_16X256B)
- return VecLen >= 4;
+static unsigned isValidVectorLength(NVVM::Tcgen05LdStShape shape,
+ unsigned vecLen) {
+ if (shape == NVVM::Tcgen05LdStShape::SHAPE_16X128B)
+ return vecLen >= 2;
+ if (shape == NVVM::Tcgen05LdStShape::SHAPE_16X256B)
+ return vecLen >= 4;
return true;
}
LogicalResult Tcgen05LdOp::verify() {
- LogicalResult Result = success();
+ LogicalResult result = success();
if (getShape() == NVVM::Tcgen05LdStShape::SHAPE_16X32BX2 && !getOffset())
- Result = emitError("shape 16x32bx2 requires offset argument");
+ result = emitError("shape 16x32bx2 requires offset argument");
- auto ResTy = getRes().getType();
- unsigned ResLen = isa<VectorType>(ResTy)
- ? llvm::cast<VectorType>(ResTy).getNumElements()
+ auto resTy = getRes().getType();
+ unsigned resLen = isa<VectorType>(resTy)
+ ? llvm::cast<VectorType>(resTy).getNumElements()
: 1;
- if (!isValidVectorLength(getShape(), ResLen))
- Result = emitError(llvm::formatv("invalid result type length {0} for shape "
+ if (!isValidVectorLength(getShape(), resLen))
+ result = emitError(llvm::formatv("invalid result type length {0} for shape "
"{1} in tcgen05.ld Op",
- ResLen, stringifyEnum(getShape())));
+ resLen, stringifyEnum(getShape())));
- return Result;
+ return result;
}
LogicalResult Tcgen05StOp::verify() {
- LogicalResult Result = success();
+ LogicalResult result = success();
if (getShape() == NVVM::Tcgen05LdStShape::SHAPE_16X32BX2 && !getOffset())
- Result = emitError("shape 16x32bx2 requires offset argument");
+ result = emitError("shape 16x32bx2 requires offset argument");
- auto ValTy = getVal().getType();
- unsigned ValLen = isa<VectorType>(ValTy)
- ? llvm::cast<VectorType>(ValTy).getNumElements()
+ auto valTy = getVal().getType();
+ unsigned valLen = isa<VectorType>(valTy)
+ ? llvm::cast<VectorType>(valTy).getNumElements()
: 1;
- if (!isValidVectorLength(getShape(), ValLen))
- Result = emitError(llvm::formatv("invalid input length {0} for shape "
+ if (!isValidVectorLength(getShape(), valLen))
+ result = emitError(llvm::formatv("invalid input length {0} for shape "
"{1} in tcgen05.st Op",
- ValLen, stringifyEnum(getShape())));
+ valLen, stringifyEnum(getShape())));
- return Result;
+ return result;
}
/// Infer the result ranges for the NVVM SpecialRangeableRegisterOp that might
@@ -1557,7 +1566,7 @@ NVVMTargetAttr::verify(function_ref<InFlightDiagnostic()> emitError,
return failure();
}
if (files && !llvm::all_of(files, [](::mlir::Attribute attr) {
- return attr && mlir::isa<StringAttr>(attr);
+ return mlir::isa_and_nonnull<StringAttr>(attr);
})) {
emitError() << "All the elements in the `link` array must be strings.";
return failure();
|
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
durga4github
reviewed
Apr 9, 2025
durga4github
reviewed
Apr 9, 2025
| @@ -1260,6 +1264,8 @@ llvm::Intrinsic::ID CpAsyncBulkTensorReduceOp::getIntrinsicID( | |||
| llvm_unreachable("Invalid Reduction Op for CpAsyncBulkTensorReduceOp"); | |||
| } | |||
|
|
|||
| #define _none | |||
There was a problem hiding this comment.
Is this one also auto-fixed by clang-tidy?
I wonder if we should name it better like macro_null_str (or something similar) to convey that this is used as an empty string in macros..
There was a problem hiding this comment.
It's not auto-generated but clang-tidy complaint about that.
durga4github
reviewed
Apr 9, 2025
| @@ -33,6 +33,7 @@ | |||
| #include "llvm/AsmParser/Parser.h" | |||
| #include "llvm/IR/Attributes.h" | |||
| #include "llvm/IR/Function.h" | |||
| #include "llvm/IR/IntrinsicsNVPTX.h" | |||
There was a problem hiding this comment.
I see this as an addition-only change.
For my understanding,
Please let me know why this header is added as part of this change..
There was a problem hiding this comment.
Ah because there was auto usage without clear type. I used a real type name, so need to add intrinsic.
durga4github
approved these changes
Apr 9, 2025
var-const
pushed a commit
to ldionne/llvm-project
that referenced
this pull request
Apr 17, 2025
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This patch fixes clang-tidy issues in NVVMDialect.cpp file.