Merged main:86bb713142c3 into amd-gfx:fd6ee38e1eb7

Local branch amd-gfx fd6ee38 Merged main:a24418375a70 into amd-gfx:923f1fe3b764
Remote branch main 86bb713 [OpenMP][FIX] Enlarge thread state array, improve test and add second

Author: SC llvm team

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 478308
+#define LLVM_MAIN_REVISION 478311
 
 /* Define if LLVM_ENABLE_DUMP is enabled */
 #cmakedefine LLVM_ENABLE_DUMP

llvm/tools/llvm-profgen/ProfiledBinary.cpp

@@ -480,12 +480,6 @@ bool ProfiledBinary::dissassembleSymbol(std::size_t SI, ArrayRef<uint8_t> Bytes,
   if (ShowDisassembly)
     outs() << '<' << SymbolName << ">:\n";
 
-  auto WarnInvalidInsts = [](uint64_t Start, uint64_t End) {
-    WithColor::warning() << "Invalid instructions at "
-                         << format("%8" PRIx64, Start) << " - "
-                         << format("%8" PRIx64, End) << "\n";
-  };
-
   uint64_t Address = StartAddress;
   // Size of a consecutive invalid instruction range starting from Address -1
   // backwards.
@@ -578,7 +572,8 @@ bool ProfiledBinary::dissassembleSymbol(std::size_t SI, ArrayRef<uint8_t> Bytes,
       }
 
       if (InvalidInstLength) {
-        WarnInvalidInsts(Address - InvalidInstLength, Address - 1);
+        AddrsWithInvalidInstruction.insert(
+            {Address - InvalidInstLength, Address - 1});
         InvalidInstLength = 0;
       }
     } else {
@@ -589,7 +584,8 @@ bool ProfiledBinary::dissassembleSymbol(std::size_t SI, ArrayRef<uint8_t> Bytes,
   }
 
   if (InvalidInstLength)
-    WarnInvalidInsts(Address - InvalidInstLength, Address - 1);
+    AddrsWithInvalidInstruction.insert(
+        {Address - InvalidInstLength, Address - 1});
 
   if (ShowDisassembly)
     outs() << "\n";
@@ -708,6 +704,19 @@ void ProfiledBinary::disassemble(const ELFObjectFileBase *Obj) {
     }
   }
 
+  if (!AddrsWithInvalidInstruction.empty()) {
+    if (ShowDetailedWarning) {
+      for (auto &Addr : AddrsWithInvalidInstruction) {
+        WithColor::warning()
+            << "Invalid instructions at " << format("%8" PRIx64, Addr.first)
+            << " - " << format("%8" PRIx64, Addr.second) << "\n";
+      }
+    }
+    WithColor::warning() << "Found " << AddrsWithInvalidInstruction.size()
+                         << " invalid instructions\n";
+    AddrsWithInvalidInstruction.clear();
+  }
+
   // Dissassemble rodata section to check if FS discriminator symbol exists.
   checkUseFSDiscriminator(Obj, AllSymbols);
 }
@@ -792,10 +801,12 @@ void ProfiledBinary::loadSymbolsFromDWARFUnit(DWARFUnit &CompilationUnit) {
         FRange.StartAddress = StartAddress;
         FRange.EndAddress = EndAddress;
       } else {
-        WithColor::warning()
-            << "Duplicated symbol start address at "
-            << format("%8" PRIx64, StartAddress) << " "
-            << R.first->second.getFuncName() << " and " << Name << "\n";
+        AddrsWithMultipleSymbols.insert(StartAddress);
+        if (ShowDetailedWarning)
+          WithColor::warning()
+              << "Duplicated symbol start address at "
+              << format("%8" PRIx64, StartAddress) << " "
+              << R.first->second.getFuncName() << " and " << Name << "\n";
       }
     }
   }
@@ -839,14 +850,18 @@ void ProfiledBinary::loadSymbolsFromDWARF(ObjectFile &Obj) {
   if (BinaryFunctions.empty())
     WithColor::warning() << "Loading of DWARF info completed, but no binary "
                             "functions have been retrieved.\n";
-
-
   // Populate the hash binary function map for MD5 function name lookup. This
   // is done after BinaryFunctions are finalized.
   for (auto &BinaryFunction : BinaryFunctions) {
     HashBinaryFunctions[MD5Hash(StringRef(BinaryFunction.first))] =
         &BinaryFunction.second;
   }
+
+  if (!AddrsWithMultipleSymbols.empty()) {
+    WithColor::warning() << "Found " << AddrsWithMultipleSymbols.size()
+                         << " start addresses with multiple symbols\n";
+    AddrsWithMultipleSymbols.clear();
+  }
 }
 
 void ProfiledBinary::populateSymbolListFromDWARF(
@@ -881,7 +896,8 @@ SampleContextFrameVector ProfiledBinary::symbolize(const InstructionPointer &IP,
   SampleContextFrameVector CallStack;
   for (int32_t I = InlineStack.getNumberOfFrames() - 1; I >= 0; I--) {
     const auto &CallerFrame = InlineStack.getFrame(I);
-    if (CallerFrame.FunctionName.empty() || (CallerFrame.FunctionName == "<invalid>"))
+    if (CallerFrame.FunctionName.empty() ||
+        (CallerFrame.FunctionName == "<invalid>"))
       break;
 
     StringRef FunctionName(CallerFrame.FunctionName);

llvm/tools/llvm-profgen/ProfiledBinary.h

@@ -230,6 +230,10 @@ class ProfiledBinary {
   // GUID to Elf symbol start address map
   DenseMap<uint64_t, uint64_t> SymbolStartAddrs;
 
+  // These maps are for temporary use of warning diagnosis.
+  DenseSet<int64_t> AddrsWithMultipleSymbols;
+  DenseSet<std::pair<uint64_t, uint64_t>> AddrsWithInvalidInstruction;
+
   // Start address to Elf symbol GUID map
   std::unordered_multimap<uint64_t, uint64_t> StartAddrToSymMap;
 
@@ -529,7 +533,7 @@ class ProfiledBinary {
 
   void flushSymbolizer() { Symbolizer.reset(); }
 
-  MissingFrameInferrer* getMissingContextInferrer() {
+  MissingFrameInferrer *getMissingContextInferrer() {
     return MissingContextInferrer.get();
   }
 

mlir/include/mlir/Dialect/Bufferization/IR/BufferizationBase.td

@@ -60,6 +60,17 @@ def Bufferization_Dialect : Dialect {
     /// arguments during One-Shot Module Bufferize.
     constexpr const static ::llvm::StringLiteral
         kBufferLayoutAttrName = "bufferization.buffer_layout";
+
+    /// An attribute that can be attached to ops with an allocation and/or
+    /// deallocation side effect. It indicates that the op is under a "manual
+    /// deallocation" scheme. In the case of an allocation op, the returned
+    /// value is *not* an automatically managed allocation and assigned an
+    /// ownership of "false". Furthermore, only deallocation ops that are
+    /// guaranteed to deallocate a buffer under "manual deallocation" are
+    /// allowed to have this attribute. (Deallocation ops without this
+    /// attribute are rejected by the ownership-based buffer deallocation pass.)
+    constexpr const static ::llvm::StringLiteral
+        kManualDeallocation = "bufferization.manual_deallocation";
   }];
   let hasOperationAttrVerify = 1;
 }

mlir/lib/Dialect/Bufferization/IR/BufferizationDialect.cpp

@@ -28,6 +28,16 @@ constexpr const ::llvm::StringLiteral BufferizationDialect::kWritableAttrName;
 constexpr const ::llvm::StringLiteral
     BufferizationDialect::kBufferLayoutAttrName;
 
+/// An attribute that can be attached to ops with an allocation and/or
+/// deallocation side effect. It indicates that the op is under a "manual
+/// deallocation" scheme. In the case of an allocation op, the returned
+/// value is *not* an automatically managed allocation and assigned an
+/// ownership of "false". Furthermore, only deallocation ops that are
+/// guaranteed to deallocate a buffer under "manual deallocation" are
+/// allowed to have this attribute. (Deallocation ops without this
+/// attribute are rejected by the ownership-based buffer deallocation pass.)
+constexpr const ::llvm::StringLiteral BufferizationDialect::kManualDeallocation;
+
 //===----------------------------------------------------------------------===//
 // Bufferization Dialect Interfaces
 //===----------------------------------------------------------------------===//
@@ -105,6 +115,16 @@ BufferizationDialect::verifyOperationAttribute(Operation *op,
                                                NamedAttribute attr) {
   using bufferization::BufferizableOpInterface;
 
+  if (attr.getName() == kManualDeallocation) {
+    if (!mlir::hasEffect<MemoryEffects::Allocate>(op) &&
+        !mlir::hasEffect<MemoryEffects::Free>(op))
+      return op->emitOpError("attribute '")
+             << kManualDeallocation
+             << "' can be used only on ops that have an allocation and/or free "
+                "side effect";
+    return success();
+  }
+
   return op->emitError()
          << "attribute '" << attr.getName()
          << "' not supported as an op attribute by the bufferization dialect";

mlir/lib/Dialect/Bufferization/Transforms/OwnershipBasedBufferDeallocation.cpp

@@ -21,6 +21,7 @@
 #include "mlir/Dialect/Bufferization/IR/BufferDeallocationOpInterface.h"
 #include "mlir/Dialect/Bufferization/IR/Bufferization.h"
 #include "mlir/Dialect/Bufferization/Transforms/Passes.h"
+#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
@@ -856,13 +857,32 @@ FailureOr<Operation *> BufferDeallocation::handleInterface(CallOpInterface op) {
 FailureOr<Operation *>
 BufferDeallocation::handleInterface(MemoryEffectOpInterface op) {
   auto *block = op->getBlock();
+  OpBuilder builder = OpBuilder::atBlockBegin(block);
 
-  for (auto operand : llvm::make_filter_range(op->getOperands(), isMemref))
-    if (op.getEffectOnValue<MemoryEffects::Free>(operand).has_value())
-      return op->emitError(
-          "memory free side-effect on MemRef value not supported!");
+  for (auto operand : llvm::make_filter_range(op->getOperands(), isMemref)) {
+    if (op.getEffectOnValue<MemoryEffects::Free>(operand).has_value()) {
+      if (!op->hasAttr(BufferizationDialect::kManualDeallocation))
+        return op->emitError(
+            "memory free side-effect on MemRef value not supported!");
+
+      // Buffers that were allocated under "manual deallocation" may be
+      // manually deallocated. We insert a runtime assertion to cover certain
+      // cases of invalid IR where an automatically managed buffer allocation
+      // is manually deallocated. This is not a bulletproof check!
+      OpBuilder::InsertionGuard g(builder);
+      builder.setInsertionPoint(op);
+      Ownership ownership = state.getOwnership(operand, block);
+      if (ownership.isUnique()) {
+        Value ownershipInverted = builder.create<arith::XOrIOp>(
+            op.getLoc(), ownership.getIndicator(),
+            buildBoolValue(builder, op.getLoc(), true));
+        builder.create<cf::AssertOp>(
+            op.getLoc(), ownershipInverted,
+            "expected that the block does not have ownership");
+      }
+    }
+  }
 
-  OpBuilder builder = OpBuilder::atBlockBegin(block);
   for (auto res : llvm::make_filter_range(op->getResults(), isMemref)) {
     auto allocEffect = op.getEffectOnValue<MemoryEffects::Allocate>(res);
     if (allocEffect.has_value()) {
@@ -880,6 +900,15 @@ BufferDeallocation::handleInterface(MemoryEffectOpInterface op) {
         continue;
       }
 
+      if (op->hasAttr(BufferizationDialect::kManualDeallocation)) {
+        // This allocation will be deallocated manually. Assign an ownership of
+        // "false", so that it will never be deallocated by the buffer
+        // deallocation pass.
+        state.resetOwnerships(res, block);
+        state.updateOwnership(res, buildBoolValue(builder, op.getLoc(), false));
+        continue;
+      }
+
       state.updateOwnership(res, buildBoolValue(builder, op.getLoc(), true));
       state.addMemrefToDeallocate(res, block);
     }

mlir/test/Dialect/Bufferization/Transforms/OwnershipBasedBufferDeallocation/dealloc-memoryeffect-interface.mlir

@@ -124,3 +124,35 @@ func.func @op_without_aliasing_and_allocation() -> memref<4xf32> {
 //       CHECK:     [[CLONE:%.+]] = bufferization.clone [[GLOBAL]]
 //       CHECK:     scf.yield [[CLONE]] :
 //       CHECK:   return [[RES]] :
+
+// -----
+
+// Allocations with "bufferization.manual_deallocation" are assigned an
+// ownership of "false".
+
+func.func @manual_deallocation(%c: i1, %f: f32, %idx: index) -> f32 {
+  %0 = memref.alloc() {bufferization.manual_deallocation} : memref<5xf32>
+  linalg.fill ins(%f : f32) outs(%0 : memref<5xf32>)
+  %1 = memref.alloc() : memref<5xf32>
+  linalg.fill ins(%f : f32) outs(%1 : memref<5xf32>)
+  %2 = arith.select %c, %0, %1 : memref<5xf32>
+  %3 = memref.load %2[%idx] : memref<5xf32>
+
+  // Only buffers that are under "manual deallocation" are allowed to be
+  // deallocated with memref.dealloc. For consistency reasons, the
+  // manual_deallocation attribute must also be specified. A runtime insertion
+  // is inserted to ensure that we do not have ownership. (This is not a
+  // bulletproof check, but covers some cases of invalid IR.)
+  memref.dealloc %0 {bufferization.manual_deallocation} : memref<5xf32>
+
+  return %3 : f32
+}
+
+// CHECK-LABEL: func @manual_deallocation(
+//       CHECK:   %[[true:.*]] = arith.constant true
+//       CHECK:   %[[manual_alloc:.*]] = memref.alloc() {bufferization.manual_deallocation} : memref<5xf32>
+//       CHECK:   %[[managed_alloc:.*]] = memref.alloc() : memref<5xf32>
+//       CHECK:   %[[selected:.*]] = arith.select
+//       CHECK:   cf.assert %[[true]], "expected that the block does not have ownership"
+//       CHECK:   memref.dealloc %[[manual_alloc]]
+//       CHECK:   bufferization.dealloc (%[[managed_alloc]] : memref<5xf32>) if (%[[true]])

mlir/test/Dialect/Bufferization/invalid.mlir

@@ -136,3 +136,10 @@ func.func @invalid_dealloc_wrong_number_of_results(%arg0: memref<2xf32>, %arg1:
   %0:3 = "bufferization.dealloc"(%arg0, %arg1, %arg2, %arg2, %arg1) <{operandSegmentSizes = array<i32: 2, 2, 1>}> : (memref<2xf32>, memref<4xi32>, i1, i1, memref<4xi32>) -> (i1, i1, i1)
   return %0#0 : i1
 }
+
+// -----
+
+func.func @invalid_manual_deallocation() {
+  // expected-error @below{{op attribute 'bufferization.manual_deallocation' can be used only on ops that have an allocation and/or free side effect}}
+  arith.constant {bufferization.manual_deallocation} 0  : index
+}

openmp/libomptarget/DeviceRTL/src/State.cpp

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

openmp/libomptarget/test/offloading/thread_state_1.c

@@ -1,5 +1,4 @@
-// Still broken "without optimizations"
-// XUN: %libomptarget-compile-run-and-check-generic
+// RUN: %libomptarget-compile-run-and-check-generic
 // RUN: %libomptarget-compileopt-run-and-check-generic
 
 #include <omp.h>
@@ -10,10 +9,10 @@ int main() {
   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)
+    num_teams(2) thread_limit(64)
   {
     if (omp_get_team_num() == 0) {
-#pragma omp parallel num_threads(128)
+#pragma omp parallel num_threads(64)
       if (omp_get_thread_num() == omp_get_num_threads() - 1) {
         o_lvl = omp_get_level();
         o_tid = omp_get_thread_num();
@@ -27,9 +26,13 @@ int main() {
       }
     }
   }
-  // CHECK: outer: lvl: 1, tid: 1, nt: 2
-  // CHECK: inner: lvl: 2, tid: 0, nt: 1
+  if (o_lvl == 1 && o_tid == o_nt - 1 && o_nt > 1 && i_lvl == 2 && i_tid == 0 &&
+      i_nt == 1) {
+    // CHECK: Success
+    printf("Success\n");
+    return 0;
+  }
   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;
+  return 1;
 }

openmp/libomptarget/test/offloading/thread_state_2.c

@@ -0,0 +1,40 @@
+// This fails when optimized for now.
+// RUN: %libomptarget-compile-run-and-check-generic
+// XUN: %libomptarget-compileopt-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  // TODO: Test all ICVs on all levels
+  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(64)
+  {
+    omp_set_max_active_levels(1);
+    if (omp_get_team_num() == 0) {
+#pragma omp parallel num_threads(64)
+      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) {
+          i_lvl = omp_get_level();
+          i_tid = omp_get_thread_num();
+          i_nt = omp_get_num_threads();
+        }
+      }
+    }
+  }
+  if (o_lvl == 1 && o_tid == o_nt - 1 && o_nt > 1 && i_lvl == 2 && i_tid == 0 &&
+      i_nt == 1) {
+    // CHECK: Success
+    printf("Success\n");
+    return 0;
+  }
+  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 1;
+}