Languages, APIs And Development Tools For GPU Computing - Nvidia

Languages, APIs and Development Toolsfor GPU ComputingWill Ramey Sr. Product Manager for GPU ComputingSan Jose Convention Center, CA September 20–23, 2010

“GPU Computing” Using all processors in the system for the things they arebest at doing— Evolution of CPUs makes them good at sequential, serial tasks— Evolution of GPUs makes them good at parallel processing

MathematicalPackagesLibrariesConsultants, Training& CertificationResearch & EducationDirectXIntegratedDevelopment EnvironmentParallel Nsight for MS Visual StudioTools & PartnersGPU Computing EcosystemLanguages & API’sFortranNVIDIA ConfidentialAll Major Platforms

CUDA - NVIDIA’s Architecture for GPU ComputingBroad AdoptionOver 250M installedCUDA-enabled GPUsGPU Computing ApplicationsOver 650k CUDA Toolkitdownloads in last 2 YrsWindows, Linux andMacOS PlatformssupportedGPU Computing spansHPC to Consumer350 Universitiesteaching GPU Computingon the CUDA ArchitectureCUDAC/C OpenCLOver 100k developersRunning in Productionsince 2008SDK Libs VisualProfiler and DebuggerCommercial OpenCLConformant DriverPublic Availabilityacross all CUDAArchitecture GPU’sSDK Visual ProfilerDirectComputeMicrosoft API forGPU ComputingSupports all CUDAArchitecture GPUs(DX10 and DX11)FortranPGI AcceleratorPGI CUDA FortranNVIDIA GPUwith the CUDA Parallel Computing ArchitectureOpenCL is a trademark of Apple Inc. used under license to the Khronos Group Inc.Python,Java, .NET, PyCUDAGPU.NETjCUDA

GPU Computing Software StackYour GPU Computing ApplicationApplication Acceleration Engines (AXEs)Middleware, Modules & Plug-insFoundation LibrariesLow-level Functional LibrariesDevelopment EnvironmentLanguages, Device APIs, Compilers, Debuggers, Profilers, etc.CUDA Architecture

Languages & APIs NVIDIA Corporation 2010

Many Different Approaches Application level integration High level, implicit parallel languages Abstraction layers & API wrappers High level, explicit language integration Low level device APIs

GPUs for MathWorks Parallel Computing Toolbox and Distributed Computing Server WorkstationMATLAB Parallel Computing Toolbox (PCT)Compute ClusterMATLAB Distributed Computing Server (MDCS) PCT enables high performance throughparallel computing on workstations MDCS allows a MATLAB PCT application to besubmitted and run on a compute cluster NVIDIA GPU acceleration now available NVIDIA GPU acceleration now availableNVIDIA Confidential

MATLAB Performance with TeslaRelative Performance, Black-Scholes DemoCompared to Single Core CPU BaselineSingle Core CPUQuad Core CPUSingle Core CPU Tesla C1060Quad Core CPU Tesla C106012.0Relative Execution Speed10.08.06.04.02.0256 K1,024 K4,096 K16,384 KInput SizeCore 2 Quad Q6600 2.4 GHz, 6 GB RAM, Windows 7 64-bit, Tesla C1060, single precision operationsNVIDIA Confidential

SUBROUTINE SAXPY (A,X,Y,N)INTEGER NREAL A,X(N),Y(N)! ACC REGIONDO I 1, NX(I) A*X(I) Y(I)ENDDO! ACC END REGIONENDPGI AcceleratorCompilerscompileAuto-generated GPU codeHost x64 asm Filesaxpy : movlmovlcall. . .call call call calllink call Unifieda.outexecute(%rbx), %eax%eax, -4(%rbp)pgi cu initpgi cu functionpgi cu allocpgi cu uploadpgi cu callpgi cu download typedef struct dim3{ unsigned int x,y,z; }dim3;typedef struct uint3{ unsigned int x,y,z; }uint3;extern uint3 const threadIdx, blockIdx;extern dim3 const blockDim, gridDim;static attribute (( global )) voidpgicuda(attribute (( shared )) int tc,attribute (( shared )) int i1,attribute (( shared )) int i2,attribute (( shared )) int n,attribute (( shared )) float* c,attribute (( shared )) float* b,attribute (( shared )) float* a ){ int i; int p1; int i;i blockIdx.x * 64 threadIdx.x;if( i tc ){a[i i2-1] (( c[i i2-1] c[i i2-1]) b[i i2-1]);b[i i2-1] c[i i2];i ( i 1);p1 (p1-1);} } no change to existing makefiles, scripts, IDEs,programming environment, etc.

PyCUDA / PyOpenCLSlide courtesy of Andreas Klöckner, Brown Universityhttp://mathema.tician.de/software/pycuda

CUDA C: C with a few keywordsvoid saxpy serial(int n, float a, float *x, float *y){for (int i 0; i n; i)y[i] a*x[i] y[i];}// Invoke serial SAXPY kernelsaxpy serial(n, 2.0, x, y);Standard C Codeglobal void saxpy parallel(int n, float a, float *x, float *y){int i blockIdx.x*blockDim.x threadIdx.x;if (i n) y[i] a*x[i] y[i];CUDA C}// Invoke parallel SAXPY kernel with 256 threads/blockint nblocks (n 255) / 256;saxpy parallel nblocks, 256 (n, 2.0, x, y);Code

Write GPU kernels in C#, F#, VB.NET, etc. Exposes a minimal API accessible from any.NET-based language— Learn a new API instead of a new language JIT compilation dynamic language support Don’t rewrite your existing code— Just give it a ―touch-up‖

OpenCL Cross-vendor open standard— Managed by the Khronos Group Low-level API for device management andlaunching kernelshttp://www.khronos.org/opencl— Close-to-the-metal programming interface— JIT compilation of kernel programs C-based language for compute kernels— Kernels must be optimized for each processor architectureNVIDIA released the first OpenCL conformant driver forWindows and Linux to thousands of developers in June 2009

DirectCompute Microsoft standard for all GPU vendors— Released with DirectX 11 / Windows 7— Runs on all 100M CUDA-enabled DirectX 10 class GPUs and later Low-level API for device management and launching kernels— Good integration with DirectX 10 and 11 Defines HLSL-based language for compute shaders— Kernels must be optimized for each processor architecture

Language & APIs for GPU ComputingApproachExamplesApplication IntegrationMATLAB, Mathematica, LabVIEWImplicit Parallel Languages PGI Accelerator, HMPPAbstraction Layer/WrapperPyCUDA, CUDA.NET, jCUDALanguage IntegrationCUDA C/C , PGI CUDA FortranLow-level Device APICUDA C/C , DirectCompute, OpenCL

Development Tools NVIDIA Corporation 2010

Parallel Nsight for Visual StudioIntegrated development for CPU and GPUBuildDebugProfile

Windows GPU Development for 2010NVIDIA Parallel Nsight 1.5nvccFX Composercuda-gdbShader Debuggercuda-memcheckPerfHUDVisual ProfilerShaderPerfcudaprofPlatform Analyzer

4 Flexible GPU Development ConfigurationsDesktopSingle machine, Single NVIDIA GPUAnalyzer, Graphics InspectorSingle machine, Dual NVIDIA GPUsAnalyzer, Graphics Inspector, Compute DebuggerNetworkedTwo machines connected over the networkAnalyzer, Graphics Inspector, Compute Debugger, Graphics DebuggerTCP/IPWorkstation SLISLI Multi OS workstation with two Quadro GPUsAnalyzer, Graphics Inspector, Compute Debugger, Graphics Debugger NVIDIA Corporation 2010

NVIDIA cuda-gdbCUDA debugging integratedinto GDB on LinuxSupported on 32bit and 64bitsystemsSeamlessly debug both thehost/CPU and device/GPU codeSet breakpoints on any sourceline or symbol nameAccess and print all CUDAmemory allocs, local, global,constant and shared varsIncluded in the CUDA ToolkitParallel SourceDebugging

Allinea DDT debugger Latest News from Allinea CUDA SDK 3.0 with DDT 2.6 Released June 2010 Fermi and Tesla support cuda-memcheck support for memory errors Combined MPI and CUDA support Stop on kernel launch feature Kernel thread control, evaluation andbreakpoints Identify thread counts, ranges and CPU/GPUthreads easily SDK 3.1 in beta with DDT 2.6.1 SDK 3.2 Coming soon: multiple GPU device support

TotalView Debugger Latest from TotalView debugger (in Beta)—Debugging of application running on the GPU device—Full visibility of both Linux threads and GPU device threads——— Device threads shown as part of the parent Unix process Correctly handle all the differences between the CPU and GPUFully represent the hierarchical memory Display data at any level (registers, local, block, global or host memory) Making it clear where data resides with type qualificationThread and Block Coordinates Built in runtime variables display threads in a warp, block and threaddimensions and indexes Displayed on the interface in the status bar, thread tab and stack frameDevice thread control —Handles CUDA function inlining —Step in to or over inlined functionsReports memory access errors —Warps advance SynchronouslyCUDA memcheckCan be used with MPI

NVIDIA Visual ProfilerAnalyze GPU HW performancesignals, kernel occupancy,instruction throughput, and moreHighly configurabletables and graphical viewsSave/load profiler sessions orexport to CSV for later analysisCompare results visuallyacross multiple sessions tosee improvementsWindows, Linux and Mac OS XOpenCL support on Windows and LinuxIncluded in the CUDA Toolkit

GPU Computing SDKHundreds of code samples forCUDA C, DirectCompute and OpenCLFinanceOil & GasVideo/Image Processing3D Volume RenderingParticle SimulationsFluid SimulationsMath Functions

ApplicationDesign Patterns 2009 NVIDIA Corporation

Trivial ApplicationDesign Rules:Serial task processing on CPUData Parallel processing on GPUCopy input data to GPUPerform parallel processingCopy results backFollow guidance in theCUDA C Best Practices GuideApplicationCPU C RuntimeCPUThe CUDA C Runtime could be substitutedwith other methods of accessing the rFortranRuntimeDriverAPIGPUGPUMemory

Basic Application“Trivial Application” plus:Maximize overlap of data transfers and computationMinimize communication required between processorsUse one CPU thread to manage each GPUApplicationCPU C RuntimeCPUMulti-GPU notebook, desktop,workstation and cluster nodeconfigurations are increasingly commonCPUMemoryCUDA C RuntimeGPUGPUMemoryGPUGPUMemory

Graphics Application“Basic Application” plus:Use graphics interop to avoid unnecessary copiesIn Multi-GPU systems, put buffers to be displayed in GPUMemory of GPU attached to the displayApplicationCPU C RuntimeCPUCPUMemoryCUDA C RuntimeGPUOpenGL / Direct3DGPUMemory

Basic Library“Basic Application” plus:Avoid unnecessary memory transfersUse data already in GPU memoryCreate and leave data in GPU memoryLibraryCPU C RuntimeCPUThese rules apply to plug-ins as wellCPUMemoryCUDA C RuntimeGPUGPUMemory

Application with Plug-ins“Basic Application” plus:Plug-in MgrAllows Application and Plug-insto (re)use same GPU memoryMulti-GPU awareFollow “Basic Library” rulesfor the Plug-insApplicationPlug-in MgrPlug-inCPU C RuntimeCPUCPUMemoryPlug-inPlug-inCUDA C RuntimeGPUGPUMemory

Database ApplicationMinimize network communicationMove analysis “upstream”to stored proceduresClient ApplicationorApplication ServerTreat each stored procedurelike a “Basic Application”App Server could also be a“Basic Application”Client Application is also a“Basic Application”Database EngineCPU C RuntimeCPUData Mining, Business Intelligence, etc.CPUMemoryStored ProcedureCUDA C RuntimeGPUGPUMemory

Multi-GPU Cluster ApplicationApplication“Basic Application” plus:CPU C RuntimeUse Shared Memory forintra-node communicationor pthreads, OpenMP, etc.Use MPI to communicatebetween nodesMPI over Ethernet, Infiniband, etc.CPUCPUMemoryCUDA C RuntimeGPUGPUMemoryGPUGPUMemoryApplicationCPU C RuntimeCPUCPUMemoryCUDA C RuntimeGPUGPUMemoryGPUGPUMemoryApplicationCPU C RuntimeCPUCPUMemoryCUDA C RuntimeGPUGPUMemoryGPUGPUMemory

Libraries 2009 NVIDIA Corporation

CUFFT 3.2: Improved Radix-3, -5, -7Radix-3 (SP, ECC off)Radix-3 (DP, ECC off )7025060200150C2070 R3.2C2070 R3.1100GFLOPSGFLOPS5040C2070 R3.2C2070 R3.130MKLMKL2050100012345678910 11 12 13 14 15log3(size)Radix-5, -7 and mixed radix improvements not shownCUFFT 3.2 & 3.1 on NVIDIA Tesla C2070 GPUMKL 10.2.3.029 on Quad-Core Intel Core i7 (Nehalem)123456789 10 11 12 13 14 15log3(size)

CUBLAS Performance12xUp to 2x average speedup over CUBLAS 3.1Speedup vs. MKL10x8xLess variation in performancefor different dimensions vs. 3.16x4x2xMKLv3.10x1024204830724096Matrix dimensions (NxN)Average speedup of {S/D/C/Z}GEMM x {NN,NT,TN,TT}CUFFT 3.2 & 3.1 on NVIDIA Tesla C2050 GPUMKL 10.2.3.029 on Quad-Core Intel Core i7 (Nehalem)512061447168v3.2