Skip to content
Draft
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -180,8 +180,8 @@ endif()

# Customized release build type with assertions: TritonRelBuildWithAsserts
if(NOT MSVC)
set(CMAKE_C_FLAGS_TRITONRELBUILDWITHASSERTS "-O2 -g")
set(CMAKE_CXX_FLAGS_TRITONRELBUILDWITHASSERTS "-O2 -g")
set(CMAKE_C_FLAGS_TRITONRELBUILDWITHASSERTS "-O2 -g -Wunused")
set(CMAKE_CXX_FLAGS_TRITONRELBUILDWITHASSERTS "-O2 -g -Wunused")
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why only for this build type?

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is mostly so we don't break release builds with assertions disabled, since many variables may become unused if they are only used inside asserts.

set(CMAKE_C_FLAGS_TRITONBUILDWITHO1 "-O1")
set(CMAKE_CXX_FLAGS_TRITONBUILDWITHO1 "-O1")
else()
Expand Down
4 changes: 2 additions & 2 deletions lib/Tools/GenericSwizzling.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -29,8 +29,8 @@ using namespace mlir::triton;

namespace {

void printBasis(const llvm::SmallVector<int32_t> &basis,
const std::string &name) {
[[maybe_unused]] void printBasis(const llvm::SmallVector<int32_t> &basis,
const std::string &name) {
llvm::errs() << name << ": ";
for (int32_t b : basis)
llvm::errs() << b << " ";
Expand Down
2 changes: 1 addition & 1 deletion lib/Tools/LinearLayout.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -54,7 +54,7 @@ BasesT makeBasesMap(
}

// Dump the matrix to stderr in a human-readable format for debugging.
void dumpMatrix(uint64_t *m, int numRows, int numCols) {
[[maybe_unused]] void dumpMatrix(uint64_t *m, int numRows, int numCols) {
assert(numCols <= 64);
for (int r = 0; r < numRows; r++) {
llvm::errs() << "0b";
Expand Down
4 changes: 3 additions & 1 deletion python/src/gluon_ir.cc
Original file line number Diff line number Diff line change
Expand Up @@ -65,7 +65,9 @@ template <typename AttrOrType>
constexpr auto hasVerifier(AttrOrType t) -> decltype(t.verifyInvariants, true) {
return true;
}
constexpr auto hasVerifier(...) { return false; }
template <typename... ArgTs> constexpr auto hasVerifier(ArgTs...) {
return false;
}

// Print a diagnostic without its location. The frontend will attach the AST
// location to the error message.
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -30,8 +30,7 @@ struct ScaledDotOpConversion
: public ConvertOpToLLVMPattern<triton::DotScaledOp> {
using ConvertOpToLLVMPattern<triton::DotScaledOp>::ConvertOpToLLVMPattern;

ScaledDotOpConversion(LLVMTypeConverter &converter, int,
PatternBenefit benefit)
ScaledDotOpConversion(LLVMTypeConverter &converter, PatternBenefit benefit)
: ConvertOpToLLVMPattern<triton::DotScaledOp>(converter, benefit) {}

LogicalResult
Expand Down Expand Up @@ -163,6 +162,5 @@ void mlir::triton::NVIDIA::populateDotOpToLLVMPatterns(
patterns.add<DotOpConversion>(typeConverter, computeCapability, benefit);
patterns.add<WarpGroupDotOpConversion>(typeConverter, benefit);
patterns.add<WarpGroupDotWaitOpConversion>(typeConverter, benefit);
patterns.add<ScaledDotOpConversion>(typeConverter, computeCapability,
benefit);
patterns.add<ScaledDotOpConversion>(typeConverter, benefit);
}
Original file line number Diff line number Diff line change
Expand Up @@ -329,10 +329,10 @@ class DotOpMmaSmemLoader : public DotOpMmaMemLoader {
}
};

static Value getOffsetedBase(Value v, gpu::MemDescType memDescTy,
const TypeConverter *typeConverter,
ConversionPatternRewriter &rewriter,
Location loc) {
static inline Value getOffsetedBase(Value v, gpu::MemDescType memDescTy,
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

static inline is a bit of an oxymoron. Should just be inline.

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

static here is not completely redundant. It still gives the function internal linkage. In practice, this likely only really makes a difference when Triton plugins are enabled, since it will affect whether the function gets exported by libtriton.so.

const TypeConverter *typeConverter,
ConversionPatternRewriter &rewriter,
Location loc) {
TritonLLVMOpBuilder tb(loc, rewriter);
auto llvmElemTy = typeConverter->convertType(memDescTy.getElementType());
auto smemObj =
Expand Down
Loading