gistfile1.txt

shader: MESA_SHADER_COMPUTE
source_sha1: {0x945f51b1, 0xcf3b29d0, 0x7186ab9c, 0xa3d0bcb8, 0x0e446417}
internal: false
workgroup_size: 0, 0, 0 (variable)
stage: 5
next_stage: 0
num_ubos: 1
float_controls_execution_mode: 0x00038011
subgroup_size: 0
divergence_analysis_run: true
first_ubo_is_default_ubo: true
writes_memory: true
ptr_size: 64
inputs: 0
outputs: 0
uniforms: 0
scratch: 20
constants: 20
decl_var uniform INTERP_MODE_NONE none uint64_t #0 (0, 0, 0)
decl_var uniform INTERP_MODE_NONE none uint64_t constant_buffer_addr (2, 8, 0)
decl_var ubo INTERP_MODE_NONE none uint8_t[16] kernel_input (0, 0, 0)
decl_var constant INTERP_MODE_NONE none struct.GFX125_3DSTATE_VERTEX_BUFFERS __const.write_3DSTATE_VERTEX_BUFFERS.v = { { 0x00000003 }, { 0x00000008 }, { 0x00000000 }, { 0x00000003 }, { 0x00000003 } }
decl_function __wrapped_write_3DSTATE_VERTEX_BUFFERS (0 params)

impl __wrapped_write_3DSTATE_VERTEX_BUFFERS {
    block b0:      // preds: 
    con 32    %0 = load_const (0x00000000)
    con 64    %1 = @load_ubo (%0 (0x0), %0 (0x0)) (access=none, align_mul=1073741824, align_offset=0, range_base=0, range=16)
    con 32    %2 = load_const (0x00000003 = 0.000000)
                   @store_global (%2 (0x3), %1) (wrmask=x, access=none, align_mul=4, align_offset=0)
                   // succs: b1 
    block b1:
}

tmp.patch

diff --git a/src/compiler/nir/nir_deref.c b/src/compiler/nir/nir_deref.c
index f0f5f140130..ff0930d06e4 100644
--- a/src/compiler/nir/nir_deref.c
+++ b/src/compiler/nir/nir_deref.c
@@ -1210,7 +1210,7 @@ opt_replace_struct_wrapper_cast(nir_builder *b, nir_deref_instr *cast)
       return false;
 
    /* we can't drop the stride information */
-   if (cast->cast.ptr_stride != glsl_get_explicit_stride(field_type))
+   if (cast->cast.ptr_stride != glsl_get_explicit_stride(field_type) && glsl_get_explicit_stride(field_type))
       return false;
 
    nir_deref_instr *replace = nir_build_deref_struct(b, parent, 0);