Merge remote-tracking branch 'upstream/main' into upgrade_triton_mlir_rocm_to_llvm_head

This commit is contained in:
Rohit Santhanam
2023-02-18 09:25:20 +00:00
98 changed files with 1873 additions and 1895 deletions

View File

@@ -49,7 +49,7 @@ llvm_update_compile_flags(triton-translate)
# MLIR core
MLIROptLib
MLIRIR
MLIRLLVMIR
MLIRLLVMDialect
MLIRPass
MLIRSupport
MLIRTransforms
@@ -59,6 +59,6 @@ llvm_update_compile_flags(triton-translate)
MLIRLLVMToLLVMIRTranslation
MLIRNVVMToLLVMIRTranslation
MLIRROCDLToLLVMIRTranslation
MLIRSCFToStandard
MLIRSCFToControlFlow
)
mlir_check_all_link_libraries(triton-translate)

View File

@@ -19,6 +19,7 @@
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/InitLLVM.h"
#include "llvm/Support/Process.h"
#include "llvm/Support/SourceMgr.h"
#include "llvm/Support/WithColor.h"
#include "llvm/Support/raw_ostream.h"
#include <cmath>
@@ -360,6 +361,8 @@ static std::string GetCheckTypeAbbreviation(Check::FileCheckType Ty) {
return "bad-not";
case Check::CheckBadCount:
return "bad-count";
case Check::CheckMisspelled:
return "misspelled";
case Check::CheckNone:
llvm_unreachable("invalid FileCheckType");
}

View File

@@ -8,7 +8,7 @@
#include "mlir/IR/Dialect.h"
#include "mlir/InitAllPasses.h"
#include "mlir/Support/MlirOptMain.h"
#include "mlir/Tools/mlir-opt/MlirOptMain.h"
namespace mlir {
namespace test {
@@ -33,8 +33,8 @@ int main(int argc, char **argv) {
// TODO: register Triton & TritonGPU passes
mlir::DialectRegistry registry;
registry.insert<mlir::triton::TritonDialect,
mlir::triton::gpu::TritonGPUDialect, mlir::math::MathDialect,
mlir::arith::ArithmeticDialect, mlir::StandardOpsDialect,
mlir::triton::gpu::TritonGPUDialect, mlir::func::FuncDialect,
mlir::math::MathDialect, mlir::arith::ArithDialect,
mlir::scf::SCFDialect, mlir::gpu::GPUDialect>();
return mlir::asMainReturnCode(mlir::MlirOptMain(

View File

@@ -5,7 +5,10 @@
#include "triton/Target/HSACO/HSACOTranslation.h"
#include "mlir/IR/AsmState.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/Parser.h"
#include "mlir/IR/Dialect.h"
#include "mlir/Parser/Parser.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Pass/PassManager.h"
#include "mlir/Support/FileUtilities.h"
#include "mlir/Support/LogicalResult.h"
#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
@@ -32,9 +35,9 @@ OwningOpRef<ModuleOp> loadMLIRModule(llvm::StringRef inputFilename,
}
mlir::DialectRegistry registry;
registry.insert<TritonDialect, triton::gpu::TritonGPUDialect,
mlir::math::MathDialect, arith::ArithmeticDialect,
StandardOpsDialect, scf::SCFDialect>();
registry
.insert<TritonDialect, triton::gpu::TritonGPUDialect,
mlir::math::MathDialect, arith::ArithDialect, scf::SCFDialect>();
context.appendDialectRegistry(registry);
@@ -46,7 +49,8 @@ OwningOpRef<ModuleOp> loadMLIRModule(llvm::StringRef inputFilename,
context.loadAllAvailableDialects();
context.allowUnregisteredDialects();
OwningOpRef<ModuleOp> module(parseSourceFile(sourceMgr, &context));
OwningOpRef<ModuleOp> module =
parseSourceFile<ModuleOp>(sourceMgr, &context);
if (!module) {
llvm::errs() << "Parse MLIR file failed.";
return nullptr;