[SYCL] Implement a builtin to mark a sycl kernel (#3894)

The unique-stable-name constraint that you can't look up the name in a
constant expression before instantiating the kernel is causing issues.

This provides a constexpr builtin so that you can mark the kernel
without having to instantiate the kernel.

Author: Erich Keane

clang/docs/LanguageExtensions.rst

@@ -2452,6 +2452,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
 ----------------------------------
 

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -6408,8 +6408,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<NullPointerArithmetic>, 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">;
 

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)
 

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;

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'

clang/lib/Sema/SemaExprCXX.cpp

@@ -4733,6 +4733,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;
   }
 }
 
@@ -5169,6 +5173,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;
   }
 }
 

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<unsigned> {
     if (const auto *RD = dyn_cast<CXXRecordDecl>(ND))
@@ -5232,9 +5233,27 @@ void Sema::AddSYCLKernelLambda(const FunctionDecl *FD) {
     return 1;
   };
 
-  QualType Ty = GetSYCLKernelObjectType(FD);
   std::unique_ptr<MangleContext> 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);
 }

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

@@ -773,20 +773,6 @@ static void instantiateDependentSYCLKernelAttr(
   // instantiation of a kernel.
   S.AddSYCLKernelLambda(cast<FunctionDecl>(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()));
 }
 

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -187,6 +187,7 @@ class accessor {
 template <int dimensions, access::mode accessmode, access::target accesstarget>
 struct opencl_image_type;
 
+#ifdef __SYCL_DEVICE_ONLY__
 #define IMAGETY_DEFINE(dim, accessmode, amsuffix, Target, ifarray_) \
   template <>                                                       \
   struct opencl_image_type<dim, access::mode::accessmode,           \
@@ -218,6 +219,8 @@ IMAGETY_WRITE_3_DIM_IMAGE
 IMAGETY_READ_2_DIM_IARRAY
 IMAGETY_WRITE_2_DIM_IARRAY
 
+#endif
+
 template <int dim, access::mode accessmode, access::target accesstarget>
 struct _ImageImplT {
 #ifdef __SYCL_DEVICE_ONLY__

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<class K>([]() {
+    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"
+  });
+}

clang/test/SemaSYCL/Inputs/sycl.hpp

@@ -228,7 +228,7 @@ template <typename Type> struct get_kernel_wrapper_name_t {
 
 #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
 template <typename KernelName = auto_name, typename KernelType>
-ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) {
+ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { // #KernelSingleTaskFunc
   kernelFunc(); // #KernelSingleTaskKernelFuncCall
 }
 template <typename KernelName = auto_name, typename KernelType>

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 <typename KN>
+struct KernelInfo {
+  static constexpr const char *c = __builtin_sycl_unique_stable_name(KN); // #KI_USN
+};
+
+template <typename KN>
+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 <template <typename> class KI,
+          typename KernelName,
+          typename KernelType>
+void wrapper(KernelType KernelFunc) {
+  (void)KI<KernelName>::c;
+  cl::sycl::kernel_single_task<KernelName>(KernelFunc); // #SingleTaskInst
+}
+
+int main() {
+  []() {
+    class KernelName1;
+    constexpr const char *C = __builtin_sycl_unique_stable_name(KernelName1);
+    // expected-error@+2 {{kernel naming changes the result of an evaluated '__builtin_sycl_unique_stable_name'}}
+    // expected-note@-2 {{'__builtin_sycl_unique_stable_name' evaluated here}}
+    (void)__builtin_sycl_mark_kernel_name(KernelName1);
+  }();
+
+  []() {
+    // expected-error@#KernelSingleTaskFunc {{kernel instantiation changes the result of an evaluated '__builtin_sycl_unique_stable_name'}}
+    // expected-note@#SingleTaskInst {{in instantiation of function template}}
+    // expected-note@+2 {{in instantiation of function template}}
+    // expected-note@#KI_USN {{'__builtin_sycl_unique_stable_name' evaluated here}}
+    wrapper<KernelInfo, class KernelName2>([]() {});
+  }();
+
+  []() {
+    wrapper<FixedKernelInfo, class KernelName3>([]() {});
+  }();
+}

sycl/include/CL/sycl/detail/kernel_desc.hpp

@@ -105,7 +105,17 @@ using make_index_sequence =
 
 template <typename T> struct KernelInfoImpl {
 private:
-  static constexpr auto n = __builtin_sycl_unique_stable_name(T);
+  // This is necessary to ensure that any kernels we get info for are properly
+  // labeled as such before we call __builtin_sycl_unique_stable_name in a
+  // constant expression, otherwise subsequent calls to a sycl_kernel function
+  // could cause the kernel name to be altered, and change the result of the
+  // builtin.
+  // Additionally, we make this a dependency of 'n' so that we can guarantee
+  // that this is evaluated first. The builtin always returns 'true', so the
+  // 'else' branch of 'n's ternary is never evaluated.
+  static constexpr bool b = __builtin_sycl_mark_kernel_name(T);
+  static constexpr auto n = b ? __builtin_sycl_unique_stable_name(T)
+                              : __builtin_sycl_unique_stable_name(T);
   template <unsigned long long... I>
   static KernelInfoData<n[I]...> impl(index_sequence<I...>) {
     return {};

sycl/test/regression/unnamed-lambda.cpp

@@ -0,0 +1,16 @@
+// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-device-only -c %s -o %t.temp
+
+// This validates that the unnamed lambda logic in the library correctly works
+// with a new implementation of __builtin_unique_stable_name, where
+// instantiation order matters.  parallel_for instantiates the KernelInfo before
+// the kernel itself, so this checks that example, which only happens when the
+// named kernel is inside another lambda.
+
+#include "CL/sycl.hpp"
+
+void foo(cl::sycl::queue queue) {
+  cl::sycl::event queue_event2 = queue.submit([&](cl::sycl::handler &cgh) {
+    cgh.parallel_for<class K1>(cl::sycl::range<1>{1},
+                                           [=](cl::sycl::item<1> id) {});
+  });
+}