diff --git a/clang/docs/LanguageExtensions.rst b/clang/docs/LanguageExtensions.rst index 129e73345b2a2..b30b6af1c17d9 100644 --- a/clang/docs/LanguageExtensions.rst +++ b/clang/docs/LanguageExtensions.rst @@ -2438,6 +2438,30 @@ their usual pattern without any special treatment. // Computes a unique stable name for the given type. constexpr const char * __builtin_sycl_unique_stable_name( type-id ); +``__builtin_sycl_mark_kernel_name`` +----------------------------------- + +``__builtin_sycl_mark_kernel_name`` is a builtin that can be used with +``__builtin_sycl_unique_stable_name`` to make sure a kernel is properly 'marked' +as a kernel without having to instantiate a sycl_kernel function. Typically, +``__builtin_sycl_unique_stable_name`` can only be called in a constant expression +context after any kernels that would change the output have been instantiated. +This is necessary, as changing the answer to the constant expression after +evaluation isn't permitted. However, in some cases it can be useful to query the +result of ``__builtin_unique_stable_name`` after we know that the name is a kernel +name, but before we are able to instantiate the kernel itself (such as when trying +to decide between two signatures at compile time). In these cases, +``__builtin_sycl_mark_kernel_name`` can be used to mark the type as a kernel name, +ensuring that ``__builtin_unique_stable_name`` gives the correct result despite the +kernel not yet being instantiated. + +**Syntax**: + +.. code-block:: c++ + + // Marks a type as the name of a sycl kernel. + constexpr bool __builtin_sycl_mark_kernel_name( type-id ); + Multiprecision Arithmetic Builtins ---------------------------------- diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 5ffbc07161f57..b32625765065a 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -6395,8 +6395,8 @@ def warn_gnu_null_ptr_arith : Warning< "arithmetic on a null pointer treated as a cast from integer to pointer is a GNU extension">, InGroup, DefaultIgnore; def err_kernel_invalidates_sycl_unique_stable_name - : Error<"kernel instantiation changes the result of an evaluated " - "'__builtin_sycl_unique_stable_name'">; + : Error<"kernel %select{naming|instantiation}0 changes the result of an " + "evaluated '__builtin_sycl_unique_stable_name'">; def note_sycl_unique_stable_name_evaluated_here : Note<"'__builtin_sycl_unique_stable_name' evaluated here">; diff --git a/clang/include/clang/Basic/TokenKinds.def b/clang/include/clang/Basic/TokenKinds.def index eadea5c9815cb..3730bea4fb84a 100644 --- a/clang/include/clang/Basic/TokenKinds.def +++ b/clang/include/clang/Basic/TokenKinds.def @@ -710,6 +710,8 @@ KEYWORD(__builtin_bit_cast , KEYALL) KEYWORD(__builtin_available , KEYALL) KEYWORD(__builtin_sycl_unique_stable_name, KEYSYCL) +TYPE_TRAIT_1(__builtin_sycl_mark_kernel_name, SYCLMarkKernelName, KEYSYCL) + // Clang-specific keywords enabled only in testing. TESTING_KEYWORD(__unknown_anytype , KEYALL) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 43b3349284f18..01217c6e07153 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -1069,8 +1069,13 @@ class Sema final { OpaqueParser = P; } + // Marks a type as a SYCL Kernel without necessarily adding it. Additionally, + // it diagnoses if this causes any of the evaluated + // __builtin_sycl_unique_stable_name values to change. + void MarkSYCLKernel(SourceLocation NewLoc, QualType Ty, bool IsInstantiation); // Does the work necessary to deal with a SYCL kernel lambda. At the moment, - // this just marks the list of lambdas required to name the kernel. + // this just marks the list of lambdas required to name the kernel. It does + // this by dispatching to MarkSYCLKernel, so it also does the diagnostics. void AddSYCLKernelLambda(const FunctionDecl *FD); class DelayedDiagnostics; diff --git a/clang/lib/Parse/ParseExpr.cpp b/clang/lib/Parse/ParseExpr.cpp index d257051549d8c..cef69c8b6491b 100644 --- a/clang/lib/Parse/ParseExpr.cpp +++ b/clang/lib/Parse/ParseExpr.cpp @@ -893,6 +893,7 @@ class CastExpressionIdValidator final : public CorrectionCandidateCallback { /// [Clang] unary-type-trait: /// '__is_aggregate' /// '__trivially_copyable' +/// '__builtin_sycl_mark_kernel_name' /// /// binary-type-trait: /// [GNU] '__is_base_of' diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp index fb82ce9bd4798..4bf3f8f30034a 100644 --- a/clang/lib/Sema/SemaExprCXX.cpp +++ b/clang/lib/Sema/SemaExprCXX.cpp @@ -4728,6 +4728,10 @@ static bool CheckUnaryTypeTraitTypeCompleteness(Sema &S, TypeTrait UTT, return !S.RequireCompleteType( Loc, ArgTy, diag::err_incomplete_type_used_in_type_trait_expr); + + // Only the type name matters, not the completeness, so always return true. + case UTT_SYCLMarkKernelName: + return true; } } @@ -5164,6 +5168,9 @@ static bool EvaluateUnaryTypeTrait(Sema &Self, TypeTrait UTT, return !T->isIncompleteType(); case UTT_HasUniqueObjectRepresentations: return C.hasUniqueObjectRepresentations(T); + case UTT_SYCLMarkKernelName: + Self.MarkSYCLKernel(KeyLoc, T, /*IsInstantiation*/ false); + return true; } } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 73165983df9b1..6f6ead1c79804 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -5222,7 +5222,8 @@ static QualType GetSYCLKernelObjectType(const FunctionDecl *KernelCaller) { return KernelParamTy; } -void Sema::AddSYCLKernelLambda(const FunctionDecl *FD) { +void Sema::MarkSYCLKernel(SourceLocation NewLoc, QualType Ty, + bool IsInstantiation) { auto MangleCallback = [](ASTContext &Ctx, const NamedDecl *ND) -> llvm::Optional { if (const auto *RD = dyn_cast(ND)) @@ -5232,9 +5233,27 @@ void Sema::AddSYCLKernelLambda(const FunctionDecl *FD) { return 1; }; - QualType Ty = GetSYCLKernelObjectType(FD); std::unique_ptr Ctx{ItaniumMangleContext::create( Context, Context.getDiagnostics(), MangleCallback)}; llvm::raw_null_ostream Out; Ctx->mangleTypeName(Ty, Out); + + // Evaluate whether this would change any of the already evaluated + // __builtin_sycl_unique_stable_name values. + for (auto &Itr : Context.SYCLUniqueStableNameEvaluatedValues) { + const std::string &CurName = Itr.first->ComputeName(Context); + if (Itr.second != CurName) { + Diag(NewLoc, diag::err_kernel_invalidates_sycl_unique_stable_name) + << IsInstantiation; + Diag(Itr.first->getLocation(), + diag::note_sycl_unique_stable_name_evaluated_here); + // Update this so future diagnostics work correctly. + Itr.second = CurName; + } + } +} + +void Sema::AddSYCLKernelLambda(const FunctionDecl *FD) { + QualType Ty = GetSYCLKernelObjectType(FD); + MarkSYCLKernel(FD->getLocation(), Ty, /*IsInstantiation*/ true); } diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index fdbdfaff2aa65..58d962d2774af 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -773,20 +773,6 @@ static void instantiateDependentSYCLKernelAttr( // instantiation of a kernel. S.AddSYCLKernelLambda(cast(New)); - // Evaluate whether this would change any of the already evaluated - // __builtin_sycl_unique_stable_name values. - for (auto &Itr : S.Context.SYCLUniqueStableNameEvaluatedValues) { - const std::string &CurName = Itr.first->ComputeName(S.Context); - if (Itr.second != CurName) { - S.Diag(New->getLocation(), - diag::err_kernel_invalidates_sycl_unique_stable_name); - S.Diag(Itr.first->getLocation(), - diag::note_sycl_unique_stable_name_evaluated_here); - // Update this so future diagnostics work correctly. - Itr.second = CurName; - } - } - New->addAttr(Attr.clone(S.getASTContext())); } diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index f709e2ed1d5f0..eba92f1aa66a8 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -187,6 +187,7 @@ class accessor { template struct opencl_image_type; +#ifdef __SYCL_DEVICE_ONLY__ #define IMAGETY_DEFINE(dim, accessmode, amsuffix, Target, ifarray_) \ template <> \ struct opencl_image_type struct _ImageImplT { #ifdef __SYCL_DEVICE_ONLY__ diff --git a/clang/test/CodeGenSYCL/mark-kernel-name.cpp b/clang/test/CodeGenSYCL/mark-kernel-name.cpp new file mode 100644 index 0000000000000..0742245cd9216 --- /dev/null +++ b/clang/test/CodeGenSYCL/mark-kernel-name.cpp @@ -0,0 +1,28 @@ +// RUN: %clang_cc1 -triple x86_64-linux-pc -fsycl-is-host -disable-llvm-passes -fdeclare-spirv-builtins -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -triple spir64 -aux-triple x86_64-linux-pc -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s + +#include "Inputs/sycl.hpp" + +// This test validates that the use of __builtin_sycl_mark_kernel_name alters +// the code-gen'ed value of __builtin_unique_stable_name. In this case, lambda1 +// emits the unmodified version like we do typically, while lambda2 is 'marked', +// so it should follow kernel naming (that is, using the E10000 naming). Note +// that the top level kernel lambda (the E10000 in common) is automatically part +// of a kernel name, since it is passed to the kernel function (which is +// necessary so that the 'device' build actually emits the builtins. + +int main() { + + cl::sycl::kernel_single_task([]() { + auto lambda1 = []() {}; + auto lambda2 = []() {}; + + (void)__builtin_sycl_unique_stable_name(decltype(lambda1)); + // CHECK: [35 x i8] c"_ZTSZZ4mainENKUlvE10000_clEvEUlvE_\00" + + // Should change the unique-stable-name of the lambda. + (void)__builtin_sycl_mark_kernel_name(decltype(lambda2)); + (void)__builtin_sycl_unique_stable_name(decltype(lambda2)); + // CHECK: [40 x i8] c"_ZTSZZ4mainENKUlvE10000_clEvEUlvE10000_\00" + }); +} diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 88dc82a4f5764..1df9a6bcfdb98 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -228,7 +228,7 @@ template struct get_kernel_wrapper_name_t { #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel)) template -ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { +ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { // #KernelSingleTaskFunc kernelFunc(); // #KernelSingleTaskKernelFuncCall } template diff --git a/clang/test/SemaSYCL/mark-kernel-name.cpp b/clang/test/SemaSYCL/mark-kernel-name.cpp new file mode 100644 index 0000000000000..d364ca962844b --- /dev/null +++ b/clang/test/SemaSYCL/mark-kernel-name.cpp @@ -0,0 +1,54 @@ +// RUN: %clang_cc1 %s -std=c++17 -triple x86_64-linux-gnu -fsycl-is-device -verify -fsyntax-only + +#include "Inputs/sycl.hpp" + +// Test to validate that __builtin_sycl_mark_kernel_name properly updates the +// constexpr checking for __builtin_sycl_unique_stable_name. We need to make +// sure that the KernelInfo change in the library both still stays broken, and +// is then 'fixed', so the definitions below help ensure that is the case. +// We also validate that this works in the event that we have a wrapper that +// first calls for the KernelInfo type, then instantiates a kernel. + +template +struct KernelInfo { + static constexpr const char *c = __builtin_sycl_unique_stable_name(KN); // #KI_USN +}; + +template +struct FixedKernelInfo { + static constexpr bool b = __builtin_sycl_mark_kernel_name(KN); + // making 'c' dependent on 'b' is necessary to ensure 'b' gets called first. + static constexpr const char *c = b + ? __builtin_sycl_unique_stable_name(KN) + : nullptr; +}; + +template