Commit Graph

12 Commits

Author SHA1 Message Date
Matt Arsenault a4451d88ee Consolidate internal denormal flushing controls
Currently there are 4 different mechanisms for controlling denormal
flushing behavior, and about as many equivalent frontend controls.

- AMDGPU uses the fp32-denormals and fp64-f16-denormals subtarget features
- NVPTX uses the nvptx-f32ftz attribute
- ARM directly uses the denormal-fp-math attribute
- Other targets indirectly use denormal-fp-math in one DAGCombine
- cl-denorms-are-zero has a corresponding denorms-are-zero attribute

AMDGPU wants a distinct control for f32 flushing from f16/f64, and as
far as I can tell the same is true for NVPTX (based on the attribute
name).

Work on consolidating these into the denormal-fp-math attribute, and a
new type specific denormal-fp-math-f32 variant. Only ARM seems to
support the two different flush modes, so this is overkill for the
other use cases. Ideally we would error on the unsupported
positive-zero mode on other targets from somewhere.

Move the logic for selecting the flush mode into the compiler driver,
instead of handling it in cc1. denormal-fp-math/denormal-fp-math-f32
are now both cc1 flags, but denormal-fp-math-f32 is not yet exposed as
a user flag.

-cl-denorms-are-zero, -fcuda-flush-denormals-to-zero and
-fno-cuda-flush-denormals-to-zero will be mapped to
-fp-denormal-math-f32=ieee or preserve-sign rather than the old
attributes.

Stop emitting the denorms-are-zero attribute for the OpenCL flag. It
has no in-tree users. The meaning would also be target dependent, such
as the AMDGPU choice to treat this as only meaning allow flushing of
f32 and not f16 or f64. The naming is also potentially confusing,
since DAZ in other contexts refers to instructions implicitly treating
input denormals as zero, not necessarily flushing output denormals to
zero.

This also does not attempt to change the behavior for the current
attribute. The LangRef now states that the default is ieee behavior,
but this is inaccurate for the current implementation. The clang
handling is slightly hacky to avoid touching the existing
denormal-fp-math uses. Fixing this will be left for a future patch.

AMDGPU is still using the subtarget feature to control the denormal
mode, but the new attribute are now emitted. A future change will
switch this and remove the subtarget features.
2020-01-17 20:09:53 -05:00
Anastasia Stulova 88ed70e247 [OpenCL] Rename lang mode flag for C++ mode
Rename lang mode flag to -cl-std=clc++/-cl-std=CLC++
or -std=clc++/-std=CLC++.

This aligns with OpenCL C conversion and removes ambiguity
with OpenCL C++. 

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

llvm-svn: 367008
2019-07-25 11:04:29 +00:00
Anastasia Stulova c645f61ada [OpenCL] Added -std/-cl-std=c++
This is std option for OpenCL C++ v1.0.

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

llvm-svn: 329911
2018-04-12 14:17:04 +00:00
Alexey Sotkin 20f65928e1 [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

llvm-svn: 325771
2018-02-22 11:54:14 +00:00
Yaxun Liu ffb60901fe [OpenCL] Handle -cl-fp32-correctly-rounded-divide-sqrt
Let the driver pass the option to frontend. Do not set precision metadata for division instructions when this option is set. Set function attribute "correctly-rounded-divide-sqrt-fp-math" based on this option.

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

llvm-svn: 278155
2016-08-09 20:10:18 +00:00
Benjamin Kramer 09a9c5537d [OpenCL] In test/Driver/opencl.cl, don't require name of Clang binary to contain "clang"
The test currently fails if the name of the Clang binary doesn't contain "clang".

This patch removes that requirement, as some environments may choose to run the test with a differently named binary. This shouldn't make the test any less strict -- the only place where the flags we're searching for can really occur is the Clang command line.

Patch by Martin Böhme!

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

llvm-svn: 275428
2016-07-14 15:06:57 +00:00
Yaxun Liu 02c3979e22 [OpenCL] Fixes failures in test/Driver/opencl.cl.
Also fixes strict-aliasing option to only be allowed when OpenCL Version 1.0. Added testcase in test/Frontend/opencl-blocks.cl.

Patch by Aaron En Ye Shi.

Differential Revision: http://reviews.llvm.org/D22170

llvm-svn: 275318
2016-07-13 21:21:05 +00:00
Yaxun Liu 79c99fb7eb [OpenCL] Add missing -cl-no-signed-zeros option into driver
Add OCL option -cl-no-signed-zeros to driver options.

Also added to opencl.cl testcases.

Patch by Aaron En Ye Shi.

Differential Revision: http://reviews.llvm.org/D22067

llvm-svn: 274923
2016-07-08 20:28:29 +00:00
Benjamin Kramer 0048964d84 Enable opencl driver tests, which never ran.
Then mark it as XFAIL because it always fails. I'll let OpenCL people
figure this out.

llvm-svn: 274221
2016-06-30 09:17:46 +00:00
Yaxun Liu b5f176e9bd [OpenCL] Allow -cl-std and other standard -cl- options in driver
Allow -cl-std and other standard -cl- options from cc1 to driver.

Added a test for the options moved.

Patch by Aaron En Ye Shi.

Differential Revision: http://reviews.llvm.org/D21031

llvm-svn: 274150
2016-06-29 19:39:32 +00:00
Cong Liu 3688c4ca78 Add -fsyntax-only to Driver/opencl.cl test.
llvm-svn: 273150
2016-06-20 11:25:26 +00:00
Yaxun Liu b5044fe421 [OpenCL] Allow -std={cl|CL}{|1.1|1.2|2.0} in driver
Fix a regression which forbids using -std=cl|CL1.1|CL1.2|CL2.0 in driver.

Allow -std and -cl-std={cl|CL}{|1.1|1.2|2.0}.

Differential Revision: http://reviews.llvm.org/D20630

llvm-svn: 273015
2016-06-17 17:19:28 +00:00