Skip to content
Open
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
19 changes: 15 additions & 4 deletions python/src/llvm.cc
Original file line number Diff line number Diff line change
Expand Up @@ -328,17 +328,28 @@ void init_triton_llvm(py::module &&m) {
});
}

PassInstrumentationCallbacks *instrCbPtr = nullptr;
PassInstrumentationCallbacks passInstrCb;
StandardInstrumentations standardInstr(mod->getContext(),
/*DebugLogging*/ true);
Copy link
Collaborator

Choose a reason for hiding this comment

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

why changing this? Does it improve compile time? If so do you have numbers?

Copy link
Author

Choose a reason for hiding this comment

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

Thanks for taking a look.

This overall change is expected to improve the logging of compile-time: mechanically, without this change LLVM_ENABLE_TIMING=1 ignores the entire pre-lower optimization pipeline. No change here is expected to improve compile-time itself, and I don't believe there's a substantive risk of regression, so I did not measure this. Maybe my PR message/title is confusing and made it sound like I expected this to improve compile time :), please let me know if I can express it better.

To the code-change highlighted: The third highlighted line (333) must be moved in order to let the second parameter /*DebuggingLogging*/ true be parameterized by which debugging flags are set. I parameterized it because the debugging information is largely orthogonal from timing, but I didn't want to change the behavior of the previously-used-flag. For the other lines, I suppose it's largely stylistic, so I'd be happy to change it if there's a more "Triton-y" way you see.

Do note, however, it seems required that standardInstr be constructed before the writes to llvm::TimePasses*---this is because those globals are read by the constructor.

- llvm/lib/Passes/StandardInstrumentations.cpp -> invokes the default constructor of its field `TimePassesHandler TimePasses;`
- llvm/lib/IR/PassTimingInfo.cpp -> defines the default constructor of TimePassesHandler as reading those globals.

bool registerCallBacks = false;
bool debugLogging = false;
if (mlir::triton::tools::getBoolEnv("LLVM_ENABLE_TIMING")) {
llvm::TimePassesIsEnabled = true;
llvm::TimePassesPerRun = true;
registerCallBacks = true;
}
if (mlir::triton::tools::getBoolEnv("LLVM_IR_ENABLE_DUMP")) {
auto optMap = llvm::cl::getRegisteredOptions();
auto optIt = optMap.find("print-after-all");
if (optIt != optMap.end()) {
auto optPtr = static_cast<llvm::cl::opt<bool> *>(optIt->second);
*optPtr = true;
}
debugLogging = true;
registerCallBacks = true;
}

PassInstrumentationCallbacks *instrCbPtr = nullptr;
PassInstrumentationCallbacks passInstrCb;
StandardInstrumentations standardInstr(mod->getContext(), debugLogging);
if (registerCallBacks) {
standardInstr.registerCallbacks(passInstrCb, &mam);
instrCbPtr = &passInstrCb;
}
Expand Down
Loading