Unknown drm_i915_query_item id=5 Unknown DRM_IOCTL_I915_GET_PARAM 58 Unknown drm_i915_query_item id=5 Unknown DRM_IOCTL_I915_GET_PARAM 58 NIR (from SPIR-V) for MESA_SHADER_COMPUTE shader: shader: MESA_SHADER_COMPUTE source_blake3: {0xc7e4853e, 0x84c91157, 0x76ae5910, 0x1481f8ee, 0x19d97bcb, 0x80de01ed, 0xd583ca35, 0xa5ccba6a} internal: false workgroup_size: 1, 1, 1 stage: 5 next_stage: 0 num_ssbos: 3 subgroup_size: 2 ptr_size: 0 inputs: 0 outputs: 0 uniforms: 0 decl_var system INTERP_MODE_NONE none uvec3 gl_GlobalInvocationID (SYSTEM_VALUE_GLOBAL_INVOCATION_ID) decl_var ssbo INTERP_MODE_NONE restrict block #0 (~0, 0, 0) decl_var ssbo INTERP_MODE_NONE restrict block #1 (~0, 0, 1) decl_var ssbo INTERP_MODE_NONE restrict block #2 (~0, 0, 2) decl_function main (0 params) impl main { con block b0: // preds: 32 %0 = deref_var &gl_GlobalInvocationID (system uvec3) 32x3 %1 = @load_deref (%0) (access=none) 32 %2 = mov %1.x 32 %3 = load_const (0x00000000 = 0.000000) 32x4 %4 = @vulkan_resource_index (%3 (0x0)) (desc_set=0, binding=0, desc_type=SSBO) 32x4 %5 = @load_vulkan_descriptor (%4) (desc_type=SSBO) 32x4 %6 = deref_cast (block *)%5 (ssbo block) (ptr_stride=0, align_mul=4, align_offset=0) 32x4 %7 = deref_struct &%6->field0 (ssbo int[]) // &((block *)%5)->field0 32x4 %8 = deref_array &(*%7)[%2] (ssbo int) // &((block *)%5)->field0[%2] 32 %9 = @load_deref (%8) (access=none) 32 %10 = load_const (0x00000000 = 0.000000) 32x4 %11 = @vulkan_resource_index (%10 (0x0)) (desc_set=0, binding=1, desc_type=SSBO) 32x4 %12 = @load_vulkan_descriptor (%11) (desc_type=SSBO) 32x4 %13 = deref_cast (block *)%12 (ssbo block) (ptr_stride=0, align_mul=4, align_offset=0) 32x4 %14 = deref_struct &%13->field0 (ssbo int[]) // &((block *)%12)->field0 32x4 %15 = deref_array &(*%14)[%2] (ssbo int) // &((block *)%12)->field0[%2] 32 %16 = @load_deref (%15) (access=none) 32 %17 = load_const (0x00000000) 32 %18 = sdot_4x8_iadd %9, %16, %17 (0x0) 32 %19 = load_const (0x00000000 = 0.000000) 32x4 %20 = @vulkan_resource_index (%19 (0x0)) (desc_set=0, binding=2, desc_type=SSBO) 32x4 %21 = @load_vulkan_descriptor (%20) (desc_type=SSBO) 32x4 %22 = deref_cast (block *)%21 (ssbo block) (ptr_stride=0, align_mul=4, align_offset=0) 32x4 %23 = deref_struct &%22->field0 (ssbo int[]) // &((block *)%21)->field0 32x4 %24 = deref_array &(*%23)[%2] (ssbo int) // &((block *)%21)->field0[%2] @store_deref (%24, %18) (wrmask=x, access=none) // succs: b1 block b1: } NIR (SSA form) for compute shader: shader: MESA_SHADER_COMPUTE source_blake3: {0xc7e4853e, 0x84c91157, 0x76ae5910, 0x1481f8ee, 0x19d97bcb, 0x80de01ed, 0xd583ca35, 0xa5ccba6a} internal: false workgroup_size: 1, 1, 1 stage: 5 next_stage: 0 num_ssbos: 3 system_values_read: 0x00000000'00001800'00000000 subgroup_size: 2 divergence_analysis_run: true bit_sizes_int: 0x20 separate_shader: true writes_memory: true ptr_size: 0 inputs: 0 outputs: 0 uniforms: 12 decl_var ssbo INTERP_MODE_NONE restrict readonly block #0 (~0, 0, 0) decl_var ssbo INTERP_MODE_NONE restrict readonly block #1 (~0, 0, 1) decl_var ssbo INTERP_MODE_NONE restrict writeonly block #2 (~0, 0, 2) decl_function main (0 params) impl main { con block b0: // preds: con 32x3 %0 = @load_workgroup_id con 32 %1 = load_const (0x00000000) con 32x3 %2 = @load_uniform (%1 (0x0)) (base=0, range=12, dest_type=uint) con 32 %3 = iadd %0.x, %2.x con 32 %4 = load_const (0xdeaddead = -6264320714451451904.000000 = -559030611 = 3735936685) con 32 %5 = load_const (0x00000002) con 32 %6 = @resource_intel (%4 (0xdeaddead), %5 (0x2), %1 (0x0), %1 (0x0)) (desc_set=0, binding=0, resource_intel=0, resource_block_intel=-1) con 32 %7 = ishl %3, %5 (0x2) con 32 %8 = @load_ssbo (%6, %7) (access=readonly|reorderable, align_mul=4, align_offset=0) con 32 %9 = load_const (0x00000003 = 0.000000) con 32 %10 = @resource_intel (%4 (0xdeaddead), %9 (0x3), %1 (0x0), %1 (0x0)) (desc_set=0, binding=1, resource_intel=0, resource_block_intel=-1) con 32 %11 = @load_ssbo (%10, %7) (access=readonly|reorderable, align_mul=4, align_offset=0) con 32 %12 = sdot_4x8_iadd %8, %11, %1 (0x0) con 32 %13 = load_const (0x00000004 = 0.000000) con 32 %14 = @resource_intel (%4 (0xdeaddead), %13 (0x4), %1 (0x0), %1 (0x0)) (desc_set=0, binding=2, resource_intel=0, resource_block_intel=-1) @store_ssbo (%12, %14, %7) (wrmask=x, access=writeonly, align_mul=4, align_offset=0) // succs: b1 block b1: } NIR (final form) for compute shader: shader: MESA_SHADER_COMPUTE source_blake3: {0xc7e4853e, 0x84c91157, 0x76ae5910, 0x1481f8ee, 0x19d97bcb, 0x80de01ed, 0xd583ca35, 0xa5ccba6a} internal: false workgroup_size: 1, 1, 1 stage: 5 next_stage: 0 num_ssbos: 3 system_values_read: 0x00000000'00001800'00000000 subgroup_size: 2 divergence_analysis_run: true bit_sizes_int: 0x20 separate_shader: true writes_memory: true ptr_size: 0 inputs: 0 outputs: 0 uniforms: 12 decl_var ssbo INTERP_MODE_NONE restrict readonly block #0 (~0, 0, 0) decl_var ssbo INTERP_MODE_NONE restrict readonly block #1 (~0, 0, 1) decl_var ssbo INTERP_MODE_NONE restrict writeonly block #2 (~0, 0, 2) decl_function main (0 params) impl main { con block b0: // preds: con 32x3 %0 = @load_workgroup_id con 32 %1 = load_const (0x00000000) con 32x3 %2 = @load_uniform (%1 (0x0)) (base=0, range=12, dest_type=uint) con 32 %3 = iadd %0.x, %2.x con 32 %4 = load_const (0xdeaddead = -6264320714451451904.000000 = -559030611 = 3735936685) con 32 %5 = load_const (0x00000002) con 32 %6 = @resource_intel (%4 (0xdeaddead), %5 (0x2), %1 (0x0), %1 (0x0)) (desc_set=0, binding=0, resource_intel=0, resource_block_intel=-1) con 32 %7 = ishl %3, %5 (0x2) con 32 %8 = @load_ssbo (%6, %7) (access=readonly|reorderable, align_mul=4, align_offset=0) con 32 %9 = load_const (0x00000003 = 0.000000) con 32 %10 = @resource_intel (%4 (0xdeaddead), %9 (0x3), %1 (0x0), %1 (0x0)) (desc_set=0, binding=1, resource_intel=0, resource_block_intel=-1) con 32 %11 = @load_ssbo (%10, %7) (access=readonly|reorderable, align_mul=4, align_offset=0) con 32 %12 = sdot_4x8_iadd %8, %11, %1 (0x0) con 32 %13 = load_const (0x00000004 = 0.000000) con 32 %14 = @resource_intel (%4 (0xdeaddead), %13 (0x4), %1 (0x0), %1 (0x0)) (desc_set=0, binding=2, resource_intel=0, resource_block_intel=-1) @store_ssbo (%12, %14, %7) (wrmask=x, access=writeonly, align_mul=4, align_offset=0) // succs: b1 block b1: } Native code for unnamed compute shader (null) (src_hash 0x00000000) (sha1 4a83bbedc93c93bfbd9ac48fabfd8034872d31d1) SIMD8 shader: 16 instructions. 0 loops. 202 cycles. 0:0 spills:fills, 4 sends, scheduled with mode top-down. Promoted 0 constants. Compacted 256 to 256 bytes (0%) START B0 (202 cycles) mov(8) g7<1>UD g0.1<0,1,0>UD { align1 WE_all 1Q }; mov(8) g2<1>D 0D { align1 WE_all 1Q }; add(8) g3<1>D g7<0,1,0>D g1<0,1,0>D { align1 WE_all 1Q @2 compacted }; shl(8) g4<1>D g3<0,1,0>D 0x00000002UD { align1 WE_all 1Q @1 }; sync nop(1) null<0,1,0>UB { align1 WE_all 1N @1 }; mov(8) g8<1>D g4<0,1,0>D { align1 1Q }; send(8) g5UD g8UD nullUD 0x02106e02 0x00000000 dp data 1 MsgDesc: (untyped surface read, Surface = 2, SIMD8, Mask = 0xe) mlen 1 ex_mlen 0 rlen 1 { align1 1Q @1 $0 }; sync nop(1) null<0,1,0>UB { align1 WE_all 1N $0.src }; send(8) g6UD g8UD nullUD 0x02106e03 0x00000000 dp data 1 MsgDesc: (untyped surface read, Surface = 3, SIMD8, Mask = 0xe) mlen 1 ex_mlen 0 rlen 1 { align1 1Q $1 }; sync nop(1) null<0,1,0>UB { align1 WE_all 1N $1.dst }; sync nop(1) null<0,1,0>UB { align1 WE_all 1N @4 }; dp4a(8) g9<1>D g2.0<0,1,0>D g5<8,8,1>D g6<1,1,1>D { align1 1Q $0.dst }; sync nop(1) null<0,1,0>UB { align1 WE_all 1N $1.src }; send(8) nullUD g8UD g9UD 0x02026e04 0x00000040 dp data 1 MsgDesc: (DC untyped surface write, Surface = 4, SIMD8, Mask = 0xe) mlen 1 ex_mlen 1 rlen 0 { align1 1Q @1 $2 }; mov(8) g126<1>UD g0<8,8,1>UD { align1 WE_all 1Q }; send(8) nullUD g126UD nullUD 0x02000000 0x00000000 thread_spawner MsgDesc: mlen 1 ex_mlen 0 rlen 0 { align1 WE_all 1Q @1 EOT }; END B0