Summary:
C++11 onwards specs the non-member functions atomic_load and atomic_load_explicit as taking the atomic<T> by const (potentially volatile) pointer. C11, in its infinite wisdom, decided to drop the const, and C17 will fix this with DR459 (the current draft forgot to fix B.16, but that’s not the normative part).
clang’s lib/Headers/stdatomic.h implements these as #define to the __c11_* equivalent, which are builtins with custom typecheck. Fix the typecheck.
D47613 takes care of the libc++ side.
Discussion: http://lists.llvm.org/pipermail/cfe-dev/2018-May/058129.html
<rdar://problem/27426936>
Reviewers: rsmith
Subscribers: cfe-commits
Differential Revision: https://reviews.llvm.org/D47618
llvm-svn: 338743
The way address space declarations for builtins currently work
is nearly useless. The code assumes the address spaces used for
builtins is a confusingly named "target address space" from user
code using __attribute__((address_space(N))) that matches
the builtin declaration. There's no way to use this to declare
a builtin that returns a language specific address space.
The terminology used is highly cofusing since it has nothing
to do with the the address space selected by the target to use
for a language address space.
This feature is essentially unused as-is. AMDGPU and NVPTX
are the only in-tree targets attempting to use this. The AMDGPU
builtins certainly do not behave as intended (i.e. all of the
builtins returning pointers can never compile because the numbered
address space never matches the expected named address space).
The NVPTX builtins are missing tests for some, and the others
seem to rely on an implicit addrspacecast.
Change the used address space for builtins based on a target
hook to allow using a language address space for a builtin.
This allows the same builtin declaration to be used for multiple
languages with similarly purposed address spaces (e.g. the same
AMDGPU builtin can be used in OpenCL and CUDA even though the
constant address spaces are arbitarily different).
This breaks the possibility of using arbitrary numbered
address spaces alongside the named address spaces for builtins.
If this is an issue we probably need to introduce another builtin
declaration character to distinguish language address spaces from
so-called "target address spaces".
llvm-svn: 338707
Summary:
Size_t, intptr_t, uintptr_t and ptrdiff_t cannot be used as kernel
arguments, according to OpenCL Specification s6.9k:
The size in bytes of these types are implementation-defined and in
addition can also be different for the OpenCL device and the host
processor making it difficult to allocate buffer objects to be passed
as arguments to a kernel declared as pointer to these types.
Patch by: Andrew Savonichev
Reviewers: Anastasia, yaxunl
Subscribers: yaxunl, Anastasia, cfe-commits
Differential Revision: https://reviews.llvm.org/D49725
llvm-svn: 338432
Summary:
OpenCL specification forbids use of several types as kernel arguments.
This patch improves existing diagnostic to look through arrays.
Patch by: Andrew Savonichev
Reviewers: Anastasia, yaxunl
Subscribers: yaxunl, Anastasia, cfe-commits
Differential Revision: https://reviews.llvm.org/D49723
llvm-svn: 338427
Summary:
As discussed in [[ https://bugs.llvm.org/show_bug.cgi?id=38166 | PR38166 ]], we need to be able to distinqush whether the cast
we are visiting is actually a cast, or part of an `ExplicitCast`.
There are at least four ways to get there:
1. Introduce a new `CastKind`, and use it instead of `IntegralCast` if we are in `ExplicitCast`.
Would work, but does not scale - what if we will need more of these cast kinds?
2. Introduce a flag in `CastExprBits`, whether this cast is part of `ExplicitCast` or not.
Would work, but it isn't immediately clear where it needs to be set.
2. Fix `ScalarExprEmitter::VisitCastExpr()` to visit these `NoOp` casts.
As pointed out by @rsmith, CodeGenFunction::EmitMaterializeTemporaryExpr calls
skipRValueSubobjectAdjustments, which steps over the CK_NoOp cast`,
which explains why we currently don't visit those.
This is probably impossible, as @efriedma points out, that is intentional as per `[class.temporary]` in the standard
3. And the simplest one, just record which NoOp casts we skip.
It just kinda works as-is afterwards.
But, the approach with a flag is the least intrusive one, and is probably the best one overall.
Reviewers: rsmith, rjmccall, majnemer, efriedma
Reviewed By: rsmith
Subscribers: cfe-commits, aaron.ballman, vsk, llvm-commits, rsmith
Differential Revision: https://reviews.llvm.org/D49508
llvm-svn: 337815
conversions are only applied to operands of class type, and the second
standard conversion sequence is not applied.
When diagnosing an invalid builtin binary operator, talk about the
original types rather than the converted types. If these differ by a
user-defined conversion, tell the user what happened.
llvm-svn: 335781
Summary:
The comment with the OpenCL clause about this clearly
says: "No type shall be qualified by qualifiers for
two or more different address spaces."
This must mean that two or more qualifiers for the
_same_ address space is allowed. However, it is
likely unintended by the programmer, so emit a
warning.
For dependent address space types, reject them like
before since we cannot know what the address space
will be.
Patch by Bevin Hansson (ebevhan).
Reviewers: Anastasia
Reviewed By: Anastasia
Subscribers: bader, cfe-commits
Differential Revision: https://reviews.llvm.org/D47630
llvm-svn: 335103
... instead of prepending it at the beginning (the original behavior
since implemented in r122535 2010-12-23). This builds up an
AttributeList in the the order in which the attributes appear in the
source.
The reverse order caused nodes for attributes in the AST (e.g. LoopHint)
to be in the reverse, and therefore printed in the wrong order by
-ast-dump. Some TODO comments mention this. The order was explicitly
reversed for enable_if attribute overload resolution and name mangling,
which is not necessary anymore with this patch.
The change unfortunately has some secondary effects, especially for
diagnostic output. In the simplest cases, the CHECK lines or expected
diagnostic were changed to the the new output. If the kind of
error/warning changed, the attribute's order was changed instead.
It also causes some 'previous occurrence here' hints to be textually
after the main marker. This typically happens when attributes are
merged, but are incompatible. Interchanging the role of the the main
and note SourceLocation will also cause the case where two different
declaration's attributes (in contrast to multiple attributes of the
same declaration) are merged to be reversed. There is no easy fix
because sometimes previous attributes are merged into a new
declaration's attribute list, sometimes new attributes are added to a
previous declaration's attribute list. Since 'previous occurrence here'
pointing to locations after the main marker is not rare, I left the
markers as-is; it is only relevant when the attributes are declared in
the same declaration anyway, which often is on the same line.
Differential Revision: https://reviews.llvm.org/D48100
llvm-svn: 335084
There are HIP applications e.g. Tensorflow 1.3 using amdgpu kernel attributes, however
currently they are only allowed on OpenCL kernel functions.
This patch will allow amdgpu kernel attributes to be applied to CUDA/HIP __global__
functions.
Differential Revision: https://reviews.llvm.org/D47958
llvm-svn: 334561
1. added restrictions to memory scope, order and volatile parameters
2. added custom processing for these builtins - currently is not used code,
needed to switch off GCCBuiltin link to the builtins (ongoing change to llvm
tree)
3. builtins renamed as requested
Differential Revision: https://reviews.llvm.org/D43281
llvm-svn: 332848
Added string literal helper function to obtain the type
attributed by a constant address space.
Also fixed predefind __func__ expr to use the helper
to constract the string literal correctly.
Differential Revision: https://reviews.llvm.org/D46049
llvm-svn: 331877
Restrict the following keywords in the OpenCL C++ language mode,
according to Sections 2.2 & 2.9 of the OpenCL C++ 1.0 Specification.
- dynamic_cast
- typeid
- register (already restricted in OpenCL C, update the diagnostic)
- thread_local
- exceptions (try/catch/throw)
- access qualifiers read_only, write_only, read_write
Support the `__global`, `__local`, `__constant`, `__private`, and
`__generic` keywords in OpenCL C++. Leave the unprefixed address
space qualifiers such as global available, i.e., do not mark them as
reserved keywords in OpenCL C++. libclcxx provides explicit address
space pointer classes such as `global_ptr` and `global<T>` that are
implemented using the `__`-prefixed qualifiers.
Differential Revision: https://reviews.llvm.org/D46022
llvm-svn: 331874
Summary:
For OpenCL 1.1 embedded profile 64 bit integers i.e. long,
ulong including the appropriate vector data types and operations
on 64-bit integers are optional. The "cles_khr_int64" extension
string will be reported if the embedded profile implementation
supports 64-bit integers.
Reviewers: Anastasia, bader
Reviewed By: Anastasia, bader
Subscribers: bader, yaxunl, Anastasia, cfe-commits
Differential Revision: https://reviews.llvm.org/D42532
llvm-svn: 323522
Summary:
Constant samplers are handled as static variables and clang's code generation
library, which leads to llvm::unreachable. We bypass emitting sampler variable
as static since it's translated to a function call later.
Reviewers: yaxunl, Anastasia
Reviewed By: yaxunl, Anastasia
Subscribers: cfe-commits
Differential Revision: https://reviews.llvm.org/D34342
llvm-svn: 318290
Changes behavior introduced in r298369 to only error out on
vector component invalid length access on OpenCL mode.
Differential Revision: https://reviews.llvm.org/D38868
rdar://problem/33568748
llvm-svn: 316016
Currently Clang uses default address space (0) to represent private address space for OpenCL
in AST. There are two issues with this:
Multiple address spaces including private address space cannot be diagnosed.
There is no mangling for default address space. For example, if private int* is emitted as
i32 addrspace(5)* in IR. It is supposed to be mangled as PUAS5i but it is mangled as
Pi instead.
This patch attempts to represent OpenCL private address space explicitly in AST. It adds
a new enum LangAS::opencl_private and adds it to the variable types which are implicitly
private:
automatic variables without address space qualifier
function parameter
pointee type without address space qualifier (OpenCL 1.2 and below)
Differential Revision: https://reviews.llvm.org/D35082
llvm-svn: 315668
Treat 'f()' as 'f(void)' rather than a function w/o a prototype.
Reviewers: Anastasia, yaxunl
Reviewed By: Anastasia, yaxunl
Subscribers: cfe-commits, echuraev, chapuni
Differential Revision: https://reviews.llvm.org/D33681
Re-apply revision 306653.
llvm-svn: 315453
Currently clang allows the following code
int a;
int b = (const int) a;
However it does not the following code
int4 a;
int4 b = (const int4) a;
This is because Clang compares the qualified types instead of unqualified types for vector type casting, which causes the inconsistency.
This patch fixes that.
Differential Revision: https://reviews.llvm.org/D38463
llvm-svn: 314802
Block captures can have different physical locations
in memory segments depending on the use case (as a function
call or as a kernel enqueue) and in different vendor
implementations.
Therefore it's unclear how to add address space to capture
addresses uniformly. Currently it has been decided to disallow
taking addresses of captured variables until further
clarifications in the spec.
Differential Revision: https://reviews.llvm.org/D36410
llvm-svn: 312728
OpenCL 2.0 atomic builtin functions have a scope argument which is ideally
represented as synchronization scope argument in LLVM atomic instructions.
Clang supports translating Clang atomic builtin functions to LLVM atomic
instructions. However it currently does not support synchronization scope
of LLVM atomic instructions. Without this, users have to use LLVM assembly
code to implement OpenCL atomic builtin functions.
This patch adds OpenCL 2.0 atomic builtin functions as Clang builtin
functions, which supports generating LLVM atomic instructions with
synchronization scope operand.
Currently only constant memory scope argument is supported. Support of
non-constant memory scope argument will be added later.
Differential Revision: https://reviews.llvm.org/D28691
llvm-svn: 310082
Summary:
does it make sense to enable K&R function declaration style for OpenCL?
clang throws following error message for the declaration w/o arguments:
```
int my_func();
error: function with no prototype cannot use the spir_function calling convention
```
Current way to fix this issue is to specify that parameter list is empty by using 'void':
```
int my_func(void);
```
Let me know what do you think about this patch.
Reviewers: Anastasia, yaxunl
Reviewed By: Anastasia
Subscribers: cfe-commits, echuraev
Differential Revision: https://reviews.llvm.org/D33681
llvm-svn: 306653
Produce an error if variables qualified with a local or
a constant address space are not declared in the outermost
scope of a kernel.
Patch by Simon Perretta.
Differential Revision: https://reviews.llvm.org/D34024
llvm-svn: 305798
Summary:
This is the fix for patch https://reviews.llvm.org/D33353
@uweigand, could you please verify that everything will be good on SystemZ?
I added triple spir-unknown-unknown.
Thank you in advance!
Reviewers: uweigand
Reviewed By: uweigand
Subscribers: yaxunl, cfe-commits, bader, Anastasia, uweigand
Differential Revision: https://reviews.llvm.org/D33648
llvm-svn: 304191
Alloca always returns a pointer in alloca address space, which may
be different from the type defined by the language. For example,
in C++ the auto variables are in the default address space. Therefore
cast alloca to the expected address space when necessary.
Differential Revision: https://reviews.llvm.org/D32248
llvm-svn: 303370
Check unqualified type for ndrange argument in device_side_enqueue so
device_side_enqueue accept const and volatile qualified ndranges.
Differential Revision: https://reviews.llvm.org/D31458
Patch by Dmitry Borisenkov!
llvm-svn: 300988
For OpenCL, the private address space qualifier is 0 in AST. Before this change, 0 address space qualifier
is always mapped to target address space 0. As now target private address space is specified by
alloca address space in data layout, address space qualifier 0 needs to be mapped to alloca addr space specified by the data layout.
This change has no impact on targets whose alloca addr space is 0.
With contributions from Matt Arsenault, Tony Tye and Wen-Heng (Jack) Chung
Differential Revision: https://reviews.llvm.org/D31404
llvm-svn: 299965
Summary:
I saw the same changes in the following review: https://reviews.llvm.org/D17438
I don't know in that way I could determine that atomic variable was initialized by macro ATOMIC_VAR_INIT. Anyway I added check that atomic variables can be initialize only in global scope.
I think that we can discuss this change.
Reviewers: Anastasia, cfe-commits
Reviewed By: Anastasia
Subscribers: bader, yaxunl
Differential Revision: https://reviews.llvm.org/D30643
llvm-svn: 299537
Fixed the assertion due to absence of source location for
implicitly defined types (using addImplicitTypedef()).
During Sema checks the source location is being expected
and therefore an assertion is triggered.
The change is not specific to OpenCL. But it is particularly
common for OpenCL types to be declared implicitly in Clang
to support the mode without the standard header.
Differential Revision: https://reviews.llvm.org/D31397
llvm-svn: 299447
Summary: I added a new rank to ImplicitConversionRank enum to resolve the function overload ambiguity with vector types. Rank of scalar types conversion is lower than vector splat. So, we can choose which function should we call. See test for more details.
Reviewers: Anastasia, cfe-commits
Reviewed By: Anastasia
Subscribers: bader, yaxunl
Differential Revision: https://reviews.llvm.org/D30816
llvm-svn: 298366
1. Reimplemented conditional operator so that it checks
compatibility of unqualified pointees of the 2nd and
the 3rd operands (C99, OpenCL v2.0 6.5.15).
Define QualTypes compatibility for OpenCL as following:
- corresponding types are compatible (C99 6.7.3)
- CVR-qualifiers are equal (C99 6.7.3)
- address spaces are equal (implementation defined)
2. Added generic address space to Itanium mangling.
Review: D30037
Patch by Dmitry Borisenkov!
llvm-svn: 297468
Removed ndrange_t as Clang builtin type and added
as a struct type in the OpenCL header.
Use type name to do the Sema checking in enqueue_kernel
and modify IR generation accordingly.
Review: D28058
Patch by Dmitry Borisenkov!
llvm-svn: 295311
Fix for bug 30217 - incorrect error given for logical
NOT operation with a pointer type: corrected sema check
and improved related tests.
Review: D29038
llvm-svn: 294313
Modify ObjC blocks impl wrt address spaces as follows:
- keep default private address space for blocks generated
as local variables (with captures);
- add global address space for global block literals (no captures);
- make the block invoke function and enqueue_kernel prototype with
the generic AS block pointer parameter to accommodate both
private and global AS cases from above;
- add block handling into default AS because it's implemented as
a special pointer type (BlockPointer) in the frontend and therefore
it is used as a pointer everywhere. This is also needed to accommodate
both private and global AS blocks for the two cases above.
- removes ObjC RT specific symbols (NSConcreteStackBlock and
NSConcreteGlobalBlock) in the OpenCL mode.
Review: https://reviews.llvm.org/D28814
llvm-svn: 293286
Prior to OpenCL 2.0, image3d_t can only be used with the write_only
access qualifier when the cl_khr_3d_image_writes extension is enabled,
see e.g. OpenCL 1.1 s6.8b.
Require the extension for write_only image3d_t types and guard uses of
write_only image3d_t in the OpenCL header.
Patch by Sven van Haastregt!
Review: https://reviews.llvm.org/D28860
llvm-svn: 293050
Summary:
We compile user opencl kernel code with spir triple. But built-ins are written in OpenCL and we compile it with triple x86_64 to be able to use x86 intrinsics. And we need address spaces to match in both cases. So, we change fake address space map in OpenCL for matching with spir.
On CPU address spaces are not really important but we'd like to preserve address space information in order to perform optimizations relying on this info like enhanced alias analysis.
Reviewers: pekka.jaaskelainen, Anastasia
Subscribers: pekka.jaaskelainen, yaxunl, bader, cfe-commits
Differential Revision: https://reviews.llvm.org/D28048
llvm-svn: 290436
This reverts commit r290171. It triggers a bunch of warnings, because
the new enumerator isn't handled in all switches. We want a warning-free
build.
Replied on the commit with more details.
llvm-svn: 290173
Summary: Enabling the compression of CLK_NULL_QUEUE to variable of type queue_t.
Reviewers: Anastasia
Subscribers: cfe-commits, yaxunl, bader
Differential Revision: https://reviews.llvm.org/D27569
llvm-svn: 290171
Added a map to associate types and declarations with extensions.
Refactored existing diagnostic for disabled types associated with extensions and extended it to declarations for generic situation.
Fixed some bugs for types associated with extensions.
Allow users to use pragma to declare types and functions for supported extensions, e.g.
#pragma OPENCL EXTENSION the_new_extension_name : begin
// declare types and functions associated with the extension here
#pragma OPENCL EXTENSION the_new_extension_name : end
Differential Revision: https://reviews.llvm.org/D21698
llvm-svn: 289979
The function SemaBuiltinFPClassification removed superfluous float to double
casts, this was changed to also remove float to float casts but this isn't
valid in all cases, for example when doing an rvaluetolvalue cast. Added a
check to only remove if this was a conventional floating cast.
Added additional tests into SemaOpenCL/extensions to cover these cases
llvm-svn: 289650
This change makes sure single-precision floating point types are used if the
cl_fp64 extension is not supported by the target.
Also removed the check to see whether the OpenCL version is >= 1.2, as this has
been incorporated into the extension setting code.
Differential Revision: https://reviews.llvm.org/D24235
llvm-svn: 289544
Summary: Although the feature was introduced only in OpenCL C v2.0 spec., it's useful for OpenCL 1.x too and doesn't require HW support.
Reviewers: Anastasia
Subscribers: yaxunl, cfe-commits, bader
Differential Revision: https://reviews.llvm.org/D27453
llvm-svn: 289535
This adds the access qualifier to the Pipe Type, rather than using a class
hierarchy.
It also fixes mergeTypes for Pipes, by disallowing merges. Only identical
pipe types can be merged. The test case in invalid-pipes-cl2.0.cl is added
to check that.
llvm-svn: 288332
Make handling integer parameters more flexible:
- For the number of events argument allow to pass larger
integers than 32 bits as soon as compiler can prove that
the range fits in 32 bits. If not, the diagnostic will be given.
- Change type of the arguments specifying the sizes of
the corresponding block arguments to be size_t.
Review: https://reviews.llvm.org/D26509
llvm-svn: 286849
- Accept NULL pointer as a valid parameter value for clk_event.
- Generate clk_event_t arguments of internal
__enqueue_kernel_XXX function as pointers in generic address space.
Review: https://reviews.llvm.org/D26507
llvm-svn: 286836
This change makes sure single-precision floating point types are used if the
cl_fp64 extension is not supported by the target.
Also removed the check to see whether the OpenCL version is >= 1.2, as this has
been incorporated into the extension setting code.
Differential Revision: https://reviews.llvm.org/D24235
llvm-svn: 286815
Certain OpenCL builtin functions are supposed to be executed by all threads in a work group or sub group. Such functions should not be made divergent during transformation. It makes sense to mark them with convergent attribute.
The adding of convergent attribute is based on Ettore Speziale's work and the original proposal and patch can be found at https://www.mail-archive.com/cfe-commits@lists.llvm.org/msg22271.html.
Differential Revision: https://reviews.llvm.org/D25343
llvm-svn: 285725
Summary:
This patch adds a command line option '-cl-ext' to control a set of
supported OpenCL extensions. Option accepts a comma-separated list
of extensions prefixed with '+' or '-'.
It can be used together with a target triple to override support for some
extensions:
// spir target supports all extensions, but we want to disable fp64
clang -cc1 -triple spir-unknown-unknown -cl-ext=-cl_khr_fp64
Special 'all' extension allows to enable or disable all possible
extensions:
// only fp64 will be supported
clang -cc1 -triple spir-unknown-unknown -cl-ext=-all,+cl_khr_fp64
Patch by asavonic (Andrew Savonichev).
Reviewers: joey, yaxunl
Subscribers: yaxunl, bader, Anastasia, cfe-commits
Differential Revision: https://reviews.llvm.org/D23712
llvm-svn: 285700
OpenCL disallows using variadic arguments (s6.9.e and s6.12.5 OpenCL v2.0)
apart from some exceptions:
- printf
- enqueue_kernel
This change adds error diagnostic for variadic functions but accepts printf
and any compiler internal function (which should cover __enqueue_kernel_XXX cases).
It also unifies diagnostic with block prototype and adds missing uncaught cases for blocks.
llvm-svn: 285395
__attribute__((amdgpu_flat_work_group_size(<min>, <max>))) - request minimum and maximum flat work group size
__attribute__((amdgpu_waves_per_eu(<min>[, <max>]))) - request minimum and/or maximum waves per execution unit
Differential Revision: https://reviews.llvm.org/D24513
llvm-svn: 282371
This patch allows us to perform incompatible pointer conversions when
resolving overloads in C. So, the following code will no longer fail to
compile (though it will still emit warnings, assuming the user hasn't
opted out of them):
```
void foo(char *) __attribute__((overloadable));
void foo(int) __attribute__((overloadable));
void callFoo() {
unsigned char bar[128];
foo(bar); // selects the char* overload.
}
```
These conversions are ranked below all others, so:
A. Any other viable conversion will win out
B. If we had another incompatible pointer conversion in the example
above (e.g. `void foo(int *)`), we would complain about
an ambiguity.
Differential Revision: https://reviews.llvm.org/D24113
llvm-svn: 280553
JF Bastien