March 2021 - mpiwg-hybridpm - mailman3.cels.anl.gov

Re: [mpiwg-hybridpm] Call for Topics
by Jim Dinan 23 Mar '21

23 Mar '21

Hi Joseph, No other topics. Meeting is all yours. ~Jim. On Tue, Mar 23, 2021 at 4:41 AM Joseph Schuchart via mpiwg-hybridpm < mpiwg-hybridpm(a)lists.mpi-forum.org> wrote: > All, > > If there are no other topics for tomorrow, I'd like to use the slot to > further discuss the callback-based completion notifications. In > particular, I'm interested in how the WG feels about the different > approaches for driving progress of outstanding callbacks, as that will > impact the design of the interface going forward. > > Cheers > Joseph > > On 3/19/21 7:38 PM, Jim Dinan via mpiwg-hybridpm wrote: > > Hi All, > > > > We will be meeting next Wednesday, March 24. Please let me know if you > > are interested in presenting or have topics you would like to discuss. > > > > Cheers, > > ~Jim. > > > > _______________________________________________ > > mpiwg-hybridpm mailing list > > mpiwg-hybridpm(a)lists.mpi-forum.org > > https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm > > > _______________________________________________ > mpiwg-hybridpm mailing list > mpiwg-hybridpm(a)lists.mpi-forum.org > https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm >

1 0

Re: [mpiwg-hybridpm] Call for Topics
by Joseph Schuchart 23 Mar '21

23 Mar '21

All, If there are no other topics for tomorrow, I'd like to use the slot to further discuss the callback-based completion notifications. In particular, I'm interested in how the WG feels about the different approaches for driving progress of outstanding callbacks, as that will impact the design of the interface going forward. Cheers Joseph On 3/19/21 7:38 PM, Jim Dinan via mpiwg-hybridpm wrote: > Hi All, > > We will be meeting next Wednesday, March 24. Please let me know if you > are interested in presenting or have topics you would like to discuss. > > Cheers, > ~Jim. > > _______________________________________________ > mpiwg-hybridpm mailing list > mpiwg-hybridpm(a)lists.mpi-forum.org > https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm >

1 0

Re: [mpiwg-hybridpm] Hybrid/Accelerator WG Meeting
by Dan Holmes 19 Mar '21

19 Mar '21

Hi Jim, This is a helpful clarification that helps to define the goals more precisely. Thanks. Cheers, Dan. — Dr Daniel Holmes PhD Executive Director Chief Technology Officer CHI Ltd danholmes(a)chi.scot > On 19 Mar 2021, at 18:56, Jim Dinan <james.dinan(a)gmail.com> wrote: > > Hi Dan, > > CUDA streams are a virtual resource and they are mapped to a fixed number of hardware channels allocated to the CUDA context. Hardware channels and streams both have FIFO execution semantics. The cuStreamWaitValue64 operation is executed by the GPU control processor (this occurs outside of the SM, thus not involving a thread block) and has the effect of blocking the hardware channel. So, operations (e.g. mem ops, copies, kernel launches, etc) that happen to land in the same hardware channel behind the cuStreamWaitValue64 are also blocked. > > For example, a user could do something like MPI_Send_on_stream(..., stream0) and MPI_Recv_on_stream(..., stream1). CUDA can map stream0 and stream1 to the same hardware channel with any interleaving of operations from the streams (as long as FIFO ordering is preserved for each stream individually). Assuming that the cuStreamWaitValue64 for the receive operation is inserted into the hardware channel first, it will block the cuStreamWriteValue64 operation corresponding to the send, leading to deadlock. > > A workaround that could be used today is for users to insert CUDA events to force the send to be performed prior to the recv. However, this will unnecessarily serialize the streams in (likely) cases when they map to separate hardware channels. The situation is obviously not ideal, and we are looking at ways to fix this. > > Cheers, > ~Jim. > > On Thu, Mar 11, 2021 at 10:57 AM Dan Holmes <danholmes(a)chi.scot> wrote: > Hi Jim, > > The CPU callback operation you describe seems to be only one-way (GPU notifying/triggering CPU), but a reverse mechanism would be needed to complete the pattern, as discussed on the call? > > If the cuStreamWaitValue64 operation works like the “wait" in Stephen’s slides, i.e. does not actually sit and wait (blocking a whole thread block), but makes the FIFO stream not runnable until the wait condition is satisfied, then that looks promising for a return path. Your hint "memops are processed from within the GPU control processor that manages stream execution” suggests this is true. Doesn’t that provide all that is needed to get this working? > The GPU has a way to signal to the CPU that the send data buffer is ready to be sent. > The CPU has a way to signal to the GPU that the send data buffer can be overwritten. > The CPU has a way to signal to the GPU that the recv data buffer is ready to be consumed. > The GPU has a way to signal to the CPU that the recv data buffer can be overwritten. > The CPU needs a helper thread to monitor the memory locations targeted by the memops. This could use general requests in MPI and rely on the MPI progress “thread”/mechanism or be a separate/dedicated CPU thread. > The GPU can pack/unpack data; the CPU can use GPUDirect. > > Nothing further seems to be necessary for a complete/functional implementation. Future changes — to MPI and/or to GPUs/CUDA — are only needed/desired to reduce/eliminate performance bottlenecks in this pattern. > > Am I getting that right? > > Cheers, > Dan. > — > Dr Daniel Holmes PhD > Executive Director > Chief Technology Officer > CHI Ltd > danholmes(a)chi.scot <mailto:[email protected]> > > > >> On 11 Mar 2021, at 15:07, Jim Dinan via mpiwg-hybridpm <mpiwg-hybridpm(a)lists.mpi-forum.org <mailto:[email protected]>> wrote: >> >> Unfortunately, CPU callbacks are not a perfect solution on their own. CUDA does not allow CUDA calls from within CPU callbacks, so for example you would not be able to launch data packing kernels or peer-to-peer copy operations from within the callback. However, you can use CPU callbacks to signal a thread in the MPI runtime to process the operation. Another option in this design space is to use CUDA memops (e.g. cuStreamWriteValue64 or cuStreamWaitValue64) to coordinate between CUDA streams and MPI communication helper threads. Because memops are processed from within the GPU control processor that manages stream execution, I would expect these to have lower overheads than CPU callbacks (although I haven't measured this). >> >> ~Jim. >> >> On Wed, Mar 10, 2021 at 10:08 PM Junchao Zhang <junchao.zhang(a)gmail.com <mailto:[email protected]>> wrote: >> Jim, >> Thanks for the slides. In Stephen's presentation today, it seems with existing techniques, i.e, CPU MPI callback nodes in CUDA graphs, one can solve the MPI GPU problem. Is my understanding correct? >> >> Thanks. >> --Junchao Zhang >> >> >> On Wed, Mar 10, 2021 at 8:34 PM Jim Dinan via mpiwg-hybridpm <mpiwg-hybridpm(a)lists.mpi-forum.org <mailto:[email protected]>> wrote: >> Hi All, >> >> I've posted Stephen's slides: https://github.com/mpiwg-hybrid/hybrid-issues/tree/master/slides <https://github.com/mpiwg-hybrid/hybrid-issues/tree/master/slides> >> >> Best, >> ~Jim. >> >> On Mon, Mar 8, 2021 at 11:21 AM Jim Dinan <james.dinan(a)gmail.com <mailto:[email protected]>> wrote: >> Hi All, >> >> We have an invited speaker this week at the HACC WG: >> >> Topic: CUDA Deep Dive For the MPI Forum HACC WG >> When: Wednesday, March 10 10-11:00am ET >> Connection Info: https://github.com/mpiwg-hybrid/hybrid-issues/wiki <https://github.com/mpiwg-hybrid/hybrid-issues/wiki> >> >> Speaker: Stephen Jones, NVIDIA >> >> Stephen Jones is one of the architects of CUDA, working on defining the language, the platform, and the hardware that it runs on, to span the needs of parallel programming from high performance computing to artificial intelligence. Prior to his present position, he lead the Simulation & Analytics group at SpaceX, working on large-scale simulation of rocket engines. He has worked in diverse other industries, including networking, CAD/CAM, and scientific computing. He has been a part of CUDA since 2008. >> >> Cheers, >> ~Jim. >> >> PS - Apologies for cross posting on the main list. If you would like to continue receiving emails relating to the Hybrid & Accelerator WG, please sign up for the mailing list here: https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm <https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm>. >> _______________________________________________ >> mpiwg-hybridpm mailing list >> mpiwg-hybridpm(a)lists.mpi-forum.org <mailto:[email protected]> >> https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm <https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm> >> _______________________________________________ >> mpiwg-hybridpm mailing list >> mpiwg-hybridpm(a)lists.mpi-forum.org <mailto:[email protected]> >> https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm <https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm> >

1 0

Re: [mpiwg-hybridpm] Hybrid/Accelerator WG Meeting
by Jim Dinan 19 Mar '21

19 Mar '21

Hi Junchao, If I understand correctly, the right solution here might be to have a submodule in the code return a CUDA graph that can then be embedded into the larger CUDA graph that you're building via cudaGraphAddChildGraphNode [1]. Have you looked at this option? Best, ~Jim. [1] https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__GRAPH.html#gro… On Thu, Mar 11, 2021 at 10:53 AM Junchao Zhang <junchao.zhang(a)gmail.com> wrote: > In PETSc, we don't use non-contiguous MPI data types, so packing is not > a problem for us. It seems if MPI can send/recv data without participation > of GPU, then it is a workable solution in limited cases. My main concern > about CUDA-graph is it is not modularizable. > 1) Wtih cudaGraphAddKernelNode() etc, one has to see the whole graph at > all. If kernels are spread and deeply wrapped in CPU routines as in PETSc, > at a higher level, one doesn't know the kernels and hence nodes in the > graph. > 2) With graph capture, again, a caller subroutine can not guarantee > a callee subroutine will execute the same path in next iteration as the one > when it was captured. > > --Junchao Zhang > > > On Thu, Mar 11, 2021 at 9:07 AM Jim Dinan <james.dinan(a)gmail.com> wrote: > >> Unfortunately, CPU callbacks are not a perfect solution on their own. >> CUDA does not allow CUDA calls from within CPU callbacks, so for example >> you would not be able to launch data packing kernels or peer-to-peer copy >> operations from within the callback. However, you can use CPU callbacks to >> signal a thread in the MPI runtime to process the operation. Another option >> in this design space is to use CUDA memops (e.g. cuStreamWriteValue64 or >> cuStreamWaitValue64) to coordinate between CUDA streams and MPI >> communication helper threads. Because memops are processed from within the >> GPU control processor that manages stream execution, I would expect these >> to have lower overheads than CPU callbacks (although I haven't measured >> this). >> >> ~Jim. >> >> On Wed, Mar 10, 2021 at 10:08 PM Junchao Zhang <junchao.zhang(a)gmail.com> >> wrote: >> >>> Jim, >>> Thanks for the slides. In Stephen's presentation today, it seems >>> with existing techniques, i.e, CPU MPI callback nodes in CUDA graphs, one >>> can solve the MPI GPU problem. Is my understanding correct? >>> >>> Thanks. >>> --Junchao Zhang >>> >>> >>> On Wed, Mar 10, 2021 at 8:34 PM Jim Dinan via mpiwg-hybridpm < >>> mpiwg-hybridpm(a)lists.mpi-forum.org> wrote: >>> >>>> Hi All, >>>> >>>> I've posted Stephen's slides: >>>> https://github.com/mpiwg-hybrid/hybrid-issues/tree/master/slides >>>> >>>> Best, >>>> ~Jim. >>>> >>>> On Mon, Mar 8, 2021 at 11:21 AM Jim Dinan <james.dinan(a)gmail.com> >>>> wrote: >>>> >>>>> Hi All, >>>>> >>>>> We have an invited speaker this week at the HACC WG: >>>>> >>>>> Topic: CUDA Deep Dive For the MPI Forum HACC WG >>>>> When: Wednesday, March 10 10-11:00am ET >>>>> Connection Info: https://github.com/mpiwg-hybrid/hybrid-issues/wiki >>>>> >>>>> Speaker: Stephen Jones, NVIDIA >>>>> >>>>> Stephen Jones is one of the architects of CUDA, working on defining >>>>> the language, the platform, and the hardware that it runs on, to span the >>>>> needs of parallel programming from high performance computing to artificial >>>>> intelligence. Prior to his present position, he lead the Simulation & >>>>> Analytics group at SpaceX, working on large-scale simulation of rocket >>>>> engines. He has worked in diverse other industries, including networking, >>>>> CAD/CAM, and scientific computing. He has been a part of CUDA since 2008. >>>>> >>>>> Cheers, >>>>> ~Jim. >>>>> >>>>> PS - Apologies for cross posting on the main list. If you would like >>>>> to continue receiving emails relating to the Hybrid & Accelerator WG, >>>>> please sign up for the mailing list here: >>>>> https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm. >>>>> >>>> _______________________________________________ >>>> mpiwg-hybridpm mailing list >>>> mpiwg-hybridpm(a)lists.mpi-forum.org >>>> https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm >>>> >>>

1 0

Re: [mpiwg-hybridpm] Hybrid/Accelerator WG Meeting
by Jim Dinan 19 Mar '21

19 Mar '21

Hi Dan, CUDA streams are a virtual resource and they are mapped to a fixed number of hardware channels allocated to the CUDA context. Hardware channels and streams both have FIFO execution semantics. The cuStreamWaitValue64 operation is executed by the GPU control processor (this occurs outside of the SM, thus not involving a thread block) and has the effect of blocking the hardware channel. So, operations (e.g. mem ops, copies, kernel launches, etc) that happen to land in the same hardware channel behind the cuStreamWaitValue64 are also blocked. For example, a user could do something like MPI_Send_on_stream(..., stream0) and MPI_Recv_on_stream(..., stream1). CUDA can map stream0 and stream1 to the same hardware channel with any interleaving of operations from the streams (as long as FIFO ordering is preserved for each stream individually). Assuming that the cuStreamWaitValue64 for the receive operation is inserted into the hardware channel first, it will block the cuStreamWriteValue64 operation corresponding to the send, leading to deadlock. A workaround that could be used today is for users to insert CUDA events to force the send to be performed prior to the recv. However, this will unnecessarily serialize the streams in (likely) cases when they map to separate hardware channels. The situation is obviously not ideal, and we are looking at ways to fix this. Cheers, ~Jim. On Thu, Mar 11, 2021 at 10:57 AM Dan Holmes <danholmes(a)chi.scot> wrote: > Hi Jim, > > The CPU callback operation you describe seems to be only one-way (GPU > notifying/triggering CPU), but a reverse mechanism would be needed to > complete the pattern, as discussed on the call? > > If the cuStreamWaitValue64 operation works like the “wait" in Stephen’s > slides, i.e. does not actually sit and wait (blocking a whole thread > block), but makes the FIFO stream not runnable until the wait condition is > satisfied, then that looks promising for a return path. Your hint "memops > are processed from within the GPU control processor that manages stream > execution” suggests this is true. Doesn’t that provide all that is needed > to get this working? > The GPU has a way to signal to the CPU that the send data buffer is ready > to be sent. > The CPU has a way to signal to the GPU that the send data buffer can be > overwritten. > The CPU has a way to signal to the GPU that the recv data buffer is ready > to be consumed. > The GPU has a way to signal to the CPU that the recv data buffer can be > overwritten. > The CPU needs a helper thread to monitor the memory locations targeted by > the memops. This could use general requests in MPI and rely on the MPI > progress “thread”/mechanism or be a separate/dedicated CPU thread. > The GPU can pack/unpack data; the CPU can use GPUDirect. > > Nothing further seems to be necessary for a complete/functional > implementation. Future changes — to MPI and/or to GPUs/CUDA — are only > needed/desired to reduce/eliminate performance bottlenecks in this pattern. > > Am I getting that right? > > Cheers, > Dan. > — > Dr Daniel Holmes PhD > Executive Director > Chief Technology Officer > CHI Ltd > danholmes(a)chi.scot > > > > On 11 Mar 2021, at 15:07, Jim Dinan via mpiwg-hybridpm < > mpiwg-hybridpm(a)lists.mpi-forum.org> wrote: > > Unfortunately, CPU callbacks are not a perfect solution on their own. CUDA > does not allow CUDA calls from within CPU callbacks, so for example you > would not be able to launch data packing kernels or peer-to-peer copy > operations from within the callback. However, you can use CPU callbacks to > signal a thread in the MPI runtime to process the operation. Another option > in this design space is to use CUDA memops (e.g. cuStreamWriteValue64 or > cuStreamWaitValue64) to coordinate between CUDA streams and MPI > communication helper threads. Because memops are processed from within the > GPU control processor that manages stream execution, I would expect these > to have lower overheads than CPU callbacks (although I haven't measured > this). > > ~Jim. > > On Wed, Mar 10, 2021 at 10:08 PM Junchao Zhang <junchao.zhang(a)gmail.com> > wrote: > >> Jim, >> Thanks for the slides. In Stephen's presentation today, it seems >> with existing techniques, i.e, CPU MPI callback nodes in CUDA graphs, one >> can solve the MPI GPU problem. Is my understanding correct? >> >> Thanks. >> --Junchao Zhang >> >> >> On Wed, Mar 10, 2021 at 8:34 PM Jim Dinan via mpiwg-hybridpm < >> mpiwg-hybridpm(a)lists.mpi-forum.org> wrote: >> >>> Hi All, >>> >>> I've posted Stephen's slides: >>> https://github.com/mpiwg-hybrid/hybrid-issues/tree/master/slides >>> >>> Best, >>> ~Jim. >>> >>> On Mon, Mar 8, 2021 at 11:21 AM Jim Dinan <james.dinan(a)gmail.com> wrote: >>> >>>> Hi All, >>>> >>>> We have an invited speaker this week at the HACC WG: >>>> >>>> Topic: CUDA Deep Dive For the MPI Forum HACC WG >>>> When: Wednesday, March 10 10-11:00am ET >>>> Connection Info: https://github.com/mpiwg-hybrid/hybrid-issues/wiki >>>> >>>> Speaker: Stephen Jones, NVIDIA >>>> >>>> Stephen Jones is one of the architects of CUDA, working on defining the >>>> language, the platform, and the hardware that it runs on, to span the needs >>>> of parallel programming from high performance computing to artificial >>>> intelligence. Prior to his present position, he lead the Simulation & >>>> Analytics group at SpaceX, working on large-scale simulation of rocket >>>> engines. He has worked in diverse other industries, including networking, >>>> CAD/CAM, and scientific computing. He has been a part of CUDA since 2008. >>>> >>>> Cheers, >>>> ~Jim. >>>> >>>> PS - Apologies for cross posting on the main list. If you would like to >>>> continue receiving emails relating to the Hybrid & Accelerator WG, please >>>> sign up for the mailing list here: >>>> https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm. >>>> >>> _______________________________________________ >>> mpiwg-hybridpm mailing list >>> mpiwg-hybridpm(a)lists.mpi-forum.org >>> https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm >>> >> _______________________________________________ > mpiwg-hybridpm mailing list > mpiwg-hybridpm(a)lists.mpi-forum.org > https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm > > >

1 0

Call for Topics
by Jim Dinan 19 Mar '21

19 Mar '21

Hi All, We will be meeting next Wednesday, March 24. Please let me know if you are interested in presenting or have topics you would like to discuss. Cheers, ~Jim.

1 0

FW: MPI Persistent/Collective/Partitioned/Accelerator WGs Weekly Meeting
by Skjellum, Anthony 16 Mar '21

16 Mar '21

Hi, the partitioned communication discussion will be 1030-11am tomorrow; Ryan Grant suggested there is open discussion 10-1030? from last week's speaker. Jim does not have an agenda planned for hybrid tomorrow. So, I suggest that open conversations happen first half hour, per Ryan; second half I will discuss some of the partitioned collective concepts to be proposed for MPI-5 (subject to feedback of course). ________________________________ From: Skjellum, Anthony Sent: Wednesday, October 11, 2017 7:04:30 AM (UTC-05:00) Eastern Time (US & Canada) To: mpiwg-persistence(a)lists.mpi-forum.org; mpiwg-coll(a)lists.mpi-forum.org Subject: MPI Persistent/Collective/Partitioned/Accelerator WGs Weekly Meeting When: Wednesday, March 17, 2021 10:00 AM-11:00 AM. Where: [Updated participant list only] Hi there, Anthony Skjellum is inviting you to a scheduled Zoom meeting. Join from PC, Mac, Linux, iOS or Android: https://tennessee.zoom.us/j/91925030900?pwd=UjlXek5pUnpmL3RuWG5xemNuSHdsUT09 Password: 136803 ---- Topic: WG Meeting for Collective, Partitioned, Persistent, and friends Time: Aug 19, 2020 10:00 AM Eastern Time (US and Canada) Every week on Wed, until Jul 28, 2021, 50 occurrence(s) Aug 19, 2020 10:00 AM Aug 26, 2020 10:00 AM Sep 2, 2020 10:00 AM Sep 9, 2020 10:00 AM Sep 16, 2020 10:00 AM Sep 23, 2020 10:00 AM Sep 30, 2020 10:00 AM Oct 7, 2020 10:00 AM Oct 14, 2020 10:00 AM Oct 21, 2020 10:00 AM Oct 28, 2020 10:00 AM Nov 4, 2020 10:00 AM Nov 11, 2020 10:00 AM Nov 18, 2020 10:00 AM Nov 25, 2020 10:00 AM Dec 2, 2020 10:00 AM Dec 9, 2020 10:00 AM Dec 16, 2020 10:00 AM Dec 23, 2020 10:00 AM Dec 30, 2020 10:00 AM Jan 6, 2021 10:00 AM Jan 13, 2021 10:00 AM Jan 20, 2021 10:00 AM Jan 27, 2021 10:00 AM Feb 3, 2021 10:00 AM Feb 10, 2021 10:00 AM Feb 17, 2021 10:00 AM Feb 24, 2021 10:00 AM Mar 3, 2021 10:00 AM Mar 10, 2021 10:00 AM Mar 17, 2021 10:00 AM Mar 24, 2021 10:00 AM Mar 31, 2021 10:00 AM Apr 7, 2021 10:00 AM Apr 14, 2021 10:00 AM Apr 21, 2021 10:00 AM Apr 28, 2021 10:00 AM May 5, 2021 10:00 AM May 12, 2021 10:00 AM May 19, 2021 10:00 AM May 26, 2021 10:00 AM Jun 2, 2021 10:00 AM Jun 9, 2021 10:00 AM Jun 16, 2021 10:00 AM Jun 23, 2021 10:00 AM Jun 30, 2021 10:00 AM Jul 7, 2021 10:00 AM Jul 14, 2021 10:00 AM Jul 21, 2021 10:00 AM Jul 28, 2021 10:00 AM Please download and import the following iCalendar (.ics) files to your calendar system. Weekly: https://tennessee.zoom.us/meeting/tJUkdeyorTsjGtR1LZNj06Oa5HHLSf6OYx0E/ics?… Join from PC, Mac, Linux, iOS or Android: https://tennessee.zoom.us/j/91925030900?pwd=UjlXek5pUnpmL3RuWG5xemNuSHdsUT09 Password: 136803 Or iPhone one-tap (US Toll): +13126266799,91925030900# or +16468769923,91925030900# Or Telephone: Dial: +1 312 626 6799 (US Toll) +1 646 876 9923 (US Toll) +1 301 715 8592 (US Toll) +1 346 248 7799 (US Toll) +1 669 900 6833 (US Toll) +1 253 215 8782 (US Toll) Meeting ID: 919 2503 0900 International numbers available: https://tennessee.zoom.us/u/as41IZiIe Or an H.323/SIP room system: H.323: 162.255.37.11 (US West) or 162.255.36.11 (US East) Meeting ID: 919 2503 0900 Password: 136803 SIP: 91925030900(a)zoomcrc.com Password: 136803

1 0

Re: [mpiwg-hybridpm] Hybrid/Accelerator WG Meeting
by Dan Holmes 11 Mar '21

11 Mar '21

Hi Jim, The CPU callback operation you describe seems to be only one-way (GPU notifying/triggering CPU), but a reverse mechanism would be needed to complete the pattern, as discussed on the call? If the cuStreamWaitValue64 operation works like the “wait" in Stephen’s slides, i.e. does not actually sit and wait (blocking a whole thread block), but makes the FIFO stream not runnable until the wait condition is satisfied, then that looks promising for a return path. Your hint "memops are processed from within the GPU control processor that manages stream execution” suggests this is true. Doesn’t that provide all that is needed to get this working? The GPU has a way to signal to the CPU that the send data buffer is ready to be sent. The CPU has a way to signal to the GPU that the send data buffer can be overwritten. The CPU has a way to signal to the GPU that the recv data buffer is ready to be consumed. The GPU has a way to signal to the CPU that the recv data buffer can be overwritten. The CPU needs a helper thread to monitor the memory locations targeted by the memops. This could use general requests in MPI and rely on the MPI progress “thread”/mechanism or be a separate/dedicated CPU thread. The GPU can pack/unpack data; the CPU can use GPUDirect. Nothing further seems to be necessary for a complete/functional implementation. Future changes — to MPI and/or to GPUs/CUDA — are only needed/desired to reduce/eliminate performance bottlenecks in this pattern. Am I getting that right? Cheers, Dan. — Dr Daniel Holmes PhD Executive Director Chief Technology Officer CHI Ltd danholmes(a)chi.scot > On 11 Mar 2021, at 15:07, Jim Dinan via mpiwg-hybridpm <mpiwg-hybridpm(a)lists.mpi-forum.org> wrote: > > Unfortunately, CPU callbacks are not a perfect solution on their own. CUDA does not allow CUDA calls from within CPU callbacks, so for example you would not be able to launch data packing kernels or peer-to-peer copy operations from within the callback. However, you can use CPU callbacks to signal a thread in the MPI runtime to process the operation. Another option in this design space is to use CUDA memops (e.g. cuStreamWriteValue64 or cuStreamWaitValue64) to coordinate between CUDA streams and MPI communication helper threads. Because memops are processed from within the GPU control processor that manages stream execution, I would expect these to have lower overheads than CPU callbacks (although I haven't measured this). > > ~Jim. > > On Wed, Mar 10, 2021 at 10:08 PM Junchao Zhang <junchao.zhang(a)gmail.com <mailto:[email protected]>> wrote: > Jim, > Thanks for the slides. In Stephen's presentation today, it seems with existing techniques, i.e, CPU MPI callback nodes in CUDA graphs, one can solve the MPI GPU problem. Is my understanding correct? > > Thanks. > --Junchao Zhang > > > On Wed, Mar 10, 2021 at 8:34 PM Jim Dinan via mpiwg-hybridpm <mpiwg-hybridpm(a)lists.mpi-forum.org <mailto:[email protected]>> wrote: > Hi All, > > I've posted Stephen's slides: https://github.com/mpiwg-hybrid/hybrid-issues/tree/master/slides <https://github.com/mpiwg-hybrid/hybrid-issues/tree/master/slides> > > Best, > ~Jim. > > On Mon, Mar 8, 2021 at 11:21 AM Jim Dinan <james.dinan(a)gmail.com <mailto:[email protected]>> wrote: > Hi All, > > We have an invited speaker this week at the HACC WG: > > Topic: CUDA Deep Dive For the MPI Forum HACC WG > When: Wednesday, March 10 10-11:00am ET > Connection Info: https://github.com/mpiwg-hybrid/hybrid-issues/wiki <https://github.com/mpiwg-hybrid/hybrid-issues/wiki> > > Speaker: Stephen Jones, NVIDIA > > Stephen Jones is one of the architects of CUDA, working on defining the language, the platform, and the hardware that it runs on, to span the needs of parallel programming from high performance computing to artificial intelligence. Prior to his present position, he lead the Simulation & Analytics group at SpaceX, working on large-scale simulation of rocket engines. He has worked in diverse other industries, including networking, CAD/CAM, and scientific computing. He has been a part of CUDA since 2008. > > Cheers, > ~Jim. > > PS - Apologies for cross posting on the main list. If you would like to continue receiving emails relating to the Hybrid & Accelerator WG, please sign up for the mailing list here: https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm <https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm>. > _______________________________________________ > mpiwg-hybridpm mailing list > mpiwg-hybridpm(a)lists.mpi-forum.org <mailto:[email protected]> > https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm <https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm> > _______________________________________________ > mpiwg-hybridpm mailing list > mpiwg-hybridpm(a)lists.mpi-forum.org > https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm

1 0

Re: [mpiwg-hybridpm] Hybrid/Accelerator WG Meeting
by Junchao Zhang 11 Mar '21

11 Mar '21

In PETSc, we don't use non-contiguous MPI data types, so packing is not a problem for us. It seems if MPI can send/recv data without participation of GPU, then it is a workable solution in limited cases. My main concern about CUDA-graph is it is not modularizable. 1) Wtih cudaGraphAddKernelNode() etc, one has to see the whole graph at all. If kernels are spread and deeply wrapped in CPU routines as in PETSc, at a higher level, one doesn't know the kernels and hence nodes in the graph. 2) With graph capture, again, a caller subroutine can not guarantee a callee subroutine will execute the same path in next iteration as the one when it was captured. --Junchao Zhang On Thu, Mar 11, 2021 at 9:07 AM Jim Dinan <james.dinan(a)gmail.com> wrote: > Unfortunately, CPU callbacks are not a perfect solution on their own. CUDA > does not allow CUDA calls from within CPU callbacks, so for example you > would not be able to launch data packing kernels or peer-to-peer copy > operations from within the callback. However, you can use CPU callbacks to > signal a thread in the MPI runtime to process the operation. Another option > in this design space is to use CUDA memops (e.g. cuStreamWriteValue64 or > cuStreamWaitValue64) to coordinate between CUDA streams and MPI > communication helper threads. Because memops are processed from within the > GPU control processor that manages stream execution, I would expect these > to have lower overheads than CPU callbacks (although I haven't measured > this). > > ~Jim. > > On Wed, Mar 10, 2021 at 10:08 PM Junchao Zhang <junchao.zhang(a)gmail.com> > wrote: > >> Jim, >> Thanks for the slides. In Stephen's presentation today, it seems >> with existing techniques, i.e, CPU MPI callback nodes in CUDA graphs, one >> can solve the MPI GPU problem. Is my understanding correct? >> >> Thanks. >> --Junchao Zhang >> >> >> On Wed, Mar 10, 2021 at 8:34 PM Jim Dinan via mpiwg-hybridpm < >> mpiwg-hybridpm(a)lists.mpi-forum.org> wrote: >> >>> Hi All, >>> >>> I've posted Stephen's slides: >>> https://github.com/mpiwg-hybrid/hybrid-issues/tree/master/slides >>> >>> Best, >>> ~Jim. >>> >>> On Mon, Mar 8, 2021 at 11:21 AM Jim Dinan <james.dinan(a)gmail.com> wrote: >>> >>>> Hi All, >>>> >>>> We have an invited speaker this week at the HACC WG: >>>> >>>> Topic: CUDA Deep Dive For the MPI Forum HACC WG >>>> When: Wednesday, March 10 10-11:00am ET >>>> Connection Info: https://github.com/mpiwg-hybrid/hybrid-issues/wiki >>>> >>>> Speaker: Stephen Jones, NVIDIA >>>> >>>> Stephen Jones is one of the architects of CUDA, working on defining the >>>> language, the platform, and the hardware that it runs on, to span the needs >>>> of parallel programming from high performance computing to artificial >>>> intelligence. Prior to his present position, he lead the Simulation & >>>> Analytics group at SpaceX, working on large-scale simulation of rocket >>>> engines. He has worked in diverse other industries, including networking, >>>> CAD/CAM, and scientific computing. He has been a part of CUDA since 2008. >>>> >>>> Cheers, >>>> ~Jim. >>>> >>>> PS - Apologies for cross posting on the main list. If you would like to >>>> continue receiving emails relating to the Hybrid & Accelerator WG, please >>>> sign up for the mailing list here: >>>> https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm. >>>> >>> _______________________________________________ >>> mpiwg-hybridpm mailing list >>> mpiwg-hybridpm(a)lists.mpi-forum.org >>> https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm >>> >>

1 0

Re: [mpiwg-hybridpm] Hybrid/Accelerator WG Meeting
by Jim Dinan 11 Mar '21

11 Mar '21

Unfortunately, CPU callbacks are not a perfect solution on their own. CUDA does not allow CUDA calls from within CPU callbacks, so for example you would not be able to launch data packing kernels or peer-to-peer copy operations from within the callback. However, you can use CPU callbacks to signal a thread in the MPI runtime to process the operation. Another option in this design space is to use CUDA memops (e.g. cuStreamWriteValue64 or cuStreamWaitValue64) to coordinate between CUDA streams and MPI communication helper threads. Because memops are processed from within the GPU control processor that manages stream execution, I would expect these to have lower overheads than CPU callbacks (although I haven't measured this). ~Jim. On Wed, Mar 10, 2021 at 10:08 PM Junchao Zhang <junchao.zhang(a)gmail.com> wrote: > Jim, > Thanks for the slides. In Stephen's presentation today, it seems > with existing techniques, i.e, CPU MPI callback nodes in CUDA graphs, one > can solve the MPI GPU problem. Is my understanding correct? > > Thanks. > --Junchao Zhang > > > On Wed, Mar 10, 2021 at 8:34 PM Jim Dinan via mpiwg-hybridpm < > mpiwg-hybridpm(a)lists.mpi-forum.org> wrote: > >> Hi All, >> >> I've posted Stephen's slides: >> https://github.com/mpiwg-hybrid/hybrid-issues/tree/master/slides >> >> Best, >> ~Jim. >> >> On Mon, Mar 8, 2021 at 11:21 AM Jim Dinan <james.dinan(a)gmail.com> wrote: >> >>> Hi All, >>> >>> We have an invited speaker this week at the HACC WG: >>> >>> Topic: CUDA Deep Dive For the MPI Forum HACC WG >>> When: Wednesday, March 10 10-11:00am ET >>> Connection Info: https://github.com/mpiwg-hybrid/hybrid-issues/wiki >>> >>> Speaker: Stephen Jones, NVIDIA >>> >>> Stephen Jones is one of the architects of CUDA, working on defining the >>> language, the platform, and the hardware that it runs on, to span the needs >>> of parallel programming from high performance computing to artificial >>> intelligence. Prior to his present position, he lead the Simulation & >>> Analytics group at SpaceX, working on large-scale simulation of rocket >>> engines. He has worked in diverse other industries, including networking, >>> CAD/CAM, and scientific computing. He has been a part of CUDA since 2008. >>> >>> Cheers, >>> ~Jim. >>> >>> PS - Apologies for cross posting on the main list. If you would like to >>> continue receiving emails relating to the Hybrid & Accelerator WG, please >>> sign up for the mailing list here: >>> https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm. >>> >> _______________________________________________ >> mpiwg-hybridpm mailing list >> mpiwg-hybridpm(a)lists.mpi-forum.org >> https://lists.mpi-forum.org/mailman/listinfo/mpiwg-hybridpm >> >

1 0