Heterogeneous Task Execution Frameworks in Charm++ Michael Robson - PowerPoint PPT Presentation

Heterogeneous Task Execution Frameworks in Charm++ Michael Robson Parallel Programming Lab Charm Workshop 2016

Charm++ GPU Frameworks 2 ¡

Accelerator Overview • NVIDIA GPUs – Programmed with CUDA – 1,000s of threads – 100s GB/s bandwidth – ~16 GB of memory – ~300 GFLOPS Double Precision 3 ¡

Charm++ GPU Frameworks 4 ¡

GPU Manager • Task Offload and Management Library • Advantages: 1. Automatic task management and synch. 2. Overlap data transfer and kernel invocation 3. Simplified workflow via callbacks 4. Reduce overhead via centralized management 5 ¡

GPU Manager • One queue of GPU requests per process • Utilize pinned memory pools • Integrated in mainline • Visualization in projections http://charm.cs.illinois.edu/manuals/html/ libraries/7.html 6 ¡

GPU Manager 7 ¡

GPU Manager 8 ¡

Using GPU Manager • Build charm with cuda target • Create and enqueue a work request – Mark/pass buffers – Give a callback to resume work • Write kernel launch functions 9 ¡

10 ¡

nodeGPU Manager • “Node-level” version of GPU Manager • One centralized queue per GPU • Enable GPU applications to run (well) in SMP mode https://charm.cs.illinois.edu/gerrit/#/c/ 802/ or branch: mprobson/nodeGPU_ff 11 ¡

nodeGPU Manager Improved API • Replace globals with functions • Register kernel launching functions • Convenience functions for marking buffers • Build with or without CUDA code 12 ¡

Improved API Example • enqueue(wrQueue, wr); -> enqueue (wr); • kernel<<…, kernel_stream>> -> • kernel<<…, getKernelStream()>> • dataInfo *info = new dataInfo; – info->hostBuffer = hapi_poolMalloc(size); – info->size = size; – memcpy(info->hostBuffer, data, size); – info->bufferID = -1; – info->transferToDevice = YES; – info->transferFromDevice = NO; – info->freeBuffer = YES; • initBuffer(info, siez, data, true, false, true); 13 ¡

Charm++ GPU Frameworks 14 ¡

[accel] Framework • Allow the runtime systems (RTS) to choose to execute on the host or device • RTS can proactively move needed data • RTS can map to various platforms • Originally targeted at cell processor 15 ¡

[accel] Framework • Builds on top of GPU manager • Annotate charm entry methods • Mark data as read, write, persistent, etc • Automatically generate accelerated code • Batch fine grained kernel launches https://charm.cs.illinois.edu/gerrit/#/c/ 824/ and branch: mprobson/accel-doc 16 ¡

[accel] Framework Example 17 ¡

[accel] Framework Example 18 ¡

[accel] Framework Usage • modifiers: – read-only, write-only, read-write – shared – one copy per batch – persist – resident in device memory • parameters: – triggered – one invocation per chare in array – splittable (int) – AEM does part of work – threadsPerBlock (int) – specify block size 19 ¡

$version • Allow users to write platform specific accelerator code • Either as two separate, equivalent kernels • Or machine specific sections/tweaks • Automatically generate multiple kernels https://charm.cs.illinois.edu/gerrit/#/c/ 1104/ 20 ¡

$version Target Specific 21 ¡

$version Two Implementations 22 ¡

Charm++ GPU Frameworks 23 ¡

NAMD GPU Acceleration • NAMD GPU code is about 5x faster than the CPU code – CPU version is becoming somewhat obsolete • General requirements – Keep data on device as much as possible – Use pinned host memory – Hide CUDA kernel launch latency • Merge all computation into few kernels • Avoid unnecessary cudaStreamSynchronize()

NAMD GPU Performance Speedup ¡vs. ¡NAMD ¡2.11 ¡ DHFR ¡(24K ¡atoms) ¡ 1.6 ¡ ApoA1 ¡(92K ¡atoms) ¡ 1.5 ¡ 1.4 ¡ 1.3 ¡ 1.2 ¡ 1.1 ¡ 1 ¡ 1 ¡ 2 ¡ 4 ¡ 8 ¡ Number ¡of ¡Titan ¡nodes ¡ • Explicit ¡solvent: ¡30% ¡-­‑ ¡57% ¡faster ¡ simula@ons ¡

NAMD GPU Performance 5.7M ¡atoms ¡ 13K ¡atoms ¡ 1.35 ¡ 4 ¡ 1.3 ¡ 3.5 ¡ 1.25 ¡ 3 ¡ 1.2 ¡ 2.5 ¡ 1.15 ¡ 2 ¡ 1.1 ¡ 1.5 ¡ 1.05 ¡ 1 ¡ 1 ¡ 1 ¡ 2 ¡ 4 ¡ 1 ¡ 2 ¡ 4 ¡ Number ¡of ¡Titan ¡nodes ¡ Number ¡of ¡Titan ¡nodes ¡ • GB ¡implicit ¡solvent: ¡Up ¡to ¡3.5x ¡faster ¡ simula@ons ¡

NAMD PME computation – case for direct GPU-GPU communication • Particle Mesh Ewald (PME) reciprocal computation requires a 3D FFT, which in turn requires repeated communications between GPUs • Communication is the bottleneck • In the current implementation, we must handle intra- and inter-node cases separately

Intra-node • Sending PE transposeDataOnGPU(d_data, ¡stream); ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡// ¡Transpose ¡data ¡locally ¡ copyDataToPeerDevice(destGPU, ¡d_data, ¡stream); ¡ ¡ ¡ ¡ ¡// ¡Copy ¡data ¡to ¡GPU ¡on ¡same ¡node ¡ cudaStreamSynchronize(stream); ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡// ¡Wait ¡for ¡CUDA ¡stream ¡to ¡finish ¡ ¡ PmeMsg* ¡msg ¡= ¡new ¡(0) ¡PmeMsg(); ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡// ¡Allocate ¡empty ¡message ¡ pmePencil.recvData(msg); ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡// ¡Send ¡message ¡to ¡PE ¡that ¡has ¡“destGPU” ¡ • Receiving ¡PE ¡ void ¡recvData(PmeMsg* ¡msg) ¡{ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡// ¡Receiving ¡empty ¡message ¡lets ¡PE ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡// ¡know ¡its ¡GPU ¡now ¡has ¡the ¡data ¡in ¡“d_data” ¡ ¡ ¡eWork(d_data, ¡stream); ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡ ¡// ¡Perform ¡work ¡on ¡data ¡ ¡ ¡… ¡ } ¡ • Requires ¡lots ¡of ¡tedious ¡work ¡from ¡the ¡user ¡ • Error ¡prone ¡