OpenCL: gate single-group launches by per-kernel work-group limit - #247

Fixes #246.

Short description
On PowerVR BXE-2-32, single-group launches sized using CL_DEVICE_MAX_WORK_GROUP_SIZE (512) exceed the per-kernel limit reported by the driver itself (480 — "the maximum work-group size that can be used to execute the kernel") and fail with CL_OUT_OF_RESOURCES; the ignored launch status turns the failure into silent data corruption.

List of changes
Gate single-group launches by CL_KERNEL_WORK_GROUP_SIZE in src/opencl/cl_reduce_by_key.hpp:cl_reduce_by_key (small path) and src/opencl/cl_sort_by_key.hpp:cl_sort_by_key_bitonic (local, pre-sort and global launches), with helper src/opencl/cl_program_builder.hpp/.cpp:kernel_group_limit; oversized requests fall back to the existing slow paths.

Type of changes

• bug-fix (change which fixes OpenCL: single-group launches exceeding per-kernel work-group limit fail silently on PowerVR #246)
• new-feature (change which proposes new feature or functionality)
• breaking-change (change which breaks compatibility of api or etc.)

Testing strategy (all runs on the PowerVR board; mechanism and numbers detailed in #246)
Standalone reproducer (gist): ./test_reduce_small — CL_OUT_OF_RESOURCES for groups of 512, success up to the reported limit of 480. In-tree sssp: ./sssp --mtxpath=belgium_osm.mtx --niters=1 yields CHECK acc with zero VERIFY lines; same for roadNet-CA. No new in-tree tests: the failure needs PowerVR BXE-2-32 hardware absent from CI.

Background

• clEnqueueNDRangeKernel: https://manpages.debian.org/unstable/opencl-1.2-man-doc/clEnqueueNDRangeKernel.3clc.en.html
• clGetKernelWorkGroupInfo: https://manpages.debian.org/unstable/opencl-1.2-man-doc/clGetKernelWorkGroupInfo.3clc.en.html