Support performance warning

[manman-ren]

This commit adds a performance warning for not selecting MMA v3 for tl.dot on Hopper.
For the added test case, we will get:

test-warning.py:24:18: remark: Warning: can't use MMA V3 for the dot op
    c = tl.dot(a, b)
                 ^
test-warning.py:24:18: note: see current operation: %39 = tt.dot %37, %38, %cst, inputPrecision = tf32 :

[Jokeren]

Where is this flag coming from "-Wno-perf-warnings"?

[Jokeren]

Adds additional warnings obtained when compiling from LLVM to assembly (e.g., register spillage, ptx performance warnings, etc.).

Can you use remark to emit ptx warnings? Maybe not?

[ptillet]

python/tutorials/test-warning.py

We shouldn't have tests as tutorials :) Could you create a unit test that exercises the new codepaths and checks for the output on std::cerr?

[manman-ren]

Thanks for the comments! As described in the summary:

• Where should we add the test case to detect the diagnostics? I can add a mlir test case with -verify-diagnostics and expected-remark. Right now, I am adding a test case under python/tutorial temporarily we will need a way to verify diagnostics are emitted at the right source line.
  @ptillet Yeah I added to tutorial temporarily for discussion, I will move to the unit directory, which will be tested via pytest and I will try to figure out how to check against stdout there. If you have a test case that does similar thing, that will be great!
• Do we want Remarks or Warnings? For llvm backend, remarks may work better.
• Support warning flags or use env variables? It is not clear to me how to support warning flags when building the .py source code.
  @Jokeren If we want to go with warning flags, I need to figure out how to support -Wno-perf-warnings when building a py code. Again if anyone has any pointer, that will be great.

Can you use remark to emit ptx warnings? Maybe not?

You mean warnings from ptxas, not sure how to get them. But we can emit remarks when lowering from llvm to ptx.

CC @joker-eph: if you have any suggestion to some of the questions, that will be great, Thanks!

[Jokeren]

You mean warnings from ptxas, not sure how to get them. But we can emit remarks when lowering from llvm to ptx.

My worry was that some warnings might be emitted from ptx to sass. I'm OK with either marks or warnings. I think marks actually seem more informative.

[joker-eph]

You mean warnings from ptxas, not sure how to get them.

You could capture the output of ptxas and regex/pattern match the output to catch these and report them as remarks?

[ptillet]

Where should we add the test case to detect the diagnostics

test/unit/warnings.py sounds like a good place

I will try to figure out how to check against stdout there

Just to be clear, this should check against stderr

[manman-ren]

This test case is not working currently. I haven't figured out how to capture the stderr in a string yet. Tried to redirect stderr, but it may be captured by pytest? The message shows:
--------------------------------------------------------------------------- Captured stderr call ---------------------------------------------------------------------------
python/test/unit/test_warning.py:21:18: remark: Warning: can't use MMA V3 for the dot op
c = tl.dot(a, b)
^
python/test/unit/test_warning.py:21:18: note: see current operation: %46 = tt.dot ...

Will clean up this once it is solved. @htyu Any suggestion here? Thanks!

[manman-ren]

@ptillet About using warning flags -Wno-perf-warnings or -Werror, I feel it is not supported in the python workflow "python test.py", right? So is the suggestion about supplying the flags for triton-opt and still use env variables for the python workflow?
CC @joker-eph in case Mehdi has some comments.

[htyu]

nit: remove the comment?

[htyu]

Actually I think instead we may want to add an enable_remarks API similar to enable_debug:

triton/python/src/ir.cc

Line 1552 in 74ad278

.def("enable_debug",

[manman-ren]

@htyu Thanks for the review! Sorry for the delay. I tried to wrap things in enable_remarks:
as part of Context

   py::class_<MLIRContext>(m, "context", py::module_local())
       .def(py::init<>())
+      .def("enable_remark",
+           [](MLIRContext &self) {
+             auto srcMgr = llvm::SourceMgr();
+             auto diag = SourceMgrDiagnosticHandler(srcMgr, &self);
+             self.printOpOnDiagnostic(true);
+           })

or as part of PassManager

   py::class_<PassManager>(m, "pass_manager", py::module_local())
       .def(py::init<MLIRContext *>())
+      .def("enable_remark",
+           [](PassManager &self) {
+             auto *context = self.getContext();
+             auto srcMgr = llvm::SourceMgr();
+             auto diag = SourceMgrDiagnosticHandler(srcMgr, context);
+             context->printOpOnDiagnostic(true);
+           })

Then call from make_ttgir:
with either
mod.context.enable_remark()
or
pm.enable_remark()

None of these worked. I think it is due to the liveness of SourceMgr and SourceMgrDiagnosticHandler.

[htyu]

Thanks for giving it a shot. Currently changes look good to me.

[manman-ren]

The tests are failing. I need to only enable the test for H100 with the env variable.

[Jokeren]

Can you document in README?