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

[BACKEND] Remove all layout conversion decomposition functions from the NVIDIA backend #5553

Merged
merged 4 commits into from
Jan 10, 2025
+1 −67
Merged

[BACKEND] Remove all layout conversion decomposition functions from the NVIDIA backend #5553

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
ffbacae
Update
Jokeren Dec 31, 2024
780dac1
Merge branch 'main' into keren/no-decompose
Jokeren Jan 8, 2025
990e9f3
Update
Jokeren Jan 8, 2025
00224aa
Update
Jokeren Jan 8, 2025
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
1 change: 0 additions & 1 deletion bin/RegisterTritonDialects.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,7 +49,6 @@ inline void registerTritonDialects(mlir::DialectRegistry &registry) {
mlir::triton::registerTritonGPUGlobalScratchAllocationPass();
mlir::triton::registerConvertTritonGPUToLLVMPass();
mlir::triton::registerConvertNVGPUToLLVMPass();
mlir::triton::registerDecomposeUnsupportedNVIDIAConversions();
mlir::registerLLVMDIScope();

// TritonAMDGPUToLLVM passes
Expand Down
10 changes: 1 addition & 9 deletions lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -356,10 +356,6 @@ struct ConvertLayoutOpUsingLinearLayoutsConversion
auto srcTy = op.getSrc().getType();
auto dstTy = op.getType();

// TODO (Keren): Currently, we handle general mma/blocked/slice/dot(ampere)
// -> mma/blocked/slice/dot(ampere) conversions. The following tasks must be
// completed before we can remove the layoutIsOK check:
// 1. Support for AMD's WMMA dot operand
std::function<bool(Attribute)> layoutIsOK = [&](Attribute layout) {
if (isa<MmaEncodingTrait>(layout)) {
return !useLegacyMMAConversion;
Expand All @@ -368,15 +364,11 @@ struct ConvertLayoutOpUsingLinearLayoutsConversion
if (isa<MmaEncodingTrait>(dotOperand.getParent())) {
return !useLegacyMMAConversion;
}
return false;
}
if (isa<BlockedEncodingAttr, LinearEncodingAttr>(layout)) {
return true;
}
if (auto slice = dyn_cast<SliceEncodingAttr>(layout)) {
return layoutIsOK(slice.getParent());
}
return false;
return true;
};
if (!layoutIsOK(srcTy.getEncoding()) || !layoutIsOK(dstTy.getEncoding())) {
return failure();
Expand Down
1 change: 0 additions & 1 deletion third_party/nvidia/backend/compiler.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -276,7 +276,6 @@ def make_llir(src, metadata, options, capability):
srcMgr = llvm.source_mgr()
diag = ir.source_mgr_diag(srcMgr, mod.context)
mod.context.printOpOnDiagnostic(True)
nvidia.passes.ttgpuir.add_decompose_unsupported_conversions(pm)
passes.ttgpuir.add_combine_tensor_select_and_if(pm)
passes.convert.add_scf_to_cf(pm)
passes.convert.add_index_to_llvmir(pm)
Expand Down
6 changes: 0 additions & 6 deletions third_party/nvidia/include/TritonNVIDIAGPUToLLVM/Passes.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,12 +17,6 @@ namespace triton {
#define GEN_PASS_DECL
#include "nvidia/include/TritonNVIDIAGPUToLLVM/Passes.h.inc"

namespace NVIDIA {
std::unique_ptr<OperationPass<ModuleOp>>
createDecomposeUnsupportedConversionsPass();

} // namespace NVIDIA

std::unique_ptr<OperationPass<ModuleOp>> createConvertTritonGPUToLLVMPass();
std::unique_ptr<OperationPass<ModuleOp>>
createConvertTritonGPUToLLVMPass(int32_t computeCapability);
Expand Down
5 changes: 0 additions & 5 deletions third_party/nvidia/include/TritonNVIDIAGPUToLLVM/Passes.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,11 +3,6 @@

include "mlir/Pass/PassBase.td"

def DecomposeUnsupportedNVIDIAConversions : Pass<"decompose-unsupported-nvidia-conversions", "mlir::ModuleOp"> {
let summary = "Decompose conversions that are not supported by TritonGPU -> LLVM";
let constructor = "mlir::triton::NVIDIA::createDecomposeUnsupportedConversionsPass()";
}

def ConvertTritonGPUToLLVM : Pass<"convert-triton-gpu-to-llvm", "mlir::ModuleOp"> {
let summary = "Convert TritonGPU to LLVM";
let description = [{
Expand Down
1 change: 0 additions & 1 deletion third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,7 +10,6 @@ add_triton_library(TritonNVIDIAGPUToLLVM
BarrierOpToLLVM.cpp
TritonGPUToLLVM.cpp
TMAToLLVM.cpp
DecomposeUnsupportedConversions.cpp
SPMDOpToLLVM.cpp
TensorPtrOpsToLLVM.cpp
ClusterOpsToLLVM.cpp
Expand Down
41 changes: 0 additions & 41 deletions third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/DecomposeUnsupportedConversions.cpp
View file
Open in desktop

This file was deleted.

3 changes: 0 additions & 3 deletions third_party/nvidia/triton_nvidia.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,9 +23,6 @@ void init_triton_nvidia_passes_ttgpuir(py::module &&m) {
pm.addPass(mlir::triton::createConvertTritonGPUToLLVMPass(
capability, ptxVersion));
});
m.def("add_decompose_unsupported_conversions", [](mlir::PassManager &pm) {
pm.addPass(NVIDIA::createDecomposeUnsupportedConversionsPass());
});
}

void init_triton_nvidia_passes_ttnvgpuir(py::module &&m) {
Expand Down
Loading