Emit remarks for SWP and vectorization failures by zengwu · Pull Request #4350 · triton-lang/triton

zengwu · 2024-07-18T05:49:27Z

Dump warning if SWP fails in the inner loop and dump option is enabled in the CL.

ThomasRaoux

Please check with @manman-ren to see if there is a more elegant way to do this.

ThomasRaoux · 2024-07-18T14:48:20Z

      int loopNumStages = getNumStagesOrDefault(forOp);
      bool pipelined = pipelineLoop(forOp, loopNumStages);
+      if (DumpSWPFailure && !pipelined) {
+        forOp->emitRemark("SWP failes in inner most loop");


@manman-ren was working on a framework to report performance warning. I don't think we should be adding command line option for every potential performance problem as this wouldn't scale.

Yeah we don't need a separate flag for each warning. Warnings are guarded via MLIR_ENABLE_REMARK

ThomasRaoux · 2024-07-18T14:48:34Z

+// RUN: triton-opt %s -split-input-file -tritongpu-pipeline=num-stages=3 -dump-swp-failure | FileCheck %s
+
+// CHECK-LABEL: @dont_pipeline_128x1
+// CHECK-NOT: local_load{{.*}}128x1


this doesn't seem to test the code added?

Tests for warnings are added here: python/test/unit/test_perf_warning.py. Can you see if you can add a case there?

manman-ren · 2024-07-19T15:59:11Z

+    if is_cuda():
+        capability = torch.cuda.get_device_capability()
+        if capability[0] < 9:
+            pytest.skip("Requires sm >= 90 to run")


I am not sure if this warning should be specific to H100. If it applies to A100, we should remove the check.

manman-ren · 2024-07-19T16:02:10Z

      int loopNumStages = getNumStagesOrDefault(forOp);
      bool pipelined = pipelineLoop(forOp, loopNumStages);
+      if (!pipelined) {
+        forOp->emitRemark("Warning: loop is not pipelined");


I wonder how hard it is to provide a reason for this. Should we check to see if the loop has num_stages > 1? I am not sure if this gets executed when num_stages is 1 and if we will emit the warning always when num_stages is 1.

manman-ren · 2024-07-22T22:49:13Z

-    if (loops.empty())
+    if (loops.empty()) {
+      auto op = getOperation();
+      op->emitRemark() << "Warning: SWP fails. There is no loop with num_stages greater than 1";


If there is no loop with num_stages greater than 1, I guess we don't need to email a warning since SWP is not requested. For generating test cases, maybe we can go through TritonBench or Triton lit tests to see if a remark will be triggered.

Can you remove this change? We don't want to emit a warning for loops that didn't request SWP.

manman-ren

Thanks for working on this!

manman-ren · 2024-08-01T18:02:58Z

-               forOp->getAttr(mlir::triton::kNumStagesAttrName))
-        .getInt();
+                         forOp->getAttr(mlir::triton::kNumStagesAttrName))
+                         .getInt();


Seems to be unrelated format change.

manman-ren · 2024-08-01T18:04:27Z

-    if (loops.empty())
+    if (loops.empty()) {
+      auto op = getOperation();
+      op->emitRemark() << "Warning: SWP fails. There is no loop with num_stages greater than 1";


Can you remove this change? We don't want to emit a warning for loops that didn't request SWP.

Jokeren · 2024-08-01T18:19:23Z

                     return !def;
-                   }))
+                   })) {
+    forOp->emitRemark() << "Warning: SWP fails due to loop distance is greater than 1";


Which operands have a distance greater than 1?

Jokeren · 2024-08-01T18:21:15Z

+        if os.environ.get("MLIR_ENABLE_REMARK", "0") == "1":
+            srcMgr = llvm.source_mgr()
+            diag = ir.source_mgr_diag(srcMgr, mod.context)
+            mod.context.printOpOnDiagnostic(True)


Is it possible to emit all diagnostic messages to a file? At least have such an option even if the default is stdout or stderr.

@manman-ren We may also want to design a warning format that could be parsed by proton

@Jokeren Yes, we need a warning format. Do you have any suggestion? Do we want to mention the pass name? Or add a severity level?

Is it possible to emit all diagnostic messages to a file? At least have such an option even if the default is stdout or stderr.

I think it is possible. Do we want an env variable to supply the file path?

I don't have a format in mind now, but I think you could use a function to concatenate the warning with a dedicated separator.

For instance, emitWarning(ir, passName/fileName, map<operands/results, problem>)

Yeah we need to talk more about writing to a file and the format of the diagnostics. Let's do that outside of this PR. Zeng and I looked at writing to a file offline: SourceMgrDiagnosticHandler can take a raw_ostream (the default is stderr).

Jokeren · 2024-08-01T18:22:24Z

    }

+    if (vec == 1 && numElems > 1)
+      op->emitRemark() << "Warning: vectorization fails vec = "


It will be more useful to emit the vec of both the pointer and the mask

Good catch! we can show getVectorSize(ptr) and the mask_alignment.

Jokeren · 2024-08-01T18:22:40Z

    }

+    if (vec == 1 && elemsPerThread > 1)
+      op->emitRemark() << "Warning: vectorization fails vec = "


manman-ren · 2024-08-12T17:18:51Z

 import pytest
 import torch
-
+import tempfile


Is this needed?

manman-ren · 2024-08-12T17:23:24Z

+        if os.environ.get("MLIR_ENABLE_REMARK", "0") == "1":
+            srcMgr = llvm.source_mgr()
+            diag = ir.source_mgr_diag(srcMgr, mod.context)
+            mod.context.printOpOnDiagnostic(True)


Yeah we need to talk more about writing to a file and the format of the diagnostics. Let's do that outside of this PR. Zeng and I looked at writing to a file offline: SourceMgrDiagnosticHandler can take a raw_ostream (the default is stderr).

manman-ren · 2024-08-12T17:24:52Z

    }

+    if (vec == 1 && numElems > 1) {
+      auto maskStr = !llMask ? "no mask" : std::to_string(getMaskAlignment(mask));


It is probably better to have an integer for "no mask" case? Also can we add the value of vec prior to the "if (llMask)" statement in the remark?

manman-ren · 2024-08-12T17:25:00Z

    }

+    if (vec == 1 && elemsPerThread > 1) {
+      auto maskStr = !llMask ? "no mask" : std::to_string(getMaskAlignment(op.getMask()));


manman-ren · 2024-08-12T21:06:04Z

Looks good to me! @adamomainz @Jokeren Let's chat more about the formats and how to write to a file.

manman-ren

waiting for tests to be green.

manman-ren · 2024-08-16T23:13:28Z

@ThomasRaoux I think it needs your approval :]
Merging can be performed automatically once the requested changes are addressed.

ThomasRaoux

Sorry for the late response.
The mechanism makes sense to me overall. I don't think the SWP warning is meaningful. The vectorization warnings look fine although may be verbose but that might be fine for now.

ThomasRaoux · 2024-08-22T01:04:50Z

          })
-          .wasInterrupted())
+          .wasInterrupted()) {
+    forOp->emitRemark() << "Warning: SWP fails on the outer loop";


that might be a bit of a noisy comment as pipelining is known to be an inner loop transformation. Do we really need to tell user that the outer loop didn't pipeline?

@ThomasRaoux, the PR is updated with your comments resolved, could you take a look? Thanks

ThomasRaoux

LGTM

Dump warning if SWP fails in the inner loop and dump option is enabled in the CL. --------- Co-authored-by: Zeng Wu <zengwu@fb.com>

Emit SWP failure warning if enabled in the commandline

f11bf68

zengwu requested a review from ptillet as a code owner July 18, 2024 05:49

ThomasRaoux requested changes Jul 18, 2024

View reviewed changes

manman-ren marked this pull request as draft July 18, 2024 16:00

zengwu force-pushed the main branch from df4a8b9 to a80e46e Compare July 19, 2024 06:07

zengwu changed the title ~~(WIP)Emit SWP failure warning if enabled in the commandline~~ (WIP)Emit SWP failure warning Jul 19, 2024

manman-ren reviewed Jul 19, 2024

View reviewed changes

Emit perf warning for the failure of SWP, vectorization

4b075ec

zengwu force-pushed the main branch from a80e46e to 4b075ec Compare July 22, 2024 06:25

manman-ren reviewed Jul 22, 2024

View reviewed changes

zengwu and others added 2 commits July 30, 2024 20:42

Merge branch 'triton-lang:main' into main

0f35316

Emit perf warning for the failure of SWP, vectorization

454693d

zengwu changed the title ~~(WIP)Emit SWP failure warning~~ Emit remarks for SWP and vectorization failures Jul 31, 2024

manman-ren reviewed Aug 1, 2024

View reviewed changes

Jokeren reviewed Aug 1, 2024

View reviewed changes

Merge branch 'triton-lang:main' into main

778fe5c

zengwu force-pushed the main branch from 644eb37 to 4b09740 Compare August 2, 2024 18:03

Emit perf warning for the failure of SWP, vectorization

76b8c3a

zengwu force-pushed the main branch from 4b09740 to 76b8c3a Compare August 2, 2024 18:10

manman-ren reviewed Aug 12, 2024

View reviewed changes

Merge branch 'triton-lang:main' into main

ae44ad5

manman-ren marked this pull request as ready for review August 12, 2024 20:59

zengwu force-pushed the main branch 6 times, most recently from a643b41 to 3c138c9 Compare August 14, 2024 04:52

Emit perf warning for the failure of SWP, vectorization

9350838

zengwu force-pushed the main branch from 3c138c9 to 9350838 Compare August 15, 2024 23:37

zengwu added 2 commits August 15, 2024 20:27

Merge branch 'main' into main

b9ed2b0

Merge branch 'triton-lang:main' into main

9004795

manman-ren approved these changes Aug 16, 2024

View reviewed changes

manman-ren requested a review from ThomasRaoux August 21, 2024 23:12

ThomasRaoux reviewed Aug 22, 2024

View reviewed changes

zengwu and others added 4 commits August 22, 2024 19:06

Merge branch 'triton-lang:main' into main

8620c7f

Merge branch 'triton-lang:main' into main

8f2d584

Emit perf warning for the failure of SWP, vectorization

85a2f93

Merge branch 't1'

5c42c19

zengwu force-pushed the main branch from 3793327 to 5c42c19 Compare August 27, 2024 16:24

ThomasRaoux approved these changes Aug 27, 2024

View reviewed changes

manman-ren merged commit cf696d4 into triton-lang:main Aug 27, 2024

jlebar mentioned this pull request Sep 3, 2024

Build LLVMAarch64CodeGen if CMAKE_OSX_ARCHITECTURES is arm64. #4637

Merged

sfzhu93 mentioned this pull request Nov 25, 2024

[WIP][Performance Remarks] Support MLIR_ENABLE_REMARK for command line tools #5250

Closed

7 tasks

Conversation

zengwu commented Jul 18, 2024

Uh oh!

ThomasRaoux left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

manman-ren left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Jokeren Aug 1, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

manman-ren commented Aug 12, 2024

Uh oh!

manman-ren left a comment

Choose a reason for hiding this comment

Uh oh!

manman-ren commented Aug 16, 2024

Uh oh!

ThomasRaoux left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

ThomasRaoux left a comment

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Jokeren Aug 1, 2024 •

edited

Loading