University:
Boston University
Course:
GRS CS 910 | Directed Study: Systems
Academic year:

2024
Views:

420

Pages:

20
Author:

Nyasia Cardenas

Metaprogramming In GPU scripting, GPU code does not need to be a compile-time constant. Metaprogramming In GPU scripting, GPU code does not need to be a compile-time constant. (Key: Code is data–it wants to be reasoned about at run time) Metaprogramming Idea In GPU scripting, GPU code does not need to be a compile-time constant. (Key: Code is data–it wants to be reasoned about at run time) Metaprogramming Idea Python Code GPU Code GPU Compiler In GPU scripting, GPU code does not need to be a compile-time constant. GPU Binary GPU Result (Key: Code is data–it wants to be reasoned about at run time) Metaprogramming Idea In GPU scripting, GPU code does not need to be a compile-time constant. Python Code GPU Code GPU Compiler GPU Binary GPU Result Machine (Key: Code is data–it wants to be reasoned about at run time) Metaprogramming Idea Human Python Code GPU Code GPU Compiler In GPU scripting, GPU code does not need to be a compile-time constant. GPU Binary GPU Result (Key: Code is data–it wants to be reasoned about at run time) Metaprogramming Idea Python Code GPU Code GPU Compiler Good for code generation In GPU scripting, GPU code does not need to be a compile-time constant. GPU Binary GPU Result (Key: Code is data–it wants to be reasoned about at run time) Metaprogramming Idea Python Code GPU Code GPU Compiler Good for code generation UDA PyCscripting, In GPU GPU code does not need to be a compile-time constant. GPU Binary GPU Result (Key: Code is data–it wants to be reasoned about at run time) Metaprogramming Idea Python Code GPU Code GPU Compiler Good for code generation A PyPOp UDCL en yCscripting, In GPU GPU code does not need to be a compile-time constant. GPU Binary GPU Result (Key: Code is data–it wants to be reasoned about at run time) Machine-generated Code Why machine-generate code? Automated Tuning (cf. ATLAS, FFTW) Data types Specialize code for given problem Constants faster than variables (→ register pressure) Loop Unrolling PyOpenCL: Support for Metaprogramming Three (main) ways of generating code: Simple %-operator substitution Combine with C preprocessor: simple, often sufficient Use a templating engine (Mako works very well) codepy: Build C syntax trees from Python Generates readable, indented C Many ways of evaluating code–most important one: Exact device timing via events How are High-Performance Codes constructed? “Traditional” Construction of High-Performance Codes: C/C++/Fortran Libraries “Alternative” Construction of High-Performance Codes: Scripting for ‘brains’ GPUs for ‘inner loops’ Play to the strengths of each programming environment. Outline 1 Leftovers 2 Code writes Code The Idea RTCG in Action How can I do it? 3 Case Study: Generic OpenCL Reduction 4 Reasoning about Generated Code 5 Automatic GPU Programming pyopencl.array: Simple Linear Algebra pyopencl.array.Array: Meant to look and feel just like numpy. p.a.to device(ctx, queue, numpy array) numpy array = ary.get() +, -, ∗, /, fill, sin, arange, exp, rand, . . . Mixed types (int32 + float32 = float64) print cl array for debugging. Allows access to raw bits Use as kernel arguments, memory maps PyOpenCL Arrays: General Usage Remember your first PyOpenCL program? Abstraction is good: 1 2 3 4 5 6 7 8 9 10 11 12 import numpy import pyopencl as cl import pyopencl.array as cl array ctx = cl. create some context () queue = cl.CommandQueue(ctx) a gpu = cl array . to device ( ctx , queue, numpy.random.randn(4,4).astype(numpy.float32)) a doubled = (2∗a gpu).get() print a doubled print a gpu PyOpenCL Arrays: General Usage Remember your first PyOpenCL program? Abstraction is good: 1 2 3 4 5 6 7 8 9 10 11 12 import numpy import pyopencl as cl import pyopencl.array as cl array ctx = cl. create some context () queue = cl.CommandQueue(ctx) a gpu = cl array . to device ( ctx , queue, numpy.random.randn(4,4).astype(numpy.float32)) a doubled = (2∗a gpu).get() print a doubled Why is code generation useful in the impleprint a gpu mentation of the array type? pyopencl.elementwise: Elementwise expressions Avoiding extra store-fetch cycles for elementwise math: n = 10000 a gpu = cl array . to ctx , queue, b gpu = cl array . to ctx , queue, device ( numpy.random.randn(n).astype(numpy.float32)) device ( numpy.random.randn(n).astype(numpy.float32)) from pyopencl.elementwise import ElementwiseKernel lin comb = ElementwiseKernel(ctx, ” float a, float ∗x, float b, float ∗y, float ∗z”, ”z[ i ] = a∗x[i ] + b∗y[i]”) c gpu = cl array . empty like (a gpu) lin comb(5, a gpu, 6, b gpu, c gpu) import numpy.linalg as la assert la .norm((c gpu − (5∗a gpu+6∗b gpu)).get()) < 1e−5 pyopencl.reduction: Reduction made easy Example: A dot product calculation from pyopencl.reduction import ReductionKernel dot = ReductionKernel(ctx, dtype out=numpy.float32, neutral=”0”, reduce expr=”a+b”, map expr=”x[i]∗y[i]”, arguments=” global const float ∗x, global const float ∗y”) import pyopencl.clrandom as cl rand x = cl rand .rand(ctx , queue, (1000∗1000), dtype=numpy.float32) y = cl rand .rand(ctx , queue, (1000∗1000), dtype=numpy.float32) x dot y = dot(x, y ). get() x dot y cpu = numpy.dot(x.get(), y. get ()) Outline 1 Leftovers 2 Code writes Code The Idea RTCG in Action How can I do it? 3 Case Study: Generic OpenCL Reduction 4 Reasoning about Generated Code 5 Automatic GPU Programming RTCG via Substitution source = (””” kernel void %(name)s(%(arguments)s) { unsigned lid = get local id (0); unsigned gsize = get global size (0); unsigned work item start = get local size (0)∗ get group id (0); for (unsigned i = work item start + lid ; i < n; i += gsize) { %(operation)s; } } ””” % { ”arguments”: ”, ”. join (arg . declarator () for arg in arguments), ”operation”: operation , ”name”: name, ”loop prep”: loop prep , }) prg = cl.Program(ctx, source ). build ()

Easy, Effective, Efficient - GPU Programming in Python with PyOpenCL and PyCUDA Part 2

Related Documents

Recommended Documents

Description

Related Documents

Free up your schedule!

Take 5 seconds to unlock