[OpenMP][AMDGPU] Introduce new matadata fields for code object 5

[saiislam]

Summary

Added new metadata fields for code object version 5. Defined various
offset and size constants.

Here are some salient differences in cov5 and cov4:

• Size of Implicit KernArgs field is now 256 (used to be 56 in cov4).
• Different metadata fields have a well defined starting location and size (cov4 used to have 7 fields of 8 bytes each).
• Cov5 uses ImplictArg ptr instead of dispatch ptr to query workgroup_size_x so __builtin_amdgcn_workgroup_size_x() does not work.

[saiislam]

Ping.

[saiislam]

Not required after D139730

openmp/libomptarget/plugins/amdgpu/impl/get_elf_mach_gfx_name.h

 //===--- amdgpu/impl/get_elf_mach_gfx_name.h ---------------------- 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 GET_ELF_MACH_GFX_NAME_H_INCLUDED
 #define GET_ELF_MACH_GFX_NAME_H_INCLUDED
 
 #include <stdint.h>
 
 const char *get_elf_mach_gfx_name(uint32_t EFlags);
 
+/// refer for details: see
+/// https://llvm.org/docs/AMDGPUUsage.html#code-object-v5-metadata
+
+enum IMPLICITARGS : uint16_t {
+  COV4_SIZE = 56,                // size of code object v4 metadata in bytes
+  COV4_HOSTCALL_PTR_OFFSET = 24, // offset of host call buffer pointer
+  HOSTCALL_PTR_SIZE =
+      8, // size of hostcall buffer pointer, same for cov4 and cov5.
+
+  COV5_SIZE = 256, // size of code object v4 metadata in bytes
+
+  COV5_BLOCK_COUNT_X_OFFSET = 0,
+  COV5_BLOCK_COUNT_X_SIZE = 4,
+
+  COV5_BLOCK_COUNT_Y_OFFSET = 4,
+  COV5_BLOCK_COUNT_Y_SIZE = 4,
+
+  COV5_BLOCK_COUNT_Z_OFFSET = 8,
+  COV5_BLOCK_COUNT_Z_SIZE = 4,
+
+  COV5_GROUP_SIZE_X_OFFSET = 12,
+  COV5_GROUP_SIZE_X_SIZE = 2,
+
+  COV5_GROUP_SIZE_Y_OFFSET = 14,
+  COV5_GROUP_SIZE_Y_SIZE = 2,
+
+  COV5_GROUP_SIZE_Z_OFFSET = 16,
+  COV5_GROUP_SIZE_Z_SIZE = 2,
+
+  COV5_REMAINDER_X_OFFSET = 18,
+  COV5_REMAINDER_X_SIZE = 2,
+
+  COV5_REMAINDER_Y_OFFSET = 20,
+  COV5_REMAINDER_Y_SIZE = 2,
+
+  COV5_REMAINDER_Z_OFFSET = 22,
+  COV5_REMAINDER_Z_SIZE = 2,
+
+  COV5_GRID_DIMS_OFFSET = 64,
+  COV5_GRID_DIMS_SIZE = 2,
+
+  COV5_HOSTCALL_PTR_OFFSET = 80,
+
+  COV5_HEAPV1_PTR_OFFSET = 96,
+  COV5_HEAPV1_PTR_SIZE = 8
+};
+
 #endif

openmp/libomptarget/plugins/amdgpu/impl/get_elf_mach_gfx_name.cpp

 //===--- amdgpu/impl/get_elf_mach_gfx_name.cpp -------------------- 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 "get_elf_mach_gfx_name.h"
 
 // This header conflicts with the system elf.h (macros vs enums of the same
 // identifier) and contains more up to date values for the enum checked here.
 // rtl.cpp uses the system elf.h.
 #include "llvm/BinaryFormat/ELF.h"
+using namespace llvm::ELF;
 
 const char *get_elf_mach_gfx_name(uint32_t EFlags) {
   using namespace llvm::ELF;
   uint32_t Gfx = (EFlags & EF_AMDGPU_MACH);
   switch (Gfx) {
   case EF_AMDGPU_MACH_AMDGCN_GFX801:
     return "gfx801";
   case EF_AMDGPU_MACH_AMDGCN_GFX802:

openmp/libomptarget/plugins/amdgpu/impl/system.cpp

enum class ValueKind {
     HiddenGlobalOffsetZ,
     HiddenNone,
     HiddenPrintfBuffer,
     HiddenDefaultQueue,
     HiddenCompletionAction,
     HiddenMultiGridSyncArg,
     HiddenHostcallBuffer,
     HiddenHeapV1,
+    HiddenBlockCountX,
+    HiddenBlockCountY,
+    HiddenBlockCountZ,
+    HiddenGroupSizeX,
+    HiddenGroupSizeY,
+    HiddenGroupSizeZ,
+    HiddenRemainderX,
+    HiddenRemainderY,
+    HiddenRemainderZ,
+    HiddenGridDims,
+    HiddenQueuePtr,
     Unknown
   };
 
   KernelArgMD()
       : name_(std::string()), size_(0), offset_(0),
         valueKind_(ValueKind::Unknown) {}
 
   // fields
static const std::map<std::string, KernelArgMD::ValueKind> ArgValueKind = {
     {"hidden_none", KernelArgMD::ValueKind::HiddenNone},
     {"hidden_printf_buffer", KernelArgMD::ValueKind::HiddenPrintfBuffer},
     {"hidden_default_queue", KernelArgMD::ValueKind::HiddenDefaultQueue},
     {"hidden_completion_action",
      KernelArgMD::ValueKind::HiddenCompletionAction},
     {"hidden_multigrid_sync_arg",
      KernelArgMD::ValueKind::HiddenMultiGridSyncArg},
     {"hidden_hostcall_buffer", KernelArgMD::ValueKind::HiddenHostcallBuffer},
-    {"hidden_heap_v1", KernelArgMD::ValueKind::HiddenHeapV1}};
+    {"hidden_heap_v1", KernelArgMD::ValueKind::HiddenHeapV1},
+    {"hidden_block_count_x", KernelArgMD::ValueKind::HiddenBlockCountX},
+    {"hidden_block_count_y", KernelArgMD::ValueKind::HiddenBlockCountY},
+    {"hidden_block_count_z", KernelArgMD::ValueKind::HiddenBlockCountZ},
+    {"hidden_group_size_x", KernelArgMD::ValueKind::HiddenGroupSizeX},
+    {"hidden_group_size_y", KernelArgMD::ValueKind::HiddenGroupSizeY},
+    {"hidden_group_size_z", KernelArgMD::ValueKind::HiddenGroupSizeZ},
+    {"hidden_remainder_x", KernelArgMD::ValueKind::HiddenRemainderX},
+    {"hidden_remainder_y", KernelArgMD::ValueKind::HiddenRemainderY},
+    {"hidden_remainder_z", KernelArgMD::ValueKind::HiddenRemainderZ},
+    {"hidden_grid_dims", KernelArgMD::ValueKind::HiddenGridDims},
+    {"hidden_queue_ptr", KernelArgMD::ValueKind::HiddenQueuePtr},
+};
 
 namespace core {
 
 hsa_status_t callbackEvent(const hsa_amd_event_t *event, void *data) {
   if (event->event_type == HSA_AMD_GPU_MEMORY_FAULT_EVENT) {
     hsa_amd_gpu_memory_fault_info_t memory_fault = event->memory_fault;
     // memory_fault.agent
     // memory_fault.virtual_address
static bool isImplicit(KernelArgMD::ValueKind value_kind) {
   case KernelArgMD::ValueKind::HiddenGlobalOffsetZ:
   case KernelArgMD::ValueKind::HiddenNone:
   case KernelArgMD::ValueKind::HiddenPrintfBuffer:
   case KernelArgMD::ValueKind::HiddenDefaultQueue:
   case KernelArgMD::ValueKind::HiddenCompletionAction:
   case KernelArgMD::ValueKind::HiddenMultiGridSyncArg:
   case KernelArgMD::ValueKind::HiddenHostcallBuffer:
   case KernelArgMD::ValueKind::HiddenHeapV1:
+  case KernelArgMD::ValueKind::HiddenBlockCountX:
+  case KernelArgMD::ValueKind::HiddenBlockCountY:
+  case KernelArgMD::ValueKind::HiddenBlockCountZ:
+  case KernelArgMD::ValueKind::HiddenGroupSizeX:
+  case KernelArgMD::ValueKind::HiddenGroupSizeY:
+  case KernelArgMD::ValueKind::HiddenGroupSizeZ:
+  case KernelArgMD::ValueKind::HiddenRemainderX:
+  case KernelArgMD::ValueKind::HiddenRemainderY:
+  case KernelArgMD::ValueKind::HiddenRemainderZ:
+  case KernelArgMD::ValueKind::HiddenGridDims:
+  case KernelArgMD::ValueKind::HiddenQueuePtr:
     return true;
   default:
     return false;
   }
 }
 
 static std::pair<const unsigned char *, const unsigned char *>
 findMetadata(const ELFObjectFile<ELF64LE> &ELFObj) {

openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

if (!ArgPool) {
       AMDGPUImplicitArgsTy *ImplArgs = reinterpret_cast<AMDGPUImplicitArgsTy *>(
           static_cast<char *>(KernArg) + ArgPool->KernargSegmentSize);
       memset(ImplArgs, 0,
              sizeof(AMDGPUImplicitArgsTy)); // may not be necessary
       ImplArgs->OffsetX = 0;
       ImplArgs->OffsetY = 0;
       ImplArgs->OffsetZ = 0;
 
+      uint64_t Buffer = 0;
       // assign a hostcall buffer for the selected Q
       if (__atomic_load_n(&DeviceInfo().HostcallRequired, __ATOMIC_ACQUIRE)) {
         // hostrpc_assign_buffer is not thread safe, and this function is
         // under a multiple reader lock, not a writer lock.
         static pthread_mutex_t HostcallInitLock = PTHREAD_MUTEX_INITIALIZER;
         pthread_mutex_lock(&HostcallInitLock);
-        uint64_t Buffer = hostrpc_assign_buffer(
+        Buffer = hostrpc_assign_buffer(
             DeviceInfo().HSAAgents[DeviceId], Queue, DeviceId);
         pthread_mutex_unlock(&HostcallInitLock);
         if (!Buffer) {
           DP("hostrpc_assign_buffer failed, gpu would dereference null and "
              "error\n");
           return OFFLOAD_FAIL;
         }
+      }
 
-        DP("Implicit argument count: %d\n",
-           KernelInfoEntry.implicit_argument_count);
-        if (KernelInfoEntry.implicit_argument_count >= 4) {
-          // Initialise pointer for implicit_argument_count != 0 ABI
-          // Guess that the right implicit argument is at offset 24 after
-          // the explicit arguments. In the future, should be able to read
-          // the offset from msgpack. Clang is not annotating it at present.
-          uint64_t Offset =
-              sizeof(void *) * (KernelInfoEntry.explicit_argument_count + 3);
-          if ((Offset + 8) > ArgPool->kernargSizeIncludingImplicit()) {
-            DP("Bad offset of hostcall: %lu, exceeds kernarg size w/ implicit "
-               "args: %d\n",
-               Offset + 8, ArgPool->kernargSizeIncludingImplicit());
-          } else {
-            memcpy(static_cast<char *>(KernArg) + Offset, &Buffer, 8);
-          }
-        }
-
-        // initialise pointer for implicit_argument_count == 0 ABI
-        ImplArgs->HostcallPtr = Buffer;
-      }
+      DP("Setting Hostcall buffer for COV4\n");
+      memcpy(&ImplArgs[IMPLICITARGS::COV4_HOSTCALL_PTR_OFFSET], &Buffer,
+              IMPLICITARGS::HOSTCALL_PTR_SIZE);
 
       Packet->kernarg_address = KernArg;
     }
 
     hsa_signal_t S = DeviceInfo().FreeSignalPool.pop();
     if (S.handle == 0) {
       DP("Failed to get signal instance\n");
       return OFFLOAD_FAIL;