# Large gap between launching cuda kernel

**URL:** <https://numba.discourse.group/t/large-gap-between-launching-cuda-kernel/1404>\
**Category:** Community Support\
**Created:** [June 23, 2022, 3:41pm UTC](https://numba.discourse.group/t/large-gap-between-launching-cuda-kernel/1404 "2022-06-23T15:41:19Z")\
**Posts on this page:** 5\
**Page:** 1

<div class="post-metadata">

**Author:** ![zhenyuwei](https://yyz2.discourse-cdn.com/free1/user_avatar/numba.discourse.group/zhenyuwei/32/631_2.png) [@zhenyuwei](https://numba.discourse.group/u/zhenyuwei)\
**Post date:** [June 23, 2022, 3:41pm UTC](https://numba.discourse.group/t/large-gap-between-launching-cuda-kernel/1404/1 "2022-06-23T15:41:19Z")

</div>

Hello Everyone.  
When I tried to benchmark the performance of my cuda kernel like:

```auto
for _ in range(500):
    my_kernel[block_per_grid, thread_per_block](args)

```

with nsight system, and I found there is always a large gap between the launching of each kernel.

 ![2022-06-23_23-38](https://global.discourse-cdn.com/free1/uploads/numba/original/1X/2f9ec8d7f2e6fe6c333ec11802b7792b01d1abc3.png)

The gap is even longer than the execution time of the kernel, which never happens when I call cupy function multiple times.

My question is, why this gap exists and how can I get rid of it?

Best,  
Zhenyu

---

<div class="post-metadata">

**Author:** ![gmarkall](https://yyz2.discourse-cdn.com/free1/user_avatar/numba.discourse.group/gmarkall/32/42_2.png) [@gmarkall](https://numba.discourse.group/u/gmarkall)\
**Post date:** [June 26, 2022, 4:08pm UTC](https://numba.discourse.group/t/large-gap-between-launching-cuda-kernel/1404/2 "2022-06-26T16:08:40Z")

</div>

Configuring the kernel is expensive in Numba. Things should be much better in the next release due to a couple of PRs that recently went in:

- [CUDA: Remove context query in launch config by gmarkall · Pull Request #8163 · numba/numba · GitHub](https://github.com/numba/numba/pull/8163)
- [lru cache for configure call by TK-21st · Pull Request #7971 · numba/numba · GitHub](https://github.com/numba/numba/pull/7971)

In the meantime, how much does changing your code to:

```python
configured_kernel = my_kernel[block_per_grid, thread_per_block]
for _ in range(500):
    configured_kernel(args)

```

close the gap?

---

<div class="post-metadata">

**Author:** ![zhenyuwei](https://yyz2.discourse-cdn.com/free1/user_avatar/numba.discourse.group/zhenyuwei/32/631_2.png) [@zhenyuwei](https://numba.discourse.group/u/zhenyuwei)\
**Post date:** [June 27, 2022, 4:52am UTC](https://numba.discourse.group/t/large-gap-between-launching-cuda-kernel/1404/3 "2022-06-27T04:52:34Z")

</div>

Hi @gmarkall, Thanks for your reply. I think this makes sense but when I tried to replace the kernel with configured one, the gap is not closed and I check the cuda API by the nsisght, I notice that the

 ![f6e08a60e0214b8fed2af92038a59f8](https://global.discourse-cdn.com/free1/uploads/numba/original/1X/fe7d3df7dbd7614562265fa2a8fdc54c1a4ef23b.jpeg)  
the `cuStreamSynchronize` is called 3 times between the gap. I think the gap exists for this reason. But I don’t understand why synchronizing is called multi-times.

Best,

---

<div class="post-metadata">

**Author:** ![gmarkall](https://yyz2.discourse-cdn.com/free1/user_avatar/numba.discourse.group/gmarkall/32/42_2.png) [@gmarkall](https://numba.discourse.group/u/gmarkall)\
**Post date:** [June 27, 2022, 10:01am UTC](https://numba.discourse.group/t/large-gap-between-launching-cuda-kernel/1404/4 "2022-06-27T10:01:43Z")

</div>

Many thanks for the info!

What are your arguments to the kernel? Can you paste some code that reproduces the issue?

A hunch - are you calling the kernel on CuPy arrays? If so, and you know you don’t need to synchronize on any prior operations in CuPy, you could try setting the environment variable [`NUMBA_CUDA_ARRAY_INTERFACE_SYNC`](https://numba.readthedocs.io/en/stable/reference/envvars.html#envvar-NUMBA_CUDA_ARRAY_INTERFACE_SYNC) to `0` (or as a config variable in Python `numba.config.CUDA_ARRAY_INTERFACE_SYNC = False`) to close the gap. For example, if I run:

```python
from numba import cuda
import cupy as cp

x = cp.arange(10)

@cuda.jit
def kernel(x):
    pass

kernel[1, 1](x)

```

with CUDA debug logging enabled, I see:

```auto
$ NUMBA_CUDA_LOG_LEVEL=DEBUG python repro.py 
== CUDA [434] INFO -- init
== CUDA [434] DEBUG -- call driver api: cuInit
== CUDA [434] DEBUG -- call driver api: cuCtxGetCurrent
== CUDA [434] DEBUG -- call driver api: cuCtxGetDevice
== CUDA [434] DEBUG -- call driver api: cuDeviceGetCount
== CUDA [434] DEBUG -- call driver api: cuDeviceGet
== CUDA [434] DEBUG -- call driver api: cuDeviceGetAttribute
== CUDA [434] DEBUG -- call driver api: cuDeviceGetAttribute
== CUDA [434] DEBUG -- call driver api: cuDeviceGetName
== CUDA [434] DEBUG -- call driver api: cuDeviceGetUuid_v2
== CUDA [434] DEBUG -- call driver api: cuDevicePrimaryCtxRetain
== CUDA [435] DEBUG -- call driver api: cuMemGetInfo_v2
== CUDA [435] DEBUG -- call driver api: cuPointerGetAttribute
== CUDA [435] DEBUG -- call driver api: cuCtxGetCurrent
== CUDA [435] DEBUG -- call driver api: cuCtxGetDevice
== CUDA [435] DEBUG -- call driver api: cuPointerGetAttribute
== CUDA [589] DEBUG -- call driver api: cuCtxGetCurrent
== CUDA [589] DEBUG -- call driver api: cuCtxGetDevice
== CUDA [590] DEBUG -- call runtime api: cudaRuntimeGetVersion
== CUDA [597] DEBUG -- call driver api: cuCtxGetCurrent
== CUDA [597] DEBUG -- call driver api: cuCtxGetDevice
== CUDA [597] DEBUG -- call driver api: cuLinkCreate_v2
== CUDA [597] DEBUG -- call driver api: cuLinkAddData_v2
== CUDA [598] DEBUG -- call driver api: cuLinkComplete
== CUDA [598] DEBUG -- call driver api: cuLinkDestroy
== CUDA [598] DEBUG -- call driver api: cuModuleLoadDataEx
== CUDA [598] DEBUG -- call driver api: cuModuleGetFunction
== CUDA [598] DEBUG -- call driver api: cuFuncGetAttribute
== CUDA [598] DEBUG -- call driver api: cuFuncGetAttribute
== CUDA [598] DEBUG -- call driver api: cuFuncGetAttribute
== CUDA [598] DEBUG -- call driver api: cuFuncGetAttribute
== CUDA [598] DEBUG -- call driver api: cuFuncGetAttribute
== CUDA [598] DEBUG -- call driver api: cuCtxGetCurrent
== CUDA [598] DEBUG -- call driver api: cuCtxGetDevice
== CUDA [598] DEBUG -- call driver api: cuCtxGetCurrent
== CUDA [598] DEBUG -- call driver api: cuCtxGetDevice
== CUDA [599] DEBUG -- call driver api: cuPointerGetAttribute
== CUDA [599] DEBUG -- call driver api: cuStreamSynchronize
== CUDA [599] DEBUG -- call driver api: cuLaunchKernel
== CUDA [599] INFO -- add pending dealloc: module_unload ? bytes

```

and with the sync disabled by environment variable, I see:

```auto
$ NUMBA_CUDA_ARRAY_INTERFACE_SYNC=0 NUMBA_CUDA_LOG_LEVEL=DEBUG python repro.py 
== CUDA [445] INFO -- init
== CUDA [445] DEBUG -- call driver api: cuInit
== CUDA [446] DEBUG -- call driver api: cuCtxGetCurrent
== CUDA [446] DEBUG -- call driver api: cuCtxGetDevice
== CUDA [446] DEBUG -- call driver api: cuDeviceGetCount
== CUDA [446] DEBUG -- call driver api: cuDeviceGet
== CUDA [446] DEBUG -- call driver api: cuDeviceGetAttribute
== CUDA [446] DEBUG -- call driver api: cuDeviceGetAttribute
== CUDA [446] DEBUG -- call driver api: cuDeviceGetName
== CUDA [446] DEBUG -- call driver api: cuDeviceGetUuid_v2
== CUDA [446] DEBUG -- call driver api: cuDevicePrimaryCtxRetain
== CUDA [446] DEBUG -- call driver api: cuMemGetInfo_v2
== CUDA [446] DEBUG -- call driver api: cuPointerGetAttribute
== CUDA [447] DEBUG -- call driver api: cuCtxGetCurrent
== CUDA [447] DEBUG -- call driver api: cuCtxGetDevice
== CUDA [447] DEBUG -- call driver api: cuPointerGetAttribute
== CUDA [600] DEBUG -- call driver api: cuCtxGetCurrent
== CUDA [600] DEBUG -- call driver api: cuCtxGetDevice
== CUDA [601] DEBUG -- call runtime api: cudaRuntimeGetVersion
== CUDA [607] DEBUG -- call driver api: cuCtxGetCurrent
== CUDA [607] DEBUG -- call driver api: cuCtxGetDevice
== CUDA [607] DEBUG -- call driver api: cuLinkCreate_v2
== CUDA [608] DEBUG -- call driver api: cuLinkAddData_v2
== CUDA [608] DEBUG -- call driver api: cuLinkComplete
== CUDA [608] DEBUG -- call driver api: cuLinkDestroy
== CUDA [609] DEBUG -- call driver api: cuModuleLoadDataEx
== CUDA [609] DEBUG -- call driver api: cuModuleGetFunction
== CUDA [609] DEBUG -- call driver api: cuFuncGetAttribute
== CUDA [609] DEBUG -- call driver api: cuFuncGetAttribute
== CUDA [609] DEBUG -- call driver api: cuFuncGetAttribute
== CUDA [609] DEBUG -- call driver api: cuFuncGetAttribute
== CUDA [609] DEBUG -- call driver api: cuFuncGetAttribute
== CUDA [609] DEBUG -- call driver api: cuCtxGetCurrent
== CUDA [609] DEBUG -- call driver api: cuCtxGetDevice
== CUDA [609] DEBUG -- call driver api: cuCtxGetCurrent
== CUDA [609] DEBUG -- call driver api: cuCtxGetDevice
== CUDA [609] DEBUG -- call driver api: cuPointerGetAttribute
== CUDA [609] DEBUG -- call driver api: cuLaunchKernel
== CUDA [609] INFO -- add pending dealloc: module_unload ? bytes

```

the difference being basically:

```diff
--- sync.txt	2022-06-27 10:56:04.357343269 +0100
+++ no-sync.txt	2022-06-27 10:56:13.741458107 +0100
@@ -35,7 +35,6 @@
 == CUDA [] DEBUG -- call driver api: cuCtxGetCurrent
 == CUDA [] DEBUG -- call driver api: cuCtxGetDevice
 == CUDA [] DEBUG -- call driver api: cuPointerGetAttribute
-== CUDA [] DEBUG -- call driver api: cuStreamSynchronize
 == CUDA [] DEBUG -- call driver api: cuLaunchKernel
 == CUDA [] INFO -- add pending dealloc: module_unload ? bytes

```

This synchronization behaviour is enabled by default to ensure correctness by default. This is explained / documented in [this section](https://numba.readthedocs.io/en/stable/cuda/cuda_array_interface.html#synchronization-in-numba) of the [CUDA Array Interface specification](https://numba.readthedocs.io/en/stable/cuda/cuda_array_interface.html). The behaviour of Numba is intended to ensure that it meets the [requirements laid out by the specification](https://numba.readthedocs.io/en/stable/cuda/cuda_array_interface.html#interface-requirements) for general correctness when sharing arrays between libraries that operate asynchronously.

---

<div class="post-metadata">

**Author:** ![zhenyuwei](https://yyz2.discourse-cdn.com/free1/user_avatar/numba.discourse.group/zhenyuwei/32/631_2.png) [@zhenyuwei](https://numba.discourse.group/u/zhenyuwei)\
**Post date:** [June 27, 2022, 1:09pm UTC](https://numba.discourse.group/t/large-gap-between-launching-cuda-kernel/1404/5 "2022-06-27T13:09:42Z")

</div>

Hi @gmarkall, really appreciate the quick reply! For the two optimization strategies you suggest, I set`numba.config.CUDA_ARRAY_INTERFACE_SYNC = False` and do the profiling with or without a configured kernel.

The performance changed from **219ms** (no configured kernel, with `sync=True`) to **158ms** (no configured kernel with `sync=False`) and **147ms** (configured kernel with `sync=True`).  
It seems that the `sync` behavior has a big influence. However, I found the kernel gap disappears for the launching of the first serval loops and then appears again.

 ![2022-06-27_21-07](https://global.discourse-cdn.com/free1/uploads/numba/original/1X/ae274db6cbcc3a1bdcc9df06ea20eda0044d39d5.png)

, while the gap shrinks from **250us** to **150us** , which is still comparable to the kernel execution time. The change of gap time is really strange for me. And there is no other Cuda API calling between the gap.

Here I give the code of the target kernel:

```auto
#!/usr/bin/env python
# -*- encoding: utf-8 -*-
"""
file : repro.py
created time : 2022/06/27
author : Zhenyu Wei
version : 1.0
contact : zhenyuwei99@gmail.com
copyright : (C)Copyright 2021-2021, Zhenyu Wei and Southeast University
"""

import numpy as np
import cupy as cp
import numba as nb
import numba.cuda as cuda
from cupy.cuda.nvtx import RangePush, RangePop

GRID_POINTS_PER_BLOCK = 8 # 8*8*8 grids point will be solved in one block
TOTAL_POINTS_PER_BLOCK = GRID_POINTS_PER_BLOCK + 2 # 10*10*10 neighbors are needed
nb.config.CUDA_ARRAY_INTERFACE_SYNC = False

@cuda.jit(
    nb.void(
        nb.float32[:, :, ::1], # relative_permittivity_map
        nb.float32[:, :, ::1], # coulombic_electric_potential_map
        nb.float32[::1], # cavity_relative_permittivity
        nb.int32[::1], # inner_grid_size
        nb.float32[:, :, ::1], # reaction_field_electric_potential_map
    )
)
def solve_equation(
    relative_permittivity_map,
    coulombic_electric_potential_map,
    cavity_relative_permittivity,
    inner_grid_size,
    reaction_field_electric_potential_map,
):
    # Local index
    local_thread_index = cuda.local.array((3), nb.int32)
    local_thread_index[0] = cuda.threadIdx.x
    local_thread_index[1] = cuda.threadIdx.y
    local_thread_index[2] = cuda.threadIdx.z
    # Global index
    global_thread_index = cuda.local.array((3), nb.int32)
    global_thread_index[0] = local_thread_index[0] + cuda.blockIdx.x * cuda.blockDim.x
    global_thread_index[1] = local_thread_index[1] + cuda.blockIdx.y * cuda.blockDim.y
    global_thread_index[2] = local_thread_index[2] + cuda.blockIdx.z * cuda.blockDim.z
    # Grid size
    grid_size = cuda.local.array((3), nb.int32)
    for i in range(3):
        grid_size[i] = inner_grid_size[i]
        if global_thread_index[i] >= grid_size[i]:
            return
    # Neighbor array index
    local_array_index = cuda.local.array((3), nb.int32)
    for i in range(3):
        local_array_index[i] = local_thread_index[i] + nb.int32(1)
    # Shared array
    shared_relative_permittivity_map = cuda.shared.array(
        (TOTAL_POINTS_PER_BLOCK, TOTAL_POINTS_PER_BLOCK, TOTAL_POINTS_PER_BLOCK),
        nb.float32,
    )
    shared_coulombic_electric_potential_map = cuda.shared.array(
        (TOTAL_POINTS_PER_BLOCK, TOTAL_POINTS_PER_BLOCK, TOTAL_POINTS_PER_BLOCK),
        nb.float32,
    )
    shared_reaction_field_electric_potential_map = cuda.shared.array(
        (TOTAL_POINTS_PER_BLOCK, TOTAL_POINTS_PER_BLOCK, TOTAL_POINTS_PER_BLOCK),
        nb.float32,
    )
    # Load self point data
    shared_relative_permittivity_map[
        local_array_index[0], local_array_index[1], local_array_index[2]
    ] = relative_permittivity_map[
        global_thread_index[0], global_thread_index[1], global_thread_index[2]
    ]
    shared_coulombic_electric_potential_map[
        local_array_index[0], local_array_index[1], local_array_index[2]
    ] = coulombic_electric_potential_map[
        global_thread_index[0], global_thread_index[1], global_thread_index[2]
    ]
    shared_reaction_field_electric_potential_map[
        local_array_index[0], local_array_index[1], local_array_index[2]
    ] = reaction_field_electric_potential_map[
        global_thread_index[0], global_thread_index[1], global_thread_index[2]
    ]
    # Load boundary point
    for i in range(3):
        if local_thread_index[i] == 0:
            tmp_local_array_index = local_array_index[i]
            tmp_global_thread_index = global_thread_index[i]
            local_array_index[i] = 0
            global_thread_index[i] -= 1
            shared_relative_permittivity_map[
                local_array_index[0], local_array_index[1], local_array_index[2]
            ] = relative_permittivity_map[
                global_thread_index[0],
                global_thread_index[1],
                global_thread_index[2],
            ]
            shared_coulombic_electric_potential_map[
                local_array_index[0], local_array_index[1], local_array_index[2]
            ] = coulombic_electric_potential_map[
                global_thread_index[0],
                global_thread_index[1],
                global_thread_index[2],
            ]
            shared_reaction_field_electric_potential_map[
                local_array_index[0], local_array_index[1], local_array_index[2]
            ] = reaction_field_electric_potential_map[
                global_thread_index[0],
                global_thread_index[1],
                global_thread_index[2],
            ]
            local_array_index[i] = tmp_local_array_index
            global_thread_index[i] = tmp_global_thread_index
        elif local_array_index[i] == GRID_POINTS_PER_BLOCK or global_thread_index[
            i
        ] == (grid_size[i] - nb.float32(1)):
            # Equivalent to local_thread_index[i] == GRID_POINTS_PER_BLOCK - 1
            tmp_local_array_index = local_array_index[i]
            tmp_global_thread_index = global_thread_index[i]
            local_array_index[i] += 1
            global_thread_index[i] += 1
            global_thread_index[i] *= nb.int32(global_thread_index[i] != grid_size[i])
            shared_relative_permittivity_map[
                local_array_index[0], local_array_index[1], local_array_index[2]
            ] = relative_permittivity_map[
                global_thread_index[0],
                global_thread_index[1],
                global_thread_index[2],
            ]
            shared_coulombic_electric_potential_map[
                local_array_index[0], local_array_index[1], local_array_index[2]
            ] = coulombic_electric_potential_map[
                global_thread_index[0],
                global_thread_index[1],
                global_thread_index[2],
            ]
            shared_reaction_field_electric_potential_map[
                local_array_index[0], local_array_index[1], local_array_index[2]
            ] = reaction_field_electric_potential_map[
                global_thread_index[0],
                global_thread_index[1],
                global_thread_index[2],
            ]
            local_array_index[i] = tmp_local_array_index
            global_thread_index[i] = tmp_global_thread_index
    cuda.syncthreads()
    # Load constant
    cavity_relative_permittivity = cavity_relative_permittivity[0]
    # Calculate
    new_val = nb.float32(0)
    denominator = nb.float32(0)
    # Self term
    self_relative_permittivity = shared_relative_permittivity_map[
        local_array_index[0],
        local_array_index[1],
        local_array_index[2],
    ]
    self_coulombic_electric_potential = shared_coulombic_electric_potential_map[
        local_array_index[0],
        local_array_index[1],
        local_array_index[2],
    ]
    new_val += (
        nb.float32(6) * cavity_relative_permittivity * self_coulombic_electric_potential
    )
    # Neighbor term
    for i in range(3):
        for j in [-1, 1]:
            local_array_index[i] += j
            relative_permittivity = nb.float32(0.5) * (
                self_relative_permittivity
                + shared_relative_permittivity_map[
                    local_array_index[0],
                    local_array_index[1],
                    local_array_index[2],
                ]
            )
            new_val += (
                relative_permittivity
                * shared_reaction_field_electric_potential_map[
                    local_array_index[0],
                    local_array_index[1],
                    local_array_index[2],
                ]
            )
            new_val += shared_coulombic_electric_potential_map[
                local_array_index[0],
                local_array_index[1],
                local_array_index[2],
            ] * (relative_permittivity - cavity_relative_permittivity)
            denominator += relative_permittivity
            local_array_index[i] -= j
    # Update
    new_val /= denominator
    new_val -= self_coulombic_electric_potential
    old_val = reaction_field_electric_potential_map[
        global_thread_index[0],
        global_thread_index[1],
        global_thread_index[2],
    ]
    cuda.atomic.add(
        reaction_field_electric_potential_map,
        (
            global_thread_index[0],
            global_thread_index[1],
            global_thread_index[2],
        ),
        nb.float32(0.9) * new_val - nb.float32(0.9) * old_val,
    )

if __name__ == " __main__":
    inner_grid_size = [100, 100, 100]
    device_inner_grid_size = cp.array(inner_grid_size)
    device_relative_permittivity_map = cp.ones(inner_grid_size, cp.float32)
    device_coulombic_electric_potential_map = cp.ones(inner_grid_size, cp.float32)
    device_cavity_relative_permittivity = cp.ones([1], cp.float32)
    device_reaction_filed_electric_potential_map = cp.zeros(inner_grid_size, cp.float32)

    thread_per_block = (
        GRID_POINTS_PER_BLOCK,
        GRID_POINTS_PER_BLOCK,
        GRID_POINTS_PER_BLOCK,
    )
    block_per_grid = (
        int(np.ceil(inner_grid_size[0] / GRID_POINTS_PER_BLOCK)),
        int(np.ceil(inner_grid_size[1] / GRID_POINTS_PER_BLOCK)),
        int(np.ceil(inner_grid_size[2] / GRID_POINTS_PER_BLOCK)),
    )
    RangePush("Configure Kernel")
    configured_kernel = solve_equation[block_per_grid, thread_per_block]
    RangePop()
    RangePush("Update")
    for _ in range(500):
        configured_kernel(
            device_relative_permittivity_map,
            device_coulombic_electric_potential_map,
            device_cavity_relative_permittivity,
            device_inner_grid_size,
            device_reaction_filed_electric_potential_map,
        )
    RangePop()

```

Hope it may help!

Best,