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

Bump llvm to llvm/llvm-project@95d993a #19811

Merged
merged 4 commits into from
Jan 25, 2025
Merged

Bump llvm to llvm/llvm-project@95d993a #19811

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
7be6a17
Bump llvm to https://github.com/iree-org/llvm-project/commit/95d993a8…
nirvedhmeshram Jan 24, 2025
5b476dc
change Util::AlingOp type to index
nirvedhmeshram Jan 24, 2025
de475bc
fix tosa tests
nirvedhmeshram Jan 25, 2025
077455c
fix tests
nirvedhmeshram Jan 25, 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
4 changes: 2 additions & 2 deletions compiler/src/iree/compiler/Codegen/Common/FlattenMemRefSubspanPass.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -96,7 +96,7 @@ struct FlattenMemRefTypeConverter final : public TypeConverter {
addConversion([](MemRefType type) -> std::optional<Type> {
int64_t offset;
SmallVector<int64_t> strides;
if (failed(getStridesAndOffset(type, strides, offset))) {
if (failed(type.getStridesAndOffset(strides, offset))) {
return nullptr;
}
// Since the memref gets linearized, use a stride 1, offset 0.
Expand Down Expand Up @@ -354,7 +354,7 @@ static Value linearizeIndices(Value sourceValue, ValueRange indices,
// dynamic.
SmallVector<int64_t> strides;
int64_t offset;
if (succeeded(getStridesAndOffset(sourceType, strides, offset))) {
if (succeeded(sourceType.getStridesAndOffset(strides, offset))) {
// The memref itself might have an offset, but we should not account for it
// when computing the linearization. The original memref might be
// `memref<?x?xf32, strided<[?, ?], offset: ?>`
Expand Down
2 changes: 1 addition & 1 deletion compiler/src/iree/compiler/Codegen/LLVMCPU/DispatchABI.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -781,7 +781,7 @@ MemRefDescriptor HALDispatchABI::loadBinding(Operation *forOp, int64_t ordinal,
// Construct the MemRefDescriptor type based on the information we have.
// NOTE: we could use the binding length to clamp this/check that the
// requested range is valid.
auto [strides, offset] = getStridesAndOffset(memRefType);
auto [strides, offset] = memRefType.getStridesAndOffset();
if (memRefType.hasStaticShape() &&
!llvm::any_of(strides, ShapedType::isDynamic) &&
!ShapedType::isDynamic(offset)) {
Expand Down
2 changes: 1 addition & 1 deletion compiler/src/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -392,7 +392,7 @@ class ConvertIREEBindingSubspanOp : public ConvertToLLVMPattern {
// Add the byte offset.
Value llvmBufferBasePtr = llvmBufferArg;

auto [strides, offset] = getStridesAndOffset(memrefType);
auto [strides, offset] = memrefType.getStridesAndOffset();
if (memrefType.hasStaticShape() &&
!llvm::any_of(strides, ShapedType::isDynamic) &&
!ShapedType::isDynamic(offset)) {
Expand Down
4 changes: 2 additions & 2 deletions compiler/src/iree/compiler/Codegen/LLVMGPU/test/nvvm_mma_sync_pipeline_test.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -71,7 +71,7 @@ hal.executable @mma_fused_fp16 {
// CHECK: llvm.br
// CHECK-COUNT-2: nvvm.ldmatrix {{.*}} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>
// CHECK-COUNT-2: nvvm.mma.sync {{.*}} {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (vector<2xf16>, vector<2xf16>, vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>
// CHECK-COUNT-2: llvm.inline_asm has_side_effects asm_dialect = att "cp.async.cg.shared.global [$0], [$1], $2, $3;\0A", "r,l,n,r" {{.*}}, {{.*}}, {{.*}}, {{.*}} : (!llvm.ptr<3>, !llvm.ptr<1>, i32, i32) -> ()
// CHECK-COUNT-2: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16, cache = cg, %{{.*}} : !llvm.ptr<3>, !llvm.ptr<1>, i32
// CHECK: nvvm.cp.async.commit.group
// CHECK: nvvm.cp.async.wait.group 2
// CHECK-COUNT-2: nvvm.ldmatrix {{.*}} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>
Expand Down Expand Up @@ -158,7 +158,7 @@ hal.executable @mma_fused_f32 {
// CHECK: nvvm.ldmatrix{{.*}} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>
// CHECK-COUNT-4: llvm.extractvalue{{.*}} : !llvm.struct<(i32, i32, i32, i32)>
// CHECK-COUNT-2: nvvm.mma.sync {{.*}} {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<tf32>, multiplicandBPtxType = #nvvm.mma_type<tf32>, shape = #nvvm.shape<m = 16, n = 8, k = 8>} : (i32, i32, f32) -> !llvm.struct<(f32, f32, f32, f32)>
// CHECK-COUNT-2: llvm.inline_asm has_side_effects asm_dialect = att "cp.async.cg.shared.global [$0], [$1], $2, $3;\0A", "r,l,n,r" {{.*}}, {{.*}}, {{.*}}, {{.*}} : (!llvm.ptr<3>, !llvm.ptr<1>, i32, i32) -> ()
// CHECK-COUNT-2: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16, cache = cg, %{{.*}} : !llvm.ptr<3>, !llvm.ptr<1>, i32
// CHECK: nvvm.cp.async.commit.group
// CHECK: nvvm.cp.async.wait.group 2
// CHECK: nvvm.ldmatrix{{.*}} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>
Expand Down
8 changes: 4 additions & 4 deletions compiler/src/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -451,7 +451,7 @@ hal.executable @mma_fused {
// SM80: nvvm.cp.async.wait.group 3
// SM80-COUNT-4: nvvm.wmma.load{{.*}} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)
// SM80-COUNT-2: nvvm.wmma.mma
// SM80-COUNT-2: llvm.inline_asm has_side_effects asm_dialect = att "cp.async.cg.shared.global [$0], [$1], $2, $3;\0A", "r,l,n,r" {{.*}}, {{.*}}, {{.*}}, {{.*}} : (!llvm.ptr<3>, !llvm.ptr<1>, i32, i32) -> ()
// SM80-COUNT-2: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16, cache = cg, %{{.*}} : !llvm.ptr<3>, !llvm.ptr<1>, i32
// SM80: nvvm.cp.async.commit.group
// SM80: llvm.br
// SM80-NOT: nvvm.wmma.mma
Expand Down Expand Up @@ -529,7 +529,7 @@ hal.executable @mma_fused_fp16 {
// SM80: nvvm.cp.async.wait.group 3
// SM80-COUNT-2: nvvm.wmma.load{{.*}} : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)
// SM80-COUNT-1: nvvm.wmma.mma
// SM80-COUNT-2: llvm.inline_asm has_side_effects asm_dialect = att "cp.async.cg.shared.global [$0], [$1], $2, $3;\0A", "r,l,n,r" {{.*}}, {{.*}}, {{.*}}, {{.*}} : (!llvm.ptr<3>, !llvm.ptr<1>, i32, i32) -> ()
// SM80-COUNT-2: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16, cache = cg, %{{.*}} : !llvm.ptr<3>, !llvm.ptr<1>, i32
// SM80: nvvm.cp.async.commit.group
// SM80: llvm.br
// SM80-NOT: nvvm.wmma.mma
Expand Down Expand Up @@ -602,7 +602,7 @@ hal.executable @mma_fused_fp16 {
// SM80: nvvm.cp.async.wait.group 3
// SM80-COUNT-4: nvvm.wmma.load{{.*}} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)
// SM80-COUNT-2: nvvm.wmma.mma
// SM80-COUNT-2: llvm.inline_asm has_side_effects asm_dialect = att "cp.async.cg.shared.global [$0], [$1], $2, $3;\0A", "r,l,n,r" {{.*}}, {{.*}}, {{.*}}, {{.*}} : (!llvm.ptr<3>, !llvm.ptr<1>, i32, i32) -> ()
// SM80-COUNT-2: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16, cache = cg, %{{.*}} : !llvm.ptr<3>, !llvm.ptr<1>, i32
// SM80: nvvm.cp.async.commit.group
// SM80: llvm.br
// SM80-NOT: nvvm.wmma.mma
Expand Down Expand Up @@ -670,7 +670,7 @@ hal.executable @mma_fused_fp16 {
// SM80: nvvm.cp.async.wait.group 3
// SM80-COUNT-4: nvvm.wmma.load{{.*}} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)
// SM80-COUNT-2: nvvm.wmma.mma
// SM80: llvm.inline_asm has_side_effects asm_dialect = att "cp.async.cg.shared.global [$0], [$1], $2, $3;\0A", "r,l,n,r" {{.*}}, {{.*}}, {{.*}}, {{.*}} : (!llvm.ptr<3>, !llvm.ptr<1>, i32, i32) -> ()
// SM80: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16, cache = cg, %{{.*}} : !llvm.ptr<3>, !llvm.ptr<1>, i32
// SM80: nvvm.cp.async.commit.group
// SM80: llvm.br
// SM80-NOT: nvvm.wmma.mma
Expand Down
12 changes: 6 additions & 6 deletions compiler/src/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -183,9 +183,9 @@ hal.executable @ceildiv_expand_dispatch {
// CDNA3-LABEL: hal.executable public @ceildiv_expand_dispatch
// CDNA3: hal.executable.variant public @rocm
// CDNA3-NOT: arith.ceildivsi
// CDNA3-COUNT-1: llvm.select {{.*}} : i1, i32
// CDNA3-COUNT-2: llvm.sdiv {{.*}} : i32
// CDNA3-COUNT-4: llvm.icmp {{.*}} : i32
// CDNA3-COUNT-2: llvm.and {{.*}} : i1
// CDNA3-COUNT-1: llvm.or {{.*}} : i1
// CDNA3-COUNT-1: llvm.select {{.*}} : i1, i32
// CDNA3-COUNT-1: llvm.select {{.*}} : vector<1xi1>, vector<1xi32>
// CDNA3-COUNT-2: llvm.sdiv {{.*}} : vector<1xi32>
// CDNA3-COUNT-4: llvm.icmp {{.*}} : vector<1xi32>
// CDNA3-COUNT-2: llvm.and {{.*}} : vector<1xi1>
// CDNA3-COUNT-1: llvm.or {{.*}} : vector<1xi1>
// CDNA3-COUNT-1: llvm.select {{.*}} : vector<1xi1>, vector<1xi32>
3 changes: 2 additions & 1 deletion compiler/src/iree/compiler/Codegen/LLVMGPU/test/vector_lowering.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,7 +12,8 @@ module {
// CHECK-LABEL: func.func @broadcast_read_lowering
// CHECK-SAME: (%[[ARG0:.+]]: memref<4096x32xf16>)
// CHECK: %[[INIT:.+]] = arith.constant dense<0.000000e+00> : vector<1x8xf16>
// CHECK: %[[ELEM:.+]] = memref.load %[[ARG0]]{{.*}} : memref<4096x32xf16>
// CHECK: %[[LOAD:.+]] = vector.load %[[ARG0]]{{.*}} : memref<4096x32xf16>
// CHECK: %[[ELEM:.+]] = vector.extract %[[LOAD]][0] : f16 from vector<1xf16>
// CHECK: %[[SPLAT:.+]] = vector.splat %[[ELEM]] : vector<8xf16>
// CHECK: %[[INSERT:.+]] = vector.insert %[[SPLAT]], %[[INIT]] [0] : vector<8xf16> into vector<1x8xf16>
// CHECK: return %[[INSERT]]
2 changes: 1 addition & 1 deletion compiler/src/iree/compiler/Codegen/VMVX/VMVXLowerLinalgMicrokernels.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -94,7 +94,7 @@ bool verifyMemRefInnerDimsContiguousRowMajor(MemRefType type) {
return true;
}

if (failed(mlir::getStridesAndOffset(type, strides, offset))) {
if (failed(type.getStridesAndOffset(strides, offset))) {
return false;
}

Expand Down
2 changes: 1 addition & 1 deletion compiler/src/iree/compiler/Dialect/Util/Conversion/MemRefToUtil/Patterns.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -58,7 +58,7 @@ static Value getByteOffsetForIndices(OpBuilder &builder, Location loc,
}
SmallVector<int64_t> strides;
int64_t offset;
if (failed(getStridesAndOffset(memrefType, strides, offset)) ||
if (failed(memrefType.getStridesAndOffset(strides, offset)) ||
strides[0] != 1) {
emitError(loc, "expected memref stride 1");
return {};
Expand Down
6 changes: 3 additions & 3 deletions compiler/src/iree/compiler/Dialect/Util/IR/UtilOps.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -346,12 +346,12 @@ def Util_AlignOp : Util_PureOp<"align", [
}];

let arguments = (ins
SignlessIntegerLike:$value,
SignlessIntegerLike:$alignment
SignlessIntegerOrIndexLike:$value,
SignlessIntegerOrIndexLike:$alignment
);

let results = (outs
SignlessIntegerLike:$result
SignlessIntegerOrIndexLike:$result
);

let assemblyFormat = [{
Expand Down
6 changes: 3 additions & 3 deletions llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/Input/InputOps.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -873,12 +873,12 @@ def IREEInput_AlignOp : IREEInput_PureOp<"align", [
}];

let arguments = (ins
SignlessIntegerLike:$value,
SignlessIntegerLike:$alignment
SignlessIntegerOrIndexLike:$value,
SignlessIntegerOrIndexLike:$alignment
);

let results = (outs
SignlessIntegerLike:$result
SignlessIntegerOrIndexLike:$result
);

let assemblyFormat = [{
Expand Down
12 changes: 6 additions & 6 deletions tests/e2e/tosa_ops/pad.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,23 +1,23 @@
func.func @pad_1D_test() {
%0 = util.unfoldable_constant dense<42> : tensor<2xi32>
%1 = "tosa.const"() { value = dense<[3, 2]> : tensor<2xi32> } : () -> (tensor<2xi32>)
%result = tosa.pad %0, %1 : (tensor<2xi32>, tensor<2xi32>) -> (tensor<7xi32>)
%1 = tosa.const_shape { value = dense<[3, 2]> : tensor<2xindex> } : () -> !tosa.shape<2>
%result = tosa.pad %0, %1 : (tensor<2xi32>, !tosa.shape<2>) -> (tensor<7xi32>)
check.expect_eq_const(%result, dense<[0, 0, 0, 42, 42, 0, 0]> : tensor<7xi32>) : tensor<7xi32>
return
}

func.func @pad_2D_test() {
%0 = util.unfoldable_constant dense<42> : tensor<2x2xi32>
%1 = "tosa.const"() { value = dense<[1, 1, 1, 1]> : tensor<4xi32> } : () -> (tensor<4xi32>)
%result = tosa.pad %0, %1 : (tensor<2x2xi32>, tensor<4xi32>) -> (tensor<4x4xi32>)
%1 = tosa.const_shape { value = dense<[1, 1, 1, 1]> : tensor<4xindex> } : () -> !tosa.shape<4>
%result = tosa.pad %0, %1 : (tensor<2x2xi32>, !tosa.shape<4>) -> (tensor<4x4xi32>)
check.expect_eq_const(%result, dense<[[0, 0, 0, 0], [0, 42, 42, 0], [0, 42, 42, 0], [0, 0, 0, 0]]> : tensor<4x4xi32>) : tensor<4x4xi32>
return
}

func.func @pad_3D_test() {
%0 = util.unfoldable_constant dense<42> : tensor<1x1x2xi32>
%1 = "tosa.const"() { value = dense<[0, 1, 1, 0, 0, 0]> : tensor<6xi32> } : () -> (tensor<6xi32>)
%result = tosa.pad %0, %1 : (tensor<1x1x2xi32>, tensor<6xi32>) -> (tensor<2x2x2xi32>)
%1 = tosa.const_shape { value = dense<[0, 1, 1, 0, 0, 0]> : tensor<6xindex> } : () -> !tosa.shape<6>
%result = tosa.pad %0, %1 : (tensor<1x1x2xi32>, !tosa.shape<6>) -> (tensor<2x2x2xi32>)
check.expect_eq_const(%result, dense<[[[0, 0], [42, 42]], [[0, 0], [0, 0]]]> : tensor<2x2x2xi32>) : tensor<2x2x2xi32>
return
}
2 changes: 1 addition & 1 deletion third_party/llvm-project
Open in desktop
Submodule llvm-project updated 3053 files
2 changes: 1 addition & 1 deletion third_party/stablehlo
Open in desktop
Submodule stablehlo updated 16 files
+2 −2 WORKSPACE.bazel
+1 −1 build_tools/llvm_version.txt
+1 −1 examples/c++/ExampleAdd.cpp
+3 −16 stablehlo/conversions/linalg/transforms/Rewriters.h
+39 −38 stablehlo/conversions/linalg/transforms/StablehloLegalizeToLinalg.cpp
+6 −4 stablehlo/conversions/tosa/tests/nullary.mlir
+6 −2 stablehlo/conversions/tosa/transforms/StablehloLegalizeToTosa.cpp
+2 −2 stablehlo/dialect/Version.cpp
+1 −1 stablehlo/dialect/Version.h
+65 −0 stablehlo/tests/chlo/chlo_legalize_to_stablehlo.mlir
+189 −0 stablehlo/tests/interpret/chlo/ragged_dot.mlir
+1 −1 stablehlo/tests/transforms/stablehlo_refine_parameters.mlir
+234 −5 stablehlo/transforms/ChloLegalizeToStablehlo.cpp
+67 −73 stablehlo/transforms/StablehloRefineArguments.cpp
+4 −0 stablehlo/transforms/StablehloRefineShapes.cpp
+19 −0 stablehlo/transforms/StablehloRefineShapes.h
2 changes: 1 addition & 1 deletion third_party/torch-mlir
Open in desktop
Submodule torch-mlir updated 40 files
+2 −2 build_tools/python_deploy/build_linux_packages.sh
+2 −2 docs/development.md
+11 −0 include/torch-mlir/Conversion/TorchToTosa/TosaLegalizeUtils.h
+28 −0 include/torch-mlir/Dialect/Torch/IR/GeneratedTorchOps.td
+45 −4 lib/Conversion/TorchOnnxToTorch/DefaultDomainGtoP.cpp
+44 −23 lib/Conversion/TorchToTosa/TorchToTosa.cpp
+4 −2 lib/Conversion/TorchToTosa/TosaLegalizeCommon.cpp
+77 −13 lib/Conversion/TorchToTosa/TosaLegalizeUtils.cpp
+49 −1 lib/Dialect/TMTensor/Transforms/Bufferize.cpp
+1 −2 lib/Dialect/TMTensor/Transforms/ConvertToLoops.cpp
+48 −0 lib/Dialect/Torch/IR/TorchOps.cpp
+30 −0 lib/Dialect/Torch/Transforms/AbstractInterpLibrary.cpp
+2 −2 lib/Dialect/Torch/Transforms/DecomposeComplexOps.cpp
+2 −2 lib/Dialect/Torch/Transforms/FuseQuantizedOps.cpp
+2 −2 lib/Dialect/Torch/Transforms/MatchQuantizedOps.cpp
+1 −1 lib/Dialect/Torch/Transforms/MaximizeValueSemantics.cpp
+2 −3 lib/Dialect/Torch/Transforms/PrepareForGlobalizeObjectGraph.cpp
+2 −2 lib/Dialect/Torch/Transforms/RecomposeComplexOps.cpp
+2 −2 lib/Dialect/Torch/Transforms/RestructureNonConstantAxes.cpp
+2 −2 lib/Dialect/Torch/Transforms/ScalarizeShapes.cpp
+2 −2 lib/Dialect/Torch/Transforms/SimplifyDtypeCalculations.cpp
+2 −2 lib/Dialect/Torch/Transforms/SimplifyShapeCalculations.cpp
+5 −4 lib/Dialect/Torch/Utils/Utils.cpp
+1 −1 lib/Dialect/TorchConversion/Transforms/BackendTypeConversionPasses.cpp
+1 −2 lib/Dialect/TorchConversion/Transforms/UnpackQuantTensor.cpp
+4 −7 lib/RefBackend/RefBackend.cpp
+6 −10 projects/pt1/e2e_testing/xfail_sets.py
+6 −3 projects/pt1/python/torch_mlir/jit_ir_importer/build_tools/abstract_interp_lib_gen.py
+4 −0 projects/pt1/python/torch_mlir/jit_ir_importer/build_tools/torch_ods_gen.py
+1 −1 projects/pt1/python/torch_mlir_e2e_test/linalg_on_tensors_backends/refbackend.py
+39 −0 projects/pt1/python/torch_mlir_e2e_test/test_suite/type_conversion.py
+27 −3 python/torch_mlir/fx.py
+1 −1 pytorch-hash.txt
+1 −1 pytorch-requirements.txt
+4 −0 setup.py
+69 −0 test/Conversion/TorchOnnxToTorch/simple_ops_g_to_p.mlir
+158 −25 test/Conversion/TorchToTosa/basic.mlir
+15 −15 test/Dialect/TMTensor/bufferize.mlir
+2 −2 test/RefBackend/mlprogram-bufferize.mlir
+1 −1 torchvision-requirements.txt
Loading