Page MenuHomePhabricator

[OpenCL] Add '-cl-uniform-work-group-size' compile option
ClosedPublic

Authored by krisb on Feb 21 2018, 5:48 AM.

Details

Summary

OpenCL 2.0 specification defines '-cl-uniform-work-group-size' option,
which requires that the global work-size be a multiple of the work-group
size specified to clEnqueueNDRangeKernel and allows optimizations that
are made possible by this restriction.

The patch introduces the support of this option.

To keep information about whether an OpenCL kernel has uniform work
group size or not, clang generates 'uniform-work-group-size' function
attribute for every kernel:

  • "uniform-work-group-size"="true" for OpenCL 1.2 and lower,
  • "uniform-work-group-size"="true" for OpenCL 2.0 and higher if '-cl-uniform-work-group-size' option was specified,
  • "uniform-work-group-size"="false" for OpenCL 2.0 and higher if no '-cl-uniform-work-group-size' options was specified.

If the function is not an OpenCL kernel, 'uniform-work-group-size'
attribute isn't generated.

Diff Detail

Repository
rL LLVM

Event Timeline

krisb created this revision.Feb 21 2018, 5:48 AM
Anastasia accepted this revision.Feb 21 2018, 7:56 AM

LGTM! Thanks!

This revision is now accepted and ready to land.Feb 21 2018, 7:56 AM
yaxunl accepted this revision.Feb 21 2018, 11:03 AM

LGTM. Thanks.

krisb updated this revision to Diff 135375.Feb 22 2018, 1:24 AM

Updated one more test where attributes became mismatched.

This revision was automatically updated to reflect the committed changes.