Skip to content

[codegen] NestedLayoutAttr::reshape changes thread ownership for non-contiguous shape casts #24955

Description

@snarang181

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.

Activity

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Labels

    No labels
    No labels

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions