Fix indentation bug in test_linker.py

`add_from_numba` and `debuggable_kernel` reference
`test_device_functions_ltoir` which is only defined inside
`if TEST_BIN_DIR:`.  Move both definitions under that block
and add `@unittest.skipUnless` to `test_debug_kernel_with_lto`
so the test is skipped when the env var is unset.

Author: leofang

numba_cuda/numba/cuda/tests/cudadrv/test_linker.py

@@ -116,17 +116,15 @@ def simple_lmem(A, B, dty):
         TEST_BIN_DIR, "test_device_functions.ltoir"
     )
 
+    add_from_numba = cuda.declare_device(
+        "add_from_numba",
+        "int32(int32, int32)",
+        link=[test_device_functions_ltoir],
+    )
 
-add_from_numba = cuda.declare_device(
-    "add_from_numba",
-    "int32(int32, int32)",
-    link=[test_device_functions_ltoir],
-)
-
-
-def debuggable_kernel(result):
-    i = cuda.grid(1)
-    result[i] = add_from_numba(i, i)
+    def debuggable_kernel(result):
+        i = cuda.grid(1)
+        result[i] = add_from_numba(i, i)
 
 
 @skip_on_cudasim("Linking unsupported in the simulator")
@@ -350,6 +348,7 @@ def test_get_local_mem_per_specialized(self):
         calc_size = np.dtype(np.float64).itemsize * LMEM_SIZE
         self.assertGreaterEqual(local_mem_size, calc_size)
 
+    @unittest.skipUnless(TEST_BIN_DIR, "NUMBA_CUDA_TEST_BIN_DIR not set")
     def test_debug_kernel_with_lto(self):
         cuda.jit("void(int32[::1])", debug=True, opt=False)(debuggable_kernel)