Skip to content

Release v3.12.0

Latest
Latest

Choose a tag to compare

View all tags
@devtbi devtbi released this 05 Oct 12:33
v3.12.0
2b05c5d
This commit was created on GitHub.com and signed with GitHub’s verified signature.
GPG key ID: B5690EEEBB952194
Verified
Learn about vigilant mode.

IREE v3.12.0 Release Notes

Release Candidate: iree-3.12.0rc20260917
Commits: ~720 PRs since v3.11.0
VMFB Bytecode Version: 17.0 (unchanged from v3.11.0), but The HAL module is now version 0.7, so previous VMFBs are incompatible.

Breaking changes

  • Recompile your VMFBs. The HAL module is now version 0.7. Modules compiled with 3.11 or earlier will not load with the 3.12 runtime. #24507

  • The HAL semaphore system was replaced. The HAL C API was changed. Each HAL semaphore is a wrapper around a shared iree_async_semaphore_t, driven by the async proactor.

    • iree_hal_semaphore_signal() and iree_hal_semaphore_list_signal() take a new const iree_async_frontier_t* argument frontier; pass NULL if you don't need it.
    • The semaphore wait functions take iree_async_wait_flags_t instead of iree_hal_wait_flags_t, and iree_hal_wait_mode_t is now iree_async_wait_mode_t.
    • iree_wait_primitive_t is now iree_async_primitive_t, so the external timepoint type WAIT_PRIMITIVE is now ASYNC_PRIMITIVE. iree_wait_source_ctl_fn_t is now iree_wait_source_resolve_fn_t.
    • iree/base/wait_handle.h, iree/base/event_pool.h and iree/hal/utils/semaphore_base.h were removed.
    • Out-of-tree HAL drivers must embed iree_async_semaphore_t.

    #23695, #23804

  • Other HAL C API changes:

    • All iree_hal_driver_create_device* functions take iree_hal_device_create_params_t, which has a required proactor_pool.
    • iree/base/loop.h and iree_loop_* moved to iree/vm/loop.h and iree_vm_loop_*.
    • create_executable_cache no longer takes a loop argument.
    • iree_hal_device_queue_alloca takes an iree_hal_pool_t* (NULL for the default pool) instead of a pool bitmask.
    • iree_hal_executable_export_* replaced with iree_hal_executable_function_*.

    #23695, #24236, #24507

  • The Vulkan driver was rewritten and now requires Vulkan 1.3 with timelineSemaphore, scalarBlockLayout and synchronization2. Older Vulkan stacks are no longer enumerated. Runtime flag defaults changed:

    • --vulkan_sparse_binding and --vulkan_sparse_residency now default to false (due to bug reports)

    #24459, #24738

  • Device profiling flags were replaced.

    • --device_profiling_file becomes --device_profiling_output=<file>.ireeprof.
    • --device_profiling_mode now takes a comma-separated list of data families, e.g. queue-events,dispatch-events, instead of queue|dispatch|executable.
    • RenderDoc/Metal captures move to --device_capture_tool/_file/_label.
    • In Python, HalDevice.begin_profiling() no longer takes a file path; use begin_external_capture() instead.

    #24240

  • Tool output comparison changed. --equality_mode accepts only approximate, the new numpy.isclose-style default, or exact; absolute and relative were removed. The new --expected_rtol (default 1e-5) sets the relative tolerance, and the --expected_f*_threshold flags are now absolute tolerances. #24394

  • Codegen pipeline names in translation_info changed. They are now backend attributes, e.g. #iree_gpu.pipeline<TileAndFuse> and #iree_cpu.pipeline<Default>, and the old keywords no longer parse. Update hand-written tuning specs and compilation_info attributes. #23816, #24398

  • Removed and renamed compiler flags:

    • The --iree-hip-* aliases were removed; use --iree-rocm-*. #24381
    • --iree-rocm-index-bits was removed. #23948
    • --iree-global-opt-enable-demote-contraction-inputs-to-bf16 was replaced by --iree-global-opt-demote-contraction-inputs-type=<f16|bf16> and --iree-global-opt-demote-contraction-inputs-operations=<all|conv|matmul>. #23850
    • --iree-global-opt-experimental-disable-conv-generalization was removed. #23377
    • --iree-vulkan-experimental-indirect-bindings was replaced by --iree-vulkan-dispatch-abi=descriptors|bda|all (default descriptors). #24459

  • The iree-dialects project was removed, together with the legacy transform.iree.match_callback-style transform ops and the iree.compiler.dialects.transform.iree_structured Python module. #24612

  • Python implicit uploads (NumPy arguments, asdevicearray) now use device-local memory instead of forcing host-visible memory. #24106

  • Packaging and build:

    • Python 3.13t (free-threaded) wheels are no longer built for Linux aarch64, because the upstream manylinux image dropped 3.13t.
    • Bazel builds now require Bazel 8.6.0. #24153
    • The old experimental/web samples were removed. #23695

Notable changes

Compiler

  • New targets:
    • WebGPU: a WGSL target that goes through SPIR-V and Tint (opt-in via IREE_TARGET_BACKEND_WEBGPU_SPIRV). #24463
    • NVIDIA: Ada (sm_89) and Blackwell (sm_120/sm_121). #24422, #24616
    • AMD: CDNA5/MI455X and RDNA3 chip details. #24923, #24419
    • Vulkan: Broadcom VideoCore VII, Arm Mali 5th-gen and llvmpipe. #24541, #24559, #24595
  • ROCm GPU codegen:
    • Virtual dense MFMA for skinny GEMMs (M≤8), up to ~24% faster on gfx950. #23677, #24103
    • DMA is now on by default for F16/BF16 GEMMs on gfx950. #24411
    • Retuned LargeGemm heuristics (~7% geomean on MI355X bf16). #23836
  • CUDA GPU codegen:
    • BF16 mma.sync lowering is enabled. #24423
    • FP8 mma.sync intrinsics for sm_89/sm_120, for intrinsic selection only. #24659
  • GPU ops: arg_compare and scan are now supported in VectorDistribute/TileAndFuse. #24003, #24187
  • Tuner: SMT-based pipeline constraint generation and verification for VectorDistribute and TileAndFuse, with Python bindings. #23795, #24539
  • CPU codegen:
    • Opt-in inner_tiled lowering for data-tiled encodings and a C-bitcode ukernel framework (--iree-llvmcpu-enable-inner-tiled, --iree-llvmcpu-enable-llvm-ukernels). #24218, #24569
    • RISC-V: scalable RVV data tiling, plus RVV bf16/int8 and SpaceMiT IME kernels. #24600, #24695, #24706
    • Data-tiled convolutions can now be lowered to a conv_nchwc ukernel. #24714, #24730
    • Ensured that SME-only targets such as Apple Silicon work alongside SME + SVE targets. #24661
    • Fixed a ~2.4x f32 matmul regression on non-AVX512 x86. #24701
  • Dispatch creation: more fusion, covering producers with multiple users (e.g. layernorm in one dispatch), multi-result producers, and scalar reductions with their consumers. #24050, #24169, #24166
  • Attention:
    • Attention now goes through online_attention throughout codegen and dispatch formation. #24110, #24068
    • New topk_v2, quantize_affine and dequantize_affine ops. #24054, #24901
  • Input conversion:
    • Torch: hop_flex_attention (including GQA) and aten.argmax/argmin lowerings. #23292, #24313, #24291
    • StableHLO: stablehlo.composite support, which fixes JAX models. #24777
    • TOSA: dynamic dims and large tensors are accepted. #24667

Runtime

  • Vulkan: the HAL driver was fully rewritten. It has native queue paths, push descriptors and a new descriptor-free buffer-device-address dispatch ABI with cached command-buffer replay. #24459, #24505
  • AMDGPU: a new native HSA HAL driver, experimental and off by default (IREE_HAL_DRIVER_AMDGPU). It replays AQL packets, includes profiling, counters and traces, and offers optional PM4 command buffers. #24298, #24452
  • CPU (local-task/local-sync): the execution path was rewritten, and the task system no longer has a poller thread or the 63-waiter limit. #24248, #23695
  • HAL: all drivers now share one async semaphore implementation integrated with the proactor, and there is a new allocation-pool substrate (TLSF, slab, fixed-block). #23695, #24236
  • WebAssembly/WebGPU:
    • A new freestanding wasm32 platform with a JavaScript proactor. #24462
    • A new webgpu HAL driver. #24463
  • Bare metal: freestanding runtime builds for IREE_PLATFORM_GENERIC, with a riscv64 QEMU sample. #24820, #24844
  • Metal and CUDA fixes:
    • Metal: several correctness fixes and a system-default device fallback. #24639, #24839
    • CUDA: export names containing $ now work. #24889
  • Python:
    • Tokenizer bindings (HuggingFace JSON and tiktoken). #24009
    • Zero-copy HalAllocator.import_host_allocation(). #23792
    • BF16/FP8 to_host(). #24263
    • create_device(params=...) now honors every param, not just the last one. #24812

Developer tools

  • iree-profile and iree-profile-render inspect HAL-native .ireeprof device profiles and export them to Perfetto. #24240
  • iree-run-replay, iree-benchmark-replay and iree-dump-replay replay HAL streams recorded with --device_replay_output. #24288
  • iree-as-module assembles .vmasm into VMFBs. iree-dump-module now prints round-trippable VMASM, so scripts that parse its output may need updating. #24508
  • iree-benchmark-module --enable_output_processing writes or checks the outputs of the last iteration. #24789
  • Build:
    • New CMake options IREE_BUILD_BENCHMARKS, IREE_ENABLE_POSIX, and the experimental IREE_EXPERIMENTAL_COMPILER_DYNAMIC_PLUGINS. #24167, #24024, #24891
    • Bazel: configurable compiler plugins and runtime Python bindings. #23212, #24553

New contributors

Full changelog

Commit history: v3.11.0...v3.12.0

Release notes were AI-generated and edited.

Contributors

jph00, dimp-pl, and 33 other contributors
Assets 27
Loading