--- OpenCL/OpenCL.pm	2012/04/19 19:36:24	1.38
+++ OpenCL/OpenCL.pm	2012/04/24 23:53:12	1.56
@@ -159,12 +159,7 @@
       }
    ';
 
-   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_log;
-
+   my $prog = $ctx->build_program ($src);
    my $kernel = $prog->kernel ("squareit");
 
 =head2 Create some input and output float buffers, then call the
@@ -262,17 +257,15 @@
      float2 z = m;
      float2 c = (float2)(sin (time * 0.05005), cos (time * 0.06001));
 
-     for (int i = 0; i < 100 && dot (z, z) < 4.f; ++i)
+     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->program_with_source ($src);
-   eval { $prog->build ($dev); 1 }
-      or die $prog->build_log ($dev);
 
+   my $prog = $ctx->build_program ($src);
    my $kernel = $prog->kernel ("juliatunnel");
 
    # program compiled, kernel ready, now draw and loop
@@ -290,7 +283,7 @@
       $queue->enqueue_release_gl_objects ([$tex]);
 
       # wait
-      $queue->flush;
+      $queue->finish;
 
       # now draw the texture, the defaults should be all right
       glTexParameterf GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST;
@@ -338,7 +331,9 @@
 
 =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
@@ -383,6 +378,155 @@
 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
@@ -410,18 +554,87 @@
 
 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, $callback->($err, $pvt) = $print_stderr
 
-Tries to create a context from a default device and platform - never worked for me.
+Tries to create a context from a default device and platform type - never worked for me.
 
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html>
 
+=item $ctx = OpenCL::context $properties, \@devices, $callback->($err, $pvt) = $print_stderr)
+
+Create a new OpenCL::Context object using the given device object(s). This
+function isn't implemented yet, use C<< $platform->context >> instead.
+
+L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContext.html>
+
 =item OpenCL::wait_for_events $wait_events...
 
 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
@@ -432,13 +645,13 @@
 
 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, $callback->($err, $pvt) = $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 = $platform->context ($properties = undef, @$devices, $notify = undef)
+=item $ctx = $platform->context ($properties, \@devices, $callback->($err, $pvt) = $print_stderr)
 
 Create a new OpenCL::Context object using the given device object(s)- a
 CL_CONTEXT_PLATFORM property is supplied automatically.
@@ -456,6 +669,13 @@
 
 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
@@ -750,9 +970,9 @@
 
 Calls C<clGetDeviceInfo> with C<CL_DEVICE_AFFINITY_DOMAINS_EXT> and returns the result.
 
-=item $uint = $device->reference_count_ext 
+=item $uint = $device->reference_count_ext
 
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_REFERENCE_COUNT_EXT > and returns the result.
+Calls C<clGetDeviceInfo> with C<CL_DEVICE_REFERENCE_COUNT_EXT> and returns the result.
 
 =item @device_partition_property_exts = $device->partition_style_ext
 
@@ -766,12 +986,46 @@
 
 =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.
@@ -790,6 +1044,13 @@
 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
@@ -811,14 +1072,21 @@
 
 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLBuffer.html
 
-=item $ctx->gl_texture2d ($flags, $target, $miplevel, $texture)
+=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 $ctx->gl_texture3d ($flags, $target, $miplevel, $texture)
+=item $img = $ctx->gl_texture3d ($flags, $target, $miplevel, $texture)
 
 Creates a new OpenCL::Image3D object that refers to the given OpenGL
 3D texture object.
@@ -891,11 +1159,17 @@
 
 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
 
@@ -949,6 +1223,22 @@
 
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>.
 
+=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.2/docs/man/xhtml/clEnqueueFillImage.html>
+
 =item $ev = $queue->enqueue_task ($kernel, $wait_events...)
 
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html>
@@ -971,10 +1261,6 @@
 
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html>
 
-=item $ev = $queue->enqueue_marker ($wait_events...)
-
-L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMarker.html>
-
 =item $ev = $queue->enqueue_acquire_gl_objects ([object, ...], $wait_events...)
 
 Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired
@@ -993,9 +1279,13 @@
 
 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
 
@@ -1120,11 +1410,13 @@
 
 =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.
 
@@ -1133,6 +1425,11 @@
 
 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
@@ -1215,12 +1512,28 @@
 
 =over 4
 
-=item $program->build ($device, $options = "")
+=item $program->build (\@devices = undef, $options = "", $cb->($program) = undef)
+
+Tries to build the program with the given options. See also the
+C<$ctx->build> convenience function.
 
-Tries to build the program with the givne options.
+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
@@ -1235,6 +1548,12 @@
 
 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)
@@ -1371,12 +1690,17 @@
 
 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
-objects of that type.
+must be an object of that type or C<undef>, local-memory arguments are
+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>
 
@@ -1395,6 +1719,13 @@
 
 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.
@@ -1464,33 +1795,15 @@
 
 =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.95';
-
-   require XSLoader;
-   XSLoader::load (__PACKAGE__, $VERSION);
-
-   @OpenCL::Buffer::ISA =
-   @OpenCL::Image::ISA      = OpenCL::Memory::;
-
-   @OpenCL::BufferObj::ISA  = OpenCL::Buffer::;
-
-   @OpenCL::Image2D::ISA    =
-   @OpenCL::Image3D::ISA    = OpenCL::Image::;
-
-   @OpenCL::UserEvent::ISA  = OpenCL::Event::;
-}
-
 1;
 
 =head1 AUTHOR