Fix indentation bug in test_linker.py (#875)

`add_from_numba` and `debuggable_kernel` reference
`test_device_functions_ltoir` which is only defined inside `if
TEST_BIN_DIR:`. When `NUMBA_CUDA_TEST_BIN_DIR` is not set, this causes a
`NameError` at module load time.

This PR:
1. Moves `add_from_numba` and `debuggable_kernel` inside the `if
TEST_BIN_DIR:` block
2. Adds `@unittest.skipUnless(TEST_BIN_DIR, ...)` to
`test_debug_kernel_with_lto` so it is skipped when the env var is unset

-- Leo's bot

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)