Support/ELF: Add R_AMDGPU_GOTPCREL relocation

[• tstellarAMD]

Summary

We will start generating this in a future patch.

Change-Id: Id5fae1e8e6887ab7bbecbf445d1834a3aed4a5c9

[• rafael]

git-clang-format the patch.

Out of curiosity, why do you need a got? Are GPU fancy enough to have
shared libraries and symbol preemption these days?

Cheers,
Rafael

[• tstellarAMD]

clang-format the patch.

[• tstellarAMD]

In D21482#461432, @rafael wrote:

git-clang-format the patch.

Out of curiosity, why do you need a got? Are GPU fancy enough to have
shared libraries and symbol preemption these days?

We only support generating shared libraries of GPU code. The shared libraries are loaded into memory, and then the CPU host code is essentially calling functions in the GPU libraries (not directly calling them directly, but using a GPU runtime library, like OpenCL, for example, to invoke the functions).

We don't support preemption, but we need to use a got for referencing external variables.

[• rafael]

We only support generating shared libraries of GPU code. The shared libraries are loaded into memory, and then the CPU host code is essentially calling functions in the GPU libraries (not directly calling them directly, but using a GPU runtime library, like OpenCL, for example, to invoke the functions).

We don't support preemption, but we need to use a got for referencing external variables.

Sorry, I am still missing something.

By "external", you mean "not in the .o" or "not in the .so". You don't
need a got to access things that are in the .so if you don't support
preemption. If they are external to the .so, where are they? Is it
possible to structure a gpu program as multiple .so files?

Cheers,
Rafael

[• tstellarAMD]

In D21482#462100, @rafael wrote:

We only support generating shared libraries of GPU code. The shared libraries are loaded into memory, and then the CPU host code is essentially calling functions in the GPU libraries (not directly calling them directly, but using a GPU runtime library, like OpenCL, for example, to invoke the functions).

We don't support preemption, but we need to use a got for referencing external variables.

Sorry, I am still missing something.

By "external", you mean "not in the .o" or "not in the .so". You don't
need a got to access things that are in the .so if you don't support
preemption. If they are external to the .so, where are they? Is it
possible to structure a gpu program as multiple .so files?

"external" means not in the .so. There can be multiple .so files. A common case for this is on a multiple GPU system where you have one .so for data that is stored in device memory. This device .so is loaded into the device memory of each GPU (so one copy of the .so per device). Then you have a second .so for data that is stored in host memory. There is only a single copy of this .so which is stored in host memory that can be accessed by all devices. The .so stored in device memory will access variables defined in the .so stored in host memory.

One question I have is if we don't support pre-emption, is there some way for the whole tool chain to know this, so the linker is able to resolve more symbols at link time?

[kzhuravl]

In D21482#462100, @rafael wrote:

We only support generating shared libraries of GPU code. The shared libraries are loaded into memory, and then the CPU host code is essentially calling functions in the GPU libraries (not directly calling them directly, but using a GPU runtime library, like OpenCL, for example, to invoke the functions).

We don't support preemption, but we need to use a got for referencing external variables.

Sorry, I am still missing something.

By "external", you mean "not in the .o" or "not in the .so". You don't
need a got to access things that are in the .so if you don't support
preemption. If they are external to the .so, where are they? Is it
possible to structure a gpu program as multiple .so files?

Cheers,
Rafael

Hi Rafael,

In additions to what Tom said, external variables can be allocated and defined by the host, at runtime, before dispatching GPU kernels. Here is a "memory copy" quick example:

GPU Kernel (OpenCL):

extern global int *ExternVar;

kernel void foo(global int *A, const int size) {
  int gid = get_global_id(0);
  if (gid < size)
    A[gid] = ExternVar[gid];
}

Host Code (using hsa runtime, pseudo code and hand waving):

int *hostA = NULL
hsa_memory_allocate(global_region, 4 * sizeof(int), &hostA)

int *hostExternVar = NULL
hsa_memory_allocate(global_region, 4 * sizeof(int), &hostExternVar)

for (i = 0; i < 4; i++) hostExternVar[i] = i

exec = hsa_executable_create
hsa_executable_global_variable_define(exec, "ExternVar", hostExternVar)
hsa_executable_load_code_object(exec, device, code object for GPU kernel above)

setup kernel arguments and dispatch kernel "foo"

[• rafael]

Cool.

Lgtm with a llvm-mc testcase.

[kzhuravl]

LGTM: add missing `` in docs and add a test.

docs/CodeGenerator.rst

2700 ↗

(On Diff #61154)

Missing ``

[kzhuravl]

LGTM: add missing `` in docs and add a test.

llvm/trunk/docs/CodeGenerator.rst

 * **word64** --- This specifies a 64-bit field occupying 8 bytes with arbitrary
   byte alignment. These values use the same byte order as other word values in
   the AMD GPU architecture
 
 Following notations are used for specifying relocation calculations:
 
 * **A** --- Represents the addend used to compute the value of the relocatable
   field
+* **G** --- Represents the offset into the global offset table at which the
+  relocation entry’s symbol will reside during execution.
+* **GOT** --- Represents the address of the global offset table.
 * **P** --- Represents the place (section offset or address) of the storage unit
   being relocated (computed using ``r_offset``)
 * **S** --- Represents the value of the symbol whose index resides in the
   relocation entry
 
 AMDGPU Backend generates *Elf64_Rela* relocation records with the following
 supported relocation types:
 
   =====================  =====  ==========  ====================
   Relocation type        Value  Field       Calculation
   =====================  =====  ==========  ====================
   ``R_AMDGPU_NONE``      0      ``none``    ``none``
   ``R_AMDGPU_ABS32_LO``  1      ``word32``  (S + A) & 0xFFFFFFFF
   ``R_AMDGPU_ABS32_HI``  2      ``word32``  (S + A) >> 32
   ``R_AMDGPU_ABS64``     3      ``word64``  S + A
   ``R_AMDGPU_REL32``     4      ``word32``  S + A - P
   ``R_AMDGPU_REL64``     5      ``word64``  S + A - P
   ``R_AMDGPU_ABS32``     6      ``word32``  S + A
+  ``R_AMDGPU_GOTPCREL``  7      ``word32``  G + GOT + A - P
   =====================  =====  ==========  ====================

llvm/trunk/include/llvm/Support/ELFRelocs/AMDGPU.def

 #ifndef ELF_RELOC
 #error "ELF_RELOC must be defined"
 #endif
 
 ELF_RELOC(R_AMDGPU_NONE,          0)
 ELF_RELOC(R_AMDGPU_ABS32_LO,      1)
 ELF_RELOC(R_AMDGPU_ABS32_HI,      2)
 ELF_RELOC(R_AMDGPU_ABS64,         3)
 ELF_RELOC(R_AMDGPU_REL32,         4)
 ELF_RELOC(R_AMDGPU_REL64,         5)
 ELF_RELOC(R_AMDGPU_ABS32,         6)
+ELF_RELOC(R_AMDGPU_GOTPCREL,      7)

llvm/trunk/lib/Target/AMDGPU/MCTargetDesc/AMDGPUELFObjectWriter.cpp

unsigned AMDGPUELFObjectWriter::getRelocType(MCContext &Ctx,
                                              bool IsPCRel) const {
   // SCRATCH_RSRC_DWORD[01] is a special global variable that represents
   // the scratch buffer.
   if (Target.getSymA()->getSymbol().getName() == "SCRATCH_RSRC_DWORD0")
     return ELF::R_AMDGPU_ABS32_LO;
   if (Target.getSymA()->getSymbol().getName() == "SCRATCH_RSRC_DWORD1")
     return ELF::R_AMDGPU_ABS32_HI;
 
+  switch (Target.getAccessVariant()) {
+  default:
+    break;
+  case MCSymbolRefExpr::VK_GOTPCREL:
+    return ELF::R_AMDGPU_GOTPCREL;
+  }
+
   switch (Fixup.getKind()) {
   default: break;
   case FK_PCRel_4:
     return ELF::R_AMDGPU_REL32;
   }
 
   llvm_unreachable("unhandled relocation type");
 }

llvm/trunk/test/MC/AMDGPU/reloc.s

 // RUN: llvm-mc -filetype=obj -triple amdgcn-- -mcpu=kaveri -show-encoding %s | llvm-readobj -relocations | FileCheck %s
 
 // CHECK: Relocations [
 // CHECK: R_AMDGPU_ABS32_LO SCRATCH_RSRC_DWORD0 0x0
 // CHECK: R_AMDGPU_ABS32_HI SCRATCH_RSRC_DWORD1 0x0
+// CHECK: R_AMDGPU_GOTPCREL global_var 0x0
 // CHECK: ]
 
 kernel:
   s_mov_b32 s0, SCRATCH_RSRC_DWORD0
   s_mov_b32 s1, SCRATCH_RSRC_DWORD1
+  s_mov_b32 s2, global_var@GOTPCREL
+
+.globl global_var
 
 .globl SCRATCH_RSRC_DWORD0