University:
Boston University
Course:
MET CS 201 | Introduction to Programming
Academic year:

2024
Views:

219

Pages:

15
Author:

Jarrett Richard

Command Queues and Buffers: A Crashy Puzzle 4 OK a = numpy.random.rand(256∗∗3).astype(numpy.float32) a dev = cl. Buffer (ctx , cl .mem flags.READ WRITE, size=a.nbytes) cl . enqueue write buffer (queue, a dev, a, is blocking =False) Command Queues and Buffers: A Crashy Puzzle 4 OK a = numpy.random.rand(256∗∗3).astype(numpy.float32) a dev = cl. Buffer (ctx , cl .mem flags.READ WRITE, size=a.nbytes) cl . enqueue write buffer (queue, a dev, a, is blocking =False) 6 Crash a dev = cl. Buffer (ctx , cl .mem flags.READ WRITE, size=256∗∗3∗4) cl . enqueue write buffer (queue, a dev, numpy.random.rand(256∗∗3).astype(numpy.float32), is blocking =False) Command Queues and Buffers: A Crashy Puzzle 4 OK a = numpy.random.rand(256∗∗3).astype(numpy.float32) a dev = cl. Buffer (ctx , cl .mem flags.READ WRITE, size=a.nbytes) cl . enqueue write buffer (queue, a dev, a, is blocking =False) 6 Crash a dev = cl. Buffer (ctx , cl .mem flags.READ WRITE, size=256∗∗3∗4) cl . enqueue write buffer (queue, a dev, numpy.random.rand(256∗∗3).astype(numpy.float32), is blocking =False) 4 OK a dev = cl. Buffer (ctx , cl .mem flags.READ WRITE, size=256∗∗3∗4) cl . enqueue write buffer (queue, a dev, numpy.random.rand(256∗∗3).astype(numpy.float32), is blocking =True) Command Queues and Buffers: A Crashy Puzzle 4 OK (usually!) a = numpy.random.rand(256∗∗3).astype(numpy.float32) a dev = cl. Buffer (ctx , cl .mem flags.READ WRITE, size=a.nbytes) cl . enqueue write buffer (queue, a dev, a, is blocking =False) 6 Crash a dev = cl. Buffer (ctx , cl .mem flags.READ WRITE, size=256∗∗3∗4) cl . enqueue write buffer (queue, a dev, numpy.random.rand(256∗∗3).astype(numpy.float32), is blocking =False) 4 OK a dev = cl. Buffer (ctx , cl .mem flags.READ WRITE, size=256∗∗3∗4) cl . enqueue write buffer (queue, a dev, numpy.random.rand(256∗∗3).astype(numpy.float32), is blocking =True) Programs and Kernels prg = cl.Program(context, src ) src: OpenCL device code Derivative of C99 Functions with kernel attribute can be invoked from host prg.build(options="", devices=None) kernel = prg.kernel name kernel(queue, (Gx , Gy , Gz ), (Lx , Ly , Lz ), arg, ..., wait for=None) Program Objects kernel (queue, (Gx,Gy,Gz), (Sx,Sy,Sz), arg , ..., wait for =None) arg may be: None (a NULL pointer) numpy sized scalars: numpy.int64,numpy.float32,... Anything with buffer interface: numpy.ndarray, str Buffer Objects Also: cl.Image, cl.Sampler, cl.LocalMemory Program Objects kernel (queue, (Gx,Gy,Gz), (Sx,Sy,Sz), arg , ..., wait for =None) Explicitly sized scalars: 6 Annoying, error-prone. Better: kernel.set scalar arg dtypes([ numpy.int32, None, numpy.float32]) Use None for non-scalars. OpenCL Object Diagram Outline 1 The OpenCL Runtime 2 Device Language Synchronization Extensions 3 OpenCL implementations OpenCL Device Language OpenCL device language is C99, with these differences: + Index getters + Memory space qualifiers + Vector data types + Many generic (‘overloaded’) math functions including fast native ... varieties. + Synchronization - Recursion - malloc() OpenCL ↔ CUDA: A dictionary OpenCL Grid Work Group Work Item kernel global local private imagend t barrier(LMF) get local id(012) get group id(012) get global id(012) CUDA Grid Block Thread global device shared local texture syncthreads() threadIdx.xyz blockIdx.xyz – (reimplement) Address Space Qualifiers Type private local global constant imagend t Per work item group grid grid grid Access R/W R/W R/W R/O R(/W) Latency 1 or 1000 2 1000 1-1000 1000 Not cached Cached Spatially cached Address Space Qualifiers Type private local global constant imagend t Per work item group grid grid grid Access R/W R/W R/W R/O R(/W) Latency 1 or 1000 2 1000 1-1000 1000 Not cached Cached Spatially cached Address Space Qualifiers Type private local global constant imagend t Per work item group grid grid grid Access R/W R/W R/W R/O R(/W) Latency 1 or 1000 2 1000 1-1000 1000 Not cached Cached Spatially cached Important Don’t “choose one” type of memory. Successful algorithms combine many types’ strengths. CL vector data types floatn vec (n=1,2,3,4,8,16) (also for double and integer types) Components: vec.s012...abcdef (or xyzw) vec.s3120 (Swizzling) vec.s024 = (float3)(1,2,3); (Lvalue, Literals) Usage: Elementwise operations (+,-,sin (generic!),...) floatn vloadn/vstoren(offset, float *) (aligned!) dot/distance Using CPU implementation: One of the sanest ways of using SSE/vector intrinsics!