diff --git a/build_tools/cmake/iree_copts.cmake b/build_tools/cmake/iree_copts.cmake
index 5b36433..10342b4 100644
--- a/build_tools/cmake/iree_copts.cmake
+++ b/build_tools/cmake/iree_copts.cmake
@@ -135,11 +135,12 @@
     # https://docs.microsoft.com/en-us/cpp/build/reference/bigobj-increase-number-of-sections-in-dot-obj-file
     "/bigobj"
 
-    # Use the modern C preprocessor to more closely match standards/clang/gcc behavior.
-    "/Zc:preprocessor"
-
     # Enable C11 standards conforming mode.
     "$<$<COMPILE_LANGUAGE:C>:/std:c11>"
+
+  MSVC
+    # Use the modern C preprocessor to more closely match standards/clang/gcc behavior.
+    "/Zc:preprocessor"
 )
 
 # Compiler diagnostics.
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/GPUDistributionPatterns.cpp b/compiler/src/iree/compiler/Codegen/Common/GPU/GPUDistributionPatterns.cpp
index 70529bf..351f22d 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GPU/GPUDistributionPatterns.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/GPU/GPUDistributionPatterns.cpp
@@ -24,22 +24,6 @@
 
 namespace {
 
-/// Given a LayoutAttr, find the shape of the given layout dimension. It is
-/// expected that the layout has at most one instance of the requested
-/// dimension. Example:
-///   LayoutAttr: <<BATCHX: 4>, <BATCHY: 4, LANEX: 4>>
-///   dim: BATCHX
-///   output: 4
-static std::optional<int64_t> findDimShape(LayoutAttr layout,
-                                           LayoutDimension dim) {
-  for (PerDimLayoutAttr dimLayout : layout.getLayouts()) {
-    if (std::optional<int64_t> shape = dimLayout.getShape(dim)) {
-      return shape;
-    }
-  }
-  return std::nullopt;
-}
-
 /// Given the state of the iterator, compute the indices of the original vector
 /// that the current iterator state is iterating over. These indices are
 /// parameterized by the thread grid.
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
index ab2f351..ac844fc 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
@@ -695,8 +695,7 @@
 /// pattern.
 struct RewriteFuncOpABI : public OpRewritePattern<LLVM::LLVMFuncOp> {
   RewriteFuncOpABI(HALDispatchABI &abi, LLVMTypeConverter &typeConverter)
-      : OpRewritePattern(&typeConverter.getContext()), abi(abi),
-        typeConverter(typeConverter) {}
+      : OpRewritePattern(&typeConverter.getContext()), abi(abi) {}
 
   LogicalResult matchAndRewrite(LLVM::LLVMFuncOp funcOp,
                                 PatternRewriter &rewriter) const override {
@@ -742,7 +741,6 @@
 
 private:
   HALDispatchABI &abi;
-  LLVMTypeConverter &typeConverter;
 };
 
 /// Lower call ops with specified ABI. The ABI to use is looked up from the
@@ -754,8 +752,7 @@
 /// pattern.
 struct RewriteCallOpABI : public OpRewritePattern<LLVM::CallOp> {
   RewriteCallOpABI(HALDispatchABI &abi, LLVMTypeConverter &typeConverter)
-      : OpRewritePattern(&typeConverter.getContext()), abi(abi),
-        typeConverter(typeConverter) {}
+      : OpRewritePattern(&typeConverter.getContext()), abi(abi) {}
 
   LogicalResult matchAndRewrite(LLVM::CallOp callOp,
                                 PatternRewriter &rewriter) const override {
@@ -790,7 +787,6 @@
 
 private:
   HALDispatchABI &abi;
-  LLVMTypeConverter &typeConverter;
 };
 
 /// Rewrites calls to extern functions to dynamic library import calls.
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/DispatchABI.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/DispatchABI.cpp
index f72e936..dd06283 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/DispatchABI.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/DispatchABI.cpp
@@ -72,7 +72,7 @@
 }
 
 ExecutableLibraryDI::ExecutableLibraryDI(const LLVMTypeConverter *typeConverter)
-    : typeConverter(typeConverter), builder(&typeConverter->getContext()) {
+    : builder(&typeConverter->getContext()) {
   auto *context = builder.getContext();
   fileAttr = LLVM::DIFileAttr::get(
       context, "runtime/src/iree/hal/local/executable_library.h", ".");
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/DispatchABI.h b/compiler/src/iree/compiler/Codegen/LLVMCPU/DispatchABI.h
index 73a4e8f..85297e5 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/DispatchABI.h
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/DispatchABI.h
@@ -95,7 +95,6 @@
   LLVM::DIDerivedTypeAttr getWorkgroupStateV0T();
 
 private:
-  const LLVMTypeConverter *typeConverter;
   Builder builder;
   LLVM::DIFileAttr fileAttr;
   unsigned ptrBitwidth;
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
index 79571b3..70186e1 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
@@ -182,14 +182,6 @@
                            [&](Range r) { return getStaticValue(r.size); });
 }
 
-/// Returns true if all the input and output tensor operands of 'op' are fully
-/// dynamic.
-static bool isFullyDynamicOp(linalg::LinalgOp op) {
-  SmallVector<int64_t> loopRanges = op.getStaticLoopRanges();
-  return llvm::all_of(loopRanges,
-                      [](int64_t size) { return ShapedType::isDynamic(size); });
-}
-
 /// Returns the vectorization pre-processing strategy (peeling, masking) for the
 /// given LinalgOp. It is based on either (in the priority order):
 ///   * user-specified value, or
@@ -983,11 +975,6 @@
   return inputAndOutputElementTypes[0];
 }
 
-static void getFullRegisterHeuristicsMatmulVectorSizes(
-    mlir::FunctionOpInterface entryPointFn, linalg::LinalgOp op,
-    int64_t vectorSize, SmallVectorImpl<int64_t> &sizes,
-    SmallVectorImpl<bool> &scalableSizeFlags) {}
-
 /// Compute or adjust existing vector sizes using a generic heuristic that will
 /// aim to fill at least one full vector register for all the element types of
 /// the matmul. For now, the current heuristics only look at the N dimension but
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerExecutableTarget.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerExecutableTarget.cpp
index b380043..028abe4 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerExecutableTarget.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerExecutableTarget.cpp
@@ -63,22 +63,6 @@
 };
 } // namespace
 
-/// The pipeline parser doesnt like strings that have `'` or `"` in them. But it
-/// is needed for demarcating the option value. So just drop them before sending
-/// it one.
-static StringRef sanitizePipelineString(StringRef input) {
-  if (input.empty())
-    return input;
-  // If first/last character is ' or ", drop them.
-  if (input.front() == '\'' || input.front() == '"') {
-    input = input.drop_front();
-  }
-  if (input.back() == '\'' || input.back() == '"') {
-    input = input.drop_back();
-  }
-  return input;
-}
-
 /// Verify that valid configuration is set for all ops within the compiled
 /// module.
 template <typename F>
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVTileAndPromote.cpp b/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVTileAndPromote.cpp
index 3259ec6..32c0663 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVTileAndPromote.cpp
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVTileAndPromote.cpp
@@ -35,8 +35,6 @@
 
 #define DEBUG_TYPE "iree-spirv-tile-and-promote"
 
-constexpr int kMaxVectorNumBits = 128;
-
 namespace mlir::iree_compiler {
 
 //====---------------------------------------------------------------------===//
diff --git a/compiler/src/iree/compiler/Codegen/Utils/LinalgOpInfo.h b/compiler/src/iree/compiler/Codegen/Utils/LinalgOpInfo.h
index eecd1c6..1bfb661 100644
--- a/compiler/src/iree/compiler/Codegen/Utils/LinalgOpInfo.h
+++ b/compiler/src/iree/compiler/Codegen/Utils/LinalgOpInfo.h
@@ -38,7 +38,6 @@
   void computeInfo(linalg::LinalgOp);
 
   TransposeMapFilter transposeMapFilter;
-  bool transposeTrait;
   bool reductionTrait;
   bool dynamicTrait;
   SmallVector<OpOperand *> transposeOperands;
diff --git a/compiler/src/iree/compiler/Dialect/Flow/Transforms/DeduplicateExecutables.cpp b/compiler/src/iree/compiler/Dialect/Flow/Transforms/DeduplicateExecutables.cpp
index f35c477..f3376be 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/Transforms/DeduplicateExecutables.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/Transforms/DeduplicateExecutables.cpp
@@ -29,15 +29,6 @@
                                    SymbolOpInterface leafOp) {
   return nestSymbolRef(baseRefAttr, FlatSymbolRefAttr::get(leafOp));
 }
-static SymbolRefAttr nestSymbolRef(SymbolOpInterface baseOp,
-                                   FlatSymbolRefAttr leafRefAttr) {
-  return nestSymbolRef(SymbolRefAttr::get(baseOp), leafRefAttr);
-}
-static SymbolRefAttr nestSymbolRef(SymbolOpInterface baseOp,
-                                   SymbolOpInterface leafOp) {
-  return nestSymbolRef(SymbolRefAttr::get(baseOp),
-                       FlatSymbolRefAttr::get(leafOp));
-}
 
 // Recursively gathers symbol->symbol replacements from the old object table
 // regions to the new object table regions into |symbolReplacements|.
diff --git a/compiler/src/iree/compiler/Dialect/Flow/Transforms/DumpDispatchGraph.cpp b/compiler/src/iree/compiler/Dialect/Flow/Transforms/DumpDispatchGraph.cpp
index 79185d3..01820cd 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/Transforms/DumpDispatchGraph.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/Transforms/DumpDispatchGraph.cpp
@@ -37,7 +37,6 @@
 static const StringRef kLineStyleDataFlow = "solid";
 static const StringRef kShapeNode = "box";
 static const StringRef kShapeBox = "box";
-static const StringRef kShapeTab = "tab";
 static const StringRef kShapeNone = "plain";
 static const StringRef kShapeEllipse = "ellipse";
 
diff --git a/compiler/src/iree/compiler/Dialect/Flow/Transforms/FormDispatchRegions.cpp b/compiler/src/iree/compiler/Dialect/Flow/Transforms/FormDispatchRegions.cpp
index fadfa37..2ceda0d 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/Transforms/FormDispatchRegions.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/Transforms/FormDispatchRegions.cpp
@@ -240,22 +240,6 @@
   return isa<IREE::LinalgExt::SetEncodingOp, tensor::PackOp>(op);
 }
 
-/// Returns the source of the pack-like operation.
-// TODO(ravishankarm): This seems like a use case for an interface.
-static Value getSourceOfPackLikeOp(Operation *op) {
-  return TypeSwitch<Operation *, Value>(op)
-      .Case<tensor::PackOp>([](auto packOp) { return packOp.getSource(); })
-      .Case<IREE::LinalgExt::SetEncodingOp>(
-          [](auto setEncodingOp) { return setEncodingOp.getSource(); })
-      .Default([](Operation *) { return nullptr; });
-}
-static RankedTensorType getSourceTypeOfPackLikeOp(Operation *op) {
-  Value source = getSourceOfPackLikeOp(op);
-  if (!source)
-    return nullptr;
-  return llvm::cast<RankedTensorType>(source.getType());
-}
-
 /// Returns true if the operation is an `unpack` op or an `unset_encoding` op,
 /// or an `extract_slice` op whose source operand matches those criteria,
 /// recursively.
diff --git a/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOpFolders.cpp b/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOpFolders.cpp
index 1201b4e..81b9be9 100644
--- a/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOpFolders.cpp
+++ b/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOpFolders.cpp
@@ -348,102 +348,6 @@
   }
 };
 
-// Materialize copy-on-write (🐄) ops where required for |rootValue|.
-// Only valid in tensor/async ops - don't use with stream.cmd.*.
-static bool materializeCOW(Location loc, Value rootValue, OpBuilder &builder) {
-  auto valueType =
-      llvm::dyn_cast<IREE::Stream::ResourceType>(rootValue.getType());
-  if (!valueType)
-    return false;
-
-  // If our rootValue is a constant then we need to ensure that we aren't
-  // tied to a constant operand. If we are we need to clone to a
-  // non-constant value.
-  bool forceClone = valueType.getLifetime() == IREE::Stream::Lifetime::Constant;
-
-  // Identify if we need to insert a copy-on-write clone.
-  // We do this per use as a single consuming op may use the result of this
-  // multiple times - some tied and some not - and if it has it tied several
-  // times each will need its own clone.
-  struct TiedUse {
-    Operation *user;
-    unsigned operandIndex;
-    Value value;
-  };
-  SmallVector<TiedUse> tiedUses;
-  unsigned untiedUses = 0;
-  for (auto &use : rootValue.getUses()) {
-    if (isa<IREE::Stream::TimepointAwaitOp>(use.getOwner()))
-      continue;
-    auto tiedOp = dyn_cast<IREE::Util::TiedOpInterface>(use.getOwner());
-    bool isTied = tiedOp && tiedOp.isOperandTied(use.getOperandNumber());
-    if (isTied) {
-      tiedUses.push_back({use.getOwner(), use.getOperandNumber(), rootValue});
-    } else {
-      ++untiedUses;
-    }
-  }
-  if (tiedUses.empty()) {
-    // All uses are as normal capturing SSA values.
-    return false;
-  } else if (tiedUses.size() == 1 && untiedUses == 0 && !forceClone) {
-    // Only one use and it's tied - we've already reserved our results for it.
-    return false;
-  }
-
-  // Mixed/multiple tied uses. Clone for each tied use but leave the untied
-  // ones referencing us.
-  IREE::Stream::AffinityAttr sourceAffinity;
-  if (auto affinityOp = dyn_cast_or_null<IREE::Stream::AffinityOpInterface>(
-          rootValue.getDefiningOp())) {
-    sourceAffinity = affinityOp.getAffinity();
-  }
-  for (auto &tiedUse : tiedUses) {
-    auto cloneLoc =
-        FusedLoc::get(builder.getContext(), {loc, tiedUse.user->getLoc()});
-
-    builder.setInsertionPoint(tiedUse.user);
-
-    auto sizeAwareType =
-        llvm::cast<IREE::Util::SizeAwareTypeInterface>(tiedUse.value.getType());
-    auto targetSize =
-        sizeAwareType.queryValueSize(cloneLoc, tiedUse.value, builder);
-
-    IREE::Stream::AffinityAttr targetAffinity;
-    if (auto affinityOp =
-            dyn_cast<IREE::Stream::AffinityOpInterface>(tiedUse.user)) {
-      targetAffinity = affinityOp.getAffinity();
-    }
-
-    auto cloneOp = builder.create<IREE::Stream::AsyncCloneOp>(
-        cloneLoc, tiedUse.value.getType(), tiedUse.value, targetSize,
-        targetSize, targetAffinity ? targetAffinity : sourceAffinity);
-    tiedUse.user->setOperand(tiedUse.operandIndex, cloneOp.getResult());
-  }
-
-  return true;
-}
-
-// Materialize copy-on-write (🐄) ops where required.
-// This models what a runtime normally does with copy-on-write but uses the
-// information we have in the SSA use-def chain to identify ties that write and
-// covering reads.
-template <typename Op>
-struct MaterializeCOW : public OpRewritePattern<Op> {
-  using OpRewritePattern<Op>::OpRewritePattern;
-  LogicalResult matchAndRewrite(Op op,
-                                PatternRewriter &rewriter) const override {
-    bool didChange = false;
-
-    // Handle results of this op (primary use case).
-    for (auto result : op->getResults()) {
-      didChange = materializeCOW(op.getLoc(), result, rewriter) || didChange;
-    }
-
-    return didChange ? success() : failure();
-  }
-};
-
 // Ties the results of execution region to their operands when the region
 // operations are tied throughout the entire body.
 //
diff --git a/runtime/src/iree/base/internal/math.h b/runtime/src/iree/base/internal/math.h
index d801bda..c3c29f9 100644
--- a/runtime/src/iree/base/internal/math.h
+++ b/runtime/src/iree/base/internal/math.h
@@ -68,7 +68,7 @@
 //==============================================================================
 
 static inline int iree_math_count_leading_zeros_u32(const uint32_t n) {
-#if defined(IREE_COMPILER_MSVC)
+#if defined(IREE_COMPILER_MSVC_COMPAT)
   unsigned long result = 0;  // NOLINT(runtime/int)
   if (_BitScanReverse(&result, n)) {
     return (int)(31 - result);
@@ -97,7 +97,7 @@
 }
 
 static inline int iree_math_count_leading_zeros_u64(uint64_t n) {
-#if defined(IREE_COMPILER_MSVC) && \
+#if defined(IREE_COMPILER_MSVC_COMPAT) && \
     (defined(IREE_ARCH_ARM_64) || defined(IREE_ARCH_X86_64))
   // MSVC does not have __buitin_clzll. Use _BitScanReverse64.
   unsigned long result = 0;  // NOLINT(runtime/int)
@@ -105,7 +105,7 @@
     return (int)(63 - result);
   }
   return 64;
-#elif defined(IREE_COMPILER_MSVC)
+#elif defined(IREE_COMPILER_MSVC_COMPAT)
   // MSVC does not have __buitin_clzll. Compose two calls to _BitScanReverse
   unsigned long result = 0;  // NOLINT(runtime/int)
   if ((n >> 32) && _BitScanReverse(&result, n >> 32)) {
@@ -141,7 +141,7 @@
 }
 
 static inline int iree_math_count_trailing_zeros_u32(uint32_t n) {
-#if defined(IREE_COMPILER_MSVC)
+#if defined(IREE_COMPILER_MSVC_COMPAT)
   unsigned long result = 0;  // NOLINT(runtime/int)
   _BitScanForward(&result, n);
   return (int)result;
@@ -160,11 +160,11 @@
 }
 
 static inline int iree_math_count_trailing_zeros_u64(uint64_t n) {
-#if defined(IREE_COMPILER_MSVC) && defined(IREE_PTR_SIZE_64)
+#if defined(IREE_COMPILER_MSVC_COMPAT) && defined(IREE_PTR_SIZE_64)
   unsigned long result = 0;  // NOLINT(runtime/int)
   _BitScanForward64(&result, n);
   return (int)result;
-#elif defined(IREE_COMPILER_MSVC) && defined(IREE_PTR_SIZE_32)
+#elif defined(IREE_COMPILER_MSVC_COMPAT) && defined(IREE_PTR_SIZE_32)
   unsigned long result = 0;  // NOLINT(runtime/int)
   if ((uint32_t)(n) == 0) {
     _BitScanForward(&result, n >> 32);
diff --git a/runtime/src/iree/base/internal/synchronization.c b/runtime/src/iree/base/internal/synchronization.c
index ccb42a4..cc484b4 100644
--- a/runtime/src/iree/base/internal/synchronization.c
+++ b/runtime/src/iree/base/internal/synchronization.c
@@ -61,7 +61,7 @@
 // Cross-platform processor yield (where supported)
 //==============================================================================
 
-#if defined(IREE_COMPILER_MSVC)
+#if defined(IREE_COMPILER_MSVC_COMPAT)
 
 // MSVC uses architecture-specific intrinsics.
 
diff --git a/runtime/src/iree/base/target_platform.h b/runtime/src/iree/base/target_platform.h
index 08310f7..a2beb8c 100644
--- a/runtime/src/iree/base/target_platform.h
+++ b/runtime/src/iree/base/target_platform.h
@@ -34,6 +34,7 @@
 // IREE_COMPILER_GCC
 // IREE_COMPILER_GCC_COMPAT
 // IREE_COMPILER_MSVC
+// IREE_COMPILER_MSVC_COMPAT
 //
 // IREE_SANITIZER_ADDRESS
 // IREE_SANITIZER_MEMORY
@@ -197,12 +198,18 @@
 
 #if defined(__clang__)
 #define IREE_COMPILER_CLANG 1
+#if defined(_MSC_VER)
+// clang-cl is msvc-like (but also still clang).
+#define IREE_COMPILER_MSVC_COMPAT 1
+#else
 #define IREE_COMPILER_GCC_COMPAT 1
+#endif  // _MSC_VER
 #elif defined(__GNUC__)
 #define IREE_COMPILER_GCC 1
 #define IREE_COMPILER_GCC_COMPAT 1
 #elif defined(_MSC_VER)
 #define IREE_COMPILER_MSVC 1
+#define IREE_COMPILER_MSVC_COMPAT 1
 #else
 #error Unrecognized compiler.
 #endif  // compiler versions
diff --git a/runtime/src/iree/builtins/ukernel/arch/x86_64/CMakeLists.txt b/runtime/src/iree/builtins/ukernel/arch/x86_64/CMakeLists.txt
index 18c8274..419afef 100644
--- a/runtime/src/iree/builtins/ukernel/arch/x86_64/CMakeLists.txt
+++ b/runtime/src/iree/builtins/ukernel/arch/x86_64/CMakeLists.txt
@@ -155,7 +155,7 @@
     "-mavx2"
     "-mfma"
     "-mf16c"
-  MSVC
+  MSVC_OR_CLANG_CL
     "/arch:AVX2"
 )
 
@@ -183,7 +183,7 @@
     "-mavx512cd"
     "-mavx512bw"
     "-mavx512dq"
-  MSVC
+  MSVC_OR_CLANG_CL
     "/arch:AVX512"
 )
 
@@ -192,7 +192,8 @@
 iree_select_compiler_opts(IREE_UK_COPTS_X86_64_AVX512_VNNI_RELATIVE
   CLANG_OR_GCC
     "-mavx512vnni"
-  MSVC
+  CLANG_CL
+    "/clang:-mavx512vnni"
 )
 set(IREE_UK_COPTS_X86_64_AVX512_VNNI
   "${IREE_UK_COPTS_X86_64_AVX512_BASE}"
@@ -204,7 +205,8 @@
 iree_select_compiler_opts(IREE_UK_COPTS_X86_64_AVX512_BF16_RELATIVE
   CLANG_OR_GCC
     "-mavx512bf16"
-  MSVC
+  CLANG_CL
+    "/clang:-mavx512bf16"
 )
 set(IREE_UK_COPTS_X86_64_AVX512_BF16
   "${IREE_UK_COPTS_X86_64_AVX512_BASE}"
@@ -224,8 +226,8 @@
 # /arch: flag is often too coarse-grained to be meaningful.
 
 # GCC version check
-if ((CMAKE_C_COMPILER_ID STREQUAL GNU) AND
-    (CMAKE_C_COMPILER_VERSION VERSION_LESS 12))
+if((CMAKE_C_COMPILER_ID STREQUAL GNU) AND
+   (CMAKE_C_COMPILER_VERSION VERSION_LESS 12))
   # Old GCC versions have incompatible x86 intrinsics. Supporting them
   # is not considered worth it. By not defining these tokens, we leave these
   # code paths out. At least GCC 9 is known to be problematic, while GCC 12 is
@@ -237,27 +239,79 @@
 endif()  # GCC version check
 
 # MSVC version check for AVX-512-BF16
-if (MSVC_VERSION AND ("${MSVC_VERSION}" VERSION_LESS 1937))
+if(MSVC_VERSION AND ("${MSVC_VERSION}" VERSION_LESS 1937))
   # Missing _mm512_cvtpbh_ps intrinsic at _MSC_VER=1930.
   set(IREE_UK_TRY_X86_64_AVX512_BF16 OFF)
 endif()  # MSVC version check for AVX-512-BF16
 
+# clang-cl version check for vnni bf16 bug.
+# Version 16-17 crash compiling the file and in clang-cl we can't use the
+# inline asm workaround: https://github.com/llvm/llvm-project/issues/68117.
+if((CMAKE_C_COMPILER_ID MATCHES "Clang" AND
+    CMAKE_C_SIMULATE_ID MATCHES "MSVC") AND
+   (CMAKE_C_COMPILER_VERSION VERSION_LESS 18))
+  set(IREE_UK_TRY_X86_64_AVX512_BF16 OFF)
+endif()
+
 # Now check compiler support for what we've decided to try.
+# Some instructions are not available with baseline arch flags and need to be
+# tested via compilation. We include the intrinsics used by the kernels whose
+# presence indicates availability of all required intrinsics (vs testing each).
 
-if (IREE_UK_TRY_X86_64_AVX2_FMA)
-  check_cxx_compiler_flag("${IREE_UK_COPTS_X86_64_AVX2_FMA}" IREE_UK_BUILD_X86_64_AVX2_FMA)
+if(IREE_UK_TRY_X86_64_AVX2_FMA)
+  check_cxx_compiler_flag(
+    "${IREE_UK_COPTS_X86_64_AVX2_FMA}"
+    IREE_UK_BUILD_X86_64_AVX2_FMA
+  )
+else()
+  set(IREE_UK_BUILD_X86_64_AVX2_FMA OFF)
 endif()
 
-if (IREE_UK_TRY_X86_64_AVX512_BASE)
-  check_cxx_compiler_flag("${IREE_UK_COPTS_X86_64_AVX512_BASE}" IREE_UK_BUILD_X86_64_AVX512_BASE)
+if(IREE_UK_TRY_X86_64_AVX512_BASE)
+  check_cxx_compiler_flag(
+    "${IREE_UK_COPTS_X86_64_AVX512_BASE}"
+    IREE_UK_BUILD_X86_64_AVX512_BASE
+  )
+else()
+  set(IREE_UK_BUILD_X86_64_AVX512_BASE OFF)
 endif()
 
-if (IREE_UK_TRY_X86_64_AVX512_VNNI)
-  check_cxx_compiler_flag("${IREE_UK_COPTS_X86_64_AVX512_VNNI}" IREE_UK_BUILD_X86_64_AVX512_VNNI)
+if(IREE_UK_TRY_X86_64_AVX512_VNNI)
+  string(REPLACE ";" " " CMAKE_REQUIRED_FLAGS "${IREE_UK_COPTS_X86_64_AVX512_VNNI}")
+  string(JOIN "\n" IREE_UK_BUILD_X86_64_AVX512_VNNI_TEST
+    "#include <immintrin.h>"
+    "int main() {"
+    "  __m512i a, b, d;"
+    "  _mm512_dpwssd_epi32(d, a, b);"
+    "  return 0;"
+    "}"
+  )
+  check_c_source_compiles(
+    "${IREE_UK_BUILD_X86_64_AVX512_VNNI_TEST}"
+    IREE_UK_BUILD_X86_64_AVX512_VNNI
+  )
+  unset(CMAKE_REQUIRED_FLAGS)
+else()
+  set(IREE_UK_BUILD_X86_64_AVX512_VNNI OFF)
 endif()
 
-if (IREE_UK_TRY_X86_64_AVX512_BF16)
-  check_cxx_compiler_flag("${IREE_UK_COPTS_X86_64_AVX512_BF16}" IREE_UK_BUILD_X86_64_AVX512_BF16)
+if(IREE_UK_TRY_X86_64_AVX512_BF16)
+  string(REPLACE ";" " " CMAKE_REQUIRED_FLAGS "${IREE_UK_COPTS_X86_64_AVX512_BF16}")
+  string(JOIN "\n" IREE_UK_BUILD_X86_64_AVX512_BF16_TEST
+    "#include <immintrin.h>"
+    "int main() {"
+    "  __m256bh a;"
+    "  _mm512_cvtneps_pbh(a);"
+    "  return 0;"
+    "}"
+  )
+  check_c_source_compiles(
+    "${IREE_UK_BUILD_X86_64_AVX512_BF16_TEST}"
+    IREE_UK_BUILD_X86_64_AVX512_BF16
+  )
+  unset(CMAKE_REQUIRED_FLAGS)
+else()
+  set(IREE_UK_BUILD_X86_64_AVX512_BF16 OFF)
 endif()
 
 # Now generate the configured header. This needs to happen after all
diff --git a/runtime/src/iree/builtins/ukernel/arch/x86_64/mmt4d_x86_64_avx512_bf16.c b/runtime/src/iree/builtins/ukernel/arch/x86_64/mmt4d_x86_64_avx512_bf16.c
index 81ec2f7..f4f0fa4 100644
--- a/runtime/src/iree/builtins/ukernel/arch/x86_64/mmt4d_x86_64_avx512_bf16.c
+++ b/runtime/src/iree/builtins/ukernel/arch/x86_64/mmt4d_x86_64_avx512_bf16.c
@@ -7,7 +7,7 @@
 #include "iree/builtins/ukernel/arch/x86_64/common_x86_64.h"
 #include "iree/builtins/ukernel/arch/x86_64/mmt4d_x86_64_internal.h"
 
-#ifdef IREE_UK_COMPILER_CLANG
+#if defined(IREE_UK_COMPILER_CLANG) && !defined(IREE_UK_COMPILER_MSVC)
 // This inline-asm function is a work-around for:
 // 1. https://github.com/llvm/llvm-project/issues/68117
 //    Summary: LLVM crash affecting Clang 16-17. Fixed in Clang 18.
@@ -38,7 +38,7 @@
       acc, bitcast_16xf32_to_32xbf16(lhs),
       bitcast_16xf32_to_32xbf16(_mm512_set1_ps(*(const float*)rhs)));
 }
-#endif
+#endif  // IREE_UK_COMPILER_CLANG
 
 static inline void
 iree_uk_mmt4d_tile_bf16bf16fXX_1x16x2_to_16x16x2_x86_64_avx512_bf16(
diff --git a/runtime/src/iree/io/formats/irpa/irpa_parser.c b/runtime/src/iree/io/formats/irpa/irpa_parser.c
index 04fac9c..dd09b3a 100644
--- a/runtime/src/iree/io/formats/irpa/irpa_parser.c
+++ b/runtime/src/iree/io/formats/irpa/irpa_parser.c
@@ -40,8 +40,8 @@
   }
   iree_io_physical_offset_t view_offset =
       base_offset + header->metadata_segment.offset + range.offset;
-  *out_view =
-      iree_make_string_view(file_contents.data + view_offset, range.length);
+  *out_view = iree_make_string_view(
+      (const char*)file_contents.data + view_offset, range.length);
   return iree_ok_status();
 }
 
diff --git a/runtime/src/iree/modules/hal/exports.inl b/runtime/src/iree/modules/hal/exports.inl
index e61638a..8100bd4 100644
--- a/runtime/src/iree/modules/hal/exports.inl
+++ b/runtime/src/iree/modules/hal/exports.inl
@@ -80,7 +80,7 @@
 
 EXPORT_FN("fence.await", iree_hal_module_fence_await, iCrD, i)
 EXPORT_FN("fence.create", iree_hal_module_fence_create, ri, r)
-EXPORT_FN("fence.fail", iree_hal_module_fence_signal, ri, v)
+EXPORT_FN("fence.fail", iree_hal_module_fence_fail, ri, v)
 EXPORT_FN("fence.join", iree_hal_module_fence_join, CrD, r)
 EXPORT_FN("fence.query", iree_hal_module_fence_query, r, i)
 EXPORT_FN("fence.signal", iree_hal_module_fence_signal, r, v)
diff --git a/runtime/src/iree/task/topology_win32.c b/runtime/src/iree/task/topology_win32.c
index b74285c..50cf48c 100644
--- a/runtime/src/iree/task/topology_win32.c
+++ b/runtime/src/iree/task/topology_win32.c
@@ -165,13 +165,13 @@
     return iree_make_status(
         iree_status_code_from_win32_error(GetLastError()),
         "failed to query logical processor information size (%08X)",
-        GetLastError());
+        (unsigned)GetLastError());
   }
   if (cache_relationships_size > 64 * 1024) {
     return iree_make_status(IREE_STATUS_RESOURCE_EXHAUSTED,
                             "logical processor information size overflow (got "
                             "%u which is large for a stack alloc)",
-                            cache_relationships_size);
+                            (unsigned)cache_relationships_size);
   }
   SYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX* cache_relationships =
       (SYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX*)iree_alloca(
@@ -182,7 +182,8 @@
                                         &cache_relationships_size)) {
     return iree_make_status(
         iree_status_code_from_win32_error(GetLastError()),
-        "failed to query logical processor information (%08X)", GetLastError());
+        "failed to query logical processor information (%08X)",
+        (unsigned)GetLastError());
   }
   SYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX* cache_relationships_end =
       (SYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX*)((uintptr_t)
@@ -222,14 +223,14 @@
     return iree_make_status(
         iree_status_code_from_win32_error(GetLastError()),
         "failed to query logical processor information size (%08X)",
-        GetLastError());
+        (unsigned)GetLastError());
   }
   if (all_relationships_size > 64 * 1024) {
     IREE_TRACE_ZONE_END(z0);
     return iree_make_status(IREE_STATUS_RESOURCE_EXHAUSTED,
                             "logical processor information size overflow (got "
                             "%u which is large for a stack alloc)",
-                            all_relationships_size);
+                            (unsigned)all_relationships_size);
   }
   SYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX* all_relationships =
       (SYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX*)iree_alloca(
@@ -241,7 +242,8 @@
     IREE_TRACE_ZONE_END(z0);
     return iree_make_status(
         iree_status_code_from_win32_error(GetLastError()),
-        "failed to query logical processor information (%08X)", GetLastError());
+        "failed to query logical processor information (%08X)",
+        (unsigned)GetLastError());
   }
   SYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX* all_relationships_end =
       (SYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX*)((uintptr_t)all_relationships +
@@ -278,8 +280,6 @@
 
   // Build an on-stack table for random access into all logical processors.
   // This isn't strictly required but makes it easier to walk the CPU table.
-  PROCESSOR_RELATIONSHIP** all_processors =
-      iree_alloca(sizeof(PROCESSOR_RELATIONSHIP*) * total_processor_count);
   iree_host_size_t global_processor_count = 0;
   for (SYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX* p = all_relationships;
        p < all_relationships_end;
@@ -361,14 +361,14 @@
     return iree_make_status(
         iree_status_code_from_win32_error(GetLastError()),
         "failed to query logical processor information size (%08X)",
-        GetLastError());
+        (unsigned)GetLastError());
   }
   if (all_relationships_size > 64 * 1024) {
     IREE_TRACE_ZONE_END(z0);
     return iree_make_status(IREE_STATUS_RESOURCE_EXHAUSTED,
                             "logical processor information size overflow (got "
                             "%u which is large for a stack alloc)",
-                            all_relationships_size);
+                            (unsigned)all_relationships_size);
   }
   SYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX* all_relationships =
       (SYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX*)iree_alloca(
@@ -380,7 +380,8 @@
     IREE_TRACE_ZONE_END(z0);
     return iree_make_status(
         iree_status_code_from_win32_error(GetLastError()),
-        "failed to query logical processor information (%08X)", GetLastError());
+        "failed to query logical processor information (%08X)",
+        (unsigned)GetLastError());
   }
   SYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX* all_relationships_end =
       (SYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX*)((uintptr_t)all_relationships +
diff --git a/runtime/src/iree/vm/bytecode/dispatch.c b/runtime/src/iree/vm/bytecode/dispatch.c
index 1c3d528..6ebbc38 100644
--- a/runtime/src/iree/vm/bytecode/dispatch.c
+++ b/runtime/src/iree/vm/bytecode/dispatch.c
@@ -923,7 +923,7 @@
       if (IREE_UNLIKELY(!buffer)) {
         return iree_make_status(IREE_STATUS_INVALID_ARGUMENT, "buffer is null");
       }
-      uint64_t* result = VM_DecResultRegI64("result");
+      uint64_t* result = (uint64_t*)VM_DecResultRegI64("result");
       *result = (uint64_t)iree_vm_buffer_length(buffer);
     });
 
@@ -1051,7 +1051,7 @@
                                 "source_buffer is null");
       }
       iree_host_size_t offset = VM_DecOperandRegI64HostSize("source_offset");
-      uint32_t* result = VM_DecResultRegI32("result");
+      int32_t* result = VM_DecResultRegI32("result");
       vm_buffer_load_i8u_inline(buffer, offset, result);
     });
     DISPATCH_OP(CORE, BufferLoadI8S, {
@@ -1064,7 +1064,7 @@
                                 "source_buffer is null");
       }
       iree_host_size_t offset = VM_DecOperandRegI64HostSize("source_offset");
-      uint32_t* result = VM_DecResultRegI32("result");
+      int32_t* result = VM_DecResultRegI32("result");
       vm_buffer_load_i8s_inline(buffer, offset, result);
     });
     DISPATCH_OP(CORE, BufferLoadI16U, {
@@ -1077,7 +1077,7 @@
                                 "source_buffer is null");
       }
       iree_host_size_t offset = VM_DecOperandRegI64HostSize("source_offset");
-      uint32_t* result = VM_DecResultRegI32("result");
+      int32_t* result = VM_DecResultRegI32("result");
       vm_buffer_load_i16u_inline(buffer, offset, result);
     });
     DISPATCH_OP(CORE, BufferLoadI16S, {
@@ -1090,7 +1090,7 @@
                                 "source_buffer is null");
       }
       iree_host_size_t offset = VM_DecOperandRegI64HostSize("source_offset");
-      uint32_t* result = VM_DecResultRegI32("result");
+      int32_t* result = VM_DecResultRegI32("result");
       vm_buffer_load_i16s_inline(buffer, offset, result);
     });
     DISPATCH_OP(CORE, BufferLoadI32, {
@@ -1103,7 +1103,7 @@
                                 "source_buffer is null");
       }
       iree_host_size_t offset = VM_DecOperandRegI64HostSize("source_offset");
-      uint32_t* result = VM_DecResultRegI32("result");
+      int32_t* result = VM_DecResultRegI32("result");
       vm_buffer_load_i32_inline(buffer, offset, result);
     });
     DISPATCH_OP(CORE, BufferLoadI64, {
@@ -1116,7 +1116,7 @@
                                 "source_buffer is null");
       }
       iree_host_size_t offset = VM_DecOperandRegI64HostSize("source_offset");
-      uint64_t* result = VM_DecResultRegI64("result");
+      int64_t* result = VM_DecResultRegI64("result");
       vm_buffer_load_i64_inline(buffer, offset, result);
     });
 
@@ -1189,7 +1189,7 @@
       iree_host_size_t source_offset =
           VM_DecOperandRegI64HostSize("source_offset");
       iree_host_size_t length = VM_DecOperandRegI64HostSize("length");
-      uint64_t* result = VM_DecResultRegI64("result");
+      int64_t* result = VM_DecResultRegI64("result");
       IREE_RETURN_IF_ERROR(
           iree_vm_buffer_hash(source_buffer, source_offset, length, result));
     });