diff --git a/clang/include/clang/AST/ASTNodeTraverser.h b/clang/include/clang/AST/ASTNodeTraverser.h index e74bb72571d64..0c64fb7d75df3 100644 --- a/clang/include/clang/AST/ASTNodeTraverser.h +++ b/clang/include/clang/AST/ASTNodeTraverser.h @@ -836,8 +836,10 @@ class ASTNodeTraverser 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) { diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h index 8f427427d71ed..974edc89073cc 100644 --- a/clang/include/clang/AST/RecursiveASTVisitor.h +++ b/clang/include/clang/AST/RecursiveASTVisitor.h @@ -2990,6 +2990,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, {}) @@ -3027,6 +3034,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; } diff --git a/clang/include/clang/AST/StmtSYCL.h b/clang/include/clang/AST/StmtSYCL.h index 28ace12d7916b..f3f4f040fb63c 100644 --- a/clang/include/clang/AST/StmtSYCL.h +++ b/clang/include/clang/AST/StmtSYCL.h @@ -28,35 +28,45 @@ 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. + /// Retrieve the original statement. CompoundStmt *getOriginalStmt() { return cast(OriginalStmt); } const CompoundStmt *getOriginalStmt() const { return cast(OriginalStmt); } + + /// Set the original statement. void setOriginalStmt(CompoundStmt *CS) { OriginalStmt = CS; } + /// Retrieve the kernel launch statement. + Stmt *getKernelLaunchStmt() { return KernelLaunchStmt; } + const Stmt *getKernelLaunchStmt() const { return KernelLaunchStmt; } + + /// Set the kernel launch statement. + void setKernelLaunchStmt(Stmt *S) { KernelLaunchStmt = S; } + /// Retrieve the outlined function declaration. OutlinedFunctionDecl *getOutlinedFunctionDecl() { return OFDecl; } const OutlinedFunctionDecl *getOutlinedFunctionDecl() const { return OFDecl; } @@ -89,6 +99,70 @@ class SYCLKernelCallStmt : public Stmt { } }; +// 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) {} + + /// Set the original statement. + void setOriginalStmt(CompoundStmt *CS) { OriginalStmt = CS; } + + /// Set the kernel launch ID expression. + 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); + } + + /// Retrieve the original statement. + CompoundStmt *getOriginalStmt() { return cast(OriginalStmt); } + const CompoundStmt *getOriginalStmt() const { + return cast(OriginalStmt); + } + + /// Retrieve the kernel launch ID expression. + 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 diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index f1dbd8af6093a..82d92437d8deb 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -532,25 +532,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 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 ``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 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. @@ -565,73 +566,84 @@ follows. namespace sycl { class handler { - template - [[ clang::sycl_kernel_entry_point(KernelNameType) ]] - static void kernel_entry_point(KernelType kernel) { - kernel(); + template + 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 + [[ 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 - void single_task(KernelType kernel) { - // Call kernel_entry_point() to trigger generation of an offload - // kernel entry point. - kernel_entry_point(kernel); - // Call functions appropriate for the desired offload backend - // (OpenCL, CUDA, HIP, Level Zero, etc...). + template + 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(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 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. - -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. - -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 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 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()``. + +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 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++ @@ -642,65 +654,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()``. + +#. The parameter and the call to ``kernelFunc()`` in the function body + correspond to the definition of ``kernel_entry_point()`` as called by + ``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 is type checked for conformance with the SYCL device + copyability and kernel parameter requirements. -#. 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. +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. -#. 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. +.. code-block:: c++ + + sycl_kernel_launch("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. }]; } diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 3e864475f22a1..0dec5dac3f1eb 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -13062,13 +13062,18 @@ 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" - " %select{non-static member function|variadic function|deleted function|" - "defaulted function|constexpr function|consteval function|" - "function declared with the 'noreturn' attribute|coroutine|" - "function defined with a function try block}1">; + " %select{variadic function|deleted function|defaulted function|" + "constructor|destructor|coroutine|" + "constexpr function|consteval function|" + "function declared with the 'noreturn' attribute|" + "function defined with a function try block|" + "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< @@ -13084,6 +13089,12 @@ 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_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 warn_cuda_maxclusterrank_sm_90 : Warning< "maxclusterrank requires sm_90 or higher, CUDA arch provided: %0, ignoring " diff --git a/clang/include/clang/Basic/StmtNodes.td b/clang/include/clang/Basic/StmtNodes.td index bf3686bb372d5..b681d3e086cd2 100644 --- a/clang/include/clang/Basic/StmtNodes.td +++ b/clang/include/clang/Basic/StmtNodes.td @@ -23,6 +23,7 @@ def CaseStmt : StmtNode; def DefaultStmt : StmtNode; def CapturedStmt : StmtNode; def SYCLKernelCallStmt : StmtNode; +def UnresolvedSYCLKernelCallStmt : StmtNode; // Break/continue. def LoopControlStmt : StmtNode; diff --git a/clang/include/clang/Sema/ScopeInfo.h b/clang/include/clang/Sema/ScopeInfo.h index 4f4d38c961140..f334f58ebd0a7 100644 --- a/clang/include/clang/Sema/ScopeInfo.h +++ b/clang/include/clang/Sema/ScopeInfo.h @@ -245,6 +245,10 @@ class FunctionScopeInfo { /// The set of GNU address of label extension "&&label". llvm::SmallVector 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. /// diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 0470645a9e7ad..5b2ea0b39b55f 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13159,6 +13159,14 @@ class Sema final : public SemaBase { /// 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. @@ -13510,6 +13518,20 @@ class Sema final : public SemaBase { } }; + /// 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 diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h index 7ae556da2bec1..5018c1093b5c6 100644 --- a/clang/include/clang/Sema/SemaSYCL.h +++ b/clang/include/clang/Sema/SemaSYCL.h @@ -66,7 +66,32 @@ class SemaSYCL : public SemaBase { 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 diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h index 5d09d5536e5ab..32b9803a2cc53 100644 --- a/clang/include/clang/Serialization/ASTBitCodes.h +++ b/clang/include/clang/Serialization/ASTBitCodes.h @@ -1615,6 +1615,9 @@ enum StmtCode { /// A SYCLKernelCallStmt record. STMT_SYCLKERNELCALL, + /// A SYCLKernelCallStmt record. + STMT_UNRESOLVED_SYCL_KERNEL_CALL, + /// A GCC-style AsmStmt record. STMT_GCCASM, diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp index ff8ca01ec5477..efc7b9d73bfb6 100644 --- a/clang/lib/AST/StmtPrinter.cpp +++ b/clang/lib/AST/StmtPrinter.cpp @@ -595,7 +595,7 @@ void StmtPrinter::VisitCapturedStmt(CapturedStmt *Node) { } void StmtPrinter::VisitSYCLKernelCallStmt(SYCLKernelCallStmt *Node) { - PrintStmt(Node->getOutlinedFunctionDecl()->getBody()); + PrintStmt(Node->getOriginalStmt()); } void StmtPrinter::VisitObjCAtTryStmt(ObjCAtTryStmt *Node) { @@ -1442,6 +1442,11 @@ void StmtPrinter::VisitSYCLUniqueStableNameExpr( OS << ")"; } +void StmtPrinter::VisitUnresolvedSYCLKernelCallStmt( + UnresolvedSYCLKernelCallStmt *Node) { + PrintStmt(Node->getOriginalStmt()); +} + void StmtPrinter::VisitPredefinedExpr(PredefinedExpr *Node) { OS << PredefinedExpr::getIdentKindName(Node->getIdentKind()); } diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp index 4a8c638c85331..de9628128ba94 100644 --- a/clang/lib/AST/StmtProfile.cpp +++ b/clang/lib/AST/StmtProfile.cpp @@ -1402,6 +1402,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())); diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp index 36be3295950b8..eb62bff00f68c 100644 --- a/clang/lib/CodeGen/CGStmt.cpp +++ b/clang/lib/CodeGen/CGStmt.cpp @@ -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" @@ -103,6 +104,7 @@ void CodeGenFunction::EmitStmt(const Stmt *S, ArrayRef 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(*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(*S)); break; } return true; diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h index 8c4c1c8c2dc95..b64d84c3f7fb6 100644 --- a/clang/lib/CodeGen/CodeGenFunction.h +++ b/clang/lib/CodeGen/CodeGenFunction.h @@ -3643,6 +3643,8 @@ class CodeGenFunction : public CodeGenTypeCache { 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); diff --git a/clang/lib/CodeGen/CodeGenSYCL.cpp b/clang/lib/CodeGen/CodeGenSYCL.cpp index 7d66d96ad0a1b..5ad5671250a92 100644 --- a/clang/lib/CodeGen/CodeGenSYCL.cpp +++ b/clang/lib/CodeGen/CodeGenSYCL.cpp @@ -17,6 +17,21 @@ using namespace clang; using namespace CodeGen; +void CodeGenFunction::EmitSYCLKernelCallStmt(const SYCLKernelCallStmt &S) { + if (getLangOpts().SYCLIsDevice) { + // A sycl_kernel_entry_point attributed function is unlikely to be emitted + // during device compilation, but might be if it is ODR-used from device + // code that is emitted. In these cases, the function is emitted with an + // empty body; the original body is emitted in the offload kernel entry + // point and the synthesized kernel launch code is only relevant for host + // compilation. + return; + } + + assert(getLangOpts().SYCLIsHost); + EmitStmt(S.getKernelLaunchStmt()); +} + static void SetSYCLKernelAttributes(llvm::Function *Fn, CodeGenFunction &CGF) { // SYCL 2020 device language restrictions require forward progress and // disallow recursion. diff --git a/clang/lib/Frontend/FrontendActions.cpp b/clang/lib/Frontend/FrontendActions.cpp index 3595bbc6c9b9e..fde47443de26f 100644 --- a/clang/lib/Frontend/FrontendActions.cpp +++ b/clang/lib/Frontend/FrontendActions.cpp @@ -476,6 +476,10 @@ class DefaultTemplateInstCallback : public TemplateInstantiationCallback { return "TypeAliasTemplateInstantiation"; case CodeSynthesisContext::PartialOrderingTTP: return "PartialOrderingTTP"; + case CodeSynthesisContext::SYCLKernelLaunchLookup: + return "SYCLKernelLaunchLookup"; + case CodeSynthesisContext::SYCLKernelLaunchOverloadResolution: + return "SYCLKernelLaunchOverloadResolution"; } return ""; } diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 086dd8ba1c670..50918bb2b1955 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -16227,6 +16227,32 @@ Decl *Sema::ActOnStartOfFunctionDef(Scope *FnBodyScope, Decl *D, maybeAddDeclWithEffects(FD); + if (FD && !FD->isInvalidDecl() && FD->hasAttr() && + 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(); + 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; } @@ -16412,28 +16438,53 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body, bool IsInstantiation, FD->getAttr(); if (FD->isDefaulted()) { Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid) - << SKEPAttr << /*defaulted function*/ 3; + << SKEPAttr << /*defaulted function*/ 2; SKEPAttr->setInvalidAttr(); } else if (FD->isDeleted()) { Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid) - << SKEPAttr << /*deleted function*/ 2; + << SKEPAttr << /*deleted function*/ 1; SKEPAttr->setInvalidAttr(); } else if (FSI->isCoroutine()) { Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid) - << SKEPAttr << /*coroutine*/ 7; + << SKEPAttr << /*coroutine*/ 5; SKEPAttr->setInvalidAttr(); } else if (Body && isa(Body)) { Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid) - << SKEPAttr << /*function defined with a function try block*/ 8; + << SKEPAttr << /*function defined with a function try block*/ 9; SKEPAttr->setInvalidAttr(); } - if (Body && !FD->isTemplated() && !SKEPAttr->isInvalidAttr()) { - StmtResult SR = - SYCL().BuildSYCLKernelCallStmt(FD, cast(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(Body), getCurFunction()->SYCLKernelLaunchIdExpr); + } else { + SR = SYCL().BuildSYCLKernelCallStmt( + FD, cast(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(); } } diff --git a/clang/lib/Sema/SemaExceptionSpec.cpp b/clang/lib/Sema/SemaExceptionSpec.cpp index a0483c3027199..8477161c38b02 100644 --- a/clang/lib/Sema/SemaExceptionSpec.cpp +++ b/clang/lib/Sema/SemaExceptionSpec.cpp @@ -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(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: @@ -1430,7 +1443,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: diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 67f3856c10615..95a0bff66b31b 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -315,43 +315,53 @@ void SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) { } } + if (isa(FD)) { + Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid) + << SKEPAttr << /*constructor*/ 3; + SKEPAttr->setInvalidAttr(); + } + if (isa(FD)) { + Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid) + << SKEPAttr << /*destructor*/ 4; + SKEPAttr->setInvalidAttr(); + } if (const auto *MD = dyn_cast(FD)) { - if (!MD->isStatic()) { + if (MD->isExplicitObjectMemberFunction()) { Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid) - << SKEPAttr << /*non-static member function*/ 0; + << SKEPAttr << /*function with an explicit object parameter*/ 10; SKEPAttr->setInvalidAttr(); } } if (FD->isVariadic()) { Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid) - << SKEPAttr << /*variadic function*/ 1; + << SKEPAttr << /*variadic function*/ 0; SKEPAttr->setInvalidAttr(); } if (FD->isDefaulted()) { Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid) - << SKEPAttr << /*defaulted function*/ 3; + << SKEPAttr << /*defaulted function*/ 2; SKEPAttr->setInvalidAttr(); } else if (FD->isDeleted()) { Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid) - << SKEPAttr << /*deleted function*/ 2; + << SKEPAttr << /*deleted function*/ 1; SKEPAttr->setInvalidAttr(); } if (FD->isConsteval()) { Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid) - << SKEPAttr << /*consteval function*/ 5; + << SKEPAttr << /*consteval function*/ 7; SKEPAttr->setInvalidAttr(); } else if (FD->isConstexpr()) { Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid) - << SKEPAttr << /*constexpr function*/ 4; + << SKEPAttr << /*constexpr function*/ 6; SKEPAttr->setInvalidAttr(); } if (FD->isNoReturn()) { Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid) - << SKEPAttr << /*function declared with the 'noreturn' attribute*/ 6; + << SKEPAttr << /*function declared with the 'noreturn' attribute*/ 8; SKEPAttr->setInvalidAttr(); } @@ -387,8 +397,162 @@ 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); + + // 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 &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); + + // 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 Stmts; + // IdExpr may be null if name lookup failed. + if (IdExpr) { + llvm::SmallVector 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 +563,10 @@ class OutlinedFunctionDeclBodyInstantiator public: using ParmDeclMap = llvm::DenseMap; - OutlinedFunctionDeclBodyInstantiator(Sema &S, ParmDeclMap &M) + OutlinedFunctionDeclBodyInstantiator(Sema &S, ParmDeclMap &M, + FunctionDecl *FD) : TreeTransform(S), SemaRef(S), - MapRef(M) {} + MapRef(M), FD(FD) {} // A new set of AST nodes is always required. bool AlwaysRebuild() { return true; } @@ -427,18 +592,61 @@ class OutlinedFunctionDeclBodyInstantiator 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(); + } + 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); const auto *SKEPAttr = FD->getAttr(); assert(SKEPAttr && "Missing sycl_kernel_entry_point attribute"); @@ -451,29 +659,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); +} diff --git a/clang/lib/Sema/SemaTemplateInstantiate.cpp b/clang/lib/Sema/SemaTemplateInstantiate.cpp index 35205f40cbcef..e00a5b9bd4aec 100644 --- a/clang/lib/Sema/SemaTemplateInstantiate.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp @@ -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,25 @@ static std::string convertCallArgsToString(Sema &S, return Result; } +static std::string +printCallArgsValueCategoryAndType(Sema &S, llvm::ArrayRef 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 +1280,31 @@ void Sema::PrintInstantiationStack(InstantiationContextDiagFuncRef DiagFunc) { << /*isTemplateTemplateParam=*/true << Active->InstantiationRange); break; + case CodeSynthesisContext::SYCLKernelLaunchLookup: { + const auto *SKEPAttr = + Active->Entity->getAttr(); + 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_kernel_launch_lookup_here) + << SKEPAttr->getKernelName()); + break; + } + case CodeSynthesisContext::SYCLKernelLaunchOverloadResolution: { + const auto *SKEPAttr = + Active->Entity->getAttr(); + 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_kernel_launch_overload_resolution_here) + << SKEPAttr->getKernelName() + << printCallArgsValueCategoryAndType( + *this, llvm::ArrayRef(Active->CallArgs, + Active->NumCallArgs))); + break; + } } } } diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index c2491489f40d2..982cfe0b03c54 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -12966,6 +12966,31 @@ ExprResult TreeTransform::TransformSYCLUniqueStableNameExpr( E->getLocation(), E->getLParenLocation(), E->getRParenLocation(), NewT); } +template +StmtResult TreeTransform::TransformUnresolvedSYCLKernelCallStmt( + UnresolvedSYCLKernelCallStmt *S) { + auto *FD = cast(SemaRef.CurContext); + const auto *SKEPAttr = FD->getAttr(); + 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(SemaRef.CurContext), cast(Body.get()), + IdExpr.get()); + if (SR.isInvalid()) + return StmtError(); + + return SR; +} + template ExprResult TreeTransform::TransformPredefinedExpr(PredefinedExpr *E) { diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp index eef97a8588f0b..234bf0f30d85a 100644 --- a/clang/lib/Serialization/ASTReaderStmt.cpp +++ b/clang/lib/Serialization/ASTReaderStmt.cpp @@ -532,6 +532,7 @@ void ASTStmtReader::VisitCapturedStmt(CapturedStmt *S) { void ASTStmtReader::VisitSYCLKernelCallStmt(SYCLKernelCallStmt *S) { VisitStmt(S); S->setOriginalStmt(cast(Record.readSubStmt())); + S->setKernelLaunchStmt(cast(Record.readSubStmt())); S->setOutlinedFunctionDecl(readDeclAs()); } @@ -597,6 +598,14 @@ void ASTStmtReader::VisitSYCLUniqueStableNameExpr(SYCLUniqueStableNameExpr *E) { E->setTypeSourceInfo(Record.readTypeSourceInfo()); } +void ASTStmtReader::VisitUnresolvedSYCLKernelCallStmt( + UnresolvedSYCLKernelCallStmt *S) { + VisitStmt(S); + + S->setOriginalStmt(cast(Record.readSubStmt())); + S->setKernelLaunchIdExpr(Record.readExpr()); +} + void ASTStmtReader::VisitPredefinedExpr(PredefinedExpr *E) { VisitExpr(E); bool HasFunctionName = Record.readInt(); @@ -3182,6 +3191,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; diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp index acf345392aa1a..477a0f90a6a7e 100644 --- a/clang/lib/Serialization/ASTWriterStmt.cpp +++ b/clang/lib/Serialization/ASTWriterStmt.cpp @@ -625,6 +625,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; @@ -683,6 +684,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); diff --git a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp index 4e472b7fc38b0..da164cc1bf160 100644 --- a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp +++ b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp @@ -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: diff --git a/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp b/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp index e3ff3dea19514..c5518d9038448 100644 --- a/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp +++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp @@ -34,6 +34,8 @@ template struct K { void operator()(Ts...) const {} }; +template +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 *)' +// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *)' lvalue Function {{.*}} 'sycl_kernel_launch' {{.*}} +// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const char *' +// 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>(K<2>); // CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} KT // CHECK-NEXT: | |-FunctionDecl {{.*}} skep2 'void (KT)' // CHECK-NEXT: | | |-ParmVarDecl {{.*}} k 'KT' -// CHECK-NEXT: | | |-CompoundStmt {{.*}} -// CHECK-NEXT: | | | `-CallExpr {{.*}} '' -// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT' +// CHECK-NEXT: | | |-UnresolvedSYCLKernelCallStmt {{.*}} +// CHECK-NEXT: | | | `-CompoundStmt {{.*}} +// CHECK-NEXT: | | | `-CallExpr {{.*}} '' +// 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>(K<2>); // CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod {{.*}} 'operator()' 'void () const' // CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const K<2>' lvalue // CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'K<2>' lvalue ParmVar {{.*}} 'k' 'K<2>' +// CHECK-NEXT: | | |-CompoundStmt {{.*}} +// CHECK-NEXT: | | | `-CallExpr {{.*}} 'void' +// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} +// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, K<2>)' lvalue Function {{.*}} 'sycl_kernel_launch' {{.*}} +// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' +// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi2EE" +// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'K<2>' 'void (K<2> &&) noexcept' +// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'K<2>' xvalue +// 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>(K<3> k) { // CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} KT // CHECK-NEXT: | |-FunctionDecl {{.*}} skep3 'void (KT)' // CHECK-NEXT: | | |-ParmVarDecl {{.*}} k 'KT' -// CHECK-NEXT: | | |-CompoundStmt {{.*}} -// CHECK-NEXT: | | | `-CallExpr {{.*}} '' -// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT' +// CHECK-NEXT: | | |-UnresolvedSYCLKernelCallStmt {{.*}} +// CHECK-NEXT: | | | `-CompoundStmt {{.*}} +// CHECK-NEXT: | | | `-CallExpr {{.*}} '' +// 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>(K<3> k) { // CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod {{.*}} 'operator()' 'void () const' // CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const K<3>' lvalue // CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'K<3>' lvalue ParmVar {{.*}} 'k' 'K<3>' +// CHECK-NEXT: | | |-CompoundStmt {{.*}} +// CHECK-NEXT: | | | `-CallExpr {{.*}} +// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(const char *, K<3>)' +// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, K<3>)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, K<3>)' {{.*}} +// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' +// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi3EE" +// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'K<3>' 'void (K<3> &&) noexcept' +// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'K<3>' xvalue +// 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' // 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)' +// 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 *' +// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi4EE" +// CHECK-NEXT: | | | |-CXXConstructExpr {{.*}} 'K<4>' 'void (K<4> &&) noexcept' +// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'K<4>' xvalue +// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'K<4>' lvalue ParmVar {{.*}} 'k' 'K<4>' +// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' xvalue +// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p1' 'int' +// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' xvalue +// 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)' +// 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 *' +// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi5EE" +// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' xvalue +// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'unused1' 'int' +// CHECK-NEXT: | | | |-CXXConstructExpr {{.*}} 'K<5>' 'void (K<5> &&) noexcept' +// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'K<5>' xvalue +// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'K<5>' lvalue ParmVar {{.*}} 'k' 'K<5>' +// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' xvalue +// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'unused2' 'int' +// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' xvalue +// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p' 'int' +// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' xvalue +// 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' // 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)' +// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, S6)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, S6)' {{.*}} +// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' +// 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 // CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'S7' lvalue ParmVar {{.*}} 'k' 'S7' +// CHECK-NEXT: | | |-CompoundStmt {{.*}} +// CHECK-NEXT: | | | `-CallExpr {{.*}} 'void' +// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(const char *, S7)' +// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, S7)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, S7)' {{.*}} +// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' +// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi7EE" +// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'S7' 'void (S7 &&) noexcept' +// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'S7' xvalue +// 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)' +// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, S8)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, S8)' {{.*}} +// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' +// 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 +// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'S8' lvalue ParmVar {{.*}} 'k' 'S8' +// CHECK: | | `-OutlinedFunctionDecl {{.*}} +// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} + +class Handler { + template + void sycl_kernel_launch(const char *, Ts...) {} +public: + template + [[clang::sycl_kernel_entry_point(KNT)]] + void skep9(KT k, int a, int b) { + k(a, b); + } +}; +void foo() { + Handler H; + H.skep9>([=] (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 {{.*}} '' +// 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' +// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} 'int (int, int) const' lvalue CXXMethod {{.*}} 'operator()' 'int (int, int) const' +// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} lvalue +// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'k' {{.*}} +// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'a' 'int' +// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'b' 'int' +// CHECK-NEXT: | | | |-CompoundStmt {{.*}} +// CHECK-NEXT: | | | | `-CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: | | | | |-MemberExpr {{.*}} '' ->sycl_kernel_launch {{.*}} +// CHECK-NEXT: | | | | | `-CXXThisExpr {{.*}} 'Handler *' implicit this +// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'const char *' +// CHECK-NEXT: | | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi9EE" +// CHECK-NEXT: | | | | |-CXXConstructExpr {{.*}} +// CHECK-NEXT: | | | | | `-ImplicitCastExpr {{.*}} xvalue +// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'k' {{.*}} +// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: | | | | | `-ImplicitCastExpr {{.*}} 'int' xvalue +// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'a' 'int' +// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' xvalue +// 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' +// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int (int, int) const' lvalue CXXMethod {{.*}} 'operator()' 'int (int, int) const' +// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} lvalue +// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} lvalue ImplicitParam {{.*}} 'k' {{.*}} +// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ImplicitParam {{.*}} 'a' 'int' +// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'int' lvalue ImplicitParam {{.*}} 'b' 'int' +// CHECK-NEXT: | | `-SYCLKernelEntryPointAttr {{.*}} struct KN<9> + void the_end() {} // CHECK: `-FunctionDecl {{.*}} the_end 'void ()' diff --git a/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp b/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp index 0171f72df0b37..011f48e91c292 100644 --- a/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp +++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp @@ -28,6 +28,9 @@ // A unique kernel name type is required for each declared kernel entry point. template struct KN; +template +void sycl_kernel_launch(const char *, Ts... Args) {} + [[clang::sycl_kernel_entry_point(KN<1>)]] void skep1() { } diff --git a/clang/test/ASTSYCL/ast-print-sycl-kernel-call.cpp b/clang/test/ASTSYCL/ast-print-sycl-kernel-call.cpp new file mode 100644 index 0000000000000..5adaa367ed9c1 --- /dev/null +++ b/clang/test/ASTSYCL/ast-print-sycl-kernel-call.cpp @@ -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 + void sycl_kernel_launch(const char *, Ts...) {} + + template + [[clang::sycl_kernel_entry_point(KernelName)]] + void sycl_kernel_entry_point(KernelType kernel) { + kernel(); + } +}; +// CHECK: template void sycl_kernel_entry_point(KernelType kernel) +// CHECK-NEXT: { +// CHECK-NEXT: kernel(); +// CHECK-NEXT: } +// CHECK: template<> void sycl_kernel_entry_point((lambda at {{.*}}) kernel) +// CHECK-NEXT: { +// CHECK-NEXT: kernel(); +// CHECK-NEXT: } + +void f(sycl_kernel_launcher skl) { + skl.sycl_kernel_entry_point([]{}); +} diff --git a/clang/test/CodeGenSYCL/function-attrs.cpp b/clang/test/CodeGenSYCL/function-attrs.cpp index 81f893644bc7c..14195fa75072c 100644 --- a/clang/test/CodeGenSYCL/function-attrs.cpp +++ b/clang/test/CodeGenSYCL/function-attrs.cpp @@ -28,6 +28,9 @@ int foo() { return 1; } +template +void sycl_kernel_launch(Ts...) {} + template [[clang::sycl_kernel_entry_point(Name)]] void kernel_single_task(const Func &kernelFunc) { kernelFunc(); diff --git a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp index cd1d4d801951d..c508cc526db72 100644 --- a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp +++ b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp @@ -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 - -// 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. +// 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 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 +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 + void sycl_kernel_launch(const char *, Ts...) {} +public: + template + [[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(lambda); + kernel_single_task<\u03b4\u03c4\u03c7>([](int){}); + handler h; + copyable c{42}; + h.kernel_entry_point([=] (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 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@?0??main@@9@V1?0??2@9@@@YAXV@?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@?0??main@@9@V1?0??2@9@@@YAXPEBDV@?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@?0??main@@9@@@YAXV@?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@?0??main@@9@@@YAXPEBDV@?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@?0??2@9@@handler@@QEAAXV@?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@?0??2@9@HH@handler@@AEAAXPEBDV@?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@?0??main@@9@QEAA@XZ"(ptr noundef nonnull align 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 } // diff --git a/clang/test/CodeGenSYCL/sycl-kernel-entry-point-exceptions.cpp b/clang/test/CodeGenSYCL/sycl-kernel-entry-point-exceptions.cpp new file mode 100644 index 0000000000000..8fe7a148a2f61 --- /dev/null +++ b/clang/test/CodeGenSYCL/sycl-kernel-entry-point-exceptions.cpp @@ -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 struct KN; + +// A generic kernel object type. +template +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 + 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 + 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 + 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 + 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: } diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp index 14366a092a1fe..c298593e2f1ab 100644 --- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp +++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp @@ -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 +void sycl_kernel_launch(const char *, Ts...) {} template [[clang::sycl_kernel_entry_point(KN)]] void kernel(Func F){ diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp index 9aba284145fcb..a17e4c1d1be0a 100644 --- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp +++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp @@ -1,5 +1,8 @@ +// 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 // 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 struct KN; +// A generic kernel launch function. +template +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 + [[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 +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 diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp index 8f81fa218c171..b1c9e270a02b8 100644 --- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp +++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp @@ -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 struct ST; // #ST-decl template using TTA = ST; // #TTA-decl +// A generic kernel launch function. +template +void sycl_kernel_launch(const char *, Ts...) {} //////////////////////////////////////////////////////////////////////////////// // Valid declarations. diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-module.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-module.cpp index 8788e147a2ae4..05a660e91e82c 100644 --- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-module.cpp +++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-module.cpp @@ -17,6 +17,10 @@ module M2 { header "m2.h" } #--- common.h template struct KN; +// A generic kernel launch function. +template +void sycl_kernel_launch(const char *, Ts...) {} + [[clang::sycl_kernel_entry_point(KN<1>)]] void common_test1() {} diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-pch.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-pch.cpp index 0575a7a5a67eb..dcea60e016d12 100644 --- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-pch.cpp +++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-pch.cpp @@ -15,6 +15,10 @@ #--- pch.h template struct KN; +// A generic kernel launch function. +template +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>(); #--- 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() {} diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name.cpp index c7b83932fefe6..2abb24cde6663 100644 --- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name.cpp +++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name.cpp @@ -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 +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), diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp index 4c61570419629..b39a77bd35878 100644 --- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp +++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp @@ -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 +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. diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp new file mode 100644 index 0000000000000..c9ab242754899 --- /dev/null +++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp @@ -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 struct remove_reference_t { + using type = T; +}; +template struct remove_reference_t { + using type = T; +}; + +namespace std { +struct type_info { + virtual ~type_info(); +}; +} // namespace std + +// A generic kernell launch function. +template +void sycl_kernel_launch(const char *, Ts...) {} + +//////////////////////////////////////////////////////////////////////////////// +// Valid declarations. +//////////////////////////////////////////////////////////////////////////////// +template 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::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 +struct S6 { + void mf() noexcept(B); + [[clang::sycl_kernel_entry_point(KN)]] void ok6() noexcept(noexcept(mf())) {} +}; +template void S6, false>::ok6(); +template void S6, true>::ok6(); + +template +struct S7 { + void mf() noexcept(B); + [[clang::sycl_kernel_entry_point(KN)]] void ok7() noexcept(noexcept(this->mf())) {} +}; +template void S7, false>::ok7(); +template void S7, true>::ok7(); + +#if __cplusplus >= 202002L +template +struct S8 { + void mf(T); + [[clang::sycl_kernel_entry_point(KN)]] void ok8() requires(requires { mf(1); }) {} +}; +template void S8, int>::ok8(); + +template +struct S9 { + void mf(T); + [[clang::sycl_kernel_entry_point(KN)]] void ok9() requires(requires { this->mf(1); }) {} +}; +template void S9, int>::ok9(); +#endif + + +//////////////////////////////////////////////////////////////////////////////// +// Invalid declarations. +//////////////////////////////////////////////////////////////////////////////// + +template 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 + [[clang::sycl_kernel_entry_point(KN)]] void bad5() { + (void)this; + } +}; +extern template void B5::bad5>(); + +// 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 + [[clang::sycl_kernel_entry_point(KN)]] void bad6() { + (void)this; + } +}; +// expected-note@+1 {{in instantiation of function template specialization 'B6::bad6>' requested here}} +template void B6::bad6>(); + +// 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 +struct B7 { + [[clang::sycl_kernel_entry_point(KN)]] void bad7() { + (void)this; + } +}; +extern template void B7>::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 +struct B8 { + [[clang::sycl_kernel_entry_point(KN)]] void bad8() { + (void)this; + } +}; +// expected-note@+1 {{in instantiation of member function 'B8>::bad8' requested here}} +template void B8>::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 diff --git a/clang/test/SemaSYCL/sycl-kernel-launch-ms-compat.cpp b/clang/test/SemaSYCL/sycl-kernel-launch-ms-compat.cpp new file mode 100644 index 0000000000000..fdcfee3860c76 --- /dev/null +++ b/clang/test/SemaSYCL/sycl-kernel-launch-ms-compat.cpp @@ -0,0 +1,86 @@ +// 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 struct KN; + +// A generic kernel object type. +template +struct KT { + void operator()() const; +}; + + +namespace ok1 { + template + struct base_handler { + protected: + // expected-note@+2 {{must qualify identifier to find this declaration in dependent base class}} + template + void sycl_kernel_launch(const char *, Ts...); + }; + template + struct handler : protected base_handler> { + // 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@+3 {{use of member 'sycl_kernel_launch' found via unqualified lookup into dependent bases of class templates is a Microsoft extension}} + // 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 struct BADKN; + +// A generic kernel object type. +template +struct BADKT { + void operator()() const; +}; + + +namespace bad1 { + template + 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 + void sycl_kernel_launch(const char *, Ts...); + }; + template + struct handler : protected base_handler> { + // 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@+4 {{use of member 'sycl_kernel_launch' found via unqualified lookup into dependent bases of class templates is a Microsoft extension}} + // 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>'}} + [[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>); +} diff --git a/clang/test/SemaSYCL/sycl-kernel-launch.cpp b/clang/test/SemaSYCL/sycl-kernel-launch.cpp new file mode 100644 index 0000000000000..00f61b8c30ccd --- /dev/null +++ b/clang/test/SemaSYCL/sycl-kernel-launch.cpp @@ -0,0 +1,541 @@ +// 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(...) +// 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 struct KN; + +// A generic kernel object type. +template +struct KT { + void operator()() const; +}; + + +// sycl_kernel_launch as function template at namespace scope. +namespace ok1 { + template + 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 + 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 + void sycl_kernel_launch(const char *); + template + 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 + 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 + 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 + 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 + struct launcher { + template + void operator()(const char *, Ts...); + }; + template + launcher 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 + struct launcher { + template + static void operator()(const char *, Ts...); + }; + template + launcher 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 + struct launcher { + template + void operator()(this launcher self, const char *, Ts...); + }; + template + launcher 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 + struct base_handler { + protected: + template + void sycl_kernel_launch(const char *, Ts...); + }; + struct handler : protected base_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 + 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 + 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 + [[clang::sycl_kernel_entry_point(KN)]] + void skep(KT k, T t) { + k(); + } + namespace nested { + template + void sycl_kernel_launch(const char *, Ts...); + struct S13 {}; + } + template void skep>(KT<13>, nested::S13); +} + + +//////////////////////////////////////////////////////////////////////////////// +// Invalid declarations. +//////////////////////////////////////////////////////////////////////////////// + +// A unique kernel name type is required for each declared kernel entry point. +template struct BADKN; + +// A generic kernel object type. +template +struct BADKT { + void operator()() const; +}; + + +// Undeclared sycl_kernel_launch identifier from non-template function. +namespace bad1 { + // expected-error@+3 {{use of undeclared identifier 'sycl_kernel_launch'}} + // 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@+4 {{use of undeclared identifier 'sycl_kernel_launch'}} + // 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 + [[clang::sycl_kernel_entry_point(KN)]] + void skep(KT k) { + k(); + } + // expected-note@+1 {{in instantiation of function template specialization 'bad2::skep, BADKT<2>>' requested here}} + template void skep>(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@+3 {{'sycl_kernel_launch' does not refer to a template}} + // 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 + void sycl_kernel_launch(const char *, KT); + // expected-error@+4 {{no matching function for call to 'sycl_kernel_launch'}} + // 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 + [[clang::sycl_kernel_entry_point(KN)]] + void skep() {} + // expected-note@+1 {{in instantiation of function template specialization 'bad4::skep>' requested here}} + template void skep>(); +} + +// 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 + void sycl_kernel_launch(const char *, KT); + // expected-error@+4 {{no matching function for call to 'sycl_kernel_launch'}} + // 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 + [[clang::sycl_kernel_entry_point(KN)]] + void skep(KT k, int i) { + k(); + } + // expected-note@+1 {{in instantiation of function template specialization 'bad5::skep, BADKT<5>>' requested here}} + template void skep>(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 + void sycl_kernel_launch(int, Ts...); + // expected-error@+4 {{no matching function for call to 'sycl_kernel_launch'}} + // 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 + [[clang::sycl_kernel_entry_point(KN)]] + void skep(KT k) { + k(); + } + // expected-note@+1 {{in instantiation of function template specialization 'bad6::skep, BADKT<6>>' requested here}} + template void skep>(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 + void sycl_kernel_launch(const char *, Ts...); + // expected-error@+3 {{no matching function for call to 'sycl_kernel_launch'}} + // 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 + void sycl_kernel_launch(const char *, KT); + // expected-error@+3 {{no matching function for call to 'sycl_kernel_launch'}} + // 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 + void sycl_kernel_launch(const char *, KT); + // expected-error@+3 {{no matching function for call to 'sycl_kernel_launch'}} + // 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 + 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 + void operator()(int, Ts...); + }; + template + launcher sycl_kernel_launch; + // expected-error@+4 {{no matching function for call to object of type 'launcher>'}} + // 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 + [[clang::sycl_kernel_entry_point(KN)]] + void skep(KT k) { + k(); + } + // expected-note@+1 {{in instantiation of function template specialization 'bad10::skep, BADKT<10>>' requested here}} + template void skep>(BADKT<10>); +} + +// No matching function for call to sycl_kernel_launch object; mismatched template parameter kind. +namespace bad11 { + template + struct launcher { + template + void operator()(int, Ts...); + }; + // expected-note@+1 {{template parameter is declared here}} + template + launcher sycl_kernel_launch; + // expected-error@+4 {{template argument for non-type template parameter must be an expression}} + // expected-note@+2 {{in implicit call to 'sycl_kernel_launch' with template argument 'KN' required here}} + template + [[clang::sycl_kernel_entry_point(KN)]] + void skep(KT k) { + k(); + } + template void skep>(BADKT<11>); +} + +// sycl_kernel_launch as variable template with private call operator template. +namespace bad12 { + template + struct launcher { + private: + // expected-note@+2 {{declared private here}} + template + void operator()(const char *, Ts...); + }; + template + launcher sycl_kernel_launch; + // expected-error@+3 {{'operator()' is a private member of 'bad12::launcher>'}} + // 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 + void sycl_kernel_launch(const char *, Ts...); + } + inline namespace in2 { + template + struct launcher { + template + void operator()(const char *, Ts...); + }; + // expected-note@+2 {{candidate found by name lookup is 'bad13::in2::sycl_kernel_launch'}} + template + launcher sycl_kernel_launch; + } + // expected-error@+4 {{reference to 'sycl_kernel_launch' is ambiguous}} + // expected-note@+2 {{in implicit call to 'sycl_kernel_launch' with template argument 'KN' required here}} + template + [[clang::sycl_kernel_entry_point(KN)]] + void skep(KT k) { + k(); + } + template void skep>(BADKT<13>); +} + +// Ambiguous call to sycl_kernel_launch. +namespace bad14 { + // expected-note@+2 {{candidate function [with KN = BADKN<14>, KT = BADKT<14>]}} + template + void sycl_kernel_launch(const char *, KT, signed char); + // expected-note@+2 {{candidate function [with KN = BADKN<14>, KT = BADKT<14>]}} + template + void sycl_kernel_launch(const char *, KT, unsigned char); + // expected-error@+3 {{call to 'sycl_kernel_launch' is ambiguous}} + // 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 + void sycl_kernel_launch(const char *, Ts...); + // expected-error@+3 {{call to non-static member function without an object argument}} + // 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 + struct base_handler { + protected: + // expected-note@+2 {{member is declared here}} + template + void sycl_kernel_launch(const char *, Ts...); + }; + template + struct handler : protected base_handler> { + // Lookup for sycl_kernel_launch fails because lookup in dependent base + // classes requires explicit qualification. + // expected-error@+3 {{explicit qualification required to use member 'sycl_kernel_launch' from dependent base class}} + // 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 + 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@+3 {{call to deleted constructor of 'bad17::non_copyable'}} + // 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@+3 {{call to deleted constructor of 'bad17::non_moveable'}} + // 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@+4 {{call to function 'sycl_kernel_launch' that is neither visible in the template definition nor found by argument-dependent lookup}} + // 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 + [[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 + void sycl_kernel_launch(Ts...) {} + // expected-note@+1 {{in instantiation of function template specialization 'bad18::skep, BADKT<18>>' requested here}} + template void skep>(BADKT<18>); +} diff --git a/clang/tools/libclang/CXCursor.cpp b/clang/tools/libclang/CXCursor.cpp index 0a43d73063c1f..9dd0582d0c97e 100644 --- a/clang/tools/libclang/CXCursor.cpp +++ b/clang/tools/libclang/CXCursor.cpp @@ -376,6 +376,7 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent, break; case Stmt::SYCLKernelCallStmtClass: + case Stmt::UnresolvedSYCLKernelCallStmtClass: K = CXCursor_UnexposedStmt; break;