From b03494e79e87e24756f108c3e34d4ac21d7e9d02 Mon Sep 17 00:00:00 2001 From: Lester Hedges Date: Wed, 16 Sep 2026 19:40:24 +0100 Subject: [PATCH] Backport fix from PR #47. [ci skip] --- src/loch/_kernels.py | 2 +- src/loch/_platforms/_opencl.py | 6 +++++- 2 files changed, 6 insertions(+), 2 deletions(-) diff --git a/src/loch/_kernels.py b/src/loch/_kernels.py index 9f6601b..07fbcb4 100644 --- a/src/loch/_kernels.py +++ b/src/loch/_kernels.py @@ -32,7 +32,7 @@ #define CONSTANT __constant #define LOCAL __local #define GET_GLOBAL_ID(dim) get_global_id(dim) - #define BLOCK_ID_Y get_group_id(1) // OpenCL: work-group ID in dimension 1 + #define BLOCK_ID_Y get_global_id(1) // OpenCL: y block size is 1, and the runtime picks the work-group shape // Map CUDA-style function names to OpenCL names #define sqrtf sqrt #define powf pow diff --git a/src/loch/_platforms/_opencl.py b/src/loch/_platforms/_opencl.py index 29f4e31..704b260 100644 --- a/src/loch/_platforms/_opencl.py +++ b/src/loch/_platforms/_opencl.py @@ -207,7 +207,11 @@ def wrapper(*args, **kwargs): grid = kwargs.get("grid", (1, 1, 1)) global_size = tuple(b * g for b, g in zip(block, grid)) - local_size = block + + # OpenCL caps the work-group size per device (e.g. 256 on AMD), + # below common CUDA block sizes. The kernels use no local memory + # or barriers, so let the runtime pick the partitioning. + local_size = None processed_args = [] for arg in args: