Merge from 'sycl' to 'sycl-web'

Author: iclsrc

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -11122,10 +11122,10 @@ def err_sycl_invalid_accessor_property_list_template_param : Error<
   "%select{parameter pack|type|non-negative integer}1">;
 def warn_sycl_pass_by_value_deprecated 
     : Warning<"Passing kernel functions by value is deprecated in SYCL 2020">,
-      InGroup<Sycl2020Compat>;
+      InGroup<Sycl2020Compat>, ShowInSystemHeader;
 def warn_sycl_pass_by_reference_future 
     : Warning<"Passing of kernel functions by reference is a SYCL 2020 extension">,
-      InGroup<Sycl2017Compat>;
+      InGroup<Sycl2017Compat>, ShowInSystemHeader;
 def warn_sycl_attibute_function_raw_ptr
     : Warning<"SYCL 1.2.1 specification does not allow %0 attribute applied "
               "to a function with a raw pointer "
@@ -11135,7 +11135,7 @@ def warn_sycl_implicit_decl
     : Warning<"SYCL 1.2.1 specification requires an explicit forward "
               "declaration for a kernel type name; your program may not "
               "be portable">,
-      InGroup<SyclStrict>, DefaultIgnore;
+      InGroup<SyclStrict>, ShowInSystemHeader, DefaultIgnore;
 def warn_sycl_restrict_recursion
     : Warning<"SYCL kernel cannot call a recursive function">,
       InGroup<SyclStrict>, DefaultError;

clang/include/clang/Driver/Options.td

@@ -498,6 +498,7 @@ def Wno_nonportable_cfstrings : Joined<["-"], "Wno-nonportable-cfstrings">, Grou
   Flags<[CC1Option]>;
 def Wnonportable_cfstrings : Joined<["-"], "Wnonportable-cfstrings">, Group<W_Group>,
   Flags<[CC1Option]>;
+def Wno_sycl_strict : Joined<["-"], "Wno-sycl-strict">, Group<W_Group>, HelpText<"Disable warnings which enforce strict SYCL language compatibility.">;
 def Wp_COMMA : CommaJoined<["-"], "Wp,">,
   HelpText<"Pass the comma separated arguments in <arg> to the preprocessor">,
   MetaVarName<"<arg>">, Group<Preprocessor_Group>;

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -281,6 +281,11 @@ ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) {
   kernelFunc();
 }
 
+template <typename KernelName = auto_name, typename KernelType>
+ATTR_SYCL_KERNEL void kernel_single_task_2017(KernelType kernelFunc) {
+  kernelFunc();
+}
+
 template <typename KernelName, typename KernelType, int Dims>
 ATTR_SYCL_KERNEL void
 kernel_parallel_for(const KernelType &KernelFunc) {
@@ -323,6 +328,16 @@ class handler {
     kernel_single_task<NameT>(kernelFunc);
 #else
     kernelFunc();
+#endif
+  }
+
+  template <typename KernelName = auto_name, typename KernelType>
+  void single_task_2017(KernelType kernelFunc) {
+    using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
+#ifdef __SYCL_DEVICE_ONLY__
+    kernel_single_task_2017<NameT>(kernelFunc);
+#else
+    kernelFunc();
 #endif
   }
 };

clang/test/CodeGenSYCL/kernel-by-reference.cpp

@@ -1,42 +1,37 @@
-// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only -sycl-std=2017 -DSYCL2017 %s
-// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only -sycl-std=2020 -DSYCL2020 %s
-// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only -Wno-sycl-strict -DNODIAG %s
-// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only -sycl-std=2020 -Wno-sycl-strict -DNODIAG %s
+// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -sycl-std=2017 -DSYCL2017 %s
+// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -sycl-std=2020 -DSYCL2020 %s
+// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -Wno-sycl-strict -DNODIAG %s
+// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -sycl-std=2020 -Wno-sycl-strict -DNODIAG %s
 
-// SYCL 1.2/2017 - kernel functions passed directly. (Also no const requirement, though mutable lambdas never supported)
-template <typename name, typename Func>
-#if defined(SYCL2020)
-// expected-warning@+2 {{Passing kernel functions by value is deprecated in SYCL 2020}}
-#endif
-__attribute__((sycl_kernel)) void sycl_2017_single_task(Func kernelFunc) {
-  kernelFunc();
-}
+#include "sycl.hpp"
 
-// SYCL 2020 - kernel functions are passed by reference.
-template <typename name, typename Func>
-#if defined(SYCL2017)
-// expected-warning@+2 {{Passing of kernel functions by reference is a SYCL 2020 extension}}
-#endif
-__attribute__((sycl_kernel)) void sycl_2020_single_task(const Func &kernelFunc) {
-  kernelFunc();
-}
+using namespace cl::sycl;
 
-int do_nothing(int i) {
+int simple_add(int i) {
   return i + 1;
 }
 
 // ensure both compile.
 int main() {
-  sycl_2017_single_task<class sycl12>([]() {
-    do_nothing(10);
+  queue q;
+#if defined(SYCL2020)
+  // expected-warning@Inputs/sycl.hpp:285 {{Passing kernel functions by value is deprecated in SYCL 2020}}
+  // expected-note@+3 {{in instantiation of function template specialization}}
+#endif
+  q.submit([&](handler &h) {
+    h.single_task_2017<class sycl2017>([]() { simple_add(10); });
   });
 
-  sycl_2020_single_task<class sycl2020>([]() {
-    do_nothing(11);
+#if defined(SYCL2017)
+  // expected-warning@Inputs/sycl.hpp:280 {{Passing of kernel functions by reference is a SYCL 2020 extension}}
+  // expected-note@+3 {{in instantiation of function template specialization}}
+#endif
+  q.submit([&](handler &h) {
+    h.single_task<class sycl2020>([]() { simple_add(11); });
   });
 
   return 0;
 }
 #if defined(NODIAG)
 // expected-no-diagnostics
-#endif
+#endif

clang/test/SemaSYCL/args-size-overflow.cpp

@@ -1,26 +1,16 @@
-// RUN: %clang_cc1 -fsycl -triple spir64 -fsycl-is-device -fsyntax-only -verify %s
-// RUN: %clang_cc1 -fsycl -triple spir64 -Werror=sycl-strict -DERROR -fsycl-is-device -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -Wsycl-strict -sycl-std=2020 -verify %s
 
-#include "Inputs/sycl.hpp"
+#include "sycl.hpp"
 class Foo;
 
-template <typename Name, typename F>
-__attribute__((sycl_kernel)) void kernel(F KernelFunc) {
-  KernelFunc();
-}
+using namespace cl::sycl;
 
-template <typename Name, typename F>
-void parallel_for(F KernelFunc) {
-#ifdef ERROR
-  // expected-error@+4 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}}
-#else
-  // expected-warning@+2 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}}
-#endif
-  kernel<Name>(KernelFunc);
-}
+queue q;
 
 using Accessor =
-    cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer>;
+    accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer>;
+
+// expected-warning@Inputs/sycl.hpp:220 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}}
 
 void use() {
   struct S {
@@ -31,6 +21,8 @@ void use() {
     int Array[1991];
   } Args;
   auto L = [=]() { (void)Args; };
-  // expected-note@+1 {{in instantiation of function template specialization 'parallel_for<Foo}}
-  parallel_for<Foo>(L);
-}
+  q.submit([&](handler &h) {
+    // expected-note@+1 {{in instantiation of function template specialization}}
+    h.single_task<class Foo>(L);
+  });
+}

clang/test/SemaSYCL/array-kernel-param.cpp

@@ -337,7 +337,7 @@ int main() {
 // CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [2][3]' lvalue
 // CHECK-NEXT: MemberExpr {{.*}} 'int [2][3]' lvalue .
 // CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class'
-// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long
+// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'int' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <ArrayToPointerDecay>
@@ -347,8 +347,8 @@ int main() {
 // CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [2][3]' lvalue
 // CHECK-NEXT: MemberExpr {{.*}} 'int [2][3]' lvalue .
 // CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class'
-// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long
-// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long
+// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned
+// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned
 
 // Check kernel_G parameters.
 // CHECK: FunctionDecl {{.*}}kernel_G{{.*}} 'void (__wrapper_class)'
@@ -369,4 +369,4 @@ int main() {
 // CHECK-NEXT: OpaqueValueExpr {{.*}} 'foo2 [2]' lvalue
 // CHECK-NEXT: MemberExpr {{.*}} 'foo2 [2]' lvalue .
 // CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class'
-// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long
+// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned

clang/test/SemaSYCL/implicit_kernel_type.cpp

@@ -1,47 +1,20 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s -Werror=sycl-strict -DERROR
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s  -Wsycl-strict -DWARN
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -verify %s  -Werror=sycl-strict
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsycl-int-header=%t.h -fsyntax-only -sycl-std=2020 -verify %s -Werror=sycl-strict -DERROR
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsycl-int-header=%t.h -fsyntax-only -sycl-std=2020 -verify %s  -Wsycl-strict -DWARN
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -sycl-std=2020 -verify %s  -Werror=sycl-strict
 
-// SYCL 1.2 Definitions
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void sycl_121_single_task(Func kernelFunc) {
-  kernelFunc();
-}
-
-class event {};
-class queue {
-public:
-  template <typename T>
-  event submit(T cgf) { return event{}; }
-};
-class auto_name {};
-template <typename Name, typename Type>
-struct get_kernel_name_t {
-  using name = Name;
-};
-class handler {
-public:
-  template <typename KernelName = auto_name, typename KernelType>
-  void single_task(KernelType kernelFunc) {
-    using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
-#ifdef __SYCL_DEVICE_ONLY__
-    sycl_121_single_task<NameT>(kernelFunc);
-#else
-    kernelFunc();
-#endif
-  }
-};
-// -- /Definitions
+#include "sycl.hpp"
 
 #ifdef __SYCL_UNNAMED_LAMBDA__
 // expected-no-diagnostics
 #endif
 
+using namespace cl::sycl;
 
+// user-defined function
 void function() {
 }
 
-// user-defined class
+// user-defined struct
 struct myWrapper {
 };
 
@@ -50,34 +23,57 @@ class myWrapper2;
 
 int main() {
   queue q;
-#ifndef __SYCL_UNNAMED_LAMBDA__
+
+#if defined(WARN)
+  // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
+  // expected-note@+7 {{InvalidKernelName1 declared here}}
+  // expected-note@+8 {{in instantiation of function template specialization}}
+#elif defined(ERROR)
+  // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
   // expected-note@+3 {{InvalidKernelName1 declared here}}
-  // expected-note@+4{{in instantiation of function template specialization}}
-  // expected-error@28 {{kernel needs to have a globally-visible name}}
+  // expected-note@+4 {{in instantiation of function template specialization}}
+#endif
   class InvalidKernelName1 {};
   q.submit([&](handler &h) {
     h.single_task<InvalidKernelName1>([]() {});
   });
-#endif
+
 #if defined(WARN)
-  // expected-warning@+6 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
-  // expected-note@+5 {{fake_kernel declared here}}
+  // expected-warning@Inputs/sycl.hpp:220 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
 #elif defined(ERROR)
-  // expected-error@+3 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
-  // expected-note@+2 {{fake_kernel declared here}}
+  // expected-error@Inputs/sycl.hpp:220 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
 #endif
-  sycl_121_single_task<class fake_kernel>([]() { function(); });
+
+  q.submit([&](handler &h) {
+#ifndef __SYCL_UNNAMED_LAMBDA__
+  // expected-note@+3 {{fake_kernel declared here}}
+  // expected-note@+2 {{in instantiation of function template specialization}}
+#endif
+    h.single_task<class fake_kernel>([]() { function(); });
+  });
+
 #if defined(WARN)
-  // expected-warning@+6 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
-  // expected-note@+5 {{fake_kernel2 declared here}}
+  // expected-warning@Inputs/sycl.hpp:220 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
 #elif defined(ERROR)
-  // expected-error@+3 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
-  // expected-note@+2 {{fake_kernel2 declared here}}
+  // expected-error@Inputs/sycl.hpp:220 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
+#endif
+
+  q.submit([&](handler &h) {
+#ifndef __SYCL_UNNAMED_LAMBDA__
+  // expected-note@+3 {{fake_kernel2 declared here}}
+  // expected-note@+2 {{in instantiation of function template specialization}}
 #endif
-  sycl_121_single_task<class fake_kernel2>([]() {
-    auto l = [](auto f) { f(); };
+    h.single_task<class fake_kernel2>([]() {
+      auto l = [](auto f) { f(); };
+    });
+  });
+
+  q.submit([&](handler &h) {
+    h.single_task<class myWrapper>([]() { function(); });
+  });
+
+  q.submit([&](handler &h) {
+    h.single_task<class myWrapper2>([]() { function(); });
   });
-  sycl_121_single_task<class myWrapper>([]() { function(); });
-  sycl_121_single_task<class myWrapper2>([]() { function(); });
   return 0;
 }

llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp

@@ -272,10 +272,13 @@ void ESIMDLowerVecArgPass::fixGlobals(Module &M) {
     if (NewTy && !G.user_empty()) {
       // Peel off ptr type that getSimdArgPtrTyOrNull applies
       NewTy = NewTy->getPointerElementType();
-      auto ZeroInit = ConstantAggregateZero::get(NewTy);
+      auto InitVal =
+          G.hasInitializer() && isa<UndefValue>(G.getInitializer())
+              ? static_cast<ConstantData *>(UndefValue::get(NewTy))
+              : static_cast<ConstantData *>(ConstantAggregateZero::get(NewTy));
       auto NewGlobalVar =
-          new GlobalVariable(NewTy, G.isConstant(), G.getLinkage(), ZeroInit,
-                             "", G.getThreadLocalMode(), G.getAddressSpace());
+          new GlobalVariable(NewTy, G.isConstant(), G.getLinkage(), InitVal, "",
+                             G.getThreadLocalMode(), G.getAddressSpace());
       NewGlobalVar->setExternallyInitialized(G.isExternallyInitialized());
       NewGlobalVar->copyAttributesFrom(&G);
       NewGlobalVar->takeName(&G);

llvm/test/SYCLLowerIR/esimd_global_undef.ll

@@ -0,0 +1,22 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py
+; RUN: opt < %s -ESIMDLowerVecArg -S | FileCheck %s
+
+; This test checks that undef initializer of a global variable is preserved
+; during ESIMDLowerVecArg transformation
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64-unknown-unknown-sycldevice"
+
+%"class.cl::sycl::INTEL::gpu::simd" = type { <2512 x i32> }
+
+; CHECK: @Global = dso_local global <2512 x i32> undef, align 16384
+@Global = dso_local global %"class.cl::sycl::INTEL::gpu::simd" undef, align 16384
+
+define void @f(<2512 x i32> %simd_val) {
+; CHECK-LABEL: @f(
+; CHECK-NEXT:    store <2512 x i32> [[SIMD_VAL:%.*]], <2512 x i32> addrspace(4)* getelementptr (%"class.cl::sycl::INTEL::gpu::simd", %"class.cl::sycl::INTEL::gpu::simd" addrspace(4)* addrspacecast (%"class.cl::sycl::INTEL::gpu::simd"* bitcast (<2512 x i32>* @Global to %"class.cl::sycl::INTEL::gpu::simd"*) to %"class.cl::sycl::INTEL::gpu::simd" addrspace(4)*), i64 0, i32 0), align 16384
+; CHECK-NEXT:    ret void
+;
+  store <2512 x i32> %simd_val, <2512 x i32> addrspace(4)* getelementptr (%"class.cl::sycl::INTEL::gpu::simd", %"class.cl::sycl::INTEL::gpu::simd" addrspace(4)* addrspacecast (%"class.cl::sycl::INTEL::gpu::simd"* @Global to %"class.cl::sycl::INTEL::gpu::simd" addrspace(4)*), i64 0, i32 0), align 16384
+  ret void
+}

sycl/doc/EnvironmentVariables.md

@@ -23,7 +23,7 @@ subject to change. Do not rely on these variables in production code.
 | SYCL_DISABLE_EXECUTION_GRAPH_CLEANUP | Any(\*) | Disable cleanup of finished command nodes at host-device synchronization points. |
 | SYCL_THROW_ON_BLOCK | Any(\*) | Throw an exception on attempt to wait for a blocked command.  |
 | SYCL_DEVICELIB_INHIBIT_NATIVE | String of device library extensions (separated by a whitespace) | Do not rely on device native support for devicelib extensions listed in this option. |
-| SYCL_DEVICE_ALLOWLIST | A list of devices and their minimum driver version following the pattern: DeviceName:{{XXX}},DriverVersion:{{X.Y.Z.W}}. Also may contain PlatformName and PlatformVersion | Filter out devices that do not match the pattern specified. Regular expression can be passed and the DPC++ runtime will select only those devices which satisfy the regex. |
+| SYCL_DEVICE_ALLOWLIST | A list of devices and their driver version following the pattern: DeviceName:{{XXX}},DriverVersion:{{X.Y.Z.W}}. Also may contain PlatformName and PlatformVersion | Filter out devices that do not match the pattern specified. Regular expression can be passed and the DPC++ runtime will select only those devices which satisfy the regex. Special characters, such as parenthesis, must be escaped. More than one device can be specified using the piping symbol "\|".|
 | SYCL_QUEUE_THREAD_POOL_SIZE | Positive integer | Number of threads in thread pool of queue. |
 | SYCL_DEVICELIB_NO_FALLBACK | Any(\*) | Disable loading and linking of device library images |
 | SYCL_PI_LEVEL0_MAX_COMMAND_LIST_CACHE | Positive integer | Maximum number of oneAPI Level Zero Command lists that can be allocated with no reuse before throwing an "out of resources" error. Default is 20000, threshold may be increased based on resource availabilty and workload demand. |