Merge from 'sycl' to 'sycl-web'

Author: iclsrc

.github/CODEOWNERS

@@ -35,7 +35,7 @@ sycl/doc/extensions/ @intel/dpcpp-specification-reviewers
 sycl/plugins/level_zero/ @intel/dpcpp-l0-pi-reviewers
 
 # ESIMD CPU emulator plug-in
-sycl/plugins/esimd_cpu/ @intel/dpcpp-esimd-reviewers
+sycl/plugins/esimd_emulator/ @intel/dpcpp-esimd-reviewers
 
 # CUDA plugin
 sycl/plugins/cuda/ @intel/llvm-reviewers-cuda

clang/lib/CodeGen/BackendUtil.cpp

@@ -62,6 +62,7 @@
 #include "llvm/Transforms/Coroutines/CoroSplit.h"
 #include "llvm/Transforms/IPO.h"
 #include "llvm/Transforms/IPO/AlwaysInliner.h"
+#include "llvm/Transforms/IPO/DeadArgumentElimination.h"
 #include "llvm/Transforms/IPO/LowerTypeTests.h"
 #include "llvm/Transforms/IPO/ThinLTOBitcodeWriter.h"
 #include "llvm/Transforms/InstCombine/InstCombine.h"
@@ -926,6 +927,8 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
   }
   if (LangOpts.SYCLIsDevice) {
     MPM.addPass(SYCLMutatePrintfAddrspacePass());
+    if (!CodeGenOpts.DisableLLVMPasses && LangOpts.EnableDAEInSpirKernels)
+      MPM.addPass(DeadArgumentEliminationSYCLPass());
   }
 
   // Add SPIRITTAnnotations pass to the pass manager if

clang/lib/Sema/SemaSYCL.cpp

@@ -3719,6 +3719,12 @@ static void CheckSYCL2020SubGroupSizes(Sema &S, FunctionDecl *SYCLKernel,
       CalcEffectiveSubGroup(S.Context, S.getLangOpts(), FD))
     return;
 
+  // No need to validate __spirv routines here since they
+  // are mapped to the equivalent SPIRV operations.
+  const IdentifierInfo *II = FD->getIdentifier();
+  if (II && II->getName().startswith("__spirv_"))
+    return;
+
   // Else we need to figure out why they don't match.
   SourceLocation FDAttrLoc = GetSubGroupLoc(FD);
   SourceLocation KernelAttrLoc = GetSubGroupLoc(SYCLKernel);

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -4,6 +4,11 @@
 
 extern "C" int printf(const char* fmt, ...);
 
+#ifdef __SYCL_DEVICE_ONLY__
+__attribute__((convergent)) extern SYCL_EXTERNAL void
+__spirv_ControlBarrier(int, int, int) noexcept;
+#endif
+
 // Dummy runtime classes to model SYCL API.
 inline namespace cl {
 namespace sycl {
@@ -399,10 +404,19 @@ kernel_parallel_for(const KernelType &KernelFunc) {
   KernelFunc(id<Dims>());
 }
 
+// Dummy parallel_for_work_item function to mimic calls from
+// parallel_for_work_group.
+void parallel_for_work_item() {
+#ifdef __SYCL_DEVICE_ONLY__
+  __spirv_ControlBarrier(0, 0, 0);
+#endif
+}
+
 template <typename KernelName, typename KernelType, int Dims>
 ATTR_SYCL_KERNEL void
 kernel_parallel_for_work_group(const KernelType &KernelFunc) {
   KernelFunc(group<Dims>());
+  parallel_for_work_item();
 }
 
 class handler {

clang/test/CodeGenSYCL/dead-elim-sycl.cpp

@@ -0,0 +1,13 @@
+// RUN: %clang -cc1 -triple spir64-unknown-unknown  -fsycl-is-device  -emit-llvm-bc  -O1 -fno-legacy-pass-manager -fdebug-pass-manager -fenable-sycl-dae -o /dev/null -x c++ < %s 2>&1 | FileCheck %s
+// RUN: %clang -cc1 -triple spir64-unknown-unknown  -fsycl-is-device  -emit-llvm-bc  -O0 -fno-legacy-pass-manager -fdebug-pass-manager  -o /dev/null -x c++ < %s 2>&1 | FileCheck %s --check-prefix DISABLE
+// RUN: %clang -cc1 -triple spir64-unknown-unknown  -fsycl-is-device  -emit-llvm-bc  -O1 -fno-legacy-pass-manager -fdebug-pass-manager -fenable-sycl-dae -disable-llvm-passes -o /dev/null -x c++ < %s 2>&1 | FileCheck %s --check-prefix DISABLE
+// RUN: %clang -cc1 -triple spir64-unknown-unknown  -fsycl-is-device  -emit-llvm-bc  -O1 -flegacy-pass-manager -mllvm -debug-pass=Structure -fenable-sycl-dae -o /dev/null -x c++ < %s 2>&1 | FileCheck %s --check-prefix OLDPM
+// RUN: %clang -cc1 -triple spir64-unknown-unknown  -fsycl-is-device  -emit-llvm-bc  -O1 -flegacy-pass-manager -mllvm -debug-pass=Structure -o /dev/null -x c++ < %s 2>&1 | FileCheck %s --check-prefix DISABLE
+// RUN: %clang -cc1 -triple spir64-unknown-unknown  -fsycl-is-device  -emit-llvm-bc  -O1 -flegacy-pass-manager -mllvm -debug-pass=Structure -fenable-sycl-dae -disable-llvm-passes -o /dev/null -x c++ < %s 2>&1 | FileCheck %s --check-prefix DISABLE
+
+// Verify that Dead Arguments Elimination for SYCL kernels is/is not added to the PM.
+
+// CHECK: Running pass: DeadArgumentEliminationSYCLPass on [module]
+// OLDPM: Dead Argument Elimination for SYCL kernels
+// DISABLE-NOT: DeadArgumentEliminationSYCLPass
+// DISABLE-NOT: Dead Argument Elimination for SYCL kernels

clang/test/CodeGenSYCL/reqd-sub-group-size-spirv-intrin.cpp

@@ -0,0 +1,28 @@
+// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -internal-isystem %S/Inputs -fdeclare-spirv-builtins %s -emit-llvm -o - | FileCheck %s
+
+// Test that when __spirv intrinsics are invoked from kernel functions
+// that have a sub_group_size specified, that such invocations don't
+// trigger the error diagnostic that the intrinsic routines must also
+// marked with the same attribute.
+
+#include "Inputs/sycl.hpp"
+
+int main() {
+  sycl::queue q;
+
+  q.submit([&](sycl::handler &cgh) {
+    auto kernel_ = [=](sycl::group<1> item) [[intel::sub_group_size(8)]] {
+    };
+
+    cgh.parallel_for_work_group<class kernel_class>(
+        cl::sycl::range<1>(), cl::sycl::range<1>(), kernel_);
+  });
+  return 0;
+}
+
+// CHECK: define dso_local spir_kernel void @{{.*}}main{{.*}}kernel_class() {{.*}} !intel_reqd_sub_group_size ![[SUBGROUPSIZE:[0-9]+]]
+// CHECK: tail call spir_func void @{{.*}}__spirv_ControlBarrier{{.*}}({{.*}})
+
+// CHECK: declare spir_func void @{{.*}}__spirv_ControlBarrier{{.*}}({{.*}})
+
+// CHECK: ![[SUBGROUPSIZE]] = !{i32 8}

devops/containers/ubuntu2004_build.Dockerfile

@@ -4,6 +4,10 @@ ENV DEBIAN_FRONTEND=noninteractive
 
 USER root
 
+# Install Nvidia keys
+# https://forums.developer.nvidia.com/t/notice-cuda-linux-repository-key-rotation/212772
+RUN apt-key adv --fetch-keys https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2004/x86_64/3bf863cc.pub
+
 # Install SYCL prerequisites
 COPY scripts/install_build_tools.sh /install.sh
 RUN /install.sh

llvm/tools/sycl-post-link/sycl-post-link.cpp

@@ -519,12 +519,17 @@ static bool removeSYCLKernelsConstRefArray(GlobalVariable *GV) {
          "Cannot remove initializer of llvm.used global");
   Initializer->destroyConstant();
   for (auto It = IOperands.begin(); It != IOperands.end(); It++) {
-    assert(llvm::isSafeToDestroyConstant(*It) &&
-           "Cannot remove an element of initializer of llvm.used global");
     auto Op = (*It)->getOperand(0);
-    (*It)->destroyConstant();
-    // Remove unused kernel declarations to avoid LLVM IR check fails.
     auto *F = dyn_cast<Function>(Op);
+    if (llvm::isSafeToDestroyConstant(*It)) {
+      (*It)->destroyConstant();
+    } else if (F) {
+      // The element in "llvm.used" array has other users. That is Ok for
+      // specialization constants, but is wrong for kernels.
+      llvm::report_fatal_error("Unexpected usage of SYCL kernel");
+    }
+
+    // Remove unused kernel declarations to avoid LLVM IR check fails.
     if (F && F->isDeclaration())
       F->eraseFromParent();
   }

sycl/cmake/modules/PreprocessBoostMp11Headers.cmake

@@ -86,12 +86,11 @@ function(preprocess_mp11_headers)
 
   # 2) Add SYCL_README.txt to the output directory root
   set(SYCL_README_TEXT "\
-  This directory contains boost/mp11 headers imported from\n\
-    ${MP11_HDRS_SRC_PATH} (${MP11_HDRS_SRC_ID})\n\
-  and adapted for use in SYCL headers in a way that does not conflict with\n\
-  potential use of boost in user code. Particularly, `BOOST_*` macros are\n\
-  replaced with `SYCL_DETAIL_BOOST_*`, APIs are moved into the top-level
-  `sycl::detail` namespace. For example, `sycl::detail::boost::mp11::mp_list`.\n")
+  This directory contains boost/mp11 headers adapted for use in SYCL headers in\n\
+  a way that does not conflict with potential use of boost in user code.\n\
+  Particularly, `BOOST_*` macros are replaced with `SYCL_DETAIL_BOOST_*`, APIs\n\
+  are moved into the top-level `sycl::detail` namespace. For example,\n\
+  `sycl::detail::boost::mp11::mp_list`.\n")
 
   set(SYCL_README_FILE_NAME "${MP11_HDRS_OUT}/README.txt")
 

sycl/doc/UsersManual.md

@@ -70,7 +70,7 @@ and not recommended to use in production environment.
     Enables (or disables) LLVM IR dead argument elimination pass to remove
     unused arguments for the kernel functions before translation to SPIR-V.
     Currently has effect only on spir64\* targets.
-    Disabled by default.
+    Enabled by default.
 
 **`-f[no-]sycl-id-queries-fit-in-int`**