Repository navigation
OpenCL: gate single-group launches by per-kernel work-group limit - #247
Open
Redkin-Aleksey wants to merge 2 commits into
Open
Redkin-Aleksey wants to merge 2 commits into
Redkin-Aleksey wants to merge 2 commits into
Conversation
…ork-group size limit
…instead of falling back to radix
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Fixes #246.
Short description
On PowerVR B-Series BXE-2-32, single-group launches sized using
CL_DEVICE_MAX_WORK_GROUP_SIZE(512) exceed the per-kernel limit reported by the driver (480, "the maximum work-group size that can be used to execute the kernel") and fail withCL_OUT_OF_RESOURCES; the ignored launch status turns the failure into silent data corruption.List of changes
Added
kernel_group_limit(declaration + definition) to queryCL_KERNEL_WORK_GROUP_SIZEandcapped_group_sizeto compute a work-group size aligned down to the wave size insrc/opencl/cl_program_builder.hpp/.cpp. Gated single-group launches by these limits insrc/opencl/cl_reduce_by_key.hpp:cl_reduce_by_key(small path) andsrc/opencl/cl_sort_by_key.hpp:cl_sort_by_key_bitonic(local, pre-sort and global launches). Bitonic groups are cut into smaller groups when the requested size exceeds the per-kernel limit. Oversized requests fall back to the existing slow paths.Type of changes
Testing strategy (all runs on the PowerVR board; mechanism and numbers detailed in #246)
Standalone reproducer (gist):
./test_reduce_smallreturnsCL_OUT_OF_RESOURCESfor groups of 512, success up to the reported limit of 480. In-treesssp:./sssp --mtxpath=belgium_osm.mtx --niters=1yieldsCHECK accwith zeroVERIFYlines; same forroadNet-CA.amazon-2008(control graph that did not trigger the failure) converges identically. 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.htmlclGetKernelWorkGroupInfo: https://manpages.debian.org/unstable/opencl-1.2-man-doc/clGetKernelWorkGroupInfo.3clc.en.html