This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

Re: [gomp4 0/8] NVPTX: initial OpenMP offloading


On Wed, 23 Sep 2015, Bernd Schmidt wrote:
> I have two major concerns here. Can I ask you how much experience you have
> with GPU programming and ptx?

I'd say I have a good understanding of the programming model and nvidia
hardware architecture, having used CUDA tools and payed attention to
r&d news for a few years.  I've discussed with OpenACC and HSA folks at the
GNU Cauldron my plans to take on this work, and I hope they can acknowledge
that I at least seemed to have a clue :)

> These patches provide stub functionality, which
> is easy enough, but I can't tell whether there's a credible plan to provide a
> full implementation. GPUs really need a different programming model than
> normal CPUs, which is something I learned the hard way, and I'm not terribly
> optimistic about porting libgomp to ptx. (I may be wrong.)

Right, libgomp running on ptx would have to do many things differently from
how it does now (and some drop entirely, like affinity).  Thankfully it can be
implemented piecemeal in config/nvptx, without #ifdef butchery in the primary
source files.  The plan towards providing a full implementation is thus to
work my way incrementally over GOMP_nn api, allowing '#pragma omp parallel' to
link successfully, then 'for', 'teams' and so on.  For 'parallel' the
intention is to either have prestarted idle threads in teams if possible, or
start another kernel via dynamic parallelism.  Exact details are to be worked
out -- I'd like to avoid introducing a hard dependency on dynamic parallelism.

> In one patch you mention newlib pthread type definitions - are you aware that
> there is no real pthreads implementation in the ptx newlib? The ptx newlib is
> really only provided for a minimal subset of libc functionality.

Sure, I'm aware.  The point was to make libgomp.h valid to be included into
the rest of to-be-ported source files, keeping modifications to it to a
minimum.  If the idea is that relying on #include <pthread.h> available on
nvptx in the first place is too much of a hack, we can discuss alternatives :)

> My other concern would be not to approve changes to the gomp-4_0-branch that
> could derail or slow down the effort to implement OpenACC, which has a much
> better chance of being in gcc-6 than this effort. You might want to make a
> private branch for your work.

I'm unclear how this work might hurt the OpenACC efforts, and in any case I
intend to be careful.  I don't imagine there will be conflicting requirements
to source code changes along the way.  In defense of the idea of working on
gomp4 branch, I expect that interleaving OpenACC and OpenMP work on a common
branch will cause less pain in case of inadvertent breakage than a merge
afterward.  Jakub, since you suggested submitting for gomp-4_0-branch, what's
your recommendation here?

Thanks
Alexander


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]