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

Missing ``

[kzhuravl]

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

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
   =====================  =====  ==========  ====================

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)

lib/Target/AMDGPU/MCTargetDesc/AMDGPUELFObjectWriter.cpp

: MCELFObjectTargetWriter(Is64Bit,
                             ELF::ELFOSABI_AMDGPU_HSA,
                             ELF::EM_AMDGPU,
                             HasRelocationAddend) { }
 
 unsigned AMDGPUELFObjectWriter::getRelocType(MCContext &Ctx,
                                              const MCValue &Target,
                                              const MCFixup &Fixup,
                                              bool IsPCRel) const {
+
+  switch (Target.getAccessVariant()) {
+  default: break;
+  case MCSymbolRefExpr::VK_GOTPCREL: return ELF::R_AMDGPU_GOTPCREL;
+  }
+
   switch (Fixup.getKind()) {
   default: break;
   case FK_Data_4:
     // 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")