HomePhabricator

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

Description

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

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.

Patch by: krisb

Reviewers: yaxunl, Anastasia, b-sumner

Reviewed By: yaxunl, Anastasia

Subscribers: nhaehnle, yaxunl, Anastasia, cfe-commits

Differential Revision: https://reviews.llvm.org/D43570

Details

Committed
AlexeySotkinFeb 22 2018, 3:54 AM
Reviewer
yaxunl
Differential Revision
D43570: [OpenCL] Add '-cl-uniform-work-group-size' compile option
Branches
Unknown
Tags
Unknown