--- OpenCL/OpenCL.pm 2011/11/16 00:36:40 1.6 +++ OpenCL/OpenCL.pm 2012/05/05 13:07:19 1.78 @@ -8,43 +8,162 @@ =head1 DESCRIPTION -This is an early release which might be useful, but hasn't seen any testing. +This is an early release which might be useful, but hasn't seen much testing. -=head1 HELPFUL RESOURCES +=head2 OpenCL FROM 10000 FEET HEIGHT -The OpenCL spec used to develop this module (1.2 spec was available, but -no implementation was available to me :). +Here is a high level overview of OpenCL: + +First you need to find one or more OpenCL::Platforms (kind of like +vendors) - usually there is only one. + +Each platform gives you access to a number of OpenCL::Device objects, e.g. +your graphics card. + +From a platform and some device(s), you create an OpenCL::Context, which is +a very central object in OpenCL: Once you have a context you can create +most other objects: + +OpenCL::Program objects, which store source code and, after building for a +specific device ("compiling and linking"), also binary programs. For each +kernel function in a program you can then create an OpenCL::Kernel object +which represents basically a function call with argument values. + +OpenCL::Memory objects of various flavours: OpenCL::Buffer objects (flat +memory areas, think arrays or structs) and OpenCL::Image objects (think 2d +or 3d array) for bulk data and input and output for kernels. + +OpenCL::Sampler objects, which are kind of like texture filter modes in +OpenGL. + +OpenCL::Queue objects - command queues, which allow you to submit memory +reads, writes and copies, as well as kernel calls to your devices. They +also offer a variety of methods to synchronise request execution, for +example with barriers or OpenCL::Event objects. + +OpenCL::Event objects are used to signal when something is complete. + +=head2 HELPFUL RESOURCES + +The OpenCL specs used to develop this module - download these and keept +hema round, they are required reference material: http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf + http://www.khronos.org/registry/cl/specs/opencl-1.2.pdf + http://www.khronos.org/registry/cl/specs/opencl-1.2-extensions.pdf OpenCL manpages: http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/ + http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/ + +If you are into UML class diagrams, the following diagram might help - if +not, it will be mildly confusing (also, the class hierarchy of this module +is much more fine-grained): + + http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/classDiagram.html + +Here's a tutorial from AMD (very AMD-centric, too), not sure how useful it +is, but at least it's free of charge: + + http://developer.amd.com/zones/OpenCLZone/courses/Documents/Introduction_to_OpenCL_Programming%20Training_Guide%20%28201005%29.pdf + +And here's NVIDIA's OpenCL Best Practises Guide: + + http://developer.download.nvidia.com/compute/cuda/3_2/toolkit/docs/OpenCL_Best_Practices_Guide.pdf + +=head1 BASIC WORKFLOW + +To get something done, you basically have to do this once (refer to the +examples below for actual code, this is just a high-level description): + +Find some platform (e.g. the first one) and some device(s) (e.g. the first +device of the platform), and create a context from those. + +Create program objects from your OpenCL source code, then build (compile) +the programs for each device you want to run them on. + +Create kernel objects for all kernels you want to use (surprisingly, these +are not device-specific). + +Then, to execute stuff, you repeat these steps, possibly resuing or +sharing some buffers: + +Create some input and output buffers from your context. Set these as +arguments to your kernel. + +Enqueue buffer writes to initialise your input buffers (when not +initialised at creation time). + +Enqueue the kernel execution. + +Enqueue buffer reads for your output buffer to read results. + +=head1 OPENCL 1.1 VS. OPENCL 1.2 + +This module supports both OpenCL version 1.1 and 1.2, although the OpenCL +1.2 interface hasn't been tested much for lack of availability of an +actual implementation. + +Every function or method in this manual page that interfaces to a +particular OpenCL function has a link to the its C manual page. + +If the link contains a F<1.1>, then this function is an OpenCL 1.1 +function. Most but not all also exist in OpenCL 1.2, and this module +tries to emulate the missing ones for you, when told to do so at +compiletime. You cna check whether a function was removed in OpenCL 1.2 by +replacing the F<1.1> component in the URL by F<1.2>. + +If the link contains a F<1.2>, then this is a OpenCL 1.2-only +function. Even if the module was compiled with OpenCL 1.2 header files +and has an 1.2 OpenCL library, calling such a function on a platform that +doesn't implement 1.2 causes undefined behaviour, usually a crash (But +this is not guaranteed). + +You can find out whether this module was compiled to prefer 1.1 +functionality by ooking at C - if it is true, then +1.1 functions generally are implemented using 1.1 OpenCL functions. If it +is false, then 1.1 functions missing from 1.2 are emulated by calling 1.2 +fucntions. + +This is a somewhat sorry state of affairs, but the Khronos group choose to +make every release of OpenCL source and binary incompatible with previous +releases. =head1 EXAMPLES =head2 Enumerate all devices and get contexts for them. +Best run this once to get a feel for the platforms and devices in your +system. + for my $platform (OpenCL::platforms) { - warn $platform->info (OpenCL::PLATFORM_NAME); - warn $platform->info (OpenCL::PLATFORM_EXTENSIONS); + printf "platform: %s\n", $platform->name; + printf "extensions: %s\n", $platform->extensions; for my $device ($platform->devices) { - warn $device->info (OpenCL::DEVICE_NAME); - my $ctx = $device->context_simple; + printf "+ device: %s\n", $device->name; + my $ctx = $platform->context (undef, [$device]); # do stuff } } =head2 Get a useful context and a command queue. - my $dev = ((OpenCL::platforms)[0]->devices)[0]; - my $ctx = $dev->context_simple; - my $queue = $ctx->command_queue_simple ($dev); +This is a useful boilerplate for any OpenCL program that only wants to use +one device, + + my ($platform) = OpenCL::platforms; # find first platform + my ($dev) = $platform->devices; # find first device of platform + my $ctx = $platform->context (undef, [$dev]); # create context out of those + my $queue = $ctx->queue ($dev); # create a command queue for the device =head2 Print all supported image formats of a context. +Best run this once for your context, to see whats available and how to +gather information. + for my $type (OpenCL::MEM_OBJECT_IMAGE2D, OpenCL::MEM_OBJECT_IMAGE3D) { - say "supported image formats for ", OpenCL::enum2str $type; + print "supported image formats for ", OpenCL::enum2str $type, "\n"; for my $f ($ctx->supported_image_formats (0, $type)) { printf " %-10s %-20s\n", OpenCL::enum2str $f->[0], OpenCL::enum2str $f->[1]; @@ -56,33 +175,30 @@ my $buf = $ctx->buffer_sv (OpenCL::MEM_COPY_HOST_PTR, "helmut"); - $queue->enqueue_read_buffer ($buf, 1, 1, 3, my $data); - warn $data; + $queue->read_buffer ($buf, 1, 1, 3, my $data); + print "$data\n"; - my $ev = $queue->enqueue_read_buffer ($buf, 0, 1, 3, my $data); + my $ev = $queue->read_buffer ($buf, 0, 1, 3, my $data); $ev->wait; - warn $data; + print "$data\n"; # prints "elm" =head2 Create and build a program, then create a kernel out of one of its functions. my $src = ' - __kernel void - squareit (__global float *input, __global float *output) + kernel void + squareit (global float *input, global float *output) { - size_t id = get_global_id (0); + $id = get_global_id (0); output [id] = input [id] * input [id]; } '; - my $prog = $ctx->program_with_source ($src); - - eval { $prog->build ($dev); 1 } - or die $prog->build_info ($dev, OpenCL::PROGRAM_BUILD_LOG); - + my $prog = $ctx->build_program ($src); my $kernel = $prog->kernel ("squareit"); -=head2 Create some input and output float buffers, then call squareit on them. +=head2 Create some input and output float buffers, then call the +'squareit' kernel on them. my $input = $ctx->buffer_sv (OpenCL::MEM_COPY_HOST_PTR, pack "f*", 1, 2, 3, 4.5); my $output = $ctx->buffer (0, OpenCL::SIZEOF_FLOAT * 5); @@ -92,25 +208,25 @@ $kernel->set_buffer (1, $output); # execute it for all 4 numbers - $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); + $queue->nd_range_kernel ($kernel, undef, [4], undef); # enqueue a synchronous read - $queue->enqueue_read_buffer ($output, 1, 0, OpenCL::SIZEOF_FLOAT * 4, my $data); + $queue->read_buffer ($output, 1, 0, OpenCL::SIZEOF_FLOAT * 4, my $data); # print the results: - say join ", ", unpack "f*", $data; + printf "%s\n", join ", ", unpack "f*", $data; =head2 The same enqueue operations as before, but assuming an out-of-order queue, showing off barriers. # execute it for all 4 numbers - $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); + $queue->nd_range_kernel ($kernel, undef, [4], undef); # enqueue a barrier to ensure in-order execution - $queue->enqueue_barrier; + $queue->barrier; # enqueue an async read - $queue->enqueue_read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data); + $queue->read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data); # wait for all requests to finish $queue->finish; @@ -119,20 +235,160 @@ showing off event objects and wait lists. # execute it for all 4 numbers - my $ev = $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); + my $ev = $queue->nd_range_kernel ($kernel, undef, [4], undef); # enqueue an async read - $ev = $queue->enqueue_read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data, $ev); + $ev = $queue->read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data, $ev); # wait for the last event to complete $ev->wait; +=head2 Use the OpenGL module to share a texture between OpenCL and OpenGL and draw some julia +set flight effect. + +This is quite a long example to get you going - you can download it from +L. + + use OpenGL ":all"; + use OpenCL; + + my $S = $ARGV[0] || 256; # window/texture size, smaller is faster + + # open a window and create a gl texture + OpenGL::glpOpenWindow width => $S, height => $S; + my $texid = glGenTextures_p 1; + glBindTexture GL_TEXTURE_2D, $texid; + glTexImage2D_c GL_TEXTURE_2D, 0, GL_RGBA8, $S, $S, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0; + + # find and use the first opencl device that let's us get a shared opengl context + my $platform; + my $dev; + my $ctx; + + for (OpenCL::platforms) { + $platform = $_; + for ($platform->devices) { + $dev = $_; + $ctx = $platform->context ([OpenCL::GLX_DISPLAY_KHR, undef, OpenCL::GL_CONTEXT_KHR, undef], [$dev]) + and last; + } + } + + $ctx + or die "cannot find suitable OpenCL device\n"; + + my $queue = $ctx->queue ($dev); + + # now attach an opencl image2d object to the opengl texture + my $tex = $ctx->gl_texture2d (OpenCL::MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, $texid); + + # now the boring opencl code + my $src = <build_program ($src); + my $kernel = $prog->kernel ("juliatunnel"); + + # program compiled, kernel ready, now draw and loop + + for (my $time; ; ++$time) { + # acquire objects from opengl + $queue->acquire_gl_objects ([$tex]); + + # configure and run our kernel + $kernel->setf ("mf", $tex, $time*2); # mf = memory object, float + $queue->nd_range_kernel ($kernel, undef, [$S, $S], undef); + + # release objects to opengl again + $queue->release_gl_objects ([$tex]); + + # wait + $queue->finish; + + # now draw the texture, the defaults should be all right + glTexParameterf GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST; + + glEnable GL_TEXTURE_2D; + glBegin GL_QUADS; + glTexCoord2f 0, 1; glVertex3i -1, -1, -1; + glTexCoord2f 0, 0; glVertex3i 1, -1, -1; + glTexCoord2f 1, 0; glVertex3i 1, 1, -1; + glTexCoord2f 1, 1; glVertex3i -1, 1, -1; + glEnd; + + glXSwapBuffers; + + select undef, undef, undef, 1/60; + } + +=head2 How to modify the previous example to not rely on GL sharing. + +For those poor souls with only a sucky CPU OpenCL implementation, you +currently have to read the image into some perl scalar, and then modify a +texture or use glDrawPixels or so). + +First, when you don't need gl sharing, you can create the context much simpler: + + $ctx = $platform->context (undef, [$dev]) + +To use a texture, you would modify the above example by creating an +OpenCL::Image manually instead of deriving it from a texture: + + my $tex = $ctx->image2d (OpenCL::MEM_WRITE_ONLY, OpenCL::RGBA, OpenCL::UNORM_INT8, $S, $S); + +And in the draw loop, intead of acquire_gl_objects/release_gl_objects, you +would read the image2d after the kernel has written it: + + $queue->read_image ($tex, 0, 0, 0, 0, $S, $S, 1, 0, 0, my $data); + +And then you would upload the pixel data to the texture (or use glDrawPixels): + + glTexSubImage2D_s GL_TEXTURE_2D, 0, 0, 0, $S, $S, GL_RGBA, GL_UNSIGNED_BYTE, $data; + +The fully modified example can be found at +L. + +=head2 Julia sets look soooo 80ies. + +Then colour them differently, e.g. using orbit traps! Replace the loop and +colour calculation from the previous examples by this: + + float2 dm = (float2)(1.f, 1.f); + + for (int i = 0; i < 25; ++i) + { + z = (float2)(z.x * z.x - z.y * z.y, 2.f * z.x * z.y) + c; + dm = fmin (dm, (float2)(fabs (dot (z, z) - 1.f), fabs (z.x - 1.f))); + } + + float3 colour = (float3)(dm.x * dm.y, dm.x * dm.y, dm.x); + +Also try C<-10.f> instead of C<-1.f>. + =head1 DOCUMENTATION =head2 BASIC CONVENTIONS -This is not a 1:1 C-style translation of OpenCL to Perl - instead I -attempted to make the interface as type-safe as possible and introducing +This is not a one-to-one C-style translation of OpenCL to Perl - instead +I attempted to make the interface as type-safe as possible by introducing object syntax where it makes sense. There are a number of important differences between the OpenCL C API and this module: @@ -142,20 +398,24 @@ to free objects explicitly (C), the release function is called automatically once all Perl references to it go away. -=item * OpenCL uses CamelCase for function names (C), -while this module uses underscores as word separator and often leaves out -prefixes (C<< $platform->info >>). +=item * OpenCL uses CamelCase for function names +(e.g. C, C), while this module +uses underscores as word separator and often leaves out prefixes +(C, C<< $platform->info >>). =item * OpenCL often specifies fixed vector function arguments as short arrays (C), while this module explicitly expects the -components as separate arguments- +components as separate arguments (C<$orig_x, $orig_y, $orig_z>) in +function calls. -=item * Where possible, the row_pitch value is calculated from the perl -scalar length and need not be specified. +=item * Structures are often specified by flattening out their components +as with short vectors, and returned as arrayrefs. =item * When enqueuing commands, the wait list is specified by adding -extra arguments to the function - everywhere a C<$wait_events...> argument -is documented this can be any number of event objects. +extra arguments to the function - anywhere a C<$wait_events...> argument +is documented this can be any number of event objects. As an extsnion +implemented by this module, C values will be ignored in the event +list. =item * When enqueuing commands, if the enqueue method is called in void context, no event is created. In all other contexts an event is returned @@ -167,23 +427,220 @@ =back +=head2 PERL AND OPENCL TYPES + +This handy(?) table lists OpenCL types and their perl, PDL and pack/unpack +format equivalents: + + OpenCL perl PDL pack/unpack + char IV - c + uchar IV byte C + short IV short s + ushort IV ushort S + int IV long? l + uint IV - L + long IV longlong q + ulong IV - Q + float NV float f + half IV ushort S + double NV double d + +=head2 GLX SUPPORT + +Due to the sad state that OpenGL support is in in Perl (mostly the OpenGL +module, which has little to no documentation and has little to no support +for glX), this module, as a special extension, treats context creation +properties C and C +specially: If either or both of these are C, then the OpenCL +module tries to dynamically resolve C and +C, call these functions and use their return values +instead. + +For this to work, the OpenGL library must be loaded, a GLX context must +have been created and be made current, and C must be available and +capable of finding the function via C. + +=head2 EVENT SYSTEM + +OpenCL can generate a number of (potentially) asynchronous events, for +example, after compiling a program, to signal a context-related error or, +perhaps most important, to signal completion of queued jobs (by setting +callbacks on OpenCL::Event objects). + +The OpenCL module converts all these callbacks into events - you can +still register callbacks, but they are not executed when your OpenCL +implementation calls the actual callback, but only later. Therefore, none +of the limitations of OpenCL callbacks apply to the perl implementation: +it is perfectly safe to make blocking operations from event callbacks, and +enqueued operations don't need to be flushed. + +To facilitate this, this module maintains an event queue - each +time an asynchronous event happens, it is queued, and perl will be +interrupted. This is implemented via the L module. In +addition, this module has L support, so it can seamlessly +integrate itself into many event loops. + +Since L is a bit hard to understand, here are some case examples: + +=head3 Don't use callbacks. + +When your program never uses any callbacks, then there will never be any +notifications you need to take care of, and therefore no need to worry +about all this. + +You can achieve a great deal by explicitly waiting for events, or using +barriers and flush calls. In many programs, there is no need at all to +tinker with asynchronous events. + +=head3 Use AnyEvent + +This module automatically registers a watcher that invokes all outstanding +event callbacks when AnyEvent is initialised (and block asynchronous +interruptions). Using this mode of operations is the safest and most +recommended one. + +To use this, simply use AnyEvent and this module normally, make sure you +have an event loop running: + + use Gtk2 -init; + use AnyEvent; + + # initialise AnyEvent, by creating a watcher, or: + AnyEvent::detect; + + my $e = $queue->marker; + $e->cb (sub { + warn "opencl is finished\n"; + }) + + main Gtk2; + +Note that this module will not initialise AnyEvent for you. Before +AnyEvent is initialised, the module will asynchronously interrupt perl +instead. To avoid any surprises, it's best to explicitly initialise +AnyEvent. + +You can temporarily enable asynchronous interruptions (see next paragraph) +by calling C<$OpenCL::INTERRUPT->unblock> and disable them again by +calling C<$OpenCL::INTERRUPT->block>. + +=head3 Let yourself be interrupted at any time + +This mode is the default unless AnyEvent is loaded and initialised. In +this mode, OpenCL asynchronously interrupts a running perl program. The +emphasis is on both I and I here. + +Asynchronously means that perl might execute your callbacks at any +time. For example, in the following code (I), +the C loop following the marker call will be interrupted by the +callback: + + my $e = $queue->marker; + my $flag; + $e->cb (sub { $flag = 1 }); + 1 until $flag; + # $flag is now 1 + +The reason why you shouldn't blindly copy the above code is that +busy waiting is a really really bad thing, and really really bad for +performance. + +While at first this asynchronous business might look exciting, it can be +really hard, because you need to be prepared for the callback code to be +executed at any time, which limits the amount of things the callback code +can do safely. + +This can be mitigated somewhat by using C<< +$OpenCL::INTERRUPT->scope_block >> (see the L +documentation for details). + +The other problem is that your program must be actively I to be +interrupted. When you calculate stuff, your program is running. When you +hang in some C functions or other block execution (by calling C, +C