Hal resolve memory properties - #2
Open
ziereis wants to merge 250 commits into
Open
Conversation
…ree-org#21047) Refactors the map_scatter folding transformations with tensor.collapse_shape and tensor.expand_shape into pattern rewrites, and adds the patterns to the BlockDynamicDimensions pass. This allows the CombineLayoutTransformation pass to be run before BlockDynamicDimensions, without creating additional reshapes in the IR. Signed-off-by: Max Dawkins <max.dawkins@gmail.com>
…org#21069) This option is useful since we don't get good performance when we fission multi-trip loops. Additionally the prefetching pass that makes use of fission is not set up to support multi-trip nested loops. Note that even if one candidate is found to be multi-trip we don't do the whole pass as we wont be doing prefetching in that case and hence no point in fissioning at all. --------- Signed-off-by: Nirvedh Meshram <nirvedh@gmail.com>
iree-org#21073) This enables elementwise op fusion and some of cases might benefit from this. Fix for the issue: iree-org#20875
This is necessary for IREE to be built against the tracy with rocprof fork. The rocprof cmake module defines this macro; causing a redefinition error in IREE. Signed-off-by: Eric Eaton <erieaton@amd.com>
API breakage due to: llvm/llvm-project#143217
…ds (iree-org#21085) Add an explicit dependency to avoid `EncodingTypes.h` inclusion failures in shared library builds. Signed-off-by: Artem Gindinson <gindinson@roofline.ai>
Enables the map_scatter op e2e tests for rocm, the `iree-llvmgpu-test-combine-layout-transformation` flag (from iree-org#21076) is not needed, since the map_scatter op is present in the input IR. Signed-off-by: Max Dawkins <max.dawkins@gmail.com>
I still need to rip out events and change them to be something command buffer local.
The signal and queue headers are derived from the HSA spec, ROCR implementation, and LLVM device library/codegen.
…rg#21043) The system is shared across all physical devices within a logical device and owns resources shared across them (it's how we avoid a cycle between logical<->physical device impls).
… options. (iree-org#21086) A uint64 version was added awhile back but the rest were missing. C's strtol and co all take a base and this just passes that through. The non-_base methods assume 0 ("interpret from string") as before. Also cleaned up hip driver int parsing code a bit (noticed while touching the string view functions).
This allows code launching threads to specify that a thread should be assigned to a processor associated with a specific group (Windows GROUP_AFFINITY, Linux NUMA node ID, etc) instead of a specific processor within the group. This is useful for threads created in various parts of the codebase that don't globally coordinate (service workers/etc), while things like the task system that explicitly layout large collections of threads will continue to explicitly assign them.
The driver will implement two buffer types to start (or more?): `iree_hal_amdgpu_external_buffer_t` and `iree_hal_amdgpu_transient_buffer_t`. The external buffer will be used for imported buffers and synchronous allocations (probably), while the transient buffer is used for all asynchronous allocations. The transient buffer maintains a reference to a device-accessible `iree_hal_amdgpu_device_allocation_handle_t` that stores the allocated pointer. The device-side allocator will update the handle as it allocates/deallocates and anyone resolving buffer pointers (e.g. `iree_hal_amdgpu_device_buffer_ref_resolve`) fetches the current pointer from the handle. The pointer is only valid between alloca/dealloca ops (which will set/reset the pointer). `iree_hal_amdgpu_device_buffer_ref_t` is a variant that allows binding tables and command buffer ops to reference typed buffers. Most code can dereference device pointers (synchronously allocated buffers, imported buffers, etc) and the transient buffer handles. Command buffers can reference binding table slots with a relative offset/length. `iree_hal_amdgpu_device_workgroup_count_buffer_ref_t` is a special case for indirect dispatches (where the length is always known) as it packs better into space-constrained structs. A simple `iree_hal_amdgpu_buffer_pool_t` handles both the host `iree_hal_buffer_t` instances and device allocation handles. The device allocation handles can be stored in device memory to avoid additional bus traffic on each dereference. In the steady state of invocations with alloca/dealloca the pool should remove all host/device allocations.
This doesn't do anything but is enough to satisfy build requirements of future PRs. This was reviewed as iree-org#21045 but incorrectly merged. This PR just merges again on main now that the base commits have landed.
ziereis
force-pushed
the
hal_resolve_memory_properties
branch
from
June 13, 2025 12:40
d075ef3 to
6ade2e6
Compare
LLVMGPUPadAndVectorDistribute was added to handle unaligned batch matmuls before TileAndFuse existed. TileAndFuse now handles unaligned matmuls much better. We don't want to do padding in VectorDistribute this way.
iree-org#20708) This patch implements support for distributing vector.constant_mask op. This is same as vector.create_mask and uses the same implementation.
…ree-org#21062) In order to implement support for block-scaled MMA and to enable future operations over other "inner tiled" operations (for example, hand-rolled CPU micro-kernels), generalize iree_gpu.multi_mma to iree_codegen.inner_tiled. This PR is the next stage of the changes begun in iree-org#21000 The main differences between iree_gpu.multi_mma and iree_codegen.inner_tiled are that: 1. inner_tiled takes an arbitrary number of inputs and accumulator outputs/inits, with the INnerTileDescAttrInterface describing the expected number of arguments. The inits are named "outputs" following the `linalg.generic` convention. 2. The permutations - when present - are now stored as an array-of-arrays, and not as one array for each argument (since there are no longer fixed "lhs", "rhs", and "acc" arguments). 3. The operation takes an InnerTileDescAttrInterface implement instead of a MmaInterfaceAttr, though checking for the old interface is still sometimes used to scope patterns down to just multi-mma-like operations. 4. We use linalg.iterator_type<> instead of iree_gpu.iterator_type<>, since trying to define iree_codegen.iterator_type<> was causing redefinition errors In terms of transformations, this PR generalizes (and, as a consequence, renames) those transformations that had clear generalizations to arbitrary intrinsics. This means that the configuration logic / lowering to multi-MMA is unaffeced, but a lot of the internal transformations (such as vectorization or tiling) have been reworked to support N inputs. These changes have not required more than mechanical (changing name + assembly format) changes to any tests and never showed end-to-end correctness issues, so it's likely that this PR causes no change in the generated code. To provide a motivating example for the generalization ```mlir iree_codegen.inner_tiled ins(%lhs, %lhs_scale, %rhs, %rhs_scale) acc(%acc) { indexing_maps = [ affine_map<(i, j, k) -> (i, k)>, affine_map<(i, j, k) -> (i, k)>, affine_map<(i, j, k) -> (j, k)>, affine_map<(i, j, k) -> (j, k)>, affine_map<(i, j, k) -> (i, j)> ], iterator_types = [parallel, parallel, reduction], kind = #iree_gpu.scaling_mma_kind<MFMA_16x16x128, lhsType = f4E2M1FN, rhsType = f8E3M3FN, acc_type = f32> } : tensor<?x?x16x128xf4E2M1FN>, tensor<?x?x16x4xf32>, // Note k dimension is a bunch smaller for scales tensor<?x?x128x16xf8E4M3FN>, tensor<?x?x4x16xf32> into tensor<?x?x16x16xf32> ``` The exact representation of scaling MMAs will be designed out later, but the above example is a demonstration of the overall motivation for this change Fixes #iree-org#20814 --------- Co-authored-by: Max191 <44243577+Max191@users.noreply.github.com> Co-authored-by: Quinn Dawkins <quinn.dawkins@gmail.com>
Signed-off-by: Ian Wood <ianwood2024@u.northwestern.edu>
Adds the `BufferizeDispatchTensorLoadStorePass` and the `GPUCombineLayoutTransformationPass` to the LLVMGPU configuration pipeline behind a new test flag, `iree-llvmgpu-test-combine-layout-transformation`. Also adds a pipeline test for the new configuration path. This will allow for enabling e2e tests for the map_scatter op, and some more complex data tiling fusions, while work for supporting map_scatter on the LLVMGPUVectorDistribute pipeline is in progress. Once the `LLVMGPUVectorDistribute` supports map_scatter, and the `LLVMGPUWarpReduction` pipeline is deprecated, then the flag can be flipped on by default. Signed-off-by: Max Dawkins <max.dawkins@gmail.com>
Adds a new data tiling encoding e2e test for a dispatch containing `unset_encoding -> elementwise -> set_encoding`. This uses the `iree-llvmgpu-test-combine-layout-transformation` test flag, and tests the ability for the LLVMGPU backend to handle the fusion case that arises when there are chained data-tiled matmuls. Signed-off-by: Max Dawkins <max.dawkins@gmail.com>
…riscv64 (iree-org#20263) ## PR Summary This PR enables the `linalg.mmt4d` operation and adds mmt4d microkernels for the riscv64 target. ## Status This is a work in progress. The current microkernel implementation is a placeholder, and we would like to discuss the tiling and microkernel design strategy with IREE authors before proceeding further ## RISC-V Vector Considerations According to the RISC-V specification, the vector register length (VLEN) can range from VLEN = 32 to VLEN = 65,536. However, in real hardware, we've mostly observed VLEN ≤ 4096. Additionally, RISC-V Vector supports a vector register group multiplier (LMUL), which allows multiple vector registers to be grouped together at runtime. This enables a single vector instruction to operate on multiple registers. ## Tile Size Selection Strategy Currently, we are considering selecting a tile size based on VLEN. The maximum tile size {7,32,1} specified [here](https://github.com/iree-org/iree/pull/20263/files#:~:text=for%20VLEN%3D256-,TileMxNxK%7B7%2C%2032%2C%201%7D,-%2C%20//%20Aim) is optimized for VLEN = 256, achieving 100% vector register utilization in mmt4d ukernel when using the vfmacc instruction with LMUL = 4: - Result Accumulation: 7 × 4 = 28 registers - RHS Tile: 1 × 4 = 4 registers For this particular tile size, the mmt4d ukernel code would look something like this: ``` if (params->flags & IREE_UK_FLAG_MMT4D_ACCUMULATE) { acc0 = __riscv_vle32_v_f32m4(out_ptr, vlmax); acc1 = __riscv_vle32_v_f32m4(out_ptr + 32, vlmax); acc2 = __riscv_vle32_v_f32m4(out_ptr + 64, vlmax); acc3 = __riscv_vle32_v_f32m4(out_ptr + 96, vlmax); acc4 = __riscv_vle32_v_f32m4(out_ptr + 128, vlmax); acc5 = __riscv_vle32_v_f32m4(out_ptr + 160, vlmax); acc6 = __riscv_vle32_v_f32m4(out_ptr + 192, vlmax); } else { acc0 = __riscv_vfmv_v_f_f32m4(0.0, vlmax); acc1 = __riscv_vfmv_v_f_f32m4(0.0, vlmax); acc2 = __riscv_vfmv_v_f_f32m4(0.0, vlmax); acc3 = __riscv_vfmv_v_f_f32m4(0.0, vlmax); acc4 = __riscv_vfmv_v_f_f32m4(0.0, vlmax); acc5 = __riscv_vfmv_v_f_f32m4(0.0, vlmax); acc6 = __riscv_vfmv_v_f_f32m4(0.0, vlmax); } for (int k = 0; k < params->K; ++k) { vfloat32m4_t rhs = __riscv_vle32_v_f32m4(rhs_ptr, vlmax); rhs_ptr += 32; float lhs0 = *lhs_ptr++; float lhs1 = *lhs_ptr++; float lhs2 = *lhs_ptr++; float lhs3 = *lhs_ptr++; float lhs4 = *lhs_ptr++; float lhs5 = *lhs_ptr++; float lhs6 = *lhs_ptr++; acc0 = __riscv_vfmacc_vf_f32m4(acc0, lhs0, rhs, vlmax); acc1 = __riscv_vfmacc_vf_f32m4(acc1, lhs1, rhs, vlmax); acc2 = __riscv_vfmacc_vf_f32m4(acc2, lhs2, rhs, vlmax); acc3 = __riscv_vfmacc_vf_f32m4(acc3, lhs3, rhs, vlmax); acc4 = __riscv_vfmacc_vf_f32m4(acc4, lhs4, rhs, vlmax); acc5 = __riscv_vfmacc_vf_f32m4(acc5, lhs5, rhs, vlmax); acc6 = __riscv_vfmacc_vf_f32m4(acc6, lhs6, rhs, vlmax); } __riscv_vse32_v_f32m4(out_ptr, acc0, vlmax); __riscv_vse32_v_f32m4(out_ptr + 32, acc1, vlmax); __riscv_vse32_v_f32m4(out_ptr + 64, acc2, vlmax); __riscv_vse32_v_f32m4(out_ptr + 96, acc3, vlmax); __riscv_vse32_v_f32m4(out_ptr + 128, acc4, vlmax); __riscv_vse32_v_f32m4(out_ptr + 160, acc5, vlmax); __riscv_vse32_v_f32m4(out_ptr + 192, acc6, vlmax); ``` --------- Signed-off-by: adeel10x <adeel.ahmad@10xengineers.ai>
Refactors memory_types and buffer_usage to values from attributes. Progress towards iree-org#20856. This will make it easier to introduce the `hal.resolve_memory_properties` in a follow up PR. --------- Signed-off-by: default <ziereis@roofline.ai> Signed-off-by: Thomas Ziereis <ziereis@roofline.ai> Co-authored-by: Ben Vanik <ben.vanik@gmail.com>
Closes: iree-org#20955 Signed-off-by: NoumanAmir657 <noumanamir453@gmail.com>
…ass (iree-org#21370) The convention of creating a pass is having a `Pass` suffix, so the revision also fixes the names. Signed-off-by: hanhanW <hanhan0912@gmail.com>
) `EmplaceAllocations` doesn't support results with multiple uses but it does support multiple results. This means that it may not always be possible to move the `updateOp` before the `dispatchOp` if `updateOp` is defined by another result of the dispatch. This wasn't handled in `tryMoveProducerBefore`. Issue originated in SDXL unet when testing some other local dispatch creation changes. [This reproducer](https://gist.github.com/IanWood1/7255096856f60bee357b48a0bd3a05d7) is from the full model and can be run with `iree-opt --run-reproducer` Signed-off-by: Ian Wood <ianwood@u.northwestern.edu>
…e-org#21335) This patch adds the `iree-rocdl-use-buffer-instructions` pass to cover some buffers not covered by: `iree-rocdl-configure-buffer-instructions`. In particular this pass will look for `tensor.extact_slices` coming from a `iree_tensor_ext.dispatch.tensor.load` and if the size is less than 2GB and the offsets are uniform, then the pass adds a `iree_gpu.buffer_resource_cast` to use raw buffer ops. This is useful when the main buffer is too big to use raw buffers, but the slices are small enough to use raw buffers. Example: ```mlir func.func @uniform_loop() { %0 = hal.interface.binding.subspan layout(#pipeline_layout) binding(0) alignment(64) : !iree_tensor_ext.dispatch.tensor<readonly:tensor<32x?xi64>> %1 = hal.interface.constant.load layout(#pipeline_layout) ordinal(0) : index %2 = util.assume.int %1<umin = 1, umax = 8589934592> : index %3 = iree_tensor_ext.dispatch.tensor.load %0, offsets = [0, 0], sizes = [32, %2], strides = [1, 1] : !iree_tensor_ext.dispatch.tensor<readonly:tensor<32x?xi64>>{%1} -> tensor<32x?xi64> %4 = hal.interface.constant.load layout(#pipeline_layout) ordinal(1) : index %5 = util.assume.int %1<umin = 1, umax = 8192> : index %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index scf.for %iv = %c0 to %5 step %c1 { %6 = tensor.extract_slice %3[%iv, 0] [32, %5] [1, 1] : tensor<32x?xi64> to tensor<32x?xi64> scf.yield } return } // Result of applying the pass. func.func @uniform_loop() { %0 = hal.interface.binding.subspan layout(#pipeline_layout) binding(0) alignment(64) : !iree_tensor_ext.dispatch.tensor<readonly:tensor<32x?xi64>> %1 = hal.interface.constant.load layout(#pipeline_layout) ordinal(0) : index %2 = util.assume.int %1<umin = 1, umax = 8589934592> : index %3 = iree_tensor_ext.dispatch.tensor.load %0, offsets = [0, 0], sizes = [32, %2], strides = [1, 1] : !iree_tensor_ext.dispatch.tensor<readonly:tensor<32x?xi64>>{%1} -> tensor<32x?xi64> %4 = hal.interface.constant.load layout(#pipeline_layout) ordinal(1) : index %5 = util.assume.int %1<umin = 1, umax = 8192> : index %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index scf.for %arg0 = %c0 to %5 step %c1 { %extracted_slice = tensor.extract_slice %3[%arg0, 0] [32, %5] [1, 1] : tensor<32x?xi64> to tensor<32x?xi64> %6 = iree_gpu.buffer_resource_cast %extracted_slice : tensor<32x?xi64> } return } ``` It's future work to merge `iree-rocdl-use-buffer-instructions` and`iree-rocdl-configure-buffer-instructions` into a single pass. Currently, they are mutually exclusive. --------- Signed-off-by: Fabian Mora <fabian.mora-cordero@amd.com>
…ee-org#21376) Resolves: iree-org#21307 When hoisting encodings for a dispatch containing encodings on a scalar tensor, the hoisted set_encoding on the scalar tensor results in a new encoding dispatch and codegen for this fails in LLVM Translation with: ``` <unknown>:0: error: LLVM Translation failed for operation: builtin.unrealized_conversion_cast ``` Check the issue above for a repro. Signed-off-by: Jorn Tuyls <jorn.tuyls@gmail.com>
Sets the default workgroup distribution to only use the X dim, and delinearize all IDs from it. This prevents bugs with trying to use too many workgroups along Y/Z, which have smaller maximum workgroup counts. --------- Signed-off-by: Max Dawkins <max.dawkins@gmail.com>
This adds end to end tests for split-reduction tiling. To facilitate testing add a new attribute that can be used to specify the reduction tile sizes for the operations on the input MLIR. This is then picked up to drive splitting reductions into two dispatches, one for the partial reduce and one for the final combination. A pipeline test to verify that this attribute works as intended is added as well. Delete the command line option that was used to turn on split-reduction using tiling. --------- Signed-off-by: MaheshRavishankar <mravisha@amd.com>
Adds null handling for indexing maps obtained via `LinalgFusionOpInterface`. Closes iree-org#21378 --------- Signed-off-by: Ian Wood <ianwood@u.northwestern.edu>
…iree-org#21137)" (iree-org#21382) Undo the revert of iree-org@887d75a and iree-org@7f35ecf now that the failure is fixed by iree-org#21334 Signed-off-by: Max Dawkins <max.dawkins@gmail.com>
…ee-org#21251) A single logical device in the HAL may have multiple virtual queues backed by multiple physical devices. Queues may either be host-side (CPU issues all queue operations from a worker thread) or device-side (GPU issues all queue operations from scheduler dispatches). Logical and physical device setup, device library loading, pool management, and routing from the top-level HAL API are the same regardless of queue implementation and it's even possible to support multiple queue types per logical device (if ganging an APU and a dGPU, for example). To start the host-side queue will be implemented in future changes and the device-side queue is present only as a placeholder for later implementation. There are some things I'm not too pleased with (like all the block pools floating around) but I'll be cleaning things up as I implement the queues and try to reduce how much state is shared between the physical device and its queue implementations.
bump to llvm/llvm-project@d9190f8 Carry revert commit iree-org/llvm-project@1aedd94 to llvm/llvm-project@1fdbe69 --------- Signed-off-by: Bangtian Liu <liubangtian@gmail.com>
This makes it easy to build all of the AMDGPU HAL tests at once as part of an IDE action. This is accomplished by adding an optional `group` to cc_test rules that appends each test to a custom target.
…ttr. (iree-org#21392) It also updates the logic in `LLVMCPUTile` pass to use the attribute interface, so it can support `IREE::CPU::LoweringConfigAttr`. --------- Signed-off-by: hanhanW <hanhan0912@gmail.com>
…-org#21391) The revision gets the lowering config interface and relies on the `TilingConfig` to get the vector reduction tile sizes. --------- Signed-off-by: hanhanW <hanhan0912@gmail.com>
There are several times that the failure can only be reproduced with a model, while the root cause is a softmax dispatch. Adding an e2e softmax test to save some time for such case. Signed-off-by: hanhanW <hanhan0912@gmail.com>
…ineLayoutTransformation pass (iree-org#21395) Credit to @Max191 for pointing out the this canonicalization patterns that fix performance degrade in convolution. ```mlir %12 = tensor.empty() : tensor<2x21x64xbf16> %13:2 = scf.forall (%arg3, %arg4, %arg5, %arg6) in (1, 1, 4, 64) shared_outs(%arg7 = %11, %arg8 = %12) -> (tensor<2x2x16x4x16xbf16>, tensor<2x21x64xbf16>) { ... %extracted_slice = tensor.extract_slice %arg8[0, 0, %15, %arg5, %14#2] [2, 2, 4, 1, 1] [1, 1, 1, 1, 1] : tensor<2x2x16x4x16xbf16> to tensor<2x2x4x1x1xbf16> %transposed = linalg.transpose ins(%20 : tensor<2x2x1x4x1xbf16>) outs(%extracted_slice : tensor<2x2x4x1x1xbf16>) permutation = [0, 1, 3, 2, 4] ``` Here `%12` is an empty tensor, so that extract_slice can be folded into a new empty tensor. Now without the canonicalization patterns, this will results in a huge private allocation ` %alloca = memref.alloca() : memref<2x2x16x4x16xbf16, #gpu.address_space<private>>` that eventually results in register spills that kills the performance. Note that this will only happen when input IR has two level nested `scf.forall`, this make it challenging to cleanly put it in as a combine layout transformation pass test. Therefore I'm putting it as a pipeline test instead. Signed-off-by: jerryyin <zhuoryin@amd.com>
…rg#21363) Now that we properly support unaligned/dynamic shapes in vector distribute pipeline, add a configuration for attention to use multiple subgroups for memory bound attention. Verified on MI300-X for accuracy and a speedup of 2x for sequence length ~4096 (unaligned, dynamic).
…g. (iree-org#21397) The revision updates the static constructor with unique_ptr because there are invalid states. E.g., the config does not recognize GPU lowering config and different structured Codegen lowering configs. In this context, we prefer factory functions to initializer methods. See https://abseil.io/tips/42 for more details. Furthermore, it updates the logic of `getVectorTileSizes`, which uses the mapping level index to get the right list. Ideally, we can implement an interface method for all the specific lowering config, but we still need the generic lowering config being able to handle the case. Thus, we only update `TilingConfig` logic now, and we will do the interface way later. --------- Signed-off-by: hanhanW <hanhan0912@gmail.com>
…onsumer pass. (iree-org#21394) The fusion order could change the IR order, e.g., softmax dispatch. It makes sure that the root op is as the same as the op that is chosen in lowering strategy selection. --------- Signed-off-by: hanhanW <hanhan0912@gmail.com>
bump to llvm/llvm-project@e0cce5c Carry revert commit iree-org/llvm-project@1aedd94 to llvm/llvm-project@1fdbe69 --------- Signed-off-by: Bangtian Liu <liubangtian@gmail.com>
Adds flag (default to enable) that can be used to optionally disable multi-use fusion. Fixes failed scatter fusion in llama 405b. --------- Signed-off-by: Ian Wood <ianwood@u.northwestern.edu>
…ree-org#21404) The pack op in the old test is fusible with fixes. Update the test to use imperfect tiling case that is known infusible atm. Signed-off-by: hanhanW <hanhan0912@gmail.com>
Signed-off-by: default <ziereis@roofline.ai>
Signed-off-by: default <ziereis@roofline.ai>
Signed-off-by: default <ziereis@roofline.ai>
ziereis
pushed a commit
that referenced
this pull request
Aug 5, 2025
feat(Broadcom): add SPIRV GPU Config fix(GPUUtils): add videocore to the target list fix: pre-commit failures (iree-org#7)
ziereis
added a commit
that referenced
this pull request
Aug 5, 2025
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
20 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
(wip) progress towards iree-org#20856.
Adds the resolve_memory_properties op and pass to resolve it statically