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

Uplift llvm project version #882

Merged
merged 7 commits into from
Oct 15, 2024
Merged

Uplift llvm project version #882

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
50ab750
Uplift llvm-version for torch-mlir project integration.
uazizTT Oct 4, 2024
f52428c
Change meshShape type to llvm::SmallVector to match llvm structs.
uazizTT Oct 7, 2024
967a421
Use ListOption with meshShape conversion.
uazizTT Oct 11, 2024
461fac8
Mark test as XFAIL failing due to missing torch package.
uazizTT Oct 15, 2024
66e3aca
Fix `emitc.variable` and `emitc.call_opque` usage
mtopalovicTT Oct 14, 2024
81436bb
Fix a formatting issue.
uazizTT Oct 15, 2024
d722d4c
Fixing missing pytorch when running mac tests
topke011 Oct 15, 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
5 changes: 3 additions & 2 deletions env/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,8 @@ cmake_minimum_required(VERSION 3.20.0)
project(ttmlir-toolchain LANGUAGES CXX C)

set(FLATBUFFERS_VERSION "fb9afbafc7dfe226b9db54d4923bfb8839635274")
set(LLVM_PROJECT_VERSION "9ddfe62f5c11e3f65f444209f514029ded2d58b9")
set(LLVM_PROJECT_VERSION "e813750354bbc08551cf23ff559a54b4a9ea1f29")
set(STABLEHLO_VERSION "d40285ef3db0687e3f1e2bb0d716d748485a9739")

include(ExternalProject)

Expand Down Expand Up @@ -69,7 +70,7 @@ ExternalProject_Add(
ExternalProject_Add(stablehlo
PREFIX ${TTMLIR_TOOLCHAIN_DIR}
GIT_REPOSITORY https://github.com/openxla/stablehlo.git
GIT_TAG v1.5.0
GIT_TAG ${STABLEHLO_VERSION}
CONFIGURE_COMMAND ""
BUILD_COMMAND ""
INSTALL_COMMAND ""
Expand Down
13 changes: 11 additions & 2 deletions lib/Conversion/TTKernelToEmitC/TTKernelToEmitC.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -115,7 +115,8 @@ class TTKernelToEmitCTypeConverter : public TypeConverter {
return Builder(ctx).getI64Type();
});
addConversion([ctx](mlir::tt::ttkernel::CBType type) -> Type {
return Builder(ctx).getType<emitc::OpaqueType>("::tt::CB");
auto cbOpaqueType = Builder(ctx).getType<emitc::OpaqueType>("::tt::CB");
return emitc::LValueType::get(cbOpaqueType);
});
}
};
Expand All @@ -138,9 +139,17 @@ class TTMetalToEmitCFuncArgsRewriter
rewriter.setInsertionPointToStart(&op.getCallableRegion()->front());
for (auto arg : blockArgs) {
auto cb = cast<ttkernel::CBType>(arg.getType());
// Get opaque type i.e emitc::LValueType<emitc::OpaqueType>
auto cbType = getTypeConverter()->convertType(cb);
auto var = rewriter.create<emitc::VariableOp>(
// Create a variable of type emitc::LValueType<emitc::OpaqueType>
auto lValueVar = rewriter.create<emitc::VariableOp>(
op.getLoc(), cbType, convertCBPort(rewriter, cb.getPort()));
// Get the emitc::OpaqueType from the emitc::LValueType<emitc::OpaqueType>
auto opaqueType = cast<emitc::LValueType>(cbType).getValueType();
// Load the value from the lvalue variable this
// will allow use to use the value
auto var =
rewriter.create<emitc::LoadOp>(op.getLoc(), opaqueType, lValueVar);
arg.replaceAllUsesWith(var);
}
op.getCallableRegion()->front().eraseArguments(0, blockArgs.size());
Expand Down
3 changes: 2 additions & 1 deletion lib/Dialect/TTMetal/Pipelines/TTMetalPipelines.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,8 @@ void createTTIRToTTMetalBackendPipeline(
OpPassManager &pm, const TTIRToTTMetalBackendPipelineOptions &options) {
pm.addPass(mlir::tt::ttir::createTTIRLoadSystemDesc());
ttir::TTIRImplicitDeviceOptions implicitDeviceOptions;
implicitDeviceOptions.meshShape = options.meshShape;
implicitDeviceOptions.meshShape = ::llvm::SmallVector<int64_t>(
options.meshShape.begin(), options.meshShape.end());
pm.addPass(mlir::tt::ttir::createTTIRImplicitDevice(implicitDeviceOptions));
pm.addPass(mlir::tt::ttir::createTTIRConstantAsFill());
pm.addPass(mlir::tt::ttir::createTTIRGenericRegion());
Expand Down
3 changes: 2 additions & 1 deletion lib/Dialect/TTNN/Pipelines/TTNNPipelines.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,7 +23,8 @@ void createTTNNPipelineTTIRPasses(
pm.addPass(mlir::tt::ttir::createTTIRLoadSystemDesc(systemDescOptions));

ttir::TTIRImplicitDeviceOptions implicitDeviceOptions;
implicitDeviceOptions.meshShape = options.meshShape;
implicitDeviceOptions.meshShape = ::llvm::SmallVector<int64_t>(
options.meshShape.begin(), options.meshShape.end());
pm.addPass(mlir::tt::ttir::createTTIRImplicitDevice(implicitDeviceOptions));
mlir::tt::ttir::TTIRLayoutOptions layoutOptions;
layoutOptions.initMemorySpace = mlir::tt::MemorySpace::System;
Expand Down
1 change: 1 addition & 0 deletions test/python/requirements.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,2 +1,3 @@
lit
pytest
torch==2.3.0 --index-url https://download.pytorch.org/whl/cpu
Copy link
Contributor

Choose a reason for hiding this comment

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

this is needed because test/python rely on torch which is installed through ttrt, but if ttrt installation doesn't happen before this test is executed it leads to test failures like here:
https://github.com/tenstorrent/tt-mlir/actions/runs/11329283307/job/31528275878

Loading