[SYCL] Diagnose non-forward declarable kernel name types (#4945)

Add additional error emission in host mode
because only on host with help of integration header possible to
differentiate case like this:
```
int main() {
  parallel_for<class KernelName>(..);
}
```
which uses technically non-forward declarable kernel name type but
still allowed by SYCL spec, from case like this:
```
int main() {
  class KernelName;
  parallel_for<class KernelName>(..);
}
```
which should be diagnosed, since the errors are emitted in runtime.

Everything works when `KernelName` typename is used directly in
`parallel_for`, because in this case when the type `KernelName` is
forward declared by integration header, host uses `::KernelName`
typename to access kernel info.
However when `KernelName` is forward declared in non-global/namespace
scope it actually produces a separate delclaration `main::KernelName`
which is not visible for runtime code that submits kernels.

Author: Fznamznon

clang/lib/Sema/SemaSYCL.cpp

@@ -3509,9 +3509,26 @@ class SYCLKernelNameTypeVisitor
             IsInvalid = true;
             return;
           }
-          // Check if the declaration is completely defined within a
-          // function or class/struct.
-          if (Tag->isCompleteDefinition()) {
+
+          // Diagnose used types without complete definition i.e.
+          //   int main() {
+          //     class KernelName1;
+          //     parallel_for<class KernelName1>(..);
+          //   }
+          // This case can only be diagnosed during host compilation because the
+          // integration header is required to distinguish between the invalid
+          // code (above) and the following valid code:
+          //   int main() {
+          //     parallel_for<class KernelName2>(..);
+          //   }
+          // The device compiler forward declares both KernelName1 and
+          // KernelName2 in the integration header as ::KernelName1 and
+          // ::KernelName2. The problem with the former case is the additional
+          // declaration 'class KernelName1' in non-global scope. Lookup in this
+          // case will resolve to ::main::KernelName1 (instead of
+          // ::KernelName1). Since this is not visible to runtime code that
+          // submits kernels, this is invalid.
+          if (Tag->isCompleteDefinition() || S.getLangOpts().SYCLIsHost) {
             S.Diag(KernelInvocationFuncLoc,
                    diag::err_sycl_kernel_incorrectly_named)
                 << /* kernel name should be forward declarable at namespace
@@ -3561,14 +3578,20 @@ class SYCLKernelNameTypeVisitor
 
 void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc,
                                ArrayRef<const Expr *> Args) {
+  QualType KernelNameType =
+      calculateKernelNameType(getASTContext(), KernelFunc);
+  SYCLKernelNameTypeVisitor KernelNameTypeVisitor(
+      *this, Args[0]->getExprLoc(), KernelNameType,
+      IsSYCLUnnamedKernel(*this, KernelFunc));
+  KernelNameTypeVisitor.Visit(KernelNameType.getCanonicalType());
+
   // FIXME: In place until the library works around its 'host' invocation
   // issues.
   if (!LangOpts.SYCLIsDevice)
     return;
+
   const CXXRecordDecl *KernelObj =
       GetSYCLKernelObjectType(KernelFunc)->getAsCXXRecordDecl();
-  QualType KernelNameType =
-      calculateKernelNameType(getASTContext(), KernelFunc);
 
   if (!KernelObj) {
     Diag(Args[0]->getExprLoc(), diag::err_sycl_kernel_not_function_object);
@@ -3609,15 +3632,10 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc,
                                             IsSIMDKernel);
 
   KernelObjVisitor Visitor{*this};
-  SYCLKernelNameTypeVisitor KernelNameTypeVisitor(
-      *this, Args[0]->getExprLoc(), KernelNameType,
-      IsSYCLUnnamedKernel(*this, KernelFunc));
 
   DiagnosingSYCLKernel = true;
 
   // Emit diagnostics for SYCL device kernels only
-  if (LangOpts.SYCLIsDevice)
-    KernelNameTypeVisitor.Visit(KernelNameType.getCanonicalType());
   Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker, DecompMarker);
   Visitor.VisitRecordFields(KernelObj, FieldChecker, UnionChecker,
                             DecompMarker);

clang/test/CodeGenSYCL/loop_fusion_host.cpp

@@ -5,6 +5,14 @@ __attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
   kernelFunc();
 }
 
+// This test uses SYCL host only mode without integration header, so
+// forward declare used kernel name class, otherwise it will be diagnosed by
+// the diagnostic implemented in https://github.com/intel/llvm/pull/4945.
+// The error happens because in host mode it is assumed that all kernel names
+// are forward declared at global or namespace scope because of integration
+// header.
+class kernel_name_1;
+
 template <int SIZE>
 class KernelFunctor5 {
 public:

clang/test/CodeGenSYCL/stall_enable_host.cpp

@@ -2,6 +2,14 @@
 
 // Tests for IR of Intel FPGA [[intel::use_stall_enable_clusters]] function attribute on Host (no-op in IR-CodeGen for host-mode).
 
+// This test uses SYCL host only mode without integration header, so
+// forward declare used kernel name class, otherwise it will be diagnosed by
+// the diagnostic implemented in https://github.com/intel/llvm/pull/4945.
+// The error happens because in host mode it is assumed that all kernel names
+// are forward declared at global or namespace scope because of integration
+// header.
+class kernel_name_1;
+
 [[intel::use_stall_enable_clusters]] void test() {}
 
 void test1() {

clang/test/SemaSYCL/Inputs/CL/sycl/detail/defines_elementary.hpp

@@ -0,0 +1,8 @@
+#pragma once
+
+#ifndef __SYCL_DISABLE_NAMESPACE_INLINE__
+#define __SYCL_INLINE_NAMESPACE(X) inline namespace X
+#else
+#define __SYCL_INLINE_NAMESPACE(X) namespace X
+#endif // __SYCL_DISABLE_NAMESPACE_INLINE__
+#define __SYCL_DLL_LOCAL

clang/test/SemaSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp

@@ -0,0 +1,51 @@
+#pragma once
+
+#include <CL/sycl/detail/defines_elementary.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+  namespace sycl {
+  namespace detail {
+
+#ifndef __SYCL_DEVICE_ONLY__
+#define _Bool bool
+#endif
+
+  // kernel parameter kinds
+  enum class kernel_param_kind_t {
+    kind_accessor = 0,
+    kind_std_layout = 1, // standard layout object parameters
+    kind_sampler = 2,
+    kind_pointer = 3,
+    kind_specialization_constants_buffer = 4,
+    kind_stream = 5,
+    kind_invalid = 0xf, // not a valid kernel kind
+  };
+
+  // describes a kernel parameter
+  struct kernel_param_desc_t {
+    // parameter kind
+    kernel_param_kind_t kind;
+    // kind == kind_std_layout
+    //   parameter size in bytes (includes padding for structs)
+    // kind == kind_accessor
+    //   access target; possible access targets are defined in access/access.hpp
+    int info;
+    // offset of the captured value of the parameter in the lambda or function
+    // object
+    int offset;
+  };
+
+  template <class KernelNameType> struct KernelInfo {
+    static constexpr unsigned getNumParams() { return 0; }
+    static const kernel_param_desc_t &getParamDesc(int) {
+      static kernel_param_desc_t Dummy;
+      return Dummy;
+    }
+    static constexpr const char *getName() { return ""; }
+    static constexpr bool isESIMD() { return 0; }
+    static constexpr bool callsThisItem() { return false; }
+    static constexpr bool callsAnyThisFreeFunction() { return false; }
+  };
+  } // namespace detail
+  } // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

clang/test/SemaSYCL/Inputs/sycl.hpp

@@ -130,6 +130,8 @@ struct opencl_image_type;
     using type = __ocl_image##dim##d_##ifarray_##amsuffix##_t;      \
   };
 
+#ifdef __SYCL_DEVICE_ONLY__
+
 #define IMAGETY_READ_3_DIM_IMAGE       \
   IMAGETY_DEFINE(1, read, ro, image, ) \
   IMAGETY_DEFINE(2, read, ro, image, ) \
@@ -154,6 +156,8 @@ IMAGETY_WRITE_3_DIM_IMAGE
 IMAGETY_READ_2_DIM_IARRAY
 IMAGETY_WRITE_2_DIM_IARRAY
 
+#endif // __SYCL_DEVICE_ONLY__
+
 template <int dim, access::mode accessmode, access::target accesstarget>
 struct _ImageImplT {
 #ifdef __SYCL_DEVICE_ONLY__
@@ -232,60 +236,59 @@ template <typename Type> struct get_kernel_wrapper_name_t {
 #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
 template <typename KernelName, typename KernelType>
 ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { // #KernelSingleTaskFunc
+#ifdef __SYCL_DEVICE_ONLY__
   kernelFunc(); // #KernelSingleTaskKernelFuncCall
+#else
+  (void)kernelFunc;
+#endif
 }
 template <typename KernelName, typename KernelType>
 ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc, kernel_handler kh) {
+#ifdef __SYCL_DEVICE_ONLY__
   kernelFunc(kh);
+#else
+  (void)kernelFunc;
+#endif
 }
 template <typename KernelName, typename KernelType>
 ATTR_SYCL_KERNEL void kernel_parallel_for(const KernelType &kernelFunc) {
+#ifdef __SYCL_DEVICE_ONLY__
   kernelFunc();
+#else
+  (void)kernelFunc;
+#endif
 }
 template <typename KernelName, typename KernelType>
 ATTR_SYCL_KERNEL void kernel_parallel_for_work_group(const KernelType &KernelFunc, kernel_handler kh) {
+#ifdef __SYCL_DEVICE_ONLY__
   KernelFunc(group<1>(), kh);
+#else
+  (void)KernelFunc;
+#endif
 }
 
 class handler {
 public:
   template <typename KernelName = auto_name, typename KernelType>
   void single_task(const KernelType &kernelFunc) {
     using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
-#ifdef __SYCL_DEVICE_ONLY__
     kernel_single_task<NameT>(kernelFunc); // #KernelSingleTask
-#else
-    kernelFunc();
-#endif
   }
   template <typename KernelName = auto_name, typename KernelType>
   void single_task(const KernelType &kernelFunc, kernel_handler kh) {
     using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
-#ifdef __SYCL_DEVICE_ONLY__
     kernel_single_task<NameT>(kernelFunc, kh);
-#else
-    kernelFunc(kh);
-#endif
   }
   template <typename KernelName = auto_name, typename KernelType>
   void parallel_for(const KernelType &kernelObj) {
     using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
     using NameWT = typename get_kernel_wrapper_name_t<NameT>::name;
-#ifdef __SYCL_DEVICE_ONLY__
     kernel_parallel_for<NameT>(kernelObj);
-#else
-    kernelObj();
-#endif
   }
   template <typename KernelName = auto_name, typename KernelType>
   void parallel_for_work_group(const KernelType &kernelFunc, kernel_handler kh) {
     using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
-#ifdef __SYCL_DEVICE_ONLY__
     kernel_parallel_for_work_group<NameT>(kernelFunc, kh);
-#else
-    group<1> G;
-    kernelFunc(G, kh);
-#endif
   }
 };
 

clang/test/SemaSYCL/non-fwd-declarable-kernel-name.cpp

@@ -0,0 +1,48 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fno-sycl-unnamed-lambda -fsyntax-only -sycl-std=2020 -fsycl-int-header=%t.h %s
+// RUN: %clang_cc1 -fsycl-is-host -internal-isystem %S/Inputs -fno-sycl-unnamed-lambda -fsyntax-only -verify -include %t.h %s
+
+// This test verifies that incorrect kernel names are diagnosed correctly.
+
+#include "sycl.hpp"
+
+using namespace cl::sycl;
+
+// user-defined function
+void function() {
+}
+
+// user-defined struct
+struct myWrapper {
+  class insideStruct;
+};
+
+template <typename KernelName> class RandomTemplate;
+
+int main() {
+  queue q;
+
+  q.submit([&](handler &h) {
+    h.single_task<class Ok>([]() { function(); });
+  });
+  q.submit([&](handler &h) {
+    h.single_task<RandomTemplate<class Ok>>([]() { function(); });
+  });
+
+  class NotOk;
+  // expected-error@#KernelSingleTask {{'NotOk' is invalid; kernel name should be forward declarable at namespace scope}}
+  // expected-note@+2 {{in instantiation of function template specialization}}
+  q.submit([&](handler &h) {
+    h.single_task<class NotOk>([]() { function(); });
+  });
+  // expected-error@#KernelSingleTask {{'myWrapper::insideStruct' is invalid; kernel name should be forward declarable at namespace scope}}
+  // expected-note@+2 {{in instantiation of function template specialization}}
+  q.submit([&](handler &h) {
+    h.single_task<class myWrapper::insideStruct>([]() { function(); });
+  });
+  // expected-error@#KernelSingleTask {{'RandomTemplate<NotOk>' is invalid; kernel name should be forward declarable at namespace scope}}
+  // expected-note@+2 {{in instantiation of function template specialization}}
+  q.submit([&](handler &h) {
+    h.single_task<RandomTemplate<NotOk>>([]() { function(); });
+  });
+  return 0;
+}

sycl/test/warnings/sycl_2020_deprecations.cpp

@@ -6,6 +6,14 @@
 #include <CL/sycl.hpp>
 #include <sycl/ext/intel/online_compiler.hpp>
 
+// This test uses SYCL host only mode without integration header, so
+// forward declare used kernel name class, otherwise it will be diagnosed by
+// the diagnostic implemented in https://github.com/intel/llvm/pull/4945.
+// The error happens because in host mode it is assumed that all kernel names
+// are forward declared at global or namespace scope because of integration
+// header.
+class Test;
+
 int main() {
   cl_context ClCtx;
   // expected-error@+1 {{no matching constructor for initialization of 'sycl::context'}}