Skip to content

Commit 79467af

Browse files
pthomadakispthomadakis
committed
Fixed warnings in COMET
Suppressed warnings coming from Triton
1 parent 1a091cf commit 79467af

File tree

6 files changed

+37
-30
lines changed

6 files changed

+37
-30
lines changed

‎lib/Conversion/GpuToTriton/GpuToTritonConversion.cpp

+2-4
Original file line numberDiff line numberDiff line change
@@ -62,7 +62,7 @@ mlir::comet::GpuConversionTarget::GpuConversionTarget(
6262

6363
for(auto oper: op->getOperands())
6464
{
65-
if(auto affineOp = llvm::isa_and_nonnull<affine::AffineApplyOp>(oper.getDefiningOp()))
65+
if(llvm::isa_and_nonnull<affine::AffineApplyOp>(oper.getDefiningOp()))
6666
{
6767
return false;
6868
}
@@ -86,10 +86,8 @@ mlir::comet::GpuConversionTarget::GpuConversionTarget(
8686
{
8787
for(auto oper : op.getCondition().getDefiningOp()->getOperands())
8888
{
89-
if(auto affineOp = llvm::isa_and_nonnull<affine::AffineApplyOp>(oper.getDefiningOp()))
89+
if(!llvm::isa_and_nonnull<affine::AffineApplyOp>(oper.getDefiningOp()))
9090
{
91-
}
92-
else {
9391
return false;
9492
}
9593
}

‎lib/Conversion/GpuToTriton/GpuToTritonPass.cpp

+12-12
Original file line numberDiff line numberDiff line change
@@ -182,8 +182,8 @@ void handleBinaryExpr(Operation* op, std::map<const void*, mlir::Value>& m, Affi
182182
auto binOp = llvm::cast<mlir::AffineBinaryOpExpr>(exp);
183183
auto lhs = m[binOp.getLHS().getAsOpaquePointer()];
184184
auto rhs = m[binOp.getRHS().getAsOpaquePointer()];
185-
auto lhsT = lhs.getType();
186-
auto rhsT = rhs.getType();
185+
// auto lhsT = lhs.getType();
186+
// auto rhsT = rhs.getType();
187187

188188
makeShapesEqual(op, lhs, rhs, rewriter);
189189

@@ -238,9 +238,9 @@ LogicalResult convertMemoryOp(Operation* op, ConversionPatternRewriter &rewriter
238238
if(op->getOperand(i).getDefiningOp() && isa<affine::AffineApplyOp>(op->getOperand(i).getDefiningOp()))
239239
{
240240
affine::AffineApplyOp affineOp = cast<affine::AffineApplyOp>(op->getOperand(i).getDefiningOp());
241-
int pidXIndex = -1;
242-
int pidYIndex = -1;
243-
int dim = -1;
241+
// int pidXIndex = -1;
242+
// int pidYIndex = -1;
243+
// int dim = -1;
244244
Value guardX = NULL, guardY = NULL, guardR = NULL;
245245
Value guardXExpr = NULL, guardYExpr = NULL, guardRExpr = NULL;
246246
std::map<const void*, mlir::Value> map, mapGuard;
@@ -411,7 +411,7 @@ LogicalResult convertMemoryOp(Operation* op, ConversionPatternRewriter &rewriter
411411
}
412412
else
413413
{
414-
auto binOp = llvm::cast<mlir::AffineBinaryOpExpr>(exp);
414+
// auto binOp = llvm::cast<mlir::AffineBinaryOpExpr>(exp);
415415
// for(auto m: {map /*mapGuard*/})
416416
handleBinaryExpr(op, map, exp, rewriter);
417417
handleBinaryExpr(op, mapGuard, exp, rewriter);
@@ -713,7 +713,7 @@ void convertBlockArgTypes(mlir::triton::FuncOp func, PatternRewriter& rewriter)
713713

714714
void isAncestor(mlir::Operation* pop, mlir::Value val, std::vector<std::vector<mlir::Operation*>>& chain, std::vector<mlir::Operation*> op_path)
715715
{
716-
bool res = false;
716+
// bool res = false;
717717
op_path.push_back(pop);
718718
// std::vector<mlir::Operation*> path;
719719
for(auto op: pop->getOperands())
@@ -762,7 +762,7 @@ class RewriteReduction : public OpConversionPattern<mlir::scf::ForOp> {
762762
std::vector<std::vector<mlir::Operation*>> local_ops_chain;
763763
std::vector<mlir::Operation*> ops_chain;
764764
mlir::Value offset = op.getPtr().getDefiningOp()->getOperand(1);
765-
mlir::Value blockPtr = op.getPtr().getDefiningOp()->getOperand(0);
765+
// mlir::Value blockPtr = op.getPtr().getDefiningOp()->getOperand(0);
766766
if(offset == forOp.getLoopRegions()[0]->getArgument(0))
767767
{
768768
// COMET_ERRS << " SAME:\n";
@@ -927,7 +927,7 @@ class RewriteReduction : public OpConversionPattern<mlir::scf::ForOp> {
927927
// COMET_ERRS << "NOT RUNNING\n";
928928
}
929929
rewriter.setInsertionPoint(forOp);
930-
if(basePtrs.size() > 0 && basePtrs.size() == numLoadOps)
930+
if(basePtrs.size() > 0 && numLoadOps > 0 && basePtrs.size() == (size_t)numLoadOps)
931931
{
932932
rewriter.replaceAllUsesWith(forOp.getLoopRegions()[0]->getArgument(0), rewriter.create<mlir::arith::ConstantIntOp>(forOp->getLoc(), 0, 32));
933933
}
@@ -1265,7 +1265,7 @@ class GpuFuncOpToTritonFuncOp : public OpConversionPattern<mlir::gpu::GPUFuncOp>
12651265
rewriter.setInsertionPointToStart(yLoop.getBody());
12661266
mlir::scf::ForOp xLoop = rewriter.create<mlir::scf::ForOp>(TTFunc->getLoc(), zero, boundX, indexNumProgramsX);
12671267
xLoop->setAttr("programs_loop_x", rewriter.getUnitAttr());
1268-
bool foundYieldOp = false;
1268+
// bool foundYieldOp = false;
12691269

12701270
auto offsetY = yLoop.getBody()->getArgument(0);
12711271
auto offsetX = xLoop.getBody()->getArgument(0);
@@ -1470,7 +1470,7 @@ class ConvertGpuToTriton
14701470

14711471
auto funcOp = *op.getOps<mlir::func::FuncOp>().begin();
14721472
auto launchOps = funcOp.getOps<mlir::gpu::LaunchFuncOp>();
1473-
bool changed = false;
1473+
// bool changed = false;
14741474
OpBuilder builder = OpBuilder(op);
14751475
auto gpuModules = op.getOps<mlir::gpu::GPUModuleOp>();
14761476

@@ -1481,7 +1481,7 @@ class ConvertGpuToTriton
14811481
{
14821482
continue;
14831483
}
1484-
changed = true;
1484+
// changed = true;
14851485
auto kernel_name = launchOp.getKernelName();
14861486
auto kernel_mod_name = launchOp.getKernelModuleName();
14871487
// COMET_ERRS <<kernel_mod_name <<"::"<< kernel_name << "\n";

‎lib/Dialect/IndexTree/Transforms/Fusion.cpp

+3-3
Original file line numberDiff line numberDiff line change
@@ -140,7 +140,7 @@ namespace
140140
}; /// class IndexTreeKernelFusionPass
141141
} /// End anonymous namespace
142142

143-
void IndexTreeKernelFusionPass::test(mlir::func::FuncOp &funcop)
143+
[[maybe_unused]] void IndexTreeKernelFusionPass::test(mlir::func::FuncOp &funcop)
144144
{
145145
int level = 0;
146146
funcop.walk([&](mlir::Operation *op)
@@ -351,7 +351,7 @@ mlir::Value IndexTreeKernelFusionPass::createNewTensorDecl(
351351
return new_dense_tensor_decl;
352352
}
353353

354-
void IndexTreeKernelFusionPass::createNewTensor(
354+
[[maybe_unused]] void IndexTreeKernelFusionPass::createNewTensor(
355355
const mlir::Value &old_tensor_alloc,
356356
const mlir::Value &old_tensor_load,
357357
uint32_t rank_base,
@@ -613,7 +613,7 @@ void IndexTreeKernelFusionPass::replaceOldTensorFillOp(
613613
}
614614
}
615615

616-
void IndexTreeKernelFusionPass::replaceOldLinalgFillOp(
616+
[[maybe_unused]] void IndexTreeKernelFusionPass::replaceOldLinalgFillOp(
617617
mlir::Value &old_tensor_alloc,
618618
mlir::Value &old_tensor_load,
619619
mlir::Value &new_tensor_load)

‎lib/Dialect/TensorAlgebra/Transforms/TensorDeclLowering.cpp

+2-2
Original file line numberDiff line numberDiff line change
@@ -477,8 +477,8 @@ namespace
477477
comet_debug() << "lowerSparseOutputTensorDec::TempSparseOutputTensorDeclOp lowering\n";
478478
}
479479

480-
assert(isa<SparseOutputTensorDeclOp>(op) || isa<TempSparseOutputTensorDeclOp>(op) &&
481-
"Op should be either SparseOutputTensorDeclOp or TempSparseOutputTensorDeclOp");
480+
assert(isa<SparseOutputTensorDeclOp>(op) || (isa<TempSparseOutputTensorDeclOp>(op) &&
481+
"Op should be either SparseOutputTensorDeclOp or TempSparseOutputTensorDeclOp"));
482482

483483
comet_vdump(op);
484484
auto loc = op.getLoc();

‎lib/ExecutionEngine/SparseUtils.cpp

+6-6
Original file line numberDiff line numberDiff line change
@@ -209,7 +209,7 @@ struct CooMatrix
209209
else if (nparsed == -1)
210210
{
211211
/// Problem description
212-
nparsed = sscanf(line, "%lld %lld %lld", &num_rows, &num_cols, &num_nonzeros);
212+
nparsed = sscanf(line, "%" PRIu64 " %" PRIu64 " %" PRIu64, &num_rows, &num_cols, &num_nonzeros);
213213
if ((!array) && (nparsed == 3))
214214
{
215215
if (symmetric)
@@ -237,7 +237,7 @@ struct CooMatrix
237237
/// Edge
238238
if (current_nz >= num_nonzeros)
239239
{
240-
fprintf(stderr, "Error parsing MARKET matrix: encountered more than %lld num_nonzeros\n", num_nonzeros);
240+
fprintf(stderr, "Error parsing MARKET matrix: encountered more than %" PRIu64 " num_nonzeros\n", num_nonzeros);
241241
exit(1);
242242
}
243243

@@ -249,7 +249,7 @@ struct CooMatrix
249249
{
250250
if (sscanf(line, "%lf", &tempVal) != 1) /// using tempVal instead of templated T val to avoid warning
251251
{
252-
fprintf(stderr, "Error parsing MARKET matrix: badly formed current_nz: '%s' at edge %lld\n", line, current_nz);
252+
fprintf(stderr, "Error parsing MARKET matrix: badly formed current_nz: '%s' at edge %" PRIu64 "\n", line, current_nz);
253253
exit(1);
254254
}
255255
val = (T)tempVal;
@@ -283,7 +283,7 @@ struct CooMatrix
283283
row = strtol(l, &t, 0);
284284
if (t == l)
285285
{
286-
fprintf(stderr, "Error parsing MARKET matrix: badly formed row at edge %lld\n", current_nz);
286+
fprintf(stderr, "Error parsing MARKET matrix: badly formed row at edge %" PRIu64 "\n", current_nz);
287287
exit(1);
288288
}
289289
l = t;
@@ -292,7 +292,7 @@ struct CooMatrix
292292
col = strtol(l, &t, 0);
293293
if (t == l)
294294
{
295-
fprintf(stderr, "Error parsing MARKET matrix: badly formed col at edge %lld\n", current_nz);
295+
fprintf(stderr, "Error parsing MARKET matrix: badly formed col at edge %" PRIu64 "\n", current_nz);
296296
exit(1);
297297
}
298298
l = t;
@@ -1028,7 +1028,7 @@ struct Coo3DTensor
10281028
if (nparsed == -1)
10291029
{
10301030
/// Problem description
1031-
nparsed = sscanf(line, "%lld %lld %lld %lld", &num_index_i, &num_index_j, &num_index_k, &num_nonzeros);
1031+
nparsed = sscanf(line, "%" PRIu64 " %" PRIu64 " %" PRIu64 " %" PRIu64, &num_index_i, &num_index_j, &num_index_k, &num_nonzeros);
10321032

10331033
if (nparsed == 4)
10341034
{

‎triton.patch

+12-3
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,16 @@
11
diff --git a/CMakeLists.txt b/CMakeLists.txt
2-
index 309855e0c..b2530e31c 100644
2+
index 309855e..4bc4db7 100644
33
--- a/CMakeLists.txt
44
+++ b/CMakeLists.txt
5+
@@ -20,7 +20,7 @@ if(NOT WIN32)
6+
endif()
7+
8+
9+
-
10+
+set_directory_properties(PROPERTIES COMPILE_OPTIONS "-w")
11+
# Options
12+
option(TRITON_BUILD_TUTORIALS "Build C++ Triton tutorials" ON)
13+
option(TRITON_BUILD_PYTHON_MODULE "Build Python Triton bindings" OFF)
514
@@ -103,7 +103,7 @@ endfunction()
615

716

@@ -22,7 +31,7 @@ index 309855e0c..b2530e31c 100644
2231
+# add_subdirectory(test)
2332
+# add_subdirectory(unittest)
2433
diff --git a/include/triton/Dialect/NVGPU/IR/CMakeLists.txt b/include/triton/Dialect/NVGPU/IR/CMakeLists.txt
25-
index f8932cdc4..37fb66417 100644
34+
index f8932cd..37fb664 100644
2635
--- a/include/triton/Dialect/NVGPU/IR/CMakeLists.txt
2736
+++ b/include/triton/Dialect/NVGPU/IR/CMakeLists.txt
2837
@@ -1,8 +1,8 @@
@@ -37,7 +46,7 @@ index f8932cdc4..37fb66417 100644
3746
mlir_tablegen(Ops.h.inc -gen-op-decls)
3847
mlir_tablegen(Ops.cpp.inc -gen-op-defs)
3948
diff --git a/include/triton/Dialect/NVGPU/IR/NVGPUDialect.td b/include/triton/Dialect/NVGPU/IR/NVGPUDialect.td
40-
index 6978173d4..3b76e10c7 100644
49+
index 6978173..3b76e10 100644
4150
--- a/include/triton/Dialect/NVGPU/IR/NVGPUDialect.td
4251
+++ b/include/triton/Dialect/NVGPU/IR/NVGPUDialect.td
4352
@@ -25,7 +25,7 @@

0 commit comments

Comments
 (0)