On Tue, 20 Oct 2015, Bernd Schmidt wrote:
> On 10/20/2015 08:34 PM, Alexander Monakov wrote:
> > 2. Make gomp_nvptx_main a device (.func) function. To have that work, we'd
> > need to additionally emit a "trampoline" of sorts in the NVPTX backend. For
> > each OpenMP target entrypoint foo$_omp_fn$0, we'd have to additionally emit
> >
> > __global__ void foo$_omp_fn$0$entry(void *args)
> > {
> > gomp_nvptx_main(foo$_omp_fn$0, args);
> > }
>
> Wouldn't it be simpler to generate a .kernel for every target region function
> (as OpenACC does)? That could be a small stub in each case which just calls
> gomp_nvptx_main with the right function pointer. We already have the machinery
> to look up the right kernel corresponding to a host address and invoke it, so
> I think we should just reuse that functionality.
As I see we are describing the same thing in different words.
In what you describe, and in my quoted paragraph, both gomp_nvptx_main and the
function originally outlined for a target region are device-only (.func)
functions. The .kernel function that the plugin looks up and launches is a
small piece of code that calls gomp_nvptx_main, passing it a pointer to the
target region function.
Unless I didn't fully catch what you say? Like I said in the email, I do like
this approach more.
Thanks.
Alexander
