PyCUDA: Even Simpler GPU Programming with Python

1 GPU Scripting PyOpenCL News RTCG ShowcasePyCUDA: Even SimplerGPU Programming with PythonAndreas Kl ocknerCourant Institute of Mathematical SciencesNew York UniversityNvidia GTC September 22, 2010 Andreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseThanksJan Hesthaven (Brown)Tim Warburton (Rice)Leslie Greengard (NYU)PyCUDA contributorsPyOpenCL contributorsNvidia CorporationAndreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOutline1 Scripting GPUs with PyCUDA2 PyOpenCL3 The News4 Run-Time Code Generation5 ShowcaseAndreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being ProductiveOutline1 Scripting GPUs with PyCUDAPyCUDA: An OverviewDo More, Faster with PyCUDA2 PyOpenCL3 The News4 Run-Time Code Generation5 ShowcaseAndreas Kl ocknerPyCUDA.

2 Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being ProductiveWhetting your as cuda2importpycuda. autoinit3importnumpy45a = (4,4).astype( )6agpu = ( ) (agpu, a)[This the PyCUDA distribution.]Andreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being ProductiveWhetting your appetite1mod = ( 2globalvoid twice ( float a)3{4int idx = + 4;5a[ idx ] = 2;6}7 )89func = ( twice )10func(agpu, block=(4,4,1))1112adoubled = (a) (adoubled, agpu)14printadoubled15printaCompute kernelAndreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being ProductiveWhetting your appetite1mod = ( 2globalvoid twice ( float a)3{4int idx = + 4;5a[ idx ] = 2.)}

3 6}7 )89func = ( twice )10func(agpu, block=(4,4,1))1112adoubled = (a) (adoubled, agpu)14printadoubled15printaCompute kernelAndreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being ProductiveWhy do Scripting for GPUs?GPUs are everything that scriptinglanguages are parallelVery architecture-sensitiveBuilt for maximum FP/memorythroughput complement each otherCPU: largely restricted to controltasks ( 1000/sec)Scripting fast enoughPython + CUDA =PyCUDAP ython + OpenCL =PyOpenCLAndreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being ProductiveScripting: PythonOne example of a scripting language.

4 PythonMatureLarge and active communityEmphasizes readabilityWritten in widely-portable CA multi-paradigm languageRich ecosystem of sci-comp relatedsoftwareAndreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being ProductiveScripting: Interpreted, not CompiledProgram creation workflow:EditCompileLinkRunAndreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being ProductiveScripting: Interpreted, not CompiledProgram creation workflow:EditCompileLinkRunAndreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being ProductiveScripting: Interpreted, not CompiledProgram creation workflow:EditCompileLinkRunAndreas Kl ocknerPyCUDA.

5 Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being ProductivePyCUDA: WorkflowEditPyCUDARunSourceModule("..")C ache? to GPURun on GPUA ndreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being ProductiveHow are High-Performance Codes constructed? Traditional Construction ofHigh-Performance Codes:C/C++/FortranLibraries Alternative Construction ofHigh-Performance Codes:Scripting for brains GPUs for inner loops Play to the strengths of eachprogramming Kl ocknerPyCUDA.

6 Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being ProductivePyCUDA PhilosophyProvide complete accessAutomatically manage resourcesProvide abstractionsCheck for and report errorsautomaticallyFull documentationIntegrate tightly withnumpyAndreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being ProductiveWhat s this numpy , anyway?Numpy: package for large,multi-dimensional , Matrices, ..A+B,sin(A),dot(A,B) (A, b), (A)cube[:, :, n-k:n+k],cube+5 All much faster than functional equivalents inPython.

7 Python s MATLAB :Basis for SciPy, plotting, ..Andreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being Productivegpuarray: Simple Linear :Meant to look and feel just (numpyarray)numpyarray = ()+, -, , /, fill, sin, exp, rand,basic indexing, norm, inner product, ..Mixed types (int32 + float32 = float64)print gpuarray for access to raw bitsUse as kernel arguments, textures, Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being ProductiveWhetting your appetite, Part II1importnumpy2importpycuda.

8 As gpuarray45agpu = ( (4,4).astype( ))7adoubled = (2 agpu).get()8printadoubled9printagpuAndre as Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being Productivegpuarray: Elementwise expressionsAvoiding extra store-fetch cycles for elementwise as curandagpu = curand((50,))bgpu = curand((50,)) = ElementwiseKernel( float a, float x, float b, float y, float z , z[ i ] = a x[i ] + b y[i ] )cgpu = gpuarray. emptylike (agpu)lincomb (5, agpu, 6, bgpu, cgpu)assert la.

9 Norm((cgpu (5 agpu+6 bgpu)).get())<1e 5 Andreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being Productivegpuarray: Reduction made easyExample: A scalar product = ReductionKernel(dtypeout= , neutral= 0 ,reduceexpr = a+b , mapexpr= x[i] y[i] ,arguments= const float x, const float y ) as curandx = curand((1000 1000), dtype= )y = curand((1000 1000), dtype= )xdoty = dot(x, y ). get ()xdotycpu = ( (), y . get ())Andreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOverview Being ProductivePyCUDA: Vital documentationMIT License(no warranty, free for all use)Requires: numpy, Python +(Win/OS X/Linux)Support via mailing listAndreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOutline1 Scripting GPUs with PyCUDA2 PyOpenCL3 The News4 Run-Time Code Generation5 ShowcaseAndreas Kl ocknerPyCUDA.

10 Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseOpenCL s perception problemOpenCL does not presently get thecredit it abstraction works well forGPUs, CPUsVendor-independenceCompute Dependency DAGA JIT C compiler baked into alibraryAndreas Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseIntroducing.. PyOpenCLPyOpenCL is PyCUDA for OpenCL Complete, mature API wrapperHas: Arrays, elementwiseoperations, RNG, ..Near feature parity with PyCUDAT ested on all availableImplementations, Kl ocknerPyCUDA: Even Simpler GPU Programming with PythonGPU Scripting PyOpenCL News RTCG ShowcaseIntroducing.