diff --git a/compiler/src/iree/compiler/Codegen/Common/ConvertToDestinationPassingStylePass.cpp b/compiler/src/iree/compiler/Codegen/Common/ConvertToDestinationPassingStylePass.cpp
index 49f0c6a..faa275f 100644
--- a/compiler/src/iree/compiler/Codegen/Common/ConvertToDestinationPassingStylePass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/ConvertToDestinationPassingStylePass.cpp
@@ -505,7 +505,7 @@
       if (!attr.isSplat()) continue;
       auto type = attr.getType().dyn_cast<RankedTensorType>();
       if (!type) continue;
-      Attribute scalarAttr = attr.getValues<Attribute>()[0];
+      TypedAttr scalarAttr = attr.getValues<TypedAttr>()[0];
 
       modifiedOutput = true;
       Value emptyTensor = rewriter.create<tensor::EmptyOp>(
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPUPatterns.cpp b/compiler/src/iree/compiler/Codegen/Common/GPUPatterns.cpp
index fb27df4..fc9542e 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GPUPatterns.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/GPUPatterns.cpp
@@ -187,9 +187,10 @@
   // Limit promotion to matmul and batch matmul, there may be generic
   // ops with more batch dimensions we didn't distribute and therefore
   // cannot find a higher bound.
-  return success(linalg::isaContractionOpInterface(op) &&
-                 linalgOp.getNumParallelLoops() >= 2 &&
-                 linalgOp.getNumParallelLoops() <= 3);
+  return success(
+      linalg::isaContractionOpInterface(cast<linalg::LinalgOp>(op)) &&
+      linalgOp.getNumParallelLoops() >= 2 &&
+      linalgOp.getNumParallelLoops() <= 3);
 }
 
 }  // namespace
diff --git a/compiler/src/iree/compiler/Codegen/Common/TransformExtensions/CommonExtensions.cpp b/compiler/src/iree/compiler/Codegen/Common/TransformExtensions/CommonExtensions.cpp
index ac7b010..fd12e4b 100644
--- a/compiler/src/iree/compiler/Codegen/Common/TransformExtensions/CommonExtensions.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/TransformExtensions/CommonExtensions.cpp
@@ -650,7 +650,7 @@
     return forallOp->emitError("mapping must be present");
   SmallVector<Attribute> blockMapping =
       llvm::to_vector(forallOp.getMapping()->getValue());
-  if (llvm::any_of(blockMapping, [](DeviceMappingAttrInterface map) {
+  if (llvm::any_of(blockMapping, [](Attribute map) {
         return !map.isa<gpu::GPUBlockMappingAttr>();
       })) {
     return forallOp->emitError("mapping must be #gpu.block<x/y/z/>");
diff --git a/compiler/src/iree/compiler/Codegen/Common/TypePropagationPass.cpp b/compiler/src/iree/compiler/Codegen/Common/TypePropagationPass.cpp
index 8266932..0d942a1 100644
--- a/compiler/src/iree/compiler/Codegen/Common/TypePropagationPass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/TypePropagationPass.cpp
@@ -137,8 +137,8 @@
     auto newAttrType = RankedTensorType::get(attrType.getShape(),
                                              legalizedElementType.value());
     auto newAttr = DenseElementsAttr::get(newAttrType, legalizedValues);
-    rewriter.replaceOpWithNewOp<arith::ConstantOp>(constantOp, newAttr,
-                                                   newAttrType);
+    rewriter.replaceOpWithNewOp<arith::ConstantOp>(constantOp, newAttrType,
+                                                   newAttr);
     return success();
   }
 };
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/convert_to_destination_passing_style.mlir b/compiler/src/iree/compiler/Codegen/Common/test/convert_to_destination_passing_style.mlir
index bc9e88a..524f13f 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/convert_to_destination_passing_style.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/convert_to_destination_passing_style.mlir
@@ -472,51 +472,6 @@
 
 // -----
 
-func.func @fill_matmul_exp() {
-  %cst = arith.constant 0.000000e+00 : f32
-  %c0 = arith.constant 0 : index
-  %c33 = arith.constant 33 : index
-  %c49 = arith.constant 49 : index
-  %0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) alignment(32) offset(%c0) : !flow.dispatch.tensor<readonly:tensor<33x16xf32>>
-  %1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) alignment(32) offset(%c0) : !flow.dispatch.tensor<readonly:tensor<16x49xf32>>
-  %2 = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer) alignment(32) offset(%c0) : !flow.dispatch.tensor<writeonly:tensor<33x49xf32>>
-  %workgroup_id_x = hal.interface.workgroup.id[0] : index
-  %workgroup_count_x = hal.interface.workgroup.count[0] : index
-  %workgroup_id_y = hal.interface.workgroup.id[1] : index
-  %workgroup_count_y = hal.interface.workgroup.count[1] : index
-  %3 = affine.apply affine_map<()[s0] -> (s0 * 16)>()[%workgroup_id_y]
-  %4 = affine.apply affine_map<()[s0] -> (s0 * 16)>()[%workgroup_count_y]
-  scf.for %arg0 = %3 to %c33 step %4 {
-    %5 = affine.apply affine_map<()[s0] -> (s0 * 16)>()[%workgroup_id_x]
-    %6 = affine.apply affine_map<()[s0] -> (s0 * 16)>()[%workgroup_count_x]
-    scf.for %arg1 = %5 to %c49 step %6 {
-      %7 = affine.min affine_map<(d0) -> (16, -d0 + 33)>(%arg0)
-      %8 = affine.min affine_map<(d0) -> (16, -d0 + 49)>(%arg1)
-      %9 = tensor.empty(%7, %8) : tensor<?x?xf32>
-      %10 = affine.min affine_map<(d0) -> (-d0 + 33, 16)>(%arg0)
-      %11 = flow.dispatch.tensor.load %0, offsets = [%arg0, 0], sizes = [%10, 16], strides = [1, 1] : !flow.dispatch.tensor<readonly:tensor<33x16xf32>> -> tensor<?x16xf32>
-      %12 = affine.min affine_map<(d0) -> (-d0 + 49, 16)>(%arg1)
-      %13 = flow.dispatch.tensor.load %1, offsets = [0, %arg1], sizes = [16, %12], strides = [1, 1] : !flow.dispatch.tensor<readonly:tensor<16x49xf32>> -> tensor<16x?xf32>
-      %14 = tensor.empty(%10, %12) : tensor<?x?xf32>
-      %15 = linalg.fill ins(%cst : f32) outs(%14 : tensor<?x?xf32>) -> tensor<?x?xf32>
-      %16 = linalg.matmul ins(%11, %13 : tensor<?x16xf32>, tensor<16x?xf32>) outs(%15 : tensor<?x?xf32>) -> tensor<?x?xf32>
-      %17 = linalg.generic {indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>], iterator_types = ["parallel", "parallel"]} ins(%16 : tensor<?x?xf32>) outs(%9 : tensor<?x?xf32>) {
-      ^bb0(%arg2: f32, %arg3: f32):
-        %18 = math.exp %arg2 : f32
-        linalg.yield %18 : f32
-      } -> tensor<?x?xf32>
-      flow.dispatch.tensor.store %17, %2, offsets = [%arg0, %arg1], sizes = [%7, %8], strides = [1, 1] : tensor<?x?xf32> -> !flow.dispatch.tensor<writeonly:tensor<33x49xf32>>
-    }
-  }
-  return
-}
-// CHECK-LABEL: func.func @fill_matmul_exp()
-//       CHECK:   %[[MATMUL:.+]] = linalg.matmul
-//       CHECK:   linalg.generic
-//  CHECK-SAME:       outs(%[[MATMUL]]
-
-// -----
-
 func.func @cumsum__2x2x2x2x2x2x2() {
   %cst = arith.constant dense<0.000000e+00> : tensor<2x2x2x2x2x2x2xf32>
   %c0 = arith.constant 0 : index
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
index 70f949d..099ba22 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
@@ -760,10 +760,11 @@
 
     Type wideType = rewriter.getIntegerType(64);
     // Shift amount necessary to extract the high bits from widened result.
-    Attribute shiftValAttr = rewriter.getI64IntegerAttr(32);
+    TypedAttr shiftValAttr = rewriter.getI64IntegerAttr(32);
     if (auto vecTy = resultType.dyn_cast<VectorType>()) {
       wideType = VectorType::get(vecTy.getShape(), wideType);
-      shiftValAttr = SplatElementsAttr::get(wideType, shiftValAttr);
+      shiftValAttr =
+          SplatElementsAttr::get(cast<ShapedType>(wideType), shiftValAttr);
     }
     Value shiftVal = rewriter.create<arith::ConstantOp>(loc, shiftValAttr);
 
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
index 9caf23a..2282595 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
@@ -839,12 +839,13 @@
     parallelTileSizes.push_back(sz);
   }
   SmallVector<int64_t> reductionTileSizes;
-  splitParallelAndReductionTiles(op.getOperation(), parallelTileSizes,
-                                 reductionTileSizes);
-
-  setVectorSizesForDynamicShapes(op.getOperation(), vecPreProcStrategy,
+  splitParallelAndReductionTiles(cast<linalg::LinalgOp>(op.getOperation()),
                                  parallelTileSizes, reductionTileSizes);
 
+  setVectorSizesForDynamicShapes(cast<linalg::LinalgOp>(op.getOperation()),
+                                 vecPreProcStrategy, parallelTileSizes,
+                                 reductionTileSizes);
+
   TileSizesListType newTileSizes;
   // Copy all the tile size levels except the workgroup one which will be split
   // into parallel and reduction.
@@ -880,8 +881,8 @@
       getMaxVectorTileSize(0, K, workgroupTileSizes.back(), vectorSize));
 
   SmallVector<int64_t> reductionTileSizes;
-  splitParallelAndReductionTiles(op.getOperation(), parallelTileSizes,
-                                 reductionTileSizes);
+  splitParallelAndReductionTiles(cast<linalg::LinalgOp>(op.getOperation()),
+                                 parallelTileSizes, reductionTileSizes);
 
   TileSizesListType tileSizes;
   tileSizes.emplace_back(flowTileSizes.begin(), flowTileSizes.end());
@@ -1092,8 +1093,8 @@
 
   SmallVector<int64_t> parallelTileSizes = getL1TileSizes();
   SmallVector<int64_t> reductionTileSizes;
-  splitParallelAndReductionTiles(mmt4dOp.getOperation(), parallelTileSizes,
-                                 reductionTileSizes);
+  splitParallelAndReductionTiles(cast<linalg::LinalgOp>(mmt4dOp.getOperation()),
+                                 parallelTileSizes, reductionTileSizes);
 
   TileSizesListType tileSizes = {getWorkgroupTileSizes(), parallelTileSizes,
                                  reductionTileSizes};
@@ -1665,8 +1666,9 @@
 
 static LogicalResult setConvNhwcRootConfigImpl(func::FuncOp entryPointFn,
                                                linalg::LinalgOp convOp) {
-  int64_t vectorSize =
-      getVectorSize(entryPointFn, convOp.getDpsInitOperand(0)->get().getType());
+  int64_t vectorSize = getVectorSize(
+      entryPointFn,
+      cast<ShapedType>(convOp.getDpsInitOperand(0)->get().getType()));
   SmallVector<int64_t> targetTileSizes =
       getConvWorkgroupSizes(entryPointFn, convOp, vectorSize);
   return setConvRootConfig(entryPointFn, convOp, targetTileSizes, vectorSize);
@@ -1706,8 +1708,9 @@
 /// operations.
 static LogicalResult setConvNchwRootConfigImpl(func::FuncOp entryPointFn,
                                                linalg::LinalgOp convOp) {
-  int64_t vectorSize =
-      getVectorSize(entryPointFn, convOp.getDpsInitOperand(0)->get().getType());
+  int64_t vectorSize = getVectorSize(
+      entryPointFn,
+      cast<ShapedType>(convOp.getDpsInitOperand(0)->get().getType()));
   SmallVector<int64_t> targetTileSizes = {1, vectorSize * 2, 1, 8, 8, 1, 1};
   return setConvRootConfig(entryPointFn, convOp, targetTileSizes, vectorSize);
 }
@@ -1731,8 +1734,8 @@
 /// operations.
 static LogicalResult setRootConfig(func::FuncOp entryPointFn,
                                    linalg::DepthwiseConv2DNhwcHwcOp convOp) {
-  int64_t vectorSize =
-      getVectorSize(entryPointFn, convOp.getResult(0).getType());
+  int64_t vectorSize = getVectorSize(
+      entryPointFn, cast<ShapedType>(convOp.getResult(0).getType()));
   SmallVector<int64_t> targetTileSizes =
       getConvWorkgroupSizes(entryPointFn, convOp, vectorSize);
   return setConvRootConfig(entryPointFn, convOp, targetTileSizes, vectorSize);
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerToUKernels.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerToUKernels.cpp
index 30e5a63..7e98448 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerToUKernels.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerToUKernels.cpp
@@ -160,7 +160,7 @@
   // If the pack op didn't have a padding_value attribute, default to 0.
   if (!paddingVal) {
     paddingVal =
-        rewriter.create<arith::ConstantOp>(loc, rewriter.getZeroAttr(i64), i64);
+        rewriter.create<arith::ConstantOp>(loc, i64, rewriter.getZeroAttr(i64));
   }
   int paddingValBitWidth = paddingVal.getType().getIntOrFloatBitWidth();
   // Non-integer element types get bitcast to integer of same bit width.
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPUSplitReduction.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPUSplitReduction.cpp
index 1f9e161..76151fb 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPUSplitReduction.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPUSplitReduction.cpp
@@ -124,8 +124,8 @@
 
   // 2) Apply splitReduction on the single vector-length array.
   // splitReduction already replaces the op.
-  FailureOr<linalg::SplitReductionResult> splitRes =
-      splitReduction(rewriter, tileResFirst->tiledOps.back(), fn);
+  FailureOr<linalg::SplitReductionResult> splitRes = splitReduction(
+      rewriter, cast<linalg::LinalgOp>(tileResFirst->tiledOps.back()), fn);
   if (failed(splitRes)) {
     LLVM_DEBUG(llvm::dbgs() << "failed on step 2 (SplitReduction)\n");
     return success();
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPUTileAndFuse.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPUTileAndFuse.cpp
index e98be0c..e4e40c0 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPUTileAndFuse.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPUTileAndFuse.cpp
@@ -77,7 +77,7 @@
   SmallVector<OpResult> yieldedValuesToOrigValues;
   SmallVector<Operation *> tiledOps;
   FailureOr<scf::SCFTilingResult> tilingResult =
-      scf::tileUsingSCFForOp(rewriter, rootOp, options);
+      scf::tileUsingSCFForOp(rewriter, cast<TilingInterface>(rootOp), options);
   if (failed(tilingResult)) {
     return failure();
   }
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPUVectorLowering.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPUVectorLowering.cpp
index 634559d..6d11daa 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPUVectorLowering.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPUVectorLowering.cpp
@@ -39,7 +39,7 @@
   auto funcOp = getOperation();
 
   // Per-function lowering pipeline.
-  auto vectorTransposeLowering = vector::VectorTransposeLowering::Shuffle;
+  auto vectorTransposeLowering = vector::VectorTransposeLowering::Shuffle1D;
   auto vectorMultiReductionLowering =
       vector::VectorMultiReductionLowering::InnerReduction;
   auto vectorContractLowering = vector::VectorContractLowering::OuterProduct;
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/VectorContractCustomKernels.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/VectorContractCustomKernels.cpp
index c14f97c..4c80a35 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/VectorContractCustomKernels.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/VectorContractCustomKernels.cpp
@@ -941,7 +941,7 @@
     VectorType flatAccVectorType =
         VectorType::get({accType.getNumElements()}, accType.getElementType());
     ;
-    Attribute resultInitializer;
+    TypedAttr resultInitializer;
     if (accElemType.isSignlessInteger()) {
       resultInitializer = DenseIntElementsAttr::get(flatAccVectorType, 0);
     } else if (accElemType.isF32()) {
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/AMDConfig.cpp b/compiler/src/iree/compiler/Codegen/SPIRV/AMDConfig.cpp
index cfa6c97..0306b85 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/AMDConfig.cpp
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/AMDConfig.cpp
@@ -89,7 +89,8 @@
     const int multipler = 32 / bitwidth;
     bool hasPaddedInput = convOp.image().getDefiningOp<tensor::PadOp>();
     const int bestTilingFactor = (hasPaddedInput ? 16 : 32) * multipler;
-    return setConvOpConfig(rootOp, subgroupSize, bestTilingFactor);
+    return setConvOpConfig(cast<linalg::LinalgOp>(rootOp), subgroupSize,
+                           bestTilingFactor);
   }
 
   return failure();
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/AdrenoConfig.cpp b/compiler/src/iree/compiler/Codegen/SPIRV/AdrenoConfig.cpp
index 6c1294c..0ef67bc 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/AdrenoConfig.cpp
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/AdrenoConfig.cpp
@@ -58,7 +58,8 @@
     linalg::detail::ConvolutionDimensions convDims;
     linalg::detail::isConvolutionInterfaceImpl(rootOp, &convDims);
     const int bestTilingFactor = (convDims.depth.empty() ? 32 : 16) * multipler;
-    return setConvOpConfig(rootOp, subgroupSize, bestTilingFactor);
+    return setConvOpConfig(cast<linalg::LinalgOp>(rootOp), subgroupSize,
+                           bestTilingFactor);
   }
 
   return failure();
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/AppleConfig.cpp b/compiler/src/iree/compiler/Codegen/SPIRV/AppleConfig.cpp
index 39fe1d3..a63ab56 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/AppleConfig.cpp
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/AppleConfig.cpp
@@ -56,7 +56,8 @@
     if (bitwidth > 32) return failure();
     const int multipler = 32 / bitwidth;
     const int bestTilingFactor = 16 * multipler;
-    return setConvOpConfig(rootOp, subgroupSize, bestTilingFactor);
+    return setConvOpConfig(cast<linalg::LinalgOp>(rootOp), subgroupSize,
+                           bestTilingFactor);
   }
 
   return failure();
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/MaliConfig.cpp b/compiler/src/iree/compiler/Codegen/SPIRV/MaliConfig.cpp
index 4cdbbe0..dfef1a2 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/MaliConfig.cpp
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/MaliConfig.cpp
@@ -60,7 +60,8 @@
     const int multipler = 32 / bitwidth;
     bool hasPaddedInput = convOp.image().getDefiningOp<tensor::PadOp>();
     const int bestTilingFactor = (hasPaddedInput ? 8 : 16) * multipler;
-    return setConvOpConfig(rootOp, subgroupSize, bestTilingFactor);
+    return setConvOpConfig(cast<linalg::LinalgOp>(rootOp), subgroupSize,
+                           bestTilingFactor);
   }
 
   return failure();
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVTileAndPromote.cpp b/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVTileAndPromote.cpp
index e80ac7e..c0dceac 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVTileAndPromote.cpp
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVTileAndPromote.cpp
@@ -319,7 +319,7 @@
   // If there are no fused elementwise ops, we can avoid promoting C matrix.
   if (linalgOps.size() <= 1) return success();
 
-  linalg::LinalgOp matmulOp = linalgOps.front();
+  auto matmulOp = cast<linalg::LinalgOp>(linalgOps.front());
   auto genericOp = cast<linalg::GenericOp>(*linalgOps.back());
 
   auto matmulType =
diff --git a/compiler/src/iree/compiler/Codegen/Utils/GPUUtils.cpp b/compiler/src/iree/compiler/Codegen/Utils/GPUUtils.cpp
index b13ce0b..b2b96df 100644
--- a/compiler/src/iree/compiler/Codegen/Utils/GPUUtils.cpp
+++ b/compiler/src/iree/compiler/Codegen/Utils/GPUUtils.cpp
@@ -385,7 +385,7 @@
 
 // List of identity elements by operation.
 // https://en.wikipedia.org/wiki/Identity_element
-static Attribute getCombiningKindIdentity(OpBuilder &builder,
+static TypedAttr getCombiningKindIdentity(OpBuilder &builder,
                                           vector::CombiningKind combiningKind,
                                           Type type) {
   switch (combiningKind) {
@@ -419,7 +419,7 @@
       return builder.getFloatAttr(type, negInfApFloat);
     }
   }
-  return Attribute();
+  return TypedAttr();
 }
 
 /// Compute the value on a single thread to get per lane reduction value.
@@ -463,11 +463,11 @@
   } else {
     // In cases where vecSize < unrollCount, we would pad the vector
     // with identity elements until it's total bit size is 32.
-    Attribute identityAttr =
+    TypedAttr identityAttr =
         getCombiningKindIdentity(builder, kind, elementType);
     identityAttr = DenseElementsAttr::get(unrolledLaneValType, identityAttr);
-    Value identity = builder.create<arith::ConstantOp>(loc, identityAttr,
-                                                       unrolledLaneValType);
+    Value identity = builder.create<arith::ConstantOp>(loc, unrolledLaneValType,
+                                                       identityAttr);
     perLaneReduction = builder.create<vector::InsertStridedSliceOp>(
         loc, input, identity, /*offsets=*/ArrayRef<int64_t>{0},
         /*strides=*/ArrayRef<int64_t>{1});
@@ -484,13 +484,13 @@
   if (vectorType) {
     elementType = vectorType.getElementType();
   }
-  Attribute identityAttr = getCombiningKindIdentity(builder, kind, elementType);
+  TypedAttr identityAttr = getCombiningKindIdentity(builder, kind, elementType);
   if (vectorType) {
     identityAttr = DenseElementsAttr::get(vectorType, identityAttr);
   }
   assert(identityAttr && "Unknown identity value for the reduction");
   Value identity =
-      builder.create<arith::ConstantOp>(loc, identityAttr, identityType);
+      builder.create<arith::ConstantOp>(loc, identityType, identityAttr);
   return identity;
 }
 
diff --git a/compiler/src/iree/compiler/ConstEval/JitGlobals.cpp b/compiler/src/iree/compiler/ConstEval/JitGlobals.cpp
index eb01e9f..cae4aa6 100644
--- a/compiler/src/iree/compiler/ConstEval/JitGlobals.cpp
+++ b/compiler/src/iree/compiler/ConstEval/JitGlobals.cpp
@@ -235,7 +235,7 @@
       }
 
       modified = true;
-      targetGlobal.setInitialValueAttr(value);
+      targetGlobal.setInitialValueAttr(cast<TypedAttr>(value));
     }
 
     // Delete any ops noted for pruning.
diff --git a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowDialect.cpp b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowDialect.cpp
index 3165824..03bb15a 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowDialect.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowDialect.cpp
@@ -79,7 +79,7 @@
 Operation *FlowDialect::materializeConstant(OpBuilder &builder, Attribute value,
                                             Type type, Location loc) {
   if (arith::ConstantOp::isBuildableWith(value, type))
-    return builder.create<arith::ConstantOp>(loc, type, value);
+    return builder.create<arith::ConstantOp>(loc, type, cast<TypedAttr>(value));
   return nullptr;
 }
 
diff --git a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOpFolders.cpp b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOpFolders.cpp
index 0e24d89..0571d12 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOpFolders.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOpFolders.cpp
@@ -749,19 +749,20 @@
   if (!value) return {};
   if (auto target = operands.getTarget().dyn_cast_or_null<ElementsAttr>()) {
     // Store into the constant target tensor.
-    if (target.getType().getRank() == 0) {
-      return DenseElementsAttr::get(target.getType(), {value});
+    auto targetType = cast<ShapedType>(target.getType());
+    if (targetType.getRank() == 0) {
+      return DenseElementsAttr::get(targetType, {value});
     }
     if (llvm::count(operands.getIndices(), nullptr) == 0) {
       uint64_t offset = getFlattenedIndex(
-          target.getType(),
+          targetType,
           llvm::to_vector<4>(
               llvm::map_range(operands.getIndices(), [](Attribute value) {
                 return value.cast<IntegerAttr>().getValue().getZExtValue();
               })));
       SmallVector<Attribute, 16> newContents(target.getValues<Attribute>());
       newContents[offset] = value;
-      return DenseElementsAttr::get(target.getType(), newContents);
+      return DenseElementsAttr::get(targetType, newContents);
     }
   }
   return {};
@@ -834,7 +835,8 @@
 // Slices tensor from start to (start + length) exclusively at dim.
 static ElementsAttr tensorSlice(ElementsAttr tensor, uint64_t dim,
                                 uint64_t start, uint64_t length) {
-  auto shape = llvm::to_vector<4>(tensor.getType().getShape());
+  auto tensorType = cast<ShapedType>(tensor.getType());
+  auto shape = llvm::to_vector<4>(tensorType.getShape());
   if (length == shape[dim]) {
     // No need to slice.
     return tensor;
@@ -851,7 +853,7 @@
                       /*init=*/1, /*op=*/std::multiplies<int64_t>());
   int64_t num = length * step / shape[dim];
   for (int64_t offset = step / shape[dim] * start,
-               numElements = tensor.getType().getNumElements();
+               numElements = tensorType.getNumElements();
        offset < numElements; offset += step) {
     newContents.append(valuesBegin + offset, valuesBegin + offset + num);
   }
diff --git a/compiler/src/iree/compiler/Dialect/Flow/Transforms/DetachElementwiseFromNamedOps.cpp b/compiler/src/iree/compiler/Dialect/Flow/Transforms/DetachElementwiseFromNamedOps.cpp
index 27d3f53..969807d 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/Transforms/DetachElementwiseFromNamedOps.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/Transforms/DetachElementwiseFromNamedOps.cpp
@@ -154,7 +154,7 @@
       Type elementType = resultType.getElementType();
       Value emptyTensorOp = rewriter.create<tensor::EmptyOp>(
           loc, resultType.getShape(), elementType);
-      Attribute constValue;
+      TypedAttr constValue;
       if (elementType.isa<IntegerType>()) {
         constValue = rewriter.getIntegerAttr(
             elementType, attr.template getSplatValue<APInt>());
diff --git a/compiler/src/iree/compiler/Dialect/Flow/Transforms/DumpDispatchGraph.cpp b/compiler/src/iree/compiler/Dialect/Flow/Transforms/DumpDispatchGraph.cpp
index b55747d..d5d9200 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/Transforms/DumpDispatchGraph.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/Transforms/DumpDispatchGraph.cpp
@@ -209,9 +209,9 @@
     // Elide "big" elements attributes.
     auto elements = attr.dyn_cast<ElementsAttr>();
     if (elements && elements.getNumElements() > largeAttrLimit) {
-      os << std::string(elements.getType().getRank(), '[') << "..."
-         << std::string(elements.getType().getRank(), ']') << " : "
-         << elements.getType();
+      auto type = cast<ShapedType>(elements.getType());
+      os << std::string(type.getRank(), '[') << "..."
+         << std::string(type.getRank(), ']') << " : " << type;
       return;
     }
 
diff --git a/compiler/src/iree/compiler/Dialect/Flow/Transforms/InitializeEmptyTensors.cpp b/compiler/src/iree/compiler/Dialect/Flow/Transforms/InitializeEmptyTensors.cpp
index d49fd9a..7cc17bb 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/Transforms/InitializeEmptyTensors.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/Transforms/InitializeEmptyTensors.cpp
@@ -20,13 +20,13 @@
 
 /// Returns a zero value attribute based on the `elementType`.
 /// Returns failure, when the type is not handled.
-static FailureOr<Attribute> getZero(OpBuilder &builder, Location loc,
+static FailureOr<TypedAttr> getZero(OpBuilder &builder, Location loc,
                                     Type elementType) {
   if (auto intType = elementType.dyn_cast<IntegerType>()) {
-    return builder.getIntegerAttr(intType, 0);
+    return cast<TypedAttr>(builder.getIntegerAttr(intType, 0));
   }
   if (auto floatType = elementType.dyn_cast<FloatType>()) {
-    return builder.getFloatAttr(floatType, 0.0);
+    return cast<TypedAttr>(builder.getFloatAttr(floatType, 0.0));
   }
   return failure();
 }
@@ -48,7 +48,7 @@
     RankedTensorType resultType = emptyTensorOp.getType();
     Type elementType = resultType.getElementType();
     Location loc = emptyTensorOp.getLoc();
-    FailureOr<Attribute> zero = getZero(rewriter, loc, elementType);
+    FailureOr<TypedAttr> zero = getZero(rewriter, loc, elementType);
     if (failed(zero)) {
       return rewriter.notifyMatchFailure(
           emptyTensorOp, "unable to get zero value for element type");
diff --git a/compiler/src/iree/compiler/Dialect/Flow/Transforms/SetEncoding.cpp b/compiler/src/iree/compiler/Dialect/Flow/Transforms/SetEncoding.cpp
index a8426e1..463d87a 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/Transforms/SetEncoding.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/Transforms/SetEncoding.cpp
@@ -47,17 +47,17 @@
 /// Returns a constant 0 of type `elementType`.
 static FailureOr<Value> getZero(OpBuilder &builder, Location loc,
                                 Type elementType) {
-  Attribute zeroVal =
-      TypeSwitch<Type, Attribute>(elementType)
+  TypedAttr zeroVal =
+      TypeSwitch<Type, TypedAttr>(elementType)
           .Case<FloatType>([&](FloatType floatType) -> Attribute {
-            return builder.getFloatAttr(floatType, 0);
+            return cast<TypedAttr>(builder.getFloatAttr(floatType, 0));
           })
           .Case<IntegerType>([&](IntegerType intType) -> Attribute {
-            return builder.getIntegerAttr(intType, 0);
+            return cast<TypedAttr>(builder.getIntegerAttr(intType, 0));
           })
           .Default([](Type type) { return nullptr; });
   if (!zeroVal) return failure();
-  return builder.create<arith::ConstantOp>(loc, zeroVal, elementType)
+  return builder.create<arith::ConstantOp>(loc, elementType, zeroVal)
       .getResult();
 }
 
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertAllocatorOps.cpp b/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertAllocatorOps.cpp
index 63c742a..2cdcec2 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertAllocatorOps.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertAllocatorOps.cpp
@@ -36,9 +36,9 @@
         ArrayRef<Value>{
             adaptor.getAllocator(),
             rewriter.createOrFold<IREE::VM::ConstI32Op>(
-                op.getLoc(), op.getMemoryTypesAttr()),
+                op.getLoc(), op.getMemoryTypesAttr().getInt()),
             rewriter.createOrFold<IREE::VM::ConstI32Op>(
-                op.getLoc(), op.getBufferUsageAttr()),
+                op.getLoc(), op.getBufferUsageAttr().getInt()),
             adaptor.getSource(),
             castToImportType(adaptor.getOffset(), rewriter.getI64Type(),
                              rewriter),
@@ -76,9 +76,9 @@
             adaptor.getAllocator(),
             rewriter.createOrFold<IREE::VM::ConstI32Op>(op.getLoc(), /*try=*/1),
             rewriter.createOrFold<IREE::VM::ConstI32Op>(
-                op.getLoc(), op.getMemoryTypesAttr()),
+                op.getLoc(), op.getMemoryTypesAttr().getInt()),
             rewriter.createOrFold<IREE::VM::ConstI32Op>(
-                op.getLoc(), op.getBufferUsageAttr()),
+                op.getLoc(), op.getBufferUsageAttr().getInt()),
             adaptor.getSource(),
             castToImportType(adaptor.getOffset(), rewriter.getI64Type(),
                              rewriter),
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertDeviceOps.cpp b/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertDeviceOps.cpp
index 5977ad0..8aafdfd 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertDeviceOps.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertDeviceOps.cpp
@@ -37,7 +37,7 @@
     auto queryOp = rewriter.create<IREE::HAL::DeviceQueryOp>(
         op.getLoc(), rewriter.getI1Type(), rewriter.getI64Type(),
         adaptor.getDevice(), op.getCategoryAttr(), op.getKeyAttr(),
-        Attribute{});
+        TypedAttr{});
     auto ok = queryOp.getOk().cast<Value>();
     auto value = queryOp.getValue();
 
diff --git a/compiler/src/iree/compiler/Dialect/HAL/IR/HALOps.cpp b/compiler/src/iree/compiler/Dialect/HAL/IR/HALOps.cpp
index fd295b4..f36173a 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/IR/HALOps.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/IR/HALOps.cpp
@@ -996,7 +996,8 @@
   ArrayRef<Type> argTypes = getArgumentTypes();
   ArrayRef<Type> resultTypes = getResultTypes();
   mlir::function_interface_impl::printFunctionSignature(
-      p, op, argTypes, /*isVariadic=*/false, resultTypes);
+      p, cast<FunctionOpInterface>(op), argTypes, /*isVariadic=*/false,
+      resultTypes);
   p << " as ";
   if (resultTypes.size() != 1) p << '(';
   llvm::interleaveComma(getKeys().getValue(), p,
diff --git a/compiler/src/iree/compiler/Dialect/HAL/IR/HALTypes.cpp b/compiler/src/iree/compiler/Dialect/HAL/IR/HALTypes.cpp
index 2f8acd7..51bf863 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/IR/HALTypes.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/IR/HALTypes.cpp
@@ -769,10 +769,10 @@
     // Empty returns false (no conditions match).
     return builder.create<arith::ConstantIntOp>(loc, /*value=*/0, /*width=*/1);
   }
-  auto conditionValues =
-      llvm::map_range(getConditions(), [&](MatchAttrInterface attr) {
-        return attr.buildConditionExpression(loc, value, builder);
-      });
+  auto conditionValues = llvm::map_range(getConditions(), [&](Attribute attr) {
+    return attr.cast<MatchAttrInterface>().buildConditionExpression(loc, value,
+                                                                    builder);
+  });
   Value resultValue;
   for (auto conditionValue : conditionValues) {
     resultValue = resultValue ? builder.createOrFold<arith::OrIOp>(
@@ -798,10 +798,10 @@
     // Empty returns true (all 0 conditions match).
     return builder.create<arith::ConstantIntOp>(loc, /*value=*/1, /*width=*/1);
   }
-  auto conditionValues =
-      llvm::map_range(getConditions(), [&](MatchAttrInterface attr) {
-        return attr.buildConditionExpression(loc, value, builder);
-      });
+  auto conditionValues = llvm::map_range(getConditions(), [&](Attribute attr) {
+    return attr.cast<MatchAttrInterface>().buildConditionExpression(loc, value,
+                                                                    builder);
+  });
   Value resultValue;
   for (auto conditionValue : conditionValues) {
     resultValue = resultValue ? builder.createOrFold<arith::AndIOp>(
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Transforms/DumpExecutableBenchmarks.cpp b/compiler/src/iree/compiler/Dialect/HAL/Transforms/DumpExecutableBenchmarks.cpp
index 2b26df2..1e274ea 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Transforms/DumpExecutableBenchmarks.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Transforms/DumpExecutableBenchmarks.cpp
@@ -49,7 +49,7 @@
   // Analyzed minimum binding sizes.
   SmallVector<Binding> bindings;
   // Push constant operands that are known constant. May be null if dynamic.
-  SmallVector<Attribute> uniformOperands;
+  SmallVector<TypedAttr> uniformOperands;
 };
 
 using DispatchParamsMap =
@@ -97,9 +97,9 @@
                             resourceLengthInt.getSExtValue()});
       }
 
-      SmallVector<Attribute> uniformOperands;
+      SmallVector<TypedAttr> uniformOperands;
       for (auto operand : dispatchOp.getUniformOperands()) {
-        Attribute uniformOperand;
+        TypedAttr uniformOperand;
         if (!matchPattern(operand, m_Constant(&uniformOperand))) {
           // Non-constant uniform operand; skip the dispatch.
           // TODO(benvanik): extract information from the executable annotations
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Transforms/MaterializeInterfaces.cpp b/compiler/src/iree/compiler/Dialect/HAL/Transforms/MaterializeInterfaces.cpp
index e85612d..928eb7e 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Transforms/MaterializeInterfaces.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Transforms/MaterializeInterfaces.cpp
@@ -434,8 +434,8 @@
 
     uint64_t dimIdx = sizeOp.getDimension().getZExtValue();
     auto dimAttr = workgroupSizeAttr[dimIdx];
-    rewriter.replaceOpWithNewOp<arith::ConstantOp>(sizeOp, dimAttr,
-                                                   rewriter.getIndexType());
+    rewriter.replaceOpWithNewOp<arith::ConstantOp>(
+        sizeOp, rewriter.getIndexType(), cast<TypedAttr>(dimAttr));
     return success();
   }
 };
diff --git a/compiler/src/iree/compiler/Dialect/Stream/IR/StreamDialect.cpp b/compiler/src/iree/compiler/Dialect/Stream/IR/StreamDialect.cpp
index da53892..559a8a2 100644
--- a/compiler/src/iree/compiler/Dialect/Stream/IR/StreamDialect.cpp
+++ b/compiler/src/iree/compiler/Dialect/Stream/IR/StreamDialect.cpp
@@ -115,7 +115,7 @@
     return builder.create<mlir::func::ConstantOp>(
         loc, type, value.cast<FlatSymbolRefAttr>());
   } else if (arith::ConstantOp::isBuildableWith(value, type)) {
-    return builder.create<arith::ConstantOp>(loc, type, value);
+    return builder.create<arith::ConstantOp>(loc, type, cast<TypedAttr>(value));
   } else if (value.isa<IREE::Stream::TimepointAttr>()) {
     return builder.create<IREE::Stream::TimepointImmediateOp>(loc);
   }
diff --git a/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOpFolders.cpp b/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOpFolders.cpp
index ef44afc..20f9370 100644
--- a/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOpFolders.cpp
+++ b/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOpFolders.cpp
@@ -1037,7 +1037,7 @@
 // to be emulated - if we can avoid that here that's a big win. Some HAL
 // implementations (such as Metal) only support 8-bit fills and anything larger
 // needs to be implemented as well.
-static Attribute tryNarrowPatternBits(Attribute patternAttr) {
+static TypedAttr tryNarrowPatternBits(TypedAttr patternAttr) {
   // Get the old pattern bitcast to an APInt. Splats are bitwise operations
   // and we don't care what the value originally was.
   APInt oldPattern;
@@ -1066,7 +1066,7 @@
   LogicalResult matchAndRewrite(TensorSplatOp splatOp,
                                 PatternRewriter &rewriter) const override {
     // Try narrowing the pattern.
-    Attribute oldPatternAttr;
+    TypedAttr oldPatternAttr;
     if (!matchPattern(splatOp.getValue(), m_Constant(&oldPatternAttr))) {
       return failure();
     }
@@ -1159,7 +1159,7 @@
   LogicalResult matchAndRewrite(TensorFillOp fillOp,
                                 PatternRewriter &rewriter) const override {
     // Try narrowing the pattern.
-    Attribute oldPatternAttr;
+    TypedAttr oldPatternAttr;
     if (!matchPattern(fillOp.getValue(), m_Constant(&oldPatternAttr))) {
       return failure();
     }
diff --git a/compiler/src/iree/compiler/Dialect/Stream/Transforms/SpecializeDispatches.cpp b/compiler/src/iree/compiler/Dialect/Stream/Transforms/SpecializeDispatches.cpp
index 0b72cfa..e3c8b0f 100644
--- a/compiler/src/iree/compiler/Dialect/Stream/Transforms/SpecializeDispatches.cpp
+++ b/compiler/src/iree/compiler/Dialect/Stream/Transforms/SpecializeDispatches.cpp
@@ -95,7 +95,7 @@
     SmallVector<TypedAttr> values;
     for (auto dispatchOp : dispatchOps) {
       auto operand = dispatchOp.getUniformOperands()[idx];
-      Attribute constantValue;
+      TypedAttr constantValue;
       matchPattern(operand, m_Constant(&constantValue));
       values.push_back(constantValue);
       set.locs.insert(operand.getLoc());
@@ -114,7 +114,7 @@
 // Builds a tensor<SITExOPERANDxTYPE> constant attribute.
 // This should probably be vector<> but that dialect has some issues with
 // expressing basic multi-dimension loads :/
-static Attribute buildConstantSetAttr(ConstantSet &set, OpBuilder &builder) {
+static TypedAttr buildConstantSetAttr(ConstantSet &set, OpBuilder &builder) {
   // TODO(benvanik): better definition of variable-width integers across HAL.
   // HACK: we can't handle index types in a few of the codegen backends (vulkan
   // at least); we convert index -> i32 here but we should probably have a
diff --git a/compiler/src/iree/compiler/Dialect/Util/IR/UtilAttrs.cpp b/compiler/src/iree/compiler/Dialect/Util/IR/UtilAttrs.cpp
index 79327cf..ca73e02 100644
--- a/compiler/src/iree/compiler/Dialect/Util/IR/UtilAttrs.cpp
+++ b/compiler/src/iree/compiler/Dialect/Util/IR/UtilAttrs.cpp
@@ -472,8 +472,9 @@
           SerializableDenseElementsAttrModel, DenseIntOrFPElementsAttr> {
   int64_t getStorageSize(Attribute baseAttr) const {
     auto attr = baseAttr.cast<ElementsAttr>();
-    return attr.getNumElements() * IREE::Util::getRoundedElementByteWidth(
-                                       attr.getType().getElementType());
+    return attr.getNumElements() *
+           IREE::Util::getRoundedElementByteWidth(
+               cast<ShapedType>(attr.getType()).getElementType());
   }
 
   LogicalResult serializeToVector(Attribute baseAttr,
diff --git a/compiler/src/iree/compiler/Dialect/Util/IR/UtilDialect.cpp b/compiler/src/iree/compiler/Dialect/Util/IR/UtilDialect.cpp
index 7c38095..e890cce 100644
--- a/compiler/src/iree/compiler/Dialect/Util/IR/UtilDialect.cpp
+++ b/compiler/src/iree/compiler/Dialect/Util/IR/UtilDialect.cpp
@@ -96,7 +96,7 @@
 Operation *UtilDialect::materializeConstant(OpBuilder &builder, Attribute value,
                                             Type type, Location loc) {
   if (arith::ConstantOp::isBuildableWith(value, type)) {
-    return builder.create<arith::ConstantOp>(loc, value, type);
+    return builder.create<arith::ConstantOp>(loc, type, cast<TypedAttr>(value));
   }
   return nullptr;
 }
diff --git a/compiler/src/iree/compiler/Dialect/Util/IR/UtilInterfaces.td b/compiler/src/iree/compiler/Dialect/Util/IR/UtilInterfaces.td
index e230d6b..3b9d0fd 100644
--- a/compiler/src/iree/compiler/Dialect/Util/IR/UtilInterfaces.td
+++ b/compiler/src/iree/compiler/Dialect/Util/IR/UtilInterfaces.td
@@ -756,7 +756,7 @@
   }];
 
   let verify = [{
-    return IREE::Util::detail::verifyTiedOp($_op);
+    return IREE::Util::detail::verifyTiedOp(cast<TiedOpInterface>($_op));
   }];
 }
 
diff --git a/compiler/src/iree/compiler/Dialect/Util/IR/UtilOpFolders.cpp b/compiler/src/iree/compiler/Dialect/Util/IR/UtilOpFolders.cpp
index 2be2f10..27a823a 100644
--- a/compiler/src/iree/compiler/Dialect/Util/IR/UtilOpFolders.cpp
+++ b/compiler/src/iree/compiler/Dialect/Util/IR/UtilOpFolders.cpp
@@ -138,9 +138,8 @@
     }
     if (constantValue != initialValue) {
       operands.insert(rewriter.create<arith::ConstantOp>(
-          op.getLoc(),
-          rewriter.getIntegerAttr(op.getResult().getType(), constantValue),
-          op.getResult().getType()));
+          op.getLoc(), op.getResult().getType(),
+          rewriter.getIntegerAttr(op.getResult().getType(), constantValue)));
     }
     rewriter.replaceOpWithNewOp<OpT>(op, op.getResult().getType(),
                                      operands.takeVector());
@@ -176,7 +175,7 @@
   return makeRangeEnd(
       loc, offset, length,
       builder.create<arith::ConstantOp>(
-          loc, builder.getIntegerAttr(offset.getType(), 1), offset.getType()),
+          loc, offset.getType(), builder.getIntegerAttr(offset.getType(), 1)),
       builder);
 }
 
@@ -224,14 +223,12 @@
     // Min/max with constant ranges. This allows for normal folding to happen
     // downstream of the op.
     auto constantMinOp = rewriter.create<arith::ConstantOp>(
-        op.getLoc(),
-        rewriter.getIntegerAttr(op.getMin().getType(), constantMin),
-        op.getMin().getType());
+        op.getLoc(), op.getMin().getType(),
+        rewriter.getIntegerAttr(op.getMin().getType(), constantMin));
     auto constantMaxOp = rewriter.create<arith::ConstantOp>(
-        op.getLoc(),
+        op.getLoc(), op.getMax().getType(),
         rewriter.getIntegerAttr(op.getMax().getType(),
-                                constantMax - constantMin + 1),
-        op.getMax().getType());
+                                constantMax - constantMin + 1));
     min = min ? rewriter.create<arith::MinUIOp>(op.getLoc(), min, constantMinOp)
                     .getResult()
               : constantMinOp.getResult();
@@ -260,8 +257,8 @@
       minValue = rewriter.create<arith::MinUIOp>(loc, op.getOffsets().front(),
                                                  op.getOffsets().back());
       auto one = rewriter.create<arith::ConstantOp>(
-          loc, rewriter.getIntegerAttr(op.getMin().getType(), 1),
-          op.getMin().getType());
+          loc, op.getMin().getType(),
+          rewriter.getIntegerAttr(op.getMin().getType(), 1));
       auto endLhs = makeRangeEnd(loc, op.getOffsets().front(),
                                  op.getLengths().front(), one, rewriter);
       auto endRhs = makeRangeEnd(loc, op.getOffsets().back(),
@@ -416,8 +413,8 @@
   using OpRewritePattern<IREE::Util::UnfoldableConstantOp>::OpRewritePattern;
   LogicalResult matchAndRewrite(UnfoldableConstantOp op,
                                 PatternRewriter &rewriter) const override {
-    auto stdConst =
-        rewriter.create<arith::ConstantOp>(op.getLoc(), op.getValue());
+    auto stdConst = rewriter.create<arith::ConstantOp>(
+        op.getLoc(), cast<TypedAttr>(op.getValue()));
     rewriter.replaceOpWithNewOp<OptimizationBarrierOp>(op,
                                                        stdConst.getResult());
     return success();
diff --git a/compiler/src/iree/compiler/Dialect/Util/Transforms/FoldGlobals.cpp b/compiler/src/iree/compiler/Dialect/Util/Transforms/FoldGlobals.cpp
index 851a750..7dafc4b 100644
--- a/compiler/src/iree/compiler/Dialect/Util/Transforms/FoldGlobals.cpp
+++ b/compiler/src/iree/compiler/Dialect/Util/Transforms/FoldGlobals.cpp
@@ -259,7 +259,7 @@
                                     OpBuilder &builder) {
   if (arith::ConstantOp::isBuildableWith(attr, type)) {
     // Common case fast-path.
-    return builder.create<arith::ConstantOp>(loc, type, attr);
+    return builder.create<arith::ConstantOp>(loc, type, cast<TypedAttr>(attr));
   } else if (mlir::func::ConstantOp::isBuildableWith(attr, type)) {
     return builder.create<mlir::func::ConstantOp>(
         loc, type, attr.cast<FlatSymbolRefAttr>());
diff --git a/compiler/src/iree/compiler/Dialect/Util/Transforms/IPO.cpp b/compiler/src/iree/compiler/Dialect/Util/Transforms/IPO.cpp
index d65feb4..759fbc5 100644
--- a/compiler/src/iree/compiler/Dialect/Util/Transforms/IPO.cpp
+++ b/compiler/src/iree/compiler/Dialect/Util/Transforms/IPO.cpp
@@ -363,8 +363,9 @@
   // themselves.
   if (arith::ConstantOp::isBuildableWith(constantValue.attr,
                                          constantValue.type)) {
-    op = builder.create<arith::ConstantOp>(
-        constantValue.loc.value(), constantValue.attr, constantValue.type);
+    op = builder.create<arith::ConstantOp>(constantValue.loc.value(),
+                                           constantValue.type,
+                                           cast<TypedAttr>(constantValue.attr));
   }
 
   // Try the attr and type dialects to see if they can materialize.
diff --git a/compiler/src/iree/compiler/Dialect/VM/Conversion/ImportUtils.cpp b/compiler/src/iree/compiler/Dialect/VM/Conversion/ImportUtils.cpp
index b69aa55..f462bfa 100644
--- a/compiler/src/iree/compiler/Dialect/VM/Conversion/ImportUtils.cpp
+++ b/compiler/src/iree/compiler/Dialect/VM/Conversion/ImportUtils.cpp
@@ -137,7 +137,8 @@
     elementValues.reserve(elementsAttr.getNumElements());
     for (auto intAttr : elementsAttr.getValues<Attribute>()) {
       elementValues.push_back(rewriter.createOrFold<mlir::arith::ConstantOp>(
-          loc, elementsAttr.getType().getElementType(), intAttr));
+          loc, elementsAttr.getType().getElementType(),
+          cast<TypedAttr>(intAttr)));
     }
     return elementValues;
   }
diff --git a/compiler/src/iree/compiler/Dialect/VM/IR/VMBase.td b/compiler/src/iree/compiler/Dialect/VM/IR/VMBase.td
index d773b5d..0567a57 100644
--- a/compiler/src/iree/compiler/Dialect/VM/IR/VMBase.td
+++ b/compiler/src/iree/compiler/Dialect/VM/IR/VMBase.td
@@ -115,7 +115,7 @@
 class VM_EncTypeOf<string name> : VM_EncEncodeExpr<
     "e.encodeType({0}())", [name]>;
 class VM_EncPrimitiveAttr<string name, int thisBitwidth> : VM_EncEncodeExpr<
-    "e.encodePrimitiveAttr(getOperation()->getAttrOfType<Attribute>(\"" # name # "\"))"> {
+    "e.encodePrimitiveAttr(getOperation()->getAttrOfType<TypedAttr>(\"" # name # "\"))"> {
   int bitwidth = thisBitwidth;
 }
 class VM_EncPrimitiveArrayAttr<string name, int thisBitwidth> : VM_EncEncodeExpr<
diff --git a/compiler/src/iree/compiler/Dialect/VM/IR/VMDialect.cpp b/compiler/src/iree/compiler/Dialect/VM/IR/VMDialect.cpp
index d919419..fa38364 100644
--- a/compiler/src/iree/compiler/Dialect/VM/IR/VMDialect.cpp
+++ b/compiler/src/iree/compiler/Dialect/VM/IR/VMDialect.cpp
@@ -261,26 +261,29 @@
 
 Operation *VMDialect::materializeConstant(OpBuilder &builder, Attribute value,
                                           Type type, Location loc) {
-  if (ConstI32Op::isBuildableWith(value, type)) {
-    auto convertedValue = ConstI32Op::convertConstValue(value);
+  auto typedValue = dyn_cast<TypedAttr>(value);
+  if (!typedValue) return nullptr;
+
+  if (ConstI32Op::isBuildableWith(typedValue, type)) {
+    auto convertedValue = ConstI32Op::convertConstValue(typedValue);
     if (convertedValue.cast<IntegerAttr>().getValue() == 0) {
       return builder.create<VM::ConstI32ZeroOp>(loc);
     }
     return builder.create<VM::ConstI32Op>(loc, convertedValue);
-  } else if (ConstI64Op::isBuildableWith(value, type)) {
-    auto convertedValue = ConstI64Op::convertConstValue(value);
+  } else if (ConstI64Op::isBuildableWith(typedValue, type)) {
+    auto convertedValue = ConstI64Op::convertConstValue(typedValue);
     if (convertedValue.cast<IntegerAttr>().getValue() == 0) {
       return builder.create<VM::ConstI64ZeroOp>(loc);
     }
     return builder.create<VM::ConstI64Op>(loc, convertedValue);
-  } else if (ConstF32Op::isBuildableWith(value, type)) {
-    auto convertedValue = ConstF32Op::convertConstValue(value);
+  } else if (ConstF32Op::isBuildableWith(typedValue, type)) {
+    auto convertedValue = ConstF32Op::convertConstValue(typedValue);
     if (convertedValue.cast<FloatAttr>().getValue().isZero()) {
       return builder.create<VM::ConstF32ZeroOp>(loc);
     }
     return builder.create<VM::ConstF32Op>(loc, convertedValue);
-  } else if (ConstF64Op::isBuildableWith(value, type)) {
-    auto convertedValue = ConstF64Op::convertConstValue(value);
+  } else if (ConstF64Op::isBuildableWith(typedValue, type)) {
+    auto convertedValue = ConstF64Op::convertConstValue(typedValue);
     if (convertedValue.cast<FloatAttr>().getValue().isZero()) {
       return builder.create<VM::ConstF64ZeroOp>(loc);
     }
diff --git a/compiler/src/iree/compiler/Dialect/VM/IR/VMOpFolders.cpp b/compiler/src/iree/compiler/Dialect/VM/IR/VMOpFolders.cpp
index 7cce235..d288c5a 100644
--- a/compiler/src/iree/compiler/Dialect/VM/IR/VMOpFolders.cpp
+++ b/compiler/src/iree/compiler/Dialect/VM/IR/VMOpFolders.cpp
@@ -486,7 +486,7 @@
     return DenseElementsAttr::get(operand.getType(), elementResult);
   } else if (auto operand = rawOperand.dyn_cast_or_null<ElementsAttr>()) {
     return operand.cast<DenseIntOrFPElementsAttr>().mapValues(
-        operand.getType().getElementType(),
+        cast<ShapedType>(operand.getType()).getElementType(),
         llvm::function_ref<ElementValueT(const ElementValueT &)>(
             [&](const ElementValueT &value) { return calculate(value); }));
   }
@@ -506,7 +506,7 @@
     return DenseElementsAttr::get(operand.getType(), elementResult);
   } else if (auto operand = rawOperand.dyn_cast_or_null<ElementsAttr>()) {
     return operand.cast<DenseIntOrFPElementsAttr>().mapValues(
-        operand.getType().getElementType(),
+        cast<ShapedType>(operand.getType()).getElementType(),
         llvm::function_ref<APInt(const APFloat &)>([&](const APFloat &value) {
           return calculate(value).bitcastToAPInt();
         }));
@@ -521,7 +521,7 @@
           class ElementValueT = typename AttrElementT::ValueType,
           class CalculationT =
               std::function<ElementValueT(ElementValueT, ElementValueT)>>
-static Attribute constFoldBinaryOp(Attribute rawLhs, Attribute rawRhs,
+static TypedAttr constFoldBinaryOp(Attribute rawLhs, Attribute rawRhs,
                                    const CalculationT &calculate) {
   if (auto lhs = rawLhs.dyn_cast_or_null<AttrElementT>()) {
     auto rhs = rawRhs.dyn_cast_or_null<AttrElementT>();
@@ -550,7 +550,7 @@
       ++lhsIt;
       ++rhsIt;
     }
-    return DenseElementsAttr::get(lhs.getType(), resultAttrs);
+    return DenseElementsAttr::get(cast<ShapedType>(lhs.getType()), resultAttrs);
   }
   return {};
 }
@@ -602,7 +602,7 @@
       ++bIt;
       ++cIt;
     }
-    return DenseElementsAttr::get(a.getType(), resultAttrs);
+    return DenseElementsAttr::get(cast<ShapedType>(a.getType()), resultAttrs);
   }
   return {};
 }
@@ -2128,7 +2128,7 @@
           class ElementValueT = typename AttrElementT::ValueType,
           class CalculationT =
               std::function<ElementValueT(ElementValueT, ElementValueT)>>
-static Attribute constFoldBinaryCmpFOp(Attribute rawLhs, Attribute rawRhs,
+static TypedAttr constFoldBinaryCmpFOp(Attribute rawLhs, Attribute rawRhs,
                                        const CalculationT &calculate) {
   if (auto lhs = rawLhs.dyn_cast_or_null<AttrElementT>()) {
     auto rhs = rawRhs.dyn_cast_or_null<AttrElementT>();
@@ -2143,8 +2143,9 @@
         lhs.getSplatValue<Attribute>(), rhs.getSplatValue<Attribute>(),
         calculate);
     if (!elementResult) return {};
-    return DenseElementsAttr::get(IntegerType::get(lhs.getContext(), 32),
-                                  elementResult);
+    auto resultType = lhs.getType().clone(
+        std::nullopt, IntegerType::get(lhs.getContext(), 32));
+    return DenseElementsAttr::get(resultType, elementResult);
   } else if (auto lhs = rawLhs.dyn_cast_or_null<ElementsAttr>()) {
     auto rhs = rawRhs.dyn_cast_or_null<ElementsAttr>();
     if (!rhs || lhs.getType() != rhs.getType()) return {};
@@ -2158,8 +2159,9 @@
       ++lhsIt;
       ++rhsIt;
     }
-    return DenseElementsAttr::get(IntegerType::get(lhs.getContext(), 32),
-                                  resultAttrs);
+    auto resultType = lhs.getShapedType().clone(
+        std::nullopt, IntegerType::get(lhs.getContext(), 32));
+    return DenseElementsAttr::get(resultType, resultAttrs);
   }
   return {};
 }
diff --git a/compiler/src/iree/compiler/Dialect/VM/IR/VMOps.cpp b/compiler/src/iree/compiler/Dialect/VM/IR/VMOps.cpp
index 5647d62..34d4a4d 100644
--- a/compiler/src/iree/compiler/Dialect/VM/IR/VMOps.cpp
+++ b/compiler/src/iree/compiler/Dialect/VM/IR/VMOps.cpp
@@ -501,7 +501,7 @@
   } else if (auto intAttr = value.dyn_cast<IntegerAttr>()) {
     return intAttr.getType().isInteger(SZ);
   } else if (auto elementsAttr = value.dyn_cast<ElementsAttr>()) {
-    return elementsAttr.getType().getElementType().isInteger(SZ);
+    return elementsAttr.getShapedType().getElementType().isInteger(SZ);
   }
   return false;
 }
@@ -520,14 +520,14 @@
   if (auto floatAttr = value.dyn_cast<FloatAttr>()) {
     elementType = floatAttr.getType();
   } else if (auto elementsAttr = value.dyn_cast<ElementsAttr>()) {
-    elementType = elementsAttr.getType().getElementType();
+    elementType = elementsAttr.getShapedType().getElementType();
   }
   if (!elementType) return false;
   return elementType.getIntOrFloatBitWidth() == SZ;
 }
 
 template <int SZ>
-static Attribute convertConstIntegerValue(TypedAttr value) {
+static TypedAttr convertConstIntegerValue(TypedAttr value) {
   assert(isConstIntegerBuildableWith<SZ>(value, value.getType()));
   Builder builder(value.getContext());
   auto integerType = builder.getIntegerType(SZ);
@@ -552,7 +552,7 @@
     }
   }
   assert(false && "unexpected attribute type");
-  return Attribute();
+  return TypedAttr();
 }
 
 static FloatType getFloatType(int bitwidth, MLIRContext *context) {
@@ -570,7 +570,7 @@
 }
 
 template <int SZ>
-static Attribute convertConstFloatValue(TypedAttr value) {
+static TypedAttr convertConstFloatValue(TypedAttr value) {
   assert(isConstFloatBuildableWith<SZ>(value, value.getType()));
   Builder builder(value.getContext());
   auto floatType = getFloatType(SZ, value.getContext());
@@ -589,7 +589,7 @@
     }
   }
   assert(false && "unexpected attribute type");
-  return Attribute();
+  return TypedAttr();
 }
 
 // static
diff --git a/compiler/src/iree/compiler/Dialect/VM/Target/Bytecode/BytecodeEncoder.cpp b/compiler/src/iree/compiler/Dialect/VM/Target/Bytecode/BytecodeEncoder.cpp
index 8256031..3d6b954 100644
--- a/compiler/src/iree/compiler/Dialect/VM/Target/Bytecode/BytecodeEncoder.cpp
+++ b/compiler/src/iree/compiler/Dialect/VM/Target/Bytecode/BytecodeEncoder.cpp
@@ -153,7 +153,7 @@
         failed(writeUint16(value.getNumElements()))) {
       return currentOp_->emitOpError() << "integer array size out of bounds";
     }
-    for (auto el : value.getValues<Attribute>()) {
+    for (auto el : value.getValues<TypedAttr>()) {
       if (failed(encodePrimitiveAttr(el))) {
         return currentOp_->emitOpError() << "failed to encode element " << el;
       }
diff --git a/compiler/src/iree/compiler/InputConversion/Common/QuantizedConvToConv.cpp b/compiler/src/iree/compiler/InputConversion/Common/QuantizedConvToConv.cpp
index def3ede..ad2836e 100644
--- a/compiler/src/iree/compiler/InputConversion/Common/QuantizedConvToConv.cpp
+++ b/compiler/src/iree/compiler/InputConversion/Common/QuantizedConvToConv.cpp
@@ -41,7 +41,7 @@
   Value empty =
       builder.create<tensor::EmptyOp>(ty.getShape(), ty.getElementType(), dyn);
 
-  Attribute attr = builder.getZeroAttr(ty.getElementType());
+  TypedAttr attr = builder.getZeroAttr(ty.getElementType());
   Value cnst = builder.create<arith::ConstantOp>(attr);
   return builder.create<linalg::FillOp>(ValueRange{cnst}, ValueRange{empty})
       .result();
diff --git a/compiler/src/iree/compiler/InputConversion/MHLO/ConvertMHLOToLinalgExt.cpp b/compiler/src/iree/compiler/InputConversion/MHLO/ConvertMHLOToLinalgExt.cpp
index 07a5948..974e04e 100644
--- a/compiler/src/iree/compiler/InputConversion/MHLO/ConvertMHLOToLinalgExt.cpp
+++ b/compiler/src/iree/compiler/InputConversion/MHLO/ConvertMHLOToLinalgExt.cpp
@@ -455,7 +455,7 @@
     Value emptyTensorOutputIndices = rewriter.create<mlir::tensor::EmptyOp>(
         loc, mixedSizes, indicesElementType);
     // Initialize indices to 0 and values to negative infinity
-    Attribute negInfAttr;
+    TypedAttr negInfAttr;
     if (auto intType = valueElementType.dyn_cast<IntegerType>()) {
       negInfAttr = rewriter.getIntegerAttr(
           intType, APInt::getSignedMinValue(intType.getWidth()));
@@ -466,7 +466,7 @@
       negInfAttr = rewriter.getFloatAttr(valueElementType, negApFloat);
     }
     Value negInf = rewriter.create<arith::ConstantOp>(loc, negInfAttr);
-    Attribute posInfAttr = rewriter.getIntegerAttr(
+    TypedAttr posInfAttr = rewriter.getIntegerAttr(
         indicesElementType, APInt::getSignedMaxValue(32));
     Value posInf = rewriter.create<arith::ConstantOp>(loc, posInfAttr);
     Value negInfTensor =
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/MapStableHLOToScalarOp.h b/compiler/src/iree/compiler/InputConversion/StableHLO/MapStableHLOToScalarOp.h
index 43e3592..2a78aa1 100644
--- a/compiler/src/iree/compiler/InputConversion/StableHLO/MapStableHLOToScalarOp.h
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/MapStableHLOToScalarOp.h
@@ -311,7 +311,7 @@
   if (VectorType vecType = t.dyn_cast<VectorType>()) {
     v = SplatElementsAttr::get(vecType, v);
   }
-  return b->create<arith::ConstantOp>(loc, t, v);
+  return b->create<arith::ConstantOp>(loc, t, cast<TypedAttr>(v));
 }
 
 template <typename PredicateType>
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/StableHLOToLinalg.cpp b/compiler/src/iree/compiler/InputConversion/StableHLO/StableHLOToLinalg.cpp
index 07b470e..b797f98 100644
--- a/compiler/src/iree/compiler/InputConversion/StableHLO/StableHLOToLinalg.cpp
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/StableHLOToLinalg.cpp
@@ -1558,8 +1558,8 @@
     }
 
     SmallVector<OpFoldResult, 3> startIndices, sizes;
-    Type originalStartIndexType =
-        dynamicSliceOp.getStartIndices().front().getType();
+    auto originalStartIndexType =
+        dynamicSliceOp.getStartIndices().front().getType().cast<ShapedType>();
     for (const auto& en : llvm::enumerate(
              llvm::zip(adaptor.getStartIndices(),
                        dynamicSliceOp.getSliceSizes().getValues<int64_t>()))) {
@@ -1636,9 +1636,9 @@
       // By mhlo.DynamicUpdateSlice definition:
       //   `start_indices[i] = clamp(start_indices[i],
       //       0, operand.dimension_size[i] - update.dimension_size[i])`
-      Value startIndex =
-          extractIndexFromTensor(rewriter, loc, en.value(),
-                                 op.getStartIndices()[en.index()].getType());
+      Value startIndex = extractIndexFromTensor(
+          rewriter, loc, en.value(),
+          cast<ShapedType>(op.getStartIndices()[en.index()].getType()));
       Value ub = rewriter.create<arith::ConstantIndexOp>(
           loc, operandType.getDimSize(en.index()) -
                    updateType.getDimSize(en.index()));
@@ -1724,7 +1724,7 @@
     for (Value operand : llvm::drop_begin(adaptor.getOperands(), 1)) {
       coercedOperands.push_back(coerceTensorShape(
           rewriter, loc, cast<TypedValue<ShapedType>>(operand),
-          operand0.getType()));
+          cast<ShapedType>(operand0.getType())));
     }
     Value output = rewriter.create<tensor::EmptyOp>(
         loc, tensor::getMixedSizes(rewriter, loc, operand0),
@@ -2208,8 +2208,8 @@
 
     // We have interior padding, which can be lowered to tensor.insert_slice.
     // Start by filling a result-sized tensor with the pad value.
-    auto emptyTensor =
-        getEmptyTensorFor(rewriter, loc, resultType, op, adaptor.getOperands());
+    auto emptyTensor = getEmptyTensorFor(
+        rewriter, loc, cast<ShapedType>(resultType), op, adaptor.getOperands());
     auto fill =
         rewriter.create<linalg::FillOp>(loc, paddingVal, emptyTensor).result();
 
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/StableHLOToLinalgPointwise.cpp b/compiler/src/iree/compiler/InputConversion/StableHLO/StableHLOToLinalgPointwise.cpp
index 589415e..e0d100b 100644
--- a/compiler/src/iree/compiler/InputConversion/StableHLO/StableHLOToLinalgPointwise.cpp
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/StableHLOToLinalgPointwise.cpp
@@ -120,7 +120,7 @@
       if (getRank(input) == maxRank) {
         mappedInputs.push_back(coerceTensorShape(
             rewriter, loc, cast<TypedValue<ShapedType>>(input),
-            emptyTensor.getType()));
+            cast<ShapedType>(emptyTensor.getType())));
         scalarInputs.push_back(nullptr);
       } else {
         scalarInputs.push_back(rewriter.create<tensor::ExtractOp>(loc, input));
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/StableHLOToLinalgReduce.cpp b/compiler/src/iree/compiler/InputConversion/StableHLO/StableHLOToLinalgReduce.cpp
index a909b21..fb1b8fc 100644
--- a/compiler/src/iree/compiler/InputConversion/StableHLO/StableHLOToLinalgReduce.cpp
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/StableHLOToLinalgReduce.cpp
@@ -111,8 +111,9 @@
       initValue = rewriter.createOrFold<tensor::ExtractOp>(loc, initValue);
 
       SmallVector<Value, 8> dynShape = getReduceOpEmptyTensorDynSizes(
-          rewriter, loc, operand, resultType, reductionDims);
-      auto emptyTensor = getEmptyTensor(rewriter, loc, resultType, dynShape);
+          rewriter, loc, operand, cast<ShapedType>(resultType), reductionDims);
+      auto emptyTensor =
+          getEmptyTensor(rewriter, loc, cast<ShapedType>(resultType), dynShape);
       Value filledTensor =
           rewriter.create<linalg::FillOp>(loc, initValue, emptyTensor).result();
       outputs.push_back(filledTensor);
diff --git a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/IR/LinalgExtOps.td b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/IR/LinalgExtOps.td
index 82b304b..fa1ef18 100644
--- a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/IR/LinalgExtOps.td
+++ b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/IR/LinalgExtOps.td
@@ -617,12 +617,12 @@
 
     // Return the output type.
     ShapedType getOutputType() {
-      return getOutput().getType();
+      return getOutput().getType().cast<ShapedType>();
     }
 
     // Return the input type.
     ShapedType getInputType() {
-      return getInput().getType();
+      return getInput().getType().cast<ShapedType>();
     }
 
     // Return the output shape.
@@ -764,12 +764,12 @@
 
     // Return the output type.
     ShapedType getOutputType() {
-      return getOutput().getType();
+      return getOutput().getType().cast<ShapedType>();
     }
 
     // Return the input type.
     ShapedType getInputType() {
-      return getInput().getType();
+      return getInput().getType().cast<ShapedType>();
     }
 
     // Return the output shape.
diff --git a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Transforms/Transforms.h b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Transforms/Transforms.h
index 812f0df..70aba5c 100644
--- a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Transforms/Transforms.h
+++ b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Transforms/Transforms.h
@@ -263,7 +263,7 @@
     // needs more investigation.
     rewriter.startRootUpdate(op);
     std::optional<linalg::LinalgOp> promotedOp =
-        promoteSubViews(rewriter, op, options);
+        promoteSubViews(rewriter, cast<linalg::LinalgOp>(op), options);
     if (!promotedOp) {
       rewriter.cancelRootUpdate(op);
       return op->emitError("subview promotion failed");
diff --git a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Passes/ConvertConv2DToWinograd.cpp b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Passes/ConvertConv2DToWinograd.cpp
index b1163b8..0680be4 100644
--- a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Passes/ConvertConv2DToWinograd.cpp
+++ b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Passes/ConvertConv2DToWinograd.cpp
@@ -62,7 +62,7 @@
 /// TODO: Codegen this as a kernel and run once at initialization
 static DenseElementsAttr
 foldFilterTransform(ArrayRef<int64_t> shape, int64_t inputTileSize,
-                    int64_t kernelSize, Type outputType, const float *G,
+                    int64_t kernelSize, ShapedType outputType, const float *G,
                     bool isSplat, float splatValue,
                     DenseElementsAttr::iterator_range<APFloat> &input,
                     FloatType floatType, bool isNchw) {
diff --git a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Passes/SplitReduction.cpp b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Passes/SplitReduction.cpp
index c99d4e8..e8a3807 100644
--- a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Passes/SplitReduction.cpp
+++ b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Passes/SplitReduction.cpp
@@ -149,7 +149,7 @@
 
   // Initialize indices to positive infinity and values to negative infinity
   // for a top (maxk) comparison.
-  Attribute negInfAttr;
+  TypedAttr negInfAttr;
   if (auto intType = valueElementType.dyn_cast<IntegerType>()) {
     negInfAttr = rewriter.getIntegerAttr(
         intType, APInt::getSignedMinValue(intType.getWidth()));
@@ -160,7 +160,7 @@
     negInfAttr = rewriter.getFloatAttr(valueElementType, negApFloat);
   }
   Value negInf = rewriter.create<arith::ConstantOp>(loc, negInfAttr);
-  Attribute posInfAttr =
+  TypedAttr posInfAttr =
       rewriter.getIntegerAttr(indicesElementType, APInt::getSignedMaxValue(32));
   Value posInf = rewriter.create<arith::ConstantOp>(loc, posInfAttr);
   Value negInfTensor =
diff --git a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/TransformOps/LinalgExtTransformOps.cpp b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/TransformOps/LinalgExtTransformOps.cpp
index 3f38e01..6723b97 100644
--- a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/TransformOps/LinalgExtTransformOps.cpp
+++ b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/TransformOps/LinalgExtTransformOps.cpp
@@ -53,7 +53,8 @@
     // Apply the pattern.
     SimplePatternRewriter rewriter(target);
     FailureOr<LinalgExt::FusionResult> result =
-        pattern.returningMatchAndRewrite(target, rewriter);
+        pattern.returningMatchAndRewrite(cast<TilingInterface>(target),
+                                         rewriter);
     if (failed(result))
       return emitDefaultDefiniteFailure(target);
 
diff --git a/llvm-external-projects/iree-dialects/test/Dialect/linalg_transform/single-tiling-full-script.mlir b/llvm-external-projects/iree-dialects/test/Dialect/linalg_transform/single-tiling-full-script.mlir
index 25ee7bb..177b43e 100644
--- a/llvm-external-projects/iree-dialects/test/Dialect/linalg_transform/single-tiling-full-script.mlir
+++ b/llvm-external-projects/iree-dialects/test/Dialect/linalg_transform/single-tiling-full-script.mlir
@@ -33,7 +33,7 @@
     lowering_strategy = "outerproduct"
       : (!pdl.operation) -> !pdl.operation
   %func_e_3 = transform.vector.lower_transpose %func_e_2
-    lowering_strategy = "shuffle"
+    lowering_strategy = "shuffle_1d"
       : (!pdl.operation) -> !pdl.operation
 
   lower_to_llvm %module_op1 : (!pdl.operation) -> !pdl.operation
diff --git a/third_party/llvm-project b/third_party/llvm-project
index 5e70d1a..be9c918 160000
--- a/third_party/llvm-project
+++ b/third_party/llvm-project
@@ -1 +1 @@
-Subproject commit 5e70d1adf85fcdd695f5d796cb0617a8fdf90f82
+Subproject commit be9c91843bab5bb46574c27836bfcd9ad6fc9ef5
diff --git a/third_party/mlir-hlo b/third_party/mlir-hlo
index 34c53ad..4d28523 160000
--- a/third_party/mlir-hlo
+++ b/third_party/mlir-hlo
@@ -1 +1 @@
-Subproject commit 34c53ad70a3cd1a64f04aff8627badc5310a7970
+Subproject commit 4d28523ba6ad600f50900f9048b06056276bf9db