Skip to content

Commit 36f0c62

Browse files
at0m741at0m741
committed
Ricci Tensor Doxygen update
1 parent 369c5b7 commit 36f0c62

File tree

851 files changed

+65066
-40677
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

851 files changed

+65066
-40677
lines changed

Doc/CauchyManifold.png renamed to Notebooks/CauchyManifold.png

File renamed without changes.

Doc/MorpheusGeneralRelativity.ipynb renamed to Notebooks/MorpheusGeneralRelativity.ipynb

File renamed without changes.

Doc/MorpheusMatMulBlock.ipynb renamed to Notebooks/MorpheusMatMulBlock.ipynb

File renamed without changes.

Doc/MorpheusMatVec.ipynb renamed to Notebooks/MorpheusMatVec.ipynb

File renamed without changes.

‎Plugins/LLVM_Handler.cpp‎

Lines changed: 49 additions & 49 deletions
Original file line numberDiff line numberDiff line change
@@ -3,52 +3,52 @@
33
#include <memory>
44

55
void extractFunctionToMLIR(clang::FunctionDecl *FD, clang::ASTContext &Context) {
6-
const std::string funcName = FD->getNameAsString();
7-
const std::string path = "output.ll";
8-
9-
llvm::LLVMContext llvmCtx;
10-
llvm::SMDiagnostic err;
11-
std::unique_ptr<llvm::Module> fullMod = llvm::parseIRFile(path, err, llvmCtx);
12-
if (!fullMod) {
13-
llvm::errs() << "could not load LLVM module from " << path << "\n";
14-
return;
15-
}
16-
17-
auto *func = fullMod->getFunction(funcName);
18-
if (!func) {
19-
llvm::errs() << "function " << funcName << " not found in " << path << "\n";
20-
return;
21-
}
22-
23-
auto newMod = std::make_unique<llvm::Module>("tensorium_gpu_kernel", llvmCtx);
24-
llvm::ValueToValueMapTy VMap;
25-
auto *clonedFunc = llvm::CloneFunction(func, VMap);
26-
newMod->getFunctionList().push_back(clonedFunc);
27-
28-
mlir::MLIRContext mlirCtx;
29-
mlirCtx.getOrLoadDialect<mlir::LLVM::LLVMDialect>();
30-
mlirCtx.getOrLoadDialect<mlir::vector::VectorDialect>();
31-
mlirCtx.getOrLoadDialect<mlir::gpu::GPUDialect>();
32-
mlirCtx.getOrLoadDialect<mlir::NVVM::NVVMDialect>();
33-
34-
auto mlirModule = mlir::translateLLVMIRToMLIR(mlirCtx, std::move(newMod));
35-
if (!mlirModule) {
36-
llvm::errs() << "MLIR translation failed\n";
37-
return;
38-
}
39-
40-
mlir::PassManager pm(&mlirCtx);
41-
pm.addPass(mlir::createConvertVectorToGPU());
42-
pm.addPass(mlir::createGPUKernelOutliningPass());
43-
pm.addPass(mlir::createConvertGPUToNVVMPass());
44-
45-
if (mlir::failed(pm.run(*mlirModule))) {
46-
llvm::errs() << "MLIR passes failed\n";
47-
return;
48-
}
49-
50-
std::string outPath = "/tmp/" + funcName + "_gpu.mlir";
51-
std::error_code ec;
52-
llvm::raw_fd_ostream out(outPath, ec);
53-
if (ec) {
54-
llvm::errs() << "could not write to " << outPath << ": " << ec.message()
6+
const std::string funcName = FD->getNameAsString();
7+
const std::string path = "output.ll";
8+
9+
llvm::LLVMContext llvmCtx;
10+
llvm::SMDiagnostic err;
11+
std::unique_ptr<llvm::Module> fullMod = llvm::parseIRFile(path, err, llvmCtx);
12+
if (!fullMod) {
13+
llvm::errs() << "could not load LLVM module from " << path << "\n";
14+
return;
15+
}
16+
17+
auto *func = fullMod->getFunction(funcName);
18+
if (!func) {
19+
llvm::errs() << "function " << funcName << " not found in " << path << "\n";
20+
return;
21+
}
22+
23+
auto newMod = std::make_unique<llvm::Module>("tensorium_gpu_kernel", llvmCtx);
24+
llvm::ValueToValueMapTy VMap;
25+
auto *clonedFunc = llvm::CloneFunction(func, VMap);
26+
newMod->getFunctionList().push_back(clonedFunc);
27+
28+
mlir::MLIRContext mlirCtx;
29+
mlirCtx.getOrLoadDialect<mlir::LLVM::LLVMDialect>();
30+
mlirCtx.getOrLoadDialect<mlir::vector::VectorDialect>();
31+
mlirCtx.getOrLoadDialect<mlir::gpu::GPUDialect>();
32+
mlirCtx.getOrLoadDialect<mlir::NVVM::NVVMDialect>();
33+
34+
auto mlirModule = mlir::translateLLVMIRToMLIR(mlirCtx, std::move(newMod));
35+
if (!mlirModule) {
36+
llvm::errs() << "MLIR translation failed\n";
37+
return;
38+
}
39+
40+
mlir::PassManager pm(&mlirCtx);
41+
pm.addPass(mlir::createConvertVectorToGPU());
42+
pm.addPass(mlir::createGPUKernelOutliningPass());
43+
pm.addPass(mlir::createConvertGPUToNVVMPass());
44+
45+
if (mlir::failed(pm.run(*mlirModule))) {
46+
llvm::errs() << "MLIR passes failed\n";
47+
return;
48+
}
49+
50+
std::string outPath = "/tmp/" + funcName + "_gpu.mlir";
51+
std::error_code ec;
52+
llvm::raw_fd_ostream out(outPath, ec);
53+
if (ec) {
54+
llvm::errs() << "could not write to " << outPath << ": " << ec.message()

‎Plugins/LLVM_Handler.hpp‎

Lines changed: 15 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -1,30 +1,29 @@
11
#pragma once
2-
#include <iostream>
3-
#include <vector>
2+
#include "mlir/Dialect/GPU/GPUDialect.h"
3+
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
4+
#include "mlir/Dialect/NVVM/IR/NVVMDialect.h"
5+
#include "mlir/Dialect/Vector/IR/VectorOps.h"
6+
#include "mlir/IR/BuiltinOps.h"
7+
#include "mlir/IR/MLIRContext.h"
8+
#include "mlir/Pass/PassManager.h"
9+
#include "mlir/Target/LLVMIR/Import.h"
410
#include "clang/AST/ASTContext.h"
511
#include "clang/AST/Decl.h"
612
#include "clang/AST/DeclCXX.h"
7-
#include "llvm/IR/LLVMContext.h"
13+
#include "llvm/AsmParser/Parser.h"
814
#include "llvm/IR/Function.h"
15+
#include "llvm/IR/LLVMContext.h"
916
#include "llvm/IR/Module.h"
1017
#include "llvm/Support/SourceMgr.h"
11-
#include "llvm/AsmParser/Parser.h"
1218
#include "llvm/Transforms/Utils/Cloning.h"
13-
#include "mlir/IR/MLIRContext.h"
14-
#include "mlir/IR/BuiltinOps.h"
15-
#include "mlir/Pass/PassManager.h"
16-
#include "mlir/Target/LLVMIR/Import.h"
17-
#include "mlir/Dialect/Vector/IR/VectorOps.h"
18-
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
19-
#include "mlir/Dialect/GPU/GPUDialect.h"
20-
#include "mlir/Dialect/NVVM/IR/NVVMDialect.h"
19+
#include <iostream>
2120
#include <string>
21+
#include <vector>
2222

2323
struct TensoriumTarget {
24-
std::string platform;
25-
std::string isa;
26-
clang::SourceLocation loc;
24+
std::string platform;
25+
std::string isa;
26+
clang::SourceLocation loc;
2727
};
2828

2929
static std::vector<TensoriumTarget> TensoriumTargetTable;
30-

‎Plugins/MLIR/ConvertMemRefToLLVMPlugin.cpp‎

Lines changed: 14 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -1,35 +1,31 @@
1-
#include "mlir/Tools/Plugins/PassPlugin.h"
2-
#include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h"
31
#include "mlir/Conversion/LLVMCommon/ConversionTarget.h"
42
#include "mlir/Conversion/LLVMCommon/TypeConverter.h"
3+
#include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h"
54
#include "mlir/IR/BuiltinOps.h"
5+
#include "mlir/Tools/Plugins/PassPlugin.h"
66
#include "mlir/Transforms/DialectConversion.h"
77

88
using namespace mlir;
99

1010
namespace {
1111
struct ConvertMemRefToLLVMPass
1212
: public PassWrapper<ConvertMemRefToLLVMPass, OperationPass<ModuleOp>> {
13-
void runOnOperation() override {
14-
RewritePatternSet patterns(&getContext());
15-
LLVMTypeConverter converter(&getContext());
16-
populateFinalizeMemRefToLLVMConversionPatterns(converter, patterns);
13+
void runOnOperation() override {
14+
RewritePatternSet patterns(&getContext());
15+
LLVMTypeConverter converter(&getContext());
16+
populateFinalizeMemRefToLLVMConversionPatterns(converter, patterns);
1717

18-
ConversionTarget target(getContext());
19-
target.addLegalDialect<LLVM::LLVMDialect>();
20-
target.addIllegalDialect<memref::MemRefDialect>();
18+
ConversionTarget target(getContext());
19+
target.addLegalDialect<LLVM::LLVMDialect>();
20+
target.addIllegalDialect<memref::MemRefDialect>();
2121

22-
if (failed(applyPartialConversion(getOperation(), target,
23-
std::move(patterns))))
24-
signalPassFailure();
25-
}
22+
if (failed(applyPartialConversion(getOperation(), target, std::move(patterns))))
23+
signalPassFailure();
24+
}
2625
};
2726
} // end anonymous namespace
2827

2928
extern "C" ::mlir::PassPluginLibraryInfo mlirGetPassPluginInfo() {
30-
return {MLIR_PLUGIN_API_VERSION, "ConvMemRef2LLVM", "v1",
31-
[](PassManager &pm) {
32-
pm.addPass(std::make_unique<ConvertMemRefToLLVMPass>());
33-
}};
29+
return {MLIR_PLUGIN_API_VERSION, "ConvMemRef2LLVM", "v1",
30+
[](PassManager &pm) { pm.addPass(std::make_unique<ConvertMemRefToLLVMPass>()); }};
3431
}
35-

‎Plugins/MLIR/tensorium-opt.cpp‎

Lines changed: 20 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -2,13 +2,13 @@
22
#include "mlir/InitAllDialects.h"
33
#include "mlir/InitAllPasses.h"
44
#include "mlir/Pass/PassManager.h"
5-
#include "mlir/Tools/Plugins/PassPlugin.h"
65
#include "mlir/Support/FileUtilities.h"
6+
#include "mlir/Tools/Plugins/PassPlugin.h"
77
#include "mlir/Tools/mlir-opt/MlirOptMain.h"
88

9-
#include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h"
109
#include "mlir/Conversion/LLVMCommon/ConversionTarget.h"
1110
#include "mlir/Conversion/LLVMCommon/TypeConverter.h"
11+
#include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h"
1212
#include "mlir/Dialect/MemRef/IR/MemRef.h"
1313
#include "mlir/IR/BuiltinOps.h"
1414
#include "mlir/Transforms/DialectConversion.h"
@@ -18,34 +18,29 @@ using namespace mlir;
1818
namespace {
1919
struct ConvertMemRefToLLVMPass
2020
: public PassWrapper<ConvertMemRefToLLVMPass, OperationPass<ModuleOp>> {
21-
void runOnOperation() override {
22-
MLIRContext &ctx = getContext();
23-
RewritePatternSet patterns(&ctx);
24-
LLVMTypeConverter converter(&ctx);
25-
populateFinalizeMemRefToLLVMConversionPatterns(converter, patterns);
21+
void runOnOperation() override {
22+
MLIRContext &ctx = getContext();
23+
RewritePatternSet patterns(&ctx);
24+
LLVMTypeConverter converter(&ctx);
25+
populateFinalizeMemRefToLLVMConversionPatterns(converter, patterns);
2626

27-
ConversionTarget target(ctx);
28-
target.addLegalDialect("llvm");
29-
target.addIllegalDialect<memref::MemRefDialect>();
27+
ConversionTarget target(ctx);
28+
target.addLegalDialect("llvm");
29+
target.addIllegalDialect<memref::MemRefDialect>();
3030

31-
if (failed(applyPartialConversion(getOperation(), target,
32-
std::move(patterns))))
33-
signalPassFailure();
34-
}
31+
if (failed(applyPartialConversion(getOperation(), target, std::move(patterns))))
32+
signalPassFailure();
33+
}
3534
};
3635
} // namespace
3736

38-
static PassPipelineRegistration<EmptyPipelineOptions> pipeline(
39-
"convert-memref-to-llvm",
40-
"Convert MemRef dialect to LLVM dialect",
41-
[](OpPassManager &pm) {
42-
pm.addPass(std::make_unique<ConvertMemRefToLLVMPass>());
43-
});
37+
static PassPipelineRegistration<EmptyPipelineOptions>
38+
pipeline("convert-memref-to-llvm", "Convert MemRef dialect to LLVM dialect",
39+
[](OpPassManager &pm) { pm.addPass(std::make_unique<ConvertMemRefToLLVMPass>()); });
4440

4541
int main(int argc, char **argv) {
46-
DialectRegistry registry;
47-
registerAllDialects(registry);
48-
registerAllPasses();
49-
return asMainReturnCode(
50-
MlirOptMain(argc, argv, "Tensorium MLIR optimizer\n", registry));
42+
DialectRegistry registry;
43+
registerAllDialects(registry);
44+
registerAllPasses();
45+
return asMainReturnCode(MlirOptMain(argc, argv, "Tensorium MLIR optimizer\n", registry));
5146
}

0 commit comments

Comments
 (0)