<p dir="ltr">Cool.</p>
<p dir="ltr">Lgtm with a llvm-mc testcase.</p>
<div class="gmail_quote">On Jun 20, 2016 12:32 PM, "Konstantin Zhuravlyov" <<a href="mailto:kzhuravl_dev@outlook.com">kzhuravl_dev@outlook.com</a>> wrote:<br type="attribution"><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">kzhuravl added a comment.<br>
<br>
In <a href="http://reviews.llvm.org/D21482#462100" rel="noreferrer" target="_blank">http://reviews.llvm.org/D21482#462100</a>, @rafael wrote:<br>
<br>
> > 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).<br>
><br>
> ><br>
><br>
> > We don't support preemption, but we need to use a got for referencing external variables.<br>
><br>
><br>
> Sorry, I am still missing something.<br>
><br>
> By "external", you mean "not in the .o" or "not in the .so". You don't<br>
>  need a got to access things that are in the .so if you don't support<br>
>  preemption. If they are external to the .so, where are they? Is it<br>
>  possible to structure a gpu program as multiple .so files?<br>
><br>
> Cheers,<br>
> Rafael<br>
<br>
<br>
Hi Rafael,<br>
<br>
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:<br>
<br>
GPU Kernel (OpenCL):<br>
<br>
  extern global int *ExternVar;<br>
<br>
  kernel void foo(global int *A, const int size) {<br>
    int gid = get_global_id(0);<br>
    if (gid < size)<br>
      A[gid] = ExternVar[gid];<br>
  }<br>
<br>
Host Code (using hsa runtime, pseudo code and hand waving):<br>
<br>
  int *hostA = NULL<br>
  hsa_memory_allocate(global_region, 4 * sizeof(int), &hostA)<br>
<br>
  int *hostExternVar = NULL<br>
  hsa_memory_allocate(global_region, 4 * sizeof(int), &hostExternVar)<br>
<br>
  for (i = 0; i < 4; i++) hostExternVar[i] = i<br>
<br>
  exec = hsa_executable_create<br>
  hsa_executable_global_variable_define(exec, "ExternVar", hostExternVar)<br>
  hsa_executable_load_code_object(exec, device, code object for GPU kernel above)<br>
<br>
  setup kernel arguments and dispatch kernel "foo"<br>
<br>
<br>
<a href="http://reviews.llvm.org/D21482" rel="noreferrer" target="_blank">http://reviews.llvm.org/D21482</a><br>
<br>
<br>
<br>
</blockquote></div>