[PATCH] D21482: Support/ELF: Add R_AMDGPU_GOTPCREL relocation

[Konstantin Zhuravlyov via llvm-commits]

kzhuravl added a comment.

In http://reviews.llvm.org/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,

By "external" variables we mean variables that are 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"


http://reviews.llvm.org/D21482