zink: nir validation failures in Sparse code
On Intel hardware, if I apply both !29197 (merged) and !29337 (merged), then run KHR-GL46.sparse_texture2_tests.UncommittedRegionsAccess
, I get:
NIR validation failed after strip_tex_ms in ../../src/gallium/drivers/zink/zink_compiler.c
8 errors:
shader: MESA_SHADER_COMPUTE
source_blake3: {0x7f1359bd, 0xda403135, 0xc92d3f61, 0x0840ecb1, 0x5f94014e, 0x3969bd0a, 0x5572d2f8, 0x975af163}
name: GLSL669
internal: false
workgroup_size: 1, 1, 1
stage: 5
next_stage: 4
num_ubos: 1
num_images: 2
system_values_read: 0x00000000'00000800'00000000
images_used: 0x00000000'00000003
msaa_images: 0x00000000'00000002
subgroup_size: 1
bit_sizes_int: 0x21
first_ubo_is_default_ubo: true
flrp_lowered: true
writes_memory: true
ptr_size: 0
inputs: 0
outputs: 0
uniforms: 1
decl_var uniform INTERP_MODE_NONE none int widthCommitted (2, 0, 0)
decl_var image INTERP_MODE_NONE writeonly r8_uint uimage2D uni_out_image (0, 0, 0)
decl_var image INTERP_MODE_NONE none r32_sint iimage2D uni_in_image (1, 1, 1)
decl_var ubo INTERP_MODE_NONE none vec4[1] uniform_0 (0, 0, 0)
decl_function main (0 params)
impl main {
con block b0: // preds:
32 %0 = load_const (0x0000007f = 127)
32 %1 = load_const (0x00000000)
32 %2 = load_const (0x0000000f = 0.000000 = 15)
32 %3 = load_const (0x00000001 = 0.000000)
32 %4 = load_const (0x00000002 = 0.000000)
32 %5 = load_const (0x00000003)
32 %6 = load_const (0x00000004 = 0.000000)
32 %7 = load_const (0x00000005 = 0.000000)
32 %8 = load_const (0x00000006 = 0.000000)
32 %9 = load_const (0x00000007)
32x4 %10 = load_const (0x00000000, 0x00000000, 0x00000000, 0x00000000) = (0.000000, 0.000000, 0.000000, 0.000000)
32x4 %11 = load_const (0x00000001, 0x00000001, 0x00000001, 0x00000001) = (0.000000, 0.000000, 0.000000, 0.000000)
32 %12 = undefined
32 %13 = undefined
32 %14 = undefined
32 %15 = undefined
32 %16 = undefined
32 %17 = undefined
32 %18 = undefined
32 %19 = undefined
32 %20 = undefined
32 %21 = undefined
32 %22 = undefined
32 %23 = undefined
32 %24 = undefined
32 %25 = undefined
32 %26 = undefined
32 %27 = undefined
32 %28 = undefined
32 %29 = undefined
32 %30 = undefined
32 %31 = undefined
32 %32 = undefined
32 %33 = undefined
32 %34 = undefined
32 %35 = undefined
32 %36 = undefined
32 %37 = undefined
32 %38 = undefined
32 %39 = undefined
32 %40 = undefined
32 %41 = undefined
32x3 %107 = @load_workgroup_id
32 %115 = @load_ubo (%1 (0x0), %1 (0x0)) (access=none, align_mul=1073741824, align_offset=0, range_base=0, range=4)
@barrier (execution_scope=NONE, memory_scope=DEVICE, mem_semantics=ACQ|REL, mem_modes=0)
1 %46 = ige %107.x, %115
// succs: b1 b32
if %46 {
con block b1: // preds: b0
32 %124 = isub %107.x, %115
32 %116 = ishl %107.y, %5 (0x3)
32 %50 = iadd %124, %116
32 %51 = iand %50, %9 (0x7)
32x2 %125 = load_const (0x00000000, 0x00000001) = (0.000000, 0.000000)
32x3 %129 = load_const (0x00000000, 0x00000001, 0x00000002) = (0.000000, 0.000000, 0.000000)
32x4 %132 = load_const (0x00000000, 0x00000001, 0x00000002, 0x00000003)
1x4 %133 = ieq %51.xxxx, %132 (0x0, 0x1, 0x2, 0x3)
1 %134 = mov %133.w
1 %131 = mov %133.z
1 %127 = mov %133.x
1 %128 = mov %133.y
// succs: b2 b3
if %127 {
con block b2: // preds: b1
32 %53 = deref_var &uni_in_image (image iimage2DMS)
error: instr->type == instr->var->type (../../src/compiler/nir/nir_validate.c:317)
32x4 %54 = vec4 %107.x, %107.y, %41, %40
32 %55 = @image_deref_atomic (%53, %54, %1 (0x0), %2 (0xf)) (image_dim=2D-MSAA, image_array=false, format=none, access=none, atomic_op=xchg)
// succs: b25
} else {
con block b3: // preds: b1, succs: b4 b5
if %128 {
con block b4: // preds: b3
32 %57 = deref_var &uni_in_image (image iimage2DMS)
error: instr->type == instr->var->type (../../src/compiler/nir/nir_validate.c:317)
32x4 %58 = vec4 %107.x, %107.y, %39, %38
32 %59 = @image_deref_atomic_swap (%57, %58, %1 (0x0), %1 (0x0), %2 (0xf)) (image_dim=2D-MSAA, image_array=false, format=none, access=none, atomic_op=cmpxchg)
// succs: b24
} else {
con block b5: // preds: b3, succs: b6 b7
if %131 {
con block b6: // preds: b5
32 %61 = deref_var &uni_in_image (image iimage2DMS)
error: instr->type == instr->var->type (../../src/compiler/nir/nir_validate.c:317)
32x4 %62 = vec4 %107.x, %107.y, %37, %36
32 %63 = @image_deref_atomic (%61, %62, %1 (0x0), %2 (0xf)) (image_dim=2D-MSAA, image_array=false, format=none, access=none, atomic_op=iadd)
// succs: b23
} else {
con block b7: // preds: b5, succs: b8 b9
if %134 {
con block b8: // preds: b7
32 %65 = deref_var &uni_in_image (image iimage2DMS)
error: instr->type == instr->var->type (../../src/compiler/nir/nir_validate.c:317)
32x4 %66 = vec4 %107.x, %107.y, %35, %34
32 %67 = @image_deref_atomic (%65, %66, %1 (0x0), %2 (0xf)) (image_dim=2D-MSAA, image_array=false, format=none, access=none, atomic_op=iand)
// succs: b22
} else {
con block b9: // preds: b7
32x2 %135 = load_const (0x00000004, 0x00000005) = (0.000000, 0.000000)
32x3 %139 = load_const (0x00000004, 0x00000005, 0x00000006) = (0.000000, 0.000000, 0.000000)
32x4 %142 = load_const (0x00000004, 0x00000005, 0x00000006, 0x00000007)
1x4 %143 = ieq %51.xxxx, %142 (0x4, 0x5, 0x6, 0x7)
1 %144 = mov %143.w
1 %141 = mov %143.z
1 %137 = mov %143.x
1 %138 = mov %143.y
// succs: b10 b11
if %137 {
con block b10: // preds: b9
32 %69 = deref_var &uni_in_image (image iimage2DMS)
error: instr->type == instr->var->type (../../src/compiler/nir/nir_validate.c:317)
32x4 %70 = vec4 %107.x, %107.y, %33, %32
32 %71 = @image_deref_atomic (%69, %70, %1 (0x0), %2 (0xf)) (image_dim=2D-MSAA, image_array=false, format=none, access=none, atomic_op=ior)
// succs: b21
} else {
con block b11: // preds: b9, succs: b12 b13
if %138 {
con block b12: // preds: b11
32 %73 = deref_var &uni_in_image (image iimage2DMS)
error: instr->type == instr->var->type (../../src/compiler/nir/nir_validate.c:317)
32x4 %74 = vec4 %107.x, %107.y, %31, %30
32 %75 = @image_deref_atomic (%73, %74, %1 (0x0), %2 (0xf)) (image_dim=2D-MSAA, image_array=false, format=none, access=none, atomic_op=ixor)
// succs: b20
} else {
con block b13: // preds: b11, succs: b14 b15
if %141 {
con block b14: // preds: b13
32 %77 = deref_var &uni_in_image (image iimage2DMS)
error: instr->type == instr->var->type (../../src/compiler/nir/nir_validate.c:317)
32x4 %78 = vec4 %107.x, %107.y, %29, %28
32 %79 = @image_deref_atomic (%77, %78, %1 (0x0), %2 (0xf)) (image_dim=2D-MSAA, image_array=false, format=none, access=none, atomic_op=imin)
// succs: b19
} else {
con block b15: // preds: b13, succs: b16 b17
if %144 {
con block b16: // preds: b15
32 %81 = deref_var &uni_in_image (image iimage2DMS)
error: instr->type == instr->var->type (../../src/compiler/nir/nir_validate.c:317)
32x4 %82 = vec4 %107.x, %107.y, %27, %26
32 %83 = @image_deref_atomic (%81, %82, %1 (0x0), %2 (0xf)) (image_dim=2D-MSAA, image_array=false, format=none, access=none, atomic_op=imax)
// succs: b18
} else {
con block b17: // preds: b15, succs: b18
}
con block b18: // preds: b16 b17
32 %84 = phi b16: %83, b17: %0 (0x7f)
// succs: b19
}
con block b19: // preds: b14 b18
32 %85 = phi b14: %79, b18: %84
// succs: b20
}
con block b20: // preds: b12 b19
32 %86 = phi b12: %75, b19: %85
// succs: b21
}
con block b21: // preds: b10 b20
32 %87 = phi b10: %71, b20: %86
// succs: b22
}
con block b22: // preds: b8 b21
32 %88 = phi b8: %67, b21: %87
// succs: b23
}
con block b23: // preds: b6 b22
32 %89 = phi b6: %63, b22: %88
// succs: b24
}
con block b24: // preds: b4 b23
32 %90 = phi b4: %59, b23: %89
// succs: b25
}
con block b25: // preds: b2 b24
32 %91 = phi b2: %55, b24: %90
32 %92 = deref_var &uni_in_image (image iimage2D)
32x4 %93 = vec4 %107.x, %107.y, %25, %24
32x4 %94 = @image_deref_load (%92, %93, %1 (0x0), %1 (0x0)) (image_dim=2D, image_array=false, format=none, access=none, dest_type=int32)
1 %95 = ieq %91, %1 (0x0)
// succs: b26 b27
if %95 {
con block b26: // preds: b25
32 %96 = deref_var &uni_out_image (image uimage2D)
32x4 %97 = vec4 %124, %107.y, %23, %22
@image_deref_store (%96, %97, %21, %10 (0x0, 0x0, 0x0, 0x0), %1 (0x0)) (image_dim=2D, image_array=false, format=none, access=writeonly, src_type=uint32)
// succs: b28
} else {
con block b27: // preds: b25
32 %98 = deref_var &uni_out_image (image uimage2D)
32x4 %99 = vec4 %124, %107.y, %20, %19
32x4 %100 = vec4 %91, %91, %91, %91
@image_deref_store (%98, %99, %18, %100, %1 (0x0)) (image_dim=2D, image_array=false, format=none, access=writeonly, src_type=uint32)
// succs: b28
}
con block b28: // preds: b26 b27
1 %101 = ieq %94.x, %1 (0x0)
// succs: b29 b30
if %101 {
con block b29: // preds: b28
32 %102 = deref_var &uni_out_image (image uimage2D)
32x4 %103 = vec4 %107.x, %107.y, %17, %16
@image_deref_store (%102, %103, %15, %10 (0x0, 0x0, 0x0, 0x0), %1 (0x0)) (image_dim=2D, image_array=false, format=none, access=writeonly, src_type=uint32)
// succs: b31
} else {
con block b30: // preds: b28
32 %104 = deref_var &uni_out_image (image uimage2D)
32x4 %105 = vec4 %107.x, %107.y, %14, %13
@image_deref_store (%104, %105, %12, %11 (0x1, 0x1, 0x1, 0x1), %1 (0x0)) (image_dim=2D, image_array=false, format=none, access=writeonly, src_type=uint32)
// succs: b31
}
con block b31: // preds: b29 b30, succs: b33
} else {
con block b32: // preds: b0, succs: b33
}
con block b33: // preds: b31 b32, succs: b34
block b34:
}
On a release build the test passes.