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

[Kristina Bessonova via Phabricator via cfe-commits]

krisb created this revision.
Herald added subscribers: cfe-commits, Anastasia, yaxunl.
krisb added reviewers: yaxunl, Anastasia.

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.


Repository:
  rC Clang

https://reviews.llvm.org/D43570

Files:
  include/clang/Driver/Options.td
  include/clang/Frontend/CodeGenOptions.def
  lib/CodeGen/CGCall.cpp
  lib/Driver/ToolChains/Clang.cpp
  lib/Frontend/CompilerInvocation.cpp
  test/CodeGenOpenCL/cl-uniform-wg-size.cl
  test/CodeGenOpenCL/convergent.cl
  test/Driver/opencl.cl

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D43570.135238.patch
Type: text/x-patch
Size: 6869 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20180221/5594d1d8/attachment.bin>