[ViewVC] Diff of: cvs/OpenCL/OpenCL.pm

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.22 by root, Sun Nov 20 10:25:17 2011 UTC vs.
Revision 1.55 by root, Tue Apr 24 22:45:38 2012 UTC

 Best run this once to get a feel for the platforms and devices in your
 system.
    for my $platform (OpenCL::platforms) {
-      printf "platform: %s\n", $platform->info (OpenCL::PLATFORM_NAME);
+      printf "platform: %s\n", $platform->name;
-      printf "extensions: %s\n", $platform->info (OpenCL::PLATFORM_EXTENSIONS);
+      printf "extensions: %s\n", $platform->extensions;
       for my $device ($platform->devices) {
-         printf "+ device: %s\n", $device->info (OpenCL::DEVICE_NAME);
+         printf "+ device: %s\n", $device->name;
-         my $ctx = $device->context;
+         my $ctx = $platform->context (undef, [$device]);
          # do stuff
       }
    }
 =head2 Get a useful context and a command queue.
 =head2 Create and build a program, then create a kernel out of one of its
 functions.
    my $src = '
-      __kernel void
+      kernel void
-      squareit (__global float *input, __global float *output)
+      squareit (global float *input, global float *output)
       {
         $id = get_global_id (0);
         output [id] = input [id] * input [id];
       }
    ';
-   my $prog = $ctx->program_with_source ($src);
+   my $prog = $ctx->build_program ($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 $kernel = $prog->kernel ("squareit");
 =head2 Create some input and output float buffers, then call the
 'squareit' kernel on them.
    # enqueue an async read
    $ev = $queue->enqueue_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 tunnel effect.
+This is quite a long example to get you going.
+   use OpenGL ":all";
+   use OpenCL;
+   # open a window and create a gl texture
+   OpenGL::glpOpenWindow width => 256, height => 256;
+   my $texid = glGenTextures_p 1;
+   glBindTexture GL_TEXTURE_2D, $texid;
+   glTexImage2D_c GL_TEXTURE_2D, 0, GL_RGBA8, 256, 256, 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 = <<EOF;
+   kernel void
+   juliatunnel (write_only image2d_t img, float time)
+   {
+     float2 p = (float2)(get_global_id (0), get_global_id (1)) / 256.f * 2.f - 1.f;
+     float2 m = (float2)(1.f, p.y) / fabs (p.x);
+     m.x = fabs (fmod (m.x + time * 0.05f, 4.f)) - 2.f;
+     float2 z = m;
+     float2 c = (float2)(sin (time * 0.05005), cos (time * 0.06001));
+     for (int i = 0; i < 25 && dot (z, z) < 4.f; ++i)
+       z = (float2)(z.x * z.x - z.y * z.y, 2.f * z.x * z.y) + c;
+     float3 colour = (float3)(z.x, z.y, z.x * z.y);
+     write_imagef (img, (int2)(get_global_id (0), get_global_id (1)), (float4)(colour * p.x * p.x, 1.));
+   }
+   EOF
+   my $prog = $ctx->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->enqueue_acquire_gl_objects ([$tex]);
+      # configure and run our kernel
+      $kernel->set_image2d (0, $tex);
+      $kernel->set_float   (1, $time);
+      $queue->enqueue_nd_range_kernel ($kernel, undef, [256, 256], undef);
+      # release objects to opengl again
+      $queue->enqueue_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;
+   }
 =head1 DOCUMENTATION
 =head2 BASIC CONVENTIONS
 =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<undef> 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
 by the method.
    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<OpenCL::GLX_DISPLAY_KHR> and C<OpenCL::GL_CONTEXT_KHR>
+specially: If either or both of these are C<undef>, then the OpenCL
+module tries to dynamically resolve C<glXGetCurrentDisplay> and
+C<glXGetCurrentContext>, 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<dlsym> must be available and
+capable of finding the function via C<RTLD_DEFAULT>.
+=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<Async::Interrupt> module. In
+addition, this module has L<AnyEvent> 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->enqueue_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<asynchronously> and I<running> here.
+Asynchronously means that perl might execute your callbacks at any
+time. For example, in the following code (I<THAT YOU SHOULD NOT COPY>),
+the C<until> loop following the marker call will be interrupted by the
+callback:
+   my $e = $queue->enqueue_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<Async::Interrupt>
+documentation for details).
+The other problem is that your program must be actively I<running> 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<sleep>,
+C<select>, running an event loop and so on), your program is waiting, not
+running.
+One way around that would be to attach a read watcher to your event loop,
+listening for events on C<< $OpenCL::INTERRUPT->pipe_fileno >>, using a
+dummy callback (C<sub { }>) to temporarily execute some perl code.
+That is then awfully close to using the built-in AnyEvent support above,
+though, so consider that one instead.
+=head3 Be creative
+OpenCL exports the L<Async::Interrupt> object it uses in the global
+variable C<$OpenCL::INTERRUPT>. You can configure it in any way you like.
+So if you want to feel like a real pro, err, wait, if you feel no risk
+menas no fun, you can experiment by implementing your own mode of
+operations.
+=cut
+package OpenCL;
+use common::sense;
+use Async::Interrupt ();
+our $POLL_FUNC; # set by XS
+BEGIN {
+   our $VERSION = '0.97';
+   require XSLoader;
+   XSLoader::load (__PACKAGE__, $VERSION);
+   @OpenCL::Platform::ISA      =
+   @OpenCL::Device::ISA        =
+   @OpenCL::Context::ISA       =
+   @OpenCL::Queue::ISA         =
+   @OpenCL::Memory::ISA        =
+   @OpenCL::Sampler::ISA       =
+   @OpenCL::Program::ISA       =
+   @OpenCL::Kernel::ISA        =
+   @OpenCL::Event::ISA         = OpenCL::Object::;
+   @OpenCL::Buffer::ISA        =
+   @OpenCL::Image::ISA         = OpenCL::Memory::;
+   @OpenCL::BufferObj::ISA     = OpenCL::Buffer::;
+   @OpenCL::Image2D::ISA       =
+   @OpenCL::Image3D::ISA       =
+   @OpenCL::Image2DArray::ISA  =
+   @OpenCL::Image1D::ISA       =
+   @OpenCL::Image1DArray::ISA  =
+   @OpenCL::Image1DBuffer::ISA = OpenCL::Image::;
+   @OpenCL::UserEvent::ISA     = OpenCL::Event::;
+}
 =head2 THE OpenCL PACKAGE
 =over 4
 =item $int = OpenCL::errno
 Comverts an error value into a human readable string.
 =item $str = OpenCL::enum2str $enum
-Converts most enum values (inof parameter names, image format constants,
+Converts most enum values (of parameter names, image format constants,
 object types, addressing and filter modes, command types etc.) into a
-human readbale string. When confronted with some random integer it can be
+human readable string. When confronted with some random integer it can be
 very helpful to pass it through this function to maybe get some readable
 string out of it.
 =item @platforms = OpenCL::platforms
 Returns all available OpenCL::Platform objects.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformIDs.html>
-=item $ctx = OpenCL::context_from_type $properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $notify = undef
+=item $ctx = OpenCL::context_from_type $properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $notify = $print_stderr
 Tries to create a context from a default device and platform - never worked for me.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html>
 Waits for all events to complete.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html>
+=item OpenCL::poll
+Checks if there are any outstanding events (see L<EVENT SYSTEM>) and
+invokes their callbacks.
+=item $OpenCL::INTERRUPT
+The L<Async::Interrupt> object used to signal asynchronous events (see
+L<EVENT SYSTEM>).
+=cut
+our $INTERRUPT = new Async::Interrupt c_cb => [$POLL_FUNC, 0];
+&_eq_initialise ($INTERRUPT->signal_func);
+=item $OpenCL::WATCHER
+The L<AnyEvent> watcher object used to watch for asynchronous events (see
+L<EVENT SYSTEM>). This variable is C<undef> until L<AnyEvent> has been
+loaded I<and> initialised (e.g. by calling C<AnyEvent::detect>).
+=cut
+our $WATCHER;
+sub _init_anyevent {
+   $INTERRUPT->block;
+   $WATCHER = AE::io ($INTERRUPT->pipe_fileno, 0, sub { $INTERRUPT->handle });
+}
+if (defined $AnyEvent::MODEL) {
+   _init_anyevent;
+} else {
+   push @AnyEvent::post_detect, \&_init_anyevent;
+}
 =back
+=head2 THE OpenCL::Object CLASS
+This is the base class for all objects in the OpenCL module. The only
+method it implements is the C<id> method, which is only useful if you want
+to interface to OpenCL on the C level.
+=over 4
+=item $iv = $obj->id
+OpenCL objects are represented by pointers or integers on the C level. If
+you want to interface to an OpenCL object directly on the C level, then
+you need this value, which is returned by this method. You should use an
+C<IV> type in your code and cast that to the correct type.
+=cut
+sub OpenCL::Object::id {
+   ref $_[0] eq "SCALAR"
+      ? ${ $_[0] }
+      : $_[0][0]
+}
+=back
 =head2 THE OpenCL::Platform CLASS
 =over 4
 =item @devices = $platform->devices ($type = OpenCL::DEVICE_TYPE_ALL)
 Returns a list of matching OpenCL::Device objects.
-=item $ctx = $platform->context_from_type ($properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $notify = undef)
+=item $ctx = $platform->context_from_type ($properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $notify = $print_stderr)
 Tries to create a context. Never worked for me, and you need devices explicitly anyway.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html>
-=item $ctx = $device->context ($properties = undef, @$devices, $notify = undef)
+=item $ctx = $platform->context ($properties, @$devices, $notify = $print_stderr)
 Create a new OpenCL::Context object using the given device object(s)- a
 CL_CONTEXT_PLATFORM property is supplied automatically.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContext.html>
 It's best to avoid this method and use one of the following convenience
 wrappers.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformInfo.html>
+=item $platform->unload_compiler
+Attempts to unload the compiler for this platform, for endless
+profit. Does nothing on OpenCL 1.1.
+L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clUnloadPlatformCompiler.html>
 =for gengetinfo begin platform
 =item $string = $platform->profile
-Calls C<clGetPlatformInfo> with C<CL_PLATFORM_PROFILE> and returns the result(s).
+Calls C<clGetPlatformInfo> with C<CL_PLATFORM_PROFILE> and returns the result.
 =item $string = $platform->version
-Calls C<clGetPlatformInfo> with C<CL_PLATFORM_VERSION> and returns the result(s).
+Calls C<clGetPlatformInfo> with C<CL_PLATFORM_VERSION> and returns the result.
 =item $string = $platform->name
-Calls C<clGetPlatformInfo> with C<CL_PLATFORM_NAME> and returns the result(s).
+Calls C<clGetPlatformInfo> with C<CL_PLATFORM_NAME> and returns the result.
 =item $string = $platform->vendor
-Calls C<clGetPlatformInfo> with C<CL_PLATFORM_VENDOR> and returns the result(s).
+Calls C<clGetPlatformInfo> with C<CL_PLATFORM_VENDOR> and returns the result.
 =item $string = $platform->extensions
-Calls C<clGetPlatformInfo> with C<CL_PLATFORM_EXTENSIONS> and returns the result(s).
+Calls C<clGetPlatformInfo> with C<CL_PLATFORM_EXTENSIONS> and returns the result.
 =for gengetinfo end platform
 =back
 =for gengetinfo begin device
 =item $device_type = $device->type
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_TYPE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_TYPE> and returns the result.
 =item $uint = $device->vendor_id
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_VENDOR_ID> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_VENDOR_ID> and returns the result.
 =item $uint = $device->max_compute_units
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_COMPUTE_UNITS> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_COMPUTE_UNITS> and returns the result.
 =item $uint = $device->max_work_item_dimensions
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS> and returns the result.
 =item $int = $device->max_work_group_size
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_WORK_GROUP_SIZE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_WORK_GROUP_SIZE> and returns the result.
 =item @ints = $device->max_work_item_sizes
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_WORK_ITEM_SIZES> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_WORK_ITEM_SIZES> and returns the result.
 =item $uint = $device->preferred_vector_width_char
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR> and returns the result.
 =item $uint = $device->preferred_vector_width_short
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT> and returns the result.
 =item $uint = $device->preferred_vector_width_int
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT> and returns the result.
 =item $uint = $device->preferred_vector_width_long
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG> and returns the result.
 =item $uint = $device->preferred_vector_width_float
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT> and returns the result.
 =item $uint = $device->preferred_vector_width_double
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE> and returns the result.
 =item $uint = $device->max_clock_frequency
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_CLOCK_FREQUENCY> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_CLOCK_FREQUENCY> and returns the result.
 =item $bitfield = $device->address_bits
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_ADDRESS_BITS> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_ADDRESS_BITS> and returns the result.
 =item $uint = $device->max_read_image_args
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_READ_IMAGE_ARGS> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_READ_IMAGE_ARGS> and returns the result.
 =item $uint = $device->max_write_image_args
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_WRITE_IMAGE_ARGS> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_WRITE_IMAGE_ARGS> and returns the result.
 =item $ulong = $device->max_mem_alloc_size
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_MEM_ALLOC_SIZE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_MEM_ALLOC_SIZE> and returns the result.
 =item $int = $device->image2d_max_width
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE2D_MAX_WIDTH> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE2D_MAX_WIDTH> and returns the result.
 =item $int = $device->image2d_max_height
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE2D_MAX_HEIGHT> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE2D_MAX_HEIGHT> and returns the result.
 =item $int = $device->image3d_max_width
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE3D_MAX_WIDTH> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE3D_MAX_WIDTH> and returns the result.
 =item $int = $device->image3d_max_height
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE3D_MAX_HEIGHT> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE3D_MAX_HEIGHT> and returns the result.
 =item $int = $device->image3d_max_depth
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE3D_MAX_DEPTH> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE3D_MAX_DEPTH> and returns the result.
 =item $uint = $device->image_support
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE_SUPPORT> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE_SUPPORT> and returns the result.
 =item $int = $device->max_parameter_size
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_PARAMETER_SIZE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_PARAMETER_SIZE> and returns the result.
 =item $uint = $device->max_samplers
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_SAMPLERS> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_SAMPLERS> and returns the result.
 =item $uint = $device->mem_base_addr_align
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_MEM_BASE_ADDR_ALIGN> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_MEM_BASE_ADDR_ALIGN> and returns the result.
 =item $uint = $device->min_data_type_align_size
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE> and returns the result.
 =item $device_fp_config = $device->single_fp_config
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_SINGLE_FP_CONFIG> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_SINGLE_FP_CONFIG> and returns the result.
 =item $device_mem_cache_type = $device->global_mem_cache_type
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_GLOBAL_MEM_CACHE_TYPE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_GLOBAL_MEM_CACHE_TYPE> and returns the result.
 =item $uint = $device->global_mem_cacheline_size
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE> and returns the result.
 =item $ulong = $device->global_mem_cache_size
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_GLOBAL_MEM_CACHE_SIZE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_GLOBAL_MEM_CACHE_SIZE> and returns the result.
 =item $ulong = $device->global_mem_size
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_GLOBAL_MEM_SIZE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_GLOBAL_MEM_SIZE> and returns the result.
 =item $ulong = $device->max_constant_buffer_size
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE> and returns the result.
 =item $uint = $device->max_constant_args
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_CONSTANT_ARGS> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_CONSTANT_ARGS> and returns the result.
 =item $device_local_mem_type = $device->local_mem_type
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_LOCAL_MEM_TYPE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_LOCAL_MEM_TYPE> and returns the result.
 =item $ulong = $device->local_mem_size
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_LOCAL_MEM_SIZE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_LOCAL_MEM_SIZE> and returns the result.
 =item $boolean = $device->error_correction_support
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_ERROR_CORRECTION_SUPPORT> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_ERROR_CORRECTION_SUPPORT> and returns the result.
 =item $int = $device->profiling_timer_resolution
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_PROFILING_TIMER_RESOLUTION> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_PROFILING_TIMER_RESOLUTION> and returns the result.
 =item $boolean = $device->endian_little
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_ENDIAN_LITTLE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_ENDIAN_LITTLE> and returns the result.
 =item $boolean = $device->available
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_AVAILABLE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_AVAILABLE> and returns the result.
 =item $boolean = $device->compiler_available
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_COMPILER_AVAILABLE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_COMPILER_AVAILABLE> and returns the result.
 =item $device_exec_capabilities = $device->execution_capabilities
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_EXECUTION_CAPABILITIES> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_EXECUTION_CAPABILITIES> and returns the result.
 =item $command_queue_properties = $device->properties
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_QUEUE_PROPERTIES> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_QUEUE_PROPERTIES> and returns the result.
 =item $ = $device->platform
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_PLATFORM> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_PLATFORM> and returns the result.
 =item $string = $device->name
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_NAME> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_NAME> and returns the result.
 =item $string = $device->vendor
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_VENDOR> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_VENDOR> and returns the result.
 =item $string = $device->driver_version
-Calls C<clGetDeviceInfo> with C<CL_DRIVER_VERSION> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DRIVER_VERSION> and returns the result.
 =item $string = $device->profile
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_PROFILE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_PROFILE> and returns the result.
 =item $string = $device->version
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_VERSION> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_VERSION> and returns the result.
 =item $string = $device->extensions
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_EXTENSIONS> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_EXTENSIONS> and returns the result.
 =item $uint = $device->preferred_vector_width_half
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF> and returns the result.
 =item $uint = $device->native_vector_width_char
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR> and returns the result.
 =item $uint = $device->native_vector_width_short
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT> and returns the result.
 =item $uint = $device->native_vector_width_int
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_INT> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_INT> and returns the result.
 =item $uint = $device->native_vector_width_long
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG> and returns the result.
 =item $uint = $device->native_vector_width_float
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT> and returns the result.
 =item $uint = $device->native_vector_width_double
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE> and returns the result.
 =item $uint = $device->native_vector_width_half
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF> and returns the result.
 =item $device_fp_config = $device->double_fp_config
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_DOUBLE_FP_CONFIG> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_DOUBLE_FP_CONFIG> and returns the result.
 =item $device_fp_config = $device->half_fp_config
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_HALF_FP_CONFIG> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_HALF_FP_CONFIG> and returns the result.
 =item $boolean = $device->host_unified_memory
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_HOST_UNIFIED_MEMORY> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_HOST_UNIFIED_MEMORY> and returns the result.
 =item $device = $device->parent_device_ext
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARENT_DEVICE_EXT> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARENT_DEVICE_EXT> and returns the result.
 =item @device_partition_property_exts = $device->partition_types_ext
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARTITION_TYPES_EXT> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARTITION_TYPES_EXT> and returns the result.
 =item @device_partition_property_exts = $device->affinity_domains_ext
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_AFFINITY_DOMAINS_EXT> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_AFFINITY_DOMAINS_EXT> and returns the result.
 =item $uint = $device->reference_count_ext
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_REFERENCE_COUNT_EXT > and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_REFERENCE_COUNT_EXT> and returns the result.
 =item @device_partition_property_exts = $device->partition_style_ext
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARTITION_STYLE_EXT> and returns the result(s).
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARTITION_STYLE_EXT> and returns the result.
 =for gengetinfo end device
 =back
 =head2 THE OpenCL::Context CLASS
 =over 4
+=item $prog = $ctx->build_program ($program, $options = "")
+This convenience function tries to build the program on all devices in
+the context. If the build fails, then the function will C<croak> with the
+build log. Otherwise ti returns the program object.
+The C<$program> can either be a C<OpenCL::Program> object or a string
+containing the program. In the latter case, a program objetc will be
+created automatically.
+=cut
+sub OpenCL::Context::build_program {
+   my ($self, $prog, $options) = @_;
+   require Carp;
+   $prog = $self->program_with_source ($prog)
+      unless ref $prog;
+   # we build separately per device so we instantly know which one failed
+   for my $dev ($self->devices) {
+      eval { $prog->build ([$dev], $options); 1 }
+         or Carp::croak ("Building OpenCL program for device '" . $dev->name . "' failed:\n"
+                         . $prog->build_log ($dev));
+   }
+   $prog
+}
 =item $queue = $ctx->queue ($device, $properties)
 Create a new OpenCL::Queue object from the context and the given device.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateCommandQueue.html>
+Example: create an out-of-order queue.
+   $queue = $ctx->queue ($device, OpenCL::QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE);
 =item $ev = $ctx->user_event
 Creates a new OpenCL::UserEvent object.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateUserEvent.html>
 =item $buf = $ctx->buffer ($flags, $len)
-Creates a new OpenCL::Buffer object with the given flags and octet-size.
+Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object with the
+given flags and octet-size.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateBuffer.html>
 =item $buf = $ctx->buffer_sv ($flags, $data)
-Creates a new OpenCL::Buffer object and initialise it with the given data values.
+Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object and
+initialise it with the given data values.
+=item $img = $ctx->image ($self, $flags, $channel_order, $channel_type, $type, $width, $height, $depth, $array_size = 0, $row_pitch = 0, $slice_pitch = 0, $num_mip_level = 0, $num_samples = 0, $*data = &PL_sv_undef)
+Creates a new OpenCL::Image object and optionally initialises it with
+the given data values.
+L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateImage.html>
 =item $img = $ctx->image2d ($flags, $channel_order, $channel_type, $width, $height, $row_pitch = 0, $data = undef)
-Creates a new OpenCL::Image2D object and optionally initialises it with the given data values.
+Creates a new OpenCL::Image2D object and optionally initialises it with
+the given data values.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateImage2D.html>
 =item $img = $ctx->image3d ($flags, $channel_order, $channel_type, $width, $height, $depth, $row_pitch = 0, $slice_pitch = 0, $data = undef)
-Creates a new OpenCL::Image3D object and optionally initialises it with the given data values.
+Creates a new OpenCL::Image3D object and optionally initialises it with
+the given data values.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateImage3D.html>
+=item $buffer = $ctx->gl_buffer ($flags, $bufobj)
+Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object that refers to the given
+OpenGL buffer object.
+http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLBuffer.html
+=item $img = $ctx->gl_texture ($flags, $target, $miplevel, $texture)
+Creates a new OpenCL::Image object that refers to the given OpenGL
+texture object or buffer.
+http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateFromGLTexture.html
+=item $img = $ctx->gl_texture2d ($flags, $target, $miplevel, $texture)
+Creates a new OpenCL::Image2D object that refers to the given OpenGL
+2D texture object.
+http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture2D.html
+=item $img = $ctx->gl_texture3d ($flags, $target, $miplevel, $texture)
+Creates a new OpenCL::Image3D object that refers to the given OpenGL
+3D texture object.
+http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture3D.html
+=item $ctx->gl_renderbuffer ($flags, $renderbuffer)
+Creates a new OpenCL::Image2D object that refers to the given OpenGL
+render buffer.
+http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLRenderbuffer.html
 =item @formats = $ctx->supported_image_formats ($flags, $image_type)
 Returns a list of matching image formats - each format is an arrayref with
 two values, $channel_order and $channel_type, in it.
 =for gengetinfo begin context
 =item $uint = $context->reference_count
-Calls C<clGetContextInfo> with C<CL_CONTEXT_REFERENCE_COUNT> and returns the result(s).
+Calls C<clGetContextInfo> with C<CL_CONTEXT_REFERENCE_COUNT> and returns the result.
 =item @devices = $context->devices
-Calls C<clGetContextInfo> with C<CL_CONTEXT_DEVICES> and returns the result(s).
+Calls C<clGetContextInfo> with C<CL_CONTEXT_DEVICES> and returns the result.
 =item @property_ints = $context->properties
-Calls C<clGetContextInfo> with C<CL_CONTEXT_PROPERTIES> and returns the result(s).
+Calls C<clGetContextInfo> with C<CL_CONTEXT_PROPERTIES> and returns the result.
 =item $uint = $context->num_devices
-Calls C<clGetContextInfo> with C<CL_CONTEXT_NUM_DEVICES> and returns the result(s).
+Calls C<clGetContextInfo> with C<CL_CONTEXT_NUM_DEVICES> and returns the result.
 =for gengetinfo end context
 =back
 for completion, unless the method is called in void context, in which case
 no event object is created.
 They also allow you to specify any number of other event objects that this
 request has to wait for before it starts executing, by simply passing the
-event objects as extra parameters to the enqueue methods.
+event objects as extra parameters to the enqueue methods. To simplify
+program design, this module ignores any C<undef> values in the list of
+events. This makes it possible to code operations such as this, without
+having to put a valid event object into C<$event> first:
+   $event = $queue->enqueue_xxx (..., $event);
 Queues execute in-order by default, without any parallelism, so in most
 cases (i.e. you use only one queue) it's not necessary to wait for or
-create event objects.
+create event objects, althoguh an our of order queue is often a bit
+faster.
 =over 4
 =item $ev = $queue->enqueue_read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...)
 =item $ev = $queue->enqueue_copy_buffer ($src, $dst, $src_offset, $dst_offset, $len, $wait_events...)
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBuffer.html>
+=item $ev = $queue->enqueue_read_buffer_rect (OpenCL::Memory buf, cl_bool blocking, $buf_x, $buf_y, $buf_z, $host_x, $host_y, $host_z, $width, $height, $depth, $buf_row_pitch, $buf_slice_pitch, $host_row_pitch, $host_slice_pitch, $data, $wait_events...)
+http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBufferRect.html
+=item $ev = $queue->enqueue_write_buffer_rect (OpenCL::Memory buf, cl_bool blocking, $buf_x, $buf_y, $buf_z, $host_x, $host_y, $host_z, $width, $height, $depth, $buf_row_pitch, $buf_slice_pitch, $host_row_pitch, $host_slice_pitch, $data, $wait_events...)
+http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBufferRect.html
 =item $ev = $queue->enqueue_read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
+L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferRect.html>
+=item $ev = $queue->enqueue_copy_buffer_to_image ($src_buffer, $dst_image, $src_offset, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $wait_events...)
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadImage.html>
 =item $ev = $queue->enqueue_write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteImage.html>
+=item $ev = $queue->enqueue_copy_image ($src_image, $dst_image, $src_x, $src_y, $src_z, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $wait_events...)
+L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html>
+=item $ev = $queue->enqueue_copy_image_to_buffer ($src_image, $dst_image, $src_x, $src_y, $src_z, $width, $height, $depth, $dst_offset, $wait_events...)
+L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html>
 =item $ev = $queue->enqueue_copy_buffer_rect ($src, $dst, $src_x, $src_y, $src_z, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $src_row_pitch, $src_slice_pitch, $dst_row_pitch, $dst_slice_pitch, $wait_event...)
 Yeah.
-L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferRect.html>
-=item $ev = $queue->enqueue_copy_buffer_to_image ($src_buffer, $dst_image, $src_offset, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $wait_events...)
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>.
-=item $ev = $queue->enqueue_copy_image ($src_image, $dst_image, $src_x, $src_y, $src_z, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $wait_events...)
+=item $ev = $queue->enqueue_fill_buffer ($mem, $pattern, $offset, $size, ...)
+Fills the given buffer object with repeated applications of C<$pattern>,
+starting at C<$offset> for C<$size> octets.
+L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillBuffer.html>
+=item $ev = $queue->enqueue_fill_image ($img, $r, $g, $b, $a, $x, $y, $z, $width, $height, $depth, ...)
+Fills the given image area with the given rgba colour components. The
+components are normally floating point values between C<0> and C<1>,
+except when the image channel data type is a signe dor unsigned
+unnormalised format, in which case the range is determined by the format.
-L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html>
+L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillImage.html>
-=item $ev = $queue->enqueue_copy_image_to_buffer ($src_image, $dst_image, $src_x, $src_y, $src_z, $width, $height, $depth, $dst_offset, $wait_events...)
-L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html>
 =item $ev = $queue->enqueue_task ($kernel, $wait_events...)
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html>
 reference to an array of local work sizes, with the same number of
 elements as @$global_work_size.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html>
-=item $ev = $queue->enqueue_marker
+=item $ev = $queue->enqueue_acquire_gl_objects ([object, ...], $wait_events...)
+Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired
+for subsequent OpenCL usage.
-L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMarker.html>
+L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueAcquireGLObjects.html>
+=item $ev = $queue->enqueue_release_gl_objects ([object, ...], $wait_events...)
+Enqueues a list (an array-ref of OpenCL::Memory objects) to be released
+for subsequent OpenGL usage.
+L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReleaseGLObjects.html>
 =item $ev = $queue->enqueue_wait_for_events ($wait_events...)
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html>
-=item $queue->enqueue_barrier
+=item $ev = $queue->enqueue_marker ($wait_events...)
+L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMarkerWithWaitList.html>
+=item $ev = $queue->enqueue_barrier ($wait_events...)
-L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueBarrier.html>
+L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueBarrierWithWaitList.html>
 =item $queue->flush
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clFlush.html>
 =for gengetinfo begin command_queue
 =item $ctx = $command_queue->context
-Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_CONTEXT> and returns the result(s).
+Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_CONTEXT> and returns the result.
 =item $device = $command_queue->device
-Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_DEVICE> and returns the result(s).
+Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_DEVICE> and returns the result.
 =item $uint = $command_queue->reference_count
-Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_REFERENCE_COUNT> and returns the result(s).
+Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_REFERENCE_COUNT> and returns the result.
 =item $command_queue_properties = $command_queue->properties
-Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_PROPERTIES> and returns the result(s).
+Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_PROPERTIES> and returns the result.
 =for gengetinfo end command_queue
 =back
 =for gengetinfo begin mem
 =item $mem_object_type = $mem->type
-Calls C<clGetMemObjectInfo> with C<CL_MEM_TYPE> and returns the result(s).
+Calls C<clGetMemObjectInfo> with C<CL_MEM_TYPE> and returns the result.
 =item $mem_flags = $mem->flags
-Calls C<clGetMemObjectInfo> with C<CL_MEM_FLAGS> and returns the result(s).
+Calls C<clGetMemObjectInfo> with C<CL_MEM_FLAGS> and returns the result.
 =item $int = $mem->size
-Calls C<clGetMemObjectInfo> with C<CL_MEM_SIZE> and returns the result(s).
+Calls C<clGetMemObjectInfo> with C<CL_MEM_SIZE> and returns the result.
 =item $ptr_value = $mem->host_ptr
-Calls C<clGetMemObjectInfo> with C<CL_MEM_HOST_PTR> and returns the result(s).
+Calls C<clGetMemObjectInfo> with C<CL_MEM_HOST_PTR> and returns the result.
 =item $uint = $mem->map_count
-Calls C<clGetMemObjectInfo> with C<CL_MEM_MAP_COUNT> and returns the result(s).
+Calls C<clGetMemObjectInfo> with C<CL_MEM_MAP_COUNT> and returns the result.
 =item $uint = $mem->reference_count
-Calls C<clGetMemObjectInfo> with C<CL_MEM_REFERENCE_COUNT> and returns the result(s).
+Calls C<clGetMemObjectInfo> with C<CL_MEM_REFERENCE_COUNT> and returns the result.
 =item $ctx = $mem->context
-Calls C<clGetMemObjectInfo> with C<CL_MEM_CONTEXT> and returns the result(s).
+Calls C<clGetMemObjectInfo> with C<CL_MEM_CONTEXT> and returns the result.
 =item $mem = $mem->associated_memobject
-Calls C<clGetMemObjectInfo> with C<CL_MEM_ASSOCIATED_MEMOBJECT> and returns the result(s).
+Calls C<clGetMemObjectInfo> with C<CL_MEM_ASSOCIATED_MEMOBJECT> and returns the result.
 =item $int = $mem->offset
-Calls C<clGetMemObjectInfo> with C<CL_MEM_OFFSET> and returns the result(s).
+Calls C<clGetMemObjectInfo> with C<CL_MEM_OFFSET> and returns the result.
 =for gengetinfo end mem
+=item ($type, $name) = $mem->gl_object_info
+Returns the OpenGL object type (e.g. OpenCL::GL_OBJECT_TEXTURE2D) and the
+object "name" (e.g. the texture name) used to create this memory object.
+L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetGLObjectInfo.html>
 =back
+=head2 THE OpenCL::Buffer CLASS
+This is a subclass of OpenCL::Memory, and the superclass of
+OpenCL::BufferObj. Its purpose is simply to distinguish between buffers
+and sub-buffers.
+=head2 THE OpenCL::BufferObj CLASS
+This is a subclass of OpenCL::Buffer and thus OpenCL::Memory. It exists
+because one cna create sub buffers of OpenLC::BufferObj objects, but not
+sub buffers from these sub buffers.
+=over 4
+=item $subbuf = $buf_obj->sub_buffer_region ($flags, $origin, $size)
+Creates an OpenCL::Buffer objects from this buffer and returns it. The
+C<buffer_create_type> is assumed to be C<CL_BUFFER_CREATE_TYPE_REGION>.
+L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateSubBuffer.html>
+=back
 =head2 THE OpenCL::Image CLASS
-This is the superclass of all image objects - OpenCL::Image2D and OpenCL::Image3D.
+This is the superclass of all image objects - OpenCL::Image1D,
+OpenCL::Image1DArray, OpenCL::Image1DBuffer, OpenCL::Image2D,
+OpenCL::Image2DArray and OpenCL::Image3D.
 =over 4
-=item $packed_value = $ev->image_info ($name)
+=item $packed_value = $image->image_info ($name)
 See C<< $platform->info >> for details.
 The reason this method is not called C<info> is that there already is an
 C<< ->info >> method inherited from C<OpenCL::Memory>.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetImageInfo.html>
+=item ($channel_order, $channel_data_type) = $image->format
+Returns the channel order and type used to create the image by calling
+C<clGetImageInfo> with C<CL_IMAGE_FORMAT>.
 =for gengetinfo begin image
 =item $int = $image->element_size
-Calls C<clGetImageInfo> with C<CL_IMAGE_ELEMENT_SIZE> and returns the result(s).
+Calls C<clGetImageInfo> with C<CL_IMAGE_ELEMENT_SIZE> and returns the result.
 =item $int = $image->row_pitch
-Calls C<clGetImageInfo> with C<CL_IMAGE_ROW_PITCH> and returns the result(s).
+Calls C<clGetImageInfo> with C<CL_IMAGE_ROW_PITCH> and returns the result.
 =item $int = $image->slice_pitch
-Calls C<clGetImageInfo> with C<CL_IMAGE_SLICE_PITCH> and returns the result(s).
+Calls C<clGetImageInfo> with C<CL_IMAGE_SLICE_PITCH> and returns the result.
 =item $int = $image->width
-Calls C<clGetImageInfo> with C<CL_IMAGE_WIDTH> and returns the result(s).
+Calls C<clGetImageInfo> with C<CL_IMAGE_WIDTH> and returns the result.
 =item $int = $image->height
-Calls C<clGetImageInfo> with C<CL_IMAGE_HEIGHT> and returns the result(s).
+Calls C<clGetImageInfo> with C<CL_IMAGE_HEIGHT> and returns the result.
 =item $int = $image->depth
-Calls C<clGetImageInfo> with C<CL_IMAGE_DEPTH> and returns the result(s).
+Calls C<clGetImageInfo> with C<CL_IMAGE_DEPTH> and returns the result.
 =for gengetinfo end image
+=for gengetinfo begin gl_texture
+=item $GLenum = $gl_texture->target
+Calls C<clGetGLTextureInfo> with C<CL_GL_TEXTURE_TARGET> and returns the result.
+=item $GLint = $gl_texture->gl_mipmap_level
+Calls C<clGetGLTextureInfo> with C<CL_GL_MIPMAP_LEVEL> and returns the result.
+=for gengetinfo end gl_texture
 =back
 =head2 THE OpenCL::Sampler CLASS
 =over 4
 =for gengetinfo begin sampler
 =item $uint = $sampler->reference_count
-Calls C<clGetSamplerInfo> with C<CL_SAMPLER_REFERENCE_COUNT> and returns the result(s).
+Calls C<clGetSamplerInfo> with C<CL_SAMPLER_REFERENCE_COUNT> and returns the result.
 =item $ctx = $sampler->context
-Calls C<clGetSamplerInfo> with C<CL_SAMPLER_CONTEXT> and returns the result(s).
+Calls C<clGetSamplerInfo> with C<CL_SAMPLER_CONTEXT> and returns the result.
 =item $addressing_mode = $sampler->normalized_coords
-Calls C<clGetSamplerInfo> with C<CL_SAMPLER_NORMALIZED_COORDS> and returns the result(s).
+Calls C<clGetSamplerInfo> with C<CL_SAMPLER_NORMALIZED_COORDS> and returns the result.
 =item $filter_mode = $sampler->addressing_mode
-Calls C<clGetSamplerInfo> with C<CL_SAMPLER_ADDRESSING_MODE> and returns the result(s).
+Calls C<clGetSamplerInfo> with C<CL_SAMPLER_ADDRESSING_MODE> and returns the result.
 =item $boolean = $sampler->filter_mode
-Calls C<clGetSamplerInfo> with C<CL_SAMPLER_FILTER_MODE> and returns the result(s).
+Calls C<clGetSamplerInfo> with C<CL_SAMPLER_FILTER_MODE> and returns the result.
 =for gengetinfo end sampler
 =back
 =head2 THE OpenCL::Program CLASS
 =over 4
-=item $program->build ($device, $options = "")
+=item $program->build (\@devices = undef, $options = "", $cb->($program) = undef)
-Tries to build the program with the givne options.
+Tries to build the program with the given options. See also the
+C<$ctx->build> convenience function.
+If a callback is specified, then it will be called when compilation is
+finished. Note that many OpenCL implementations block your program while
+compiling whether you use a callback or not. See C<build_async> if you
+want to make sure the build is done in the background.
+Note that some OpenCL implementations atc up badly, and don't call the
+callback in some error cases (but call it in others). This implementation
+assumes the callback will always be called, and leaks memory if this is
+not so. So best make sure you don't pass in invalid values.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clBuildProgram.html>
+=item $program->build_async (\@devices = undef, $options = "", $cb->($program) = undef)
+Similar to C<< ->build >>, except it starts a thread, and never fails (you
+need to check the compilation status form the callback, or by polling).
 =item $packed_value = $program->build_info ($device, $name)
 Similar to C<< $platform->info >>, but returns build info for a previous
 build attempt for the given device.
 Creates an OpenCL::Kernel object out of the named C<__kernel> function in
 the program.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernel.html>
+=item @kernels = $program->kernels_in_program
+Returns all kernels successfully compiled for all devices in program.
+http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernelsInProgram.html
 =for gengetinfo begin program_build
 =item $build_status = $program->build_status ($device)
-Calls C<clGetProgramBuildInfo> with C<CL_PROGRAM_BUILD_STATUS> and returns the result(s).
+Calls C<clGetProgramBuildInfo> with C<CL_PROGRAM_BUILD_STATUS> and returns the result.
 =item $string = $program->build_options ($device)
-Calls C<clGetProgramBuildInfo> with C<CL_PROGRAM_BUILD_OPTIONS> and returns the result(s).
+Calls C<clGetProgramBuildInfo> with C<CL_PROGRAM_BUILD_OPTIONS> and returns the result.
 =item $string = $program->build_log ($device)
-Calls C<clGetProgramBuildInfo> with C<CL_PROGRAM_BUILD_LOG> and returns the result(s).
+Calls C<clGetProgramBuildInfo> with C<CL_PROGRAM_BUILD_LOG> and returns the result.
 =for gengetinfo end program_build
 =item $packed_value = $program->info ($name)
 =for gengetinfo begin program
 =item $uint = $program->reference_count
-Calls C<clGetProgramInfo> with C<CL_PROGRAM_REFERENCE_COUNT> and returns the result(s).
+Calls C<clGetProgramInfo> with C<CL_PROGRAM_REFERENCE_COUNT> and returns the result.
 =item $ctx = $program->context
-Calls C<clGetProgramInfo> with C<CL_PROGRAM_CONTEXT> and returns the result(s).
+Calls C<clGetProgramInfo> with C<CL_PROGRAM_CONTEXT> and returns the result.
 =item $uint = $program->num_devices
-Calls C<clGetProgramInfo> with C<CL_PROGRAM_NUM_DEVICES> and returns the result(s).
+Calls C<clGetProgramInfo> with C<CL_PROGRAM_NUM_DEVICES> and returns the result.
 =item @devices = $program->devices
-Calls C<clGetProgramInfo> with C<CL_PROGRAM_DEVICES> and returns the result(s).
+Calls C<clGetProgramInfo> with C<CL_PROGRAM_DEVICES> and returns the result.
 =item $string = $program->source
-Calls C<clGetProgramInfo> with C<CL_PROGRAM_SOURCE> and returns the result(s).
+Calls C<clGetProgramInfo> with C<CL_PROGRAM_SOURCE> and returns the result.
 =item @ints = $program->binary_sizes
-Calls C<clGetProgramInfo> with C<CL_PROGRAM_BINARY_SIZES> and returns the result(s).
+Calls C<clGetProgramInfo> with C<CL_PROGRAM_BINARY_SIZES> and returns the result.
 =for gengetinfo end program
+=item @blobs = $program->binaries
+Returns a string for the compiled binary for every device associated with
+the program, empty strings indicate missing programs, and an empty result
+means no program binaries are available.
+These "binaries" are often, in fact, informative low-level assembly
+sources.
+L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetProgramInfo.html>
 =back
 =head2 THE OpenCL::Kernel CLASS
 =over 4
 =for gengetinfo begin kernel
 =item $string = $kernel->function_name
-Calls C<clGetKernelInfo> with C<CL_KERNEL_FUNCTION_NAME> and returns the result(s).
+Calls C<clGetKernelInfo> with C<CL_KERNEL_FUNCTION_NAME> and returns the result.
 =item $uint = $kernel->num_args
-Calls C<clGetKernelInfo> with C<CL_KERNEL_NUM_ARGS> and returns the result(s).
+Calls C<clGetKernelInfo> with C<CL_KERNEL_NUM_ARGS> and returns the result.
 =item $uint = $kernel->reference_count
-Calls C<clGetKernelInfo> with C<CL_KERNEL_REFERENCE_COUNT> and returns the result(s).
+Calls C<clGetKernelInfo> with C<CL_KERNEL_REFERENCE_COUNT> and returns the result.
 =item $ctx = $kernel->context
-Calls C<clGetKernelInfo> with C<CL_KERNEL_CONTEXT> and returns the result(s).
+Calls C<clGetKernelInfo> with C<CL_KERNEL_CONTEXT> and returns the result.
 =item $program = $kernel->program
-Calls C<clGetKernelInfo> with C<CL_KERNEL_PROGRAM> and returns the result(s).
+Calls C<clGetKernelInfo> with C<CL_KERNEL_PROGRAM> and returns the result.
 =for gengetinfo end kernel
 =item $packed_value = $kernel->work_group_info ($device, $name)
 =for gengetinfo begin kernel_work_group
 =item $int = $kernel->work_group_size ($device)
-Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_WORK_GROUP_SIZE> and returns the result(s).
+Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_WORK_GROUP_SIZE> and returns the result.
 =item @ints = $kernel->compile_work_group_size ($device)
-Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_COMPILE_WORK_GROUP_SIZE> and returns the result(s).
+Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_COMPILE_WORK_GROUP_SIZE> and returns the result.
 =item $ulong = $kernel->local_mem_size ($device)
-Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_LOCAL_MEM_SIZE> and returns the result(s).
+Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_LOCAL_MEM_SIZE> and returns the result.
 =item $int = $kernel->preferred_work_group_size_multiple ($device)
-Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE> and returns the result(s).
+Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE> and returns the result.
 =item $ulong = $kernel->private_mem_size ($device)
-Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PRIVATE_MEM_SIZE> and returns the result(s).
+Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PRIVATE_MEM_SIZE> and returns the result.
 =for gengetinfo end kernel_work_group
 =item $kernel->set_TYPE ($index, $value)
 This is a family of methods to set the kernel argument with the number C<$index> to the give C<$value>.
 TYPE is one of C<char>, C<uchar>, C<short>, C<ushort>, C<int>, C<uint>,
 C<long>, C<ulong>, C<half>, C<float>, C<double>, C<memory>, C<buffer>,
-C<image2d>, C<image3d>, C<sampler> or C<event>.
+C<image2d>, C<image3d>, C<sampler>, C<local> or C<event>.
 Chars and integers (including the half type) are specified as integers,
 float and double as floating point values, memory/buffer/image2d/image3d
-must be an object of that type or C<undef>, and sampler and event must be
+must be an object of that type or C<undef>, local-memory arguments are
-objects of that type.
+set by specifying the size, and sampler and event must be objects of that
+type.
+Setting an argument for a kernel does NOT keep a reference to the object -
+for example, if you set an argument to some image object, free the image,
+and call the kernel, you will run into undefined behaviour.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html>
 =back
 Waits for the event to complete.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html>
+=item $ev->cb ($exec_callback_type, $callback->($event, $event_command_exec_status))
+Adds a callback to the callback stack for the given event type. There is
+no way to remove a callback again.
+L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetEventCallback.html>
 =item $packed_value = $ev->info ($name)
 See C<< $platform->info >> for details.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetEventInfo.html>
 =for gengetinfo begin event
 =item $queue = $event->command_queue
-Calls C<clGetEventInfo> with C<CL_EVENT_COMMAND_QUEUE> and returns the result(s).
+Calls C<clGetEventInfo> with C<CL_EVENT_COMMAND_QUEUE> and returns the result.
 =item $command_type = $event->command_type
-Calls C<clGetEventInfo> with C<CL_EVENT_COMMAND_TYPE> and returns the result(s).
+Calls C<clGetEventInfo> with C<CL_EVENT_COMMAND_TYPE> and returns the result.
 =item $uint = $event->reference_count
-Calls C<clGetEventInfo> with C<CL_EVENT_REFERENCE_COUNT> and returns the result(s).
+Calls C<clGetEventInfo> with C<CL_EVENT_REFERENCE_COUNT> and returns the result.
 =item $uint = $event->command_execution_status
-Calls C<clGetEventInfo> with C<CL_EVENT_COMMAND_EXECUTION_STATUS> and returns the result(s).
+Calls C<clGetEventInfo> with C<CL_EVENT_COMMAND_EXECUTION_STATUS> and returns the result.
 =item $ctx = $event->context
-Calls C<clGetEventInfo> with C<CL_EVENT_CONTEXT> and returns the result(s).
+Calls C<clGetEventInfo> with C<CL_EVENT_CONTEXT> and returns the result.
 =for gengetinfo end event
 =item $packed_value = $ev->profiling_info ($name)
 =for gengetinfo begin profiling
 =item $ulong = $event->profiling_command_queued
-Calls C<clGetEventProfilingInfo> with C<CL_PROFILING_COMMAND_QUEUED> and returns the result(s).
+Calls C<clGetEventProfilingInfo> with C<CL_PROFILING_COMMAND_QUEUED> and returns the result.
 =item $ulong = $event->profiling_command_submit
-Calls C<clGetEventProfilingInfo> with C<CL_PROFILING_COMMAND_SUBMIT> and returns the result(s).
+Calls C<clGetEventProfilingInfo> with C<CL_PROFILING_COMMAND_SUBMIT> and returns the result.
 =item $ulong = $event->profiling_command_start
-Calls C<clGetEventProfilingInfo> with C<CL_PROFILING_COMMAND_START> and returns the result(s).
+Calls C<clGetEventProfilingInfo> with C<CL_PROFILING_COMMAND_START> and returns the result.
 =item $ulong = $event->profiling_command_end
-Calls C<clGetEventProfilingInfo> with C<CL_PROFILING_COMMAND_END> and returns the result(s).
+Calls C<clGetEventProfilingInfo> with C<CL_PROFILING_COMMAND_END> and returns the result.
 =for gengetinfo end profiling
 =back
 =over 4
 =item $ev->set_status ($execution_status)
+Sets the execution status of the user event. Can only be called once,
+either with OpenCL::COMPLETE or a negative number as status.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetUserEventStatus.html>
 =back
 =cut
-package OpenCL;
-use common::sense;
-BEGIN {
-   our $VERSION = '0.55';
-   require XSLoader;
-   XSLoader::load (__PACKAGE__, $VERSION);
-   @OpenCL::Buffer::ISA    =
-   @OpenCL::Image::ISA     = OpenCL::Memory::;
-   @OpenCL::Image2D::ISA   =
-   @OpenCL::Image3D::ISA   = OpenCL::Image::;
-   @OpenCL::UserEvent::ISA = OpenCL::Event::;
-}
 1;
 =head1 AUTHOR

Diff Legend

-–
+Removed lines
-+
+Added lines
-<
+Changed lines
->
+Changed lines

Comparing OpenCL/OpenCL.pm (file contents): Revision 1.22 by root, Sun Nov 20 10:25:17 2011 UTC vs. Revision 1.55 by root, Tue Apr 24 22:45:38 2012 UTC

Diff Legend

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.22 by root, Sun Nov 20 10:25:17 2011 UTC vs.
Revision 1.55 by root, Tue Apr 24 22:45:38 2012 UTC