This patch is part of the upstreaming effort for supporting SYCL
language front end.
It makes the following changes:
1. Adds sycl_external attribute for functions with external linkage,
which is intended for use to implement the SYCL_EXTERNAL macro as
specified by the SYCL 2020 specification
2. Adds checks to avoid emitting device code when sycl_external and
sycl_kernel_entry_point attributes are not enabled
3. Fixes test failures caused by the above changes
This patch is missing diagnostics for the following diagnostics listed
in the SYCL 2020 specification's section 5.10.1, which will be addressed
in a subsequent PR:
Functions that are declared using SYCL_EXTERNAL have the following
additional restrictions beyond those imposed on other device functions:
1. If the SYCL backend does not support the generic address space then
the function cannot use raw pointers as parameter or return types.
Explicit pointer classes must be used instead;
2. The function cannot call group::parallel_for_work_item;
3. The function cannot be called from a parallel_for_work_group scope.
In addition to that, the subsequent PR will also implement diagnostics
for inline functions including virtual functions defined as inline.
---------
Co-authored-by: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Previously, the `sycl_kernel_entry_point` attribute could be specified
using either the GNU or C++11 spelling styles. Future SYCL attributes
are expected to support only the C++11 spelling style, so support for
the GNU style is being removed.
In order to ensure consistent presentation of the attribute in
diagnostic messages, diagnostics specific to this attribute now require
the attribute to be provided as an argument. This delegates formatting
of the attribute name to the diagnostic engine.
As an additional nicety, "the" is added to some diagnostic messages so
that they read more like proper sentences.
The MSVC STL includes specializations of `_Is_memfunptr` for every
function pointer type, including every calling convention.
The problem is the AMDGPU target doesn't support the x86 `vectorcall`
calling convention so clang sets it to the default CC. This ends up
clashing with the already-existing overload for the default CC, so we
get a duplicate definition error when including `type_traits` (which we
heavily use in the SYCL STL) and compiling for AMDGPU on Windows.
This doesn't happen for pure AMDGPU non-SYCL because it doesn't include
the C++ STL, and it doesn't happen for CUDA/HIP because a similar
workaround was done
[here](fa49c3a888).
I am not an expert in Sema, so I did a kinda of hardcoded fix, please
let me know if there is a better way to fix this.
As far as I can tell we can't do exactly the same fix that was done for
CUDA because we can't differentiate between device and host code so
easily.
---------
Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
This diagnostic is disabled for device compilation as float128 is not
supported on the device side.
Other diagnostics are already covering the cases where float128 is
actually used in the kernel code, and it's already tested for in the
existing test.
This is expanding on the patch 318bff6 that handled this for cuda
compilation.
A SYCL kernel entry point function is a non-member function or a static
member function declared with the `sycl_kernel_entry_point` attribute.
Such functions define a pattern for an offload kernel entry point
function to be generated to enable execution of a SYCL kernel on a
device. A SYCL library implementation orchestrates the invocation of
these functions with corresponding SYCL kernel arguments in response to
calls to SYCL kernel invocation functions specified by the SYCL 2020
specification.
The offload kernel entry point function (sometimes referred to as the
SYCL kernel caller function) is generated from the SYCL kernel entry
point function by a transformation of the function parameters followed
by a transformation of the function body to replace references to the
original parameters with references to the transformed ones. Exactly how
parameters are transformed will be explained in a future change that
implements non-trivial transformations. For now, it suffices to state
that a given parameter of the SYCL kernel entry point function may be
transformed to multiple parameters of the offload kernel entry point as
needed to satisfy offload kernel argument passing requirements.
Parameters that are decomposed in this way are reconstituted as local
variables in the body of the generated offload kernel entry point
function.
For example, given the following SYCL kernel entry point function
definition:
```
template<typename KernelNameType, typename KernelType>
[[clang::sycl_kernel_entry_point(KernelNameType)]]
void sycl_kernel_entry_point(KernelType kernel) {
kernel();
}
```
and the following call:
```
struct Kernel {
int dm1;
int dm2;
void operator()() const;
};
Kernel k;
sycl_kernel_entry_point<class kernel_name>(k);
```
the corresponding offload kernel entry point function that is generated
might look as follows (assuming `Kernel` is a type that requires
decomposition):
```
void offload_kernel_entry_point_for_kernel_name(int dm1, int dm2) {
Kernel kernel{dm1, dm2};
kernel();
}
```
Other details of the generated offload kernel entry point function, such
as its name and calling convention, are implementation details that need
not be reflected in the AST and may differ across target devices. For
that reason, only the transformation described above is represented in
the AST; other details will be filled in during code generation.
These transformations are represented using new AST nodes introduced
with this change. `OutlinedFunctionDecl` holds a sequence of
`ImplicitParamDecl` nodes and a sequence of statement nodes that
correspond to the transformed parameters and function body.
`SYCLKernelCallStmt` wraps the original function body and associates it
with an `OutlinedFunctionDecl` instance. For the example above, the AST
generated for the `sycl_kernel_entry_point<kernel_name>` specialization
would look as follows:
```
FunctionDecl 'sycl_kernel_entry_point<kernel_name>(Kernel)'
TemplateArgument type 'kernel_name'
TemplateArgument type 'Kernel'
ParmVarDecl kernel 'Kernel'
SYCLKernelCallStmt
CompoundStmt
<original statements>
OutlinedFunctionDecl
ImplicitParamDecl 'dm1' 'int'
ImplicitParamDecl 'dm2' 'int'
CompoundStmt
VarDecl 'kernel' 'Kernel'
<initialization of 'kernel' with 'dm1' and 'dm2'>
<transformed statements with redirected references of 'kernel'>
```
Any ODR-use of the SYCL kernel entry point function will (with future
changes) suffice for the offload kernel entry point to be emitted. An
actual call to the SYCL kernel entry point function will result in a
call to the function. However, evaluation of a `SYCLKernelCallStmt`
statement is a no-op, so such calls will have no effect other than to
trigger emission of the offload kernel entry point.
Additionally, as a related change inspired by code review feedback,
these changes disallow use of the `sycl_kernel_entry_point` attribute
with functions defined with a _function-try-block_. The SYCL 2020
specification prohibits the use of C++ exceptions in device functions.
Even if exceptions were not prohibited, it is unclear what the semantics
would be for an exception that escapes the SYCL kernel entry point
function; the boundary between host and device code could be an implicit
noexcept boundary that results in program termination if violated, or
the exception could perhaps be propagated to host code via the SYCL
library. Pending support for C++ exceptions in device code and clear
semantics for handling them at the host-device boundary, this change
makes use of the `sycl_kernel_entry_point` attribute with a function
defined with a _function-try-block_ an error.
Commit 1a73654b3212b623ac21b9deb3aeaadc6906b7e4 added a missing
diagnostic for incorrect placement of an attribute in a namespace
declaration. This change corrects a SYCL test that inadvertently
exercised the `sycl_kernel_entry_point` attribute in the wrong
declaration location.
The `sycl_kernel_entry_point` attribute is used to declare a function that
defines a pattern for an offload kernel entry point. The attribute requires
a single type argument that specifies a class type that meets the requirements
for a SYCL kernel name as described in section 5.2, "Naming of kernels", of
the SYCL 2020 specification. A unique kernel name type is required for each
function declared with the attribute. The attribute may not first appear on a
declaration that follows a definition of the function. The function is
required to have a non-deduced `void` return type. The function must not be
a non-static member function, be deleted or defaulted, be declared with the
`constexpr` or `consteval` specifiers, be declared with the `[[noreturn]]`
attribute, be a coroutine, or accept variadic arguments.
Diagnostics are not yet provided for the following:
- Use of a type as a kernel name that does not satisfy the forward
declarability requirements specified in section 5.2, "Naming of kernels",
of the SYCL 2020 specification.
- Use of a type as a parameter of the attributed function that does not
satisfy the kernel parameter requirements specified in section 4.12.4,
"Rules for parameter passing to kernels", of the SYCL 2020 specification
(each such function parameter constitutes a kernel parameter).
- Use of language features that are not permitted in device functions as
specified in section 5.4, "Language restrictions for device functions",
of the SYCL 2020 specification.
There are several issues noted by various FIXME comments.
- The diagnostic generated for kernel name conflicts needs additional work
to better detail the relevant source locations; such as the location of
each declaration as well as the original source of each kernel name.
- A number of the tests illustrate spurious errors being produced due to
attributes that appertain to function templates being instantiated too
early (during overload resolution as opposed to after an overload is
selected).
Included changes allow the `SYCLKernelEntryPointAttr` attribute to be
marked as invalid if a `sycl_kernel_entry_point` attribute is used incorrectly.
This is intended to prevent trying to emit an offload kernel entry point
without having to mark the associated function as invalid since doing so
would affect overload resolution; which this attribute should not do.
Unfortunately, Clang eagerly instantiates attributes that appertain to
functions with the result that errors might be issued for function
declarations that are never selected by overload resolution. Tests have
been added to demonstrate this. Further work will be needed to address
these issues (for this and other attributes).
The `sycl_kernel_entry_point` attribute is used to declare a function that
defines a pattern for an offload kernel to be emitted. The attribute requires
a single type argument that specifies the type used as a SYCL kernel name as
described in section 5.2, "Naming of kernels", of the SYCL 2020 specification.
Properties of the offload kernel are collected when a function declared with
the `sycl_kernel_entry_point` attribute is parsed or instantiated. These
properties, such as the kernel name type, are stored in the AST context where
they are (or will be) used for diagnostic purposes and to facilitate reflection
to a SYCL run-time library. These properties are not serialized with the AST
but are recreated upon deserialization.
The `sycl_kernel_entry_point` attribute is intended to replace the existing
`sycl_kernel` attribute which is intended to be deprecated in a future change
and removed following an appropriate deprecation period. The new attribute
differs in that it is enabled for both SYCL host and device compilation, may
be used with non-template functions, explicitly indicates the type used as
the kernel name type, and will impact AST generation.
This change adds the basic infrastructure for the new attribute. Future
changes will add diagnostics and new AST support that will be used to drive
generation of the corresponding offload kernel.
Reapplication of 7339c0f782d5c70e0928f8991b0c05338a90c84c with a fix
for a crash involving arrays without a size expression.
Clang supports VLAs in C++ as an extension, but we currently only warn
on their use when you pass -Wvla, -Wvla-extension, or -pedantic.
However, VLAs as they're expressed in C have been considered by WG21
and rejected, are easy to use accidentally to the surprise of users
(e.g., https://ddanilov.me/default-non-standard-features/), and they
have potential security implications beyond constant-size arrays
(https://wiki.sei.cmu.edu/confluence/display/c/ARR32-C.+Ensure+size+arguments+for+variable+length+arrays+are+in+a+valid+range).
C++ users should strongly consider using other functionality such as
std::vector instead.
This seems like sufficiently compelling evidence to warn users about
VLA use by default in C++ modes. This patch enables the -Wvla-extension
diagnostic group in C++ language modes by default, and adds the warning
group to -Wall in GNU++ language modes. The warning is still opt-in in
C language modes, where support for VLAs is somewhat less surprising to
users.
RFC: https://discourse.llvm.org/t/rfc-diagnosing-use-of-vlas-in-c/73109
Fixes https://github.com/llvm/llvm-project/issues/62836
Differential Revision: https://reviews.llvm.org/D156565
This patch uses existing deferred diagnostics framework to emit error
for unsupported type __bf16 in device code. Error is not emitted in
host code.
Differential Revision: https://reviews.llvm.org/D141375
Without this patch, clang will not wrap in an ElaboratedType node types written
without a keyword and nested name qualifier, which goes against the intent that
we should produce an AST which retains enough details to recover how things are
written.
The lack of this sugar is incompatible with the intent of the type printer
default policy, which is to print types as written, but to fall back and print
them fully qualified when they are desugared.
An ElaboratedTypeLoc without keyword / NNS uses no storage by itself, but still
requires pointer alignment due to pre-existing bug in the TypeLoc buffer
handling.
---
Troubleshooting list to deal with any breakage seen with this patch:
1) The most likely effect one would see by this patch is a change in how
a type is printed. The type printer will, by design and default,
print types as written. There are customization options there, but
not that many, and they mainly apply to how to print a type that we
somehow failed to track how it was written. This patch fixes a
problem where we failed to distinguish between a type
that was written without any elaborated-type qualifiers,
such as a 'struct'/'class' tags and name spacifiers such as 'std::',
and one that has been stripped of any 'metadata' that identifies such,
the so called canonical types.
Example:
```
namespace foo {
struct A {};
A a;
};
```
If one were to print the type of `foo::a`, prior to this patch, this
would result in `foo::A`. This is how the type printer would have,
by default, printed the canonical type of A as well.
As soon as you add any name qualifiers to A, the type printer would
suddenly start accurately printing the type as written. This patch
will make it print it accurately even when written without
qualifiers, so we will just print `A` for the initial example, as
the user did not really write that `foo::` namespace qualifier.
2) This patch could expose a bug in some AST matcher. Matching types
is harder to get right when there is sugar involved. For example,
if you want to match a type against being a pointer to some type A,
then you have to account for getting a type that is sugar for a
pointer to A, or being a pointer to sugar to A, or both! Usually
you would get the second part wrong, and this would work for a
very simple test where you don't use any name qualifiers, but
you would discover is broken when you do. The usual fix is to
either use the matcher which strips sugar, which is annoying
to use as for example if you match an N level pointer, you have
to put N+1 such matchers in there, beginning to end and between
all those levels. But in a lot of cases, if the property you want
to match is present in the canonical type, it's easier and faster
to just match on that... This goes with what is said in 1), if
you want to match against the name of a type, and you want
the name string to be something stable, perhaps matching on
the name of the canonical type is the better choice.
3) This patch could expose a bug in how you get the source range of some
TypeLoc. For some reason, a lot of code is using getLocalSourceRange(),
which only looks at the given TypeLoc node. This patch introduces a new,
and more common TypeLoc node which contains no source locations on itself.
This is not an inovation here, and some other, more rare TypeLoc nodes could
also have this property, but if you use getLocalSourceRange on them, it's not
going to return any valid locations, because it doesn't have any. The right fix
here is to always use getSourceRange() or getBeginLoc/getEndLoc which will dive
into the inner TypeLoc to get the source range if it doesn't find it on the
top level one. You can use getLocalSourceRange if you are really into
micro-optimizations and you have some outside knowledge that the TypeLocs you are
dealing with will always include some source location.
4) Exposed a bug somewhere in the use of the normal clang type class API, where you
have some type, you want to see if that type is some particular kind, you try a
`dyn_cast` such as `dyn_cast<TypedefType>` and that fails because now you have an
ElaboratedType which has a TypeDefType inside of it, which is what you wanted to match.
Again, like 2), this would usually have been tested poorly with some simple tests with
no qualifications, and would have been broken had there been any other kind of type sugar,
be it an ElaboratedType or a TemplateSpecializationType or a SubstTemplateParmType.
The usual fix here is to use `getAs` instead of `dyn_cast`, which will look deeper
into the type. Or use `getAsAdjusted` when dealing with TypeLocs.
For some reason the API is inconsistent there and on TypeLocs getAs behaves like a dyn_cast.
5) It could be a bug in this patch perhaps.
Let me know if you need any help!
Signed-off-by: Matheus Izvekov <mizvekov@gmail.com>
Differential Revision: https://reviews.llvm.org/D112374
This reverts commit 7c51f02effdbd0d5e12bfd26f9c3b2ab5687c93f because it
stills breaks the LLDB tests. This was re-landed without addressing the
issue or even agreement on how to address the issue. More details and
discussion in https://reviews.llvm.org/D112374.
Without this patch, clang will not wrap in an ElaboratedType node types written
without a keyword and nested name qualifier, which goes against the intent that
we should produce an AST which retains enough details to recover how things are
written.
The lack of this sugar is incompatible with the intent of the type printer
default policy, which is to print types as written, but to fall back and print
them fully qualified when they are desugared.
An ElaboratedTypeLoc without keyword / NNS uses no storage by itself, but still
requires pointer alignment due to pre-existing bug in the TypeLoc buffer
handling.
---
Troubleshooting list to deal with any breakage seen with this patch:
1) The most likely effect one would see by this patch is a change in how
a type is printed. The type printer will, by design and default,
print types as written. There are customization options there, but
not that many, and they mainly apply to how to print a type that we
somehow failed to track how it was written. This patch fixes a
problem where we failed to distinguish between a type
that was written without any elaborated-type qualifiers,
such as a 'struct'/'class' tags and name spacifiers such as 'std::',
and one that has been stripped of any 'metadata' that identifies such,
the so called canonical types.
Example:
```
namespace foo {
struct A {};
A a;
};
```
If one were to print the type of `foo::a`, prior to this patch, this
would result in `foo::A`. This is how the type printer would have,
by default, printed the canonical type of A as well.
As soon as you add any name qualifiers to A, the type printer would
suddenly start accurately printing the type as written. This patch
will make it print it accurately even when written without
qualifiers, so we will just print `A` for the initial example, as
the user did not really write that `foo::` namespace qualifier.
2) This patch could expose a bug in some AST matcher. Matching types
is harder to get right when there is sugar involved. For example,
if you want to match a type against being a pointer to some type A,
then you have to account for getting a type that is sugar for a
pointer to A, or being a pointer to sugar to A, or both! Usually
you would get the second part wrong, and this would work for a
very simple test where you don't use any name qualifiers, but
you would discover is broken when you do. The usual fix is to
either use the matcher which strips sugar, which is annoying
to use as for example if you match an N level pointer, you have
to put N+1 such matchers in there, beginning to end and between
all those levels. But in a lot of cases, if the property you want
to match is present in the canonical type, it's easier and faster
to just match on that... This goes with what is said in 1), if
you want to match against the name of a type, and you want
the name string to be something stable, perhaps matching on
the name of the canonical type is the better choice.
3) This patch could exposed a bug in how you get the source range of some
TypeLoc. For some reason, a lot of code is using getLocalSourceRange(),
which only looks at the given TypeLoc node. This patch introduces a new,
and more common TypeLoc node which contains no source locations on itself.
This is not an inovation here, and some other, more rare TypeLoc nodes could
also have this property, but if you use getLocalSourceRange on them, it's not
going to return any valid locations, because it doesn't have any. The right fix
here is to always use getSourceRange() or getBeginLoc/getEndLoc which will dive
into the inner TypeLoc to get the source range if it doesn't find it on the
top level one. You can use getLocalSourceRange if you are really into
micro-optimizations and you have some outside knowledge that the TypeLocs you are
dealing with will always include some source location.
4) Exposed a bug somewhere in the use of the normal clang type class API, where you
have some type, you want to see if that type is some particular kind, you try a
`dyn_cast` such as `dyn_cast<TypedefType>` and that fails because now you have an
ElaboratedType which has a TypeDefType inside of it, which is what you wanted to match.
Again, like 2), this would usually have been tested poorly with some simple tests with
no qualifications, and would have been broken had there been any other kind of type sugar,
be it an ElaboratedType or a TemplateSpecializationType or a SubstTemplateParmType.
The usual fix here is to use `getAs` instead of `dyn_cast`, which will look deeper
into the type. Or use `getAsAdjusted` when dealing with TypeLocs.
For some reason the API is inconsistent there and on TypeLocs getAs behaves like a dyn_cast.
5) It could be a bug in this patch perhaps.
Let me know if you need any help!
Signed-off-by: Matheus Izvekov <mizvekov@gmail.com>
Differential Revision: https://reviews.llvm.org/D112374
This reverts commit bdc6974f92304f4ed542241b9b89ba58ba6b20aa because it
breaks all the LLDB tests that import the std module.
import-std-module/array.TestArrayFromStdModule.py
import-std-module/deque-basic.TestDequeFromStdModule.py
import-std-module/deque-dbg-info-content.TestDbgInfoContentDequeFromStdModule.py
import-std-module/forward_list.TestForwardListFromStdModule.py
import-std-module/forward_list-dbg-info-content.TestDbgInfoContentForwardListFromStdModule.py
import-std-module/list.TestListFromStdModule.py
import-std-module/list-dbg-info-content.TestDbgInfoContentListFromStdModule.py
import-std-module/queue.TestQueueFromStdModule.py
import-std-module/stack.TestStackFromStdModule.py
import-std-module/vector.TestVectorFromStdModule.py
import-std-module/vector-bool.TestVectorBoolFromStdModule.py
import-std-module/vector-dbg-info-content.TestDbgInfoContentVectorFromStdModule.py
import-std-module/vector-of-vectors.TestVectorOfVectorsFromStdModule.py
https://green.lab.llvm.org/green/view/LLDB/job/lldb-cmake/45301/
Without this patch, clang will not wrap in an ElaboratedType node types written
without a keyword and nested name qualifier, which goes against the intent that
we should produce an AST which retains enough details to recover how things are
written.
The lack of this sugar is incompatible with the intent of the type printer
default policy, which is to print types as written, but to fall back and print
them fully qualified when they are desugared.
An ElaboratedTypeLoc without keyword / NNS uses no storage by itself, but still
requires pointer alignment due to pre-existing bug in the TypeLoc buffer
handling.
Signed-off-by: Matheus Izvekov <mizvekov@gmail.com>
Differential Revision: https://reviews.llvm.org/D112374
Special classes such as accessor, sampler, and stream need additional
implementation when they are passed from host to device.
This patch is adding a new attribute “sycl_special_class” used to mark
SYCL classes/struct that need the additional compiler handling.
Adds diagnosing on attempt to use zero length arrays, pointers, refs, arrays
of them and structs/classes containing all of it.
In case a struct/class with zero length array is used this emits a set
of notes pointing out how zero length array got into used struct, like
this:
```
struct ContainsArr {
int A[0]; // note: field of illegal type declared here
};
struct Wrapper {
ContainsArr F; // note: within field of type ContainsArr declared here
// ...
}
// Device code
Wrapper W;
W.use(); // error: zero-length arrays are not permitted
```
Total deep check of each used declaration may result in double
diagnosing at the same location.
Reviewed By: aaron.ballman
Differential Revision: https://reviews.llvm.org/D114080
This patch attempts to fix a compiler crash that occurs when long
double type is used with -mno-x87 compiler option.
The option disables x87 target feature, which in turn disables x87
registers, so CG cannot select them for x86_fp80 LLVM IR type. Long
double is lowered as x86_fp80 for some targets, so it leads to a
crash.
The option seems to contradict the SystemV ABI, which requires long
double to be represented as a 80-bit floating point, and it also
requires to use x87 registers.
To avoid that, `long double` type is disabled when -mno-x87 option is
set. In addition to that, `float` and `double` also use x87 registers
for return values on 32-bit x86, so they are disabled as well.
Differential Revision: https://reviews.llvm.org/D98895
This was committed as ec6c847179fd, but then reverted after a failure
in: https://lab.llvm.org/buildbot/#/builders/84/builds/13983
I was not able to reproduce the problem, but I added an extra check
for a NULL QualType just in case.
Original comit message:
The patch adds missing diagnostics for cases like:
float F3 = ((__float128)F1 * (__float128)F2) / 2.0f;
Sema::checkDeviceDecl (renamed to checkTypeSupport) is changed to work
with a type without the corresponding ValueDecl. It is also refactored
so that host diagnostics for unsupported types can be added here as
well.
Differential Revision: https://reviews.llvm.org/D109315
After significant problems in our downstream with the previous
implementation, the SYCL standard has opted to make using macros/etc to
change kernel-naming-lambdas in any way UB (even passively). As a
result, we are able to just emit the itanium mangling.
However, this DOES require a little work in the CXXABI, as the microsoft
and itanium mangler use different numbering schemes for lambdas. This
patch adds a pair of mangling contexts that use the normal 'itanium'
mangling strategy to fill in the "DeviceManglingNumber" used previously
by CUDA.
Differential Revision: https://reviews.llvm.org/D110281
The patch adds missing diagnostics for cases like:
float F3 = ((__float128)F1 * (__float128)F2) / 2.0f;
Sema::checkDeviceDecl (renamed to checkTypeSupport) is changed to work
with a type without the corresponding ValueDecl. It is also refactored
so that host diagnostics for unsupported types can be added here as
well.
Differential Revision: https://reviews.llvm.org/D109315
I discovered when merging the __builtin_sycl_unique_stable_name into my
downstream that it is actually possible for the cc1 invocation to have
more than 1 Sema instance, if you pass it multiple input files, each
gets its own Sema instance and thus ASTContext instance. The result was
that the call to Filter the SYCL kernels was using an
ItaniumMangleContext stored via a 'magic static', so it had an invalid
reference to ASTContext when processing the 2nd failure.
The failure is unfortunately flakey/transient, but the test that fails
was added anyway.
The magic-static was switched to a unique_ptr member variable in
ASTContext that is initialized when needed.
The original version of this was reverted, and @rjmcall provided some
advice to architect a new solution. This is that solution.
This implements a builtin to provide a unique name that is stable across
compilations of this TU for the purposes of implementing the library
component of the unnamed kernel feature of SYCL. It does this by
running the Itanium mangler with a few modifications.
Because it is somewhat common to wrap non-kernel-related lambdas in
macros that aren't present on the device (such as for logging), this
uniquely generates an ID for all lambdas involved in the naming of a
kernel. It uses the lambda-mangling number to do this, except replaces
this with its own number (starting at 10000 for readabililty reasons)
for lambdas used to name a kernel.
Additionally, this implements itself as constexpr with a slight catch:
if a name would be invalidated by the use of this lambda in a later
kernel invocation, it is diagnosed as an error (see the Sema tests).
Differential Revision: https://reviews.llvm.org/D103112
Default address space (applies when no explicit address space was
specified) maps to generic (4) address space.
Added SYCL named address spaces `sycl_global`, `sycl_local` and
`sycl_private` defined as sub-sets of the default address space.
Static variables without address space now reside in global address
space when compile for SPIR target, unless they have an explicit address
space qualifier in source code.
Differential Revision: https://reviews.llvm.org/D89909
SYCL compilations initiated by the driver will spawn off one or more
frontend compilation jobs (one for device and one for host). This patch
reworks the driver options to make upstreaming this from the downstream
SYCL fork easier.
This patch introduces a language option to identify host executions
(SYCLIsHost) and a -cc1 frontend option to enable this mode. -fsycl and
-fno-sycl become driver-only options that are rejected when passed to
-cc1. This is because the frontend and beyond should be looking at
whether the user is doing a device or host compilation specifically.
Because the frontend should only ever be in one mode or the other,
-fsycl-is-device and -fsycl-is-host are mutually exclusive options.
Emit error for use of 128-bit integer inside device code had been
already implemented in https://reviews.llvm.org/D74387. However,
the error is not emitted for SPIR64, because for SPIR64, hasInt128Type
return true.
hasInt128Type: is also used to control generation of certain 128-bit
predefined macros, initializer predefined 128-bit integer types and
build 128-bit ArithmeticTypes. Except predefined macros, only the
device target is considered, since error only emit when 128-bit
integer is used inside device code, the host target (auxtarget) also
needs to be considered.
The change address:
1. (SPIR.h) Correct hasInt128Type() for SPIR targets.
2. Sema.cpp and SemaOverload.cpp: Add additional check to consider host
target(auxtarget) when call to hasInt128Type. So that __int128_t
and __int128() are allowed to avoid error when they used outside
device code.
3. SemaType.cpp: add check for SYCLIsDevice to delay the error message.
The error will be emitted if the use of 128-bit integer in the device
code.
Reviewed By: Johannes Doerfert and Aaron Ballman
Differential Revision: https://reviews.llvm.org/D92439
Summary:
SYCL and OpenMP prohibits thread local storage in device code,
so this commit ensures that error is emitted for device code and not
emitted for host code when host target supports it.
Reviewers: jdoerfert, erichkeane, bader
Reviewed By: jdoerfert, erichkeane
Subscribers: guansong, riccibruno, ABataev, yaxunl, ebevhan, Anastasia, sstefan1, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D81641
Summary:
Diagnostic is emitted if some declaration of unsupported type
declaration is used inside device code.
Memcpy operations for structs containing member with unsupported type
are allowed. Fixed crash on attempt to emit diagnostic outside of the
functions.
The approach is generalized between SYCL and OpenMP.
CUDA/OMP deferred diagnostic interface is going to be used for SYCL device.
Reviewers: rsmith, rjmccall, ABataev, erichkeane, bader, jdoerfert, aaron.ballman
Reviewed By: jdoerfert
Subscribers: guansong, sstefan1, yaxunl, mgorny, bader, ebevhan, Anastasia, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D74387
Summary:
User can select the version of SYCL the compiler will
use via the flag -sycl-std, similar to -cl-std.
The flag defines the LangOpts.SYCLVersion option to the
version of SYCL. The default value is undefined.
If driver is building SYCL code, flag is set to the default SYCL
version (1.2.1)
The preprocessor uses this variable to define CL_SYCL_LANGUAGE_VERSION macro,
which should be defined according to SYCL 1.2.1 standard.
Only valid value at this point for the flag is 1.2.1.
Co-Authored-By: David Wood <Q0KPU0H1YOEPHRY1R2SN5B5RL@david.davidtw.co>
Signed-off-by: Ruyman Reyes <ruyman@codeplay.com>
Subscribers: ebevhan, Anastasia, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D72857
Summary:
User can select the version of SYCL the compiler will
use via the flag -sycl-std, similar to -cl-std.
The flag defines the LangOpts.SYCLVersion option to the
version of SYCL. The default value is undefined.
If driver is building SYCL code, flag is set to the default SYCL
version (1.2.1)
The preprocessor uses this variable to define CL_SYCL_LANGUAGE_VERSION macro,
which should be defined according to SYCL 1.2.1 standard.
Only valid value at this point for the flag is 1.2.1.
Co-Authored-By: David Wood <Q0KPU0H1YOEPHRY1R2SN5B5RL@david.davidtw.co>
Signed-off-by: Ruyman Reyes <ruyman@codeplay.com>
Subscribers: ebevhan, Anastasia, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D72857
Signed-off-by: Alexey Bader <alexey.bader@intel.com>
SYCL is single source offload programming model relying on compiler to
separate device code (i.e. offloaded to an accelerator) from the code
executed on the host.
Here is code example of the SYCL program to demonstrate compiler
outlining work:
```
int foo(int x) { return ++x; }
int bar(int x) { throw std::exception("CPU code only!"); }
...
using namespace cl::sycl;
queue Q;
buffer<int, 1> a(range<1>{1024});
Q.submit([&](handler& cgh) {
auto A = a.get_access<access::mode::write>(cgh);
cgh.parallel_for<init_a>(range<1>{1024}, [=](id<1> index) {
A[index] = index[0] + foo(42);
});
}
...
```
SYCL device compiler must compile lambda expression passed to
cl::sycl::handler::parallel_for method and function foo called from this
lambda expression for an "accelerator". SYCL device compiler also must
ignore bar function as it's not required for offloaded code execution.
This patch adds the sycl_kernel attribute, which is used to mark code
passed to cl::sycl::handler::parallel_for as "accelerated code".
Attribute must be applied to function templates which parameters include
at least "kernel name" and "kernel function object". These parameters
will be used to establish an ABI between the host application and
offloaded part.
Reviewers: jlebar, keryell, Naghasan, ABataev, Anastasia, bader, aaron.ballman, rjmccall, rsmith
Reviewed By: keryell, bader
Subscribers: mgorny, OlegM, ArturGainullin, agozillon, aaron.ballman, ebevhan, Anastasia, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D60455
Signed-off-by: Alexey Bader <alexey.bader@intel.com>
