[PATCH] D87841: [openmp][libomptarget] Include header from LLVM source tree

[Jon Chesterfield via Phabricator via llvm-commits]

JonChesterfield added subscribers: saiislam, gregrodgers.
JonChesterfield added inline comments.


================
Comment at: llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h:3
 //
-//                     The LLVM Compiler Infrastructure
-//
-// This file is distributed under the University of Illinois Open Source
-// License. See LICENSE.TXT for details.
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
----------------
Somewhat unfortunately for this change, this header was added with the old license. Correcting here. Introduced by @saiislam and written by @gregrodgers from my team at AMD.


================
Comment at: openmp/libomptarget/plugins/amdgpu/src/rtl.cpp:41
 
-// Get static gpu grid values from clang target-specific constants managed
-// in the header file llvm/Frontend/OpenMP/OMPGridValues.h
-// Copied verbatim to meet the requirement that libomptarget builds without
-// a copy of llvm checked out nearby
-namespace llvm {
-namespace omp {
-enum GVIDX {
-  /// The maximum number of workers in a kernel.
-  /// (THREAD_ABSOLUTE_LIMIT) - (GV_Warp_Size), might be issue for blockDim.z
-  GV_Threads,
-  /// The size reserved for data in a shared memory slot.
-  GV_Slot_Size,
-  /// The default value of maximum number of threads in a worker warp.
-  GV_Warp_Size,
-  /// Alternate warp size for some AMDGCN architectures. Same as GV_Warp_Size
-  /// for NVPTX.
-  GV_Warp_Size_32,
-  /// The number of bits required to represent the max number of threads in warp
-  GV_Warp_Size_Log2,
-  /// GV_Warp_Size * GV_Slot_Size,
-  GV_Warp_Slot_Size,
-  /// the maximum number of teams.
-  GV_Max_Teams,
-  /// Global Memory Alignment
-  GV_Mem_Align,
-  /// (~0u >> (GV_Warp_Size - GV_Warp_Size_Log2))
-  GV_Warp_Size_Log2_Mask,
-  // An alternative to the heavy data sharing infrastructure that uses global
-  // memory is one that uses device __shared__ memory.  The amount of such space
-  // (in bytes) reserved by the OpenMP runtime is noted here.
-  GV_SimpleBufferSize,
-  // The absolute maximum team size for a working group
-  GV_Max_WG_Size,
-  // The default maximum team size for a working group
-  GV_Default_WG_Size,
-  // This is GV_Max_WG_Size / GV_WarpSize. 32 for NVPTX and 16 for AMDGCN.
-  GV_Max_Warp_Number,
-  /// The slot size that should be reserved for a working warp.
-  /// (~0u >> (GV_Warp_Size - GV_Warp_Size_Log2))
-  GV_Warp_Size_Log2_MaskL
-};
-
-static constexpr unsigned AMDGPUGpuGridValues[] = {
-    448,       // GV_Threads
-    256,       // GV_Slot_Size
-    64,        // GV_Warp_Size
-    32,        // GV_Warp_Size_32
-    6,         // GV_Warp_Size_Log2
-    64 * 256,  // GV_Warp_Slot_Size
-    128,       // GV_Max_Teams
-    256,       // GV_Mem_Align
-    63,        // GV_Warp_Size_Log2_Mask
-    896,       // GV_SimpleBufferSize
-    1024,      // GV_Max_WG_Size,
-    256,       // GV_Defaut_WG_Size
-    1024 / 64, // GV_Max_WG_Size / GV_WarpSize
-    63         // GV_Warp_Size_Log2_MaskL
-};
-} // namespace omp
-} // namespace llvm
+#include "llvm/Frontend/OpenMP/OMPGridValues.h"
 
----------------
Lines on the left are byte for byte identical to the middle of this header, as that's where they were copied from. At some point in the future they're likely to accidentally diverge and I don't want to find that as the root cause while debugging.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D87841/new/

https://reviews.llvm.org/D87841