Merge from 'sycl' to 'sycl-web'

Author: iclsrc

buildbot/dependency.conf

@@ -3,10 +3,10 @@
 ocl_cpu_rt_ver=2020.11.11.0.04
 # https://github.com/intel/llvm/releases/download/2020-WW45/win-oclcpuexp-2020.11.11.0.04_rel.zip
 ocl_cpu_rt_ver_win=2020.11.11.0.04
-# Same GPU driver supports Level Zero and OpenCL:
-# https://github.com/intel/compute-runtime/releases/tag/20.45.18403
-ocl_gpu_rt_ver=20.45.18403
-# Same GPU driver supports Level Zero and OpenCL:
+# Same GPU driver supports Level Zero and OpenCL
+# https://github.com/intel/compute-runtime/releases/tag/20.46.18421
+ocl_gpu_rt_ver=20.46.18421
+# Same GPU driver supports Level Zero and OpenCL
 # https://downloadmirror.intel.com/29988/a08/igfx_win10_100.8935.zip
 ocl_gpu_rt_ver_win=27.20.100.8935
 intel_sycl_ver=build
@@ -25,7 +25,7 @@ fpga_ver_win=20201022_000005
 [DRIVER VERSIONS]
 cpu_driver_lin=2020.11.11.0.04
 cpu_driver_win=2020.11.11.0.04
-gpu_driver_lin=20.45.18403
+gpu_driver_lin=20.46.18421
 gpu_driver_win=27.20.100.8935
 fpga_driver_lin=2020.11.11.0.04
 fpga_driver_win=2020.11.11.0.04

clang/include/clang/AST/Type.h

@@ -384,8 +384,8 @@ class Qualifiers {
          | (((uint32_t) space) << AddressSpaceShift);
   }
   void removeAddressSpace() { setAddressSpace(LangAS::Default); }
-  void addAddressSpace(LangAS space) {
-    assert(space != LangAS::Default);
+  void addAddressSpace(LangAS space, bool AllowDefaultAddrSpace = false) {
+    assert(space != LangAS::Default || AllowDefaultAddrSpace);
     setAddressSpace(space);
   }
 

clang/lib/Driver/Driver.cpp

@@ -2791,7 +2791,7 @@ bool Driver::checkForOffloadStaticLib(Compilation &C,
   if (Args.hasArg(options::OPT_offload_lib_Group))
     return true;
   SmallVector<const char *, 16> OffloadLibArgs(getLinkerArgs(C, Args));
-  for (const StringRef &OLArg : OffloadLibArgs)
+  for (StringRef OLArg : OffloadLibArgs)
     if (isStaticArchiveFile(OLArg) && hasOffloadSections(C, OLArg, Args)) {
       // FPGA binaries with AOCX or AOCR sections are not considered fat
       // static archives.
@@ -4443,7 +4443,7 @@ class OffloadingActionBuilder final {
       // incorporated into the aoc compilation
       if (SYCLfpgaTriple) {
         SmallVector<const char *, 16> LinkArgs(getLinkerArgs(C, Args));
-        for (const StringRef &LA : LinkArgs) {
+        for (StringRef LA : LinkArgs) {
           if (isStaticArchiveFile(LA) && hasOffloadSections(C, LA, Args)) {
             const llvm::opt::OptTable &Opts = C.getDriver().getOpts();
             Arg *InputArg = MakeInputArg(Args, Opts, Args.MakeArgString(LA));
@@ -5124,7 +5124,7 @@ void Driver::BuildActions(Compilation &C, DerivedArgList &Args,
       UnbundlerInputs.push_back(Current);
     };
     bool IsWholeArchive = false;
-    for (const StringRef &LA : LinkArgs) {
+    for (StringRef LA : LinkArgs) {
       if (isStaticArchiveFile(LA)) {
         addUnbundlerInput(
             IsWholeArchive ? types::TY_WholeArchive : types::TY_Archive, LA);

clang/lib/Sema/SemaInit.cpp

@@ -4905,7 +4905,8 @@ static void TryReferenceInitializationCore(Sema &S,
     // Add addr space conversion if required.
     if (T1Quals.getAddressSpace() != T2Quals.getAddressSpace()) {
       auto T4Quals = cv1T4.getQualifiers();
-      T4Quals.addAddressSpace(T1Quals.getAddressSpace());
+      T4Quals.addAddressSpace(T1Quals.getAddressSpace(),
+                              S.getLangOpts().SYCLIsDevice);
       QualType cv1T4WithAS = S.Context.getQualifiedType(T2, T4Quals);
       Sequence.AddQualificationConversionStep(cv1T4WithAS, ValueKind);
       cv1T4 = cv1T4WithAS;

clang/lib/Sema/SemaSYCL.cpp

@@ -816,7 +816,7 @@ class KernelObjVisitor {
   template <typename... Tn>
   bool handleField(FieldDecl *FD, QualType FDTy, Tn &&... tn) {
     bool result = true;
-    std::initializer_list<int>{(result = result && tn(FD, FDTy), 0)...};
+    (void)std::initializer_list<int>{(result = result && tn(FD, FDTy), 0)...};
     return result;
   }
   template <typename... Tn>

clang/test/Driver/clang-offload-bundler-tgtsym.c

@@ -0,0 +1,33 @@
+// REQUIRES: x86-registered-target
+
+// This test check that clang-offload-bundler adds .tgtsym section to the output
+// file when creating a fat object. This section contains names of the external
+// symbols defined in the embdedded target objects with target prefixes.
+
+// RUN: %clang -target %itanium_abi_triple -c %s -o %t.o
+// RUN: %clang -target x86_64-pc-linux-gnu -c %s -o %t.tgt1
+// RUN: %clang -target spir64 -emit-llvm   -c %s -o %t.tgt2
+
+// RUN: clang-offload-bundler -type=o -targets=host-%itanium_abi_triple,openmp-x86_64-pc-linux-gnu,sycl-spir64 -inputs=%t.o,%t.tgt1,%t.tgt2 -outputs=%t.fat.o
+// RUN: llvm-readobj --string-dump=.tgtsym %t.fat.o | FileCheck %s
+
+// CHECK: String dump of section '.tgtsym':
+// CHECK-DAG: openmp-x86_64-pc-linux-gnu.foo
+// CHECK-DAG: openmp-x86_64-pc-linux-gnu.bar
+// CHECK-DAG: sycl-spir64.foo
+// CHECK-DAG: sycl-spir64.bar
+// CHECK-NOT: undefined_func
+// CHECK-NOT: static_func
+
+extern void undefined_func(void);
+
+void foo(void) {
+  undefined_func();
+}
+
+static void static_func(void) __attribute__((noinline));
+static void static_func(void) {}
+
+void bar(void) {
+  static_func();
+}

clang/test/SemaSYCL/sycl-templ-addr-space.cpp

@@ -0,0 +1,52 @@
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64 -verify %s
+
+// Test that the compiler no longer asserts while processing this test case.
+
+template <int N>
+struct integral_constant {
+  static constexpr int val = N;
+};
+template <typename T>
+T &&convert(int);
+template <typename T>
+auto declval(convert<T>(0));
+template <typename Op1, typename Op2>
+class sub_group {
+  template <typename T, decltype(declval<T>)>
+  // expected-note@+1 {{possible target for call}}
+  static integral_constant<true> binary_op();
+  // expected-error@+1 {{reference to overloaded function could not be resolved; did you mean to call it?}}
+  decltype(binary_op<Op1, Op2>) i;
+};
+// expected-note-re@+2 {{in instantiation {{.*}} requested here}}
+template <typename n>
+struct group : sub_group<n, int> {
+};
+template <typename T>
+struct wrapper {
+  typedef T val;
+};
+class element_type {
+};
+struct Container {
+  using __element_type = __attribute__((opencl_global)) element_type;
+};
+template <bool, class Base, class> using BaseImpl = Base;
+// expected-note-re@+1 2{{candidate constructor {{.*}} not viable: {{.*}}}}
+template <int> class id_type {
+  template <class...> struct Base;
+  // expected-note-re@+1 {{in instantiation {{.*}} requested here}}
+  template <class Der> struct Base<Der> : BaseImpl<Der::val, Base<>, Der> {};
+  // expected-note-re@+1 {{in instantiation {{.*}} requested here}}
+  template <typename T> using Base2 = wrapper<typename Base<group<T>>::val>;
+  // expected-note-re@+3 {{in instantiation {{.*}} requested here}}
+  // expected-note-re@+2 {{in instantiation {{.*}} required here}}
+  // expected-note-re@+1 {{candidate template ignored: {{.*}}}}
+  template <typename T, typename = Base2<T>> id_type(T &);
+};
+id_type<1> get() {
+  Container::__element_type *ElemPtr;
+  // expected-error@+2 {{no viable conversion from returned value of type 'Container::__element_type' (aka '__global element_type') to function return type 'id_type<1>'}}
+  // expected-note-re@+1 {{while substituting {{.*}}}}
+  return ElemPtr[0];
+}

clang/tools/clang-offload-bundler/CMakeLists.txt

@@ -1,4 +1,4 @@
-set(LLVM_LINK_COMPONENTS Object Support)
+set(LLVM_LINK_COMPONENTS Core Object Support)
 
 add_clang_tool(clang-offload-bundler
   ClangOffloadBundler.cpp

clang/tools/clang-offload-bundler/ClangOffloadBundler.cpp

@@ -23,6 +23,7 @@
 #include "llvm/ADT/StringSet.h"
 #include "llvm/ADT/StringSwitch.h"
 #include "llvm/ADT/Triple.h"
+#include "llvm/IR/LLVMContext.h"
 #include "llvm/Object/Archive.h"
 #include "llvm/Object/ArchiveWriter.h"
 #include "llvm/Object/Binary.h"
@@ -121,6 +122,9 @@ static cl::opt<unsigned>
 /// Prefix of an added section name with bundle size.
 #define SIZE_SECTION_PREFIX "__CLANG_OFFLOAD_BUNDLE_SIZE__"
 
+/// Section name which holds target symbol names.
+#define SYMBOLS_SECTION_NAME ".tgtsym"
+
 /// The index of the host input in the list of inputs.
 static unsigned HostInputIndex = ~0u;
 
@@ -546,6 +550,62 @@ class ObjectFileHandler final : public FileHandler {
   /// Input sizes.
   SmallVector<uint64_t, 16u> InputSizes;
 
+  // Return a buffer with symbol names that are defined in target objects.
+  // Each symbol name is prefixed by a target name <kind>-<triple> to uniquely
+  // identify the target it belongs to, and symbol names are separated from each
+  // other by '\0' characters.
+  Expected<SmallVector<char, 0>> makeTargetSymbolTable() {
+    SmallVector<char, 0> SymbolsBuf;
+    raw_svector_ostream SymbolsOS(SymbolsBuf);
+
+    for (unsigned I = 0; I < NumberOfInputs; ++I) {
+      if (I == HostInputIndex)
+        continue;
+
+      // Get the list of symbols defined in the target object. Open file and
+      // check if it is a symbolic file.
+      ErrorOr<std::unique_ptr<MemoryBuffer>> BufOrErr =
+          MemoryBuffer::getFileOrSTDIN(InputFileNames[I]);
+      if (!BufOrErr)
+        return createFileError(InputFileNames[I], BufOrErr.getError());
+
+      LLVMContext Context;
+      Expected<std::unique_ptr<Binary>> BinOrErr =
+          createBinary(BufOrErr.get()->getMemBufferRef(), &Context);
+
+      // If it is not a symbolic file just ignore it since we cannot do anything
+      // with it.
+      if (!BinOrErr) {
+        if (auto Err = isNotObjectErrorInvalidFileType(BinOrErr.takeError()))
+          return std::move(Err);
+        continue;
+      }
+      auto *SF = dyn_cast<SymbolicFile>(&**BinOrErr);
+      if (!SF)
+        continue;
+
+      for (BasicSymbolRef Symbol : SF->symbols()) {
+        Expected<uint32_t> FlagsOrErr = Symbol.getFlags();
+        if (!FlagsOrErr)
+          return FlagsOrErr.takeError();
+
+        // We are interested in externally visible and defined symbols only, so
+        // ignore it if this is not such a symbol.
+        bool Undefined = *FlagsOrErr & SymbolRef::SF_Undefined;
+        bool Global = *FlagsOrErr & SymbolRef::SF_Global;
+        if (Undefined || !Global)
+          continue;
+
+        // Add symbol name with the target prefix to the buffer.
+        SymbolsOS << TargetNames[I] << ".";
+        if (Error Err = Symbol.printName(SymbolsOS))
+          return std::move(Err);
+        SymbolsOS << '\0';
+      }
+    }
+    return SymbolsBuf;
+  }
+
 public:
   ObjectFileHandler(std::unique_ptr<ObjectFile> ObjIn)
       : FileHandler(), Obj(std::move(ObjIn)),
@@ -793,6 +853,25 @@ class ObjectFileHandler final : public FileHandler {
                                     SIZE_SECTION_PREFIX + TargetNames[I] + "=" +
                                     *SizeFileOrErr));
     }
+
+    // Add a section with symbol names that are defined in target objects to the
+    // output fat object.
+    Expected<SmallVector<char, 0>> SymbolsOrErr = makeTargetSymbolTable();
+    if (!SymbolsOrErr)
+      return SymbolsOrErr.takeError();
+
+    if (!SymbolsOrErr->empty()) {
+      // Add section with symbols names to fat object.
+      Expected<StringRef> SymbolsFileOrErr =
+          TempFiles.Create(makeArrayRef(*SymbolsOrErr));
+      if (!SymbolsFileOrErr)
+        return SymbolsFileOrErr.takeError();
+
+      ObjcopyArgs.push_back(SS.save(Twine("--add-section=") +
+                                    SYMBOLS_SECTION_NAME + "=" +
+                                    *SymbolsFileOrErr));
+    }
+
     ObjcopyArgs.push_back(InputFileNames[HostInputIndex]);
     ObjcopyArgs.push_back(IntermediateObj);
 

llvm/lib/Support/PropertySetIO.cpp

@@ -182,8 +182,7 @@ PropertyValue &PropertyValue::operator=(PropertyValue &&P) {
 
 PropertyValue &PropertyValue::operator=(const PropertyValue &P) {
   if (P.getType() == BYTE_ARRAY)
-    *this =
-        std::move(PropertyValue(P.asByteArray(), P.getByteArraySizeInBits()));
+    *this = PropertyValue(P.asByteArray(), P.getByteArraySizeInBits());
   else
     copy(P);
   return *this;

llvm/lib/Support/SimpleTable.cpp

@@ -76,7 +76,7 @@ Error SimpleTable::addColumn(const Twine &Title, ArrayRef<std::string> Cells) {
   if (!Rows.empty() && (Rows.size() != N))
     return makeError("column size mismatch for " + Title);
   if (Error Err = addColumnName(Title.str()))
-    return std::move(Err);
+    return Err;
   if (Rows.empty()) {
     Rows.resize(Cells.size());
     for (auto &R : Rows)
@@ -114,7 +114,7 @@ Error SimpleTable::renameColumn(StringRef OldName, StringRef NewName) {
 
   if (I < 0)
     return makeError("column not found: " + OldName);
-  *ColumnNum2Name[I] = std::move(NewName.str());
+  *ColumnNum2Name[I] = NewName.str();
   ColumnName2Num.erase(OldName);
   ColumnName2Num[StringRef(*ColumnNum2Name[I])] = I;
   return Error::success();

sycl/include/CL/sycl/ONEAPI/experimental/spec_constant.hpp

@@ -17,6 +17,7 @@
 
 #pragma once
 
+#include <CL/sycl/detail/stl_type_traits.hpp>
 #include <CL/sycl/detail/sycl_fe_intrins.hpp>
 #include <CL/sycl/exception.hpp>
 
@@ -41,11 +42,15 @@ template <typename T, typename ID = T> class spec_constant {
   spec_constant(T Cst) : Val(Cst) {}
 
   T Val;
-#endif
+#else
+  char padding[sizeof(T)];
+#endif // __SYCL_DEVICE_ONLY__
   friend class cl::sycl::program;
 
 public:
-  T get() const { // explicit access.
+  template <typename V = T>
+  typename sycl::detail::enable_if_t<std::is_arithmetic<V>::value, V>
+  get() const { // explicit access.
 #ifdef __SYCL_DEVICE_ONLY__
     const char *TName = __builtin_unique_stable_name(ID);
     return __sycl_getSpecConstantValue<T>(TName);
@@ -54,6 +59,19 @@ template <typename T, typename ID = T> class spec_constant {
 #endif // __SYCL_DEVICE_ONLY__
   }
 
+  template <typename V = T>
+  typename sycl::detail::enable_if_t<std::is_class<V>::value &&
+                                         std::is_pod<V>::value,
+                                     V>
+  get() const { // explicit access.
+#ifdef __SYCL_DEVICE_ONLY__
+    const char *TName = __builtin_unique_stable_name(ID);
+    return __sycl_getCompositeSpecConstantValue<T>(TName);
+#else
+    return Val;
+#endif // __SYCL_DEVICE_ONLY__
+  }
+
   operator T() const { // implicit conversion.
     return get();
   }