Fix issue #43

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
 --------------------------------------------------------------------------------------

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

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: