freedreno/ir3/ra: Fix validate crash with shared reg file
Fixes a crash I was hitting with clover, with a shared reg accessed from a block other than the one it was defined in:
ra validation fail: wrong definition reaches source ssa_13:3 + 0
expected: ssa_12:1(r48.x) + 0
actual: test_basic: ../src/freedreno/ir3/ir3_ra_validate.c:429: void dump_reg_state(struct reg_state *): Assertion `state->def != UNKNOWN' failed.
Thread 1 "test_basic" received signal SIGABRT, Aborted.
0x0000007ff7123750 in ?? () from /lib64/libc.so.6
(gdb) set pagination off
(gdb) bt
#0 0x0000007ff7123750 in ?? () from /lib64/libc.so.6
#1 0x0000007ff710f7dc in ?? () from /lib64/libc.so.6
#2 0x0000007ff711c9c8 in ?? () from /lib64/libc.so.6
#3 0x0000007ff711ca40 in ?? () from /lib64/libc.so.6
#4 0x0000007fedbb35bc in dump_reg_state (state=0x7fffffd0e0) at ../src/freedreno/ir3/ir3_ra_validate.c:429
#5 0x0000007fedbb32a4 in check_reaching_src (ctx=0x1b77170, instr=0x1b25510, src=0x1b256c0) at ../src/freedreno/ir3/ir3_ra_validate.c:462
#6 0x0000007fedbb3148 in check_reaching_instr (ctx=0x1b77170, instr=0x1b25510) at ../src/freedreno/ir3/ir3_ra_validate.c:480
#7 0x0000007fedbb2498 in check_reaching_block (ctx=0x1b77170, block=0x1b350e0) at ../src/freedreno/ir3/ir3_ra_validate.c:490
#8 0x0000007fedbb1dc0 in check_reaching_defs (ctx=0x1b77170, ir=0x1b2af40) at ../src/freedreno/ir3/ir3_ra_validate.c:532
#9 0x0000007fedbb19bc in ir3_ra_validate (v=0x21a0f00, full_size=384, half_size=0, block_count=10) at ../src/freedreno/ir3/ir3_ra_validate.c:567
#10 0x0000007fedbab0fc in ir3_ra (v=0x21a0f00) at ../src/freedreno/ir3/ir3_ra.c:2211
#11 0x0000007fedb85a84 in ir3_compile_shader_nir (compiler=0x1a4b220, so=0x21a0f00) at ../src/freedreno/ir3/ir3_compiler_nir.c:4297
#12 0x0000007fedb79378 in compile_variant (v=0x21a0f00) at ../src/freedreno/ir3/ir3_shader.c:293
#13 0x0000007fedb775ac in create_variant (shader=0x1b29bc0, key=0x7fffffd6c0, write_disasm=false) at ../src/freedreno/ir3/ir3_shader.c:388
#14 0x0000007fedb772e0 in ir3_shader_get_variant (shader=0x1b29bc0, key=0x7fffffd6c0, binning_pass=false, write_disasm=false, created=0x7fffffd69f) at ../src/freedreno/ir3/ir3_shader.c:425
#15 0x0000007feda66040 in ir3_shader_variant (shader=0x1b29bc0, key=..., binning_pass=false, debug=0x1a5e3c0) at ../src/gallium/drivers/freedreno/ir3/ir3_gallium.c:133
#16 0x0000007feda66754 in ir3_shader_compute_state_create (pctx=0x1a596c0, cso=0x1ec3ea0) at ../src/gallium/drivers/freedreno/ir3/ir3_gallium.c:326
#17 0x0000007ff768f834 in clover::kernel::exec_context::bind (this=0x1ec3dd0, _q=..., grid_offset=std::vector of length 1, capacity 1 = {...}) at ../src/gallium/frontends/clover/core/kernel.cpp:286
#18 0x0000007ff768e6a4 in clover::kernel::launch (this=0x1ec3d50, q=..., grid_offset=std::vector of length 1, capacity 1 = {...}, grid_size=std::vector of length 1, capacity 1 = {...}, block_size=std::vector of length 1, capacity 1 = {...}) at ../src/gallium/frontends/clover/core/kernel.cpp:74
#19 0x0000007ff7625730 in clEnqueueNDRangeKernel::$_1::operator() (this=0x1b2b520) at ../src/gallium/frontends/clover/api/kernel.cpp:333
#20 0x0000007ff76256e0 in std::__invoke_impl<void, clEnqueueNDRangeKernel::$_1&, clover::event&> (__f=..., __args=...) at /usr/lib/gcc/aarch64-redhat-linux/11/../../../../include/c++/11/bits/invoke.h:61
#21 0x0000007ff7625674 in std::__invoke_r<void, clEnqueueNDRangeKernel::$_1&, clover::event&> (__fn=..., __args=...) at /usr/lib/gcc/aarch64-redhat-linux/11/../../../../include/c++/11/bits/invoke.h:154
#22 0x0000007ff7625520 in std::_Function_handler<void (clover::event&), clEnqueueNDRangeKernel::$_1>::_M_invoke(std::_Any_data const&, clover::event&) (__functor=..., __args=...) at /usr/lib/gcc/aarch64-redhat-linux/11/../../../../include/c++/11/bits/std_function.h:291
#23 0x0000007ff7684eb0 in std::function<void (clover::event&)>::operator()(clover::event&) const (this=0x1f7b7e0, __args=...) at /usr/lib/gcc/aarch64-redhat-linux/11/../../../../include/c++/11/bits/std_function.h:560
#24 0x0000007ff7684d10 in clover::event::trigger (this=0x1f7b7a0) at ../src/gallium/frontends/clover/core/event.cpp:54
#25 0x0000007ff7685790 in clover::hard_event::hard_event(clover::command_queue&, unsigned int, clover::ref_vector<clover::event> const&, std::function<void (clover::event&)>) (this=0x1f7b7a0, q=..., command=4592, deps=..., action=...) at ../src/gallium/frontends/clover/core/event.cpp:138
#26 0x0000007ff76148f0 in clover::create<clover::hard_event, clover::command_queue&, int, clover::ref_vector<clover::event>&, clEnqueueNDRangeKernel::$_1> (as=..., as=..., as=..., as=...) at ../src/gallium/frontends/clover/util/pointer.hpp:241
#27 0x0000007ff7613d28 in clEnqueueNDRangeKernel (d_q=0x1a356d8, d_kern=0x1ec3d58, dims=1, d_grid_offset=0x0, d_grid_size=0x7fffffe420, d_block_size=0x7fffffe428, num_deps=0, d_deps=0x0, rd_ev=0x0) at ../src/gallium/frontends/clover/api/kernel.cpp:330
#28 0x0000000000479234 in test_local_kernel_scope (device=0x1a48458, context=0x1a59528, queue=0x1a356d8, num_elements=16384) at /home/robclark/src/OpenCL-CTS/test_conformance/basic/test_local_kernel_scope.cpp:105
#29 0x00000000004abbd8 in callSingleTestFunction (test=..., deviceToUse=0x1a48458, forceNoContextCreation=0, numElementsToUse=16384, queueProps=0) at /home/robclark/src/OpenCL-CTS/test_common/harness/testHarness.cpp:838
#30 0x00000000004ab8b8 in callTestFunctions (testList=0x543c30 <test_list>, selectedTestList=0x1a592c0 "", resultTestList=0x1a59340, testNum=117, deviceToUse=0x1a48458, forceNoContextCreation=0, numElementsToUse=16384, queueProps=0) at /home/robclark/src/OpenCL-CTS/test_common/harness/testHarness.cpp:749
#31 0x00000000004ab718 in parseAndCallCommandLineTests (argc=2, argv=0x7fffffed78, device=0x1a48458, testNum=117, testList=0x543c30 <test_list>, forceNoContextCreation=0, queueProps=0, num_elements=16384) at /home/robclark/src/OpenCL-CTS/test_common/harness/testHarness.cpp:703
#32 0x00000000004aafd8 in runTestHarnessWithCheck (argc=2, argv=0x7fffffed78, testNum=117, testList=0x543c30 <test_list>, forceNoContextCreation=0, queueProps=0, deviceCheckFn=0x0) at /home/robclark/src/OpenCL-CTS/test_common/harness/testHarness.cpp:518
#33 0x00000000004a9f10 in runTestHarness (argc=2, argv=0x7fffffed78, testNum=117, testList=0x543c30 <test_list>, forceNoContextCreation=0, queueProps=0) at /home/robclark/src/OpenCL-CTS/test_common/harness/testHarness.cpp:67
#34 0x0000000000405b98 in main (argc=2, argv=0x7fffffed78) at /home/robclark/src/OpenCL-CTS/test_conformance/basic/main.cpp:170
(gdb)