[SYCL] SYCL host kernel launch support for the sycl_kernel_entry_point attribute. (#152403)
The `sycl_kernel_entry_point` attribute facilitates the generation of an
offload kernel entry point function based on the parameters and body
of the attributed function. This change extends the behavior of that
attribute to support integration with a SYCL runtime library through
an interface that communicates symbol names and kernel arguments
for the generated offload kernel entry point functions.
Consider the following function declared with the
`sycl_kernel_entry_point` attribute with a call to this function
occurring in the implementation of a SYCL kernel invocation function
such as `sycl::handler::single_task()`.
```c++
template<typename KernelName, typename KernelType>
[[clang::sycl_kernel_entry_point(KernelName)]]
void kernel_entry_point(KernelType kernel) {
kernel();
}
```
The body of the above function specifies the parameters and body of the
generated offload kernel entry point. Clearly, a call to the above
function by a SYCL kernel invocation function is not intended to execute
the body as written. Previously, code generation emitted an empty
function body so that calls to the function had no effect other than to
trigger the generation of the offload kernel entry point. The function
body is therefore available to hook for SYCL library support and is now
substituted with a call to a (SYCL library provided) function template
or variable template named `sycl_kernel_launch()` with the kernel
name type passed as the first template argument, the symbol name
of the offload kernel entry point passed as a string literal for the first
function argument, and the function parameters passed as the
remaining explicit function arguments. Given a call like this:
```c++
kernel_entry_point<struct KN>([]{})
```
the body of the instantiated `kernel_entry_point()` specialization would
be substituted as follows with "kernel-symbol-name" substituted for the
generated symbol name and `kernel` forwarded.
```c++
sycl_kernel_launch<KN>("kernel-symbol-name", kernel)
```
Name lookup and overload resolution for the `sycl_kernel_launch()`
function is performed at the point of definition of the
`sycl_kernel_entry_point` attributed function (or the point of
instantiation for an instantiated function template specialization). If
overload resolution fails, the program is ill-formed.
Implementation of the `sycl_kernel_launch()` function might require
additional information provided by the SYCL library. This is facilitated
by removing the previous prohibition against use of the
`sycl_kernel_entry_point` attribute with a non-static member function.
If the `sycl_kernel_entry_point` attributed function is a non-static
member function, then overload resolution for the `sycl_kernel_launch()`
function template may select a non-static member function in which case,
`this` will be implicitly passed as the implicit object argument.
If a `sycl_kernel_entry_point` attributed function is a non-static
member function, use of `this` in a potentially evaluated expression is
prohibited in the definition since `this` is not a kernel argument and
will not be available within the generated offload kernel entry point
function. The attribute cannot be applied to a function with an
explicit object parameter.
---------
Co-authored-by: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
This commit is contained in:
Tom Honermann
GitHub
parent
41ac6ebf0e
commit
23e4fe040b
@ -839,8 +839,10 @@ public:
|
||||
|
||||
void VisitSYCLKernelCallStmt(const SYCLKernelCallStmt *Node) {
|
||||
Visit(Node->getOriginalStmt());
|
||||
if (Traversal != TK_IgnoreUnlessSpelledInSource)
|
||||
if (Traversal != TK_IgnoreUnlessSpelledInSource) {
|
||||
Visit(Node->getKernelLaunchStmt());
|
||||
Visit(Node->getOutlinedFunctionDecl());
|
||||
}
|
||||
}
|
||||
|
||||
void VisitOMPExecutableDirective(const OMPExecutableDirective *Node) {
|
||||
|
||||
@ -3001,6 +3001,13 @@ DEF_TRAVERSE_STMT(ParenListExpr, {})
|
||||
DEF_TRAVERSE_STMT(SYCLUniqueStableNameExpr, {
|
||||
TRY_TO(TraverseTypeLoc(S->getTypeSourceInfo()->getTypeLoc()));
|
||||
})
|
||||
DEF_TRAVERSE_STMT(UnresolvedSYCLKernelCallStmt, {
|
||||
if (getDerived().shouldVisitImplicitCode()) {
|
||||
TRY_TO(TraverseStmt(S->getOriginalStmt()));
|
||||
TRY_TO(TraverseStmt(S->getKernelLaunchIdExpr()));
|
||||
ShouldVisitChildren = false;
|
||||
}
|
||||
})
|
||||
DEF_TRAVERSE_STMT(OpenACCAsteriskSizeExpr, {})
|
||||
DEF_TRAVERSE_STMT(PredefinedExpr, {})
|
||||
DEF_TRAVERSE_STMT(ShuffleVectorExpr, {})
|
||||
@ -3038,6 +3045,7 @@ DEF_TRAVERSE_STMT(CapturedStmt, { TRY_TO(TraverseDecl(S->getCapturedDecl())); })
|
||||
DEF_TRAVERSE_STMT(SYCLKernelCallStmt, {
|
||||
if (getDerived().shouldVisitImplicitCode()) {
|
||||
TRY_TO(TraverseStmt(S->getOriginalStmt()));
|
||||
TRY_TO(TraverseStmt(S->getKernelLaunchStmt()));
|
||||
TRY_TO(TraverseDecl(S->getOutlinedFunctionDecl()));
|
||||
ShouldVisitChildren = false;
|
||||
}
|
||||
|
||||
@ -28,40 +28,44 @@ namespace clang {
|
||||
/// of such a function specifies the statements to be executed on a SYCL device
|
||||
/// to invoke a SYCL kernel with a particular set of kernel arguments. The
|
||||
/// SYCLKernelCallStmt associates an original statement (the compound statement
|
||||
/// that is the function body) with an OutlinedFunctionDecl that holds the
|
||||
/// kernel parameters and the transformed body. During code generation, the
|
||||
/// OutlinedFunctionDecl is used to emit an offload kernel entry point suitable
|
||||
/// for invocation from a SYCL library implementation. If executed, the
|
||||
/// SYCLKernelCallStmt behaves as a no-op; no code generation is performed for
|
||||
/// it.
|
||||
/// that is the function body) with a kernel launch statement to execute on a
|
||||
/// SYCL host and an OutlinedFunctionDecl that holds the kernel parameters and
|
||||
/// the transformed body to execute on a SYCL device. During code generation,
|
||||
/// the OutlinedFunctionDecl is used to emit an offload kernel entry point
|
||||
/// suitable for invocation from a SYCL library implementation.
|
||||
class SYCLKernelCallStmt : public Stmt {
|
||||
friend class ASTStmtReader;
|
||||
friend class ASTStmtWriter;
|
||||
|
||||
private:
|
||||
Stmt *OriginalStmt = nullptr;
|
||||
Stmt *KernelLaunchStmt = nullptr;
|
||||
OutlinedFunctionDecl *OFDecl = nullptr;
|
||||
|
||||
public:
|
||||
/// Construct a SYCL kernel call statement.
|
||||
SYCLKernelCallStmt(CompoundStmt *CS, OutlinedFunctionDecl *OFD)
|
||||
: Stmt(SYCLKernelCallStmtClass), OriginalStmt(CS), OFDecl(OFD) {}
|
||||
SYCLKernelCallStmt(CompoundStmt *CS, Stmt *S, OutlinedFunctionDecl *OFD)
|
||||
: Stmt(SYCLKernelCallStmtClass), OriginalStmt(CS), KernelLaunchStmt(S),
|
||||
OFDecl(OFD) {}
|
||||
|
||||
/// Construct an empty SYCL kernel call statement.
|
||||
SYCLKernelCallStmt(EmptyShell Empty) : Stmt(SYCLKernelCallStmtClass, Empty) {}
|
||||
|
||||
/// Retrieve the model statement.
|
||||
CompoundStmt *getOriginalStmt() { return cast<CompoundStmt>(OriginalStmt); }
|
||||
const CompoundStmt *getOriginalStmt() const {
|
||||
return cast<CompoundStmt>(OriginalStmt);
|
||||
}
|
||||
|
||||
void setOriginalStmt(CompoundStmt *CS) { OriginalStmt = CS; }
|
||||
|
||||
/// Retrieve the outlined function declaration.
|
||||
Stmt *getKernelLaunchStmt() { return KernelLaunchStmt; }
|
||||
const Stmt *getKernelLaunchStmt() const { return KernelLaunchStmt; }
|
||||
|
||||
void setKernelLaunchStmt(Stmt *S) { KernelLaunchStmt = S; }
|
||||
|
||||
OutlinedFunctionDecl *getOutlinedFunctionDecl() { return OFDecl; }
|
||||
const OutlinedFunctionDecl *getOutlinedFunctionDecl() const { return OFDecl; }
|
||||
|
||||
/// Set the outlined function declaration.
|
||||
void setOutlinedFunctionDecl(OutlinedFunctionDecl *OFD) { OFDecl = OFD; }
|
||||
|
||||
SourceLocation getBeginLoc() const LLVM_READONLY {
|
||||
@ -89,6 +93,66 @@ public:
|
||||
}
|
||||
};
|
||||
|
||||
// UnresolvedSYCLKernelCallStmt represents an invocation of a SYCL kernel in
|
||||
// a dependent context for which lookup of the sycl_kernel_launch identifier
|
||||
// cannot be performed. These statements are transformed to SYCLKernelCallStmt
|
||||
// during template instantiation.
|
||||
class UnresolvedSYCLKernelCallStmt : public Stmt {
|
||||
friend class ASTStmtReader;
|
||||
friend class ASTStmtWriter;
|
||||
|
||||
private:
|
||||
Stmt *OriginalStmt = nullptr;
|
||||
// KernelLaunchIdExpr stores an UnresolvedLookupExpr or UnresolvedMemberExpr
|
||||
// corresponding to the SYCL kernel launch function for which a call
|
||||
// will be synthesized during template instantiation.
|
||||
Expr *KernelLaunchIdExpr = nullptr;
|
||||
|
||||
UnresolvedSYCLKernelCallStmt(CompoundStmt *CS, Expr *IdExpr)
|
||||
: Stmt(UnresolvedSYCLKernelCallStmtClass), OriginalStmt(CS),
|
||||
KernelLaunchIdExpr(IdExpr) {}
|
||||
|
||||
void setOriginalStmt(CompoundStmt *CS) { OriginalStmt = CS; }
|
||||
|
||||
void setKernelLaunchIdExpr(Expr *IdExpr) { KernelLaunchIdExpr = IdExpr; }
|
||||
|
||||
public:
|
||||
static UnresolvedSYCLKernelCallStmt *Create(const ASTContext &C,
|
||||
CompoundStmt *CS, Expr *IdExpr) {
|
||||
return new (C) UnresolvedSYCLKernelCallStmt(CS, IdExpr);
|
||||
}
|
||||
|
||||
static UnresolvedSYCLKernelCallStmt *CreateEmpty(const ASTContext &C) {
|
||||
return new (C) UnresolvedSYCLKernelCallStmt(nullptr, nullptr);
|
||||
}
|
||||
|
||||
CompoundStmt *getOriginalStmt() { return cast<CompoundStmt>(OriginalStmt); }
|
||||
const CompoundStmt *getOriginalStmt() const {
|
||||
return cast<CompoundStmt>(OriginalStmt);
|
||||
}
|
||||
|
||||
Expr *getKernelLaunchIdExpr() { return KernelLaunchIdExpr; }
|
||||
const Expr *getKernelLaunchIdExpr() const { return KernelLaunchIdExpr; }
|
||||
|
||||
SourceLocation getBeginLoc() const LLVM_READONLY {
|
||||
return getOriginalStmt()->getBeginLoc();
|
||||
}
|
||||
|
||||
SourceLocation getEndLoc() const LLVM_READONLY {
|
||||
return getOriginalStmt()->getEndLoc();
|
||||
}
|
||||
static bool classof(const Stmt *T) {
|
||||
return T->getStmtClass() == UnresolvedSYCLKernelCallStmtClass;
|
||||
}
|
||||
child_range children() {
|
||||
return child_range(&OriginalStmt, &OriginalStmt + 1);
|
||||
}
|
||||
|
||||
const_child_range children() const {
|
||||
return const_child_range(&OriginalStmt, &OriginalStmt + 1);
|
||||
}
|
||||
};
|
||||
|
||||
} // end namespace clang
|
||||
|
||||
#endif // LLVM_CLANG_AST_STMTSYCL_H
|
||||
|
||||
@ -580,25 +580,26 @@ The following examples demonstrate the use of this attribute:
|
||||
def SYCLKernelEntryPointDocs : Documentation {
|
||||
let Category = DocCatFunction;
|
||||
let Content = [{
|
||||
The ``sycl_kernel_entry_point`` attribute facilitates the generation of an
|
||||
offload kernel entry point, sometimes called a SYCL kernel caller function,
|
||||
suitable for invoking a SYCL kernel on an offload device. The attribute is
|
||||
intended for use in the implementation of SYCL kernel invocation functions
|
||||
like the ``single_task`` and ``parallel_for`` member functions of the
|
||||
``sycl::handler`` class specified in section 4.9.4, "Command group ``handler``
|
||||
class", of the SYCL 2020 specification.
|
||||
The ``sycl_kernel_entry_point`` attribute facilitates the launch of a SYCL
|
||||
kernel and the generation of an offload kernel entry point, sometimes called
|
||||
a SYCL kernel caller function, suitable for invoking a SYCL kernel on an
|
||||
offload device. The attribute is intended for use in the implementation of
|
||||
SYCL kernel invocation functions like the ``single_task`` and ``parallel_for``
|
||||
member functions of the ``sycl::handler`` class specified in section 4.9.4,
|
||||
"Command group ``handler`` class", of the SYCL 2020 specification.
|
||||
|
||||
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 attribute requires a single type argument 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 attribute only appertains to functions and only those that meet the
|
||||
following requirements.
|
||||
|
||||
* Has a non-deduced ``void`` return type.
|
||||
* Is not a non-static member function, constructor, or destructor.
|
||||
* Is not a constructor or destructor.
|
||||
* Is not a non-static member function with an explicit object parameter.
|
||||
* Is not a C variadic function.
|
||||
* Is not a coroutine.
|
||||
* Is not defined as deleted or as defaulted.
|
||||
@ -613,73 +614,84 @@ follows.
|
||||
|
||||
namespace sycl {
|
||||
class handler {
|
||||
template<typename KernelNameType, typename KernelType>
|
||||
[[ clang::sycl_kernel_entry_point(KernelNameType) ]]
|
||||
static void kernel_entry_point(KernelType kernel) {
|
||||
kernel();
|
||||
template<typename KernelName, typename... Ts>
|
||||
void sycl_kernel_launch(const char* kernelSymbol, Ts&&... kernelArgs) {
|
||||
// This code will run on the host and is responsible for calling functions
|
||||
// appropriate for the desired offload backend (OpenCL, CUDA, HIP,
|
||||
// Level Zero, etc...) to copy the kernel arguments denoted by kernelArgs
|
||||
// to a device and to schedule an invocation of the offload kernel entry
|
||||
// point denoted by kernelSymbol with the copied arguments.
|
||||
}
|
||||
|
||||
template<typename KernelName, typename KernelType>
|
||||
[[ clang::sycl_kernel_entry_point(KernelName) ]]
|
||||
void kernel_entry_point(KernelType kernelFunc) {
|
||||
// This code will run on the device. The call to kernelFunc() invokes
|
||||
// the SYCL kernel.
|
||||
kernelFunc();
|
||||
}
|
||||
|
||||
public:
|
||||
template<typename KernelNameType, typename KernelType>
|
||||
void single_task(KernelType kernel) {
|
||||
// Call kernel_entry_point() to trigger generation of an offload
|
||||
// kernel entry point.
|
||||
kernel_entry_point<KernelNameType>(kernel);
|
||||
// Call functions appropriate for the desired offload backend
|
||||
// (OpenCL, CUDA, HIP, Level Zero, etc...).
|
||||
template<typename KernelName, typename KernelType>
|
||||
void single_task(const KernelType& kernelFunc) {
|
||||
// This code will run on the host. kernel_entry_point() is called to
|
||||
// trigger generation of an offload kernel entry point and to schedule
|
||||
// an invocation of it on a device with kernelFunc (a SYCL kernel object)
|
||||
// passed as a kernel argument. This call will result in an implicit call
|
||||
// to sycl_kernel_launch() with the symbol name for the generated offload
|
||||
// kernel entry point passed as the first function argument followed by
|
||||
// kernelFunc.
|
||||
kernel_entry_point<KernelName>(kernelFunc);
|
||||
}
|
||||
};
|
||||
} // namespace sycl
|
||||
|
||||
A SYCL kernel is a callable object of class type that is constructed on a host,
|
||||
often via a lambda expression, and then passed to a SYCL kernel invocation
|
||||
function to be executed on an offload device. A SYCL kernel invocation function
|
||||
is responsible for copying the provided SYCL kernel object to an offload
|
||||
device and initiating a call to it. The SYCL kernel object and its data members
|
||||
constitute the parameters of an offload kernel.
|
||||
A SYCL kernel object is a callable object of class type that is constructed on
|
||||
a host, often via a lambda expression, and then passed to a SYCL kernel
|
||||
invocation function to be executed on an offload device. The ``kernelFunc``
|
||||
parameters in the example code above correspond to SYCL kernel objects.
|
||||
|
||||
A SYCL kernel type is required to satisfy the device copyability requirements
|
||||
specified in section 3.13.1, "Device copyable", of the SYCL 2020 specification.
|
||||
Additionally, any data members of the kernel object type are required to satisfy
|
||||
section 4.12.4, "Rules for parameter passing to kernels". For most types, these
|
||||
rules require that the type is trivially copyable. However, the SYCL
|
||||
specification mandates that certain special SYCL types, such as
|
||||
``sycl::accessor`` and ``sycl::stream`` be device copyable even if they are not
|
||||
trivially copyable. These types require special handling because they cannot
|
||||
be copied to device memory as if by ``memcpy()``. Additionally, some offload
|
||||
backends, OpenCL for example, require objects of some of these types to be
|
||||
passed as individual arguments to the offload kernel.
|
||||
A SYCL kernel object type is required to satisfy the device copyability
|
||||
requirements specified in section 3.13.1, "Device copyable", of the SYCL 2020
|
||||
specification. Additionally, any data members of the kernel object type are
|
||||
required to satisfy section 4.12.4, "Rules for parameter passing to kernels".
|
||||
For most types, these rules require that the type is trivially copyable.
|
||||
However, the SYCL specification mandates that certain special SYCL types, such
|
||||
as ``sycl::accessor`` and ``sycl::stream``, be device copyable even if they are
|
||||
not trivially copyable. These types require special handling because they cannot
|
||||
necessarily be copied to device memory as if by ``memcpy()``.
|
||||
|
||||
An offload kernel consists of an entry point function that declares the
|
||||
parameters of the offload kernel and the set of all functions and variables that
|
||||
are directly or indirectly used by the entry point function.
|
||||
The SYCL kernel object and its data members constitute the parameters of an
|
||||
offload kernel. An offload kernel consists of an offload entry point function
|
||||
and the set of all functions and variables that are directly or indirectly used
|
||||
by the entry point function.
|
||||
|
||||
A SYCL kernel invocation function invokes a SYCL kernel on a device by
|
||||
performing the following tasks (likely with the help of an offload backend
|
||||
like OpenCL):
|
||||
A SYCL kernel invocation function is responsible for performing the following
|
||||
tasks (likely with the help of an offload backend like OpenCL):
|
||||
|
||||
#. Identifying the offload kernel entry point to be used for the SYCL kernel.
|
||||
|
||||
#. Deconstructing the SYCL kernel object, if necessary, to produce the set of
|
||||
offload kernel arguments required by the offload kernel entry point.
|
||||
#. Validating that the SYCL kernel object type and its data members meet the
|
||||
SYCL device copyability and kernel parameter requirements noted above.
|
||||
|
||||
#. Copying the offload kernel arguments to device memory.
|
||||
#. Copying the SYCL kernel object and any other kernel arguments to device
|
||||
memory including any special handling required for SYCL special types.
|
||||
|
||||
#. Initiating execution of the offload kernel entry point.
|
||||
|
||||
The offload kernel entry point for a SYCL kernel performs the following tasks:
|
||||
|
||||
#. Reconstituting the SYCL kernel object, if necessary, using the offload
|
||||
kernel parameters.
|
||||
#. Calling the ``operator()`` member function of the SYCL kernel object.
|
||||
|
||||
#. Calling the ``operator()`` member function of the (reconstituted) SYCL kernel
|
||||
object.
|
||||
The ``sycl_kernel_entry_point`` attribute facilitates or automates these tasks
|
||||
by providing generation of an offload kernel entry point with a unique symbol
|
||||
name, type checking of kernel argument requirements, and initiation of kernel
|
||||
execution via synthesized calls to a ``sycl_kernel_launch`` template.
|
||||
|
||||
The ``sycl_kernel_entry_point`` attribute automates generation of an offload
|
||||
kernel entry point that performs those latter tasks. The parameters and body of
|
||||
a function declared with the ``sycl_kernel_entry_point`` attribute specify a
|
||||
pattern from which the parameters and body of the entry point function are
|
||||
derived. Consider the following call to a SYCL kernel invocation function.
|
||||
A function declared with the ``sycl_kernel_entry_point`` attribute specifies
|
||||
the parameters and body of an offload entry point function. Consider the
|
||||
following call to the ``single_task()`` SYCL kernel invocation function assuming
|
||||
an implementation similar to the one shown above.
|
||||
|
||||
.. code-block:: c++
|
||||
|
||||
@ -690,65 +702,87 @@ derived. Consider the following call to a SYCL kernel invocation function.
|
||||
});
|
||||
}
|
||||
|
||||
The SYCL kernel object is the result of the lambda expression. It has two
|
||||
data members corresponding to the captures of ``sout`` and ``s``. Since one
|
||||
of these data members corresponds to a special SYCL type that must be passed
|
||||
individually as an offload kernel parameter, it is necessary to decompose the
|
||||
SYCL kernel object into its constituent parts; the offload kernel will have
|
||||
two kernel parameters. Given a SYCL implementation that uses a
|
||||
``sycl_kernel_entry_point`` attributed function like the one shown above, an
|
||||
offload kernel entry point function will be generated that looks approximately
|
||||
The SYCL kernel object is the result of the lambda expression. The call to
|
||||
``kernel_entry_point()`` via the call to ``single_task()`` triggers the
|
||||
generation of an offload kernel entry point function that looks approximately
|
||||
as follows.
|
||||
|
||||
.. code-block:: c++
|
||||
|
||||
void sycl-kernel-caller-for-KN(sycl::stream sout, S s) {
|
||||
kernel-type kernel = { sout, s );
|
||||
kernel();
|
||||
void sycl-kernel-caller-for-KN(kernel-type kernelFunc) {
|
||||
kernelFunc();
|
||||
}
|
||||
|
||||
There are a few items worthy of note:
|
||||
|
||||
#. The name of the generated function incorporates the SYCL kernel name,
|
||||
``KN``, that was passed as the ``KernelNameType`` template parameter to
|
||||
``kernel_entry_point()`` and provided as the argument to the
|
||||
``sycl_kernel_entry_point`` attribute. There is a one-to-one correspondence
|
||||
between SYCL kernel names and offload kernel entry points.
|
||||
#. ``sycl-kernel-caller-for-KN`` is an exposition only name; the actual name
|
||||
generated for an entry point is an implementation detail and subject to
|
||||
change. However, the name will incorporate the SYCL kernel name, ``KN``,
|
||||
that was passed as the ``KernelName`` template parameter to
|
||||
``single_task()`` and eventually provided as the argument to the
|
||||
``sycl_kernel_entry_point`` attribute in order to ensure that a unique
|
||||
name is generated for each entry point. There is a one-to-one correspondence
|
||||
between SYCL kernel names and offload kernel entry points.
|
||||
|
||||
#. The SYCL kernel is a lambda closure type and therefore has no name;
|
||||
``kernel-type`` is substituted above and corresponds to the ``KernelType``
|
||||
template parameter deduced in the call to ``kernel_entry_point()``.
|
||||
Lambda types cannot be declared and initialized using the aggregate
|
||||
initialization syntax used above, but the intended behavior should be clear.
|
||||
template parameter deduced in the call to ``single_task()``.
|
||||
|
||||
#. ``S`` is a device copyable type that does not directly or indirectly contain
|
||||
a data member of a SYCL special type. It therefore does not need to be
|
||||
decomposed into its constituent members to be passed as a kernel argument.
|
||||
#. The parameter and the call to ``kernelFunc()`` in the function body
|
||||
correspond to the definition of ``kernel_entry_point()`` as called by
|
||||
``single_task()``.
|
||||
|
||||
#. The depiction of the ``sycl::stream`` parameter as a single self contained
|
||||
kernel parameter is an oversimplification. SYCL special types may require
|
||||
additional decomposition such that the generated function might have three
|
||||
or more parameters depending on how the SYCL library implementation defines
|
||||
these types.
|
||||
#. The parameter is type checked for conformance with the SYCL device
|
||||
copyability and kernel parameter requirements.
|
||||
|
||||
#. The call to ``kernel_entry_point()`` has no effect other than to trigger
|
||||
emission of the entry point function. The statments that make up the body
|
||||
of the function are not executed when the function is called; they are
|
||||
only used in the generation of the entry point function.
|
||||
Within ``single_task()``, the call to ``kernel_entry_point()`` is effectively
|
||||
replaced with a synthesized call to a ''sycl_kernel_launch`` template that
|
||||
looks approximately as follows.
|
||||
|
||||
.. code-block:: c++
|
||||
|
||||
sycl_kernel_launch<KN>("sycl-kernel-caller-for-KN", kernelFunc);
|
||||
|
||||
There are a few items worthy of note:
|
||||
|
||||
#. Lookup for the ``sycl_kernel_launch`` template is performed as if from the
|
||||
body of the (possibly instantiated) definition of ``kernel_entry_point()``.
|
||||
If name lookup or overload resolution fails, the program is ill-formed.
|
||||
If the selected overload is a non-static member function, then ``this`` is
|
||||
passed as the implicit object parameter.
|
||||
|
||||
#. Function arguments passed to ``sycl_kernel_launch()`` are passed
|
||||
as if by ``std::move(x)``.
|
||||
|
||||
#. The ``sycl_kernel_launch`` template is expected to be provided by the SYCL
|
||||
library implementation. It is responsible for copying the kernel arguments
|
||||
to device memory and for scheduling execution of the generated offload
|
||||
kernel entry point identified by the symbol name passed as the first
|
||||
function argument. ``sycl-kernel-caller-for-KN`` is substituted above for
|
||||
the actual symbol name that would be generated for the offload kernel entry
|
||||
point.
|
||||
|
||||
It is not necessary for a function declared with the ``sycl_kernel_entry_point``
|
||||
attribute to be called for the offload kernel entry point to be emitted. For
|
||||
inline functions and function templates, any ODR-use will suffice. For other
|
||||
functions, an ODR-use is not required; the offload kernel entry point will be
|
||||
emitted if the function is defined.
|
||||
emitted if the function is defined. In any case, a call to the function is
|
||||
required for the synthesized call to ``sycl_kernel_launch()`` to occur.
|
||||
|
||||
A function declared with the ``sycl_kernel_entry_point`` attribute may include
|
||||
an exception specification. If a non-throwing exception specification is
|
||||
present, an exception propagating from the implicit call to the
|
||||
``sycl_kernel_launch`` template will result in a call to ``std::terminate()``.
|
||||
Otherwise, such an exception will propagate normally.
|
||||
|
||||
Functions declared with the ``sycl_kernel_entry_point`` attribute are not
|
||||
limited to the simple example shown above. They may have additional template
|
||||
parameters, declare additional function parameters, and have complex control
|
||||
flow in the function body. Function parameter decomposition and reconstitution
|
||||
is performed for all function parameters. The function must abide by the
|
||||
language feature restrictions described in section 5.4, "Language restrictions
|
||||
for device functions" in the SYCL 2020 specification.
|
||||
flow in the function body. The function must abide by the language feature
|
||||
restrictions described in section 5.4, "Language restrictions for device
|
||||
functions" in the SYCL 2020 specification. If the function is a non-static
|
||||
member function, ``this`` shall not be used in a potentially evaluated
|
||||
expression.
|
||||
}];
|
||||
}
|
||||
|
||||
|
||||
@ -13352,19 +13352,24 @@ def warn_sycl_external_missing_on_first_decl : Warning<
|
||||
// SYCL kernel entry point diagnostics
|
||||
def err_sycl_entry_point_invalid : Error<
|
||||
"the %0 attribute cannot be applied to a %enum_select<InvalidSKEPReason>{"
|
||||
"%NonStaticMemberFn{non-static member function}|"
|
||||
"%VariadicFn{variadic function}|"
|
||||
"%DeletedFn{deleted function}|"
|
||||
"%DefaultedFn{defaulted function}|"
|
||||
"%Constructor{constructor}|"
|
||||
"%Destructor{destructor}|"
|
||||
"%Coroutine{coroutine}|"
|
||||
"%ConstexprFn{constexpr function}|"
|
||||
"%ConstevalFn{consteval function}|"
|
||||
"%NoreturnFn{function declared with the 'noreturn' attribute}|"
|
||||
"%Coroutine{coroutine}|"
|
||||
"%FunctionTryBlock{function defined with a function try block}"
|
||||
"%FunctionTryBlock{function defined with a function try block}|"
|
||||
"%ExplicitObjectFn{function with an explicit object parameter}|"
|
||||
"}1">;
|
||||
def err_sycl_entry_point_invalid_redeclaration : Error<
|
||||
"the %0 kernel name argument does not match prior"
|
||||
" declaration%diff{: $ vs $|}1,2">;
|
||||
def err_sycl_entry_point_invalid_this : Error<
|
||||
"'this' cannot be%select{| implicitly}0 used in a potentially evaluated"
|
||||
" expression in the body of a function declared with the %1 attribute">;
|
||||
def err_sycl_kernel_name_conflict : Error<
|
||||
"the %0 kernel name argument conflicts with a previous declaration">;
|
||||
def warn_sycl_kernel_name_not_a_class_type : Warning<
|
||||
@ -13380,6 +13385,18 @@ def err_sycl_entry_point_return_type : Error<
|
||||
def err_sycl_entry_point_deduced_return_type : Error<
|
||||
"the %0 attribute only applies to functions with a non-deduced 'void' return"
|
||||
" type">;
|
||||
def note_sycl_runtime_defect : Note<
|
||||
"this indicates a problem with the SYCL runtime header files; please consider"
|
||||
" reporting this to your SYCL runtime provider">;
|
||||
def note_sycl_kernel_launch_lookup_here : Note<
|
||||
"in implicit call to 'sycl_kernel_launch' with template argument %0 required"
|
||||
" here">;
|
||||
def note_sycl_kernel_launch_overload_resolution_here : Note<
|
||||
"in implicit call to 'sycl_kernel_launch' with template argument %0 and"
|
||||
" function arguments %1 required here">;
|
||||
def err_sycl_entry_point_device_use : Error<
|
||||
"function %0 cannot be used in device code because it is declared with the"
|
||||
" %1 attribute">;
|
||||
|
||||
def warn_cuda_maxclusterrank_sm_90 : Warning<
|
||||
"maxclusterrank requires sm_90 or higher, CUDA arch provided: %0, ignoring "
|
||||
|
||||
@ -24,6 +24,7 @@ def CaseStmt : StmtNode<SwitchCase>;
|
||||
def DefaultStmt : StmtNode<SwitchCase>;
|
||||
def CapturedStmt : StmtNode<Stmt>;
|
||||
def SYCLKernelCallStmt : StmtNode<Stmt>;
|
||||
def UnresolvedSYCLKernelCallStmt : StmtNode<Stmt>;
|
||||
|
||||
// Break/continue.
|
||||
def LoopControlStmt : StmtNode<Stmt, 1>;
|
||||
|
||||
@ -245,6 +245,10 @@ public:
|
||||
/// The set of GNU address of label extension "&&label".
|
||||
llvm::SmallVector<AddrLabelExpr *, 4> AddrLabels;
|
||||
|
||||
/// An unresolved identifier lookup expression for an implicit call
|
||||
/// to a SYCL kernel launch function in a dependent context.
|
||||
Expr *SYCLKernelLaunchIdExpr = nullptr;
|
||||
|
||||
public:
|
||||
/// Represents a simple identification of a weak object.
|
||||
///
|
||||
|
||||
@ -1430,7 +1430,8 @@ public:
|
||||
|
||||
/// Diagnostics that are emitted only if we discover that the given function
|
||||
/// must be codegen'ed. Because handling these correctly adds overhead to
|
||||
/// compilation, this is currently only enabled for CUDA compilations.
|
||||
/// compilation, this is currently only used for offload languages like CUDA,
|
||||
/// OpenMP, and SYCL.
|
||||
SemaDiagnosticBuilder::DeferredDiagnosticsType DeviceDeferredDiags;
|
||||
|
||||
/// CurContext - This is the current declaration context of parsing.
|
||||
@ -13275,6 +13276,14 @@ public:
|
||||
|
||||
/// We are performing partial ordering for template template parameters.
|
||||
PartialOrderingTTP,
|
||||
|
||||
/// We are performing name lookup for a function template or variable
|
||||
/// template named 'sycl_kernel_launch'.
|
||||
SYCLKernelLaunchLookup,
|
||||
|
||||
/// We are performing overload resolution for a call to a function
|
||||
/// template or variable template named 'sycl_kernel_launch'.
|
||||
SYCLKernelLaunchOverloadResolution,
|
||||
} Kind;
|
||||
|
||||
/// Whether we're substituting into constraints.
|
||||
@ -13630,6 +13639,20 @@ public:
|
||||
operator=(const SynthesizedFunctionScope &) = delete;
|
||||
};
|
||||
|
||||
/// RAII object to ensure that a code synthesis context is popped on scope
|
||||
/// exit.
|
||||
class ScopedCodeSynthesisContext {
|
||||
Sema &S;
|
||||
|
||||
public:
|
||||
ScopedCodeSynthesisContext(Sema &S, const CodeSynthesisContext &Ctx)
|
||||
: S(S) {
|
||||
S.pushCodeSynthesisContext(Ctx);
|
||||
}
|
||||
|
||||
~ScopedCodeSynthesisContext() { S.popCodeSynthesisContext(); }
|
||||
};
|
||||
|
||||
/// List of active code synthesis contexts.
|
||||
///
|
||||
/// This vector is treated as a stack. As synthesis of one entity requires
|
||||
|
||||
@ -64,9 +64,38 @@ public:
|
||||
void handleKernelAttr(Decl *D, const ParsedAttr &AL);
|
||||
void handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL);
|
||||
|
||||
/// Issues a deferred diagnostic if use of the declaration designated
|
||||
/// by 'ND' is invalid in a device context.
|
||||
void CheckDeviceUseOfDecl(NamedDecl *ND, SourceLocation Loc);
|
||||
|
||||
void CheckSYCLExternalFunctionDecl(FunctionDecl *FD);
|
||||
void CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD);
|
||||
StmtResult BuildSYCLKernelCallStmt(FunctionDecl *FD, CompoundStmt *Body);
|
||||
|
||||
/// Builds an expression for the lookup of a 'sycl_kernel_launch' template
|
||||
/// with 'KernelName' as an explicit template argument. Lookup is performed
|
||||
/// as if from the first statement of the body of 'FD' and thus requires
|
||||
/// searching the scopes that exist at parse time. This function therefore
|
||||
/// requires the current semantic context to be the definition of 'FD'. In a
|
||||
/// dependent context, the returned expression will be an UnresolvedLookupExpr
|
||||
/// or an UnresolvedMemberExpr. In a non-dependent context, the returned
|
||||
/// expression will be a DeclRefExpr or MemberExpr. If lookup fails, a null
|
||||
/// error result is returned. The resulting expression is intended to be
|
||||
/// passed as the 'LaunchIdExpr' argument in a call to either
|
||||
/// BuildSYCLKernelCallStmt() or BuildUnresolvedSYCLKernelCallStmt() after
|
||||
/// the function body has been parsed.
|
||||
ExprResult BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD, QualType KernelName);
|
||||
|
||||
/// Builds a SYCLKernelCallStmt to wrap 'Body' and to be used as the body of
|
||||
/// 'FD'. 'LaunchIdExpr' specifies the lookup result returned by a previous
|
||||
/// call to BuildSYCLKernelLaunchIdExpr().
|
||||
StmtResult BuildSYCLKernelCallStmt(FunctionDecl *FD, CompoundStmt *Body,
|
||||
Expr *LaunchIdExpr);
|
||||
|
||||
/// Builds an UnresolvedSYCLKernelCallStmt to wrap 'Body'. 'LaunchIdExpr'
|
||||
/// specifies the lookup result returned by a previous call to
|
||||
/// BuildSYCLKernelLaunchIdExpr().
|
||||
StmtResult BuildUnresolvedSYCLKernelCallStmt(CompoundStmt *Body,
|
||||
Expr *LaunchIdExpr);
|
||||
};
|
||||
|
||||
} // namespace clang
|
||||
|
||||
@ -1618,6 +1618,9 @@ enum StmtCode {
|
||||
/// A SYCLKernelCallStmt record.
|
||||
STMT_SYCLKERNELCALL,
|
||||
|
||||
/// An UnresolvedSYCLKernelCallStmt record.
|
||||
STMT_UNRESOLVED_SYCL_KERNEL_CALL,
|
||||
|
||||
/// A GCC-style AsmStmt record.
|
||||
STMT_GCCASM,
|
||||
|
||||
|
||||
@ -600,7 +600,7 @@ void StmtPrinter::VisitCapturedStmt(CapturedStmt *Node) {
|
||||
}
|
||||
|
||||
void StmtPrinter::VisitSYCLKernelCallStmt(SYCLKernelCallStmt *Node) {
|
||||
PrintStmt(Node->getOutlinedFunctionDecl()->getBody());
|
||||
PrintStmt(Node->getOriginalStmt());
|
||||
}
|
||||
|
||||
void StmtPrinter::VisitObjCAtTryStmt(ObjCAtTryStmt *Node) {
|
||||
@ -1447,6 +1447,11 @@ void StmtPrinter::VisitSYCLUniqueStableNameExpr(
|
||||
OS << ")";
|
||||
}
|
||||
|
||||
void StmtPrinter::VisitUnresolvedSYCLKernelCallStmt(
|
||||
UnresolvedSYCLKernelCallStmt *Node) {
|
||||
PrintStmt(Node->getOriginalStmt());
|
||||
}
|
||||
|
||||
void StmtPrinter::VisitPredefinedExpr(PredefinedExpr *Node) {
|
||||
OS << PredefinedExpr::getIdentKindName(Node->getIdentKind());
|
||||
}
|
||||
|
||||
@ -1410,6 +1410,11 @@ void StmtProfiler::VisitSYCLUniqueStableNameExpr(
|
||||
VisitType(S->getTypeSourceInfo()->getType());
|
||||
}
|
||||
|
||||
void StmtProfiler::VisitUnresolvedSYCLKernelCallStmt(
|
||||
const UnresolvedSYCLKernelCallStmt *S) {
|
||||
VisitStmt(S);
|
||||
}
|
||||
|
||||
void StmtProfiler::VisitPredefinedExpr(const PredefinedExpr *S) {
|
||||
VisitExpr(S);
|
||||
ID.AddInteger(llvm::to_underlying(S->getIdentKind()));
|
||||
|
||||
@ -19,6 +19,7 @@
|
||||
#include "clang/AST/Attr.h"
|
||||
#include "clang/AST/Expr.h"
|
||||
#include "clang/AST/Stmt.h"
|
||||
#include "clang/AST/StmtSYCL.h"
|
||||
#include "clang/AST/StmtVisitor.h"
|
||||
#include "clang/Basic/Builtins.h"
|
||||
#include "clang/Basic/DiagnosticSema.h"
|
||||
@ -99,6 +100,7 @@ void CodeGenFunction::EmitStmt(const Stmt *S, ArrayRef<const Attr *> Attrs) {
|
||||
case Stmt::SEHExceptStmtClass:
|
||||
case Stmt::SEHFinallyStmtClass:
|
||||
case Stmt::MSDependentExistsStmtClass:
|
||||
case Stmt::UnresolvedSYCLKernelCallStmtClass:
|
||||
llvm_unreachable("invalid statement class to emit generically");
|
||||
case Stmt::NullStmtClass:
|
||||
case Stmt::CompoundStmtClass:
|
||||
@ -543,21 +545,7 @@ bool CodeGenFunction::EmitSimpleStmt(const Stmt *S,
|
||||
EmitSEHLeaveStmt(cast<SEHLeaveStmt>(*S));
|
||||
break;
|
||||
case Stmt::SYCLKernelCallStmtClass:
|
||||
// SYCL kernel call statements are generated as wrappers around the body
|
||||
// of functions declared with the sycl_kernel_entry_point attribute. Such
|
||||
// functions are used to specify how a SYCL kernel (a function object) is
|
||||
// to be invoked; the SYCL kernel call statement contains a transformed
|
||||
// variation of the function body and is used to generate a SYCL kernel
|
||||
// caller function; a function that serves as the device side entry point
|
||||
// used to execute the SYCL kernel. The sycl_kernel_entry_point attributed
|
||||
// function is invoked by host code in order to trigger emission of the
|
||||
// device side SYCL kernel caller function and to generate metadata needed
|
||||
// by SYCL run-time library implementations; the function is otherwise
|
||||
// intended to have no effect. As such, the function body is not evaluated
|
||||
// as part of the invocation during host compilation (and the function
|
||||
// should not be called or emitted during device compilation); the SYCL
|
||||
// kernel call statement is thus handled as a null statement for the
|
||||
// purpose of code generation.
|
||||
EmitSYCLKernelCallStmt(cast<SYCLKernelCallStmt>(*S));
|
||||
break;
|
||||
}
|
||||
return true;
|
||||
|
||||
@ -3675,6 +3675,8 @@ public:
|
||||
LValue EmitCoyieldLValue(const CoyieldExpr *E);
|
||||
RValue EmitCoroutineIntrinsic(const CallExpr *E, unsigned int IID);
|
||||
|
||||
void EmitSYCLKernelCallStmt(const SYCLKernelCallStmt &S);
|
||||
|
||||
void EnterCXXTryStmt(const CXXTryStmt &S, bool IsFnTryBlock = false);
|
||||
void ExitCXXTryStmt(const CXXTryStmt &S, bool IsFnTryBlock = false);
|
||||
|
||||
|
||||
@ -13,10 +13,23 @@
|
||||
|
||||
#include "CodeGenFunction.h"
|
||||
#include "CodeGenModule.h"
|
||||
#include <cassert>
|
||||
|
||||
using namespace clang;
|
||||
using namespace CodeGen;
|
||||
|
||||
void CodeGenFunction::EmitSYCLKernelCallStmt(const SYCLKernelCallStmt &S) {
|
||||
// SYCLKernelCallStmt instances are only injected in the definitions of
|
||||
// functions declared with the sycl_kernel_entry_point attribute. ODR-use of
|
||||
// such a function in code emitted during device compilation should be
|
||||
// diagnosed. Thus, any attempt to emit a SYCLKernelCallStmt during device
|
||||
// compilation indicates a missing diagnostic.
|
||||
assert(!getLangOpts().SYCLIsDevice &&
|
||||
"Attempt to emit a SYCL kernel call statement during device"
|
||||
" compilation");
|
||||
EmitStmt(S.getKernelLaunchStmt());
|
||||
}
|
||||
|
||||
static void SetSYCLKernelAttributes(llvm::Function *Fn, CodeGenFunction &CGF) {
|
||||
// SYCL 2020 device language restrictions require forward progress and
|
||||
// disallow recursion.
|
||||
|
||||
@ -476,6 +476,10 @@ private:
|
||||
return "TypeAliasTemplateInstantiation";
|
||||
case CodeSynthesisContext::PartialOrderingTTP:
|
||||
return "PartialOrderingTTP";
|
||||
case CodeSynthesisContext::SYCLKernelLaunchLookup:
|
||||
return "SYCLKernelLaunchLookup";
|
||||
case CodeSynthesisContext::SYCLKernelLaunchOverloadResolution:
|
||||
return "SYCLKernelLaunchOverloadResolution";
|
||||
}
|
||||
return "";
|
||||
}
|
||||
|
||||
@ -16360,6 +16360,32 @@ Decl *Sema::ActOnStartOfFunctionDef(Scope *FnBodyScope, Decl *D,
|
||||
|
||||
maybeAddDeclWithEffects(FD);
|
||||
|
||||
if (FD && !FD->isInvalidDecl() && FD->hasAttr<SYCLKernelEntryPointAttr>() &&
|
||||
FnBodyScope) {
|
||||
// An implicit call expression is synthesized for functions declared with
|
||||
// the sycl_kernel_entry_point attribute. The call may resolve to a
|
||||
// function template, a member function template, or a call operator
|
||||
// of a variable template depending on the results of unqualified lookup
|
||||
// for 'sycl_kernel_launch' from the beginning of the function body.
|
||||
// Performing that lookup requires the stack of parsing scopes active
|
||||
// when the definition is parsed and is thus done here; the result is
|
||||
// cached in FunctionScopeInfo and used to synthesize the (possibly
|
||||
// unresolved) call expression after the function body has been parsed.
|
||||
const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
|
||||
if (!SKEPAttr->isInvalidAttr()) {
|
||||
ExprResult LaunchIdExpr =
|
||||
SYCL().BuildSYCLKernelLaunchIdExpr(FD, SKEPAttr->getKernelName());
|
||||
// Do not mark 'FD' as invalid if construction of `LaunchIDExpr` produces
|
||||
// an invalid result. Name lookup failure for 'sycl_kernel_launch' is
|
||||
// treated as an error in the definition of 'FD'; treating it as an error
|
||||
// of the declaration would affect overload resolution which would
|
||||
// potentially result in additional errors. If construction of
|
||||
// 'LaunchIDExpr' failed, then 'SYCLKernelLaunchIdExpr' will be assigned
|
||||
// a null pointer value below; that is expected.
|
||||
getCurFunction()->SYCLKernelLaunchIdExpr = LaunchIdExpr.get();
|
||||
}
|
||||
}
|
||||
|
||||
return D;
|
||||
}
|
||||
|
||||
@ -16561,12 +16587,37 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body, bool IsInstantiation,
|
||||
SKEPAttr->setInvalidAttr();
|
||||
}
|
||||
|
||||
if (Body && !FD->isTemplated() && !SKEPAttr->isInvalidAttr()) {
|
||||
StmtResult SR =
|
||||
SYCL().BuildSYCLKernelCallStmt(FD, cast<CompoundStmt>(Body));
|
||||
if (SR.isInvalid())
|
||||
return nullptr;
|
||||
Body = SR.get();
|
||||
// Build an unresolved SYCL kernel call statement for a function template,
|
||||
// validate that a SYCL kernel call statement was instantiated for an
|
||||
// (implicit or explicit) instantiation of a function template, or otherwise
|
||||
// build a (resolved) SYCL kernel call statement for a non-templated
|
||||
// function or an explicit specialization.
|
||||
if (Body && !SKEPAttr->isInvalidAttr()) {
|
||||
StmtResult SR;
|
||||
if (FD->isTemplateInstantiation()) {
|
||||
// The function body should already be a SYCLKernelCallStmt in this
|
||||
// case, but might not be if there were previous errors.
|
||||
SR = Body;
|
||||
} else if (!getCurFunction()->SYCLKernelLaunchIdExpr) {
|
||||
// If name lookup for a template named sycl_kernel_launch failed
|
||||
// earlier, don't try to build a SYCL kernel call statement as that
|
||||
// would cause additional errors to be issued; just proceed with the
|
||||
// original function body.
|
||||
SR = Body;
|
||||
} else if (FD->isTemplated()) {
|
||||
SR = SYCL().BuildUnresolvedSYCLKernelCallStmt(
|
||||
cast<CompoundStmt>(Body), getCurFunction()->SYCLKernelLaunchIdExpr);
|
||||
} else {
|
||||
SR = SYCL().BuildSYCLKernelCallStmt(
|
||||
FD, cast<CompoundStmt>(Body),
|
||||
getCurFunction()->SYCLKernelLaunchIdExpr);
|
||||
}
|
||||
// If construction of the replacement body fails, just continue with the
|
||||
// original function body. An early error return here is not valid; the
|
||||
// current declaration context and function scopes must be popped before
|
||||
// returning.
|
||||
if (SR.isUsable())
|
||||
Body = SR.get();
|
||||
}
|
||||
}
|
||||
|
||||
@ -21037,7 +21088,9 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD,
|
||||
|
||||
// SYCL functions can be template, so we check if they have appropriate
|
||||
// attribute prior to checking if it is a template.
|
||||
if (LangOpts.SYCLIsDevice && FD->hasAttr<SYCLKernelAttr>())
|
||||
if (LangOpts.SYCLIsDevice && (FD->hasAttr<SYCLKernelAttr>() ||
|
||||
FD->hasAttr<SYCLKernelEntryPointAttr>() ||
|
||||
FD->hasAttr<SYCLExternalAttr>()))
|
||||
return FunctionEmissionStatus::Emitted;
|
||||
|
||||
// Templates are emitted when they're instantiated.
|
||||
|
||||
@ -15,6 +15,7 @@
|
||||
#include "clang/AST/Expr.h"
|
||||
#include "clang/AST/ExprCXX.h"
|
||||
#include "clang/AST/StmtObjC.h"
|
||||
#include "clang/AST/StmtSYCL.h"
|
||||
#include "clang/AST/TypeLoc.h"
|
||||
#include "clang/Basic/Diagnostic.h"
|
||||
#include "clang/Basic/SourceManager.h"
|
||||
@ -1250,6 +1251,18 @@ CanThrowResult Sema::canThrow(const Stmt *S) {
|
||||
return CT;
|
||||
}
|
||||
|
||||
case Stmt::SYCLKernelCallStmtClass: {
|
||||
auto *SKCS = cast<SYCLKernelCallStmt>(S);
|
||||
if (getLangOpts().SYCLIsDevice)
|
||||
return canSubStmtsThrow(*this,
|
||||
SKCS->getOutlinedFunctionDecl()->getBody());
|
||||
assert(getLangOpts().SYCLIsHost);
|
||||
return canSubStmtsThrow(*this, SKCS->getKernelLaunchStmt());
|
||||
}
|
||||
|
||||
case Stmt::UnresolvedSYCLKernelCallStmtClass:
|
||||
return CT_Dependent;
|
||||
|
||||
// ObjC message sends are like function calls, but never have exception
|
||||
// specs.
|
||||
case Expr::ObjCMessageExprClass:
|
||||
@ -1433,7 +1446,6 @@ CanThrowResult Sema::canThrow(const Stmt *S) {
|
||||
case Stmt::AttributedStmtClass:
|
||||
case Stmt::BreakStmtClass:
|
||||
case Stmt::CapturedStmtClass:
|
||||
case Stmt::SYCLKernelCallStmtClass:
|
||||
case Stmt::CaseStmtClass:
|
||||
case Stmt::CompoundStmtClass:
|
||||
case Stmt::ContinueStmtClass:
|
||||
|
||||
@ -406,6 +406,9 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
|
||||
targetDiag(*Locs.begin(), diag::err_thread_unsupported);
|
||||
}
|
||||
|
||||
if (LangOpts.SYCLIsDevice && isa<FunctionDecl>(D))
|
||||
SYCL().CheckDeviceUseOfDecl(D, Loc);
|
||||
|
||||
return false;
|
||||
}
|
||||
|
||||
|
||||
@ -30,15 +30,25 @@ SemaSYCL::SemaSYCL(Sema &S) : SemaBase(S) {}
|
||||
Sema::SemaDiagnosticBuilder SemaSYCL::DiagIfDeviceCode(SourceLocation Loc,
|
||||
unsigned DiagID) {
|
||||
assert(getLangOpts().SYCLIsDevice &&
|
||||
"Should only be called during SYCL compilation");
|
||||
FunctionDecl *FD = dyn_cast<FunctionDecl>(SemaRef.getCurLexicalContext());
|
||||
SemaDiagnosticBuilder::Kind DiagKind = [this, FD] {
|
||||
if (!FD)
|
||||
return SemaDiagnosticBuilder::K_Nop;
|
||||
if (SemaRef.getEmissionStatus(FD) == Sema::FunctionEmissionStatus::Emitted)
|
||||
return SemaDiagnosticBuilder::K_ImmediateWithCallStack;
|
||||
return SemaDiagnosticBuilder::K_Deferred;
|
||||
}();
|
||||
"Device diagnostics Should only be issued during device compilation");
|
||||
SemaDiagnosticBuilder::Kind DiagKind = SemaDiagnosticBuilder::K_Nop;
|
||||
FunctionDecl *FD = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
|
||||
if (FD) {
|
||||
Sema::FunctionEmissionStatus FES = SemaRef.getEmissionStatus(FD);
|
||||
switch (FES) {
|
||||
case Sema::FunctionEmissionStatus::Emitted:
|
||||
DiagKind = SemaDiagnosticBuilder::K_ImmediateWithCallStack;
|
||||
break;
|
||||
case Sema::FunctionEmissionStatus::Unknown:
|
||||
case Sema::FunctionEmissionStatus::TemplateDiscarded:
|
||||
DiagKind = SemaDiagnosticBuilder::K_Deferred;
|
||||
break;
|
||||
case Sema::FunctionEmissionStatus::OMPDiscarded:
|
||||
llvm_unreachable("OMPDiscarded unexpected in SYCL device compilation");
|
||||
case Sema::FunctionEmissionStatus::CUDADiscarded:
|
||||
llvm_unreachable("CUDADiscarded unexpected in SYCL device compilation");
|
||||
}
|
||||
}
|
||||
return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, FD, SemaRef);
|
||||
}
|
||||
|
||||
@ -211,6 +221,23 @@ void SemaSYCL::handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL) {
|
||||
SYCLKernelEntryPointAttr(SemaRef.Context, AL, TSI));
|
||||
}
|
||||
|
||||
void SemaSYCL::CheckDeviceUseOfDecl(NamedDecl *ND, SourceLocation Loc) {
|
||||
assert(getLangOpts().SYCLIsDevice &&
|
||||
"Should only be called during SYCL device compilation");
|
||||
|
||||
// Function declarations with the sycl_kernel_entry_point attribute cannot
|
||||
// be ODR-used in a potentially evaluated context.
|
||||
if (FunctionDecl *FD = dyn_cast<FunctionDecl>(ND)) {
|
||||
if (const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>()) {
|
||||
if (SemaRef.currentEvaluationContext().isPotentiallyEvaluated()) {
|
||||
DiagIfDeviceCode(Loc, diag::err_sycl_entry_point_device_use)
|
||||
<< FD << SKEPAttr;
|
||||
DiagIfDeviceCode(SKEPAttr->getLocation(), diag::note_attribute) << FD;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// Given a potentially qualified type, SourceLocationForUserDeclaredType()
|
||||
// returns the source location of the canonical declaration of the unqualified
|
||||
// desugared user declared type, if any. For non-user declared types, an
|
||||
@ -315,10 +342,20 @@ void SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) {
|
||||
}
|
||||
}
|
||||
|
||||
if (isa<CXXConstructorDecl>(FD)) {
|
||||
Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
|
||||
<< SKEPAttr << diag::InvalidSKEPReason::Constructor;
|
||||
SKEPAttr->setInvalidAttr();
|
||||
}
|
||||
if (isa<CXXDestructorDecl>(FD)) {
|
||||
Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
|
||||
<< SKEPAttr << diag::InvalidSKEPReason::Destructor;
|
||||
SKEPAttr->setInvalidAttr();
|
||||
}
|
||||
if (const auto *MD = dyn_cast<CXXMethodDecl>(FD)) {
|
||||
if (!MD->isStatic()) {
|
||||
if (MD->isExplicitObjectMemberFunction()) {
|
||||
Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
|
||||
<< SKEPAttr << diag::InvalidSKEPReason::NonStaticMemberFn;
|
||||
<< SKEPAttr << diag::InvalidSKEPReason::ExplicitObjectFn;
|
||||
SKEPAttr->setInvalidAttr();
|
||||
}
|
||||
}
|
||||
@ -387,8 +424,165 @@ void SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) {
|
||||
}
|
||||
}
|
||||
|
||||
ExprResult SemaSYCL::BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD,
|
||||
QualType KNT) {
|
||||
// The current context must be the function definition context to ensure
|
||||
// that name lookup is performed within the correct scope.
|
||||
assert(SemaRef.CurContext == FD && "The current declaration context does not "
|
||||
"match the requested function context");
|
||||
|
||||
// An appropriate source location is required to emit diagnostics if
|
||||
// lookup fails to produce an overload set. The desired location is the
|
||||
// start of the function body, but that is not yet available since the
|
||||
// body of the function has not yet been set when this function is called.
|
||||
// The general location of the function is used instead.
|
||||
SourceLocation Loc = FD->getLocation();
|
||||
|
||||
ASTContext &Ctx = SemaRef.getASTContext();
|
||||
IdentifierInfo &SYCLKernelLaunchID =
|
||||
Ctx.Idents.get("sycl_kernel_launch", tok::TokenKind::identifier);
|
||||
|
||||
// Establish a code synthesis context for the implicit name lookup of
|
||||
// a template named 'sycl_kernel_launch'. In the event of an error, this
|
||||
// ensures an appropriate diagnostic note is issued to explain why the
|
||||
// lookup was performed.
|
||||
Sema::CodeSynthesisContext CSC;
|
||||
CSC.Kind = Sema::CodeSynthesisContext::SYCLKernelLaunchLookup;
|
||||
CSC.Entity = FD;
|
||||
Sema::ScopedCodeSynthesisContext ScopedCSC(SemaRef, CSC);
|
||||
|
||||
// Perform ordinary name lookup for a function or variable template that
|
||||
// accepts a single type template argument.
|
||||
LookupResult Result(SemaRef, &SYCLKernelLaunchID, Loc,
|
||||
Sema::LookupOrdinaryName);
|
||||
CXXScopeSpec EmptySS;
|
||||
if (SemaRef.LookupTemplateName(Result, SemaRef.getCurScope(), EmptySS,
|
||||
/*ObjectType*/ QualType(),
|
||||
/*EnteringContext*/ false,
|
||||
Sema::TemplateNameIsRequired))
|
||||
return ExprError();
|
||||
if (Result.isAmbiguous())
|
||||
return ExprError();
|
||||
|
||||
TemplateArgumentListInfo TALI{Loc, Loc};
|
||||
TemplateArgument KNTA = TemplateArgument(KNT);
|
||||
TemplateArgumentLoc TAL =
|
||||
SemaRef.getTrivialTemplateArgumentLoc(KNTA, QualType(), Loc);
|
||||
TALI.addArgument(TAL);
|
||||
|
||||
ExprResult IdExpr;
|
||||
if (SemaRef.isPotentialImplicitMemberAccess(EmptySS, Result,
|
||||
/*IsAddressOfOperand*/ false)) {
|
||||
// The lookup result allows for a possible implicit member access that
|
||||
// would require an implicit or explicit 'this' argument.
|
||||
IdExpr = SemaRef.BuildPossibleImplicitMemberExpr(
|
||||
EmptySS, SourceLocation(), Result, &TALI, SemaRef.getCurScope());
|
||||
} else {
|
||||
IdExpr = SemaRef.BuildTemplateIdExpr(EmptySS, SourceLocation(), Result,
|
||||
/*RequiresADL*/ true, &TALI);
|
||||
}
|
||||
|
||||
// The resulting expression may be invalid if, for example, 'FD' is a
|
||||
// non-static member function and sycl_kernel_launch lookup selects a
|
||||
// member function (which would require a 'this' argument which is
|
||||
// not available).
|
||||
if (IdExpr.isInvalid())
|
||||
return ExprError();
|
||||
|
||||
return IdExpr;
|
||||
}
|
||||
|
||||
namespace {
|
||||
|
||||
// Constructs the arguments to be passed for the SYCL kernel launch call.
|
||||
// The first argument is a string literal that contains the SYCL kernel
|
||||
// name. The remaining arguments are the parameters of 'FD' passed as
|
||||
// move-elligible xvalues. Returns true on error and false otherwise.
|
||||
bool BuildSYCLKernelLaunchCallArgs(Sema &SemaRef, FunctionDecl *FD,
|
||||
const SYCLKernelInfo *SKI,
|
||||
SmallVectorImpl<Expr *> &Args,
|
||||
SourceLocation Loc) {
|
||||
// The current context must be the function definition context to ensure
|
||||
// that parameter references occur within the correct scope.
|
||||
assert(SemaRef.CurContext == FD && "The current declaration context does not "
|
||||
"match the requested function context");
|
||||
|
||||
// Prepare a string literal that contains the kernel name.
|
||||
ASTContext &Ctx = SemaRef.getASTContext();
|
||||
const std::string &KernelName = SKI->GetKernelName();
|
||||
QualType KernelNameCharTy = Ctx.CharTy.withConst();
|
||||
llvm::APInt KernelNameSize(Ctx.getTypeSize(Ctx.getSizeType()),
|
||||
KernelName.size() + 1);
|
||||
QualType KernelNameArrayTy = Ctx.getConstantArrayType(
|
||||
KernelNameCharTy, KernelNameSize, nullptr, ArraySizeModifier::Normal, 0);
|
||||
Expr *KernelNameExpr =
|
||||
StringLiteral::Create(Ctx, KernelName, StringLiteralKind::Ordinary,
|
||||
/*Pascal*/ false, KernelNameArrayTy, Loc);
|
||||
Args.push_back(KernelNameExpr);
|
||||
|
||||
// Forward all parameters of 'FD' to the SYCL kernel launch function as if
|
||||
// by std::move().
|
||||
for (ParmVarDecl *PVD : FD->parameters()) {
|
||||
QualType ParamType = PVD->getOriginalType().getNonReferenceType();
|
||||
ExprResult E = SemaRef.BuildDeclRefExpr(PVD, ParamType, VK_LValue, Loc);
|
||||
if (E.isInvalid())
|
||||
return true;
|
||||
if (!PVD->getType()->isLValueReferenceType())
|
||||
E = ImplicitCastExpr::Create(SemaRef.Context, E.get()->getType(), CK_NoOp,
|
||||
E.get(), nullptr, VK_XValue,
|
||||
FPOptionsOverride());
|
||||
if (E.isInvalid())
|
||||
return true;
|
||||
Args.push_back(E.get());
|
||||
}
|
||||
|
||||
return false;
|
||||
}
|
||||
|
||||
// Constructs the SYCL kernel launch call.
|
||||
StmtResult BuildSYCLKernelLaunchCallStmt(Sema &SemaRef, FunctionDecl *FD,
|
||||
const SYCLKernelInfo *SKI,
|
||||
Expr *IdExpr, SourceLocation Loc) {
|
||||
SmallVector<Stmt *> Stmts;
|
||||
// IdExpr may be null if name lookup failed.
|
||||
if (IdExpr) {
|
||||
llvm::SmallVector<Expr *, 12> Args;
|
||||
|
||||
// Establish a code synthesis context for construction of the arguments
|
||||
// for the implicit call to 'sycl_kernel_launch'.
|
||||
{
|
||||
Sema::CodeSynthesisContext CSC;
|
||||
CSC.Kind = Sema::CodeSynthesisContext::SYCLKernelLaunchLookup;
|
||||
CSC.Entity = FD;
|
||||
Sema::ScopedCodeSynthesisContext ScopedCSC(SemaRef, CSC);
|
||||
|
||||
if (BuildSYCLKernelLaunchCallArgs(SemaRef, FD, SKI, Args, Loc))
|
||||
return StmtError();
|
||||
}
|
||||
|
||||
// Establish a code synthesis context for the implicit call to
|
||||
// 'sycl_kernel_launch'.
|
||||
{
|
||||
Sema::CodeSynthesisContext CSC;
|
||||
CSC.Kind = Sema::CodeSynthesisContext::SYCLKernelLaunchOverloadResolution;
|
||||
CSC.Entity = FD;
|
||||
CSC.CallArgs = Args.data();
|
||||
CSC.NumCallArgs = Args.size();
|
||||
Sema::ScopedCodeSynthesisContext ScopedCSC(SemaRef, CSC);
|
||||
|
||||
ExprResult LaunchResult =
|
||||
SemaRef.BuildCallExpr(SemaRef.getCurScope(), IdExpr, Loc, Args, Loc);
|
||||
if (LaunchResult.isInvalid())
|
||||
return StmtError();
|
||||
|
||||
Stmts.push_back(SemaRef.MaybeCreateExprWithCleanups(LaunchResult).get());
|
||||
}
|
||||
}
|
||||
|
||||
return CompoundStmt::Create(SemaRef.getASTContext(), Stmts,
|
||||
FPOptionsOverride(), Loc, Loc);
|
||||
}
|
||||
|
||||
// The body of a function declared with the [[sycl_kernel_entry_point]]
|
||||
// attribute is cloned and transformed to substitute references to the original
|
||||
// function parameters with references to replacement variables that stand in
|
||||
@ -399,9 +593,10 @@ class OutlinedFunctionDeclBodyInstantiator
|
||||
public:
|
||||
using ParmDeclMap = llvm::DenseMap<ParmVarDecl *, VarDecl *>;
|
||||
|
||||
OutlinedFunctionDeclBodyInstantiator(Sema &S, ParmDeclMap &M)
|
||||
OutlinedFunctionDeclBodyInstantiator(Sema &S, ParmDeclMap &M,
|
||||
FunctionDecl *FD)
|
||||
: TreeTransform<OutlinedFunctionDeclBodyInstantiator>(S), SemaRef(S),
|
||||
MapRef(M) {}
|
||||
MapRef(M), FD(FD) {}
|
||||
|
||||
// A new set of AST nodes is always required.
|
||||
bool AlwaysRebuild() { return true; }
|
||||
@ -427,18 +622,62 @@ public:
|
||||
return DRE;
|
||||
}
|
||||
|
||||
// Diagnose CXXThisExpr in a potentially evaluated expression.
|
||||
ExprResult TransformCXXThisExpr(CXXThisExpr *CTE) {
|
||||
if (SemaRef.currentEvaluationContext().isPotentiallyEvaluated()) {
|
||||
SemaRef.Diag(CTE->getExprLoc(), diag::err_sycl_entry_point_invalid_this)
|
||||
<< (CTE->isImplicitCXXThis() ? /* implicit */ 1 : /* empty */ 0)
|
||||
<< FD->getAttr<SYCLKernelEntryPointAttr>();
|
||||
}
|
||||
return CTE;
|
||||
}
|
||||
|
||||
private:
|
||||
Sema &SemaRef;
|
||||
ParmDeclMap &MapRef;
|
||||
FunctionDecl *FD;
|
||||
};
|
||||
|
||||
OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
|
||||
FunctionDecl *FD,
|
||||
CompoundStmt *Body) {
|
||||
using ParmDeclMap = OutlinedFunctionDeclBodyInstantiator::ParmDeclMap;
|
||||
ParmDeclMap ParmMap;
|
||||
|
||||
OutlinedFunctionDecl *OFD = OutlinedFunctionDecl::Create(
|
||||
SemaRef.getASTContext(), FD, FD->getNumParams());
|
||||
unsigned i = 0;
|
||||
for (ParmVarDecl *PVD : FD->parameters()) {
|
||||
ImplicitParamDecl *IPD = ImplicitParamDecl::Create(
|
||||
SemaRef.getASTContext(), OFD, SourceLocation(), PVD->getIdentifier(),
|
||||
PVD->getType(), ImplicitParamKind::Other);
|
||||
OFD->setParam(i, IPD);
|
||||
ParmMap[PVD] = IPD;
|
||||
++i;
|
||||
}
|
||||
|
||||
OutlinedFunctionDeclBodyInstantiator OFDBodyInstantiator(SemaRef, ParmMap,
|
||||
FD);
|
||||
Stmt *OFDBody = OFDBodyInstantiator.TransformStmt(Body).get();
|
||||
OFD->setBody(OFDBody);
|
||||
OFD->setNothrow();
|
||||
|
||||
return OFD;
|
||||
}
|
||||
|
||||
} // unnamed namespace
|
||||
|
||||
StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD,
|
||||
CompoundStmt *Body) {
|
||||
CompoundStmt *Body,
|
||||
Expr *LaunchIdExpr) {
|
||||
assert(!FD->isInvalidDecl());
|
||||
assert(!FD->isTemplated());
|
||||
assert(FD->hasPrototype());
|
||||
// The current context must be the function definition context to ensure
|
||||
// that name lookup and parameter and local variable creation are performed
|
||||
// within the correct scope.
|
||||
assert(SemaRef.CurContext == FD && "The current declaration context does not "
|
||||
"match the requested function context");
|
||||
|
||||
const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
|
||||
assert(SKEPAttr && "Missing sycl_kernel_entry_point attribute");
|
||||
@ -451,29 +690,28 @@ StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD,
|
||||
getASTContext().getSYCLKernelInfo(SKEPAttr->getKernelName());
|
||||
assert(declaresSameEntity(SKI.getKernelEntryPointDecl(), FD) &&
|
||||
"SYCL kernel name conflict");
|
||||
(void)SKI;
|
||||
|
||||
using ParmDeclMap = OutlinedFunctionDeclBodyInstantiator::ParmDeclMap;
|
||||
ParmDeclMap ParmMap;
|
||||
|
||||
assert(SemaRef.CurContext == FD);
|
||||
// Build the outline of the synthesized device entry point function.
|
||||
OutlinedFunctionDecl *OFD =
|
||||
OutlinedFunctionDecl::Create(getASTContext(), FD, FD->getNumParams());
|
||||
unsigned i = 0;
|
||||
for (ParmVarDecl *PVD : FD->parameters()) {
|
||||
ImplicitParamDecl *IPD = ImplicitParamDecl::Create(
|
||||
getASTContext(), OFD, SourceLocation(), PVD->getIdentifier(),
|
||||
PVD->getType(), ImplicitParamKind::Other);
|
||||
OFD->setParam(i, IPD);
|
||||
ParmMap[PVD] = IPD;
|
||||
++i;
|
||||
}
|
||||
BuildSYCLKernelEntryPointOutline(SemaRef, FD, Body);
|
||||
assert(OFD);
|
||||
|
||||
OutlinedFunctionDeclBodyInstantiator OFDBodyInstantiator(SemaRef, ParmMap);
|
||||
Stmt *OFDBody = OFDBodyInstantiator.TransformStmt(Body).get();
|
||||
OFD->setBody(OFDBody);
|
||||
OFD->setNothrow();
|
||||
Stmt *NewBody = new (getASTContext()) SYCLKernelCallStmt(Body, OFD);
|
||||
// Build the host kernel launch statement. An appropriate source location
|
||||
// is required to emit diagnostics.
|
||||
SourceLocation Loc = Body->getLBracLoc();
|
||||
StmtResult LaunchResult =
|
||||
BuildSYCLKernelLaunchCallStmt(SemaRef, FD, &SKI, LaunchIdExpr, Loc);
|
||||
if (LaunchResult.isInvalid())
|
||||
return StmtError();
|
||||
|
||||
Stmt *NewBody =
|
||||
new (getASTContext()) SYCLKernelCallStmt(Body, LaunchResult.get(), OFD);
|
||||
|
||||
return NewBody;
|
||||
}
|
||||
|
||||
StmtResult SemaSYCL::BuildUnresolvedSYCLKernelCallStmt(CompoundStmt *Body,
|
||||
Expr *LaunchIdExpr) {
|
||||
return UnresolvedSYCLKernelCallStmt::Create(SemaRef.getASTContext(), Body,
|
||||
LaunchIdExpr);
|
||||
}
|
||||
|
||||
@ -10,7 +10,6 @@
|
||||
//===----------------------------------------------------------------------===/
|
||||
|
||||
#include "TreeTransform.h"
|
||||
#include "clang/AST/ASTConcept.h"
|
||||
#include "clang/AST/ASTConsumer.h"
|
||||
#include "clang/AST/ASTContext.h"
|
||||
#include "clang/AST/ASTLambda.h"
|
||||
@ -593,6 +592,8 @@ bool Sema::CodeSynthesisContext::isInstantiationRecord() const {
|
||||
case BuildingDeductionGuides:
|
||||
case TypeAliasTemplateInstantiation:
|
||||
case PartialOrderingTTP:
|
||||
case SYCLKernelLaunchLookup:
|
||||
case SYCLKernelLaunchOverloadResolution:
|
||||
return false;
|
||||
|
||||
// This function should never be called when Kind's value is Memoization.
|
||||
@ -898,6 +899,26 @@ static std::string convertCallArgsToString(Sema &S,
|
||||
return Result;
|
||||
}
|
||||
|
||||
static std::string
|
||||
convertCallArgsValueCategoryAndTypeToString(Sema &S,
|
||||
llvm::ArrayRef<const Expr *> Args) {
|
||||
std::string Result;
|
||||
llvm::raw_string_ostream OS(Result);
|
||||
llvm::ListSeparator Comma;
|
||||
OS << "(";
|
||||
for (const Expr *Arg : Args) {
|
||||
ExprValueKind EVK = Arg->getValueKind();
|
||||
const char *ValueCategory =
|
||||
(EVK == VK_LValue ? "lvalue"
|
||||
: (EVK == VK_XValue ? "xvalue" : "prvalue"));
|
||||
OS << Comma << ValueCategory << " of type '";
|
||||
Arg->getType().print(OS, S.getPrintingPolicy());
|
||||
OS << "'";
|
||||
}
|
||||
OS << ")";
|
||||
return Result;
|
||||
}
|
||||
|
||||
void Sema::PrintInstantiationStack(InstantiationContextDiagFuncRef DiagFunc) {
|
||||
// Determine which template instantiations to skip, if any.
|
||||
unsigned SkipStart = CodeSynthesisContexts.size(), SkipEnd = SkipStart;
|
||||
@ -1260,6 +1281,33 @@ void Sema::PrintInstantiationStack(InstantiationContextDiagFuncRef DiagFunc) {
|
||||
<< /*isTemplateTemplateParam=*/true
|
||||
<< Active->InstantiationRange);
|
||||
break;
|
||||
case CodeSynthesisContext::SYCLKernelLaunchLookup: {
|
||||
const auto *SKEPAttr =
|
||||
Active->Entity->getAttr<SYCLKernelEntryPointAttr>();
|
||||
assert(SKEPAttr && "Missing sycl_kernel_entry_point attribute");
|
||||
assert(!SKEPAttr->isInvalidAttr() &&
|
||||
"sycl_kernel_entry_point attribute is invalid");
|
||||
DiagFunc(SKEPAttr->getLocation(), PDiag(diag::note_sycl_runtime_defect));
|
||||
DiagFunc(SKEPAttr->getLocation(),
|
||||
PDiag(diag::note_sycl_kernel_launch_lookup_here)
|
||||
<< SKEPAttr->getKernelName());
|
||||
break;
|
||||
}
|
||||
case CodeSynthesisContext::SYCLKernelLaunchOverloadResolution: {
|
||||
const auto *SKEPAttr =
|
||||
Active->Entity->getAttr<SYCLKernelEntryPointAttr>();
|
||||
assert(SKEPAttr && "Missing sycl_kernel_entry_point attribute");
|
||||
assert(!SKEPAttr->isInvalidAttr() &&
|
||||
"sycl_kernel_entry_point attribute is invalid");
|
||||
DiagFunc(SKEPAttr->getLocation(), PDiag(diag::note_sycl_runtime_defect));
|
||||
DiagFunc(SKEPAttr->getLocation(),
|
||||
PDiag(diag::note_sycl_kernel_launch_overload_resolution_here)
|
||||
<< SKEPAttr->getKernelName()
|
||||
<< convertCallArgsValueCategoryAndTypeToString(
|
||||
*this, llvm::ArrayRef(Active->CallArgs,
|
||||
Active->NumCallArgs)));
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@ -13076,6 +13076,31 @@ ExprResult TreeTransform<Derived>::TransformSYCLUniqueStableNameExpr(
|
||||
E->getLocation(), E->getLParenLocation(), E->getRParenLocation(), NewT);
|
||||
}
|
||||
|
||||
template <typename Derived>
|
||||
StmtResult TreeTransform<Derived>::TransformUnresolvedSYCLKernelCallStmt(
|
||||
UnresolvedSYCLKernelCallStmt *S) {
|
||||
auto *FD = cast<FunctionDecl>(SemaRef.CurContext);
|
||||
const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
|
||||
if (!SKEPAttr || SKEPAttr->isInvalidAttr())
|
||||
return StmtError();
|
||||
|
||||
ExprResult IdExpr = getDerived().TransformExpr(S->getKernelLaunchIdExpr());
|
||||
if (IdExpr.isInvalid())
|
||||
return StmtError();
|
||||
|
||||
StmtResult Body = getDerived().TransformStmt(S->getOriginalStmt());
|
||||
if (Body.isInvalid())
|
||||
return StmtError();
|
||||
|
||||
StmtResult SR = SemaRef.SYCL().BuildSYCLKernelCallStmt(
|
||||
cast<FunctionDecl>(SemaRef.CurContext), cast<CompoundStmt>(Body.get()),
|
||||
IdExpr.get());
|
||||
if (SR.isInvalid())
|
||||
return StmtError();
|
||||
|
||||
return SR;
|
||||
}
|
||||
|
||||
template <typename Derived>
|
||||
ExprResult TreeTransform<Derived>::TransformCXXReflectExpr(CXXReflectExpr *E) {
|
||||
// TODO(reflection): Implement its transform
|
||||
|
||||
@ -543,6 +543,7 @@ void ASTStmtReader::VisitCXXReflectExpr(CXXReflectExpr *E) {
|
||||
void ASTStmtReader::VisitSYCLKernelCallStmt(SYCLKernelCallStmt *S) {
|
||||
VisitStmt(S);
|
||||
S->setOriginalStmt(cast<CompoundStmt>(Record.readSubStmt()));
|
||||
S->setKernelLaunchStmt(cast<Stmt>(Record.readSubStmt()));
|
||||
S->setOutlinedFunctionDecl(readDeclAs<OutlinedFunctionDecl>());
|
||||
}
|
||||
|
||||
@ -608,6 +609,14 @@ void ASTStmtReader::VisitSYCLUniqueStableNameExpr(SYCLUniqueStableNameExpr *E) {
|
||||
E->setTypeSourceInfo(Record.readTypeSourceInfo());
|
||||
}
|
||||
|
||||
void ASTStmtReader::VisitUnresolvedSYCLKernelCallStmt(
|
||||
UnresolvedSYCLKernelCallStmt *S) {
|
||||
VisitStmt(S);
|
||||
|
||||
S->setOriginalStmt(cast<CompoundStmt>(Record.readSubStmt()));
|
||||
S->setKernelLaunchIdExpr(Record.readExpr());
|
||||
}
|
||||
|
||||
void ASTStmtReader::VisitPredefinedExpr(PredefinedExpr *E) {
|
||||
VisitExpr(E);
|
||||
bool HasFunctionName = Record.readInt();
|
||||
@ -3212,6 +3221,10 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
|
||||
S = SYCLUniqueStableNameExpr::CreateEmpty(Context);
|
||||
break;
|
||||
|
||||
case STMT_UNRESOLVED_SYCL_KERNEL_CALL:
|
||||
S = UnresolvedSYCLKernelCallStmt::CreateEmpty(Context);
|
||||
break;
|
||||
|
||||
case EXPR_OPENACC_ASTERISK_SIZE:
|
||||
S = OpenACCAsteriskSizeExpr::CreateEmpty(Context);
|
||||
break;
|
||||
|
||||
@ -637,6 +637,7 @@ void ASTStmtWriter::VisitCapturedStmt(CapturedStmt *S) {
|
||||
void ASTStmtWriter::VisitSYCLKernelCallStmt(SYCLKernelCallStmt *S) {
|
||||
VisitStmt(S);
|
||||
Record.AddStmt(S->getOriginalStmt());
|
||||
Record.AddStmt(S->getKernelLaunchStmt());
|
||||
Record.AddDeclRef(S->getOutlinedFunctionDecl());
|
||||
|
||||
Code = serialization::STMT_SYCLKERNELCALL;
|
||||
@ -695,6 +696,16 @@ void ASTStmtWriter::VisitSYCLUniqueStableNameExpr(SYCLUniqueStableNameExpr *E) {
|
||||
Code = serialization::EXPR_SYCL_UNIQUE_STABLE_NAME;
|
||||
}
|
||||
|
||||
void ASTStmtWriter::VisitUnresolvedSYCLKernelCallStmt(
|
||||
UnresolvedSYCLKernelCallStmt *S) {
|
||||
VisitStmt(S);
|
||||
|
||||
Record.AddStmt(S->getOriginalStmt());
|
||||
Record.AddStmt(S->getKernelLaunchIdExpr());
|
||||
|
||||
Code = serialization::STMT_UNRESOLVED_SYCL_KERNEL_CALL;
|
||||
}
|
||||
|
||||
void ASTStmtWriter::VisitPredefinedExpr(PredefinedExpr *E) {
|
||||
VisitExpr(E);
|
||||
|
||||
|
||||
@ -1825,6 +1825,7 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred,
|
||||
case Stmt::OMPTargetParallelGenericLoopDirectiveClass:
|
||||
case Stmt::CapturedStmtClass:
|
||||
case Stmt::SYCLKernelCallStmtClass:
|
||||
case Stmt::UnresolvedSYCLKernelCallStmtClass:
|
||||
case Stmt::OpenACCComputeConstructClass:
|
||||
case Stmt::OpenACCLoopConstructClass:
|
||||
case Stmt::OpenACCCombinedConstructClass:
|
||||
|
||||
@ -34,6 +34,8 @@ template<int> struct K {
|
||||
void operator()(Ts...) const {}
|
||||
};
|
||||
|
||||
template<typename KernelName, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...) {}
|
||||
|
||||
[[clang::sycl_kernel_entry_point(KN<1>)]]
|
||||
void skep1() {
|
||||
@ -41,6 +43,12 @@ void skep1() {
|
||||
// CHECK: |-FunctionDecl {{.*}} skep1 'void ()'
|
||||
// CHECK-NEXT: | |-SYCLKernelCallStmt {{.*}}
|
||||
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | | `-CallExpr {{.*}}
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(const char *)' <FunctionToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *)' lvalue Function {{.*}} 'sycl_kernel_launch' {{.*}}
|
||||
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
|
||||
// CHECK-NEXT: | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi1EE"
|
||||
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
|
||||
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<1>
|
||||
@ -57,9 +65,10 @@ void skep2<KN<2>>(K<2>);
|
||||
// CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} KT
|
||||
// CHECK-NEXT: | |-FunctionDecl {{.*}} skep2 'void (KT)'
|
||||
// CHECK-NEXT: | | |-ParmVarDecl {{.*}} k 'KT'
|
||||
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | | `-CallExpr {{.*}} '<dependent type>'
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
|
||||
// CHECK-NEXT: | | |-UnresolvedSYCLKernelCallStmt {{.*}}
|
||||
// CHECK-NEXT: | | | `-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | | `-CallExpr {{.*}} '<dependent type>'
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
|
||||
// CHECK-NEXT: | | `-SYCLKernelEntryPointAttr {{.*}} KNT
|
||||
|
||||
// CHECK-NEXT: | `-FunctionDecl {{.*}} skep2 'void (K<2>)' explicit_instantiation_definition instantiated_from 0x{{.+}}
|
||||
@ -77,6 +86,15 @@ void skep2<KN<2>>(K<2>);
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod {{.*}} 'operator()' 'void () const'
|
||||
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const K<2>' lvalue <NoOp>
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'K<2>' lvalue ParmVar {{.*}} 'k' 'K<2>'
|
||||
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | | `-CallExpr {{.*}} 'void'
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} <FunctionToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, K<2>)' lvalue Function {{.*}} 'sycl_kernel_launch' {{.*}}
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi2EE"
|
||||
// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'K<2>' 'void (K<2> &&) noexcept'
|
||||
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'K<2>' xvalue <NoOp>
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'K<2>' lvalue ParmVar {{.*}} 'k' 'K<2>'
|
||||
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
|
||||
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'K<2>'
|
||||
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
|
||||
@ -102,9 +120,10 @@ void skep3<KN<3>>(K<3> k) {
|
||||
// CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} KT
|
||||
// CHECK-NEXT: | |-FunctionDecl {{.*}} skep3 'void (KT)'
|
||||
// CHECK-NEXT: | | |-ParmVarDecl {{.*}} k 'KT'
|
||||
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | | `-CallExpr {{.*}} '<dependent type>'
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
|
||||
// CHECK-NEXT: | | |-UnresolvedSYCLKernelCallStmt {{.*}}
|
||||
// CHECK-NEXT: | | | `-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | | `-CallExpr {{.*}} '<dependent type>'
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
|
||||
// CHECK-NEXT: | | `-SYCLKernelEntryPointAttr {{.*}} KNT
|
||||
|
||||
// CHECK-NEXT: | `-Function {{.*}} 'skep3' 'void (K<3>)'
|
||||
@ -123,6 +142,15 @@ void skep3<KN<3>>(K<3> k) {
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod {{.*}} 'operator()' 'void () const'
|
||||
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const K<3>' lvalue <NoOp>
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'K<3>' lvalue ParmVar {{.*}} 'k' 'K<3>'
|
||||
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | | `-CallExpr {{.*}}
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(const char *, K<3>)' <FunctionToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, K<3>)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, K<3>)' {{.*}}
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi3EE"
|
||||
// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'K<3>' 'void (K<3> &&) noexcept'
|
||||
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'K<3>' xvalue <NoOp>
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'K<3>' lvalue ParmVar {{.*}} 'k' 'K<3>'
|
||||
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
|
||||
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'K<3>'
|
||||
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
|
||||
@ -152,6 +180,21 @@ void skep4(K<4> k, int p1, int p2) {
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p1' 'int'
|
||||
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p2' 'int'
|
||||
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | | `-CallExpr {{.*}} 'void'
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(const char *, K<4>, int, int)' <FunctionToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, K<4>, int, int)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, K<4>, int, int)' {{.*}}
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi4EE"
|
||||
// CHECK-NEXT: | | | |-CXXConstructExpr {{.*}} 'K<4>' 'void (K<4> &&) noexcept'
|
||||
// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'K<4>' xvalue <NoOp>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'K<4>' lvalue ParmVar {{.*}} 'k' 'K<4>'
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
|
||||
// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' xvalue <NoOp>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p1' 'int'
|
||||
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
|
||||
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' xvalue <NoOp>
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p2' 'int'
|
||||
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
|
||||
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'K<4>'
|
||||
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used p1 'int'
|
||||
@ -182,7 +225,28 @@ void skep5(int unused1, K<5> k, int unused2, int p, int unused3) {
|
||||
// CHECK-NEXT: | |-ParmVarDecl {{.*}} unused3 'int'
|
||||
// CHECK-NEXT: | |-SYCLKernelCallStmt {{.*}}
|
||||
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
|
||||
// CHECK: | | `-OutlinedFunctionDecl {{.*}}
|
||||
// CHECK: | | |-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | | `-CallExpr {{.*}} 'void'
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(const char *, int, K<5>, int, int, int)' <FunctionToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, int, K<5>, int, int, int)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, int, K<5>, int, int, int)' {{.*}}
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi5EE"
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
|
||||
// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' xvalue <NoOp>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'unused1' 'int'
|
||||
// CHECK-NEXT: | | | |-CXXConstructExpr {{.*}} 'K<5>' 'void (K<5> &&) noexcept'
|
||||
// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'K<5>' xvalue <NoOp>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'K<5>' lvalue ParmVar {{.*}} 'k' 'K<5>'
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
|
||||
// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' xvalue <NoOp>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'unused2' 'int'
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
|
||||
// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' xvalue <NoOp>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p' 'int'
|
||||
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
|
||||
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' xvalue <NoOp>
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'unused3' 'int'
|
||||
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
|
||||
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit unused1 'int'
|
||||
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'K<5>'
|
||||
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit unused2 'int'
|
||||
@ -227,6 +291,14 @@ void skep6(const S6 &k) {
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)() const' <FunctionToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod {{.*}} 'operator()' 'void () const'
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'const S6' lvalue ParmVar {{.*}} 'k' 'const S6 &'
|
||||
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | | `-CallExpr {{.*}} 'void'
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(const char *, S6)' <FunctionToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, S6)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, S6)' {{.*}}
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi6EE"
|
||||
// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'S6' 'void (const S6 &) noexcept'
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'const S6' lvalue ParmVar {{.*}} 'k' 'const S6 &'
|
||||
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
|
||||
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'const S6 &'
|
||||
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
|
||||
@ -260,6 +332,15 @@ void skep7(S7 k) {
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod {{.*}} 'operator()' 'void () const'
|
||||
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const S7' lvalue <NoOp>
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'S7' lvalue ParmVar {{.*}} 'k' 'S7'
|
||||
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | | `-CallExpr {{.*}} 'void'
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(const char *, S7)' <FunctionToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, S7)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, S7)' {{.*}}
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi7EE"
|
||||
// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'S7' 'void (S7 &&) noexcept'
|
||||
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'S7' xvalue <NoOp>
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'S7' lvalue ParmVar {{.*}} 'k' 'S7'
|
||||
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
|
||||
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'S7'
|
||||
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
|
||||
@ -270,6 +351,114 @@ void skep7(S7 k) {
|
||||
// CHECK-NEXT: | | `-DeclRefExpr {{.*}} 'S7' lvalue ImplicitParam {{.*}} 'k' 'S7'
|
||||
// CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<7>
|
||||
|
||||
// Symbol names generated for the kernel entry point function should be
|
||||
// representable in the ordinary literal encoding even when the kernel name
|
||||
// type is named with esoteric characters.
|
||||
struct \u03b4\u03c4\u03c7; // Delta Tau Chi (δτχ)
|
||||
struct S8 {
|
||||
void operator()() const;
|
||||
};
|
||||
[[clang::sycl_kernel_entry_point(\u03b4\u03c4\u03c7)]]
|
||||
void skep8(S8 k) {
|
||||
k();
|
||||
}
|
||||
// CHECK: |-FunctionDecl {{.*}} skep8 'void (S8)'
|
||||
// CHECK-NEXT: | |-ParmVarDecl {{.*}} used k 'S8'
|
||||
// CHECK-NEXT: | |-SYCLKernelCallStmt {{.*}}
|
||||
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
|
||||
// CHECK: | | |-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | | `-CallExpr {{.*}} 'void'
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(const char *, S8)' <FunctionToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, S8)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, S8)' {{.*}}
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[12]' lvalue "_ZTS6\316\264\317\204\317\207"
|
||||
// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'S8' 'void (S8 &&) noexcept'
|
||||
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'S8' xvalue <NoOp>
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'S8' lvalue ParmVar {{.*}} 'k' 'S8'
|
||||
// CHECK: | | `-OutlinedFunctionDecl {{.*}}
|
||||
// CHECK: | `-SYCLKernelEntryPointAttr {{.*}}
|
||||
|
||||
class Handler {
|
||||
template <typename KNT, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...) {}
|
||||
public:
|
||||
template<typename KNT, typename KT>
|
||||
[[clang::sycl_kernel_entry_point(KNT)]]
|
||||
void skep9(KT k, int a, int b) {
|
||||
k(a, b);
|
||||
}
|
||||
};
|
||||
void foo() {
|
||||
Handler H;
|
||||
H.skep9<KN<9>>([=] (int a, int b) { return a+b; }, 1, 2);
|
||||
}
|
||||
|
||||
// CHECK: | |-FunctionTemplateDecl {{.*}} skep9
|
||||
// CHECK-NEXT: | | |-TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 0 KNT
|
||||
// CHECK-NEXT: | | |-TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 1 KT
|
||||
// CHECK-NEXT: | | |-CXXMethodDecl {{.*}} skep9 'void (KT, int, int)' implicit-inline
|
||||
// CHECK-NEXT: | | | |-ParmVarDecl {{.*}} referenced k 'KT'
|
||||
// CHECK-NEXT: | | | |-ParmVarDecl {{.*}} referenced a 'int'
|
||||
// CHECK-NEXT: | | | |-ParmVarDecl {{.*}} referenced b 'int'
|
||||
// CHECK-NEXT: | | | |-UnresolvedSYCLKernelCallStmt {{.*}}
|
||||
// CHECK-NEXT: | | | | `-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | | | `-CallExpr {{.*}} '<dependent type>'
|
||||
// CHECK-NEXT: | | | | |-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
|
||||
// CHECK-NEXT: | | | | |-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'a' 'int'
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'b' 'int'
|
||||
// CHECK-NEXT: | | | `-SYCLKernelEntryPointAttr {{.*}} KNT
|
||||
// CHECK-NEXT: | | `-CXXMethodDecl {{.*}} used skep9 {{.*}} implicit_instantiation implicit-inline instantiated_from 0x{{.*}}
|
||||
// CHECK-NEXT: | | |-TemplateArgument type 'KN<9>'
|
||||
// CHECK-NEXT: | | | `-RecordType {{.*}} 'KN<9>' canonical
|
||||
// CHECK-NEXT: | | | `-ClassTemplateSpecialization {{.*}}'KN'
|
||||
// CHECK-NEXT: | | |-TemplateArgument type {{.*}}
|
||||
// CHECK-NEXT: | | | `-RecordType {{.*}}
|
||||
// CHECK-NEXT: | | | `-CXXRecord {{.*}}
|
||||
// CHECK-NEXT: | | |-ParmVarDecl {{.*}} used k {{.*}}
|
||||
// CHECK-NEXT: | | |-ParmVarDecl {{.*}} used a 'int'
|
||||
// CHECK-NEXT: | | |-ParmVarDecl {{.*}} used b 'int'
|
||||
// CHECK-NEXT: | | |-SYCLKernelCallStmt {{.*}}
|
||||
// CHECK-NEXT: | | | |-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | | | `-CXXOperatorCallExpr {{.*}} 'int' '()'
|
||||
// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'int (*)(int, int) const' <FunctionToPointerDecay>
|
||||
// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} 'int (int, int) const' lvalue CXXMethod {{.*}} 'operator()' 'int (int, int) const'
|
||||
// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} lvalue <NoOp>
|
||||
// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'k' {{.*}}
|
||||
// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
|
||||
// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'a' 'int'
|
||||
// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'b' 'int'
|
||||
// CHECK-NEXT: | | | |-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | | | `-CXXMemberCallExpr {{.*}} 'void'
|
||||
// CHECK-NEXT: | | | | |-MemberExpr {{.*}} '<bound member function type>' ->sycl_kernel_launch {{.*}}
|
||||
// CHECK-NEXT: | | | | | `-CXXThisExpr {{.*}} 'Handler *' implicit this
|
||||
// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
|
||||
// CHECK-NEXT: | | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi9EE"
|
||||
// CHECK-NEXT: | | | | |-CXXConstructExpr {{.*}}
|
||||
// CHECK-NEXT: | | | | | `-ImplicitCastExpr {{.*}} xvalue <NoOp>
|
||||
// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'k' {{.*}}
|
||||
// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
|
||||
// CHECK-NEXT: | | | | | `-ImplicitCastExpr {{.*}} 'int' xvalue <NoOp>
|
||||
// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'a' 'int'
|
||||
// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
|
||||
// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' xvalue <NoOp>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'b' 'int'
|
||||
// CHECK-NEXT: | | | `-OutlinedFunctionDecl {{.*}}
|
||||
// CHECK-NEXT: | | | |-ImplicitParamDecl {{.*}} implicit used k {{.*}}
|
||||
// CHECK-NEXT: | | | |-ImplicitParamDecl {{.*}} implicit used a 'int'
|
||||
// CHECK-NEXT: | | | |-ImplicitParamDecl {{.*}} implicit used b 'int'
|
||||
// CHECK-NEXT: | | | `-CompoundStmt {{.*}}
|
||||
// CHECK-NEXT: | | | `-CXXOperatorCallExpr {{.*}} 'int' '()'
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int (*)(int, int) const' <FunctionToPointerDecay>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int (int, int) const' lvalue CXXMethod {{.*}} 'operator()' 'int (int, int) const'
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} lvalue <NoOp>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} lvalue ImplicitParam {{.*}} 'k' {{.*}}
|
||||
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
|
||||
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ImplicitParam {{.*}} 'a' 'int'
|
||||
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
|
||||
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'int' lvalue ImplicitParam {{.*}} 'b' 'int'
|
||||
// CHECK-NEXT: | | `-SYCLKernelEntryPointAttr {{.*}} struct KN<9>
|
||||
|
||||
|
||||
void the_end() {}
|
||||
// CHECK: `-FunctionDecl {{.*}} the_end 'void ()'
|
||||
|
||||
@ -28,6 +28,9 @@
|
||||
// A unique kernel name type is required for each declared kernel entry point.
|
||||
template<int, int=0> struct KN;
|
||||
|
||||
template<typename KernelName, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts... Args) {}
|
||||
|
||||
[[clang::sycl_kernel_entry_point(KN<1>)]]
|
||||
void skep1() {
|
||||
}
|
||||
|
||||
25
clang/test/ASTSYCL/ast-print-sycl-kernel-call.cpp
Normal file
25
clang/test/ASTSYCL/ast-print-sycl-kernel-call.cpp
Normal file
@ -0,0 +1,25 @@
|
||||
// RUN: %clang_cc1 -fsycl-is-host -ast-print %s -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -ast-print %s -o - | FileCheck %s
|
||||
|
||||
struct sycl_kernel_launcher {
|
||||
template<typename KernelName, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...) {}
|
||||
|
||||
template<typename KernelName, typename KernelType>
|
||||
[[clang::sycl_kernel_entry_point(KernelName)]]
|
||||
void sycl_kernel_entry_point(KernelType kernel) {
|
||||
kernel();
|
||||
}
|
||||
};
|
||||
// CHECK: template <typename KernelName, typename KernelType> void sycl_kernel_entry_point(KernelType kernel)
|
||||
// CHECK-NEXT: {
|
||||
// CHECK-NEXT: kernel();
|
||||
// CHECK-NEXT: }
|
||||
// CHECK: template<> void sycl_kernel_entry_point<KN, (lambda at {{.*}})>((lambda at {{.*}}) kernel)
|
||||
// CHECK-NEXT: {
|
||||
// CHECK-NEXT: kernel();
|
||||
// CHECK-NEXT: }
|
||||
|
||||
void f(sycl_kernel_launcher skl) {
|
||||
skl.sycl_kernel_entry_point<struct KN>([]{});
|
||||
}
|
||||
@ -26,6 +26,9 @@ int foo() {
|
||||
return 1;
|
||||
}
|
||||
|
||||
template <typename Name, typename... Ts>
|
||||
void sycl_kernel_launch(Ts...) {}
|
||||
|
||||
template <typename Name, typename Func>
|
||||
[[clang::sycl_kernel_entry_point(Name)]] void kernel_single_task(const Func &kernelFunc) {
|
||||
kernelFunc();
|
||||
|
||||
@ -2,31 +2,36 @@
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-unknown-linux-gnu -triple amdgcn-amd-amdhsa -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-AMDGCN %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-unknown-linux-gnu -triple nvptx-nvidia-cuda -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-NVPTX %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-unknown-linux-gnu -triple nvptx64-nvidia-cuda -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-NVPTX %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-unknown-linux-gnu -triple spir-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-unknown-linux-gnu -triple spir64-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-unknown-linux-gnu -triple spirv32-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-unknown-linux-gnu -triple spirv64-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-unknown-linux-gnu -triple spir-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR,CHECK-SPIRNV %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-unknown-linux-gnu -triple spir64-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR,CHECK-SPIRNV %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-unknown-linux-gnu -triple spirv32-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR,CHECK-SPIRV %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-unknown-linux-gnu -triple spirv64-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR,CHECK-SPIRV %s
|
||||
// RUN: %clang_cc1 -fsycl-is-host -emit-llvm -triple x86_64-pc-windows-msvc -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-HOST,CHECK-HOST-WINDOWS %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-pc-windows-msvc -triple amdgcn-amd-amdhsa -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-AMDGCN %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-pc-windows-msvc -triple nvptx-nvidia-cuda -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-NVPTX %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-pc-windows-msvc -triple nvptx64-nvidia-cuda -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-NVPTX %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-pc-windows-msvc -triple spir-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-pc-windows-msvc -triple spir64-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-pc-windows-msvc -triple spirv32-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-pc-windows-msvc -triple spirv64-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-pc-windows-msvc -triple spir-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR,CHECK-SPIRNV %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-pc-windows-msvc -triple spir64-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR,CHECK-SPIRNV %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-pc-windows-msvc -triple spirv32-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR,CHECK-SPIRV %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-pc-windows-msvc -triple spirv64-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR,CHECK-SPIRV %s
|
||||
// RUN: %clang_cc1 -fsycl-is-host -emit-llvm -triple x86_64-uefi -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-HOST,CHECK-HOST-WINDOWS %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-uefi -triple amdgcn-amd-amdhsa -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-AMDGCN %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-uefi -triple nvptx-nvidia-cuda -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-NVPTX %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-uefi -triple nvptx64-nvidia-cuda -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-NVPTX %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-uefi -triple spir-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-uefi -triple spir64-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-uefi -triple spirv32-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-uefi -triple spirv64-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-uefi -triple spir-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR,CHECK-SPIRNV %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-uefi -triple spir64-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR,CHECK-SPIRNV %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-uefi -triple spirv32-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR,CHECK-SPIRV %s
|
||||
// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-uefi -triple spirv64-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE,CHECK-SPIR,CHECK-SPIRV %s
|
||||
|
||||
// Test the generation of SYCL kernel caller functions. These functions are
|
||||
// generated from functions declared with the sycl_kernel_entry_point attribute
|
||||
// and emited during device compilation. They are not emitted during device
|
||||
// compilation.
|
||||
// Test code generation for functions declared with the sycl_kernel_entry_point
|
||||
// attribute. During host compilation, the bodies of such functions are replaced
|
||||
// with calls to a function template or variable template (with suitable call
|
||||
// operator) named sycl_kernel_launch. During device compilation, the bodies of
|
||||
// these functions are used to generate offload kernel entry points (SYCL kernel
|
||||
// caller functions).
|
||||
|
||||
template <typename KernelName, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...) {}
|
||||
|
||||
struct single_purpose_kernel_name;
|
||||
struct single_purpose_kernel {
|
||||
@ -44,57 +49,169 @@ void kernel_single_task(KernelType kernelFunc) {
|
||||
kernelFunc(42);
|
||||
}
|
||||
|
||||
// Exercise code gen with kernel name types named with esoteric characters.
|
||||
struct \u03b4\u03c4\u03c7; // Delta Tau Chi (δτχ)
|
||||
|
||||
class handler {
|
||||
template <typename KernelName, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...) {}
|
||||
public:
|
||||
template <typename KernelName, typename KernelType>
|
||||
[[clang::sycl_kernel_entry_point(KernelName)]]
|
||||
void kernel_entry_point(KernelType k, int a, int b) {
|
||||
k(a, b);
|
||||
}
|
||||
};
|
||||
|
||||
struct copyable {
|
||||
int i;
|
||||
~copyable();
|
||||
};
|
||||
|
||||
int main() {
|
||||
single_purpose_kernel obj;
|
||||
single_purpose_kernel_task(obj);
|
||||
int capture;
|
||||
auto lambda = [=](auto) { (void) capture; };
|
||||
kernel_single_task<decltype(lambda)>(lambda);
|
||||
kernel_single_task<\u03b4\u03c4\u03c7>([](int){});
|
||||
handler h;
|
||||
copyable c{42};
|
||||
h.kernel_entry_point<struct KN>([=] (int a, int b) { return c.i + a + b; }, 1, 2);
|
||||
}
|
||||
|
||||
// Verify that SYCL kernel caller functions are not emitted during host
|
||||
// compilation.
|
||||
//
|
||||
// CHECK-HOST-NOT: _ZTS26single_purpose_kernel_name
|
||||
// CHECK-HOST-NOT: _ZTSZ4mainE18lambda_kernel_name
|
||||
// CHECK-HOST-NOT: define {{.*}} @_ZTS26single_purpose_kernel_name
|
||||
// CHECK-HOST-NOT: define {{.*}} @_ZTSZ4mainEUlT_E_
|
||||
// CHECK-HOST-NOT: define {{.*}} @"_ZTS6\CE\B4\CF\84\CF\87"
|
||||
// CHECK-HOST-NOT: define {{.*}} @_ZTSZ4mainE2KN
|
||||
|
||||
// Verify that sycl_kernel_entry_point attributed functions are not emitted
|
||||
// during device compilation.
|
||||
//
|
||||
// CHECK-DEVICE-NOT: single_purpose_kernel_task
|
||||
// CHECK-DEVICE-NOT: kernel_single_task
|
||||
// CHECK-DEVICE-NOT: kernel_entry_point
|
||||
|
||||
// Verify that no code is generated for the bodies of sycl_kernel_entry_point
|
||||
// attributed functions during host compilation. ODR-use of these functions may
|
||||
// require them to be emitted, but they have no effect if called.
|
||||
// Verify that kernel launch code is generated for sycl_kernel_entry_point
|
||||
// attributed functions during host compilation.
|
||||
//
|
||||
// CHECK-HOST-LINUX: @.str = private unnamed_addr constant [33 x i8] c"_ZTS26single_purpose_kernel_name\00", align 1
|
||||
// CHECK-HOST-LINUX: @.str.1 = private unnamed_addr constant [18 x i8] c"_ZTSZ4mainEUlT_E_\00", align 1
|
||||
// CHECK-HOST-LINUX: @.str.2 = private unnamed_addr constant [12 x i8] c"_ZTS6\CE\B4\CF\84\CF\87\00", align 1
|
||||
//
|
||||
// CHECK-HOST-LINUX: define dso_local void @_Z26single_purpose_kernel_task21single_purpose_kernel() #{{[0-9]+}} {
|
||||
// CHECK-HOST-LINUX-NEXT: entry:
|
||||
// CHECK-HOST-LINUX-NEXT: %kernelFunc = alloca %struct.single_purpose_kernel, align 1
|
||||
// CHECK-HOST-LINUX-NEXT: %agg.tmp = alloca %struct.single_purpose_kernel, align 1
|
||||
// CHECK-HOST-LINUX-NEXT: call void @_Z18sycl_kernel_launchI26single_purpose_kernel_nameJ21single_purpose_kernelEEvPKcDpT0_(ptr noundef @.str)
|
||||
// CHECK-HOST-LINUX-NEXT: ret void
|
||||
// CHECK-HOST-LINUX-NEXT: }
|
||||
//
|
||||
// CHECK-HOST-LINUX: define internal void @_Z18kernel_single_taskIZ4mainEUlT_E_S1_EvT0_(i32 %kernelFunc.coerce) #{{[0-9]+}} {
|
||||
// CHECK-HOST-LINUX-NEXT: entry:
|
||||
// CHECK-HOST-LINUX-NEXT: %kernelFunc = alloca %class.anon, align 4
|
||||
// CHECK-HOST-LINUX-NEXT: %agg.tmp = alloca %class.anon, align 4
|
||||
// CHECK-HOST-LINUX-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon, ptr %kernelFunc, i32 0, i32 0
|
||||
// CHECK-HOST-LINUX-NEXT: store i32 %kernelFunc.coerce, ptr %coerce.dive, align 4
|
||||
// CHECK-HOST-LINUX-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %kernelFunc, i64 4, i1 false)
|
||||
// CHECK-HOST-LINUX-NEXT: %coerce.dive1 = getelementptr inbounds nuw %class.anon, ptr %agg.tmp, i32 0, i32 0
|
||||
// CHECK-HOST-LINUX-NEXT: %0 = load i32, ptr %coerce.dive1, align 4
|
||||
// CHECK-HOST-LINUX-NEXT: call void @_Z18sycl_kernel_launchIZ4mainEUlT_E_JS1_EEvPKcDpT0_(ptr noundef @.str.1, i32 %0)
|
||||
// CHECK-HOST-LINUX-NEXT: ret void
|
||||
// CHECK-HOST-LINUX-NEXT: }
|
||||
//
|
||||
// CHECK-HOST-LINUX: define internal void @"_Z18kernel_single_taskI6\CE\B4\CF\84\CF\87Z4mainEUliE_EvT0_"() #{{[0-9]+}} {
|
||||
// CHECK-HOST-LINUX-NEXT: entry:
|
||||
// CHECK-HOST-LINUX-NEXT: %kernelFunc = alloca %class.anon.0, align 1
|
||||
// CHECK-HOST-LINUX-NEXT: %agg.tmp = alloca %class.anon.0, align 1
|
||||
// CHECK-HOST-LINUX-NEXT: call void @"_Z18sycl_kernel_launchI6\CE\B4\CF\84\CF\87JZ4mainEUliE_EEvPKcDpT0_"(ptr noundef @.str.2)
|
||||
// CHECK-HOST-LINUX-NEXT: ret void
|
||||
// CHECK-HOST-LINUX-NEXT: }
|
||||
|
||||
|
||||
// CHECK-HOST-LINUX: define internal void @_ZN7handler18kernel_entry_pointIZ4mainE2KNZ4mainEUliiE_EEvT0_ii(ptr noundef nonnull align 1 dereferenceable(1) %this, ptr noundef %k, i32 noundef %a, i32 noundef %b) #{{[0-9]+}} align 2 {
|
||||
// CHECK-HOST-LINUX-NEXT: entry:
|
||||
// CHECK-HOST-LINUX-NEXT: %this.addr = alloca ptr, align 8
|
||||
// CHECK-HOST-LINUX-NEXT: %k.indirect_addr = alloca ptr, align 8
|
||||
// CHECK-HOST-LINUX-NEXT: %a.addr = alloca i32, align 4
|
||||
// CHECK-HOST-LINUX-NEXT: %b.addr = alloca i32, align 4
|
||||
// CHECK-HOST-LINUX-NEXT: %agg.tmp = alloca %class.anon.1, align 4
|
||||
// CHECK-HOST-LINUX-NEXT: store ptr %this, ptr %this.addr, align 8
|
||||
// CHECK-HOST-LINUX-NEXT: store ptr %k, ptr %k.indirect_addr, align 8
|
||||
// CHECK-HOST-LINUX-NEXT: store i32 %a, ptr %a.addr, align 4
|
||||
// CHECK-HOST-LINUX-NEXT: store i32 %b, ptr %b.addr, align 4
|
||||
// CHECK-HOST-LINUX-NEXT: %this1 = load ptr, ptr %this.addr, align 8
|
||||
// CHECK-HOST-LINUX-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %k, i64 4, i1 false)
|
||||
// CHECK-HOST-LINUX-NEXT: %0 = load i32, ptr %a.addr, align 4
|
||||
// CHECK-HOST-LINUX-NEXT: %1 = load i32, ptr %b.addr, align 4
|
||||
// CHECK-HOST-LINUX-NEXT: call void @_ZN7handler18sycl_kernel_launchIZ4mainE2KNJZ4mainEUliiE_iiEEEvPKcDpT0_(ptr noundef nonnull align 1 dereferenceable(1) %this1, ptr noundef @.str.3, ptr noundef %agg.tmp, i32 noundef %0, i32 noundef %1)
|
||||
// CHECK-HOST-LINUX-NEXT: call void @_ZZ4mainENUliiE_D1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) %agg.tmp) #{{[0-9]+}}
|
||||
// CHECK-HOST-LINUX-NEXT: ret void
|
||||
// CHECK-HOST-LINUX-NEXT: }
|
||||
|
||||
// CHECK-HOST-WINDOWS: define dso_local void @"?single_purpose_kernel_task@@YAXUsingle_purpose_kernel@@@Z"(i8 %kernelFunc.coerce) #{{[0-9]+}} {
|
||||
// CHECK-HOST-WINDOWS-NEXT: entry:
|
||||
// CHECK-HOST-WINDOWS-NEXT: %kernelFunc = alloca %struct.single_purpose_kernel, align 1
|
||||
// CHECK-HOST-WINDOWS-NEXT: %agg.tmp = alloca %struct.single_purpose_kernel, align 1
|
||||
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive = getelementptr inbounds nuw %struct.single_purpose_kernel, ptr %kernelFunc, i32 0, i32 0
|
||||
// CHECK-HOST-WINDOWS-NEXT: store i8 %kernelFunc.coerce, ptr %coerce.dive, align 1
|
||||
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive1 = getelementptr inbounds nuw %struct.single_purpose_kernel, ptr %agg.tmp, i32 0, i32 0
|
||||
// CHECK-HOST-WINDOWS-NEXT: %0 = load i8, ptr %coerce.dive1, align 1
|
||||
// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch@Usingle_purpose_kernel_name@@Usingle_purpose_kernel@@@@YAXPEBDUsingle_purpose_kernel@@@Z"(ptr noundef @"??_C@_0CB@KFIJOMLB@_ZTS26single_purpose_kernel_name@", i8 %0)
|
||||
// CHECK-HOST-WINDOWS-NEXT: ret void
|
||||
// CHECK-HOST-WINDOWS-NEXT: }
|
||||
//
|
||||
// CHECK-HOST-WINDOWS: define internal void @"??$kernel_single_task@V<lambda_1>@?0??main@@9@V1?0??2@9@@@YAXV<lambda_1>@?0??main@@9@@Z"(i32 %kernelFunc.coerce) #{{[0-9]+}} {
|
||||
// CHECK-HOST-WINDOWS-NEXT: entry:
|
||||
// CHECK-HOST-WINDOWS-NEXT: %kernelFunc = alloca %class.anon, align 4
|
||||
// CHECK-HOST-WINDOWS-NEXT: %agg.tmp = alloca %class.anon, align 4
|
||||
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon, ptr %kernelFunc, i32 0, i32 0
|
||||
// CHECK-HOST-WINDOWS-NEXT: store i32 %kernelFunc.coerce, ptr %coerce.dive, align 4
|
||||
// CHECK-HOST-WINDOWS-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %kernelFunc, i64 4, i1 false)
|
||||
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive1 = getelementptr inbounds nuw %class.anon, ptr %agg.tmp, i32 0, i32 0
|
||||
// CHECK-HOST-WINDOWS-NEXT: %0 = load i32, ptr %coerce.dive1, align 4
|
||||
// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch@V<lambda_1>@?0??main@@9@V1?0??2@9@@@YAXPEBDV<lambda_1>@?0??main@@9@@Z"(ptr noundef @"??_C@_0BC@NHCDOLAA@_ZTSZ4mainEUlT_E_?$AA@", i32 %0)
|
||||
//
|
||||
// CHECK-HOST-WINDOWS-NEXT: ret void
|
||||
// CHECK-HOST-WINDOWS-NEXT: }
|
||||
//
|
||||
// CHECK-HOST-WINDOWS: define internal void @"??$kernel_single_task@U\CE\B4\CF\84\CF\87@@V<lambda_2>@?0??main@@9@@@YAXV<lambda_2>@?0??main@@9@@Z"(i8 %kernelFunc.coerce) #{{[0-9]+}} {
|
||||
// CHECK-HOST-WINDOWS-NEXT: entry:
|
||||
// CHECK-HOST-WINDOWS-NEXT: %kernelFunc = alloca %class.anon.0, align 1
|
||||
// CHECK-HOST-WINDOWS-NEXT: %agg.tmp = alloca %class.anon.0, align 1
|
||||
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon.0, ptr %kernelFunc, i32 0, i32 0
|
||||
// CHECK-HOST-WINDOWS-NEXT: store i8 %kernelFunc.coerce, ptr %coerce.dive, align 1
|
||||
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive1 = getelementptr inbounds nuw %class.anon.0, ptr %agg.tmp, i32 0, i32 0
|
||||
// CHECK-HOST-WINDOWS-NEXT: %0 = load i8, ptr %coerce.dive1, align 1
|
||||
// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch@U\CE\B4\CF\84\CF\87@@V<lambda_2>@?0??main@@9@@@YAXPEBDV<lambda_2>@?0??main@@9@@Z"(ptr noundef @"??_C@_0M@BCGAEMBE@_ZTS6?N?$LE?O?$IE?O?$IH?$AA@", i8 %0)
|
||||
// CHECK-HOST-WINDOWS-NEXT: ret void
|
||||
// CHECK-HOST-WINDOWS-NEXT: }
|
||||
|
||||
// CHECK-HOST-WINDOWS: define internal void @"??$kernel_entry_point@UKN@?1??main@@9@V<lambda_3>@?0??2@9@@handler@@QEAAXV<lambda_3>@?0??main@@9@HH@Z"(ptr noundef nonnull align 1 dereferenceable(1) %this, i32 %k.coerce, i32 noundef %a, i32 noundef %b) #{{[0-9]+}} align 2
|
||||
// CHECK-HOST-WINDOWS-NEXT: entry:
|
||||
// CHECK-HOST-WINDOWS-NEXT: %k = alloca %class.anon.1, align 4
|
||||
// CHECK-HOST-WINDOWS-NEXT: %b.addr = alloca i32, align 4
|
||||
// CHECK-HOST-WINDOWS-NEXT: %a.addr = alloca i32, align 4
|
||||
// CHECK-HOST-WINDOWS-NEXT: %this.addr = alloca ptr, align 8
|
||||
// CHECK-HOST-WINDOWS-NEXT: %agg.tmp = alloca %class.anon.1, align 4
|
||||
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon.1, ptr %k, i32 0, i32 0
|
||||
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive1 = getelementptr inbounds nuw %struct.copyable, ptr %coerce.dive, i32 0, i32 0
|
||||
// CHECK-HOST-WINDOWS-NEXT: store i32 %k.coerce, ptr %coerce.dive1, align 4
|
||||
// CHECK-HOST-WINDOWS-NEXT: store i32 %b, ptr %b.addr, align 4
|
||||
// CHECK-HOST-WINDOWS-NEXT: store i32 %a, ptr %a.addr, align 4
|
||||
// CHECK-HOST-WINDOWS-NEXT: store ptr %this, ptr %this.addr, align 8
|
||||
// CHECK-HOST-WINDOWS-NEXT: %this2 = load ptr, ptr %this.addr, align 8
|
||||
// CHECK-HOST-WINDOWS-NEXT: %0 = load i32, ptr %b.addr, align 4
|
||||
// CHECK-HOST-WINDOWS-NEXT: %1 = load i32, ptr %a.addr, align 4
|
||||
// CHECK-HOST-WINDOWS-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %k, i64 4, i1 false)
|
||||
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive3 = getelementptr inbounds nuw %class.anon.1, ptr %agg.tmp, i32 0, i32 0
|
||||
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive4 = getelementptr inbounds nuw %struct.copyable, ptr %coerce.dive3, i32 0, i32 0
|
||||
// CHECK-HOST-WINDOWS-NEXT: %2 = load i32, ptr %coerce.dive4, align 4
|
||||
// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch@UKN@?1??main@@9@V<lambda_3>@?0??2@9@HH@handler@@AEAAXPEBDV<lambda_3>@?0??main@@9@HH@Z"(ptr noundef nonnull align 1 dereferenceable(1) %this2, ptr noundef @"??_C@_0P@DLGHPODL@_ZTSZ4mainE2KN?$AA@", i32 %2, i32 noundef %1, i32 noundef %0)
|
||||
// CHECK-HOST-WINDOWS-NEXT: call void @"??1<lambda_3>@?0??main@@9@QEAA@XZ"(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) %k) #{{[0-9]+}}
|
||||
// CHECK-HOST-WINDOWS-NEXT: ret void
|
||||
// CHECK-HOST-WINDOWS-NEXT: }
|
||||
|
||||
@ -179,6 +296,122 @@ int main() {
|
||||
// CHECK-SPIR-NEXT: }
|
||||
// CHECK-SPIR: define internal spir_func void @_ZZ4mainENKUlT_E_clIiEEDaS_
|
||||
|
||||
// IR for the SYCL kernel caller function generated for kernel_single_task with
|
||||
// the Delta Tau Chi type as the SYCL kernel name type.
|
||||
//
|
||||
// CHECK-AMDGCN: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
|
||||
// CHECK-AMDGCN-NEXT: define dso_local amdgpu_kernel void @"_ZTS6\CE\B4\CF\84\CF\87"
|
||||
// CHECK-AMDGCN-SAME: (ptr addrspace(4) noundef byref(%class.anon.0) align 1 %0) #[[AMDGCN_ATTR0]] {
|
||||
// CHECK-AMDGCN-NEXT: entry:
|
||||
// CHECK-AMDGCN-NEXT: %coerce = alloca %class.anon.0, align 1, addrspace(5)
|
||||
// CHECK-AMDGCN-NEXT: %kernelFunc = addrspacecast ptr addrspace(5) %coerce to ptr
|
||||
// CHECK-AMDGCN-NEXT: call void @llvm.memcpy.p0.p4.i64(ptr align 1 %kernelFunc, ptr addrspace(4) align 1 %0, i64 1, i1 false)
|
||||
// CHECK-AMDGCN-NEXT: call void @_ZZ4mainENKUliE_clEi
|
||||
// CHECK-AMDGCN-SAME: (ptr noundef nonnull align 1 dereferenceable(1) %kernelFunc, i32 noundef 42) #[[AMDGCN_ATTR1:[0-9]+]]
|
||||
// CHECK-AMDGCN-NEXT: ret void
|
||||
// CHECK-AMDGCN-NEXT: }
|
||||
// CHECK-AMDGCN: define internal void @_ZZ4mainENKUliE_clEi
|
||||
//
|
||||
// CHECK-NVPTX: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
|
||||
// CHECK-NVPTX-NEXT: define dso_local ptx_kernel void @"_ZTS6\CE\B4\CF\84\CF\87"
|
||||
// CHECK-NVPTX-SAME: (ptr noundef byval(%class.anon.0) align 1 %kernelFunc) #[[NVPTX_ATTR0:[0-9]+]] {
|
||||
// CHECK-NVPTX-NEXT: entry:
|
||||
// CHECK-NVPTX-NEXT: call void @_ZZ4mainENKUliE_clEi
|
||||
// CHECK-NVPTX-SAME: (ptr noundef nonnull align 1 dereferenceable(1) %kernelFunc, i32 noundef 42) #[[NVPTX_ATTR1:[0-9]+]]
|
||||
// CHECK-NVPTX-NEXT: ret void
|
||||
// CHECK-NVPTX-NEXT: }
|
||||
// CHECK-NVPTX: define internal void @_ZZ4mainENKUliE_clEi
|
||||
//
|
||||
// CHECK-SPIR: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
|
||||
// CHECK-SPIR-NEXT: define {{[a-z_ ]*}}spir_kernel void @"_ZTS6\CE\B4\CF\84\CF\87"
|
||||
// CHECK-SPIR-SAME: (ptr noundef byval(%class.anon.0) align 1 %kernelFunc) #[[SPIR_ATTR0:[0-9]+]] {
|
||||
// CHECK-SPIR-NEXT: entry:
|
||||
// CHECK-SPIR-NEXT: %kernelFunc.ascast = addrspacecast ptr %kernelFunc to ptr addrspace(4)
|
||||
// CHECK-SPIR-NEXT: call spir_func void @_ZZ4mainENKUliE_clEi
|
||||
// CHECK-SPIR-SAME: (ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %kernelFunc.ascast, i32 noundef 42) #[[SPIR_ATTR1:[0-9]+]]
|
||||
// CHECK-SPIR-NEXT: ret void
|
||||
// CHECK-SPIR-NEXT: }
|
||||
// CHECK-SPIR: define internal spir_func void @_ZZ4mainENKUliE_clEi
|
||||
|
||||
// IR for the SYCL kernel caller function generated for
|
||||
// handler::kernel_entry_point with main::KN as the SYCL kernel name type.
|
||||
//
|
||||
// CHECK-AMDGCN: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
|
||||
// CHECK-AMDGCN-NEXT: define dso_local amdgpu_kernel void @_ZTSZ4mainE2KN
|
||||
// CHECK-AMDGCN-SAME: (i32 %k.coerce, i32 noundef %a, i32 noundef %b) #[[AMDGCN_ATTR0]] {
|
||||
// CHECK-AMDGCN-NEXT: entry:
|
||||
// CHECK-AMDGCN-NEXT: %k = alloca %class.anon.1, align 4, addrspace(5)
|
||||
// CHECK-AMDGCN-NEXT: %a.addr = alloca i32, align 4, addrspace(5)
|
||||
// CHECK-AMDGCN-NEXT: %b.addr = alloca i32, align 4, addrspace(5)
|
||||
// CHECK-AMDGCN-NEXT: %k2 = addrspacecast ptr addrspace(5) %k to ptr
|
||||
// CHECK-AMDGCN-NEXT: %a.addr.ascast = addrspacecast ptr addrspace(5) %a.addr to ptr
|
||||
// CHECK-AMDGCN-NEXT: %b.addr.ascast = addrspacecast ptr addrspace(5) %b.addr to ptr
|
||||
// CHECK-AMDGCN-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon.1, ptr %k2, i32 0, i32 0
|
||||
// CHECK-AMDGCN-NEXT: %coerce.dive1 = getelementptr inbounds nuw %struct.copyable, ptr %coerce.dive, i32 0, i32 0
|
||||
// CHECK-AMDGCN-NEXT: store i32 %k.coerce, ptr %coerce.dive1, align 4
|
||||
// CHECK-AMDGCN-NEXT: store i32 %a, ptr %a.addr.ascast, align 4
|
||||
// CHECK-AMDGCN-NEXT: store i32 %b, ptr %b.addr.ascast, align 4
|
||||
// CHECK-AMDGCN-NEXT: %0 = load i32, ptr %a.addr.ascast, align 4
|
||||
// CHECK-AMDGCN-NEXT: %1 = load i32, ptr %b.addr.ascast, align 4
|
||||
// CHECK-AMDGCN-NEXT: %call = call noundef i32 @_ZZ4mainENKUliiE_clEii
|
||||
// CHECK-AMDGCN-SAME: (ptr noundef nonnull align 4 dereferenceable(4) %k2, i32 noundef %0, i32 noundef %1) #[[AMDGCN_ATTR1:[0-9]+]]
|
||||
// CHECK-AMDGCN-NEXT: ret void
|
||||
// CHECK-AMDGCN-NEXT: }
|
||||
//
|
||||
// CHECK-NVPTX: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
|
||||
// CHECK-NVPTX-NEXT: define dso_local ptx_kernel void @_ZTSZ4mainE2KN
|
||||
// CHECK-NVPTX-SAME: (ptr noundef byval(%class.anon.1) align 4 %k, i32 noundef %a, i32 noundef %b) #[[NVPTX_ATTR0:[0-9]+]] {
|
||||
// CHECK-NVPTX-NEXT: entry:
|
||||
// CHECK-NVPTX-NEXT: %a.addr = alloca i32, align 4
|
||||
// CHECK-NVPTX-NEXT: %b.addr = alloca i32, align 4
|
||||
// CHECK-NVPTX-NEXT: store i32 %a, ptr %a.addr, align 4
|
||||
// CHECK-NVPTX-NEXT: store i32 %b, ptr %b.addr, align 4
|
||||
// CHECK-NVPTX-NEXT: %0 = load i32, ptr %a.addr, align 4
|
||||
// CHECK-NVPTX-NEXT: %1 = load i32, ptr %b.addr, align 4
|
||||
// CHECK-NVPTX-NEXT: %call = call noundef i32 @_ZZ4mainENKUliiE_clEii
|
||||
// CHECK-NVPTX-SAME: (ptr noundef nonnull align 4 dereferenceable(4) %k, i32 noundef %0, i32 noundef %1) #[[NVPTX_ATTR1:[0-9]+]]
|
||||
// CHECK-NVPTX-NEXT: ret void
|
||||
// CHECK-NVPTX-NEXT: }
|
||||
//
|
||||
// CHECK-SPIRNV: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
|
||||
// CHECK-SPIRNV-NEXT: define dso_local spir_kernel void @_ZTSZ4mainE2KN
|
||||
// CHECK-SPIRNV-SAME: (ptr noundef %k, i32 noundef %a, i32 noundef %b) #[[SPIR_ATTR0:[0-9]+]] {
|
||||
// CHECK-SPIRNV-NEXT: entry:
|
||||
// CHECK-SPIRNV-NEXT: %k.indirect_addr = alloca ptr addrspace(4), align {{[48]}}
|
||||
// CHECK-SPIRNV-NEXT: %a.addr = alloca i32, align 4
|
||||
// CHECK-SPIRNV-NEXT: %b.addr = alloca i32, align 4
|
||||
// CHECK-SPIRNV-NEXT: %k.indirect_addr.ascast = addrspacecast ptr %k.indirect_addr to ptr addrspace(4)
|
||||
// CHECK-SPIRNV-NEXT: %a.addr.ascast = addrspacecast ptr %a.addr to ptr addrspace(4)
|
||||
// CHECK-SPIRNV-NEXT: %b.addr.ascast = addrspacecast ptr %b.addr to ptr addrspace(4)
|
||||
// CHECK-SPIRNV-NEXT: store ptr %k, ptr addrspace(4) %k.indirect_addr.ascast, align {{[48]}}
|
||||
// CHECK-SPIRNV-NEXT: %k.ascast = addrspacecast ptr %k to ptr addrspace(4)
|
||||
// CHECK-SPIRNV-NEXT: store i32 %a, ptr addrspace(4) %a.addr.ascast, align 4
|
||||
// CHECK-SPIRNV-NEXT: store i32 %b, ptr addrspace(4) %b.addr.ascast, align 4
|
||||
// CHECK-SPIRNV-NEXT: %0 = load i32, ptr addrspace(4) %a.addr.ascast, align 4
|
||||
// CHECK-SPIRNV-NEXT: %1 = load i32, ptr addrspace(4) %b.addr.ascast, align 4
|
||||
// CHECK-SPIRNV-NEXT: %call = call spir_func noundef i32 @_ZZ4mainENKUliiE_clEii
|
||||
// CHECK-SPIRNV-SAME: (ptr addrspace(4) noundef align 4 dereferenceable_or_null(4) %k.ascast, i32 noundef %0, i32 noundef %1) #[[SPIR_ATTR1:[0-9]+]]
|
||||
// CHECK-SPIRNV-NEXT: ret void
|
||||
// CHECK-SPIRNV-NEXT: }
|
||||
//
|
||||
// CHECK-SPIRV: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
|
||||
// CHECK-SPIRV-NEXT: define spir_kernel void @_ZTSZ4mainE2KN
|
||||
// CHECK-SPIRV-SAME: (ptr noundef byval(%class.anon.1) align 4 %k, i32 noundef %a, i32 noundef %b) #[[SPIR_ATTR0:[0-9]+]] {
|
||||
// CHECK-SPIRV-NEXT: entry:
|
||||
// CHECK-SPIRV-NEXT: %a.addr = alloca i32, align 4
|
||||
// CHECK-SPIRV-NEXT: %b.addr = alloca i32, align 4
|
||||
// CHECK-SPIRV-NEXT: %a.addr.ascast = addrspacecast ptr %a.addr to ptr addrspace(4)
|
||||
// CHECK-SPIRV-NEXT: %b.addr.ascast = addrspacecast ptr %b.addr to ptr addrspace(4)
|
||||
// CHECK-SPIRV-NEXT: %k.ascast = addrspacecast ptr %k to ptr addrspace(4)
|
||||
// CHECK-SPIRV-NEXT: store i32 %a, ptr addrspace(4) %a.addr.ascast, align 4
|
||||
// CHECK-SPIRV-NEXT: store i32 %b, ptr addrspace(4) %b.addr.ascast, align 4
|
||||
// CHECK-SPIRV-NEXT: %0 = load i32, ptr addrspace(4) %a.addr.ascast, align 4
|
||||
// CHECK-SPIRV-NEXT: %1 = load i32, ptr addrspace(4) %b.addr.ascast, align 4
|
||||
// CHECK-SPIRV-NEXT: %call = call spir_func noundef i32 @_ZZ4mainENKUliiE_clEii
|
||||
// CHECK-SPIRV-SAME: (ptr addrspace(4) noundef align 4 dereferenceable_or_null(4) %k.ascast, i32 noundef %0, i32 noundef %1) #[[SPIR_ATTR1:[0-9]+]]
|
||||
// CHECK-SPIRV-NEXT: ret void
|
||||
// CHECK-SPIRV-NEXT: }
|
||||
|
||||
// CHECK-AMDGCN: #[[AMDGCN_ATTR0]] = { convergent mustprogress noinline norecurse nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
|
||||
// CHECK-AMDGCN: #[[AMDGCN_ATTR1]] = { convergent nounwind }
|
||||
//
|
||||
|
||||
@ -0,0 +1,95 @@
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fcxx-exceptions -fexceptions -fsycl-is-host -emit-llvm -o - %s | FileCheck %s
|
||||
|
||||
// Validate generation of exception handling code for functions declared
|
||||
// with the sycl_kernel_entry_point attribute that implicitly call a
|
||||
// sycl_kernel_launch function that may throw an exception. Exception
|
||||
// handling is not relevant for the generated offload kernel entry point
|
||||
// function, so device compilation is intentionally not exercised.
|
||||
|
||||
// A unique kernel name type is required for each declared kernel entry point.
|
||||
template<int> struct KN;
|
||||
|
||||
// A generic kernel object type.
|
||||
template<int, int = 0>
|
||||
struct KT {
|
||||
void operator()() const;
|
||||
};
|
||||
|
||||
|
||||
// Validate that exception handling instructions are omitted when a
|
||||
// potentially throwing sycl_kernel_entry_point attributed function
|
||||
// calls a potentially throwing sycl_kernel_launch function (a thrown
|
||||
// exception will propagate with no explicit handling required).
|
||||
namespace ns1 {
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...);
|
||||
[[clang::sycl_kernel_entry_point(KN<1>)]]
|
||||
void skep(KT<1> k) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
// CHECK: ; Function Attrs: mustprogress noinline optnone
|
||||
// CHECK: define dso_local void @_ZN3ns14skepE2KTILi1ELi0EE() #{{[0-9]+}} {
|
||||
// CHECK: call void @_ZN3ns118sycl_kernel_launchI2KNILi1EEJ2KTILi1ELi0EEEEEvPKcDpT0_(ptr noundef @.str)
|
||||
// CHECK: ret void
|
||||
// CHECK: }
|
||||
|
||||
|
||||
// Validate that exception handling instructions are emitted when a
|
||||
// non-throwing sycl_kernel_entry_point attributed function calls
|
||||
// a potentially throwing sycl_kernel_launch function.
|
||||
namespace ns2 {
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...);
|
||||
[[clang::sycl_kernel_entry_point(KN<2>)]]
|
||||
void skep(KT<2> k) noexcept {
|
||||
k();
|
||||
}
|
||||
}
|
||||
// CHECK: ; Function Attrs: mustprogress noinline nounwind optnone
|
||||
// CHECK: define dso_local void @_ZN3ns24skepE2KTILi2ELi0EE() #{{[0-9]+}} personality ptr @__gxx_personality_v0 {
|
||||
// CHECK: invoke void @_ZN3ns218sycl_kernel_launchI2KNILi2EEJ2KTILi2ELi0EEEEEvPKcDpT0_(ptr noundef @.str.1)
|
||||
// CHECK: to label %invoke.cont unwind label %terminate.lpad
|
||||
// CHECK: invoke.cont:
|
||||
// CHECK: ret void
|
||||
// CHECK: terminate.lpad:
|
||||
// CHECK: call void @__clang_call_terminate(ptr %1) #{{[0-9]+}}
|
||||
// CHECK: unreachable
|
||||
// CHECK: }
|
||||
|
||||
|
||||
// Validate that exception handling instructions are omitted when a
|
||||
// potentially throwing sycl_kernel_entry_point attributed function
|
||||
// calls a non-throwing sycl_kernel_launch function (a thrown
|
||||
// exception will terminate within sycl_kernel_launch).
|
||||
namespace ns3 {
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...) noexcept;
|
||||
[[clang::sycl_kernel_entry_point(KN<3>)]]
|
||||
void skep(KT<3> k) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
// CHECK: ; Function Attrs: mustprogress noinline nounwind optnone
|
||||
// CHECK: define dso_local void @_ZN3ns34skepE2KTILi3ELi0EE() #{{[0-9]+}} {
|
||||
// CHECK: call void @_ZN3ns318sycl_kernel_launchI2KNILi3EEJ2KTILi3ELi0EEEEEvPKcDpT0_(ptr noundef @.str.2)
|
||||
// CHECK: ret void
|
||||
// CHECK: }
|
||||
|
||||
|
||||
// Validate that exception handling instructions are omitted when a
|
||||
// non-throwing sycl_kernel_entry_point attributed function calls a
|
||||
// non-throwing sycl_kernel_launch function.
|
||||
namespace ns4 {
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...) noexcept;
|
||||
[[clang::sycl_kernel_entry_point(KN<4>)]]
|
||||
void skep(KT<4> k) noexcept {
|
||||
k();
|
||||
}
|
||||
}
|
||||
// CHECK: ; Function Attrs: mustprogress noinline nounwind optnone
|
||||
// CHECK: define dso_local void @_ZN3ns44skepE2KTILi4ELi0EE() #{{[0-9]+}} {
|
||||
// CHECK: call void @_ZN3ns418sycl_kernel_launchI2KNILi4EEJ2KTILi4ELi0EEEEEvPKcDpT0_(ptr noundef @.str.3)
|
||||
// CHECK: ret void
|
||||
// CHECK: }
|
||||
@ -1,6 +1,8 @@
|
||||
// RUN: %clang_cc1 -triple spir64-unknown-unknown -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE=addrspace(1) '
|
||||
// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsycl-is-host -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE='
|
||||
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...) {}
|
||||
|
||||
template<typename KN, typename Func>
|
||||
[[clang::sycl_kernel_entry_point(KN)]] void kernel(Func F){
|
||||
|
||||
@ -1,6 +1,9 @@
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-device -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only -fsycl-is-device -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++23 -fsyntax-only -fsycl-is-device -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-host -fcxx-exceptions -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-device -fcxx-exceptions -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only -fsycl-is-host -fcxx-exceptions -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only -fsycl-is-device -fcxx-exceptions -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++23 -fsyntax-only -fsycl-is-host -fcxx-exceptions -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++23 -fsyntax-only -fsycl-is-device -fcxx-exceptions -verify %s
|
||||
|
||||
// These tests validate appertainment for the sycl_kernel_entry_point attribute.
|
||||
|
||||
@ -37,6 +40,9 @@ struct coroutine_traits {
|
||||
// A unique kernel name type is required for each declared kernel entry point.
|
||||
template<int, int = 0> struct KN;
|
||||
|
||||
// A generic kernel launch function.
|
||||
template<typename KNT, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...) {}
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// Valid declarations.
|
||||
@ -131,6 +137,16 @@ struct S15 {
|
||||
static T ok15();
|
||||
};
|
||||
|
||||
struct S16 {
|
||||
// Non-static member function declaration.
|
||||
[[clang::sycl_kernel_entry_point(KN<16>)]]
|
||||
void ok16();
|
||||
};
|
||||
|
||||
#if __cplusplus >= 202302L
|
||||
auto ok17 = [] [[clang::sycl_kernel_entry_point(KN<17>)]] -> void {};
|
||||
#endif
|
||||
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// Invalid declarations.
|
||||
@ -163,13 +179,6 @@ struct B2 {
|
||||
static int bad2;
|
||||
};
|
||||
|
||||
struct B3 {
|
||||
// Non-static member function declaration.
|
||||
// expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute cannot be applied to a non-static member function}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<3>)]]
|
||||
void bad3();
|
||||
};
|
||||
|
||||
// expected-error@+1 {{'clang::sycl_kernel_entry_point' attribute only applies to functions}}
|
||||
namespace [[clang::sycl_kernel_entry_point(BADKN<4>)]] bad4 {}
|
||||
|
||||
@ -244,13 +253,13 @@ void bad19() {
|
||||
#endif
|
||||
|
||||
struct B20 {
|
||||
// expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute cannot be applied to a non-static member function}}
|
||||
// expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute cannot be applied to a constructor}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<20>)]]
|
||||
B20();
|
||||
};
|
||||
|
||||
struct B21 {
|
||||
// expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute cannot be applied to a non-static member function}}
|
||||
// expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute cannot be applied to a destructor}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<21>)]]
|
||||
~B21();
|
||||
};
|
||||
@ -337,11 +346,6 @@ struct B34 {
|
||||
[[noreturn]] friend void bad34() {}
|
||||
};
|
||||
|
||||
#if __cplusplus >= 202302L
|
||||
// expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute cannot be applied to a non-static member function}}
|
||||
auto bad35 = [] [[clang::sycl_kernel_entry_point(BADKN<35>)]] -> void {};
|
||||
#endif
|
||||
|
||||
#if __cplusplus >= 202302L
|
||||
// expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute only applies to functions with a non-deduced 'void' return type}}
|
||||
auto bad36 = [] [[clang::sycl_kernel_entry_point(BADKN<36>)]] static {};
|
||||
@ -373,3 +377,29 @@ struct B42 {
|
||||
// expected-warning@+1 {{declaration does not declare anything}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<42>)]];
|
||||
};
|
||||
|
||||
#if __cplusplus >= 202302L
|
||||
struct B43 {
|
||||
// expected-error@+2 {{the 'clang::sycl_kernel_entry_point' attribute cannot be applied to a function with an explicit object parameter}}
|
||||
template<typename KNT>
|
||||
[[clang::sycl_kernel_entry_point(KNT)]]
|
||||
void bad43(this B43) {}
|
||||
};
|
||||
#endif
|
||||
|
||||
#if __cplusplus >= 202302L
|
||||
struct B44 {
|
||||
// expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute cannot be applied to a function with an explicit object parameter}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<44>)]]
|
||||
void bad44(this B44);
|
||||
};
|
||||
#endif
|
||||
|
||||
#if __cplusplus >= 202302L
|
||||
template<typename KNT>
|
||||
struct B45 {
|
||||
// expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute cannot be applied to a function with an explicit object parameter}}
|
||||
[[clang::sycl_kernel_entry_point(KNT)]]
|
||||
void bad45(this B45);
|
||||
};
|
||||
#endif
|
||||
|
||||
@ -0,0 +1,142 @@
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsycl-is-host -verify=host %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsycl-is-device -verify=device %s
|
||||
|
||||
// These tests validate that a diagnostic is issued if a function declared with
|
||||
// the sycl_kernel_entry_point attribute is ODR-used from code that is emitted
|
||||
// during device compilation. Such uses are ill-formed because such functions
|
||||
// are used to define an offload kernel entry point function; they aren't
|
||||
// available for ordinary function use.
|
||||
|
||||
// host-no-diagnostics
|
||||
|
||||
// Emulate inclusion of <typeinfo>.
|
||||
namespace std {
|
||||
struct type_info {
|
||||
virtual ~type_info();
|
||||
};
|
||||
} // namespace std
|
||||
|
||||
// A generic kernel launch function.
|
||||
template<typename KernelName, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...) {}
|
||||
|
||||
// A kernel name type template.
|
||||
template<int> struct KN;
|
||||
|
||||
// SYCL kernel entry point functions. These are used to both trigger the
|
||||
// emission of a function during device compilation (but not during host
|
||||
// compilation) and to trigger a diagnostic if ODR-used from a function
|
||||
// emitted during device compilation.
|
||||
// device-note@+1 4 {{attribute is here}}
|
||||
[[clang::sycl_kernel_entry_point(KN<1>)]]
|
||||
void skep();
|
||||
struct SKL {
|
||||
// device-note@+1 6 {{attribute is here}}
|
||||
[[clang::sycl_kernel_entry_point(KN<2>)]]
|
||||
void mskep();
|
||||
// device-note@+1 6 {{attribute is here}}
|
||||
[[clang::sycl_kernel_entry_point(KN<3>)]]
|
||||
static void smskep();
|
||||
// device-note@+1 2 {{attribute is here}}
|
||||
[[clang::sycl_kernel_entry_point(KN<4>)]]
|
||||
void operator()() const;
|
||||
};
|
||||
|
||||
// A function that is emitted on the device due to usage reachable from a
|
||||
// SYCL kernel entry point function. ODR-uses of sycl_kernel_entry_point
|
||||
// attributed functions within this function require a diagnostic during
|
||||
// device compilation.
|
||||
void df() {
|
||||
// Not ODR-uses; ok.
|
||||
decltype(&skep) p1 = nullptr;
|
||||
decltype(&SKL::mskep) p2 = nullptr;
|
||||
decltype(&SKL::smskep) p3 = nullptr;
|
||||
|
||||
// Not ODR-uses; ok.
|
||||
(void)noexcept(skep());
|
||||
(void)noexcept(SKL{}.mskep());
|
||||
(void)noexcept(SKL::smskep());
|
||||
|
||||
// Not ODR-uses; ok.
|
||||
(void)typeid(&skep);
|
||||
(void)typeid(&SKL::mskep);
|
||||
(void)typeid(&SKL::smskep);
|
||||
|
||||
// device-error@+1 2 {{function 'skep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
skep();
|
||||
// device-error@+1 2 {{function 'mskep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
SKL{}.mskep();
|
||||
// device-error@+1 2 {{function 'smskep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
SKL::smskep();
|
||||
|
||||
// device-error@+1 2 {{function 'skep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
(void)&skep;
|
||||
// device-error@+1 2 {{function 'mskep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
(void)&SKL::mskep;
|
||||
// device-error@+1 2 {{function 'smskep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
(void)&SKL::smskep;
|
||||
|
||||
SKL sklo;
|
||||
// device-error@+1 2 {{function 'operator()' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
sklo();
|
||||
}
|
||||
|
||||
// device-note@+1 5 {{attribute is here}}
|
||||
[[clang::sycl_kernel_entry_point(KN<1>)]]
|
||||
void skep() {
|
||||
// device-note@+1 {{called by 'skep'}}
|
||||
df();
|
||||
// device-error@+1 {{function 'skep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
skep();
|
||||
// device-error@+1 2 {{function 'mskep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
SKL{}.mskep();
|
||||
// device-error@+1 {{function 'smskep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
SKL::smskep();
|
||||
}
|
||||
|
||||
// device-note@+1 7 {{attribute is here}}
|
||||
[[clang::sycl_kernel_entry_point(KN<2>)]]
|
||||
void SKL::mskep() {
|
||||
df();
|
||||
// device-error@+1 {{function 'skep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
skep();
|
||||
// device-error@+1 2 {{function 'mskep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
SKL{}.mskep();
|
||||
// device-error@+1 {{function 'smskep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
SKL::smskep();
|
||||
}
|
||||
|
||||
// device-note@+1 3 {{attribute is here}}
|
||||
[[clang::sycl_kernel_entry_point(KN<3>)]]
|
||||
void SKL::smskep() {
|
||||
df();
|
||||
// device-error@+1 {{function 'skep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
skep();
|
||||
// device-error@+1 2 {{function 'mskep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
SKL{}.mskep();
|
||||
// device-error@+1 {{function 'smskep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
SKL::smskep();
|
||||
}
|
||||
|
||||
[[clang::sycl_kernel_entry_point(KN<4>)]]
|
||||
void SKL::operator()() const {
|
||||
df();
|
||||
// device-error@+1 {{function 'skep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
skep();
|
||||
// device-error@+1 2 {{function 'mskep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
SKL{}.mskep();
|
||||
// device-error@+1 {{function 'smskep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
SKL::smskep();
|
||||
}
|
||||
|
||||
[[clang::sycl_external]]
|
||||
void sedf() {
|
||||
// device-note@+1 {{called by 'sedf'}}
|
||||
df();
|
||||
// device-error@+1 {{function 'skep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
skep();
|
||||
// device-error@+1 {{function 'mskep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
SKL{}.mskep();
|
||||
// device-error@+1 {{function 'smskep' cannot be used in device code because it is declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
SKL::smskep();
|
||||
}
|
||||
@ -1,4 +1,6 @@
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-host -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-device -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only -fsycl-is-host -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only -fsycl-is-device -verify %s
|
||||
|
||||
// These tests validate parsing of the sycl_kernel_entry_point argument list
|
||||
@ -8,6 +10,9 @@
|
||||
template<int> struct ST; // #ST-decl
|
||||
template<int N> using TTA = ST<N>; // #TTA-decl
|
||||
|
||||
// A generic kernel launch function.
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...) {}
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// Valid declarations.
|
||||
|
||||
@ -17,6 +17,10 @@ module M2 { header "m2.h" }
|
||||
#--- common.h
|
||||
template<int> struct KN;
|
||||
|
||||
// A generic kernel launch function.
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...) {}
|
||||
|
||||
[[clang::sycl_kernel_entry_point(KN<1>)]]
|
||||
void common_test1() {}
|
||||
|
||||
|
||||
@ -15,6 +15,10 @@
|
||||
#--- pch.h
|
||||
template<int> struct KN;
|
||||
|
||||
// A generic kernel launch function.
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...) {}
|
||||
|
||||
[[clang::sycl_kernel_entry_point(KN<1>)]]
|
||||
void pch_test1() {} // << expected previous declaration note here.
|
||||
|
||||
@ -26,11 +30,11 @@ template void pch_test2<KN<2>>();
|
||||
|
||||
#--- test.cpp
|
||||
// expected-error@+3 {{the 'clang::sycl_kernel_entry_point' kernel name argument conflicts with a previous declaration}}
|
||||
// expected-note@pch.h:4 {{previous declaration is here}}
|
||||
// expected-note@pch.h:8 {{previous declaration is here}}
|
||||
[[clang::sycl_kernel_entry_point(KN<1>)]]
|
||||
void test1() {}
|
||||
|
||||
// expected-error@+3 {{the 'clang::sycl_kernel_entry_point' kernel name argument conflicts with a previous declaration}}
|
||||
// expected-note@pch.h:8 {{previous declaration is here}}
|
||||
// expected-note@pch.h:12 {{previous declaration is here}}
|
||||
[[clang::sycl_kernel_entry_point(KN<2>)]]
|
||||
void test2() {}
|
||||
|
||||
@ -1,4 +1,6 @@
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-host -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-device -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only -fsycl-is-host -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only -fsycl-is-device -verify %s
|
||||
|
||||
// These tests validate that the kernel name type argument provided to the
|
||||
@ -7,6 +9,11 @@
|
||||
// specification.
|
||||
|
||||
struct S1;
|
||||
|
||||
// A generic kernel launch function.
|
||||
template<typename KernelName, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...) {}
|
||||
|
||||
// expected-warning@+3 {{redundant 'clang::sycl_kernel_entry_point' attribute}}
|
||||
// expected-note@+1 {{previous attribute is here}}
|
||||
[[clang::sycl_kernel_entry_point(S1),
|
||||
|
||||
@ -1,4 +1,6 @@
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-host -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-device -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only -fsycl-is-host -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only -fsycl-is-device -verify %s
|
||||
|
||||
// These tests are intended to validate that a sycl_kernel_entry_point attribute
|
||||
@ -8,6 +10,10 @@
|
||||
// attribute during instantiation of a specialization unless that specialization
|
||||
// is selected by overload resolution.
|
||||
|
||||
// A generic kernel launch function.
|
||||
template<typename KernelName, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...) {}
|
||||
|
||||
// FIXME: C++23 [temp.expl.spec]p12 states:
|
||||
// FIXME: ... Similarly, attributes appearing in the declaration of a template
|
||||
// FIXME: have no effect on an explicit specialization of that template.
|
||||
|
||||
188
clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp
Normal file
188
clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp
Normal file
@ -0,0 +1,188 @@
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsyntax-only -std=c++17 -fsycl-is-host -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsyntax-only -std=c++17 -fsycl-is-device -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsyntax-only -std=c++20 -fsycl-is-host -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsyntax-only -std=c++20 -fsycl-is-device -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsyntax-only -std=c++23 -fsycl-is-host -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsyntax-only -std=c++23 -fsycl-is-device -verify %s
|
||||
|
||||
// These tests validate diagnostics for invalid use of 'this' in the body of
|
||||
// a function declared with the sycl_kernel_entry_point attribute.
|
||||
|
||||
|
||||
template<typename T> struct remove_reference_t {
|
||||
using type = T;
|
||||
};
|
||||
template<typename T> struct remove_reference_t<T&> {
|
||||
using type = T;
|
||||
};
|
||||
|
||||
namespace std {
|
||||
struct type_info {
|
||||
virtual ~type_info();
|
||||
};
|
||||
} // namespace std
|
||||
|
||||
// A generic kernel launch function.
|
||||
template<typename KernelName, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...) {}
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// Valid declarations.
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
template<int, int=0> struct KN;
|
||||
|
||||
struct S1 {
|
||||
[[clang::sycl_kernel_entry_point(KN<1>)]] void ok1() {
|
||||
(void)sizeof(this);
|
||||
}
|
||||
};
|
||||
|
||||
struct S2 {
|
||||
[[clang::sycl_kernel_entry_point(KN<2>)]] void ok2() {
|
||||
(void)noexcept(this);
|
||||
}
|
||||
};
|
||||
|
||||
struct S3 {
|
||||
[[clang::sycl_kernel_entry_point(KN<3>)]] void ok3() {
|
||||
decltype(this) x = nullptr;
|
||||
}
|
||||
};
|
||||
|
||||
struct S4 {
|
||||
static void smf();
|
||||
[[clang::sycl_kernel_entry_point(KN<4>)]] void ok4() {
|
||||
remove_reference_t<decltype(*this)>::type::smf();
|
||||
}
|
||||
};
|
||||
|
||||
struct S5 {
|
||||
int dm;
|
||||
void mf();
|
||||
[[clang::sycl_kernel_entry_point(KN<5>)]] void ok5() {
|
||||
(void)typeid(*this); // S5 is not abstract, so 'this' is not evaluated.
|
||||
(void)typeid(dm); // 'int' is not an abstract class type; implicit 'this' is not evaluated.
|
||||
(void)typeid(mf()); // 'void' is not an abstract class type; implicit 'this' is not evaluated.
|
||||
}
|
||||
};
|
||||
|
||||
template<typename KN, bool B>
|
||||
struct S6 {
|
||||
void mf() noexcept(B);
|
||||
[[clang::sycl_kernel_entry_point(KN)]] void ok6() noexcept(noexcept(mf())) {}
|
||||
};
|
||||
template void S6<KN<6,0>, false>::ok6();
|
||||
template void S6<KN<6,1>, true>::ok6();
|
||||
|
||||
template<typename KN, bool B>
|
||||
struct S7 {
|
||||
void mf() noexcept(B);
|
||||
[[clang::sycl_kernel_entry_point(KN)]] void ok7() noexcept(noexcept(this->mf())) {}
|
||||
};
|
||||
template void S7<KN<7,0>, false>::ok7();
|
||||
template void S7<KN<7,1>, true>::ok7();
|
||||
|
||||
#if __cplusplus >= 202002L
|
||||
template<typename KN, typename T>
|
||||
struct S8 {
|
||||
void mf(T);
|
||||
[[clang::sycl_kernel_entry_point(KN)]] void ok8() requires(requires { mf(1); }) {}
|
||||
};
|
||||
template void S8<KN<8>, int>::ok8();
|
||||
|
||||
template<typename KN, typename T>
|
||||
struct S9 {
|
||||
void mf(T);
|
||||
[[clang::sycl_kernel_entry_point(KN)]] void ok9() requires(requires { this->mf(1); }) {}
|
||||
};
|
||||
template void S9<KN<9>, int>::ok9();
|
||||
#endif
|
||||
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// Invalid declarations.
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
template<int, int=0> struct BADKN;
|
||||
|
||||
// expected-error@+3 {{'this' cannot be used in a potentially evaluated expression in the body of a function declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
struct B1 {
|
||||
[[clang::sycl_kernel_entry_point(BADKN<1>)]] void bad1() {
|
||||
(void)this;
|
||||
}
|
||||
};
|
||||
|
||||
// expected-error@+4 {{'this' cannot be implicitly used in a potentially evaluated expression in the body of a function declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
struct B2 {
|
||||
int dm;
|
||||
[[clang::sycl_kernel_entry_point(BADKN<2>)]] void bad2() {
|
||||
(void)dm;
|
||||
}
|
||||
};
|
||||
|
||||
// expected-error@+4 {{'this' cannot be implicitly used in a potentially evaluated expression in the body of a function declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
struct B3 {
|
||||
void mf();
|
||||
[[clang::sycl_kernel_entry_point(BADKN<3>)]] void bad3() {
|
||||
(void)mf();
|
||||
}
|
||||
};
|
||||
|
||||
// expected-error@+4 {{'this' cannot be used in a potentially evaluated expression in the body of a function declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
struct B4 {
|
||||
virtual void vmf() = 0;
|
||||
[[clang::sycl_kernel_entry_point(BADKN<4>)]] void bad4() {
|
||||
(void)typeid(*this); // B4 is abstract, so 'this' is evaluated.
|
||||
}
|
||||
};
|
||||
|
||||
// A diagnostic is not currently issued for uninstantiated definitions. In this
|
||||
// case, a declaration is instantiated, but a definition isn't. A diagnostic
|
||||
// will be issued if a definition is instantiated (as the next test exercises).
|
||||
struct B5 {
|
||||
template<typename KN>
|
||||
[[clang::sycl_kernel_entry_point(KN)]] void bad5() {
|
||||
(void)this;
|
||||
}
|
||||
};
|
||||
extern template void B5::bad5<BADKN<5>>();
|
||||
|
||||
// expected-error@+4 {{'this' cannot be used in a potentially evaluated expression in the body of a function declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
struct B6 {
|
||||
template<typename KN>
|
||||
[[clang::sycl_kernel_entry_point(KN)]] void bad6() {
|
||||
(void)this;
|
||||
}
|
||||
};
|
||||
// expected-note@+1 {{in instantiation of function template specialization 'B6::bad6<BADKN<6>>' requested here}}
|
||||
template void B6::bad6<BADKN<6>>();
|
||||
|
||||
// A diagnostic is not currently issued for uninstantiated definitions. In this
|
||||
// case, a declaration is instantiated, but a definition isn't. A diagnostic
|
||||
// will be issued if a definition is instantiated (as the next test exercises).
|
||||
template<typename KN>
|
||||
struct B7 {
|
||||
[[clang::sycl_kernel_entry_point(KN)]] void bad7() {
|
||||
(void)this;
|
||||
}
|
||||
};
|
||||
extern template void B7<BADKN<7>>::bad7();
|
||||
|
||||
// expected-error@+4 {{'this' cannot be used in a potentially evaluated expression in the body of a function declared with the 'clang::sycl_kernel_entry_point' attribute}}
|
||||
template<typename KN>
|
||||
struct B8 {
|
||||
[[clang::sycl_kernel_entry_point(KN)]] void bad8() {
|
||||
(void)this;
|
||||
}
|
||||
};
|
||||
// expected-note@+1 {{in instantiation of member function 'B8<BADKN<8>>::bad8' requested here}}
|
||||
template void B8<BADKN<8>>::bad8();
|
||||
|
||||
#if __cplusplus >= 202302L
|
||||
struct B9 {
|
||||
// expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute cannot be applied to a function with an explicit object parameter}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<9>)]] void bad9(this B9 self) {
|
||||
(void)self;
|
||||
}
|
||||
};
|
||||
#endif
|
||||
88
clang/test/SemaSYCL/sycl-kernel-launch-ms-compat.cpp
Normal file
88
clang/test/SemaSYCL/sycl-kernel-launch-ms-compat.cpp
Normal file
@ -0,0 +1,88 @@
|
||||
// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -std=c++20 -fsyntax-only -fsycl-is-host -fms-compatibility -fcxx-exceptions -verify=host,expected %s
|
||||
// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -std=c++20 -fsyntax-only -fsycl-is-device -fms-compatibility -verify=device,expected %s
|
||||
|
||||
// Test Microsoft extensions for lookup of a sycl_kernel_launch member template
|
||||
// in a dependent base class.
|
||||
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// Valid declarations.
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
// A unique kernel name type is required for each declared kernel entry point.
|
||||
template<int> struct KN;
|
||||
|
||||
// A generic kernel object type.
|
||||
template<int>
|
||||
struct KT {
|
||||
void operator()() const;
|
||||
};
|
||||
|
||||
|
||||
namespace ok1 {
|
||||
template<typename Derived>
|
||||
struct base_handler {
|
||||
protected:
|
||||
// expected-note@+2 {{must qualify identifier to find this declaration in dependent base class}}
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...);
|
||||
};
|
||||
template<int N>
|
||||
struct handler : protected base_handler<handler<N>> {
|
||||
// A warning is issued because, in standard C++, unqualified lookup for
|
||||
// sycl_kernel_launch would not consider dependent base classes. Such
|
||||
// lookups are allowed as a Microsoft compatible extension.
|
||||
// expected-warning@+4 {{use of member 'sycl_kernel_launch' found via unqualified lookup into dependent bases of class templates is a Microsoft extension}}
|
||||
// expected-note@+2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+1 {{in implicit call to 'sycl_kernel_launch' with template argument 'KN<1>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'KT<1>') required here}}
|
||||
[[clang::sycl_kernel_entry_point(KN<1>)]]
|
||||
void skep(KT<1> k) {
|
||||
k();
|
||||
}
|
||||
};
|
||||
// expected-note@+1 {{in instantiation of member function 'ok1::handler<1>::skep' requested here}}
|
||||
template void handler<1>::skep(KT<1>);
|
||||
}
|
||||
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// Invalid declarations.
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
// A unique kernel name type is required for each declared kernel entry point.
|
||||
template<int> struct BADKN;
|
||||
|
||||
// A generic kernel object type.
|
||||
template<int>
|
||||
struct BADKT {
|
||||
void operator()() const;
|
||||
};
|
||||
|
||||
|
||||
namespace bad1 {
|
||||
template<typename Derived>
|
||||
struct base_handler {
|
||||
private:
|
||||
// expected-note@+3 {{must qualify identifier to find this declaration in dependent base class}}
|
||||
// expected-note@+2 {{declared private here}}
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...);
|
||||
};
|
||||
template<int N>
|
||||
struct handler : protected base_handler<handler<N>> {
|
||||
// In standard C++, unqualified lookup for sycl_kernel_launch would not
|
||||
// consider dependent base classes. Such lookups are allowed as a Microsoft
|
||||
// compatible extension, but access checks are still performed which makes
|
||||
// this case an error.
|
||||
// expected-warning@+5 {{use of member 'sycl_kernel_launch' found via unqualified lookup into dependent bases of class templates is a Microsoft extension}}
|
||||
// expected-note@+3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+2 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<1>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<1>') required here}}
|
||||
// expected-error@+2 {{'sycl_kernel_launch' is a private member of 'bad1::base_handler<bad1::handler<1>>'}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<1>)]]
|
||||
void skep(BADKT<1> k) {
|
||||
k();
|
||||
}
|
||||
};
|
||||
// expected-note@+1 {{in instantiation of member function 'bad1::handler<1>::skep' requested here}}
|
||||
template void handler<1>::skep(BADKT<1>);
|
||||
}
|
||||
560
clang/test/SemaSYCL/sycl-kernel-launch.cpp
Normal file
560
clang/test/SemaSYCL/sycl-kernel-launch.cpp
Normal file
@ -0,0 +1,560 @@
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-host -fcxx-exceptions -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-device -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only -fsycl-is-host -fcxx-exceptions -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only -fsycl-is-device -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++23 -fsyntax-only -fsycl-is-host -fcxx-exceptions -verify %s
|
||||
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++23 -fsyntax-only -fsycl-is-device -verify %s
|
||||
|
||||
// Test overload resolution for implicit calls to sycl_kernel_launch<KN>(...)
|
||||
// synthesized for functions declared with the sycl_kernel_entry_point
|
||||
// attribute.
|
||||
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// Valid declarations.
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
// A unique kernel name type is required for each declared kernel entry point.
|
||||
template<int, int = 0> struct KN;
|
||||
|
||||
// A generic kernel object type.
|
||||
template<int, int = 0>
|
||||
struct KT {
|
||||
void operator()() const;
|
||||
};
|
||||
|
||||
|
||||
// sycl_kernel_launch as function template at namespace scope.
|
||||
namespace ok1 {
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...);
|
||||
[[clang::sycl_kernel_entry_point(KN<1>)]]
|
||||
void skep(KT<1> k) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
|
||||
// sycl_kernel_launch as function template at namespace scope with default
|
||||
// template arguments and default function arguments..
|
||||
namespace ok2 {
|
||||
template<typename KN, typename T = int>
|
||||
void sycl_kernel_launch(const char *, KT<2>, T = 2);
|
||||
[[clang::sycl_kernel_entry_point(KN<2>)]]
|
||||
void skep(KT<2> k) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
|
||||
// sycl_kernel_launch as overload set.
|
||||
namespace ok3 {
|
||||
template<typename KN>
|
||||
void sycl_kernel_launch(const char *);
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...);
|
||||
[[clang::sycl_kernel_entry_point(KN<3>)]]
|
||||
void skep(KT<3> k) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
|
||||
// sycl_kernel_launch as static member function template.
|
||||
namespace ok4 {
|
||||
struct handler {
|
||||
private:
|
||||
template<typename KN, typename... Ts>
|
||||
static void sycl_kernel_launch(const char *, Ts...);
|
||||
public:
|
||||
[[clang::sycl_kernel_entry_point(KN<4,0>)]]
|
||||
static void skep(KT<4,0> k) {
|
||||
k();
|
||||
}
|
||||
[[clang::sycl_kernel_entry_point(KN<4,1>)]]
|
||||
void skep(KT<4,1> k) {
|
||||
k();
|
||||
}
|
||||
};
|
||||
}
|
||||
|
||||
// sycl_kernel_launch as non-static member function template.
|
||||
namespace ok5 {
|
||||
struct handler {
|
||||
private:
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...);
|
||||
public:
|
||||
[[clang::sycl_kernel_entry_point(KN<5>)]]
|
||||
void skep(KT<5> k) {
|
||||
k();
|
||||
}
|
||||
};
|
||||
}
|
||||
|
||||
#if __cplusplus >= 202302L
|
||||
// sycl_kernel_launch as non-static member function template with explicit
|
||||
// object parameter.
|
||||
namespace ok6 {
|
||||
struct handler {
|
||||
private:
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(this handler self, const char *, Ts...);
|
||||
public:
|
||||
[[clang::sycl_kernel_entry_point(KN<6>)]]
|
||||
void skep(KT<6> k) {
|
||||
k();
|
||||
}
|
||||
};
|
||||
}
|
||||
#endif
|
||||
|
||||
// sycl_kernel_launch as variable template.
|
||||
namespace ok7 {
|
||||
template<typename KN>
|
||||
struct launcher {
|
||||
template<typename... Ts>
|
||||
void operator()(const char *, Ts...);
|
||||
};
|
||||
template<typename KN>
|
||||
launcher<KN> sycl_kernel_launch;
|
||||
[[clang::sycl_kernel_entry_point(KN<7>)]]
|
||||
void skep(KT<7> k) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
|
||||
#if __cplusplus >= 202302L
|
||||
// sycl_kernel_launch as variable template with static call operator template.
|
||||
namespace ok8 {
|
||||
template<typename KN>
|
||||
struct launcher {
|
||||
template<typename... Ts>
|
||||
static void operator()(const char *, Ts...);
|
||||
};
|
||||
template<typename KN>
|
||||
launcher<KN> sycl_kernel_launch;
|
||||
[[clang::sycl_kernel_entry_point(KN<8>)]]
|
||||
void skep(KT<8> k) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
#if __cplusplus >= 202302L
|
||||
// sycl_kernel_launch as variable template with call operator template with
|
||||
// explicit object parameter.
|
||||
namespace ok9 {
|
||||
template<typename KN>
|
||||
struct launcher {
|
||||
template<typename... Ts>
|
||||
void operator()(this launcher self, const char *, Ts...);
|
||||
};
|
||||
template<typename KN>
|
||||
launcher<KN> sycl_kernel_launch;
|
||||
[[clang::sycl_kernel_entry_point(KN<9>)]]
|
||||
void skep(KT<9> k) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
// sycl_kernel_launch as base class non-static member function template.
|
||||
namespace ok10 {
|
||||
template<typename Derived>
|
||||
struct base_handler {
|
||||
protected:
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...);
|
||||
};
|
||||
struct handler : protected base_handler<handler> {
|
||||
public:
|
||||
[[clang::sycl_kernel_entry_point(KN<10>)]]
|
||||
void skep(KT<10> k) {
|
||||
k();
|
||||
}
|
||||
};
|
||||
}
|
||||
|
||||
// sycl_kernel_launch with non-reference parameters.
|
||||
namespace ok11 {
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...);
|
||||
struct move_only {
|
||||
move_only(move_only&&) = default;
|
||||
};
|
||||
[[clang::sycl_kernel_entry_point(KN<11>)]]
|
||||
void skep(KT<11> k, move_only) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
|
||||
// sycl_kernel_launch with forward reference parameters.
|
||||
namespace ok12 {
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts &&...);
|
||||
struct non_copyable {
|
||||
non_copyable(const non_copyable&) = delete;
|
||||
};
|
||||
struct non_moveable {
|
||||
non_moveable(non_moveable&&) = delete;
|
||||
};
|
||||
struct move_only {
|
||||
move_only(move_only&&) = default;
|
||||
};
|
||||
[[clang::sycl_kernel_entry_point(KN<12>)]]
|
||||
void skep(KT<12> k, non_copyable, non_moveable, move_only) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
|
||||
// ADL for sycl_kernel_launch.
|
||||
namespace ok13 {
|
||||
template<typename KN, typename KT, typename T>
|
||||
[[clang::sycl_kernel_entry_point(KN)]]
|
||||
void skep(KT k, T t) {
|
||||
k();
|
||||
}
|
||||
namespace nested {
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...);
|
||||
struct S13 {};
|
||||
}
|
||||
template void skep<KN<13>>(KT<13>, nested::S13);
|
||||
}
|
||||
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// Invalid declarations.
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
// A unique kernel name type is required for each declared kernel entry point.
|
||||
template<int, int = 0> struct BADKN;
|
||||
|
||||
// A generic kernel object type.
|
||||
template<int, int = 0>
|
||||
struct BADKT {
|
||||
void operator()() const;
|
||||
};
|
||||
|
||||
|
||||
// Undeclared sycl_kernel_launch identifier from non-template function.
|
||||
namespace bad1 {
|
||||
// expected-error@+4 {{use of undeclared identifier 'sycl_kernel_launch'}}
|
||||
// expected-note@+2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<1>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<1>') required here}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<1>)]]
|
||||
void skep(BADKT<1> k) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
|
||||
// Undeclared sycl_kernel_launch identifier from function template.
|
||||
namespace bad2 {
|
||||
// expected-error@+5 {{use of undeclared identifier 'sycl_kernel_launch'}}
|
||||
// expected-note@+3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+2 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<2>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<2>') required here}}
|
||||
template<typename KN, typename KT>
|
||||
[[clang::sycl_kernel_entry_point(KN)]]
|
||||
void skep(KT k) {
|
||||
k();
|
||||
}
|
||||
// expected-note@+1 {{in instantiation of function template specialization 'bad2::skep<BADKN<2>, BADKT<2>>' requested here}}
|
||||
template void skep<BADKN<2>>(BADKT<2>);
|
||||
}
|
||||
|
||||
// No matching function for call to sycl_kernel_launch; not a template.
|
||||
namespace bad3 {
|
||||
// expected-note@+1 {{declared as a non-template here}}
|
||||
void sycl_kernel_launch(const char *, BADKT<3>);
|
||||
// expected-error@+4 {{'sycl_kernel_launch' does not refer to a template}}
|
||||
// expected-note@+2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note@+1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<3>' required here}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<3>)]]
|
||||
void skep(BADKT<3> k) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
|
||||
// No matching function for call to sycl_kernel_launch; not enough arguments.
|
||||
namespace bad4 {
|
||||
// expected-note@+2 {{candidate function template not viable: requires 2 arguments, but 1 was provided}}
|
||||
template<typename KN, typename KT>
|
||||
void sycl_kernel_launch(const char *, KT);
|
||||
// expected-error@+5 {{no matching function for call to 'sycl_kernel_launch'}}
|
||||
// expected-note@+3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+2 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<4>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]') required here}}
|
||||
template<typename KN>
|
||||
[[clang::sycl_kernel_entry_point(KN)]]
|
||||
void skep() {}
|
||||
// expected-note@+1 {{in instantiation of function template specialization 'bad4::skep<BADKN<4>>' requested here}}
|
||||
template void skep<BADKN<4>>();
|
||||
}
|
||||
|
||||
// No matching function for call to sycl_kernel_launch; too many arguments.
|
||||
namespace bad5 {
|
||||
// expected-note@+2 {{candidate function template not viable: requires 2 arguments, but 3 were provided}}
|
||||
template<typename KN, typename KT>
|
||||
void sycl_kernel_launch(const char *, KT);
|
||||
// expected-error@+5 {{no matching function for call to 'sycl_kernel_launch'}}
|
||||
// expected-note@+3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+2 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<5>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<5>', xvalue of type 'int') required here}}
|
||||
template<typename KN, typename KT>
|
||||
[[clang::sycl_kernel_entry_point(KN)]]
|
||||
void skep(KT k, int i) {
|
||||
k();
|
||||
}
|
||||
// expected-note@+1 {{in instantiation of function template specialization 'bad5::skep<BADKN<5>, BADKT<5>>' requested here}}
|
||||
template void skep<BADKN<5>>(BADKT<5>, int);
|
||||
}
|
||||
|
||||
// No matching function for call to sycl_kernel_launch; mismatched function parameter type.
|
||||
namespace bad6 {
|
||||
// expected-note-re@+2 {{candidate function template not viable: no known conversion from 'const char[{{[0-9]*}}]' to 'int' for 1st argument}}
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(int, Ts...);
|
||||
// expected-error@+5 {{no matching function for call to 'sycl_kernel_launch'}}
|
||||
// expected-note@+3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+2 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<6>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<6>') required here}}
|
||||
template<typename KN, typename KT>
|
||||
[[clang::sycl_kernel_entry_point(KN)]]
|
||||
void skep(KT k) {
|
||||
k();
|
||||
}
|
||||
// expected-note@+1 {{in instantiation of function template specialization 'bad6::skep<BADKN<6>, BADKT<6>>' requested here}}
|
||||
template void skep<BADKN<6>>(BADKT<6>);
|
||||
}
|
||||
|
||||
// No matching function for call to sycl_kernel_launch; mismatched template parameter kind.
|
||||
namespace bad7 {
|
||||
// expected-note@+2 {{candidate template ignored: invalid explicitly-specified argument for 1st template parameter}}
|
||||
template<int, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...);
|
||||
// expected-error@+4 {{no matching function for call to 'sycl_kernel_launch'}}
|
||||
// expected-note@+2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<7>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<7>') required here}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<7>)]]
|
||||
void skep(BADKT<7> k) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
|
||||
// No matching function for call to sycl_kernel_launch; substitution failure.
|
||||
namespace bad8 {
|
||||
// expected-note@+2 {{candidate template ignored: substitution failure [with KN = BADKN<8>, KT = BADKT<8>]: no type named 'no_such_type' in 'BADKT<8>'}}
|
||||
template<typename KN, typename KT, typename T = typename KT::no_such_type>
|
||||
void sycl_kernel_launch(const char *, KT);
|
||||
// expected-error@+4 {{no matching function for call to 'sycl_kernel_launch'}}
|
||||
// expected-note@+2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<8>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<8>') required here}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<8>)]]
|
||||
void skep(BADKT<8> k) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
|
||||
// No matching function for call to sycl_kernel_launch; deduction failure.
|
||||
namespace bad9 {
|
||||
// expected-note@+2 {{candidate template ignored: couldn't infer template argument 'T'}}
|
||||
template<typename KN, typename KT, typename T>
|
||||
void sycl_kernel_launch(const char *, KT);
|
||||
// expected-error@+4 {{no matching function for call to 'sycl_kernel_launch'}}
|
||||
// expected-note@+2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<9>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<9>') required here}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<9>)]]
|
||||
void skep(BADKT<9> k) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
|
||||
// No matching function for call to sycl_kernel_launch object; mismatched function parameter type.
|
||||
namespace bad10 {
|
||||
template<typename KN>
|
||||
struct launcher {
|
||||
// expected-note-re@+2 {{candidate function template not viable: no known conversion from 'const char[{{[0-9]*}}]' to 'int' for 1st argument}}
|
||||
template<typename... Ts>
|
||||
void operator()(int, Ts...);
|
||||
};
|
||||
template<typename KN>
|
||||
launcher<KN> sycl_kernel_launch;
|
||||
// expected-error@+5 {{no matching function for call to object of type 'launcher<BADKN<10, 0>>'}}
|
||||
// expected-note@+3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+2 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<10>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<10>') required here}}
|
||||
template<typename KN, typename KT>
|
||||
[[clang::sycl_kernel_entry_point(KN)]]
|
||||
void skep(KT k) {
|
||||
k();
|
||||
}
|
||||
// expected-note@+1 {{in instantiation of function template specialization 'bad10::skep<BADKN<10>, BADKT<10>>' requested here}}
|
||||
template void skep<BADKN<10>>(BADKT<10>);
|
||||
}
|
||||
|
||||
// No matching function for call to sycl_kernel_launch object; mismatched template parameter kind.
|
||||
namespace bad11 {
|
||||
template<int KN>
|
||||
struct launcher {
|
||||
template<typename... Ts>
|
||||
void operator()(int, Ts...);
|
||||
};
|
||||
// expected-note@+1 {{template parameter is declared here}}
|
||||
template<int KN>
|
||||
launcher<KN> sycl_kernel_launch;
|
||||
// expected-error@+5 {{template argument for non-type template parameter must be an expression}}
|
||||
// expected-note@+3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note@+2 {{in implicit call to 'sycl_kernel_launch' with template argument 'KN' required here}}
|
||||
template<typename KN, typename KT>
|
||||
[[clang::sycl_kernel_entry_point(KN)]]
|
||||
void skep(KT k) {
|
||||
k();
|
||||
}
|
||||
template void skep<BADKN<11>>(BADKT<11>);
|
||||
}
|
||||
|
||||
// sycl_kernel_launch as variable template with private call operator template.
|
||||
namespace bad12 {
|
||||
template<typename KN>
|
||||
struct launcher {
|
||||
private:
|
||||
// expected-note@+2 {{declared private here}}
|
||||
template<typename... Ts>
|
||||
void operator()(const char *, Ts...);
|
||||
};
|
||||
template<typename KN>
|
||||
launcher<KN> sycl_kernel_launch;
|
||||
// expected-error@+4 {{'operator()' is a private member of 'bad12::launcher<BADKN<12>>'}}
|
||||
// expected-note@+2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<12>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<12>') required here}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<12>)]]
|
||||
void skep(BADKT<12> k) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
|
||||
// Ambiguous reference to sycl_kernel_launch.
|
||||
namespace bad13 {
|
||||
inline namespace in1 {
|
||||
// expected-note@+2 {{candidate found by name lookup is 'bad13::in1::sycl_kernel_launch'}}
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...);
|
||||
}
|
||||
inline namespace in2 {
|
||||
template<typename KN>
|
||||
struct launcher {
|
||||
template<typename KT, typename... Ts>
|
||||
void operator()(const char *, Ts...);
|
||||
};
|
||||
// expected-note@+2 {{candidate found by name lookup is 'bad13::in2::sycl_kernel_launch'}}
|
||||
template<typename KN>
|
||||
launcher<KN> sycl_kernel_launch;
|
||||
}
|
||||
// expected-error@+5 {{reference to 'sycl_kernel_launch' is ambiguous}}
|
||||
// expected-note@+3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note@+2 {{in implicit call to 'sycl_kernel_launch' with template argument 'KN' required here}}
|
||||
template<typename KN, typename KT>
|
||||
[[clang::sycl_kernel_entry_point(KN)]]
|
||||
void skep(KT k) {
|
||||
k();
|
||||
}
|
||||
template void skep<BADKN<13>>(BADKT<13>);
|
||||
}
|
||||
|
||||
// Ambiguous call to sycl_kernel_launch.
|
||||
namespace bad14 {
|
||||
// expected-note@+2 {{candidate function [with KN = BADKN<14>, KT = BADKT<14>]}}
|
||||
template<typename KN, typename KT>
|
||||
void sycl_kernel_launch(const char *, KT, signed char);
|
||||
// expected-note@+2 {{candidate function [with KN = BADKN<14>, KT = BADKT<14>]}}
|
||||
template<typename KN, typename KT>
|
||||
void sycl_kernel_launch(const char *, KT, unsigned char);
|
||||
// expected-error@+4 {{call to 'sycl_kernel_launch' is ambiguous}}
|
||||
// expected-note@+2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<14>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<14>', xvalue of type 'int') required here}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<14>)]]
|
||||
void skep(BADKT<14> k, int i) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
|
||||
// Call to member sycl_kernel_launch from non-static member.
|
||||
namespace bad15 {
|
||||
struct S {
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...);
|
||||
// expected-error@+4 {{call to non-static member function without an object argument}}
|
||||
// expected-note@+2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note@+1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<15>' required here}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<15>)]]
|
||||
static void skep(BADKT<15> k) {
|
||||
k();
|
||||
}
|
||||
};
|
||||
}
|
||||
|
||||
// sycl_kernel_launch as dependent base class non-static member function
|
||||
// template.
|
||||
namespace bad16 {
|
||||
template<typename Derived>
|
||||
struct base_handler {
|
||||
protected:
|
||||
// expected-note@+2 {{member is declared here}}
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...);
|
||||
};
|
||||
template<int N>
|
||||
struct handler : protected base_handler<handler<N>> {
|
||||
// Lookup for sycl_kernel_launch fails because lookup in dependent base
|
||||
// classes requires explicit qualification.
|
||||
// expected-error@+4 {{explicit qualification required to use member 'sycl_kernel_launch' from dependent base class}}
|
||||
// expected-note@+2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<16>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<16>') required here}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<16>)]]
|
||||
void skep(BADKT<16> k) {
|
||||
k();
|
||||
}
|
||||
};
|
||||
// expected-note@+1 {{in instantiation of member function 'bad16::handler<16>::skep' requested here}}
|
||||
template void handler<16>::skep(BADKT<16>);
|
||||
}
|
||||
|
||||
// sycl_kernel_launch with non-reference parameters and non-moveable arguments.
|
||||
namespace bad17 {
|
||||
// expected-note@+2 2 {{passing argument to parameter here}}
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(const char *, Ts...);
|
||||
struct non_copyable {
|
||||
// expected-note@+1 {{'non_copyable' has been explicitly marked deleted here}}
|
||||
non_copyable(const non_copyable&) = delete;
|
||||
};
|
||||
// expected-error@+4 {{call to deleted constructor of 'bad17::non_copyable'}}
|
||||
// expected-note@+2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<17, 0>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<17, 0>', xvalue of type 'non_copyable') required here}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<17,0>)]]
|
||||
void skep(BADKT<17,0> k, non_copyable) {
|
||||
k();
|
||||
}
|
||||
struct non_moveable {
|
||||
// expected-note@+1 {{'non_moveable' has been explicitly marked deleted here}}
|
||||
non_moveable(non_moveable&&) = delete;
|
||||
};
|
||||
// expected-error@+4 {{call to deleted constructor of 'bad17::non_moveable'}}
|
||||
// expected-note@+2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<17, 1>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<17, 1>', xvalue of type 'non_moveable') required here}}
|
||||
[[clang::sycl_kernel_entry_point(BADKN<17,1>)]]
|
||||
void skep(BADKT<17,1> k, non_moveable) {
|
||||
k();
|
||||
}
|
||||
}
|
||||
|
||||
// sycl_kernel_launch declared after use and not found by ADL.
|
||||
namespace bad18 {
|
||||
// expected-error@+5 {{call to function 'sycl_kernel_launch' that is neither visible in the template definition nor found by argument-dependent lookup}}
|
||||
// expected-note@+3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
|
||||
// expected-note-re@+2 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<18>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<18>') required here}}
|
||||
template<typename KN, typename KT>
|
||||
[[clang::sycl_kernel_entry_point(KN)]]
|
||||
void skep(KT k) {
|
||||
k();
|
||||
}
|
||||
// expected-note@+2 {{'sycl_kernel_launch' should be declared prior to the call site or in the global namespace}}
|
||||
template<typename KN, typename... Ts>
|
||||
void sycl_kernel_launch(Ts...) {}
|
||||
// expected-note@+1 {{in instantiation of function template specialization 'bad18::skep<BADKN<18>, BADKT<18>>' requested here}}
|
||||
template void skep<BADKN<18>>(BADKT<18>);
|
||||
}
|
||||
@ -383,6 +383,7 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent,
|
||||
break;
|
||||
|
||||
case Stmt::SYCLKernelCallStmtClass:
|
||||
case Stmt::UnresolvedSYCLKernelCallStmtClass:
|
||||
K = CXCursor_UnexposedStmt;
|
||||
break;
|
||||
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user