SWDEV-179954 - OpenCL/LC - Merge branch amd-master into amd-common

Change-Id: Ia35fd55e8612ec5c14b5730b97a9c1115c552c5b

include/llvm/BinaryFormat/Wasm.h

@@ -242,7 +242,9 @@ enum : unsigned {
 enum : unsigned {
   WASM_OPCODE_END = 0x0b,
   WASM_OPCODE_CALL = 0x10,
+  WASM_OPCODE_LOCAL_GET = 0x20,
   WASM_OPCODE_GLOBAL_GET = 0x23,
+  WASM_OPCODE_GLOBAL_SET = 0x24,
   WASM_OPCODE_I32_STORE = 0x36,
   WASM_OPCODE_I32_CONST = 0x41,
   WASM_OPCODE_I64_CONST = 0x42,

include/llvm/CodeGen/GlobalISel/CallLowering.h

@@ -27,6 +27,7 @@
 
 namespace llvm {
 
+class CCState;
 class DataLayout;
 class Function;
 class MachineIRBuilder;
@@ -163,7 +164,10 @@ class CallLowering {
   /// \return True if everything has succeeded, false otherwise.
   bool handleAssignments(MachineIRBuilder &MIRBuilder, ArrayRef<ArgInfo> Args,
                          ValueHandler &Handler) const;
-
+  bool handleAssignments(CCState &CCState,
+                         SmallVectorImpl<CCValAssign> &ArgLocs,
+                         MachineIRBuilder &MIRBuilder, ArrayRef<ArgInfo> Args,
+                         ValueHandler &Handler) const;
 public:
   CallLowering(const TargetLowering *TLI) : TLI(TLI) {}
   virtual ~CallLowering() = default;

include/llvm/DebugInfo/PDB/Native/HashTable.h

@@ -72,6 +72,12 @@ class HashTableIterator
     assert(Map->Present.test(Index));
     return Map->Buckets[Index];
   }
+
+  // Implement postfix op++ in terms of prefix op++ by using the superclass
+  // implementation.
+  using iterator_facade_base<HashTableIterator<ValueT>,
+                             std::forward_iterator_tag,
+                             const std::pair<uint32_t, ValueT>>::operator++;
   HashTableIterator &operator++() {
     while (Index < Map->Buckets.size()) {
       ++Index;
@@ -94,9 +100,6 @@ class HashTableIterator
 
 template <typename ValueT>
 class HashTable {
-  using const_iterator = HashTableIterator<ValueT>;
-  friend const_iterator;
-
   struct Header {
     support::ulittle32_t Size;
     support::ulittle32_t Capacity;
@@ -105,6 +108,9 @@ class HashTable {
   using BucketList = std::vector<std::pair<uint32_t, ValueT>>;
 
 public:
+  using const_iterator = HashTableIterator<ValueT>;
+  friend const_iterator;
+
   HashTable() { Buckets.resize(8); }
   explicit HashTable(uint32_t Capacity) {
     Buckets.resize(Capacity);

include/llvm/DebugInfo/PDB/Native/InjectedSourceStream.h

@@ -0,0 +1,44 @@
+//===- InjectedSourceStream.h - PDB Headerblock Stream Access ---*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_DEBUGINFO_PDB_RAW_PDBINJECTEDSOURCESTREAM_H
+#define LLVM_DEBUGINFO_PDB_RAW_PDBINJECTEDSOURCESTREAM_H
+
+#include "llvm/DebugInfo/PDB/Native/HashTable.h"
+#include "llvm/DebugInfo/PDB/Native/RawTypes.h"
+#include "llvm/Support/Error.h"
+
+namespace llvm {
+namespace msf {
+class MappedBlockStream;
+}
+namespace pdb {
+class PDBFile;
+class PDBStringTable;
+
+class InjectedSourceStream {
+public:
+  InjectedSourceStream(std::unique_ptr<msf::MappedBlockStream> Stream);
+  Error reload(const PDBStringTable &Strings);
+
+  using const_iterator = HashTable<SrcHeaderBlockEntry>::const_iterator;
+  const_iterator begin() const { return InjectedSourceTable.begin(); }
+  const_iterator end() const { return InjectedSourceTable.end(); }
+
+  uint32_t size() const { return InjectedSourceTable.size(); }
+
+private:
+  std::unique_ptr<msf::MappedBlockStream> Stream;
+
+  const SrcHeaderBlockHeader* Header;
+  HashTable<SrcHeaderBlockEntry> InjectedSourceTable;
+};
+}
+}
+
+#endif

include/llvm/DebugInfo/PDB/Native/NativeEnumInjectedSources.h

@@ -0,0 +1,43 @@
+//==- NativeEnumInjectedSources.cpp - Native Injected Source Enumerator --*-==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_DEBUGINFO_PDB_NATIVE_NATIVEENUMINJECTEDSOURCES_H
+#define LLVM_DEBUGINFO_PDB_NATIVE_NATIVEENUMINJECTEDSOURCES_H
+
+#include "llvm/DebugInfo/PDB/IPDBEnumChildren.h"
+#include "llvm/DebugInfo/PDB/IPDBInjectedSource.h"
+#include "llvm/DebugInfo/PDB/Native/InjectedSourceStream.h"
+
+namespace llvm {
+namespace pdb {
+
+class InjectedSourceStream;
+class PDBStringTable;
+
+class NativeEnumInjectedSources : public IPDBEnumChildren<IPDBInjectedSource> {
+public:
+  NativeEnumInjectedSources(PDBFile &File, const InjectedSourceStream &IJS,
+                            const PDBStringTable &Strings);
+
+  uint32_t getChildCount() const override;
+  std::unique_ptr<IPDBInjectedSource>
+  getChildAtIndex(uint32_t Index) const override;
+  std::unique_ptr<IPDBInjectedSource> getNext() override;
+  void reset() override;
+
+private:
+  PDBFile &File;
+  const InjectedSourceStream &Stream;
+  const PDBStringTable &Strings;
+  InjectedSourceStream::const_iterator Cur;
+};
+
+} // namespace pdb
+} // namespace llvm
+
+#endif

include/llvm/DebugInfo/PDB/Native/PDBFile.h

@@ -32,6 +32,7 @@ namespace pdb {
 class DbiStream;
 class GlobalsStream;
 class InfoStream;
+class InjectedSourceStream;
 class PDBStringTable;
 class PDBFileBuilder;
 class PublicsStream;
@@ -87,6 +88,8 @@ class PDBFile : public msf::IMSFFile {
   createIndexedStream(uint16_t SN) const;
   Expected<std::unique_ptr<msf::MappedBlockStream>>
   safelyCreateIndexedStream(uint32_t StreamIndex) const;
+  Expected<std::unique_ptr<msf::MappedBlockStream>>
+  safelyCreateNamedStream(StringRef Name);
 
   msf::MSFStreamLayout getStreamLayout(uint32_t StreamIdx) const;
   msf::MSFStreamLayout getFpmStreamLayout() const;
@@ -102,6 +105,7 @@ class PDBFile : public msf::IMSFFile {
   Expected<PublicsStream &> getPDBPublicsStream();
   Expected<SymbolStream &> getPDBSymbolStream();
   Expected<PDBStringTable &> getStringTable();
+  Expected<InjectedSourceStream &> getInjectedSourceStream();
 
   BumpPtrAllocator &getAllocator() { return Allocator; }
 
@@ -113,6 +117,7 @@ class PDBFile : public msf::IMSFFile {
   bool hasPDBSymbolStream();
   bool hasPDBTpiStream() const;
   bool hasPDBStringTable();
+  bool hasPDBInjectedSourceStream();
 
   uint32_t getPointerSize();
 
@@ -133,6 +138,7 @@ class PDBFile : public msf::IMSFFile {
   std::unique_ptr<SymbolStream> Symbols;
   std::unique_ptr<msf::MappedBlockStream> DirectoryStream;
   std::unique_ptr<msf::MappedBlockStream> StringTableStream;
+  std::unique_ptr<InjectedSourceStream> InjectedSources;
   std::unique_ptr<PDBStringTable> Strings;
 };
 }

include/llvm/IR/IntrinsicsAMDGPU.td

@@ -296,29 +296,33 @@ def int_amdgcn_fract : Intrinsic<
   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
 >;
 
-def int_amdgcn_cvt_pkrtz : Intrinsic<
-  [llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty],
-  [IntrNoMem, IntrSpeculatable]
+def int_amdgcn_cvt_pkrtz : GCCBuiltin<"__builtin_amdgcn_cvt_pkrtz">,
+  Intrinsic<[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty],
+            [IntrNoMem, IntrSpeculatable]
 >;
 
-def int_amdgcn_cvt_pknorm_i16 : Intrinsic<
-  [llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty],
-  [IntrNoMem, IntrSpeculatable]
+def int_amdgcn_cvt_pknorm_i16 :
+  GCCBuiltin<"__builtin_amdgcn_cvt_pknorm_i16">,
+  Intrinsic<[llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty],
+            [IntrNoMem, IntrSpeculatable]
 >;
 
-def int_amdgcn_cvt_pknorm_u16 : Intrinsic<
-  [llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty],
-  [IntrNoMem, IntrSpeculatable]
+def int_amdgcn_cvt_pknorm_u16 :
+  GCCBuiltin<"__builtin_amdgcn_cvt_pknorm_u16">,
+  Intrinsic<[llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty],
+            [IntrNoMem, IntrSpeculatable]
 >;
 
-def int_amdgcn_cvt_pk_i16 : Intrinsic<
+def int_amdgcn_cvt_pk_i16 :
+    GCCBuiltin<"__builtin_amdgcn_cvt_pk_i16">,
+    Intrinsic<
   [llvm_v2i16_ty], [llvm_i32_ty, llvm_i32_ty],
   [IntrNoMem, IntrSpeculatable]
 >;
 
-def int_amdgcn_cvt_pk_u16 : Intrinsic<
-  [llvm_v2i16_ty], [llvm_i32_ty, llvm_i32_ty],
-  [IntrNoMem, IntrSpeculatable]
+def int_amdgcn_cvt_pk_u16 : GCCBuiltin<"__builtin_amdgcn_cvt_pk_u16">,
+  Intrinsic<[llvm_v2i16_ty], [llvm_i32_ty, llvm_i32_ty],
+    [IntrNoMem, IntrSpeculatable]
 >;
 
 def int_amdgcn_class : Intrinsic<
@@ -1246,13 +1250,13 @@ def int_amdgcn_ds_swizzle :
             [IntrNoMem, IntrConvergent, ImmArg<1>]>;
 
 def int_amdgcn_ubfe : Intrinsic<[llvm_anyint_ty],
-  [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty],
-  [IntrNoMem, IntrSpeculatable]
+    [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty],
+    [IntrNoMem, IntrSpeculatable]
 >;
 
 def int_amdgcn_sbfe : Intrinsic<[llvm_anyint_ty],
-  [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty],
-  [IntrNoMem, IntrSpeculatable]
+    [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty],
+    [IntrNoMem, IntrSpeculatable]
 >;
 
 def int_amdgcn_lerp :
@@ -1340,13 +1344,14 @@ def int_amdgcn_writelane :
   [IntrNoMem, IntrConvergent]
 >;
 
-def int_amdgcn_alignbit : Intrinsic<[llvm_i32_ty],
+def int_amdgcn_alignbit :
+  GCCBuiltin<"__builtin_amdgcn_alignbit">, Intrinsic<[llvm_i32_ty],
   [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
   [IntrNoMem, IntrSpeculatable]
 >;
 
-def int_amdgcn_alignbyte : Intrinsic<[llvm_i32_ty],
-  [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
+def int_amdgcn_alignbyte : GCCBuiltin<"__builtin_amdgcn_alignbyte">,
+  Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
   [IntrNoMem, IntrSpeculatable]
 >;
 
@@ -1515,13 +1520,13 @@ def int_amdgcn_ds_bpermute :
 //===----------------------------------------------------------------------===//
 
 // llvm.amdgcn.permlane16 <old> <src0> <src1> <src2> <fi> <bound_control>
-def int_amdgcn_permlane16 :
+def int_amdgcn_permlane16 : GCCBuiltin<"__builtin_amdgcn_permlane16">,
   Intrinsic<[llvm_i32_ty],
             [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i1_ty],
             [IntrNoMem, IntrConvergent, ImmArg<4>, ImmArg<5>]>;
 
 // llvm.amdgcn.permlanex16 <old> <src0> <src1> <src2> <fi> <bound_control>
-def int_amdgcn_permlanex16 :
+def int_amdgcn_permlanex16 : GCCBuiltin<"__builtin_amdgcn_permlanex16">,
   Intrinsic<[llvm_i32_ty],
             [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i1_ty],
             [IntrNoMem, IntrConvergent, ImmArg<4>, ImmArg<5>]>;

include/llvm/IR/IntrinsicsARM.td

@@ -19,7 +19,7 @@ let TargetPrefix = "arm" in {  // All intrinsics start with "llvm.arm.".
 // A space-consuming intrinsic primarily for testing ARMConstantIslands. The
 // first argument is the number of bytes this "instruction" takes up, the second
 // and return value are essentially chains, used to force ordering during ISel.
-def int_arm_space : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], []>;
+def int_arm_space : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [ImmArg<0>]>;
 
 // 16-bit multiplications
 def int_arm_smulbb : GCCBuiltin<"__builtin_arm_smulbb">,

include/llvm/IR/IntrinsicsWebAssembly.td

@@ -124,4 +124,13 @@ def int_wasm_data_drop :
             [llvm_i32_ty],
             [IntrNoDuplicate, IntrHasSideEffects, ImmArg<0>]>;
 
+//===----------------------------------------------------------------------===//
+// Thread-local storage intrinsics
+//===----------------------------------------------------------------------===//
+
+def int_wasm_tls_size :
+  Intrinsic<[llvm_anyint_ty],
+            [],
+            [IntrNoMem, IntrSpeculatable]>;
+
 } // TargetPrefix = "wasm"

include/llvm/MC/MCSectionWasm.h

@@ -66,7 +66,8 @@ class MCSectionWasm final : public MCSection {
   bool isVirtualSection() const override;
 
   bool isWasmData() const {
-    return Kind.isGlobalWriteableData() || Kind.isReadOnly();
+    return Kind.isGlobalWriteableData() || Kind.isReadOnly() ||
+           Kind.isThreadLocal();
   }
 
   bool isUnique() const { return UniqueID != ~0U; }

lib/CodeGen/AsmPrinter/DwarfCompileUnit.cpp

@@ -543,7 +543,8 @@ DIE *DwarfCompileUnit::constructInlinedScopeDIE(LexicalScope *Scope) {
   addUInt(*ScopeDIE, dwarf::DW_AT_call_file, None,
           getOrCreateSourceID(IA->getFile()));
   addUInt(*ScopeDIE, dwarf::DW_AT_call_line, None, IA->getLine());
-  addUInt(*ScopeDIE, dwarf::DW_AT_call_column, None, IA->getColumn());
+  if (IA->getColumn())
+    addUInt(*ScopeDIE, dwarf::DW_AT_call_column, None, IA->getColumn());
   if (IA->getDiscriminator() && DD->getDwarfVersion() >= 4)
     addUInt(*ScopeDIE, dwarf::DW_AT_GNU_discriminator, None,
             IA->getDiscriminator());

lib/CodeGen/GlobalISel/CallLowering.cpp

@@ -163,10 +163,19 @@ bool CallLowering::handleAssignments(MachineIRBuilder &MIRBuilder,
                                      ValueHandler &Handler) const {
   MachineFunction &MF = MIRBuilder.getMF();
   const Function &F = MF.getFunction();
-  const DataLayout &DL = F.getParent()->getDataLayout();
-
   SmallVector<CCValAssign, 16> ArgLocs;
   CCState CCInfo(F.getCallingConv(), F.isVarArg(), MF, ArgLocs, F.getContext());
+  return handleAssignments(CCInfo, ArgLocs, MIRBuilder, Args, Handler);
+}
+
+bool CallLowering::handleAssignments(CCState &CCInfo,
+                                     SmallVectorImpl<CCValAssign> &ArgLocs,
+                                     MachineIRBuilder &MIRBuilder,
+                                     ArrayRef<ArgInfo> Args,
+                                     ValueHandler &Handler) const {
+  MachineFunction &MF = MIRBuilder.getMF();
+  const Function &F = MF.getFunction();
+  const DataLayout &DL = F.getParent()->getDataLayout();
 
   unsigned NumArgs = Args.size();
   for (unsigned i = 0; i != NumArgs; ++i) {

lib/DebugInfo/PDB/CMakeLists.txt

@@ -47,9 +47,11 @@ add_pdb_impl_folder(Native
   Native/HashTable.cpp
   Native/InfoStream.cpp
   Native/InfoStreamBuilder.cpp
+  Native/InjectedSourceStream.cpp
   Native/ModuleDebugStream.cpp
   Native/NativeCompilandSymbol.cpp
   Native/NativeEnumGlobals.cpp
+  Native/NativeEnumInjectedSources.cpp
   Native/NativeEnumModules.cpp
   Native/NativeEnumTypes.cpp
   Native/NativeExeSymbol.cpp