Summary
Vector layout propagation reshapes a nested_layout through
vector.shape_cast even when one thread's row-major source elements become a
non-contiguous target set. Vector distribution then performs only a local
shape cast and stores a contiguous tile, silently misplacing eight of sixteen elements.
Reproducer
#l = #iree_vector_ext.nested_layout<
subgroup_tile=[1,1],batch_tile=[1,1],outer_tile=[1,1],
thread_tile=[2,2],element_tile=[2,2],
subgroup_strides=[0,0],thread_strides=[1,2]>
#t = #iree_codegen.translation_info<
pipeline=#iree_codegen.no_pipeline workgroup_size=[4,1,1] subgroup_size=4>
func.func @case(%src: memref<4x4xf32>,%dst: memref<2x8xf32>)
attributes {translation_info=#t} {
%c0=arith.constant 0:index
%p=arith.constant 0.0:f32
%v=vector.transfer_read %src[%c0,%c0],%p {in_bounds=[true,true]} : memref<4x4xf32>,vector<4x4xf32>
%x=iree_vector_ext.to_layout %v to layout(#l) : vector<4x4xf32>
%r=vector.shape_cast %x : vector<4x4xf32> to vector<2x8xf32>
vector.transfer_write %r,%dst[%c0,%c0] {in_bounds=[true,true]} : vector<2x8xf32>,memref<2x8xf32>
return
}
Verified with iree-opt at a5ca3e2beaa8 (LLVM 24.0.0git, optimized with assertions):
$ iree-opt '--pass-pipeline=builtin.module(func.func(iree-llvmgpu-vector-distribute))' repro.mlir | grep -E 'affine.linearize_index|vector.shape_cast' | head -7
%0 = affine.linearize_index disjoint [%thread_id_z, %thread_id_y, %thread_id_x] by (1, 1, 4) : index
%3 = affine.linearize_index disjoint [%1#2, %c0, %c0, %2#2, %c0] by (1, 1, 1, 2, 2) : index
%4 = affine.linearize_index disjoint [%1#1, %c0, %c0, %2#1, %c0] by (1, 1, 1, 2, 2) : index
%7 = vector.shape_cast %6 : vector<1x1x1x1x2x2xf32> to vector<1x1x1x1x1x4xf32>
%12 = affine.linearize_index disjoint [%8#2, %c0, %c0, %9#2, %c0] by (1, 1, 1, 2, 1) : index
%13 = affine.linearize_index disjoint [%8#1, %c0, %c0, %9#1, %c0] by (1, 1, 1, 2, 4) : index
Expected thread 2 target columns are {2,3,6,7}; the inferred layout assigns
{4,5,6,7}. For input 0..15, rows become [0,1,4,5,2,3,6,7] and
[8,9,12,13,10,11,14,15] instead of row-major order.
Agent Analysis
At a5ca3e2beaa8, compiler/src/iree/compiler/Codegen/Dialect/VectorExt/IR/ VectorExtAttrs.cpp:105 implements NestedLayoutAttr::reshape; forward
propagation calls it from VectorLayoutAnalysis.cpp:319 without checking
ownership contiguity. Decline the reshape or insert an explicit redistribution.
AI tool use disclosure
Adding this to comply with the project's
AI tool use policy
Assisted-by: Claude Code
Reproducer and root-cause analysis were prepared with LLM assistance and verified manually on the referenced commit.
Summary
Vector layout propagation reshapes a
nested_layoutthroughvector.shape_casteven when one thread's row-major source elements become anon-contiguous target set. Vector distribution then performs only a local
shape cast and stores a contiguous tile, silently misplacing eight of sixteen elements.
Reproducer
Verified with
iree-optata5ca3e2beaa8(LLVM 24.0.0git, optimized with assertions):Expected thread 2 target columns are
{2,3,6,7}; the inferred layout assigns{4,5,6,7}. For input0..15, rows become[0,1,4,5,2,3,6,7]and[8,9,12,13,10,11,14,15]instead of row-major order.Agent Analysis
At
a5ca3e2beaa8,compiler/src/iree/compiler/Codegen/Dialect/VectorExt/IR/ VectorExtAttrs.cpp:105implementsNestedLayoutAttr::reshape; forwardpropagation calls it from
VectorLayoutAnalysis.cpp:319without checkingownership contiguity. Decline the reshape or insert an explicit redistribution.
AI tool use disclosure
Adding this to comply with the project's
AI tool use policy
Assisted-by: Claude Code
Reproducer and root-cause analysis were prepared with LLM assistance and verified manually on the referenced commit.