Release v3.10.0

IREE Release v3.10.0

1. Compiler

1.1 Codegen & Dispatch Improvements

• Introduced Producer-Consumer Fusion (PCF) infrastructure for improved fusion of producer and consumer operations across dispatch boundaries:
  • Added LowerStructuralPCFPass for lowering structural PCF patterns. (#22915)
  • Added FusePCFWritesPass for fusing writes in PCF patterns. (#22914)
  • Added FuseConsumersPass to improve consumer fusion in codegen pipelines. (#22912)
  • Added ConvertForallToLoopsPass for converting scf.forall operations to standard loops. (#22911)

1.2 LinalgExt Enhancements

• Added map_gather op and verifier to complement existing map_scatter functionality. (#23041)
• Simplified linearize-delinearize pairs in map_scatter operations. (#22958)
• Avoided unnecessary masking in map_scatter for improved performance. (#22959)

1.3 Dispatch Creation & Encoding

• Create more multi-use dispatches for improved dispatch efficiency. (#22011)
• Added FoldExtractSliceOfBroadcast pattern for dispatch creation. (#22694)
• Sink bit-extend producers into split-reduction forall loop. (#21850)
• Added UnifyEncodingForGlobals pass for unified encoding handling. (#22767)

1.4 AMDGPU & LDS Improvements

• Unified DMA transfer linearization based on destination contiguity for better memory access patterns. (#22995)
• Preferred contiguous subviews in GPUConvertToCoalescedDMA for improved memory coalescence. (#23011)
• Added 3-stage pipelining with hipblaslt compute->write->read ordering (#22788)

1.5 Data Tiling (CPU)

• Adjusted scalable tile sizes and flags for data-tiling fusion on CPU backends. (#22878)
• Adjusted tile sizes for unpack operations on SVE. (#21728)
• Redefined intrinsics M/N interleaving for GPU data tiling. (#22812)

1.6 Bug Fixes & Improvements

• Fixed HAL dealloca verifier to avoid segfault. (#22988)
• Fixed VM integer to f64 cast folding. (#22997)
• Fixed Flatbuffers to not canonicalize empty strings to null strings. (#23039)
• Fixed RISC-V toolchain prefix usage and added RISCV_TOOLCHAIN_TRIPLE. (#22853)

2. Runtime

2.1 VM Improvements

• Added new linear-scan register allocator for improved register allocation. (#23027)
• Added DiscardRefs and AssignRef VM ops for better reference management. (#23021)
• Added --iree-vm-materialize-ref-discards pass for reference discard materialization. (#23026)
• Added iree_vm_list_get_ref_retain_or_move for flexible reference handling in VM lists. (#23043)
• Added VMFB disassembly support and cleaned up move handling. (#23018)
• Fixed potential stack unwinding leak in VM. (#23019)
• Supported float type conversion between f32 and f64 in VM lists. (#22893)
• Un-packed VM ABI shims and added proper alignment in C and C++ modules. (#22845)
• Cleaned up bytecode ISA decoding macros (NFC). (#23042)
• Cleaned up move bit handling for disassembly/verifier/dispatch. (#23020)

3. Infrastructure & CI

• Upgraded to GCC 11 to fix stablehlo template instantiation error. (#23008)
• Added ci-extra flags for scheduled builds. (#23033)
• Allowed ci-extra to be set on integrates. (#23034)

---

Change Log

Git History

What's Changed

• [Codegen] Test Cleanup 7/8: SPIRV tests by @qedawkins in #22750
• [Codegen] Test Cleanup 6/8: LLVMGPU tests by @qedawkins in #22749
• [Codegen] Test Cleanup 5/8: LLVMCPU tests by @qedawkins in #22748
• [Codegen] Test Cleanup 2/8: Common GPU tests by @qedawkins in #22745
• Use scf::tileAndFuseConsumer in GPUFuseAndHoistParallelLoops by @MaheshRavishankar in #22617
• Integrate llvm/llvm-project@c582688b by @yzhang93 in #22758
• [Codegen] Test Cleanup 3/8: Common tests by @qedawkins in #22746
• [DispatchCreation] Add FoldExtractSliceOfBroadcast Pattern by @bangtianliu in #22694
• Update gfx1250 LDS size by @Hardcode84 in #22760
• Bump version to 3.10 after 3.9 release. by @sa-faizal in #22759
• Integrate llvm/llvm-project@ebf5d9ef by @yzhang93 in #22761
• [Encoding] Implement compatibility check for packed_storage by @sommerlukas in #22757
• [Input] Register IREETensorExtDialect for Torch plugin by @IanWood1 in #22719
• Integrate LLVM @ 356479191ca0 by @lialan in #22772
• [Codegen][Tuner] Expose the python bindings for LinalgExt::inferScaledContractionDims and LinalgExt::isaScaledContractionOpInterface by @Muzammiluddin-Syed-ECE in #22763
• [tests][e2e] Add custom mxfp4 gemm tests to verify shape of interest. by @Muzammiluddin-Syed-ECE in #22775
• [Bazel] Migrate to bzlmod for LLVM compatibility by @maxbartel in #22771
• ScheduleExecution enhancements for timeline-aware scheduling and SCF. by @benvanik in #22483
• [Flow] Annotate scaled matmul dispatches by @Yu-Zhewen in #22773
• build_tools: fix: ensure that iree-flatcc-cli and iree-c-embed-data are build for the target during cross-compilation by @Manewing in #22755
• [Codegen][GPU] Add fusion barrier after result promotion by @qedawkins in #21709
• [CI] Update iree-org/iree-test-suites@132f91e4 by @efric in #22784
• [Codegen][GPU] Enable 3-stage pipelining with hipblaslt compute->write->read ordering by @jerryyin in #22788
• [LLVMGPU][Codegen] Reland "Emit packed chain FMA from select multi_reductions and contracts" by @efric in #22789
• Bump llvm-project to @a7c1f467339abd1942c89f2ef8b79083e89e7dad by @Max191 in #22787
• [Codegen] Support dynamic offsets in collapse_shape fusion to interface stores by @qedawkins in #22800
• [HAL] fix IREE_HAL_MAX_QUEUES to be number of bits in queue affinity type by @schuermans-roofline in #22702
• [Codegen][LLVMGPU] Replace TransposeSharedMem pipeline by @qedawkins in #21661
• Add iree_status_t stack trace support on Linux. by @benvanik in #22796
• [LLVMGPU] Fix lowering strategy for direct convolution by @yzhang93 in #22802
• [GPU] Add M dimension constraints for pingpong ukernel by @IanWood1 in #22801
• [Encoding] Remove unneeded command line option by @Muzammiluddin-Syed-ECE in #22816
• Add SCF support and fence coverage to ElideTimepointsPass. by @benvanik in #22611
• [GPU] Fix alignment check for scaled matmul by @Yu-Zhewen in #22737
• [Codegen][GPU] Replace prefetch_shared_memory with prefetch_num_stages in IREEGPUAttrs by @jerryyin in #22818
• [GlobalOpt] Fix rank-reduced permutation in SinkTransposeThroughExtractSlice by @ziliangzl in #22754
• [LLVMGPU] Update seeds for scaled gemm by @Muzammiluddin-Syed-ECE in #22798
• [Dispatch Creation] Don't fuse if there are no common parallel loops by @IanWood1 in #22819
• Reland "[LLVMGPU] Unroll elementwise operations #21665" by @lialan in #22828
• [Codegen][GPU] Preserve loop domain when collapsing dims in Conv to Matmul conversion by @yzhang93 in #22821
• Add jtuyls and Yu-Zhewen to CODEOWNERS for ROCM plugin and Encoding by @jtuyls in #22810
• [Encoding] Use struct directive for TestingAttr assembly format by @jtuyls in #22826
• Fix tests: noubsan was not being honored, and MXFP4 matmul tests are static-shape-only by @bjacob in #22836
• ElideAsyncCopiesPass refactoring for SCF/transfer support. by @benvanik in #22739
• [Runtime][HIP] Correct O(log n) bound search logic and eliminate O(n) loop(#22733) by @swote-git in #22734
• [e2e][ukernel] Remove dead/duplicate tests by @Yu-Zhewen in #22834
• [Codegen][GPU] Allow channel first layouts to lower through direct convolution path by @yzhang93 in #22840
• [Codegen] Add vector.to/from_elements to bf16 -> i16 conversion by @qedawkins in #22846
• Fix deadlock in ROCMDialect::getMlirUKernels by @bjacob in #22843
• [SPIRV][Codegen] Use single subgroup when reduction consumer has non-distributable broadcast by @efric in #22832
• [Encoding] Improve specialized encoding usage in lit test by @jtuyls in #22851
• [Stream] Encode packed_storage device and host tensors by @sommerlukas in #22722
• [Encoding] Add verifier for iree_encoding.layouts by @jtuyls in #22850
• [Codegen] Add WorkgroupCountHintOp to defer populating the workgroup count by @qedawkins in #22533
• Unify RISC-V toolchain environment variables and remove default path by @HanKuanChen in #22710
• [tests][e2e] Add more llama related shapes by @Muzammiluddin-Syed-ECE in #22831
• [Util] Implement InferIntDivisibilityOpInterface for affine ops by @Max191 in #22723
• Add CDNA3 test filtering to CI by @lialan in #22848
• [Torch Models] Fix SDXL golden dispatch counts by @IanWood1 in #22855
• [Dispatch Creation] Create more multi-use dispatches by @IanWood1 in #22011
• [LDS] Add AMDGPULowerCoalescedDMAToGatherLDS pass for direct global to LDS loads by @lialan in #22356
• Revert "[Util] Implement InferIntDivisibilityOpInterface for affine o… by @IanWood1 in #22860
• Integrates/llvm 20251208 by @bangtianliu in #22856
• [Codegen] Use workgroup_count_hint for most code paths by @qedawkins in #22549
• [Stream] Add UnifyEncodingForGlobals pass. 1/n by @hanhanW in #22767
• [Encoding] Add verifier for gpu/cpu/vmvx_encoding_resolver by @jtuyls in #22838
• [libcall] Update musl Makefile to include missing files. by @lialan in #22859
• Integrates/llvm 20251209 by @bangtianliu in #22864
• Bump the github-actions group across 1 directory with 3 updates by @dependabot[bot] in #22858
• [Codegen][GPU]Skip prologue pipeline barriers only for non-nested pipelined loops by @jerryyin in #22868
• [Pkgci] Update golden dispatch counts by @IanWood1 in #22869
• [DT][GPU] Redefine intrinsics M/N interleaving by @Yu-Zhewen in #22812
• [bindings] Add iree_tensor_ext to python bindings by @rkayaith in #22872
• [DT][SVE] adjust tile sizes for mmt4d & disable transposition of narrow-N matmuls by @egebeysel in #21701
• [Encoding] Add resolver swizzle verification by @jtuyls in #22867
• [Codegen] Do not swap extract_slice and collapse_shape for a special case by @yzhang93 in #22870
• Integrates/llvm 20251210 by @bangtianliu in #22876
• [LDS] Improve multiple transfers per lane by @lialan in #22879
• Integrates/llvm 20251211 by @bangtianliu in #22885
• [Stream] Handle AsyncCloneOp in UnifyEncodingForGlobals tracing. by @hanhanW in #22895
• Integrates/llvm 20251212 by @bangtianliu in #22897
• Update shark/SHARK to amd-shark/AMD-SHARK in documentation and URLs in IREE by @bangtianliu in #22883
• [Codegen] Add PCF dialect by @qedawkins in #22804
• Bump torch-mlir by @rkayaith in #22894
• [ROCM][DT] Add architecture matching to ukernel_info attribute by @Yu-Zhewen in #22899
• [Codegen] Add PCF bufferization interfaces by @qedawkins in #22805
• Bump to llvm/torch-mlir@3cebce2 by @zjgarvey in #22902
• [Stream] Update dispatch sites encodings in UnifyEncodingForGlobals pass 2/n by @hanhanW in #22886
• [LDS] Remove GPULowerToGlobalLoadsPass and GlobalLoadDMAOp by @lialan in #22862
• Integrate LLVM at b3ec8be by @bjacob in #22908
• [Dispatch Creation] Ignore unfusable consumer by @IanWood1 in #22874
• Add Codegen/Dialect CODEOWNERS entries for myself by @qedawkins in #22918
• [Stream] Handle tied operand result encodings in UnifyEncodingForGlobals 3/n by @hanhanW in #22891
• [Codegen] Drop certain encoding attributes after type propagation by @sommerlukas in #22740
• [DT][SVE] adjust tile sizes for unpack by @egebeysel in #21728
• [Stream] Selecting unified encodings from encoding resolvers. by @hanhanW in #22898
• [Encoding] Add EncodingProperties retrieval to SerializableAttr by @jtuyls in #22905
• [GlobalOpt] Fuse transpose into matmul-looking linalg.generic by @IanWood1 in #22901
• Bump the github-actions group with 3 updates by @dependabot[bot] in #22913
• Bump stablehlo to 6fabd27 by @kuhar in #22927
• [Encoding] Test i1 mask attention with packed_storage by @sommerlukas in #22756
• [bindings] Link CAPIAMDGPU into libIREECompiler by @tgymnich in #22929
• Integrate LLVM at 0729a74 by @bjacob in #22925
• [LinalgExt] Added support for LinalgExt::IndexOp(s) inside of attention regions by @keshavvinayak01 in #22768
• [CI][torch models] Increase acceptable tolerance in punet by @efric in #22935
• [DispatchCreation] Sink bit-extend producers into split-reduction forall loop by @rkayaith in #21850
• Integrate LLVM at 5f15fee by @bjacob in #22933
• [Stream] Prioritize NamedParameterAttr in UnifyEncodingForGlobals source tracing. 4/n by @hanhanW in #22922
• Bump torch-mlir to ac7b5f5 by @zjgarvey in #22934
• Add myself to CODEOWNERS for tests/e2e by @bjacob in #22944
• e2e matmul tests: fix some recent CDNA3-only tests by @bjacob in #22945
• [Input][Torch] Direct lowering of torch convolution_backward to linalg by @a-sidorova in #22904
• [CI][Torch Models] Tighten golden times by @efric in #22936
• [LinalgExt] Added DecomposeAggregatedOp Pass by @hhkit in #22779
• [LinalgExt] Added Decomposition of ExpReductionOp to LinAlg by @hhkit in #22561
• [Codegen][GPU] Improve lowering strategy for direct convolution by @yzhang93 in #22937
• Integrate LLVM at 1c023cb by @bjacob in #22943
• [GPU] Implement dummy getUnifiedEncoding method for GPU encoding resolver. by @hanhanW in #22923
• Integrate LLVM at 60e7c47 by @bjacob in #22952
• [CI] Update golden dispatch counts by @efric in #22950
• [CI][Torch Models] Update golden binary size by @efric in #22951
• Disable simplify regions after scf-to-cf by @nirvedhmeshram in #22946
• [GPU][Codegen] Expand iteration space based on new expand_dims attribute by @efric in #22342
• Revert "[GPU][Codegen] Expand iteration space based on new expand_dims attribute (#22342) by @efric in #22982
• Integrate LLVM at c0f4a8a by @jtuyls in #22979
• Integrate LLVM at 75a0347 by @jtuyls in #22985
• [Codegen][ROCDL] Improve dynamic dimension bounds handling in ROCDLConfigureBufferInstructions by @Yu-Zhewen in #22892
• [Codegen] add FoldExtractSliceOfBroadcast pattern to TileAndDistributeToWorkgroups by @bangtianliu in #22953
• [LDS] CoalescedGatherDMAOp support even more transfer sizes. by @lialan in #22890
• Integrate llvm-project/llvm@66c65f0132d05 by @kuhar in #22990
• [Codegen] Fix ConvertAccGemmToGemm to handle read-write arguments correctly. by @MaheshRavishankar in #22975
• [CI] Enable split dwarf and debug info in CI builds by @kuhar in #22993
• Bump dawidd6/action-download-artifact from 11 to 12 in the github-actions group by @dependabot[bot] in #22989
• [CI] Bump MacOS x64 version used in CI by @kuhar in #22991
• [CI] Reclaim disk space on macos arm runners by @kuhar in #22998
• [CI] Use build_all.sh in clang checks by @kuhar in #22999
• Integrate llvm/llvm-project@292c9e3d198249 by @kuhar in #23000
• [CI] Allow macos checks to be triggered with ci-extra by @kuhar in #23001
• [Codegen][ROCDL] Fix int64 overflow in getSpannedBytes by @Yu-Zhewen in #23002
• [CI] Drop macos job concurrency block by @kuhar in #23004
• Integrate llvm-project/llvm@0db04963d34be by @kuhar in #23006
• [CI] Do not run schedule-only jobs on integrates by @kuhar in #23009
• [TensorExt] Add Interface methods for lowering to loops and estimating bounds. by @MaheshRavishankar in #22268
• Integrate llvm/llvm-project@7976ac990000 by @kuhar in #23010
• Bump Torch-MLIR to 9a1f70ddd46c1f217522d83c0dfc9d36d65e9eda by @vivekkhandelwal1 in #22980
• [NFC] Adding iree_vm_bytecode_module_create flag enum arg. by @benvanik in #23017
• Support float type conversion between f32 and f64 in VM lists. by @willghatch in #22893
• Adding VMFB disassembly and cleaning up move handling. by @benvanik in #23018
• Fixing potential stack unwinding leak in VM. by @benvanik in #23019
• [AMDGPU][LDS] Unify DMA transfer linearization based on destination contiguity by @lialan in #22995
• [CI] Upgrade to GCC 11 to fix stablehlo template instantiation error by @Yu-Zhewen in #23008
• Un-packing VM ABI shims and adding proper alignment in C and C++ modules. by @benvanik in #22845
• Cleaning up move bit handling for disassembly/verifier/dispatch. by @benvanik in #23020
• Adding DiscardRefs and AssignRef VM ops. by @benvanik in #23021
• Adding --iree-vm-materialize-ref-discards pass. by @benvanik in #23026
• Integrate llvm/llvm-project@10a245bd02 by @Abhishek-Varma in #23022
• [Codegen] Add ConvertForallToLoopsPass by @qedawkins in #22911
• [AMDGPU][LDS] Prefer contiguous subviews in GPUConvertToCoalescedDMA by @lialan in #23011
• Adding new linear-scan register allocator. by @benvanik in #23027
• [CI] Add ci-extra flags for scheduled builds. by @amd-eochoalo in #23033
• [Codegen] Add FuseConsumersPass by @qedawkins in #22912
• [NFC] Cleanup of bytecode ISA decoding macros. by @benvanik in #23042
• [Codegen] Add FusePCFWritesPass by @qedawkins in #22914
• Adding iree_vm_list_get_ref_retain_or_move. by @benvanik in #23043

Contributors

@Abhishek-Varma, @bangtianliu, @benvanik, @bjacob, @dependabot, @efric, @egebeysel, @fwalbroel, @gymlet, @hanhanW, @HanKuanChen, @ihowell, @IanWood1, @IvanButygin, @jerryyin, @jtuyls, @keshavvinayak01, @kuhar, @lialan, @Max191, @maxbartel, @MaheshRavishankar, @Muzammiluddin-Syed-ECE, @nirvedhmeshram, @nohyeonkwon, @qedawkins, @rkayaith, @sa-faizal, @schuermans-roofline, @sommerlukas, @asidorova, @vivekkhandelwal1, @willghatch, @yzhang93, @Yu-Zhewen, @ziliangzl, @zjgarvey, @amd-eochoalo

New Contributors

• @swote-git made their first contribution in #22734
• @a-sidorova made their first contribution in #22904

Full Changelog: v3.9.0...v3.10.0