Část 4
NOVÉ RYSY CUDA, NOVINKY V OBLASTI NÁSTROJŮ PRO CUDA
Základní nástroje pro CUDA Linux: • CUDA-gdb (debug) • Memcheck (i racecheck) • Visual Profiler MS Windows • Nsight (integruje výše zmíněné) Společná nevýhoda: debugging i profiling ovlivňují běh aplikace
Nástup architektury Pascal TESLA P100 : • 16nm GPU "GP100" 15,3 miliardy tranzistorů • HBM2 paměti s propustností 1TB/s • Má mít v první verzi 3584 výpočetních jednotek FP32 a 1792 jednotek FP64. • Výkon: – 21,2TFLOPS v HP (FP16) – 10,6TFLOPS v SP (FP32) – 5,304TFLOPS v DP (FP64)
GPU datové typy Sizeof in bytes
Exponent in bits
Mantisa in bits
valid decimal digits
Maximal number
Single
4
8
23+1
6-7
3.39*10^38
Double
8
11
52+1
15-16
1.8*10^308
Half
2
5
10+1
3.3
65504
NVIDIA DGX-1 •DEEP LEARNING SYSTÉM •170 TFLOPS (GPU FP16), 3 TFLOPS (CPU FP32) •16 GB per GPU •Dual 20-core Intel® Xeon® E5-2698 v4 2.2 GHz •512 GB 2133 MHz DDR4 •GPU cores 28672 •Ubuntu Server Linux OS
NVIDIA NVLINK HIGH-SPEED INTERCONNECT The technology allows data sharing at rates 5 to 12 times faster than the traditional PCIe Gen3 interconnect
Kepler unified memory • GPU nemůže přistupovat do paměti CPU a opačně • Práce s paměťovými bloky existujícími na CPU i GPU – Dříve nutno explicitně alokovat a synchronizovat – Unified memory ulehčuje tyto operace
• Dostupná od CUDA SDK 6.0 a CC 3.0
Tento paměťový blok bude alokován jak na CPU tak na GPU
double* mA; cudaMallocManaged(&mA, ARRAY_SIZE * sizeof(double)); for (int i = 0; i < ARRAY_SIZE; i++) Zápis pomocí CPU mA[i] = 1.0*i; increment<<<2, 512>>>(mA, inc_val, ARRAY_SIZE); cudaDeviceSynchronize(); // práce s mA Zápis pomocí GPU cudaFree(mA);
Čekání na dokončení kernelu
Pascal GP100 Unified Memory GP100 extends GPU: • enable 49-bit virtual addressing. (48-bit virtual address spaces of modern CPUs+ GPU’s own memory) •allows programs to access the full address spaces of all CPUs and GPUs in the system as a single virtual address space, unlimited by the physical memory size of any one processor.
Možnost dynamického paralelismu • Od Compute capability 3.5 • Kernel může volat sub/kernel •Může vést k lepšímu load-balancingu •Ale díky větší režii volání jsou GPU rek. stromy volání nižší a širší než CPU stromy
Dynamické volání kernelu I • Nová možnost dynamického paralelismu • Vhodné pro neregulární situace, kdy „klasické“ řešení není efektivní • Od CUDA CC 3.5 může kernel volat jiné kernely včetně (případně) jiného nastavení počtu bloků a vláken • Kernel je pak ukončen teprve po ukončení všech volaných vnořených kernelů (volaných jednotlivými vlákny)
Dynamické volání kernelu II • CUDA Device Runtime garantuje, že otec a syn (kernely) mají fully consistent view of global memory, ale jen v určitých situacích • Pro zajištění ale často nutno explicitně synchronizovat pomocí __syncthreads a cudaDeviceSynchronize • Pozor! Ač hloubka vnoření není většinou omezena, pro efektivní běh by i tyto dynamicky volané kernely měly obsahovat dostatečný počet bloků a vláken
CUDA proudy • proud (stream) – fronta GPU operací (spouštění kernelů, kopírování dat, události), které se provádějí v pořadí v jakém byly do fronty vloženy • některé operace asynchronní s CPU např. asynchronní kopírování dat vyžaduje page-locked paměť • položka deviceOverlap ve vlastnostech zařízení znamená schopnost současně spouštět kernel a kopírovat data z/do GPU • položka concurrentKernels ve vlastnostech zařízení znamená schopnost současně spouštět více kernelů a kopírovat data z/do GPU současně (CC 2.0) • Od CC ? možno prioritizovat proudy
Knihovny v rámci SDK cuFFT cuBLAS cuSPARSE cuSOLVER cuRAND NPP Thrust math.h cuDNN Fast Fourier Transforms Library Complete BLAS Library
Sparse Matrix Library Linear Solver Library Random Number Generation (RNG) Library Performance Primitives for Image & Video Processing Templated Parallel Algorithms & Data Structures C99 floating-point Library Deep Neural Net building blocks
CUDA C++ Parallel Template Library Parallel Template library for CUDA C++ Host and Device Containers that mimic the C++ STL Optimized Parallel Algorithms for sort, reduce, scan, etc. TBB and OpenMP CPU Backends Performance portable
NPP: NVIDIA Performance Primitives Over 5000 image and signal processing routines: color transforms, geometric transforms, move operations, linear filters, image & signal statistics, image & signal arithmetic, JPEG building blocks, image segmentation, median filter, BGR/YUV conversion, 3D LUT color conversion
