MVAPICH2-GPU: Op0mized GPU to GPU Communica0on for - - PowerPoint PPT Presentation

MVAPICH2-­‑GPU: ¡Op0mized ¡GPU ¡to ¡GPU ¡ Communica0on ¡for ¡InfiniBand ¡Clusters ¡

¡ ¡

• H. ¡Wang, ¡S. ¡Potluri, ¡M. ¡Luo, ¡A. ¡K. ¡Singh, ¡S. ¡Sur ¡
• D. ¡K. ¡Panda ¡

¡ Network-­‑Based ¡Compu0ng ¡Laboratory ¡ The ¡Ohio ¡State ¡University ¡

1 ISC ¡2011 ¡Hamburg ¡

Outline ¡

ISC 2011 Hamburg 2

• Introduc0on ¡
• Problem ¡Statement ¡
• Our ¡Solu0on: ¡MVAPICH2-­‑GPU ¡ ¡
• Design ¡Considera0ons ¡
• Performance ¡Evalua0on ¡
• Conclusion ¡& ¡Future ¡Work ¡

InfiniBand ¡Clusters ¡in ¡Top500 ¡

• Percentage ¡share ¡of ¡InfiniBand ¡is ¡steadily ¡increasing ¡ ¡
• 41% ¡of ¡systems ¡in ¡TOP ¡500 ¡using ¡InfiniBand ¡(June ¡’11) ¡
• 61% ¡of ¡systems ¡in ¡TOP ¡100 ¡using ¡InfiniBand ¡(June ¡‘11) ¡

3 ISC 2011 Hamburg

Growth ¡in ¡GPGPUs ¡

• GPGPUs ¡are ¡gaining ¡significance ¡on ¡clusters ¡for ¡data-­‑centric ¡

applica0ons: ¡

– Word ¡Occurrence, ¡Sparse ¡Integer ¡Occurrence ¡ – K-­‑means ¡clustering, ¡Linear ¡regression ¡

• GPGPUs ¡+ ¡InfiniBand ¡are ¡gaining ¡momentum ¡for ¡large ¡clusters ¡

– #2 ¡(Tianhe-­‑1A), ¡#4 ¡(Nebulae) ¡and ¡#5 ¡(Tsubame) ¡Petascale ¡systems ¡

• GPGPUs ¡programming ¡

– CUDA ¡or ¡OpenCL ¡+ ¡MPI ¡ – Dr. ¡Sumit ¡Gupta ¡briefed ¡industry ¡users ¡at ¡NVIDIA ¡mee0ng ¡yesterday ¡on ¡ programmability ¡advances ¡on ¡GPUs ¡

• Big ¡issues: ¡performance ¡of ¡data ¡movement ¡ ¡

– Latency ¡ – Bandwidth ¡ – Overlap ¡ ¡

¡

4 ISC 2011 Hamburg

Data ¡movement ¡in ¡GPU ¡clusters ¡

• Data ¡movement ¡in ¡InfiniBand ¡clusters ¡with ¡GPUs ¡

– CUDA: ¡Device ¡memory ¡à ¡Main ¡memory ¡ ¡[at ¡source ¡process] ¡ – MPI: ¡Source ¡rank ¡à ¡Des0na0on ¡process ¡ – CUDA: ¡Main ¡memory ¡à ¡Device ¡memory ¡ ¡[at ¡des0na0on ¡process] ¡

• GPU ¡and ¡InfiniBand ¡require ¡separate ¡memory ¡registra0on ¡

Main Memory GPU

IB

Adapter GPU

IB

Adapter Main Memory PCI-E Hub PCI-E Hub PCI-E PCI-E IB Network PCI-E PCI-E 5 ISC 2011 Hamburg

GPU ¡Direct ¡

without GPU Direct with GPU Direct

• Collabora0on ¡between ¡Mellanox ¡and ¡NVIDIA ¡to ¡converge ¡on ¡
• ne ¡memory ¡registra0on ¡technique ¡
• Both ¡devices ¡can ¡register ¡same ¡host ¡memory: ¡

– GPU ¡and ¡network ¡adapters ¡can ¡access ¡the ¡buffer ¡

6 ISC 2011 Hamburg

Outline ¡

ISC 2011 Hamburg 7

• Introduc0on ¡
• Problem ¡Statement ¡
• Our ¡Solu0on: ¡MVAPICH2-­‑GPU ¡ ¡
• Design ¡Considera0ons ¡
• Performance ¡Evalua0on ¡
• Conclusion ¡& ¡Future ¡Work ¡

Problem ¡Statement ¡

• Data ¡movement ¡from/to ¡GPGPUs ¡

– Performance ¡bojleneck ¡ – Reduced ¡programmer ¡produc0vity ¡

• Hard ¡to ¡op0mize ¡at ¡the ¡applica0on ¡level ¡

– CUDA ¡and ¡MPI ¡exper0se ¡required ¡for ¡efficient ¡implementa0on ¡ – Hardware ¡dependent ¡latency ¡characteris0cs ¡ – Hard ¡to ¡support ¡and ¡op0mize ¡collec0ves ¡ – Hard ¡to ¡support ¡advanced ¡features ¡like ¡one-­‑sided ¡communica0on ¡

8 ISC 2011 Hamburg

Outline ¡

ISC 2011 Hamburg 9

• Introduc0on ¡
• Problem ¡Statement ¡
• Our ¡Solu0on: ¡MVAPICH2-­‑GPU ¡ ¡
• Design ¡Considera0ons ¡
• Performance ¡Evalua0on ¡
• Conclusion ¡& ¡Future ¡Work ¡

MVAPICH2-­‑GPU: ¡Design ¡Goals ¡

• Support ¡GPU ¡to ¡GPU ¡communica0on ¡through ¡standard ¡MPI ¡

interfaces ¡

– e.g. ¡enable ¡MPI_Send, ¡MPI_Recv ¡from/to ¡GPU ¡memory ¡

• Provide ¡high ¡performance ¡without ¡exposing ¡low ¡level ¡details ¡

to ¡the ¡programmer ¡

– Pipelined ¡data ¡transfer ¡which ¡automa:cally ¡provides ¡op0miza0ons ¡ inside ¡MPI ¡library ¡without ¡user ¡tuning ¡

• Available ¡to ¡work ¡with ¡ ¡

– GPU ¡Direct ¡ – Without ¡GPU ¡Direct ¡

¡

10 ISC 2011 Hamburg

Sample ¡Code ¡-­‑ ¡without ¡MPI ¡integra0on ¡

• Naïve ¡implementa0on ¡with ¡MPI ¡and ¡CUDA ¡
• High ¡produc:vity ¡but ¡poor ¡performance ¡

At ¡Sender: ¡ ¡

¡cudaMemcpy(s_buf, ¡s_device, ¡size, ¡cudaMemcpyDeviceToHost); ¡ ¡MPI_Send(s_buf, ¡size, ¡MPI_CHAR, ¡1, ¡1, ¡MPI_COMM_WORLD); ¡

¡ At ¡Receiver: ¡

¡MPI_Recv(r_buf, ¡size, ¡MPI_CHAR, ¡0, ¡1, ¡MPI_COMM_WORLD, ¡&req); ¡ ¡cudaMemcpy(r_device, ¡r_buf, ¡size, ¡cudaMemcpyHostToDevice); ¡

11 ISC 2011 Hamburg

Sample ¡Code ¡– ¡User ¡op0mized ¡code ¡ ¡

• Pipelining ¡at ¡user ¡level ¡with ¡non-­‑blocking ¡MPI ¡and ¡CUDA ¡interfaces ¡
• Code ¡repeated ¡at ¡receiver ¡side ¡ ¡
• Good ¡performance ¡but ¡poor ¡produc:vity ¡

At ¡Sender: ¡ ¡

¡ ¡ ¡for ¡(j ¡= ¡0; ¡j ¡< ¡pipeline_len; ¡j++) ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡cudaMemcpyAsync(s_buf ¡+ ¡j ¡* ¡block_sz, ¡s_device ¡+ ¡j ¡* ¡block_sz, ¡…); ¡ ¡ ¡ ¡for ¡(j ¡= ¡0; ¡j ¡< ¡pipeline_len; ¡j++) ¡{ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡while ¡(result ¡!= ¡cudaSucess) ¡{ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡result ¡= ¡cudaStreamQuery(…); ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡if(j ¡> ¡0) ¡MPI_Test(…); ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡} ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡MPI_Isend(s_buf ¡+ ¡j ¡* ¡block_sz, ¡block_sz, ¡MPI_CHAR, ¡1, ¡1, ¡….); ¡ ¡ ¡ ¡ ¡ ¡} ¡ ¡ ¡ ¡ ¡MPI_Waitall(); ¡

12 ISC 2011 Hamburg

Sample ¡Code ¡– ¡MVAPICH2-­‑GPU ¡

• MVAPICH2-­‑GPU: ¡provides ¡standard ¡MPI ¡interfaces ¡for ¡GPU ¡
• High productivity and high performance!

At ¡Sender: ¡ ¡ ¡ ¡ ¡MPI_Send(s_device, ¡size, ¡…); ¡// ¡s_device ¡is ¡data ¡buffer ¡in ¡GPU ¡ ¡ ¡ ¡ At ¡Receiver: ¡ ¡ ¡ ¡ ¡ ¡MPI_Recv(r_device, ¡size, ¡…); ¡// ¡r_device ¡is ¡data ¡buffer ¡in ¡GPU ¡

13 ISC 2011 Hamburg

Outline ¡

ISC 2011 Hamburg 14

• Introduc0on ¡
• Problem ¡Statement ¡
• Our ¡Solu0on: ¡MVAPICH2-­‑GPU ¡ ¡
• Design ¡Considera0ons ¡
• Performance ¡Evalua0on ¡
• Conclusion ¡& ¡Future ¡Work ¡

Design ¡considera0ons ¡

• Memory ¡detec0on ¡

– CUDA ¡4.0 ¡introduces ¡Unified ¡Virtual ¡Addressing ¡(UVA) ¡ ¡ – MPI ¡library ¡can ¡differen0ate ¡between ¡device ¡memory ¡and ¡ host ¡memory ¡without ¡any ¡hints ¡from ¡the ¡user ¡

• Overlap ¡CUDA ¡copy ¡and ¡RDMA ¡transfer ¡

– Pipeline ¡DMA ¡of ¡data ¡from ¡GPU ¡and ¡InfiniBand ¡RDMA ¡ – Allow ¡for ¡progressing ¡DMAs ¡individual ¡data ¡chunks ¡

15 ISC 2011 Hamburg

Pipelined ¡Design ¡

– Data ¡is ¡divided ¡into ¡chunks ¡ – Pipeline ¡CUDA ¡copies ¡with ¡RDMA ¡transfers ¡ – If ¡system ¡does ¡not ¡have ¡GPU-­‑Direct, ¡an ¡extra ¡copy ¡is ¡required ¡

GPU ¡Device ¡ Memory ¡ Host ¡Main ¡ Memory ¡ MPI_Send ¡ … ¡

cudaMemcpy ¡

Async() ¡ … ¡

cudaStream ¡

Query() ¡ GPU ¡Device ¡ Memory ¡ Host ¡Main ¡ Memory ¡ MPI_Recv ¡ … ¡

cudaMemcpy ¡

Async() ¡ … ¡

cudaStream ¡

Query() ¡ … ¡

RTS/CTS ¡

RDMA ¡Write ¡ & ¡Finish ¡MSG ¡ with GPU-Direct 16 ISC 2011 Hamburg

Pipeline ¡Design ¡(Cont.) ¡

• Chunk ¡size ¡depends ¡on ¡CUDA ¡copy ¡cost ¡and ¡RDMA ¡latency ¡over ¡

the ¡network ¡

• Automa0c ¡tuning ¡of ¡chunk ¡size ¡

– Detects ¡CUDA ¡copy ¡and ¡RDMA ¡latencies ¡during ¡installa0on ¡ – Chunk ¡size ¡can ¡be ¡stored ¡in ¡configura0on ¡file ¡(mvapich.conf) ¡

• User ¡transparent ¡to ¡deliver ¡the ¡best ¡performance ¡

¡

17 ISC 2011 Hamburg

Outline ¡

ISC 2011 Hamburg 18

• Introduc0on ¡
• Problem ¡Statement ¡
• Our ¡Solu0on: ¡MVAPICH2-­‑GPU ¡ ¡
• Design ¡Considera0ons ¡
• Performance ¡Evalua0on ¡
• Conclusion ¡& ¡Future ¡Work ¡

Performance ¡Evalua0on ¡

• Experimental ¡environment ¡

– NVIDIA ¡Tesla ¡C2050 ¡ – Mellanox ¡QDR ¡InfiniBand ¡HCA ¡MT26428 ¡ – Intel ¡Westmere ¡processor ¡with ¡12 ¡GB ¡main ¡memory ¡ – MVAPICH2 ¡1.6, ¡CUDA ¡Toolkit ¡4.0, ¡GPUDirect ¡

• OSU ¡Micro ¡Benchmarks ¡

– The ¡source ¡and ¡des0na0on ¡addresses ¡are ¡in ¡GPU ¡device ¡memory ¡

19 ISC 2011 Hamburg

Ping ¡Pong ¡Latency ¡

• With ¡GPU-­‑Direct ¡

– 45% ¡improvement ¡compared ¡to ¡Memcpy+Send ¡(4MB ¡messages) ¡ – 24% ¡improvement ¡compared ¡to ¡MemcpyAsync+Isend ¡(4MB ¡messages) ¡

• Without ¡GPU-­‑Direct ¡

– 38% ¡improvement ¡compared ¡to ¡Memcpy+send ¡(4MB ¡messages) ¡ – 33% ¡improvement ¡compared ¡to ¡MemcpyAsync+Isend ¡(4MB ¡messages) ¡

with GPU Direct without GPU Direct

0 ¡ 500 ¡ 1000 ¡ 1500 ¡ 2000 ¡ 2500 ¡ 3000 ¡

32K ¡ 64K ¡ 128K ¡256K ¡512K ¡ 1M ¡ 2M ¡ 4M ¡ Time ¡(us) ¡ Message ¡Size ¡(bytes) ¡ Memcpy+Send ¡ MemcpyAsync+Isend ¡ MVAPICH2-­‑GPU ¡

0 ¡ 500 ¡ 1000 ¡ 1500 ¡ 2000 ¡ 2500 ¡ 3000 ¡

32K ¡ 64K ¡ 128K ¡256K ¡512K ¡ 1M ¡ 2M ¡ 4M ¡ Time ¡(us) ¡ Message ¡Size ¡(bytes) ¡ Memcpy-­‑Send ¡ MemcpyAsync+Isend ¡ MVAPICH2-­‑GPU ¡

20 ISC 2011 Hamburg

One-­‑sided ¡Communica0on ¡ ¡

• 45% ¡improvement ¡compared ¡to ¡Memcpy+Put ¡(with ¡GPU ¡Direct) ¡
• 39% ¡improvement ¡compared ¡with ¡Memcpy+Put ¡(without ¡GPU ¡Direct) ¡
• Similar ¡improvement ¡for ¡Get ¡opera0on ¡
• Major ¡improvement ¡in ¡programming ¡ ¡

– One ¡sided ¡communica0on ¡not ¡feasible ¡without ¡MPI+CUDA ¡support ¡ 0 ¡ 500 ¡ 1000 ¡ 1500 ¡ 2000 ¡ 2500 ¡ 3000 ¡ 3500 ¡ 32K ¡ 64K ¡ 128K ¡256K ¡512K ¡ 1M ¡ 2M ¡ 4M ¡ Time ¡(us) ¡ Message ¡Size ¡(bytes) ¡ Mempy+Put ¡w/ ¡GPU ¡Direct ¡ MVAPICH2-­‑GPU ¡w/ ¡GPU ¡Direct ¡ Mempy+Put ¡w/o ¡GPU ¡Direct ¡ MVAPICH2-­‑GPU ¡w/o ¡GPU ¡Direct ¡ 0 ¡ 500 ¡ 1000 ¡ 1500 ¡ 2000 ¡ 2500 ¡ 3000 ¡ 32K ¡ 64K ¡ 128K ¡256K ¡512K ¡ 1M ¡ 2M ¡ 4M ¡ Time ¡(us) ¡ Message ¡Size ¡(bytes) ¡ Mempy+Get ¡w/ ¡GPU ¡Direct ¡ MVAPICH2-­‑GPU ¡w/ ¡GPU ¡Direct ¡ Mempy+Get ¡w/o ¡GPU ¡Direct ¡ MVAPICH2-­‑GPU ¡w/o ¡GPU ¡Direct ¡

21 ISC 2011 Hamburg

Collec0ve ¡Communica0on ¡

0 ¡ 5 ¡ 10 ¡ 15 ¡ 20 ¡ 25 ¡ 32K ¡ 64K ¡ 128K ¡256K ¡512K ¡ 1M ¡ 2M ¡ 4M ¡ Time ¡(ms) ¡ Message ¡Size ¡(bytes) ¡ Mempy+Gather ¡w/ ¡GPU ¡Direct ¡ MVAPICH2-­‑GPU ¡w/ ¡GPU ¡Direct ¡ Mempy+Gather ¡w/o ¡GPU ¡Direct ¡ MVAPICH2-­‑GPU ¡w/o ¡GPU ¡Direct ¡ 0 ¡ 5 ¡ 10 ¡ 15 ¡ 20 ¡ 25 ¡ 32K ¡ 64K ¡ 128K ¡256K ¡512K ¡ 1M ¡ 2M ¡ 4M ¡ Time ¡(ms) ¡ Message ¡Size ¡(bytes) ¡ Mempy+Scajer ¡w/ ¡GPU ¡Direct ¡ MVAPICH2-­‑GPU ¡w/ ¡GPU ¡Direct ¡ Mempy+Scajer ¡w/o ¡GPU ¡Direct ¡ MVAPICH2-­‑GPU ¡w/o ¡GPU ¡Direct ¡

22 ISC 2011 Hamburg

• With ¡GPU ¡Direct ¡

– 37% ¡improvement ¡for ¡MPI_Gather ¡(1MB) ¡and ¡32% ¡improvement ¡for ¡MPI_Scajer ¡(4MB) ¡

• Without ¡GPU ¡Direct ¡

– 33% ¡improvement ¡for ¡MPI_Gather ¡(1MB) ¡and ¡23% ¡ ¡improvement ¡MPI_Scajer ¡ ¡(4MB) ¡

• With ¡GPU ¡Direct ¡

– 30% ¡improvement ¡for ¡4MB ¡messages ¡ 0 ¡ 10 ¡ 20 ¡ 30 ¡ 40 ¡ 50 ¡ 32K ¡ 64K ¡ 128K ¡256K ¡512K ¡ 1M ¡ 2M ¡ 4M ¡ Time ¡(ms) ¡ Message ¡Size ¡(bytes) ¡ Mempy+Alltoall ¡w/ ¡GPU ¡Direct ¡ MVAPICH2-­‑GPU ¡w/ ¡GPU ¡Direct ¡ Mempy+Alltoall ¡w/o ¡GPU ¡Direct ¡ MVAPICH2-­‑GPU ¡w/o ¡GPU ¡Direct ¡

Collective Communication: Alltoall

23 ISC 2011 Hamburg

• Without ¡GPU ¡Direct ¡

– 26% ¡improvement ¡for ¡4MB ¡messages ¡

Comparison ¡with ¡Other ¡MPI ¡Stack: ¡ Latency ¡

ISC 2011 Hamburg 24

• Open ¡MPI ¡implementa0on ¡by ¡Rolf ¡vandeVaart(NVIDIA) ¡

– hjps://bitbucket.org/rolfv/ompi-­‑trunk-­‑cuda-­‑3 ¡ ¡

• Small ¡message ¡latency ¡around ¡10% ¡bejer ¡(1B ¡– ¡512B) ¡
• Large ¡message ¡latency ¡around ¡40% ¡bejer ¡(4M) ¡

0 ¡ 10 ¡ 20 ¡ 30 ¡ 40 ¡ 50 ¡ 60 ¡

1 ¡ 4 ¡ 16 ¡ 64 ¡ 256 ¡ 1K ¡ 4K ¡ Latnecy ¡(us) ¡

Message ¡Size ¡(Bytes) ¡

Small ¡Messages ¡

MVAPICH2-­‑GPU ¡ OpenMPI-­‑CUDA ¡

0 ¡ 500 ¡ 1000 ¡ 1500 ¡ 2000 ¡ 2500 ¡

8K ¡ 32K ¡ 128K ¡ 512K ¡ 2M ¡

Latency ¡(us) ¡

Message ¡Size ¡(Bytes) ¡

Large ¡Messages ¡

MVAPICH2-­‑GPU ¡ OpenMPI-­‑CUDA ¡

Comparison ¡with ¡Other ¡MPI ¡Stack: ¡ Bandwidth ¡

ISC 2011 Hamburg 25

• Bandwidth ¡for ¡large ¡messages ¡512K ¡improved ¡by ¡almost ¡factor ¡of ¡three ¡

¡

500 1000 1500 2000 2500 3000 3500 1 4 16 64 256 1K 4K 16K 64K 256K 1M 4M Bandwidth (MBytes/s) Message Size (Bytes)

MVAPICH2-GPU OpenMPI-CUDA

Conclusion ¡& ¡Future ¡Work ¡

• GPUs ¡are ¡key ¡for ¡data-­‑centric ¡applica0ons ¡
• However, ¡data ¡movement ¡is ¡the ¡key ¡bojleneck ¡in ¡current ¡genera0on ¡

clusters ¡with ¡GPUs ¡

• Asynchronous ¡CUDA ¡Copy ¡and ¡RDMA ¡provide ¡opportunity ¡for ¡latency ¡

hiding ¡

• Op0miza0ons ¡at ¡applica0on ¡level ¡reduce ¡produc0vity ¡and ¡achieve ¡sub-­‑
• p0mal ¡performance ¡
• MPI ¡Library ¡can ¡handle ¡this ¡for ¡the ¡users ¡efficiently ¡
• MVAPICH2-­‑GPU ¡shows ¡substan0al ¡improvements ¡for ¡MPI ¡opera0ons ¡

– Will ¡be ¡available ¡in ¡future ¡MVAPICH2 ¡release ¡

• Future ¡work ¡includes ¡collec0ve ¡specific ¡op0miza0ons ¡and ¡applica0ons-­‑

level ¡studies ¡with ¡MVAPICH2-­‑GPU ¡

• Upcoming ¡paper ¡at ¡Cluster ¡2011: ¡“Op:mized ¡Non-­‑con:guous ¡MPI ¡Datatype ¡

Communica:on ¡for ¡GPU ¡Clusters: ¡Design, ¡Implementa:on ¡and ¡Evalua:on ¡with ¡ MVAPICH2”, ¡H. ¡Wang, ¡S. ¡Potluri, ¡M. ¡Luo, ¡A. ¡K. ¡Singh, ¡X. ¡Ouyang, ¡S. ¡Sur, ¡D. ¡K. ¡Panda ¡

26 ISC 2011 Hamburg

¡Thank ¡You! ¡

{wangh, ¡potluri, ¡luom, ¡singhas, ¡surs, ¡panda}@cse.ohio-­‑state.edu ¡

¡

¡

¡ ¡ Network-­‑Based ¡Compu0ng ¡Laboratory ¡ hjp://nowlab.cse.ohio-­‑state.edu/ ¡ MVAPICH ¡Web ¡Page ¡ hjp://mvapich.cse.ohio-­‑state.edu/

27