Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[AMD] Fix slow compilation issue due to inline print calls #5153

Merged
merged 6 commits into from
Nov 22, 2024
Merged

[AMD] Fix slow compilation issue due to inline print calls #5153

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
5d1b07b
[AMD] Fix slow compilation issue of test_scan_layouts
binarman Nov 14, 2024
a2da0c4
Uppercase comment
binarman Nov 18, 2024
53f1af2
Merge branch 'main' into enable_test_scan_layouts
binarman Nov 18, 2024
17d3f11
replace explicit function names with prefixes
binarman Nov 22, 2024
4650586
Merge remote-tracking branch 'origin/main' into enable_test_scan_layouts
antiagainst Nov 22, 2024
052fd0a
Light touches over implementation
antiagainst Nov 22, 2024
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
2 changes: 0 additions & 2 deletions python/src/llvm.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -139,8 +139,6 @@ std::string translateLLVMIRToASM(llvm::Module &module,
{
llvm::raw_string_ostream stream(result);
llvm::buffer_ostream pstream(stream);
for (llvm::Function &f : module.functions())
f.addFnAttr(llvm::Attribute::AlwaysInline);
Comment on lines -142 to -143
Copy link
Contributor Author

Choose a reason for hiding this comment

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

Removed this, because this attribute insertion already done on lines 106-108.
Seems that this duplication happened when Hopper support was merged.

llvm::legacy::PassManager pass;
// emit
auto fileType = isObject ? llvm::CodeGenFileType::ObjectFile
Expand Down
2 changes: 0 additions & 2 deletions python/test/unit/language/test_core.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2563,8 +2563,6 @@ def kernel(X, Y, N, BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, NUM_PID_N: tl.
@pytest.mark.parametrize("axis", [0, 1])
@pytest.mark.parametrize("add_overflow_check", [False, True])
def test_scan_layouts(M, N, src_layout, axis, add_overflow_check, device, tmp_path: pathlib.Path):
if add_overflow_check is True and is_hip():
pytest.skip("overflow check disabled on HIP while fixing issues")

overflow_check = """
%17 = arith.extsi %arg2 : i32 to i64
Expand Down
3 changes: 3 additions & 0 deletions third_party/amd/backend/compiler.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -342,6 +342,9 @@ def make_llir(src, metadata, options):
metadata["shared"] = src.get_int_attr("triton_gpu.shared")

amd.cleanup_bitcode_metadata(llvm_mod)
# Disable inlining of print related functions,
# because inlining of these function could slow down compilation significantly
amd.disable_print_inline(llvm_mod)
return str(llvm_mod)

@staticmethod
Expand Down
18 changes: 18 additions & 0 deletions third_party/amd/python/triton_amd.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -161,6 +161,24 @@ void init_triton_amd(py::module &&m) {
module->eraseNamedMetadata(openclVersion);
});

m.def("disable_print_inline", [](llvm::Module *module) {
antiagainst marked this conversation as resolved.
Show resolved Hide resolved
// List of functions name prefixes we want to forbid inline.
std::array<const char *, 2> prefixes = {"__ockl_fprintf", "__ockl_printf"};

for (llvm::Function &f : module->functions()) {
if (!f.hasName())
continue;
llvm::StringRef name = f.getName();

auto isNamePrefixed = [&name](const char *prefix) {
return name.starts_with(prefix);
};

if (llvm::any_of(prefixes, isNamePrefixed))
f.addFnAttr(llvm::Attribute::NoInline);
}
});

m.def(
"assemble_amdgcn",
[](const std::string &assembly, const std::string &arch,
Expand Down
Loading