Merge from 'sycl' to 'sycl-web' (29 commits)

Author: iclsrc

.github/workflows/sycl-linux-build.yml

@@ -296,6 +296,9 @@ jobs:
         cmake --build $GITHUB_WORKSPACE/build --target install-llvm-cov
         cmake --build $GITHUB_WORKSPACE/build --target install-llvm-profdata
         cmake --build $GITHUB_WORKSPACE/build --target install-compiler-rt
+        # This is required to perform the DeviceConfigFile consistency test, see
+        # sycl/test-e2e/Basic/device_config_file_consistency.cpp.
+        cmake --install $GITHUB_WORKSPACE/build --component DeviceConfigFile
     - name: Additional Install for "--shared-libs" build
       if: ${{ always() && !cancelled() && steps.build.conclusion == 'success' && contains(inputs.build_configure_extra_args, '--shared-libs') }}
       run: |

buildbot/configure.py

@@ -66,7 +66,6 @@ def do_configure(args, passthrough_args):
     xpti_enable_werror = "OFF"
     llvm_enable_zstd = "OFF"
     spirv_enable_dis = "OFF"
-    sycl_install_device_config_file = "OFF"
 
     if sys.platform != "darwin":
         # For more info on the enablement of level_zero_v2 refer to this document:
@@ -162,7 +161,6 @@ def do_configure(args, passthrough_args):
                 libclc_targets_to_build += libclc_nvidia_target_names
             libclc_gen_remangled_variants = "ON"
             spirv_enable_dis = "ON"
-            sycl_install_device_config_file = "ON"
 
     if args.enable_backends:
         sycl_enabled_backends += args.enable_backends
@@ -211,7 +209,6 @@ def do_configure(args, passthrough_args):
         "-DSYCL_ENABLE_EXTENSION_JIT={}".format(sycl_enable_jit),
         "-DSYCL_ENABLE_MAJOR_RELEASE_PREVIEW_LIB={}".format(sycl_preview_lib),
         "-DBUG_REPORT_URL=https://github.com/intel/llvm/issues",
-        "-DSYCL_INSTALL_DEVICE_CONFIG_FILE={}".format(sycl_install_device_config_file),
     ]
 
     if libclc_enabled:

clang/lib/AST/ASTContext.cpp

@@ -12949,7 +12949,8 @@ static GVALinkage adjustGVALinkageForAttributes(const ASTContext &Context,
     if (Context.shouldExternalize(D))
       return GVA_StrongExternal;
   } else if (Context.getLangOpts().SYCLIsDevice &&
-             D->hasAttr<DeviceKernelAttr>()) {
+             (D->hasAttr<DeviceKernelAttr>() &&
+              D->getAttr<DeviceKernelAttr>()->isImplicit())) {
     if (L == GVA_DiscardableODR)
       return GVA_StrongODR;
   }

clang/lib/Sema/SemaSYCL.cpp

@@ -4099,10 +4099,13 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
   }
 
   void addArrayElementInit(FieldDecl *FD, QualType T) {
-    Expr *RCE = createReinterpretCastExpr(
-        createGetAddressOf(ArrayParamBases.pop_back_val()),
-        SemaSYCLRef.getASTContext().getPointerType(T));
-    Expr *Initializer = createDerefOp(RCE);
+    Expr *Initializer = ArrayParamBases.pop_back_val();
+    if (!T->isPointerType()) {
+      Expr *RCE = createReinterpretCastExpr(
+          createGetAddressOf(Initializer),
+          SemaSYCLRef.getASTContext().getPointerType(T));
+      Initializer = createDerefOp(RCE);
+    }
     addFieldInit(FD, T, Initializer);
   }
 
@@ -5447,9 +5450,13 @@ void SemaSYCL::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
   ESIMDKernelDiagnostics esimdKernel(*this, KernelObj->getLocation(),
                                      IsSIMDKernel);
 
-  SyclKernelDeclCreator kernel_decl(*this, KernelObj->getLocation(),
-                                    KernelCallerFunc->isInlined(), IsSIMDKernel,
-                                    KernelCallerFunc);
+  // In case of syntax errors in input programs we are not able to access
+  // CallOperator. In this case the value of IsInlined doesn't matter, because
+  // compilation will fail with errors anyways.
+  const bool IsInlined =
+      CallOperator ? CallOperator->isInlined() : /* placeholder */ false;
+  SyclKernelDeclCreator kernel_decl(*this, KernelObj->getLocation(), IsInlined,
+                                    IsSIMDKernel, KernelCallerFunc);
   SyclKernelBodyCreator kernel_body(*this, kernel_decl, KernelObj,
                                     KernelCallerFunc, IsSIMDKernel,
                                     CallOperator);

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

@@ -0,0 +1,66 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+
+#include "sycl.hpp"
+
+sycl::queue myQueue;
+
+using namespace sycl;
+
+// CHECK: %struct.__wrapper_class = type { [2 x i32] }
+// CHECK: %class.anon = type { [2 x i32] }
+// CHECK: %struct.__wrapper_class.0 = type { [2 x ptr addrspace(1)] }
+// CHECK: %class.anon.1 = type { [2 x ptr addrspace(4)] }
+
+int main() {
+  int Array[2];
+  myQueue.submit([&](sycl::handler &h) {
+    h.single_task<class IntArray>(
+        [=] {
+          int local = Array[1];
+        });
+  });
+
+// CHECK-LABEL: @{{.*}}IntArray(ptr {{.*}}byval(%struct.__wrapper_class)
+// CHECK:   %__SYCLKernel = alloca %class.anon, align 4
+// CHECK:   %__SYCLKernel.ascast = addrspacecast ptr %__SYCLKernel to ptr addrspace(4)
+// CHECK:   %_arg_Array.ascast = addrspacecast ptr %_arg_Array to ptr addrspace(4)
+// CHECK:   %Array = getelementptr inbounds nuw %class.anon, ptr addrspace(4) %__SYCLKernel.ascast, i32 0, i32 0
+// CHECK:   %0 = getelementptr inbounds nuw %struct.__wrapper_class, ptr addrspace(4) %_arg_Array.ascast, i32 0, i32 0
+// CHECK:   %arrayinit.begin = getelementptr inbounds [2 x i32], ptr addrspace(4) %Array, i64 0, i64 0
+// CHECK:   br label %arrayinit.body
+// CHECK: arrayinit.body:                                   ; preds = %arrayinit.body, %entry
+// CHECK:   %arrayinit.index = phi i64 [ 0, %entry ], [ %arrayinit.next, %arrayinit.body ]
+// CHECK:   %1 = getelementptr inbounds i32, ptr addrspace(4) %arrayinit.begin, i64 %arrayinit.index
+// CHECK:   %arrayidx = getelementptr inbounds nuw [2 x i32], ptr addrspace(4) %0, i64 0, i64 %arrayinit.index
+// CHECK:   %2 = load i32, ptr addrspace(4) %arrayidx, align 4
+// CHECK:   store i32 %2, ptr addrspace(4) %1, align 4
+// CHECK:   %arrayinit.next = add nuw i64 %arrayinit.index, 1
+// CHECK:   %arrayinit.done = icmp eq i64 %arrayinit.next, 2
+// CHECK:   br i1 %arrayinit.done, label %arrayinit.end, label %arrayinit.body
+// CHECK: arrayinit.end:                                    ; preds = %arrayinit.body
+
+  int *ArrayOfPointers[2];
+  myQueue.submit([&](sycl::handler &h) {
+    h.single_task<class PtrArray>(
+        [=] {
+          int local = *ArrayOfPointers[1];
+        });
+  });
+// CHECK-LABEL: @{{.*}}PtrArray(ptr {{.*}}byval(%struct.__wrapper_class.0)
+// CHECK:   %__SYCLKernel = alloca %class.anon.1, align 8
+// CHECK:   %__SYCLKernel.ascast = addrspacecast ptr %__SYCLKernel to ptr addrspace(4)
+// CHECK:   %_arg_ArrayOfPointers.ascast = addrspacecast ptr %_arg_ArrayOfPointers to ptr addrspace(4)
+// CHECK:   %ArrayOfPointers = getelementptr inbounds nuw %class.anon.1, ptr addrspace(4) %__SYCLKernel.ascast, i32 0, i32 0
+// CHECK:   %0 = getelementptr inbounds nuw %struct.__wrapper_class.0, ptr addrspace(4) %_arg_ArrayOfPointers.ascast, i32 0, i32 0
+// CHECK:   %arrayidx = getelementptr inbounds nuw [2 x ptr addrspace(1)], ptr addrspace(4) %0, i64 0, i64 0
+// CHECK:   %1 = load ptr addrspace(1), ptr addrspace(4) %arrayidx, align 8
+// CHECK:   %2 = addrspacecast ptr addrspace(1) %1 to ptr addrspace(4)
+// CHECK:   store ptr addrspace(4) %2, ptr addrspace(4) %ArrayOfPointers, align 8
+// CHECK:   %arrayinit.element = getelementptr inbounds ptr addrspace(4), ptr addrspace(4) %ArrayOfPointers, i64 1
+// CHECK:   %3 = getelementptr inbounds nuw %struct.__wrapper_class.0, ptr addrspace(4) %_arg_ArrayOfPointers.ascast, i32 0, i32 0
+// CHECK:   %arrayidx1 = getelementptr inbounds nuw [2 x ptr addrspace(1)], ptr addrspace(4) %3, i64 0, i64 1
+// CHECK:   %4 = load ptr addrspace(1), ptr addrspace(4) %arrayidx1, align 8
+// CHECK:   %5 = addrspacecast ptr addrspace(1) %4 to ptr addrspace(4)
+// CHECK:   store ptr addrspace(4) %5, ptr addrspace(4) %arrayinit.element, align 8
+}

clang/test/CodeGenSYCL/device_has.cpp

@@ -6,7 +6,7 @@
 using namespace sycl;
 queue q;
 
-// CHECK-DAG: define dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]]
+// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]]
 
 // CHECK-DAG: define {{.*}}spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] {
 [[sycl::device_has(sycl::aspect::cpu)]] void func1() {}
@@ -67,7 +67,7 @@ void foo() {
   q.submit([&](handler &h) {
     KernelFunctor f1;
     h.single_task<class kernel_name_1>(f1);
-    // CHECK-DAG: define dso_local spir_kernel void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]] !srcloc ![[SRCLOC8:[0-9]+]]
+    // CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]] !srcloc ![[SRCLOC8:[0-9]+]]
     h.single_task<class kernel_name_2>([]() [[sycl::device_has(sycl::aspect::gpu)]] {});
   });
 }

clang/test/CodeGenSYCL/dynamic_local_accessor.cpp

@@ -7,7 +7,7 @@
 // The first two RUN commands verify that the init call is generated with the correct arguments in LLVM IR
 // and the second two RUN commands verify the contents of the integration header produced by the frontend.
 //
-// CHECK-IR: define dso_local spir_kernel void @
+// CHECK-IR: define {{.*}}spir_kernel void @
 // CHECK-IR-SAME: ptr addrspace(3) noundef align 4 [[PTR:%[a-zA-Z0-9_]+]]
 //
 // CHECK-IR: [[PTR]].addr = alloca ptr addrspace(3), align 8

clang/test/CodeGenSYCL/dynamic_work_group_memory.cpp

@@ -7,7 +7,7 @@
 // The first two RUN commands verify that the init call is generated with the correct arguments in LLVM IR
 // and the second two RUN commands verify the contents of the integration header produced by the frontend.
 //
-// CHECK-IR: define dso_local spir_kernel void @
+// CHECK-IR: define {{.*}}spir_kernel void @
 // CHECK-IR-SAME: ptr addrspace(3) noundef align 4 [[PTR:%[a-zA-Z0-9_]+]]
 //
 // CHECK-IR: [[PTR]].addr = alloca ptr addrspace(3), align 8

clang/test/CodeGenSYCL/generated-types-initialization.cpp

@@ -38,7 +38,7 @@ int main() {
   });
   return 0;
 }
-// CHECK: define dso_local spir_kernel void @{{.*}}basic(ptr noundef byval(%struct.__generated_B) align 8 %_arg_Obj)
+// CHECK: define {{.*}}spir_kernel void @{{.*}}basic(ptr noundef byval(%struct.__generated_B) align 8 %_arg_Obj)
 //
 // Kernel object clone.
 // CHECK: %[[K:[a-zA-Z0-9_.]+]] = alloca %class.anon
@@ -54,7 +54,7 @@ int main() {
 // Kernel body call.
 // CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) %[[K_as_cast]])
 
-// CHECK: define dso_local spir_kernel void @{{.*}}nns(ptr noundef byval(%struct.__generated_B.0) align 8 %_arg_NNSObj)
+// CHECK: define {{.*}}spir_kernel void @{{.*}}nns(ptr noundef byval(%struct.__generated_B.0) align 8 %_arg_NNSObj)
 //
 // Kernel object clone.
 // CHECK: %[[NNSK:[a-zA-Z0-9_.]+]] = alloca %class.anon.2

clang/test/CodeGenSYCL/kernel-handler.cpp

@@ -22,7 +22,7 @@ void test(int val) {
   });
 }
 
-// ALL: define dso_local{{ spir_kernel | ptx_kernel | }}void @{{.*}}test_kernel_handler{{[^(]*}}
+// ALL: define {{.*}}{{ spir_kernel | ptx_kernel | }}void @{{.*}}test_kernel_handler{{[^(]*}}
 // NONATIVESUPPORT-SAME: (ptr noundef byval(%class.anon) align 4 %_arg__sycl_functor, ptr addrspace(1) noundef align 1 %_arg__specialization_constants_buffer)
 // NATIVESUPPORT-SAME: (i32 noundef %_arg_a, ptr addrspace(1) noundef align 1 %_arg__specialization_constants_buffer)
 // ALL: %kh = alloca %"class.sycl::_V1::kernel_handler", align 1