merge main into amd-stg-open

  locally : reverts d3921e4 [OpenMP] Basic BumpAllocator for (AMD)GPUs (llvm#69806)

Change-Id: Id512e729870279855744ce65bfc69e2155fb68ee

Author: ronlieb

openmp/libomptarget/DeviceRTL/include/Allocator.h

@@ -0,0 +1,44 @@
+//===-------- Allocator.h - OpenMP memory allocator interface ---- 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 OMPTARGET_ALLOCATOR_H
+#define OMPTARGET_ALLOCATOR_H
+
+#include "Types.h"
+
+// Forward declaration.
+struct KernelEnvironmentTy;
+
+#pragma omp begin declare target device_type(nohost)
+
+namespace ompx {
+
+namespace allocator {
+
+static uint64_t constexpr ALIGNMENT = 16;
+
+/// Initialize the allocator according to \p KernelEnvironment
+void init(bool IsSPMD, KernelEnvironmentTy &KernelEnvironment);
+
+/// Allocate \p Size bytes.
+[[gnu::alloc_size(1), gnu::assume_aligned(ALIGNMENT), gnu::malloc]] void *
+alloc(uint64_t Size);
+
+/// Free the allocation pointed to by \p Ptr.
+void free(void *Ptr);
+
+} // namespace allocator
+
+} // namespace ompx
+
+#pragma omp end declare target
+
+#endif

openmp/libomptarget/DeviceRTL/src/Allocator.cpp

@@ -0,0 +1,80 @@
+//===------ State.cpp - OpenMP State & ICV interface ------------- 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
+//
+//===----------------------------------------------------------------------===//
+//
+//===----------------------------------------------------------------------===//
+
+#include "Allocator.h"
+#include "Configuration.h"
+#include "Environment.h"
+#include "Mapping.h"
+#include "Synchronization.h"
+#include "Types.h"
+#include "Utils.h"
+
+using namespace ompx;
+
+#pragma omp begin declare target device_type(nohost)
+
+[[gnu::used, gnu::retain, gnu::weak,
+  gnu::visibility(
+      "protected")]] DeviceMemoryPoolTy __omp_rtl_device_memory_pool;
+[[gnu::used, gnu::retain, gnu::weak,
+  gnu::visibility("protected")]] DeviceMemoryPoolTrackingTy
+    __omp_rtl_device_memory_pool_tracker;
+
+/// Stateless bump allocator that uses the __omp_rtl_device_memory_pool
+/// directly.
+struct BumpAllocatorTy final {
+
+  void *alloc(uint64_t Size) {
+    Size = utils::roundUp(Size, uint64_t(allocator::ALIGNMENT));
+
+    if (config::isDebugMode(DeviceDebugKind::AllocationTracker)) {
+      atomic::add(&__omp_rtl_device_memory_pool_tracker.NumAllocations, 1,
+                  atomic::seq_cst);
+      atomic::add(&__omp_rtl_device_memory_pool_tracker.AllocationTotal, Size,
+                  atomic::seq_cst);
+      atomic::min(&__omp_rtl_device_memory_pool_tracker.AllocationMin, Size,
+                  atomic::seq_cst);
+      atomic::max(&__omp_rtl_device_memory_pool_tracker.AllocationMax, Size,
+                  atomic::seq_cst);
+    }
+
+    uint64_t *Data =
+        reinterpret_cast<uint64_t *>(&__omp_rtl_device_memory_pool.Ptr);
+    uint64_t End =
+        reinterpret_cast<uint64_t>(Data) + __omp_rtl_device_memory_pool.Size;
+
+    uint64_t OldData = atomic::add(Data, Size, atomic::seq_cst);
+    if (OldData + Size > End)
+      __builtin_trap();
+
+    return reinterpret_cast<void *>(OldData);
+  }
+
+  void free(void *) {}
+};
+
+BumpAllocatorTy BumpAllocator;
+
+/// allocator namespace implementation
+///
+///{
+
+void allocator::init(bool IsSPMD, KernelEnvironmentTy &KernelEnvironment) {
+  // TODO: Check KernelEnvironment for an allocator choice as soon as we have
+  // more than one.
+}
+
+void *allocator::alloc(uint64_t Size) { return BumpAllocator.alloc(Size); }
+
+void allocator::free(void *Ptr) { BumpAllocator.free(Ptr); }
+
+///}
+
+#pragma omp end declare target

openmp/libomptarget/test/offloading/malloc.c

@@ -0,0 +1,37 @@
+// RUN: %libomptarget-compile-generic && %libomptarget-run-generic
+// RUN: %libomptarget-compileopt-generic && %libomptarget-run-generic
+
+#include <stdio.h>
+#include <stdlib.h>
+
+int main() {
+  long unsigned *DP = 0;
+  int N = 128;
+  int Threads = 128;
+  int Teams = 440;
+
+  // Allocate ~55MB on the device.
+#pragma omp target map(from : DP)
+  DP = (long unsigned *)malloc(sizeof(long unsigned) * N * Threads * Teams);
+
+#pragma omp target teams distribute parallel for num_teams(Teams)              \
+    thread_limit(Threads) is_device_ptr(DP)
+  for (int i = 0; i < Threads * Teams; ++i) {
+    for (int j = 0; j < N; ++j) {
+      DP[i * N + j] = i + j;
+    }
+  }
+
+  long unsigned s = 0;
+#pragma omp target teams distribute parallel for num_teams(Teams)              \
+    thread_limit(Threads) reduction(+ : s)
+  for (int i = 0; i < Threads * Teams; ++i) {
+    for (int j = 0; j < N; ++j) {
+      s += DP[i * N + j];
+    }
+  }
+
+  // CHECK: Sum: 203458478080
+  printf("Sum: %li\n", s);
+  return 0;
+}

openmp/libomptarget/test/offloading/malloc_parallel.c

@@ -0,0 +1,42 @@
+// RUN: %libomptarget-compile-generic && %libomptarget-run-generic
+// RUN: %libomptarget-compileopt-generic && %libomptarget-run-generic
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int main() {
+  long unsigned **DP = 0;
+  int N = 128;
+  int Threads = 128;
+  int Teams = 440;
+
+#pragma omp target map(from : DP)
+  DP = (long unsigned **)malloc(sizeof(long unsigned *) * Threads * Teams);
+
+#pragma omp target teams distribute parallel for num_teams(Teams)              \
+    thread_limit(Threads)
+  for (int i = 0; i < Threads * Teams; ++i)
+    DP[i] = (long unsigned *)malloc(sizeof(long unsigned) * N);
+
+#pragma omp target teams distribute parallel for num_teams(Teams)              \
+    thread_limit(Threads)
+  for (int i = 0; i < Threads * Teams; ++i) {
+    for (int j = 0; j < N; ++j) {
+      DP[i][j] = i + j;
+    }
+  }
+
+  long unsigned s = 0;
+#pragma omp target teams distribute parallel for num_teams(Teams)              \
+    thread_limit(Threads) reduction(+ : s)
+  for (int i = 0; i < Threads * Teams; ++i) {
+    for (int j = 0; j < N; ++j) {
+      s += DP[i][j];
+    }
+  }
+
+  // CHECK: Sum: 203458478080
+  printf("Sum: %li\n", s);
+  return 0;
+}

revert_patches.txt

@@ -5,6 +5,9 @@ Johannes: breakage
 Revert "[OpenMP] Disable early vectorization of loads/stores in the runtime " 
   asserts in  snap and MI-Teams
 
+Revert "[OpenMP] Basic BumpAllocator for (AMD)GPUs (#69806) "
+  breaks flang-272343 flang-272343-3 flang-305553
+
 Nicole and Saiyed:
 Revert " [OpenMP][DeviceRTL][AMDGPU] Support code object version 5"