Merged main:508a697acd18 into amd-gfx:57bd89831925

Local branch amd-gfx 57bd898 Merged main:747e0d9f0aad into amd-gfx:fdf184dabd64
Remote branch main 508a697 [JITLink] Allow multiple relocations at same offset in EHFrameEdgeFixer (llvm#68252)

Author: SC llvm team

clang/lib/Format/TokenAnnotator.cpp

@@ -4197,8 +4197,7 @@ bool TokenAnnotator::spaceRequiredBetween(const AnnotatedLine &Line,
   if ((Left.is(tok::l_brace) && Left.isNot(BK_Block)) ||
       (Right.is(tok::r_brace) && Right.MatchingParen &&
        Right.MatchingParen->isNot(BK_Block))) {
-    return Style.Cpp11BracedListStyle ? Style.SpacesInParensOptions.Other
-                                      : true;
+    return !Style.Cpp11BracedListStyle || Style.SpacesInParensOptions.Other;
   }
   if (Left.is(TT_BlockComment)) {
     // No whitespace in x(/*foo=*/1), except for JavaScript.

llvm/include/llvm/Config/llvm-config.h.cmake

@@ -16,7 +16,7 @@
 
 /* Indicate that this is LLVM compiled from the amd-gfx branch. */
 #define LLVM_HAVE_BRANCH_AMD_GFX
-#define LLVM_MAIN_REVISION 478294
+#define LLVM_MAIN_REVISION 478303
 
 /* Define if LLVM_ENABLE_DUMP is enabled */
 #cmakedefine LLVM_ENABLE_DUMP

llvm/lib/ExecutionEngine/JITLink/EHFrameSupport.cpp

@@ -126,16 +126,23 @@ Error EHFrameEdgeFixer::processBlock(ParseContext &PC, Block &B) {
   }
 
   // Find the offsets of any existing edges from this block.
-  BlockEdgeMap BlockEdges;
+  BlockEdgesInfo BlockEdges;
   for (auto &E : B.edges())
     if (E.isRelocation()) {
-      if (BlockEdges.count(E.getOffset()))
-        return make_error<JITLinkError>(
-            "Multiple relocations at offset " +
-            formatv("{0:x16}", E.getOffset()) + " in " + EHFrameSectionName +
-            " block at address " + formatv("{0:x16}", B.getAddress()));
-
-      BlockEdges[E.getOffset()] = EdgeTarget(E);
+      // Check if we already saw more than one relocation at this offset.
+      if (BlockEdges.Multiple.contains(E.getOffset()))
+        continue;
+
+      // Otherwise check if we previously had exactly one relocation at this
+      // offset. If so, we now have a second one and move it from the TargetMap
+      // into the Multiple set.
+      auto It = BlockEdges.TargetMap.find(E.getOffset());
+      if (It != BlockEdges.TargetMap.end()) {
+        BlockEdges.TargetMap.erase(It);
+        BlockEdges.Multiple.insert(E.getOffset());
+      } else {
+        BlockEdges.TargetMap[E.getOffset()] = EdgeTarget(E);
+      }
     }
 
   BinaryStreamReader BlockReader(
@@ -172,7 +179,7 @@ Error EHFrameEdgeFixer::processBlock(ParseContext &PC, Block &B) {
 
 Error EHFrameEdgeFixer::processCIE(ParseContext &PC, Block &B,
                                    size_t CIEDeltaFieldOffset,
-                                   const BlockEdgeMap &BlockEdges) {
+                                   const BlockEdgesInfo &BlockEdges) {
 
   LLVM_DEBUG(dbgs() << "    Record is CIE\n");
 
@@ -285,7 +292,7 @@ Error EHFrameEdgeFixer::processCIE(ParseContext &PC, Block &B,
 Error EHFrameEdgeFixer::processFDE(ParseContext &PC, Block &B,
                                    size_t CIEDeltaFieldOffset,
                                    uint32_t CIEDelta,
-                                   const BlockEdgeMap &BlockEdges) {
+                                   const BlockEdgesInfo &BlockEdges) {
   LLVM_DEBUG(dbgs() << "    Record is FDE\n");
 
   orc::ExecutorAddr RecordAddress = B.getAddress();
@@ -303,12 +310,17 @@ Error EHFrameEdgeFixer::processFDE(ParseContext &PC, Block &B,
 
   {
     // Process the CIE pointer field.
-    auto CIEEdgeItr = BlockEdges.find(CIEDeltaFieldOffset);
+    if (BlockEdges.Multiple.contains(CIEDeltaFieldOffset))
+      return make_error<JITLinkError>(
+          "CIE pointer field already has multiple edges at " +
+          formatv("{0:x16}", RecordAddress + CIEDeltaFieldOffset));
+
+    auto CIEEdgeItr = BlockEdges.TargetMap.find(CIEDeltaFieldOffset);
 
     orc::ExecutorAddr CIEAddress =
         RecordAddress + orc::ExecutorAddrDiff(CIEDeltaFieldOffset) -
         orc::ExecutorAddrDiff(CIEDelta);
-    if (CIEEdgeItr == BlockEdges.end()) {
+    if (CIEEdgeItr == BlockEdges.TargetMap.end()) {
       LLVM_DEBUG({
         dbgs() << "        Adding edge at "
                << (RecordAddress + CIEDeltaFieldOffset)
@@ -497,7 +509,7 @@ Error EHFrameEdgeFixer::skipEncodedPointer(uint8_t PointerEncoding,
 }
 
 Expected<Symbol *> EHFrameEdgeFixer::getOrCreateEncodedPointerEdge(
-    ParseContext &PC, const BlockEdgeMap &BlockEdges, uint8_t PointerEncoding,
+    ParseContext &PC, const BlockEdgesInfo &BlockEdges, uint8_t PointerEncoding,
     BinaryStreamReader &RecordReader, Block &BlockToFix,
     size_t PointerFieldOffset, const char *FieldName) {
   using namespace dwarf;
@@ -508,8 +520,8 @@ Expected<Symbol *> EHFrameEdgeFixer::getOrCreateEncodedPointerEdge(
   // If there's already an edge here then just skip the encoded pointer and
   // return the edge's target.
   {
-    auto EdgeI = BlockEdges.find(PointerFieldOffset);
-    if (EdgeI != BlockEdges.end()) {
+    auto EdgeI = BlockEdges.TargetMap.find(PointerFieldOffset);
+    if (EdgeI != BlockEdges.TargetMap.end()) {
       LLVM_DEBUG({
         dbgs() << "      Existing edge at "
                << (BlockToFix.getAddress() + PointerFieldOffset) << " to "
@@ -522,6 +534,10 @@ Expected<Symbol *> EHFrameEdgeFixer::getOrCreateEncodedPointerEdge(
         return std::move(Err);
       return EdgeI->second.Target;
     }
+
+    if (BlockEdges.Multiple.contains(PointerFieldOffset))
+      return make_error<JITLinkError>("Multiple relocations at offset " +
+                                      formatv("{0:x16}", PointerFieldOffset));
   }
 
   // Switch absptr to corresponding udata encoding.

llvm/lib/ExecutionEngine/JITLink/EHFrameSupportImpl.h

@@ -60,7 +60,11 @@ class EHFrameEdgeFixer {
     Edge::AddendT Addend = 0;
   };
 
-  using BlockEdgeMap = DenseMap<Edge::OffsetT, EdgeTarget>;
+  struct BlockEdgesInfo {
+    DenseMap<Edge::OffsetT, EdgeTarget> TargetMap;
+    DenseSet<Edge::OffsetT> Multiple;
+  };
+
   using CIEInfosMap = DenseMap<orc::ExecutorAddr, CIEInformation>;
 
   struct ParseContext {
@@ -82,9 +86,9 @@ class EHFrameEdgeFixer {
 
   Error processBlock(ParseContext &PC, Block &B);
   Error processCIE(ParseContext &PC, Block &B, size_t CIEDeltaFieldOffset,
-                   const BlockEdgeMap &BlockEdges);
+                   const BlockEdgesInfo &BlockEdges);
   Error processFDE(ParseContext &PC, Block &B, size_t CIEDeltaFieldOffset,
-                   uint32_t CIEDelta, const BlockEdgeMap &BlockEdges);
+                   uint32_t CIEDelta, const BlockEdgesInfo &BlockEdges);
 
   Expected<AugmentationInfo>
   parseAugmentationString(BinaryStreamReader &RecordReader);
@@ -94,9 +98,9 @@ class EHFrameEdgeFixer {
   Error skipEncodedPointer(uint8_t PointerEncoding,
                            BinaryStreamReader &RecordReader);
   Expected<Symbol *> getOrCreateEncodedPointerEdge(
-      ParseContext &PC, const BlockEdgeMap &BlockEdges, uint8_t PointerEncoding,
-      BinaryStreamReader &RecordReader, Block &BlockToFix,
-      size_t PointerFieldOffset, const char *FieldName);
+      ParseContext &PC, const BlockEdgesInfo &BlockEdges,
+      uint8_t PointerEncoding, BinaryStreamReader &RecordReader,
+      Block &BlockToFix, size_t PointerFieldOffset, const char *FieldName);
 
   Expected<Symbol &> getOrCreateSymbol(ParseContext &PC,
                                        orc::ExecutorAddr Addr);

mlir/examples/toy/Ch1/parser/AST.cpp

@@ -12,9 +12,12 @@
 
 #include "toy/AST.h"
 
+#include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/Twine.h"
 #include "llvm/ADT/TypeSwitch.h"
+#include "llvm/Support/Casting.h"
 #include "llvm/Support/raw_ostream.h"
+#include <string>
 
 using namespace toy;
 

mlir/examples/toy/Ch1/toyc.cpp

@@ -10,13 +10,18 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include "toy/AST.h"
+#include "toy/Lexer.h"
 #include "toy/Parser.h"
 
 #include "llvm/ADT/StringRef.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/ErrorOr.h"
 #include "llvm/Support/MemoryBuffer.h"
 #include "llvm/Support/raw_ostream.h"
+#include <memory>
+#include <string>
+#include <system_error>
 
 using namespace toy;
 namespace cl = llvm::cl;

mlir/examples/toy/Ch2/mlir/Dialect.cpp

@@ -13,10 +13,22 @@
 
 #include "toy/Dialect.h"
 
+#include "mlir/IR/Attributes.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinTypes.h"
 #include "mlir/IR/OpImplementation.h"
+#include "mlir/IR/Operation.h"
+#include "mlir/IR/OperationSupport.h"
+#include "mlir/IR/Value.h"
 #include "mlir/Interfaces/FunctionImplementation.h"
+#include "mlir/Support/LLVM.h"
+#include "mlir/Support/LogicalResult.h"
+#include "llvm/ADT/ArrayRef.h"
+#include "llvm/ADT/STLExtras.h"
+#include "llvm/ADT/StringRef.h"
+#include "llvm/Support/Casting.h"
+#include <algorithm>
+#include <string>
 
 using namespace mlir;
 using namespace mlir::toy;

mlir/examples/toy/Ch2/mlir/MLIRGen.cpp

@@ -12,20 +12,31 @@
 //===----------------------------------------------------------------------===//
 
 #include "toy/MLIRGen.h"
+#include "mlir/IR/Block.h"
+#include "mlir/IR/Diagnostics.h"
+#include "mlir/IR/Value.h"
+#include "mlir/Support/LogicalResult.h"
 #include "toy/AST.h"
 #include "toy/Dialect.h"
 
-#include "mlir/IR/Attributes.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/BuiltinTypes.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/IR/Verifier.h"
+#include "toy/Lexer.h"
 
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/ScopedHashTable.h"
-#include "llvm/Support/raw_ostream.h"
+#include "llvm/ADT/SmallVector.h"
+#include "llvm/ADT/StringRef.h"
+#include "llvm/ADT/Twine.h"
+#include <cassert>
+#include <cstdint>
+#include <functional>
 #include <numeric>
+#include <optional>
+#include <vector>
 
 using namespace mlir::toy;
 using namespace toy;

mlir/examples/toy/Ch2/parser/AST.cpp

@@ -12,9 +12,12 @@
 
 #include "toy/AST.h"
 
+#include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/Twine.h"
 #include "llvm/ADT/TypeSwitch.h"
+#include "llvm/Support/Casting.h"
 #include "llvm/Support/raw_ostream.h"
+#include <string>
 
 using namespace toy;
 

openmp/libomptarget/DeviceRTL/src/State.cpp

@@ -265,7 +265,7 @@ void state::enterDataEnvironment(IdentTy *Ident) {
     uint32_t Bytes = sizeof(ThreadStates[0]) * mapping::getMaxTeamThreads();
     void *ThreadStatesPtr =
         memory::allocGlobal(Bytes, "Thread state array allocation");
-    memset(ThreadStatesPtr, '0', Bytes);
+    memset(ThreadStatesPtr, 0, Bytes);
     if (!atomic::cas(ThreadStatesBitsPtr, uintptr_t(0),
                      reinterpret_cast<uintptr_t>(ThreadStatesPtr),
                      atomic::seq_cst, atomic::seq_cst))

openmp/libomptarget/test/offloading/thread_state_1.c

@@ -1,30 +1,35 @@
-// RUN: %libomptarget-compile-run-and-check-generic
+// Still broken "without optimizations"
+// XUN: %libomptarget-compile-run-and-check-generic
 // RUN: %libomptarget-compileopt-run-and-check-generic
 
-// UNSUPPORTED: amdgcn-amd-amdhsa
-
 #include <omp.h>
 #include <stdio.h>
 
 int main() {
   // TODO: Test all ICVs on all levels
-  int lvl = 333, tid = 666, nt = 999;
-#pragma omp target teams map(tofrom : lvl, tid, nt) num_teams(2)
+  int o_lvl = 111, i_lvl = 222, o_tid = 333, i_tid = 333, o_nt = 444,
+      i_nt = 555;
+#pragma omp target teams map(tofrom : o_lvl, i_lvl, o_tid, i_tid, o_nt, i_nt)  \
+    num_teams(2) thread_limit(2)
   {
     if (omp_get_team_num() == 0) {
 #pragma omp parallel num_threads(128)
-      if (omp_get_thread_num() == 17) {
+      if (omp_get_thread_num() == omp_get_num_threads() - 1) {
+        o_lvl = omp_get_level();
+        o_tid = omp_get_thread_num();
+        o_nt = omp_get_num_threads();
 #pragma omp parallel num_threads(64)
         if (omp_get_thread_num() == omp_get_num_threads() - 1) {
-          lvl = omp_get_level();
-          tid = omp_get_thread_num();
-          nt = omp_get_num_threads();
+          i_lvl = omp_get_level();
+          i_tid = omp_get_thread_num();
+          i_nt = omp_get_num_threads();
         }
       }
     }
   }
-  // TODO: This is wrong, but at least it doesn't crash
-  // CHECK: lvl: 2, tid: 0, nt: 1
-  printf("lvl: %i, tid: %i, nt: %i\n", lvl, tid, nt);
+  // CHECK: outer: lvl: 1, tid: 1, nt: 2
+  // CHECK: inner: lvl: 2, tid: 0, nt: 1
+  printf("outer: lvl: %i, tid: %i, nt: %i\n", o_lvl, o_tid, o_nt);
+  printf("inner: lvl: %i, tid: %i, nt: %i\n", i_lvl, i_tid, i_nt);
   return 0;
 }