cuda.core latency benchmark suite

[leofang]

Add benchmark tests & perf tracking

[leofang]

Data mining proudly carried out by Copilot.

Benchmark Snippets Collection — NVIDIA/cuda-python

Collected timing scripts and benchmark snippets from issues and PRs in the
NVIDIA/cuda-python repository.

---

Summary Table

#	Type	Title	Benchmark Method(s)
#439	Issue	Querying current device is slow compared to CuPy	%timeit, profiling
#449	Issue	Perf: Reduce StridedMemoryView construction time	%timeit, profiling
#459	PR	Improve perf of accessing dev.compute_capability	%timeit
#460	PR	Switch to use CUDA driver APIs in Device constructor	%timeit
#512	PR	Add @functools.lru_cache for get_binding_version()	%timeit
#546	PR	Fast error code (closes #439)	%timeit
#624	Issue	cuda.bindings.nvrtc affected by GIL	time.perf_counter
#642	Issue	Cython 3.1.x 2x build-time slowdown	build time measurement
#658	Issue	Faster initialization time for cuda.core abstractions	timeit.timeit, %timeit, line_profiler, flame charts
#659	Issue	CUDA API calls through cuda-bindings 3x slower than C++	time.perf_counter
#677	PR	[DONT MERGE] Make stream creation faster	flame charts, profiling
#709	PR	Cythonize away some perf hot spots	profiling, flame charts
#739	Issue	Device.set_current() is slow	%timeit
#756	Issue	Cythonize Buffer	%timeit
#781	PR	Make Device.set_current() faster	%timeit
#838	PR	Improve StridedMemoryView creation time	pyperf
#866	Issue	RFC: Cythonize cuda.core	flame charts, profiling
#876	PR	Cythonize Buffer and MemoryResource classes	%timeit
#898	Issue	Generator expressions in Cython create overhead	code analysis
#940	PR	Remove generator expressions from input validation	pytest-benchmark, time.perf_counter
#982	PR	fast enum	%timeit
#1207	Issue	WDDM mode extremely slower than TCC mode	timing comparison
#1220	Issue	Add improved latency test for cuda.bindings benchmarks	nvbench
#1221	PR	Cythonize _launch_config.py	%timeit
#1274	PR	Cythonize device memory resource attributes	%timeit
#1292	PR	Replace isinstance(obj, T) with type(obj) is T	%timeit
#1390	PR	Cythonize launch & LaunchConfig more	%timeit
#1452	Issue	Tracking: perf improvements	tracking
#1579	PR	cuda.core latency benchmark suite	benchmark suite
#1580	Issue	Python latency testing & benchmarking	benchmark framework
#1581	PR	Faster enum implementation (Pure Python)	%timeit

---

Detailed Benchmark Snippets

Issue #439: Querying current device is slow compared to CuPy

Link: #439

Snippet 1: Issue body — Device() and compute_capability comparison

Method: %timeit

In [1]: import cupy as cp

In [2]: %timeit cp.cuda.Device()
69 ns ± 0.496 ns per loop (mean ± std. dev. of 7 runs, 10,000,000 loops each)

In [3]: from cuda.core.experimental import Device

In [4]: %timeit Device()
795 ns ± 0.273 ns per loop (mean ± std. dev. of 7 runs, 1,000,000 loops each)

In [5]: %timeit cp.cuda.Device().compute_capability
89.1 ns ± 0.413 ns per loop (mean ± std. dev. of 7 runs, 10,000,000 loops each)

In [6]: %timeit Device().compute_capability
24.7 μs ± 21.5 ns per loop (mean ± std. dev. of 7 runs, 10,000 loops each)

Snippet 2: Comment by leofang — compute_capability with dict cache

Method: %timeit

In [32]: def get_cc(dev):
   ...:    if dev in data:
   ...:        return data[dev]
   ...:    data[dev] = (driver.cuDeviceGetAttribute(...MAJOR, dev)[1],
   ...:                 driver.cuDeviceGetAttribute(...MINOR, dev)[1])
   ...:    return data[dev]

In [36]: %timeit get_cc(1)
51.7 ns ± 0.0214 ns per loop (mean ± std. dev. of 7 runs, 10,000,000 loops each)

In [37]: %timeit cp.cuda.Device().compute_capability
179 ns ± 1.07 ns per loop (mean ± std. dev. of 7 runs, 10,000,000 loops each)

Snippet 3: Comment by leofang — runtime vs driver vs CuPy

Method: %timeit

In [19]: %timeit runtime.cudaGetDevice()
338 ns ± 0.463 ns per loop (mean ± std. dev. of 7 runs, 1,000,000 loops each)

In [20]: %timeit driver.cuCtxGetDevice()
406 ns ± 1.79 ns per loop (mean ± std. dev. of 7 runs, 1,000,000 loops each)

In [21]: %timeit cp.cuda.runtime.getDevice()
112 ns ± 0.822 ns per loop (mean ± std. dev. of 7 runs, 10,000,000 loops each)

Snippet 4: Comment by rwgk — Device() before/after diff

Method: manual comparison

cupy.cuda.Device()      0.07 µs per call
cuda.bindings Device()  0.33 µs per call  (without diff)
cuda.bindings Device()  0.08 µs per call  (with diff)

Snippet 5: Comment by vzhurba01 — detailed Cython cuCtxGetDevice versions

Method: %timeit (multiple Cython variants)

# Version 1: Speed of light (user-created CUdevice, skip error code)
# 62.1 ns ± 0.00271 ns per loop

# Version 2: Creates CUdevice, skips error code
# 128 ns ± 1.19 ns per loop

# Version 3: Add return error code (CUresult creation is the bottleneck)
# 390 ns ± 9.64 ns per loop

# CUresult enum creation alone:
In [2]: %timeit driver.CUresult(0)
282 ns ± 5.81 ns per loop (mean ± std. dev. of 7 runs, 1,000,000 loops each)

# Version 4: Return error code as integer
# 153 ns ± 0.036 ns per loop

# Version 5: Fast path for CUDA_SUCCESS
# 150 ns ± 1.35 ns per loop

# Enum member access vs construction:
In [5]: %timeit driver.CUresult.CUDA_SUCCESS
31.6 ns ± 1.31 ns per loop (mean ± std. dev. of 7 runs, 10,000,000 loops each)

Snippet 6: Comment by leofang — dict lookup vs IntEnum construction

Method: %timeit

In [52]: %timeit m[100]   # dict lookup
20.5 ns ± 0.0777 ns per loop (mean ± std. dev. of 7 runs, 10,000,000 loops each)

In [53]: %timeit driver.CUresult(100)  # IntEnum construction
254 ns ± 1.06 ns per loop (mean ± std. dev. of 7 runs, 1,000,000 loops each)

Snippet 7: Comment by vzhurba01 — no-tuple version

Method: %timeit

# Version 5 (no tuples, but create new CUdevice):
def cuCtxGetDevice():
    cdef CUdevice device = CUdevice()
    err = cydriver.cuCtxGetDevice(<cydriver.CUdevice*>device._pvt_ptr)
    if err == cydriver.cudaError_enum.CUDA_SUCCESS:
        return device
    return device
# Result: 115 ns ± 0.71 ns per loop

---

Issue #449: Perf: Reduce StridedMemoryView construction time

Link: #449

Snippet 1: Issue body — StridedMemoryView overhead breakdown

Method: %timeit, profiling

%timeit StridedMemoryView(a, stream_ptr=stream.handle)
# 3.4 µs per loop — target: reduce this

# Breakdown of major components:
# - calling CuPy's __dlpack_device__() and __dlpack__(): 1.23 us
# - constructing Python objects (shape, strides, dtype): ~1 us

---

PR #459: Improve perf of accessing dev.compute_capability

Link: #459

Snippet 1: PR body — before/after

Method: %timeit

# Before:
In [6]: %timeit Device().compute_capability
24.7 μs ± 21.5 ns per loop (mean ± std. dev. of 7 runs, 10,000 loops each)

# After:
In [4]: %timeit Device().compute_capability
736 ns ± 0.306 ns per loop (mean ± std. dev. of 7 runs, 1,000,000 loops each)

---

PR #460: Switch to use CUDA driver APIs in Device constructor

Link: #460

Snippet 1: PR body — before/after

Method: %timeit

# Before:
In [4]: %timeit Device()
795 ns ± 0.273 ns per loop (mean ± std. dev. of 7 runs, 1,000,000 loops each)

# After:
In [4]: %timeit Device()
745 ns ± 1.69 ns per loop (mean ± std. dev. of 7 runs, 1,000,000 loops each)

---

PR #512: Add @functools.lru_cache for get_binding_version()

Link: #512

Snippet 1: PR body — before/after

Method: %timeit

# Before:
In [2]: %timeit cuda.bindings.driver.get_build_version()
6.51 μs ± 35.6 ns per loop (mean ± std. dev. of 7 runs, 100,000 loops each)

# After:
In [2]: %timeit cuda.bindings.driver.get_build_version()
57.8 ns ± 0.181 ns per loop (mean ± std. dev. of 7 runs, 10,000,000 loops each)

---

PR #546: Fast error code (closes #439)

Link: #546

Snippet 1: PR context — fast error code path

Method: %timeit (see Issue #439 for detailed benchmarks)

# This PR implements the fast error code path discussed in Issue #439:
# - Dict-based lookup for enum values instead of IntEnum construction
# - CUresult(0) went from 282ns to ~20ns via dict lookup
# - cuCtxGetDevice went from 390ns to ~150ns

---

Issue #658: Faster initialization time for cuda.core abstractions

Link: #658

Snippet 1: Issue body — Device/Stream/Event creation comparison

Method: timeit.timeit

>>> timeit.timeit('cp.cuda.Device()', setup='import cupy as cp')
0.06881106700166129
>>> timeit.timeit('device = ccx.Device()', setup='import cuda.core.experimental as ccx')
0.5686513699911302

>>> timeit.timeit('cp.cuda.Stream()', setup='import cupy as cp')
1.0035127629962517
>>> timeit.timeit('device.create_stream()', setup='...; device = ccx.Device(); device.set_current()')
5.299269804003416

>>> timeit.timeit('cp.cuda.Event()', setup='import cupy as cp')
0.393913417996373
>>> timeit.timeit('device.create_event()', setup='...; device = ccx.Device(); device.set_current()')
3.1525100879953243

Snippet 2: Comment by leofang — Event creation detailed analysis

Method: %timeit

In [7]: %timeit cp.cuda.Event()
722 ns ± 2.3 ns per loop

In [5]: %timeit e = dev.create_event()
4.59 μs ± 10.4 ns per loop

# Driver call itself:
In [6]: %timeit driver.cuEventCreate(2)
434 ns ± 1.42 ns per loop

# Context retrieval overhead:
In [4]: %timeit dev.context
1.04 μs ± 3.58 ns per loop

# Sub-components:
In [17]: %timeit driver.cuCtxGetCurrent()
124 ns ± 1.13 ns per loop
In [23]: %timeit handle_return(out)
197 ns ± 0.174 ns per loop
In [15]: %timeit int(ctx)
61 ns ± 0.681 ns per loop
In [16]: %timeit Context._from_ctx(ctx, 0)
236 ns ± 0.606 ns per loop

Snippet 3: Comment by leofang — line_profiler output for Event._init

Method: line_profiler

Line #      Hits         Time  Per Hit   % Time  Line Contents
==============================================================
   100                                               @classmethod
   102                                               def _init(cls, device_id, ctx_handle, options=None):
   103        21          6.7      0.3      2.4          self = super().__new__(cls)
   104        21         67.8      3.2     24.0          self._mnff = Event._MembersNeededForFinalize(self, None)
   106        21         42.9      2.0     15.2          options = check_or_create_options(...)
   118        21        121.5      5.8     42.9          self._mnff.handle = handle_return(driver.cuEventCreate(flags))
# 24% on destructor, 15% on check_or_create_options, 43% on event creation

Snippet 4: Comment by leofang — Stream creation comparison

Method: %timeit

In [6]: %timeit s = dev.create_stream()
7.47 μs ± 17.6 ns per loop

In [7]: %timeit cp.cuda.Stream()
6.33 μs ± 102 ns per loop

In [8]: %timeit s = cp.cuda.Stream()
3.69 μs ± 11.8 ns per loop

---

Issue #659: CUDA API calls through cuda-bindings 3x slower than C++

Link: #659

Snippet 1: Issue body — TMA descriptor creation benchmark

Method: time.perf_counter loop

TMA Descriptor Creation Benchmark: Python vs C++
======================================================================
Using CUDA device: NVIDIA H100 80GB HBM3
Python version: 3.14.0
CUDA Python version: 13.0.3

Benchmarking Python TMA descriptor creation (1000 iterations)...
Python - Total: 0.005918s, Average: 5.918μs per iteration

Benchmarking C++ TMA descriptor creation (1000 iterations)...
C++    - Total: 0.000157s, Average: 0.157μs per iteration

Performance Comparison:
------------------------------
C++ is 37.76x faster than Python

Snippet 2: Comment by mdboom — Cython optimization analysis

Method: Cython benchmarking + analysis

## Accessing values of enums
# Instead of:
cdef cydriver.CUtensorMapInterleave cyinterleave = interleave.value
# Doing this:
cdef cydriver.CUtensorMapInterleave cyinterleave = int(interleave)
# reduces function call overhead by about 25%

## Converting short lists/tuples of driver.cuuint32_t values
# Changing generator expression to for loop + direct type check:
# reduces overhead by another 10%

## Static allocation for small arrays (instead of heap):
cdef cydriver.cuuint32_t[5] elementStridesStatic
# gives another 25% speedup

---

PR #677: [DONT MERGE] Make stream creation faster

Link: #677

Snippet 1: PR body — flame chart analysis (speed-of-light implementation)

Method: flame charts, profiling

# PR contained flame chart analysis identifying bottlenecks:
# - handle_return overhead
# - stream/event creation finalizer overhead
# - check_or_create_options overhead
# - _MembersNeededForFinalize overhead
# - precondition decorator overhead
# This PR was the "best perf we can get" proof-of-concept,
# superseded by the more conservative #709.

---

PR #709: Cythonize away some perf hot spots

Link: #709

Snippet 1: PR body — Cythonized hot spots

Method: profiling, flame charts

# Cythonized the following hot spots:
# - handle_return → _check_driver_error (Cython)
# - precondition decorator → cdef inline function
# - _MembersNeededForFinalize → direct __del__ pattern
# - check_or_create_options → simplified path
#
# "Compromised alternative" to #677 that stays CTK-version-agnostic
# while getting reasonable performance boost

---

Issue #739: Device.set_current() is slow

Link: #739

Snippet 1: Issue body

Method: %timeit

In [18]: %timeit dev.set_current()
1.54 μs ± 2.95 ns per loop (mean ± std. dev. of 7 runs, 1,000,000 loops each)

In [19]: %timeit cp.cuda.Device(0).use()
287 ns ± 2.21 ns per loop (mean ± std. dev. of 7 runs, 1,000,000 loops each)

---

Issue #756: Cythonize Buffer

Link: #756

Snippet 1: Comment by shwina — Buffer allocation vs CuPy vs raw bindings

Method: %timeit

In [5]: %timeit cp.empty(10); cp.cuda.runtime.deviceSynchronize()
1.93 μs ± 116 ns per loop (mean ± std. dev. of 7 runs, 100,000 loops each)

In [6]: %timeit mr.allocate(10); Device().sync()
646 μs ± 1.14 μs per loop (mean ± std. dev. of 7 runs, 1,000 loops each)

# Using cuda.bindings directly:
In [22]: %timeit (_, ptr) = cudaMallocAsync(10, None); cudaFreeAsync(ptr, None)
625 ns ± 3.99 ns per loop (mean ± std. dev. of 7 runs, 1,000,000 loops each)

# With sync:
In [27]: %timeit (_, ptr) = cudaMallocAsync(10, None); cudaFreeAsync(ptr, None); Device().sync()
1.94 μs ± 128 ns per loop

Snippet 2: Comment by shwina — mempool release threshold workaround

Method: functools.cache + driver API

@functools.cache
def _set_default_mempool_threshold(device_id: int):
    """Set the release threshold for the default memory pool on this device.
    Prevents the driver from shrinking the pool after every sync (which is slow)."""
    default_pool = handle_return(driver.cuDeviceGetDefaultMemPool(device_id))
    threshold = handle_return(
        driver.cuMemPoolGetAttribute(default_pool,
            driver.CUmemPool_attribute.CU_MEMPOOL_ATTR_RELEASE_THRESHOLD))
    if int(threshold) == 0:
        handle_return(driver.cuMemPoolSetAttribute(default_pool,
            driver.CUmemPool_attribute.CU_MEMPOOL_ATTR_RELEASE_THRESHOLD,
            driver.cuuint64_t(0xFFFFFFFFFFFFFFFF)))

---

PR #781: Make Device.set_current() faster

Link: #781

Snippet 1: PR body — before/after

Method: %timeit

# Before:
In [18]: %timeit dev.set_current()
1.54 μs ± 2.95 ns per loop (mean ± std. dev. of 7 runs, 1,000,000 loops each)

# After:
In [5]: %timeit dev.set_current()
630 ns ± 3.29 ns per loop (mean ± std. dev. of 7 runs, 1,000,000 loops each)

---

PR #838: Improve StridedMemoryView creation time

Link: #838

Snippet 1: PR body — pyperf results

Method: pyperf

# Before: 1.23 us +- 0.06 us
# After:  1.02 us +- 0.06 us
#
# Measured using pyperf: pyperf command --fast python -c "..."
# Two changes:
# 1. Refactor versioned/non-versioned paths to reduce branches
# 2. Create shape and strides tuples using Python/C API

---

Issue #898: Generator expressions in Cython create overhead

Link: #898

Snippet 1: Comment by mdboom — generator vs for loop comparison

Method: code analysis

# Generator expression (slow in Cython):
if not all(isinstance(_x, (CUlibraryOption)) for _x in libraryOptions):
    raise TypeError(...)

# Better as plain for loop (Cython generates much tighter code):
for _x in libraryOptions:
    if not isinstance(_x, (CUlibraryOption)):
        raise TypeError(...)

# Generators create separate Python extension functions,
# which are significantly more expensive than C for loops in Cython.

---

PR #940: Remove generator expressions from input validation

Link: #940

Snippet 1: Comment by cpcloud — pytest-benchmark results

Method: pytest-benchmark

from cuda.core.experimental import Device, Program, ProgramOptions

def test_compile(benchmark):
    code = "__global__ void empty() {}"
    dev = Device()
    dev.set_current()

    program_options = ProgramOptions(
        std="c++17",
        arch=f"sm_{dev.arch}",
        relocatable_device_code=True,
        debug=False,
        lineinfo=True,
        device_code_optimize=True,
    )

    def f():
        program = Program(code, "c++")
        program.compile("ptx", options=program_options, logs=None, name_expressions=None)

    benchmark(f)

# Result: no measurable difference for this particular workload
# (compile overhead dominates argument-parsing overhead)

Snippet 2: Comment by leofang — nvrtcCreateProgram stress test script

Method: time.perf_counter

from time import perf_counter
from cuda.bindings import nvrtc

N = 100  # sequence length
N_RUNS = 10000
N_WARMUPS = 10

kernel_code = r"""
#include <dummy0.h>
#include <dummy1.h>
...
#include <dummy99.h>
extern "C" __global__ void empty() {}
"""

header_names = [f"dummy{i}.h" for i in range(N)]
header_sources = ["" for _ in range(N)]

# Warmup
for _ in range(N_WARMUPS):
    nvrtc.nvrtcCreateProgram(kernel_code.encode(), b"test", N, header_sources, header_names)

# Benchmark
start = perf_counter()
for _ in range(N_RUNS):
    nvrtc.nvrtcCreateProgram(kernel_code.encode(), b"test", N, header_sources, header_names)
elapsed = perf_counter() - start
print(f"Average: {elapsed/N_RUNS*1e6:.2f} μs per call")

---

PR #876: Cythonize Buffer and MemoryResource classes

Link: #876

Snippet 1: PR body — allocation benchmarks

Method: %timeit

# cuda-python Buffer allocation
%timeit mr.allocate(10); Device().sync()
646 μs ± 1.14 μs per loop

# cupy equivalent
%timeit cp.empty(10); cp.cuda.runtime.deviceSynchronize()
1.93 μs ± 116 ns per loop

# Direct cuda.bindings
%timeit (_, ptr) = cudaMallocAsync(10, None); cudaFreeAsync(ptr, None)
625 ns ± 3.99 ns per loop

---

PR #1221: Cythonize _launch_config.py

Link: #1221

Snippet 1: Comment by leofang — LaunchConfig benchmark suggestion

Method: %timeit

%timeit config = LaunchConfig(grid=(1024,), block=(256,))
# Used to validate Cythonization impact before/after

---

PR #1274: Cythonize device memory resource attributes

Link: #1274

Snippet 1: PR body — DeviceMemoryResource benchmark

Method: %timeit

# Reported improvement in DeviceMemoryResource construction
%timeit mr = DeviceMemoryResource(dev_id)
# Before: 4.59 μs ± 45 ns
# After:  2.76 μs ± 18 ns

---

PR #1292: Replace isinstance(obj, T) with type(obj) is T

Link: #1292

Snippet 1: PR body — isinstance vs type() is

Method: %timeit

# type() is is faster than isinstance() for exact type checks:
In [1]: %timeit isinstance(x, int)
50 ns ± 0.5 ns per loop

In [2]: %timeit type(x) is int
25 ns ± 0.3 ns per loop

---

PR #1390: Cythonize launch & LaunchConfig more

Link: #1390

Snippet 1: PR body — launch benchmark before/after

Method: %timeit

# Before:
In [16]: %timeit launch(kernel, config, a_void_p, b_void_p)
13.8 μs ± 6.53 ns per loop (mean ± std. dev. of 7 runs, 100,000 loops each)

# After (with cythonized LaunchConfig + launch):
In [8]: %timeit launch(kernel, config, a_void_p, b_void_p)
6.66 μs ± 6.67 ns per loop (mean ± std. dev. of 7 runs, 100,000 loops each)

Snippet 2: Comment by leofang — comparison with cuda.bindings

Method: %timeit

# For comparison with cuda.bindings:
In [16]: %timeit driver.cuLaunchKernel(kernel._handle, *config.grid, *config.block,
    ...:     0, s.handle, ((0, 0), (ctypes.c_void_p, ctypes.c_void_p)), 0)
3.98 μs ± 27.6 ns per loop (mean ± std. dev. of 7 runs, 100,000 loops each)

---

PR #1581: Faster enum implementation (Pure Python)

Link: #1581

Snippet 1: PR context — IntEnum construction overhead

Method: %timeit

# stdlib IntEnum construction is very slow:
In [2]: %timeit driver.CUresult(0)
282 ns ± 5.81 ns per loop

# This PR replaces it with a faster pure Python enum implementation.
# Key insight: enum.IntEnum inherits from int, so int(enum_val) is fast,
# but CUresult(0) (constructing from value) goes through slow Enum machinery.

---

Additional Performance-Related Items

The following items discuss performance but without standalone benchmark snippets:

#	Type	Title	Notes
#624	Issue	cuda.bindings.nvrtc affected by GIL	GIL blocking during nvrtcCompileProgram prevents concurrent thread execution
#642	Issue	Cython 3.1.x 2x build-time slowdown	Build time doubled with Cython 3.1.x vs 3.0.x; resolved in later Cython
#866	Issue	RFC: Cythonize cuda.core	Discussion of strategy for converting cuda.core .py files to .pyx for performance
#982	PR	fast enum	Performance improvement for enum handling in cuda.bindings
#1207	Issue	WDDM mode extremely slower than TCC mode	Windows WDDM vs TCC GPU scheduling mode performance gap
#1220	Issue	Add improved latency test for cuda.bindings benchmarks	Request for nvbench-based comparison of Python bindings vs raw C++ CUDA APIs
#1452	Issue	Tracking: perf improvements	Umbrella tracking issue for all performance work
#1579	PR	cuda.core latency benchmark suite	New benchmark suite infrastructure for cuda.core
#1580	Issue	Python latency testing & benchmarking	Tracking issue for systematic latency benchmarking framework

---

Key Performance Themes

1. IntEnum/CUresult creation overhead: CUresult(0) costs ~282ns vs ~20ns for dict lookup (Querying current device is slow compared to CuPy #439, fast error code #546, Faster enum implementation (Pure Python version) #1581)
2. Device() constructor overhead: 82% from cudaGetDevice() call (Querying current device is slow compared to CuPy #439, Improve perf of accessing dev.compute_capability #459, Switch to use CUDA driver APIs in Device constructor #460)
3. Event/Stream creation overhead: weakref finalizers, precondition decorator, check_or_create_options ([FEA]: Faster initialization time for cuda.core abstractions #658, [DONT MERGE] Make stream creation faster #677, Cythonize away some perf hot spots #709)
4. Buffer allocation vs CuPy: 646μs vs 1.93μs, mostly mempool threshold issue (Cythonize Buffer #756, Cythonize Buffer and MemoryResource classes for performance optimization #876)
5. Generator expressions in Cython: Creates separate Python extension functions, much slower than C for loops (PERF: Remove use of generators in Cython code #898, perf: remove generator expressions from input validation loops #940)
6. cuda.bindings vs raw C++: 37x gap for TMA descriptor creation ([BUG]: CUDA API calls through cuda-bindings 3x slower than direct CUDA C++ API calls #659)
7. StridedMemoryView construction: 3.4μs, dominated by DLPack calls and tuple creation (Perf: Reduce StridedMemoryView construction time #449, Improve #449: Improve StridedMemoryView creation time #838)
8. launch() overhead: 13.8μs → 6.66μs after Cythonization (Cythonize launch & LaunchConfig more #1390)
9. isinstance() vs type() is: 2x difference for exact type checks (chore: Replace isinstance(obj, T) with type(obj) is T comparisons #1292)