Kernel launch overhead

[99991]

Describe the bug

Launching a kernel via PyOpenCL is 20 times slower than launching the same kernel with the corresponding OpenCL functions via ctypes.

To Reproduce

I discovered this problem while trying to implement bitonic sort. Here is a demonstration which first invokes the kernel with PyOpenCL's call syntax and then switches to calling the same kernel directly via the OpenCL functions loaded from libOpenCL.so using ctypes.

Output with Intel Core i5-2400 CPU:

231.444 ms with PyOpenCL dispatch method
221.348 ms with PyOpenCL dispatch method
221.905 ms with PyOpenCL dispatch method
220.082 ms with PyOpenCL dispatch method
219.206 ms with PyOpenCL dispatch method
223.196 ms with PyOpenCL dispatch method
246.064 ms with PyOpenCL dispatch method
252.709 ms with PyOpenCL dispatch method
255.384 ms with PyOpenCL dispatch method
249.199 ms with PyOpenCL dispatch method
  8.495 ms with ctypes dispatch method
  8.350 ms with ctypes dispatch method
  8.540 ms with ctypes dispatch method
  7.898 ms with ctypes dispatch method
  7.979 ms with ctypes dispatch method
  7.343 ms with ctypes dispatch method
  7.300 ms with ctypes dispatch method
  9.271 ms with ctypes dispatch method
  7.459 ms with ctypes dispatch method
  7.111 ms with ctypes dispatch method

Output with NVIDIA V100 GPU:

 25.175 ms with PyOpenCL dispatch method
 24.560 ms with PyOpenCL dispatch method
 24.169 ms with PyOpenCL dispatch method
 23.987 ms with PyOpenCL dispatch method
 24.396 ms with PyOpenCL dispatch method
 23.612 ms with PyOpenCL dispatch method
 23.937 ms with PyOpenCL dispatch method
 24.184 ms with PyOpenCL dispatch method
 23.541 ms with PyOpenCL dispatch method
 23.887 ms with PyOpenCL dispatch method
  1.951 ms with ctypes dispatch method
  1.872 ms with ctypes dispatch method
  1.876 ms with ctypes dispatch method
  1.885 ms with ctypes dispatch method
  1.879 ms with ctypes dispatch method
  1.872 ms with ctypes dispatch method
  1.914 ms with ctypes dispatch method
  1.962 ms with ctypes dispatch method
  1.869 ms with ctypes dispatch method
  1.863 ms with ctypes dispatch method

Expected behavior

Should be faster. 🚀

Environment (please complete the following information):

Computer 1:

• OS: Debian 12
• ICD Loader and version: 2.3.1-1
• ICD and version: pocl 1.5
• CPU/GPU: Intel Core i5-2400 CPU
• Python version: 3.10.12
• PyOpenCL version: 2024.2.7

Computer 2:

• OS: Debian 12
• ICD Loader and version: 2.3.1-1
• ICD and version: NVIDIA driver 535.247.01
• CPU/GPU: NVIDIA V100 GPU
• Python version: 3.12.3
• PyOpenCL version: 2025.1

Additional context

A similar issue had been created in 2016, but the answer links to a now non-existent mailing list. Unfortunately, the mailing list had not been archived on https://web.archive.org

[inducer]

Repeatedly calling enqueue_nd_range_kernel should be emphatically faster than ctypes, not slower. This is the case simply on account of the fact that nanobind is categorically faster than ctypes. Could you provide a reproducer/a benchmark test?

A few questions ahead of time:

• Do you repeatedly retrieve the kernel? (I.e. prg.sum(...) instead of sum(...))? If so, don't do that---that's very costly.
• Did you call set_scalar_arg_dtypes? This triggers the use of a different, somewhat more efficient invocation path?

[99991]

Oops, I must have deleted the reproducer by accident while formatting the issue. Here it is:

# https://github.com/inducer/pyopencl/issues/830
import pyopencl as cl
import numpy as np
import ctypes
import time
import sys

if sys.platform == "linux":
    dll = ctypes.cdll.LoadLibrary("libOpenCL.so")
elif sys.platform == "win32":
    dll = ctypes.cdll.LoadLibrary("OpenCL.dll")
else:
    raise NotImplementedError(f"{sys.platform} not supported")

clSetKernelArg = dll.clSetKernelArg
clSetKernelArg.argtypes = [
    ctypes.c_void_p,
    ctypes.c_uint,
    ctypes.c_size_t,
    ctypes.c_void_p,
]

clEnqueueNDRangeKernel = dll.clEnqueueNDRangeKernel
clEnqueueNDRangeKernel.argtypes = [
    ctypes.c_void_p,
    ctypes.c_void_p,
    ctypes.c_uint,
    ctypes.POINTER(ctypes.c_size_t),
    ctypes.POINTER(ctypes.c_size_t),
    ctypes.POINTER(ctypes.c_size_t),
    ctypes.c_uint,
    ctypes.c_void_p,
    ctypes.c_void_p,
]

kernel_src = """
__kernel void swap_pairs_kernel(__global int* arr, int x, int n) {
    int i = get_global_id(0);
    int j = i ^ x;

    if (i < j && j < n && arr[j] < arr[i]) {
        int tmp = arr[i];
        arr[i] = arr[j];
        arr[j] = tmp;
    }
}
"""

context = cl.create_some_context()
queue = cl.CommandQueue(context)
program = cl.Program(context, kernel_src).build()
kernel = program.swap_pairs_kernel

def div_ceil(a, b):
    return (a + b - 1) // b

def bitonic_sort(arr, n, dispatch_method):
    if n <= 1: return

    num_shifts = (n - 1).bit_length()

    assert (1 << (num_shifts - 1)) < n <= (1 << num_shifts)

    global_size = (div_ceil(n, 256) * 256,)
    local_size = (256,)

    if dispatch_method == "set_scalar_arg_dtypes":
        # CuPy's default kernel dispatch method seems to be very slow
        program.swap_pairs_kernel.set_scalar_arg_dtypes([None, np.int32, np.int32])

        def swap_pairs(x):
            program.swap_pairs_kernel(queue, global_size, local_size, arr, np.int32(x), np.int32(n))

    elif dispatch_method == "PyOpenCL":
        # CuPy's default kernel dispatch method seems to be very slow
        def swap_pairs(x):
            program.swap_pairs_kernel(queue, global_size, local_size, arr, np.int32(x), np.int32(n))

    elif dispatch_method == "ctypes":
        # Calling the OpenCL API directly is much faster
        def swap_pairs(x):
            arr_ptr = ctypes.c_void_p(arr.int_ptr)
            clSetKernelArg(kernel.int_ptr, 0, 8, ctypes.pointer(arr_ptr))

            x_int32 = ctypes.c_int32(x)
            clSetKernelArg(kernel.int_ptr, 1, 4, ctypes.pointer(x_int32))

            n_int32 = ctypes.c_int32(n)
            clSetKernelArg(kernel.int_ptr, 2, 4, ctypes.pointer(n_int32))

            global_size = ctypes.c_size_t(div_ceil(n, 256) * 256)
            local_size = ctypes.c_size_t(256)
            global_size_ptr = ctypes.pointer(global_size)
            local_size_ptr = ctypes.pointer(local_size)

            clEnqueueNDRangeKernel(queue.int_ptr, kernel.int_ptr, 1, None, global_size_ptr, local_size_ptr, 0, None, None)

    for shift in range(num_shifts):
        k = 1 << (shift + 1)
        swap_pairs(k - 1)

        for shift2 in range(shift, -1, -1):
            swap_pairs(1 << shift2)

def test(dispatch_method):
    for _ in range(10):
        n = 64 * 1024

        arr_np = np.random.randint(0, 2**31 - 1, size=n, dtype=np.int32)

        sorted_arr = np.sort(arr_np)

        arr_ocl = cl.Buffer(context, cl.mem_flags.READ_WRITE | cl.mem_flags.COPY_HOST_PTR, hostbuf=arr_np)

        queue.finish()

        start_time = time.perf_counter()

        bitonic_sort(arr_ocl, n, dispatch_method)
        queue.finish()

        elapsed_time = time.perf_counter() - start_time

        print(f"{elapsed_time * 1000:7.3f} ms with {dispatch_method} dispatch method")

        cl.enqueue_copy(queue, arr_np, arr_ocl).wait()

        assert np.allclose(arr_np, sorted_arr)

    print()

if __name__ == "__main__":
    test(dispatch_method="set_scalar_arg_dtypes")
    test(dispatch_method="PyOpenCL")
    test(dispatch_method="ctypes")

I'll check whether set_scalar_arg_dtypes makes a difference.

EDIT: Unfortunately, set_scalar_arg_dtypes does not make it much faster. I have added it to the reproducer above and the corresponding output is below:

EDIT2: The code above is incorrect, because program.swap_pairs_kernel creates a new kernel every time it is accessed, so calling set_scalar_arg_dtypes on it does nothing.

Output with Intel Core i5-2400 CPU:

231.193 ms with set_scalar_arg_dtypes dispatch method
212.818 ms with set_scalar_arg_dtypes dispatch method
226.454 ms with set_scalar_arg_dtypes dispatch method
224.414 ms with set_scalar_arg_dtypes dispatch method
218.675 ms with set_scalar_arg_dtypes dispatch method
214.710 ms with set_scalar_arg_dtypes dispatch method
238.942 ms with set_scalar_arg_dtypes dispatch method
253.583 ms with set_scalar_arg_dtypes dispatch method
237.087 ms with set_scalar_arg_dtypes dispatch method
244.114 ms with set_scalar_arg_dtypes dispatch method

254.649 ms with PyOpenCL dispatch method
241.925 ms with PyOpenCL dispatch method
238.379 ms with PyOpenCL dispatch method
245.392 ms with PyOpenCL dispatch method
188.989 ms with PyOpenCL dispatch method
258.210 ms with PyOpenCL dispatch method
233.849 ms with PyOpenCL dispatch method
239.077 ms with PyOpenCL dispatch method
229.708 ms with PyOpenCL dispatch method
225.514 ms with PyOpenCL dispatch method

  7.919 ms with ctypes dispatch method
  7.527 ms with ctypes dispatch method
  7.270 ms with ctypes dispatch method
  7.267 ms with ctypes dispatch method
  7.598 ms with ctypes dispatch method
  7.458 ms with ctypes dispatch method
  7.158 ms with ctypes dispatch method
  7.145 ms with ctypes dispatch method
  7.808 ms with ctypes dispatch method
  7.212 ms with ctypes dispatch method

[inducer]

As I suspected above, you're repeatedly retrieving the kernel. With this change, things even out:

--- 830-orig.py 2025-05-27 08:53:01.334050177 -0500
+++ 830.py      2025-05-27 08:52:34.934132773 -0500
@@ -66,15 +66,15 @@

     if dispatch_method == "set_scalar_arg_dtypes":
         # CuPy's default kernel dispatch method seems to be very slow
-        program.swap_pairs_kernel.set_scalar_arg_dtypes([None, np.int32, np.int32])
+        kernel.set_scalar_arg_dtypes([None, np.int32, np.int32])

         def swap_pairs(x):
-            program.swap_pairs_kernel(queue, global_size, local_size, arr, np.int32(x), np.int32(n))
+            kernel(queue, global_size, local_size, arr, np.int32(x), np.int32(n))

     elif dispatch_method == "PyOpenCL":
         # CuPy's default kernel dispatch method seems to be very slow
         def swap_pairs(x):
-            program.swap_pairs_kernel(queue, global_size, local_size, arr, np.int32(x), np.int32(n))
+            kernel(queue, global_size, local_size, arr, np.int32(x), np.int32(n))

     elif dispatch_method == "ctypes":
         # Calling the OpenCL API directly is much faster

This seems to be a sufficiently big gotcha that we should probably warn about it (e.g. if the same kernel gets retrieved more than... 3 times?). See #831.

 93.862 ms with set_scalar_arg_dtypes dispatch method
 13.640 ms with set_scalar_arg_dtypes dispatch method
 11.917 ms with set_scalar_arg_dtypes dispatch method
 13.449 ms with set_scalar_arg_dtypes dispatch method
 13.062 ms with set_scalar_arg_dtypes dispatch method
 13.236 ms with set_scalar_arg_dtypes dispatch method
  9.711 ms with set_scalar_arg_dtypes dispatch method
 10.248 ms with set_scalar_arg_dtypes dispatch method
  7.326 ms with set_scalar_arg_dtypes dispatch method
  7.624 ms with set_scalar_arg_dtypes dispatch method

  8.871 ms with PyOpenCL dispatch method
  9.375 ms with PyOpenCL dispatch method
  8.529 ms with PyOpenCL dispatch method
  8.115 ms with PyOpenCL dispatch method
  8.633 ms with PyOpenCL dispatch method
  6.813 ms with PyOpenCL dispatch method
  7.300 ms with PyOpenCL dispatch method
  7.368 ms with PyOpenCL dispatch method
  7.937 ms with PyOpenCL dispatch method
  9.054 ms with PyOpenCL dispatch method

 10.665 ms with ctypes dispatch method
  6.869 ms with ctypes dispatch method
  5.748 ms with ctypes dispatch method
  6.046 ms with ctypes dispatch method
  6.977 ms with ctypes dispatch method
  6.454 ms with ctypes dispatch method
  5.871 ms with ctypes dispatch method
  6.210 ms with ctypes dispatch method
  7.208 ms with ctypes dispatch method
  8.253 ms with ctypes dispatch method

[inducer]

For context, program.swap_pairs_kernel is clCreateKernel. Plus a bunch of caches in pyopencl are attached to kernel objects, they are rebuilt every time if the kernel object is not reused.

[99991]

Ah, I see. Thank you! I can confirm that moving kernel = program.swap_pairs_kernel out of the loops gives very good performance!

Warning about repeated clCreateKernel calls would indeed have been informative, but I fear that there might be a large number of programs which use PyOpenCL like that and might be annoyed by the warning.

In hindsight, an API with a function call instead of an "invisible" function call might have been better, but it would not have looked as nice.

Alternatively, would it be possible to cache the kernel so only the first call will call clCreateKernel? Or are there people who really want new kernels every time?

[inducer]

I fear that there might be a large number of programs which use PyOpenCL like that and might be annoyed by the warning.

Without the warning, those programs will be subject to increase launch overhead. In the end, I think they might appreciate the warning. And even if not, it's straightforward to disable.

In hindsight, an API with a function call instead of an "invisible" function call might have been better, but it would not have looked as nice.

Such are the trade-offs in API design. 🙂

Alternatively, would it be possible to cache the kernel so only the first call will call clCreateKernel? Or are there people who really want new kernels every time?

Going in this direction now would break compatibility. Consider the case

knl_a = prg.mykernel
knl_b = prg.mykernel

In present-day pyopencl, it is perfectly reasonable to expect that knl_a.set_args() will not affect knl_b. Caching would break that assumption.

In addition, the inclusion of such a cache would require weak references, to avoid making kernels and programs effectively immortal through a reference loop that crosses between Python's and OpenCL's reference counting. And weak references are... really slow in Python, so I think little would be gained in the end.

[99991]

knl_a = prg.mykernel
knl_b = prg.mykernel

In present-day pyopencl, it is perfectly reasonable to expect that knl_a.set_args() will not affect knl_b. Caching would break that assumption.

Earlier today, I was in a very similar situation as in your example and was wondering why my launched kernels did nothing.

program.my_kernel.set_args(...)
queue.enqueue_nd_range_kernel(program.my_kernel, ...)

This was confusing at the time, but now makes perfect sense, knowing that both attempts to access program.my_kernel result in different objects. It appeared like set_args had been "forgotten" when the kernel was launched, but the arguments just went to a different instance of the kernel.

In any case, breaking compatibility is definitely a bad idea, so I suppose it has to stay the way it is.

[inducer]

was wondering why my launched kernels did nothing.

This suggests we should complain at the second kernel retrieval.