Merge remote-tracking branch 'intel_llvm/sycl' into llvmspirv_pulldow…

…n_ww46-47

clang/lib/CodeGen/BackendUtil.cpp

@@ -42,6 +42,7 @@
 #include "llvm/Passes/PassBuilder.h"
 #include "llvm/Passes/PassPlugin.h"
 #include "llvm/Passes/StandardInstrumentations.h"
+#include "llvm/SYCLLowerIR/ESIMDVerifier.h"
 #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h"
 #include "llvm/Support/BuryPointer.h"
 #include "llvm/Support/CommandLine.h"
@@ -856,6 +857,9 @@ void EmitAssemblyHelper::CreatePasses(legacy::PassManager &MPM,
     FPM.add(createVerifierPass());
 
   // Set up the per-module pass manager.
+  if (LangOpts.SYCLIsDevice)
+    MPM.add(createESIMDVerifierPass());
+
   if (!CodeGenOpts.RewriteMapFiles.empty())
     addSymbolRewriterPass(CodeGenOpts, &MPM);
 

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/Driver/sycl-offload-with-split.c

@@ -310,7 +310,7 @@
 // RUN:    | FileCheck %s -check-prefixes=CHK-NO-SPLIT
 // RUN:   %clang_cl -### -fsycl -fsycl-device-code-split -fsycl-device-code-split=off %s 2>&1 \
 // RUN:    | FileCheck %s -check-prefixes=CHK-NO-SPLIT
-// CHK-NO-SPLIT-NOT: sycl-post-link{{.*}} -split{{.*}}
+// CHK-NO-SPLIT-NOT: sycl-post-link{{.*}} "-split={{.*}}
 
 // Check no device code split mode is passed to sycl-post-link when -fsycl-device-code-split is not set and the target is FPGA
 // RUN:   %clang -### -fsycl -fsycl-targets=spir64_fpga-unknown-unknown %s 2>&1 | FileCheck %s -check-prefixes=CHK-NO-SPLIT

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;
+}

libclc/CMakeLists.txt

@@ -322,6 +322,8 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
 		endforeach()
 	endforeach()
 
+        # Please NOTE that variables in the foreach are not local and thus need
+        # to be reset every iteration.
 	foreach( d ${${t}_devices} )
 		# Some targets don't have a specific GPU to target
 		if( ${d} STREQUAL "none" OR ${ARCH} STREQUAL "spirv" OR ${ARCH} STREQUAL "spirv64" )
@@ -330,6 +332,8 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
 				# Disables NVVM reflection to defer to after linking
 				set( flags "SHELL:-Xclang -target-feature" "SHELL:-Xclang +ptx72"
 						 "SHELL:-march=sm_86" "SHELL:-mllvm --nvvm-reflect-enable=false")
+                        else()
+                                set ( flags )
 			endif()
 			set( arch_suffix "${t}" )
 		else()