diff --git a/CHANGELOG.md b/CHANGELOG.md index 419b0a7..74e78ab 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -6,6 +6,7 @@ Changelog * Please add an item to this CHANGELOG for any new features or bug fixes when creating a PR. * Fix OpenCL kernel address-space errors on AMD GPUs by placing program-scope constants in `__constant` memory and copying the GCMC target into private memory before calling `distance2` [#45](https://github.com/OpenBioSim/loch/pull/45). +* Let the OpenCL runtime choose the work-group size rather than passing the CUDA block size, which exceeds the limit on AMD GPUs [#47](https://github.com/OpenBioSim/loch/pull/47). [2026.2.0](https://github.com/openbiosim/loch/compare/2026.1.0...2026.2.0) - Sep 2026 -------------------------------------------------------------------------------------- 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: