[BUG] `KeyError` with `inline="always"` and some interaction with DCE

[gmarkall]

Maybe similar / related to #624.

The following:

import numpy as np
from numba import cuda
from numba.extending import overload

def get_42():
    raise NotImplementedError()

@overload(get_42, target="cuda", inline="always")
def ol_blas_get_accumulator():
    def impl():
        a = cuda.local.array(1, dtype=np.float32)  # Dummy use of cuda to ensure this runs on device
        a[0] = 42.0
        return a[0]

    return impl

@cuda.jit
def kernel(a):
    a[0] = get_42()

def main():
    a = np.empty(1, dtype=np.float32)

    d_a = cuda.to_device(a)
    kernel[1, 1](d_a)
    d_a.copy_to_host(a)

    assert a[0] == 42.0

if __name__ == "__main__":
    main()

results in

Traceback (most recent call last):
  File "/home/gmarkall/numbadev/issues/numba-cuda-624/repro2.py", line 31, in <module>
    main()
  File "/home/gmarkall/numbadev/issues/numba-cuda-624/repro2.py", line 25, in main
    kernel[1, 1](d_a)
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/dispatcher.py", line 693, in __call__
    return self.dispatcher.call(
           ^^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/dispatcher.py", line 1616, in call
    kernel = _dispatcher.Dispatcher._cuda_call(self, *args)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/dispatcher.py", line 1624, in _compile_for_args
    return self.compile(tuple(argtypes))
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/core/compiler_lock.py", line 74, in _acquire_compile_lock
    return func(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/dispatcher.py", line 1880, in compile
    kernel = _Kernel(self.py_func, argtypes, **self.targetoptions)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/core/compiler_lock.py", line 74, in _acquire_compile_lock
    return func(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/dispatcher.py", line 154, in __init__
    cres = compile_cuda(
           ^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/core/compiler_lock.py", line 74, in _acquire_compile_lock
    return func(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/compiler.py", line 745, in compile_cuda
    cres = compile_extra(
           ^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/compiler.py", line 589, in compile_extra
    return pipeline.compile_extra(func)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/core/compiler.py", line 123, in compile_extra
    return self._compile_bytecode()
           ^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/core/compiler.py", line 191, in _compile_bytecode
    return self._compile_core()
           ^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/core/compiler.py", line 167, in _compile_core
    raise e
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/core/compiler.py", line 159, in _compile_core
    pm.run(self.state)
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/core/compiler_machinery.py", line 393, in run
    raise e
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/core/compiler_machinery.py", line 386, in run
    self._runPass(idx, pass_inst, state)
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/core/compiler_lock.py", line 74, in _acquire_compile_lock
    return func(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/core/compiler_machinery.py", line 338, in _runPass
    mutated |= check(pss.run_pass, internal_state)
               ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/core/compiler_machinery.py", line 292, in check
    mangled = func(compiler_state)
              ^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/core/typed_passes.py", line 550, in run_pass
    dead_code_elimination(state.func_ir, typemap=state.typemap)
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/core/ir_utils.py", line 615, in dead_code_elimination
    while remove_dead(
          ^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/core/ir_utils.py", line 665, in remove_dead
    removed |= remove_dead_block(
               ^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/core/ir_utils.py", line 777, in remove_dead_block
    lives.remove(lhs.name)
KeyError: 'ol_blas_get_accumulator__locals__impl_v2_a'

Something not quite right is happening in the pipeline. Using numba.cuda.extending works around the issue, so there is some discrepancy when we interact with IR or parts of the pipeline on the Numba side.

cc @ZzEeKkAa