Merge from 'sycl' to 'sycl-web'

Author: iclsrc

buildbot/testlist.cfg

@@ -0,0 +1,45 @@
+:test_accessor
+:test_address_space
+:test_atomic
+:test_buffer
+:test_context
+:test_device
+:test_device_event
+:test_device_selector
+:test_error
+:test_event
+:test_exception_handling
+:test_group
+:test_h_item
+:test_handler
+:test_header
+:test_hierarchical
+:test_id
+:test_image
+:test_invoke
+:test_item
+:test_kernel
+:test_kernel_args
+:test_math_builtin_api
+:test_multi_ptr
+:test_nd_item
+:test_nd_range
+:test_opencl_interop
+:test_platform
+:test_pointers
+:test_program
+:test_queue
+:test_range
+:test_sampler
+:test_scalars
+:test_std_classes
+:test_stream
+:test_vector_alias
+:test_vector_api
+:test_vector_constructors
+:test_vector_load_store
+# Disable test to speedup testing until JIT is optimized
+#:test_vector_operators
+:test_vector_swizzle_assignment
+:test_vector_swizzles
+:test_vector_swizzles_opencl

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10490,8 +10490,9 @@ def err_sycl_virtual_types : Error<
   "No class with a vtable can be used in a SYCL kernel or any code included in the kernel">;
 def note_sycl_used_here : Note<"used here">;
 def note_sycl_recursive_function_declared_here: Note<"function implemented using recursion declared here">;
-def err_sycl_non_trivially_copyable_type : Error<
-  "kernel parameter has non-trivially copyable class/struct type %0">;
+def err_sycl_non_trivially_copy_ctor_dtor_type
+    : Error<"kernel parameter has non-trivially %select{copy "
+            "constructible|destructible}0 class/struct type %1">;
 def err_sycl_non_std_layout_type : Error<
   "kernel parameter has non-standard layout class/struct type %0">;
 def err_conflicting_sycl_kernel_attributes : Error<

clang/lib/CodeGen/SYCLLowerIR/LowerWGScope.cpp

@@ -40,10 +40,10 @@
 // local "shadow" variable. Before each PFWI invocation leader WI stores its
 // private copy of the variable into the shadow (under "is leader" guard), then
 // all WIs (ouside of "is leader" guard) load the shadow value into their
-// private copies ("materialize" the private copy). This works becase these
+// private copies ("materialize" the private copy). This works because these
 // variables are uniform - i.e. have the same value in all WIs and are not
 // changed within PFWI. The only exceptions are captures of private_memory
-// isntances - see next.
+// instances - see next.
 // ** Kind 1:
 // Even though WG-scope locals are supposed to be uniform, there is one
 // exception - capture of local of kind 1. It is always captured by non-const
@@ -52,7 +52,7 @@
 // of kind 1 variable's alloca is stored within the PFWI lambda.
 // Materialization of the lambda object value writes result of alloca of the
 // leader WI's private variable into the private copy of the lambda object,
-// which is wrong. So for tese variables this pass adds a write of the private
+// which is wrong. So for these variables this pass adds a write of the private
 // variable's address into the private copy of the lambda object right after its
 // materialization:
 //     if (is_leader())
@@ -120,9 +120,6 @@ class SYCLLowerWGScopeLegacyPass : public FunctionPass {
 
   // run the LowerWGScope pass on the specified module
   bool runOnFunction(Function &F) override {
-    if (skipFunction(F))
-      return false;
-
     FunctionAnalysisManager FAM;
     auto PA = Impl.run(F, FAM);
     return !PA.areAllPreserved();
@@ -479,7 +476,7 @@ static void materializeLocalsInWIScopeBlocksImpl(
 // Checks if there is a need to materialize value of given local in given work
 // item-scope basic block.
 static bool localMustBeMaterialized(const AllocaInst *L, const BasicBlock &BB) {
-  // TODO this is overly convervative - see speculations below.
+  // TODO this is overly conservative - see speculations below.
   return true;
 }
 
@@ -745,9 +742,9 @@ PreservedAnalyses SYCLLowerWGScopePass::run(Function &F,
     for (; I->getOpcode() == Instruction::Alloca; I = I->getNextNode()) {
       auto *AllocaI = dyn_cast<AllocaInst>(I);
       // Allocas marked with "work_item_scope" are those originating from
-      // cl::sycl::private_memory<T> variables, which must in private. No
-      // shadows/materialization is needed for them because they can be updated
-      // only within PFWIs
+      // cl::sycl::private_memory<T> variables, which must be in private memory.
+      // No shadows/materialization is needed for them because they can be
+      // updated only within PFWIs
       if (!AllocaI->getMetadata(WI_SCOPE_MD))
         Allocas.insert(AllocaI);
     }
@@ -801,7 +798,7 @@ PreservedAnalyses SYCLLowerWGScopePass::run(Function &F,
   }
   // There can be allocas not corresponding to any variable declared in user
   // code but generated by the compiler - e.g. for non-trivially typed
-  // parameters passed by by value. There can be WG scope stores into such
+  // parameters passed by value. There can be WG scope stores into such
   // allocas, which need to be made visible to all WIs. This is done via
   // creating a "shadow" workgroup-shared variable and using it to propagate
   // the value of the alloca'ed variable to worker WIs from the leader.
@@ -815,7 +812,7 @@ PreservedAnalyses SYCLLowerWGScopePass::run(Function &F,
   // Now materialize the locals:
   materializeLocalsInWIScopeBlocks(Allocas, WIScopeBBs);
 
-  // Fixup captured addresses of private_memory isntances in current WI
+  // Fixup captured addresses of private_memory instances in current WI
   for (auto *PFWICall : PFWICalls)
     fixupPrivateMemoryPFWILambdaCaptures(PFWICall);
 

clang/lib/Sema/SemaDecl.cpp

@@ -7057,6 +7057,13 @@ NamedDecl *Sema::ActOnVariableDeclarator(
       NewVD->setTSCSpec(TSCS);
   }
 
+  // Static variables declared inside SYCL device code must be const or
+  // constexpr
+  if (getLangOpts().SYCLIsDevice && SCSpec == DeclSpec::SCS_static &&
+      !R.isConstant(Context))
+    SYCLDiagIfDeviceCode(D.getIdentifierLoc(), diag::err_sycl_restrict)
+        << Sema::KernelNonConstStaticDataVariable;
+
   switch (D.getDeclSpec().getConstexprSpecifier()) {
   case CSK_unspecified:
     break;

clang/lib/Sema/SemaSYCL.cpp

@@ -1069,10 +1069,22 @@ static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
           continue;
         }
       }
-      if (!ArgTy.isTriviallyCopyableType(Context)) {
+
+      CXXRecordDecl *RD =
+          cast<CXXRecordDecl>(ArgTy->getAs<RecordType>()->getDecl());
+      if (!RD->hasTrivialCopyConstructor()) {
+        Context.getDiagnostics().Report(
+            Fld->getLocation(),
+            diag::err_sycl_non_trivially_copy_ctor_dtor_type)
+            << 0 << ArgTy;
+        AllArgsAreValid = false;
+        continue;
+      }
+      if (!RD->hasTrivialDestructor()) {
         Context.getDiagnostics().Report(
-            Fld->getLocation(), diag::err_sycl_non_trivially_copyable_type)
-            << ArgTy;
+            Fld->getLocation(),
+            diag::err_sycl_non_trivially_copy_ctor_dtor_type)
+            << 1 << ArgTy;
         AllArgsAreValid = false;
         continue;
       }

clang/test/CodeGenSYCL/address-space-of-returns.cpp

@@ -10,13 +10,14 @@ const char *ret_char() {
 // CHECK: ret i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([2 x i8], [2 x i8]* @.str, i64 0, i64 0) to i8 addrspace(4)*)
 
 const char *ret_arr() {
-  static char Arr[42];
+  const static char Arr[36] = "Carrots, cabbage, radish, potatoes!";
   return Arr;
 }
-// CHECK: ret i8 addrspace(4)* getelementptr inbounds ([42 x i8], [42 x i8] addrspace(4)* addrspacecast ([42 x i8] addrspace(1)* @{{.*}}ret_arr{{.*}}Arr to [42 x i8] addrspace(4)*), i64 0, i64 0)
+
+// CHECK: ret i8 addrspace(4)* getelementptr inbounds ([36 x i8], [36 x i8] addrspace(4)* addrspacecast ([36 x i8] addrspace(1)* @{{.*}}ret_arr{{.*}}Arr to [36 x i8] addrspace(4)*), i64 0, i64 0)
 
 const char &ret_ref() {
-  static char a = 'A';
+  const static char a = 'A';
   return a;
 }
 // CHECK: ret i8 addrspace(4)* addrspacecast (i8 addrspace(1)* @{{.*}}ret_ref{{.*}} to i8 addrspace(4)*)

clang/test/CodeGenSYCL/address-space-swap.cpp

@@ -1,35 +1,33 @@
 // RUN: %clang -fsycl-device-only -S -emit-llvm %s -o - | FileCheck %s
 #include <algorithm>
 
-
 void test() {
-  static int foo = 0x42;
-// CHECK: @[[FOO:[a-zA-Z0-9_]+]] = internal addrspace(1) global i32 66, align 4
+  int foo = 0x42;
+  // CHECK: %[[FOO:[a-zA-Z0-9_]+]] = alloca i32, align 4
   int i = 43;
-// CHECK: %[[I:[a-zA-Z0-9_]+]] = alloca i32, align 4
+  // CHECK: %[[I:[a-zA-Z0-9_]+]] = alloca i32, align 4
 
-  int* p1 = &foo;
-  int* p2 = &i;
-// CHECK: %[[P1:[a-zA-Z0-9_]+]] = alloca i32 addrspace(4)*, align 8
-// CHECK: %[[P2:[a-zA-Z0-9_]+]] = alloca i32 addrspace(4)*, align 8
-// CHECK: %[[P1GEN:[a-zA-Z0-9_]+]] = addrspacecast i32 addrspace(4)** %[[P1]] to i32 addrspace(4)* addrspace(4)*
-// CHECK: %[[P2GEN:[a-zA-Z0-9_]+]] = addrspacecast i32 addrspace(4)** %[[P2]] to i32 addrspace(4)* addrspace(4)*
+  int *p1 = &foo;
+  int *p2 = &i;
+  // CHECK: %[[P1:[a-zA-Z0-9_]+]] = alloca i32 addrspace(4)*, align 8
+  // CHECK: %[[P2:[a-zA-Z0-9_]+]] = alloca i32 addrspace(4)*, align 8
+  // CHECK: %[[P1GEN:[a-zA-Z0-9_]+]] = addrspacecast i32 addrspace(4)** %[[P1]] to i32 addrspace(4)* addrspace(4)*
+  // CHECK: %[[P2GEN:[a-zA-Z0-9_]+]] = addrspacecast i32 addrspace(4)** %[[P2]] to i32 addrspace(4)* addrspace(4)*
 
   std::swap(p1, p2);
-// CHECK: call spir_func void @_ZSt4swap{{.*}}(i32 addrspace(4)* addrspace(4)* dereferenceable(8) %[[P1GEN]], i32 addrspace(4)* addrspace(4)* dereferenceable(8) %[[P2GEN]])
+  // CHECK: call spir_func void @_ZSt4swap{{.*}}(i32 addrspace(4)* addrspace(4)* dereferenceable(8) %[[P1GEN]], i32 addrspace(4)* addrspace(4)* dereferenceable(8) %[[P2GEN]])
 
   std::swap(foo, i);
-// CHECK: %[[ICAST:[a-zA-Z0-9_]+]] = addrspacecast i32* %[[I]] to i32 addrspace(4)*
-// CHECK: call spir_func void @_ZSt4swap{{.*}}(i32 addrspace(4)* dereferenceable(4) addrspacecast (i32 addrspace(1)* @[[FOO]] to i32 addrspace(4)*), i32 addrspace(4)* dereferenceable(4) %[[ICAST]])
+  // CHECK: %[[FOOCAST:[a-zA-Z0-9_]+]] = addrspacecast i32* %[[FOO]] to i32 addrspace(4)*
+  // CHECK: %[[ICAST:[a-zA-Z0-9_]+]] = addrspacecast i32* %[[I]] to i32 addrspace(4)*
+  // CHECK: call spir_func void @_ZSt4swap{{.*}}(i32 addrspace(4)* dereferenceable(4) %[[FOOCAST]], i32 addrspace(4)* dereferenceable(4) %[[ICAST]])
 }
 
-
 template <typename name, typename Func>
 __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
   kernelFunc();
 }
 
-
 int main() {
   kernel_single_task<class fake_kernel>([]() { test(); });
   return 0;

clang/test/CodeGenSYCL/intel-fpga-local.cpp

@@ -1,7 +1,9 @@
 // RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s -check-prefixes CHECK-DEVICE,CHECK-BOTH
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s -check-prefixes CHECK-HOST,CHECK-BOTH
 
+// CHECK-BOTH: @_ZZ15attrs_on_staticvE15static_numbanks = internal{{.*}}constant i32 20, align 4
 // CHECK-DEVICE:  [[ANN_numbanks_4:@.str]] = {{.*}}{memory:DEFAULT}{sizeinfo:4}{numbanks:4}
+// CHECK-BOTH: @_ZZ15attrs_on_staticvE15static_annotate = internal{{.*}}constant i32 30, align 4
 // CHECK-BOTH:    [[ANN_annotate:@.str[.0-9]*]] = {{.*}}foobar
 // CHECK-DEVICE:  [[ANN_register:@.str.[0-9]*]] = {{.*}}{register:1}
 // CHECK-DEVICE:  [[ANN_memory_default:@.str.[0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4}
@@ -25,24 +27,17 @@
 
 // CHECK-BOTH: @llvm.global.annotations
 // CHECK-DEVICE-SAME: { i8* addrspacecast (i8 addrspace(1)* bitcast (i32 addrspace(1)* @_ZZ15attrs_on_staticvE15static_numbanks to i8 addrspace(1)*) to i8*)
-// CHECK-DEVICE-SAME: [[ANN_numbanks_4]]{{.*}}i32 38
+// CHECK-DEVICE-SAME: [[ANN_numbanks_4]]{{.*}} i32 39
 // CHECK-DEVICE-SAME: { i8* addrspacecast (i8 addrspace(1)* bitcast (i32 addrspace(1)* @_ZZ15attrs_on_staticvE15static_annotate to i8 addrspace(1)*) to i8*)
 // CHECK-HOST-SAME: { i8* bitcast (i32* @_ZZ15attrs_on_staticvE15static_annotate to i8*)
-// CHECK-BOTH-SAME: [[ANN_annotate]]{{.*}}i32 42
+// CHECK-BOTH-SAME: [[ANN_annotate]]{{.*}} i32 40
 
 // CHECK-HOST-NOT: llvm.var.annotation
 // CHECK-HOST-NOT: llvm.ptr.annotation
 
 void attrs_on_static() {
-  int a = 42;
-  static int static_numbanks [[intelfpga::numbanks(4)]];
-  // CHECK-BOTH: load{{.*}}static_numbanks
-  // CHECK-BOTH: store{{.*}}static_numbanks
-  static_numbanks = static_numbanks + a;
-  static int static_annotate [[clang::annotate("foobar")]];
-  // CHECK-BOTH: load{{.*}}static_annotate
-  // CHECK-BOTH: store{{.*}}static_annotate
-  static_annotate = static_annotate + a;
+  const static int static_numbanks [[intelfpga::numbanks(4)]] = 20;
+  const static int static_annotate [[clang::annotate("foobar")]] = 30;
 }
 
 void attrs_on_var() {

clang/test/SemaSYCL/non-trivially-copyable-kernel-param.cpp

@@ -11,6 +11,16 @@ struct B {
   B (const B& x) : i(x.i) {}
 };
 
+struct C : A {
+  const A C2;
+  C() : A{0}, C2{2}{}
+};
+
+struct D {
+  int i;
+  ~D();
+};
+
 template <typename Name, typename Func>
 __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
   kernelFunc();
@@ -20,9 +30,14 @@ void test() {
   A IamGood;
   IamGood.i = 0;
   B IamBad(1);
+  C IamAlsoGood;
+  D IamAlsoBad{0};
   kernel_single_task<class kernel_capture_refs>([=] {
     int a = IamGood.i;
-    // expected-error@+1 {{kernel parameter has non-trivially copyable class/struct type}}
+    // expected-error@+1 {{kernel parameter has non-trivially copy constructible class/struct type}}
     int b = IamBad.i;
+    int c = IamAlsoGood.i;
+    // expected-error@+1 {{kernel parameter has non-trivially destructible class/struct type}}
+    int d = IamAlsoBad.i;
   });
 }

clang/test/SemaSYCL/sycl-device-const-static.cpp

@@ -0,0 +1,29 @@
+// RUN: %clang_cc1 -verify -fsyntax-only -fsycl-is-device %s
+
+void usage() {
+  // expected-error@+1{{SYCL kernel cannot use a non-const static data variable}}
+  static int s1;
+  const static int cs = 0;
+  constexpr static int ces = 0;
+}
+
+template <typename Name, typename Func>
+__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+  // expected-error@+1{{SYCL kernel cannot use a non-const static data variable}}
+  static int z;
+  // expected-note-re@+2{{called by 'kernel_single_task<fake_kernel, (lambda at {{.*}})>}}
+  // expected-note-re@+1{{called by 'kernel_single_task<fake_kernel, (lambda at {{.*}})>}}
+  kernelFunc();
+}
+
+int main() {
+  static int s2;
+  kernel_single_task<class fake_kernel>([]() {
+    //  expected-note@+1{{called by 'operator()'}}
+    usage();
+    // expected-error@+1{{SYCL kernel cannot use a non-const static data variable}}
+    static int s3;
+  });
+
+  return 0;
+}