Summary:
This is the first part of work announced in
"[RFC] Adding lifetime analysis to clang" [0],
i.e. the addition of the [[gsl::Owner(T)]] and
[[gsl::Pointer(T)]] attributes, which
will enable user-defined types to participate in
the lifetime analysis (which will be part of the
next PR).
The type `T` here is called "DerefType" in the paper,
and denotes the type that an Owner owns and a Pointer
points to. E.g. `std::vector<int>` should be annotated
with `[[gsl::Owner(int)]]` and
a `std::vector<int>::iterator` with `[[gsl::Pointer(int)]]`.
[0] http://lists.llvm.org/pipermail/cfe-dev/2018-November/060355.html
Reviewers: gribozavr
Subscribers: xazax.hun, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D63954
llvm-svn: 367040
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
Defining CLK_NULL_EVENT with a `(void*)` cast has the (unintended?)
side-effect that the address space will be fixed (as generic in OpenCL
2.0 mode). The consequence is that any target specific address space
for the clk_event_t type will not be applied.
It is not clear why the void pointer cast was needed in the first
place, and it seems we can do without it.
Differential Revision: https://reviews.llvm.org/D63876
llvm-svn: 366546
Using the -fdeclare-opencl-builtins option will require a way to
predefine types and macros such as `int4`, `CLK_GLOBAL_MEM_FENCE`,
etc. Move these out of opencl-c.h into opencl-c-base.h such that the
latter can be shared by -fdeclare-opencl-builtins and
-finclude-default-header.
This changes the behaviour of -finclude-default-header when
-fdeclare-opencl-builtins is specified: instead of including the full
header, it will include the header with only the base definitions.
Differential revision: https://reviews.llvm.org/D63256
llvm-svn: 363794
Summary:
I've found that most often the proper way to fix this warning is to add
`static`, because if the code otherwise compiles and links, the function
or variable is apparently not needed outside of the TU.
We can't provide a fix-it hint for variable declarations, because
multiple VarDecls can share the same type, and if we put static in front
of that, we affect all declared variables, some of which might have
previous declarations.
We also provide no fix-it hint for the rare case of an `extern` function
definition, because that would require removing `extern` and I have no
idea how to get the source location of the storage class specifier from
a FunctionDecl. I believe this information is only available earlier in
the AST construction from DeclSpec::getStorageClassSpecLoc(), but we
don't have that here.
Reviewed By: aaron.ballman
Differential Revision: https://reviews.llvm.org/D59402
llvm-svn: 363749
Summary:
Remove unnecessary definition (otherwise the extension will be defined
where it's not supposed to be defined).
Consider the code:
#pragma OPENCL EXTENSION cl_intel_planar_yuv : begin
// some declarations
#pragma OPENCL EXTENSION cl_intel_planar_yuv : end
is enough for extension to become known for clang.
Patch by: Dmitry Sidorov <dmitry.sidorov@intel.com>
Reviewers: Anastasia, yaxunl
Reviewed By: Anastasia
Tags: #clang
Differential Revision: https://reviews.llvm.org/D58666
llvm-svn: 362398
This patch adds a `-fdeclare-opencl-builtins` command line option to
the clang frontend. This enables clang to verify OpenCL C builtin
function declarations using a fast StringMatcher lookup, instead of
including the opencl-c.h file with the `-finclude-default-header`
option. This avoids the large parse time penalty of the header file.
This commit only adds the basic infrastructure and some of the OpenCL
builtins. It does not cover all builtins defined by the various OpenCL
specifications. As such, it is not a replacement for
`-finclude-default-header` yet.
RFC: http://lists.llvm.org/pipermail/cfe-dev/2018-November/060041.html
Co-authored-by: Pierre Gondois
Co-authored-by: Joey Gouly
Co-authored-by: Sven van Haastregt
Differential Revision: https://reviews.llvm.org/D60763
llvm-svn: 362371
OpenCL spec v2.0 s6.13.14:
Samplers can also be declared as global constants in the program
source using the following syntax.
const sampler_t <sampler name> = <value>
This works fine for OpenCL 1.2 but fails for 2.0, because clang duduces
address space of file-scope const sampler variable to be in global address
space whereas spec v2.0 s6.9.b forbids file-scope sampler variable to be
in global address space.
The fix is not to deduce address space for file-scope sampler variables.
Differential Revision: https://reviews.llvm.org/D62197
llvm-svn: 361757
Support queue_t and clk_event_t comparisons in C++ for OpenCL mode, to
preserve backwards compatibility with OpenCL C.
Differential Revision: https://reviews.llvm.org/D62208
llvm-svn: 361467
Support the OpenCL C pipe feature in C++ for OpenCL mode, to preserve
backwards compatibility with OpenCL C.
Various changes had to be made in Parse and Sema to enable
pipe-specific diagnostics, so enable a SemaOpenCL test for C++.
Differential Revision: https://reviews.llvm.org/D62181
llvm-svn: 361382
The semantics for converting nested pointers between address
spaces are not very well defined. Some conversions which do not
really carry any meaning only produce warnings, and in some cases
warnings hide invalid conversions, such as 'global int*' to
'local float*'!
This patch changes the logic in checkPointerTypesForAssignment
and checkAddressSpaceCast to fail properly on implicit conversions
that should definitely not be permitted. We also dig deeper into the
pointer types and warn on explicit conversions where the address
space in a nested pointer changes, regardless of whether the address
space is compatible with the corresponding pointer nesting level
on the destination type.
Fixes PR39674!
Patch by ebevhan (Bevin Hansson)!
Differential Revision: https://reviews.llvm.org/D58236
llvm-svn: 360258
Improved some checks and moved testing of the default header
in C++ mode into the Headers folder.
Differential Revision: https://reviews.llvm.org/D59486
llvm-svn: 356450
As for OpenCL C, we need to allow using printf and toolchain variadic
functions (prefixed by "__") in C++ mode.
Differential Revision: https://reviews.llvm.org/D59219
llvm-svn: 355915
This change adds a new diagnostic for mismatching address spaces
to be used for C++ casts (only enabled in C style cast for now,
the rest will follow!).
The change extends C-style cast rules to account for address spaces.
It also adds a separate function for address space cast checking that
can be used to map from a separate address space cast operator
addrspace_cast (to be added as a follow up patch).
Note, that after this change clang will no longer allows arbitrary
address space conversions in reinterpret_casts because they can lead
to accidental errors. The implicit safe conversions would still be
allowed.
Differential Revision: https://reviews.llvm.org/D58346
llvm-svn: 355609
Valid OpenCL C code should still compile in C++ mode.
This change enables extensions and OpenCL types.
Differential Revision: https://reviews.llvm.org/D57824
llvm-svn: 353431
Re-enable format string warnings on printf.
The warnings are still incomplete. Apparently it is undefined to use a
vector specifier without a length modifier, which is not currently
warned on. Additionally, type warnings appear to not be working with
the hh modifier, and aren't warning on all of the special restrictions
from c99 printf.
llvm-svn: 352540
This reverts r348083. This was based on a misreading of the spec
for printf specifiers.
Also revert r343653, as without a subsequent patch, a correctly
specified format for a vector will incorrectly warn.
Fixes bug 40491.
llvm-svn: 352539
The vector modifier is considered separate, so
don't treat it as a conversion specifier.
This is still not warning on some cases, like
using a type that isn't a valid vector element.
Fixes bug 39652
llvm-svn: 348084
This avoids spurious warnings, but could use
a lot of work. For example the number of vector
elements is not verified, and the passed
value type is not checked.
Fixes bug 39486
llvm-svn: 346806
This patch breaks Index/opencl-types.cl LIT test:
Script:
--
: 'RUN: at line 1'; stage1/bin/c-index-test -test-print-type llvm/tools/clang/test/Index/opencl-types.cl -cl-std=CL2.0 | stage1/bin/FileCheck llvm/tools/clang/test/Index/opencl-types.cl
--
Command Output (stderr):
--
llvm/tools/clang/test/Index/opencl-types.cl:3:26: warning: unsupported OpenCL extension 'cl_khr_fp16' - ignoring [-Wignored-pragmas]
llvm/tools/clang/test/Index/opencl-types.cl:4:26: warning: unsupported OpenCL extension 'cl_khr_fp64' - ignoring [-Wignored-pragmas]
llvm/tools/clang/test/Index/opencl-types.cl:8:9: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
llvm/tools/clang/test/Index/opencl-types.cl:11:8: error: declaring variable of type 'half' is not allowed
llvm/tools/clang/test/Index/opencl-types.cl:15:3: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
llvm/tools/clang/test/Index/opencl-types.cl:16:3: error: use of type 'double4' (vector of 4 'double' values) requires cl_khr_fp64 extension to be enabled
llvm/tools/clang/test/Index/opencl-types.cl:26:26: warning: unsupported OpenCL extension 'cl_khr_gl_msaa_sharing' - ignoring [-Wignored-pragmas]
llvm/tools/clang/test/Index/opencl-types.cl:35:44: error: use of type '__read_only image2d_msaa_t' requires cl_khr_gl_msaa_sharing extension to be enabled
llvm/tools/clang/test/Index/opencl-types.cl:36:49: error: use of type '__read_only image2d_array_msaa_t' requires cl_khr_gl_msaa_sharing extension to be enabled
llvm/tools/clang/test/Index/opencl-types.cl:37:49: error: use of type '__read_only image2d_msaa_depth_t' requires cl_khr_gl_msaa_sharing extension to be enabled
llvm/tools/clang/test/Index/opencl-types.cl:38:54: error: use of type '__read_only image2d_array_msaa_depth_t' requires cl_khr_gl_msaa_sharing extension to be enabled
llvm-svn: 346338
Summary:
I wonder if there are some extension which need to be disabled to get
overloadable candidate available.
Reviewers: asavonic, Anastasia
Reviewed By: Anastasia
Subscribers: yaxunl, sidorovd, cfe-commits
Differential Revision: https://reviews.llvm.org/D54152
llvm-svn: 346311
Also rename `invalid-clk-events-cl2.0.cl` to `clk_event_t.cl` and
repurpose it to include both positive and negative clk_event_t tests.
Differential Revision: https://reviews.llvm.org/D53871
llvm-svn: 345825
Summary:
I recently discovered that adding the following code into `opencl-c.h` causes
failure of `test/Headers/opencl-c-header.cl`:
```
#pragma OPENCL EXTENSION cl_my_ext : begin
void cl_my_ext_foobarbaz();
#pragma OPENCL EXTENSIOn cl_my_ext : end
```
Clang crashes at the assertion is `ASTReader::getGlobalSubmoduleID()`:
```
assert(I != M.SubmoduleRemap.end() && "Invalid index into submodule index remap");
```
The root cause of the problem that to deserialize `OPENCL_EXTENSION_DECLS`
section `ASTReader` needs to deserialize a Decl contained in it. In turn,
deserializing a Decl requires information about whether this declaration is
part of a (sub)module, but this information is not read yet because it is
located further in a module file.
Reviewers: Anastasia, yaxunl, JDevlieghere
Reviewed By: Anastasia
Subscribers: sidorovd, cfe-commits, asavonic
Differential Revision: https://reviews.llvm.org/D53200
llvm-svn: 345497
The get_kernel_* functions used in cl20-device-side-enqueue.cl all return
unsigned integers. This patch avoids undesired implicit conversions on the
returned values.
Differential Revision: https://reviews.llvm.org/D52873
llvm-svn: 344778
Emit llvm.amdgcn.update.dpp for both __builtin_amdgcn_mov_dpp and
__builtin_amdgcn_update_dpp. The first argument to
llvm.amdgcn.update.dpp will be undef for __builtin_amdgcn_mov_dpp.
Differential Revision: https://reviews.llvm.org/D52320
llvm-svn: 344665
Summary:
Allowed extension name (that ought to be disabled) printing in the note message.
This diagnostic was proposed here: https://reviews.llvm.org/D51341
Reviewers: Anastasia, yaxunl
Reviewed By: Anastasia
Subscribers: cfe-commits, asavonic, bader
Differential Revision: https://reviews.llvm.org/D52292
llvm-svn: 344246
C style cast in OpenCL C++ was ignoring the address space
conversions from OpenCL C and as a result accepting incorrect
code to compile. This commit adds special function for checking
correctness of address spaces that is shared between C and C++
casts.
llvm-svn: 344148
This reverts r326937 as it broke block argument handling in OpenCL.
See the discussion on https://reviews.llvm.org/D43783 .
The next commit will add a test case that revealed the issue.
llvm-svn: 343582
Add cl_khr_depth_images to extension-version.cl.
Extend to_addr_builtin.cl to additionally test the built-in methods
to_private and to_local, and test assignment with to_global to
incorrect types.
Patch by Alistair Davies.
Differential Revision: https://reviews.llvm.org/D52020
llvm-svn: 343207
This change allows for zero assignment and comparison of queue_t
type variables, and extends null_queue.cl to test this.
Patch by Alistair Davies.
Differential Revision: https://reviews.llvm.org/D51727
llvm-svn: 342968
Discussed on cfe-commits (Week-of-Mon-20180820), this change leads to
the generation of invalid IR for OpenCL without giving an error.
Therefore, the conclusion was to revert.
llvm-svn: 342885
Add a warning if a parameter with a named address space is passed
to a to_addr builtin.
For example:
int i;
to_private(&i); // generate warning as conversion from private to private is redundant.
Patch by Alistair Davies.
Differential Revision: https://reviews.llvm.org/D51411
llvm-svn: 342638
Summary:
Merged the recently added `err_attribute_argument_negative` diagnostic
with existing `err_attribute_requires_positive_integer` diagnostic:
the former allows only strictly positive integer, while the latter
also allows zero.
Reviewers: aaron.ballman
Reviewed By: aaron.ballman
Subscribers: cfe-commits
Differential Revision: https://reviews.llvm.org/D51853
llvm-svn: 342367
Summary:
Negative arguments in kernel attributes are silently bitcast'ed to
unsigned, for example:
__attribute__((reqd_work_group_size(1, -1, 1)))
__kernel void k() {}
is a complete equivalent of:
__attribute__((reqd_work_group_size(1, 4294967294, 1)))
__kernel void k() {}
This is likely an error, so the patch forbids negative arguments in
several OpenCL attributes. Users who really want 4294967294 can still
use it as an unsigned representation.
Reviewers: Anastasia, yaxunl, bader
Reviewed By: Anastasia, yaxunl, bader
Subscribers: bader, cfe-commits
Differential Revision: https://reviews.llvm.org/D50259
llvm-svn: 341539
Summary:
Given the following kernel:
__kernel void foo() {
double d;
double4 dd;
}
and cl_khr_fp64 is disabled, the compilation would fail due to
the presence of 'double d', but when removed, it passes.
The expectation is that extended vector types of unsupported types
will also be unsupported.
The patch adds the check for this scenario.
Patch by: Ofir Cohen
Reviewers: bader, Anastasia, AlexeySotkin, yaxunl
Reviewed By: Anastasia
Subscribers: cfe-commits
Differential Revision: https://reviews.llvm.org/D51296
llvm-svn: 341309
Recommit of r335084 after revert in r335516.
... 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 order, and therefore printed in the wrong order in
-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 effect, especially on
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 attributes' order was changed instead.
This unfortunately causes some 'previous occurrence here' hints to be
textually after the main marker. This typically happens when attributes
are merged, but are incompatible to each other. 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 reverse. 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.
Differential Revision: https://reviews.llvm.org/D48100
llvm-svn: 338800
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
Matthias Gehre