Re: [RFC PATCH] Coalesce host to device transfers in libgomp

[Jakub Jelinek <jakub at redhat dot com>]

On Tue, Oct 24, 2017 at 08:47:39AM -0700, Cesar Philippidis wrote:
> On 10/24/2017 02:55 AM, Jakub Jelinek wrote:
> 
> > Poeple from NVidia reported privately unexpected amount of host2dev
> > transfers for #pragma omp target*.
> 
> Did they mention which program they were testing?

No.  Just the nvprof counted counts from GCC and LLVM.

> > The following patch implements coalescing of transfers (only those that are
> > copied into the freshly allocated device buffer) into one or multiple larger
> > transfers.  The patch doesn't coalesce > 32KB transfers or transfers where
> > the gap is 4KB or more.  I guess it would be not too hard to do similar
> > coalescing for the dev2host transfers that are from a single device mapping,
> > though probably far less important than the more common host2dev transfers.
> 
> Why did you chose the 32KB and 4KB limits? I wonder if that would have
> any impact on firstprivate_int values. If this proves to be effective,
> it seems like we should be able to eliminate GOMP_MAP_FIRSTPRIVATE_INT
> altogether.

The thing is that this is a generic code, so it is hard to come up with
reasonable limits.  We could even have some limits e.g. in *devicep
if we get different needs for different offloading targets.

The 32KB and 4KB just come from some discussions with Alexander on IRC
that larger copies saturate the PCI and the overhead isn't significant, so
in that case copying e.g. megabyte into another memory and then to the
device would likely not be beneficial.

I'd prefer to keep GOMP_MAP_FIRSTPRIVATE_INT, I think it is a useful
optimization for the most common case, even if it is not 2 separate host2dev
copies for it compared to 1 for GOMP_MAP_FIRSTPRIVATE_INT, it is still extra
memory dereferences both on the host and on the target.

> > +struct gomp_map_cache
> > +{
> > +  void *buf;
> > +  struct target_mem_desc *tgt;
> > +  size_t *chunks;
> > +  long chunk_cnt;
> > +  long use_cnt;
> > +};
> > +
> 
> Maybe include a comment here stating that you want to restrict caching
> to 32KB with variables with no gaps larger than 4KB?

Sure.  Maybe even better to turn those for now into defines and add comments
to those.

> One other minor optimization, would be to change arguments to offloaded
> functions from a single struct to individual arguments. At least for
> nvptx, cuLaunchKernel accepts variable arguments for PTX kernels. There
> are two advantages of this. 1) At least with nvptx, nvptx_exec wouldn't
> need to reserve a block of device memory for struct argument. 2) This
> would eliminate one level of indirection for each offloaded argument
> (although SRA probably takes care of the majority of this already).

At least for OpenMP, we are now using a wrapper around the generated code
which sets stuff up, so not sure if that would be possible.  The wrapper
among other things sets up the soft-stack.  Not sure if it wouldn't be
possible to replace it with a magic call at the begining of OpenMP kernel
starts.

	Jakub