Merge remote-tracking branch 'intel_llvm/sycl' into llvmspirv_pulldown

Author: vmaksimo

.github/CODEOWNERS

@@ -50,6 +50,6 @@ clang/tools/clang-offload-*/ @intel/dpcpp-tools-reviewers
 # Explicit SIMD
 ESIMD/ @intel/dpcpp-esimd-reviewers
 esimd/ @intel/dpcpp-esimd-reviewers
-sycl/include/sycl/ext/intel/experimental/esimd.hpp @intel/dpcpp-esimd-reviewers
+sycl/include/sycl/ext/intel/esimd.hpp @intel/dpcpp-esimd-reviewers
 sycl/doc/extensions/experimental/sycl_ext_intel_esimd/ @intel/dpcpp-esimd-reviewers
 llvm/lib/SYCLLowerIR/CMakeLists.txt @intel/dpcpp-tools-reviewers @intel/dpcpp-esimd-reviewers

.github/workflows/sycl_containers.yaml

@@ -10,11 +10,15 @@ on:
     paths:
       - 'devops/containers/**'
       - 'devops/dependencies.json'
+      - 'devops/scripts/install_drivers.sh'
+      - 'devops/scripts/install_build_tools.sh'
       - '.github/workflows/sycl_containers.yaml'
   pull_request:
     paths:
       - 'devops/containers/**'
       - 'devops/dependencies.json'
+      - 'devops/scripts/install_drivers.sh'
+      - 'devops/scripts/install_build_tools.sh'
       - '.github/workflows/sycl_containers.yaml'
 
 jobs:

.github/workflows/sycl_linux_build_and_test.yml

@@ -15,10 +15,6 @@ on:
         type: string
         required: false
         default: "ghcr.io/intel/llvm/ubuntu2004_build:latest"
-      build_runs_on:
-        type: string
-        required: false
-        default: "build"
       build_ref:
         type: string
         required: false
@@ -68,7 +64,7 @@ on:
 jobs:
   build:
     name: Build + LIT
-    runs-on: ${{ inputs.build_runs_on }}
+    runs-on: [Linux, build]
     container:
       image: ${{ inputs.build_image }}
       options: -u 1001:1001

.github/workflows/sycl_nightly.yml

@@ -64,7 +64,6 @@ jobs:
     if: github.repository == 'intel/llvm'
     uses: ./.github/workflows/sycl_linux_build_and_test.yml
     with:
-      build_runs_on: build
       build_cache_root: "/__w/"
       build_cache_suffix: new_pm
       build_artifact_suffix: new_pm

.github/workflows/sycl_post_commit.yml

@@ -15,14 +15,12 @@ jobs:
     name: Linux Default
     uses: ./.github/workflows/sycl_linux_build_and_test.yml
     with:
-      build_runs_on: build
       build_cache_root: "/__w/llvm"
       build_artifact_suffix: default
   linux_no_assert:
     name: Linux (no assert)
     uses: ./.github/workflows/sycl_linux_build_and_test.yml
     with:
-      build_runs_on: build
       build_cache_root: "/__w/llvm"
       build_cache_suffix: gcc_no_assertions
       build_artifact_suffix: gcc_no_assertions

.github/workflows/sycl_precommit.yml

@@ -4,6 +4,16 @@ on:
   pull_request:
     branches:
     - sycl
+    # Do not run builds if changes are only in the following locations
+    paths-ignore:
+    - 'devops/containers/**'
+    - 'devops/scripts/install_drivers.sh'
+    - 'devops/scripts/install_build_tools.sh'
+    - 'sycl/doc/**'
+    - 'sycl/gdb/**'
+    - 'clang/docs/**'
+    - '**.md'
+    - '**.rst'
 
 jobs:
   lint:
@@ -25,7 +35,6 @@ jobs:
     if: always() && (success() || contains(github.event.pull_request.labels.*.name, 'ignore-lint'))
     uses: ./.github/workflows/sycl_linux_build_and_test.yml
     with:
-      build_runs_on: "build"
       build_cache_root: "/__w/"
       build_cache_size: "8G"
       build_artifact_suffix: "default"

clang/include/clang/Basic/BuiltinsNVPTX.def

@@ -300,6 +300,12 @@ BUILTIN(__nvvm_sin_approx_f, "ff", "")
 BUILTIN(__nvvm_cos_approx_ftz_f, "ff", "")
 BUILTIN(__nvvm_cos_approx_f, "ff", "")
 
+// Tanh
+
+TARGET_BUILTIN(__nvvm_tanh_approx_f, "ff", "", AND(SM_75,PTX70))
+TARGET_BUILTIN(__nvvm_tanh_approx_f16, "hh", "", AND(SM_75, PTX70))
+TARGET_BUILTIN(__nvvm_tanh_approx_f16x2, "V2hV2h", "", AND(SM_75, PTX70))
+
 // Fma
 
 TARGET_BUILTIN(__nvvm_fma_rn_f16, "hhhh", "", AND(SM_53, PTX42))

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -7755,6 +7755,8 @@ let CategoryName = "Lambda Issue" in {
     "%select{| explicitly}1 captured here">;
   def err_implicit_this_capture : Error<
     "implicit capture of 'this' is not allowed for kernel functions">;
+  def err_lambda_member_access : Error<
+    "invalid attempt to access member of lambda">;
 
   // C++14 lambda init-captures.
   def warn_cxx11_compat_init_capture : Warning<

clang/lib/Driver/Driver.cpp

@@ -2992,9 +2992,10 @@ static SmallVector<const char *, 16>
 getLinkerArgs(Compilation &C, DerivedArgList &Args, bool IncludeObj = false) {
   SmallVector<const char *, 16> LibArgs;
   SmallVector<std::string, 8> LibPaths;
-  // Add search directories from LIBRARY_PATH env variable
+  bool IsMSVC = C.getDefaultToolChain().getTriple().isWindowsMSVCEnvironment();
+  // Add search directories from LIBRARY_PATH/LIB env variable
   llvm::Optional<std::string> LibPath =
-      llvm::sys::Process::GetEnv("LIBRARY_PATH");
+      llvm::sys::Process::GetEnv(IsMSVC ? "LIB" : "LIBRARY_PATH");
   if (LibPath) {
     SmallVector<StringRef, 8> SplitPaths;
     const char EnvPathSeparatorStr[] = {llvm::sys::EnvPathSeparator, '\0'};
@@ -3028,12 +3029,36 @@ getLinkerArgs(Compilation &C, DerivedArgList &Args, bool IncludeObj = false) {
   };
   for (const auto *A : Args) {
     std::string FileName = A->getAsString(Args);
+    auto addLibArg = [&](StringRef LibName) -> bool {
+      if (isStaticArchiveFile(LibName) ||
+          (IncludeObj && isObjectFile(LibName.str()))) {
+        LibArgs.push_back(Args.MakeArgString(LibName));
+        return true;
+      }
+      return false;
+    };
     if (A->getOption().getKind() == Option::InputClass) {
-      StringRef Value(A->getValue());
-      if (isStaticArchiveFile(Value) ||
-          (IncludeObj && isObjectFile(Value.str()))) {
-        LibArgs.push_back(Args.MakeArgString(FileName));
+      if (addLibArg(FileName))
         continue;
+    }
+    // Evaluate any libraries passed along after /link. These are typically
+    // ignored by the driver and sent directly to the linker. When performing
+    // offload, we should evaluate them at the driver level.
+    if (A->getOption().matches(options::OPT__SLASH_link)) {
+      for (const std::string &Value : A->getValues()) {
+        // Add any libpath values.
+        StringRef OptCheck(Value);
+        if (OptCheck.startswith_insensitive("-libpath:") ||
+            OptCheck.startswith_insensitive("/libpath:"))
+          LibPaths.emplace_back(Value.substr(std::string("-libpath:").size()));
+        if (addLibArg(Value))
+          continue;
+        for (auto LPath : LibPaths) {
+          SmallString<128> FullLib(LPath);
+          llvm::sys::path::append(FullLib, Value);
+          if (addLibArg(FullLib))
+            continue;
+        }
       }
     }
     if (A->getOption().matches(options::OPT_Wl_COMMA) ||
@@ -3075,9 +3100,8 @@ getLinkerArgs(Compilation &C, DerivedArgList &Args, bool IncludeObj = false) {
           llvm::StringSaver S(A);
           llvm::cl::ExpandResponseFiles(
               S,
-              C.getDefaultToolChain().getTriple().isWindowsMSVCEnvironment()
-                  ? llvm::cl::TokenizeWindowsCommandLine
-                  : llvm::cl::TokenizeGNUCommandLine,
+              IsMSVC ? llvm::cl::TokenizeWindowsCommandLine
+                     : llvm::cl::TokenizeGNUCommandLine,
               ExpandArgs);
           for (StringRef EA : ExpandArgs)
             addKnownValues(EA);
@@ -8350,11 +8374,10 @@ bool clang::driver::isStaticArchiveFile(const StringRef &FileName) {
   if (!llvm::sys::path::has_extension(FileName))
     // Any file with no extension should not be considered an Archive.
     return false;
-  StringRef Ext(llvm::sys::path::extension(FileName).drop_front());
   llvm::file_magic Magic;
   llvm::identify_magic(FileName, Magic);
   // Only .lib and archive files are to be considered.
-  return (Ext == "lib" || Magic == llvm::file_magic::archive);
+  return (Magic == llvm::file_magic::archive);
 }
 
 bool clang::driver::willEmitRemarks(const ArgList &Args) {

clang/lib/Sema/SemaAccess.cpp

@@ -1903,7 +1903,13 @@ void Sema::CheckLookupAccess(const LookupResult &R) {
       AccessTarget Entity(Context, AccessedEntity::Member,
                           R.getNamingClass(), I.getPair(),
                           R.getBaseObjectType());
-      Entity.setDiag(diag::err_access);
+      // This is to avoid leaking implementation details of lambda object.
+      // We do not want to generate 'private member access' diagnostic for
+      // lambda object.
+      if ((R.getNamingClass())->isLambda())
+        Diag(R.getNameLoc(), diag::err_lambda_member_access);
+      else
+        Entity.setDiag(diag::err_access);
       CheckAccess(*this, R.getNameLoc(), Entity);
     }
   }

clang/lib/Sema/SemaLambda.cpp

@@ -1702,12 +1702,23 @@ FieldDecl *Sema::BuildCaptureField(RecordDecl *RD,
                                    const sema::Capture &Capture) {
   SourceLocation Loc = Capture.getLocation();
   QualType FieldType = Capture.getCaptureType();
+  IdentifierInfo *Id = nullptr;
 
   TypeSourceInfo *TSI = nullptr;
   if (Capture.isVariableCapture()) {
     auto *Var = Capture.getVariable();
     if (Var->isInitCapture())
       TSI = Capture.getVariable()->getTypeSourceInfo();
+
+    // TODO: Upstream this behavior to LLVM project to save
+    // user speciifed names for all lambdas.
+    // For SYCL compilations, save user specified names for
+    // lambda capture.
+    if (getLangOpts().SYCLIsDevice || getLangOpts().SYCLIsHost) {
+      StringRef CaptureName = Var->getName();
+      if (!CaptureName.empty())
+        Id = &Context.Idents.get(CaptureName.str());
+    }
   }
 
   // FIXME: Should we really be doing this? A null TypeSourceInfo seems more
@@ -1717,8 +1728,8 @@ FieldDecl *Sema::BuildCaptureField(RecordDecl *RD,
 
   // Build the non-static data member.
   FieldDecl *Field =
-      FieldDecl::Create(Context, RD, /*StartLoc=*/Loc, /*IdLoc=*/Loc,
-                        /*Id=*/nullptr, FieldType, TSI, /*BW=*/nullptr,
+      FieldDecl::Create(Context, RD, /*StartLoc=*/Loc, /*IdLoc=*/Loc, Id,
+                        FieldType, TSI, /*BW=*/nullptr,
                         /*Mutable=*/false, ICIS_NoInit);
   // If the variable being captured has an invalid type, mark the class as
   // invalid as well.

clang/test/CXX/expr/expr.prim/expr.prim.lambda/p11-1y.cpp

@@ -1,6 +1,12 @@
-// RUN: %clang_cc1 -std=c++1y %s -verify
+// TODO: The SYCL changes in this test can be removed if/when changes (i.e.
+// PR to save user-specified names in lambda class) is upstreamed to LLVM
+// project.
 
-const char *has_no_member = [x("hello")] {}.x; // expected-error {{no member named 'x'}}
+// RUN: %clang_cc1 -std=c++1y %s -verify=notsycl,expected
+// RUN: %clang_cc1 -fsycl-is-device -std=c++1y %s -verify=sycl,expected
+
+const char *has_no_member = [x("hello")] {}.x; // notsycl-error {{no member named 'x'}}
+// sycl-error@-1 {{invalid attempt to access member of lambda}}
 
 double f;
 auto with_float = [f(1.0f)] {

clang/test/CodeGenSYCL/accessor-readonly.cpp

@@ -14,8 +14,8 @@ void f0(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::bu
 
 // CHECK: spir_kernel{{.*}}f1_kernel
 // CHECK-NOT: readonly
-// CHECK-SAME: %_arg_{{.*}}%_arg_1{{.*}}%_arg_2{{.*}}%_arg_3
-// CHECK-SAME:  readonly align 4 %_arg_4
+// CHECK-SAME: %_arg_write_acc{{.*}}%_arg_write_acc1{{.*}}%_arg_write_acc2{{.*}}%_arg_write_acc3
+// CHECK-SAME:  readonly align 4 %_arg_read_acc
 void f1(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::buffer<int, 1> &out_buf) {
   myQueue.submit([&](cl::sycl::handler &cgh) {
     auto write_acc = out_buf.get_access<cl::sycl::access::mode::write>(cgh);
@@ -25,9 +25,9 @@ void f1(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::bu
 }
 
 // CHECK: spir_kernel{{.*}}f2_kernel
-// CHECK-SAME: readonly align 4 %_arg_
+// CHECK-SAME: readonly align 4 %_arg_read_acc
 // CHECK-NOT: readonly
-// CHECK-SAME: %_arg_8
+// CHECK-SAME: %_arg_write_acc
 void f2(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::buffer<int, 1> &out_buf) {
   myQueue.submit([&](cl::sycl::handler &cgh) {
     auto read_acc = in_buf.get_access<cl::sycl::access::mode::read>(cgh);

clang/test/CodeGenSYCL/accessor_no_alias_property.cpp

@@ -20,7 +20,7 @@ int main() {
       accessorB;
 
   // Check that noalias parameter attribute is emitted when no_alias accessor property is used
-  // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE16kernel_function1({{.*}} noalias {{.*}} %_arg_, {{.*}})
+  // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE16kernel_function1({{.*}} noalias {{.*}} %_arg_accessorA, {{.*}})
   cl::sycl::kernel_single_task<class kernel_function1>(
       [=]() {
         accessorA.use();

clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp

@@ -21,8 +21,8 @@ int main() {
 
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_function
 // CHECK-SAME: i32 addrspace(1)* noundef align 4 [[MEM_ARG:%[a-zA-Z0-9_]+]],
-// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[ACC_RANGE:%[a-zA-Z0-9_]+_1]],
-// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[MEM_RANGE:%[a-zA-Z0-9_]+_2]],
+// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[ACC_RANGE:%[a-zA-Z0-9_]+1]],
+// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[MEM_RANGE:%[a-zA-Z0-9_]+2]],
 // CHECK-SAME: %"struct.cl::sycl::id"* noundef byval{{.*}}align 4 [[OFFSET:%[a-zA-Z0-9_]+]])
 // Check alloca for pointer argument
 // CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(1)*

clang/test/CodeGenSYCL/device-variables.cpp

@@ -32,10 +32,10 @@ int main() {
     // CHECK: store i32 1, i32 addrspace(4)* %b
     foo(local_value);
     // Local variables and constexprs captured by lambda
-    // CHECK:  [[GEP:%[0-9]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{.*}}, i32 0, i32 0
+    // CHECK:  [[GEP:%[a-z_]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{.*}}, i32 0, i32 0
     // CHECK: call spir_func void @{{.*}}foo{{.*}}(i32 addrspace(4)* noundef align 4 dereferenceable(4) [[GEP]])
     int some_device_local_var = some_local_var;
-    // CHECK:  [[GEP1:%[0-9]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{.*}}, i32 0, i32 1
+    // CHECK:  [[GEP1:%[a-z_]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{.*}}, i32 0, i32 1
     // CHECK:  [[LOAD1:%[0-9]+]] = load i32, i32 addrspace(4)* [[GEP1]]
     // CHECK:  store i32 [[LOAD1]], i32 addrspace(4)* %some_device_local_var
   });

clang/test/CodeGenSYCL/esimd-accessor-ptr-md.cpp

@@ -30,7 +30,7 @@ void test(int val) {
   // --- Attributes
   // CHECK: {{.*}} !kernel_arg_accessor_ptr ![[ACC_PTR_ATTR:[0-9]+]] !sycl_explicit_simd !{{[0-9]+}} {{.*}}{
   // --- init_esimd call is expected instead of __init:
-  // CHECK:   call spir_func void @{{.*}}__init_esimd{{.*}}(%"{{.*}}sycl::accessor" addrspace(4)* {{[^,]*}} %{{[0-9]+}}, i32 addrspace(1)* noundef %{{[0-9]+}})
+  // CHECK:   call spir_func void @{{.*}}__init_esimd{{.*}}(%"{{.*}}sycl::accessor" addrspace(4)* {{[^,]*}} %{{[a-zA-Z0-9_]+}}, i32 addrspace(1)* noundef %{{[0-9]+}})
   // CHECK-LABEL: }
   // CHECK: ![[ACC_PTR_ATTR]] = !{i1 true, i1 false, i1 true}
 }

clang/test/CodeGenSYCL/image_accessor.cpp

@@ -8,27 +8,27 @@
 //
 // CHECK-1DRO: %opencl.image1d_ro_t = type opaque
 // CHECK-1DRO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image1d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_ro_t addrspace(1)* %{{[0-9]+}})
+// CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image1d_ro_t addrspace(1)* %{{[0-9]+}})
 //
 // CHECK-2DRO: %opencl.image2d_ro_t = type opaque
 // CHECK-2DRO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image2d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_ro_t addrspace(1)* %{{[0-9]+}})
+// CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image2d_ro_t addrspace(1)* %{{[0-9]+}})
 //
 // CHECK-3DRO: %opencl.image3d_ro_t = type opaque
 // CHECK-3DRO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image3d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_ro_t addrspace(1)* %{{[0-9]+}})
+// CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image3d_ro_t addrspace(1)* %{{[0-9]+}})
 //
 // CHECK-1DWO: %opencl.image1d_wo_t = type opaque
 // CHECK-1DWO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image1d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_wo_t addrspace(1)* %{{[0-9]+}})
+// CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image1d_wo_t addrspace(1)* %{{[0-9]+}})
 //
 // CHECK-2DWO: %opencl.image2d_wo_t = type opaque
 // CHECK-2DWO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image2d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_wo_t addrspace(1)* %{{[0-9]+}})
+// CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image2d_wo_t addrspace(1)* %{{[0-9]+}})
 //
 // CHECK-3DWO: %opencl.image3d_wo_t = type opaque
 // CHECK-3DWO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image3d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_wo_t addrspace(1)* %{{[0-9]+}})
+// CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image3d_wo_t addrspace(1)* %{{[0-9]+}})
 //
 // TODO: Add tests for the image_array opencl datatype support.
 #include "Inputs/sycl.hpp"