--- OpenCL/OpenCL.pm 2011/11/17 03:02:25 1.14 +++ OpenCL/OpenCL.pm 2012/05/04 14:56:50 1.73 @@ -29,9 +29,9 @@ 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::Buffers objects (flat -memory areas, think array) and OpenCL::Image objects (think 2d or 3d -array) for bulk data and input and output for kernels. +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. @@ -45,14 +45,32 @@ =head2 HELPFUL RESOURCES -The OpenCL spec used to develop this module (1.2 spec was available, but -no implementation was available to me :). +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 @@ -89,11 +107,11 @@ system. for my $platform (OpenCL::platforms) { - printf "platform: %s\n", $platform->info (OpenCL::PLATFORM_NAME); - printf "extensions: %s\n", $platform->info (OpenCL::PLATFORM_EXTENSIONS); + printf "platform: %s\n", $platform->name; + printf "extensions: %s\n", $platform->extensions; for my $device ($platform->devices) { - printf "+ device: %s\n", $device->info (OpenCL::DEVICE_NAME); - my $ctx = $device->context; + printf "+ device: %s\n", $device->name; + my $ctx = $platform->context (undef, [$device]); # do stuff } } @@ -126,10 +144,10 @@ my $buf = $ctx->buffer_sv (OpenCL::MEM_COPY_HOST_PTR, "helmut"); - $queue->enqueue_read_buffer ($buf, 1, 1, 3, my $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; print "$data\n"; # prints "elm" @@ -137,20 +155,15 @@ 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); - - # build croaks on compile errors, so catch it and print the compile errors - 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 the @@ -164,10 +177,10 @@ $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: printf "%s\n", join ", ", unpack "f*", $data; @@ -176,13 +189,13 @@ 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; @@ -191,14 +204,154 @@ 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 @@ -214,23 +367,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 * Structures are often specified with their components, and returned -as arrayrefs. - -=item * Where possible, one of the pitch values 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 - anywhere a C<$wait_events...> argument -is documented this can be any number of event objects. +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 @@ -260,6 +414,176 @@ 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). + +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 this module 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