On Tue, Oct 20, 2015 at 11:19:06PM +0200, Bernd Schmidt wrote:
> On 10/20/2015 11:13 PM, Alexander Monakov wrote:
> >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.
>
> Could be that we're talking about the same thing. I think I was confused by
> a reference to .func vs .kernel and sm_30 vs sm_35 in patch 2/14. So let's
> go for this approach.
But you'd better then rename the .kernel stub to the name that libgomp looks
up and rename the actual outlined body to some other name (add some suffix).
Or maybe another possibility is to use the normal outlined target body
function, and just inject some gomp_nvptx_start and gomp_nvptx_end functions
to the beginning and end of "omp target entrypoint" function.
Jakub
