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

[Rafael Espíndola via llvm-commits]

Cool.

Lgtm with a llvm-mc testcase.
On Jun 20, 2016 12:32 PM, "Konstantin Zhuravlyov" <kzhuravl_dev at outlook.com>
wrote:

> 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
>
>
>
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20160620/84efb871/attachment.html>