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

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.58 by root, Wed Apr 25 20:29:03 2012 UTC vs.
Revision 1.66 by root, Tue May 1 16:37:23 2012 UTC

 =head2 Create a buffer with some predefined data, read it back synchronously,
 then asynchronously.
    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"
 =head2 Create and build a program, then create a kernel out of one of its
 functions.
    # set buffer
    $kernel->set_buffer (0, $input);
    $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;
 =head2 The same enqueue operations as before, but assuming an out-of-order queue,
 showing off barriers.
    # execute it for all 4 numbers
-   $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef);
+   $queue->nd_range_kernel ($kernel, undef, [4], undef);
    # enqueue a barrier to ensure in-order execution
-   $queue->enqueue_barrier;
+   $queue->barrier;
    # enqueue an async read
-   $queue->enqueue_read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data);
+   $queue->read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data);
    # wait for all requests to finish
    $queue->finish;
 =head2 The same enqueue operations as before, but assuming an out-of-order queue,
 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 tunnel effect.
-This is quite a long example to get you going.
+This is quite a long example to get you going - you can download it from
+L<http://cvs.schmorp.de/OpenCL/examples/juliaflight>.
    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 => 256, height => 256;
+   OpenGL::glpOpenWindow width => $S, height => $S;
    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;
+   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;
    # 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;
+     int2 xy = (int2)(get_global_id (0), get_global_id (1));
+     float2 p = convert_float2 (xy) / $S.f * 2.f - 1.f;
-     float2 m = (float2)(1.f, p.y) / fabs (p.x);
+     float2 m = (float2)(1.f, p.y) / fabs (p.x); // tunnel
-     m.x = fabs (fmod (m.x + time * 0.05f, 4.f)) - 2.f;
+     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));
+     float2 c = (float2)(sin (time * 0.01133f), cos (time * 0.02521f));
-     for (int i = 0; i < 25 && dot (z, z) < 4.f; ++i)
+     for (int i = 0; i < 25 && dot (z, z) < 4.f; ++i) // standard julia
        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);
+     float3 colour = (float3)(z.x, z.y, atan2 (z.y, z.x));
-     write_imagef (img, (int2)(get_global_id (0), get_global_id (1)), (float4)(colour * p.x * p.x, 1.));
+     write_imagef (img, xy, (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]);
+      $queue->acquire_gl_objects ([$tex]);
       # configure and run our kernel
-      $kernel->set_image2d (0, $tex);
+      $kernel->setf ("mf", $tex, $time*2); # mf = memory object, float
-      $kernel->set_float   (1, $time);
-      $queue->enqueue_nd_range_kernel ($kernel, undef, [256, 256], undef);
+      $queue->nd_range_kernel ($kernel, undef, [$S, $S], undef);
       # release objects to opengl again
-      $queue->enqueue_release_gl_objects ([$tex]);
+      $queue->release_gl_objects ([$tex]);
       # wait
       $queue->finish;
       # now draw the texture, the defaults should be all right
       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 darw 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<http://cvs.schmorp.de/OpenCL/examples/juliaflight-nosharing>.
 =head1 DOCUMENTATION
 =head2 BASIC CONVENTIONS
    use AnyEvent;
    # initialise AnyEvent, by creating a watcher, or:
    AnyEvent::detect;
-   my $e = $queue->enqueue_marker;
+   my $e = $queue->marker;
    $e->cb (sub {
       warn "opencl is finished\n";
    })
    main Gtk2;
 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 $e = $queue->marker;
    my $flag;
    $e->cb (sub { $flag = 1 });
    1 until $flag;
    # $flag is now 1
 =cut
 package OpenCL;
 use common::sense;
+use Carp ();
 use Async::Interrupt ();
 our $POLL_FUNC; # set by XS
 BEGIN {
-   our $VERSION = '0.97';
+   our $VERSION = '0.98';
    require XSLoader;
    XSLoader::load (__PACKAGE__, $VERSION);
    @OpenCL::Platform::ISA      =
    @OpenCL::Image1D::ISA       =
    @OpenCL::Image1DArray::ISA  =
    @OpenCL::Image1DBuffer::ISA = OpenCL::Image::;
    @OpenCL::UserEvent::ISA     = OpenCL::Event::;
+   @OpenCL::MappedBuffer       =
+   @OpenCL::MappedImage        = OpenCL::Mapped::;
 }
 =head2 THE OpenCL PACKAGE
 =over 4
 =item $int = OpenCL::errno
 The last error returned by a function - it's only valid after an error occured
 and before calling another OpenCL function.
-=item $str = OpenCL::err2str $errval
+=item $str = OpenCL::err2str [$errval]
-Comverts an error value into a human readable string.
+Converts an error value into a human readable string. IF no error value is
+given, then the last error will be used (as returned by OpenCL::errno).
 =item $str = OpenCL::enum2str $enum
 Converts most enum values (of parameter names, image format constants,
 object types, addressing and filter modes, command types etc.) into a
 =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
+   eval { $prog->build (undef, $options); 1 }
+      or errno == BUILD_PROGRAM_FAILURE
+      or errno == INVALID_BINARY # workaround nvidia bug
+      or Carp::croak "OpenCL::Context->build_program: " . err2str;
+   # we check status for all devices
    for my $dev ($self->devices) {
-      eval { $prog->build ([$dev], $options); 1 }
+      $prog->build_status ($dev) == BUILD_SUCCESS
-         or Carp::croak ("Building OpenCL program for device '" . $dev->name . "' failed:\n"
+         or Carp::croak "Building OpenCL program for device '" . $dev->name . "' failed:\n"
-                         . $prog->build_log ($dev));
+                        . $prog->build_log ($dev);
    }
    $prog
 }
 =item $buf = $ctx->buffer_sv ($flags, $data)
 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)
+=item $img = $ctx->image ($self, $flags, $channel_order, $channel_type, $type, $width, $height, $depth = 0, $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>
 =back
 =head2 THE OpenCL::Queue CLASS
 An OpenCL::Queue represents an execution queue for OpenCL. You execute
-requests by calling their respective C<enqueue_xxx> method and waitinf for
+requests by calling their respective method and waiting for it to complete
-it to complete in some way.
+in some way.
-All the enqueue methods return an event object that can be used to wait
+Most methods that enqueue some request return an event object that can
-for completion, unless the method is called in void context, in which case
+be used to wait for completion (optionally using a callback), unless
-no event object is created.
+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. 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);
+   $event = $queue->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, 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->read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...)
 Reads data from buffer into the given string.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBuffer.html>
-=item $ev = $queue->enqueue_write_buffer ($buffer, $blocking, $offset, $data, $wait_events...)
+=item $ev = $queue->write_buffer ($buffer, $blocking, $offset, $data, $wait_events...)
 Writes data to buffer from the given string.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBuffer.html>
-=item $ev = $queue->enqueue_copy_buffer ($src, $dst, $src_offset, $dst_offset, $len, $wait_events...)
+=item $ev = $queue->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...)
+=item $ev = $queue->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...)
+=item $ev = $queue->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...)
+=item $ev = $queue->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->read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
+C<$row_pitch> (and C<$slice_pitch>) can be C<0>, in which case the OpenCL
+module uses the image width (and height) to supply default values.
 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...)
+=item $ev = $queue->write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
+C<$row_pitch> (and C<$slice_pitch>) can be C<0>, in which case the OpenCL
+module uses the image width (and height) to supply default values.
 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...)
+=item $ev = $queue->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...)
+=item $ev = $queue->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...)
+=item $ev = $queue->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/clEnqueueCopyBufferToImage.html>.
-=item $ev = $queue->enqueue_fill_buffer ($mem, $pattern, $offset, $size, ...)
+=item $ev = $queue->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, ...)
+=item $ev = $queue->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...)
+=item $ev = $queue->task ($kernel, $wait_events...)
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html>
-=item $ev = $queue->enqueue_nd_range_kernel ($kernel, \@global_work_offset, \@global_work_size, \@local_work_size, $wait_events...)
+=item $ev = $queue->nd_range_kernel ($kernel, \@global_work_offset, \@global_work_size, \@local_work_size, $wait_events...)
 Enqueues a kernel execution.
 \@global_work_size must be specified as a reference to an array of
 integers specifying the work sizes (element counts).
 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_acquire_gl_objects ([object, ...], $wait_events...)
+=item $ev = $queue->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/clEnqueueAcquireGLObjects.html>
-=item $ev = $queue->enqueue_release_gl_objects ([object, ...], $wait_events...)
+=item $ev = $queue->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...)
+=item $ev = $queue->wait_for_events ($wait_events...)
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html>
-=item $ev = $queue->enqueue_marker ($wait_events...)
+=item $ev = $queue->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...)
+=item $ev = $queue->barrier ($wait_events...)
 L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueBarrierWithWaitList.html>
 =item $queue->flush
 =item $command_queue_properties = $command_queue->properties
 Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_PROPERTIES> and returns the result.
 =for gengetinfo end command_queue
+=back
+=head3 MEMORY MAPPED BUFFERS
+OpenCL allows you to map buffers and images to host memory (read: perl
+scalars). This is done much like reading or copying a buffer, by enqueuing
+a map or unmap operation on the command queue.
+The map operations return a C<OpenCL::Mapped> object - see L<THE
+OpenCL::Mapped CLASS> section for details on what to do with these
+objects.
+The object will be unmapped automatically when the mapped object is
+destroyed (you can use a barrier to make sure the unmap has finished,
+before using the buffer in a kernel), but you can also enqueue an unmap
+operation manually.
+=over 4
+=item $mapped_buffer = $queue->map_buffer ($buf, $data, $blocking=1, $map_flags=OpenCL::MAP_READ|OpenCL::MAP_WRITE, $offset=0, $size=0, $wait_events...)
+Maps the given buffer into host memory and returns a C<OpenCL::MappedBuffer> object.
+L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMapBuffer.html>
+=item $mapped_image = $queue->map_image ($img, $data, $blocking=1, $map_flags=OpenCL::MAP_READ|OpenCL::MAP_WRITE, $x=0, $y=0, $z=0, $width=0, $height=0, $depth=0, $wait_events...)
+Maps the given image area into host memory and return a
+C<OpenCL::MappedImage> object. Although there are default values for most
+arguments, you currently have to specify all arguments, otherwise the call
+will fail.
+L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMapImage.html>
+=item $ev = $queue->unmap ($mapped, $wait_events...)
+Unmaps the data from host memory. You must not call any methods that
+modify the data, or modify the data scalar directly, after calling this
+method.
+The mapped event object will always be passed as part of the
+$wait_events. The mapped event object will be replaced by the new event
+object that this request creates.
 =back
 =head2 THE OpenCL::Memory CLASS
 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
+Note that some OpenCL implementations act 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.
+Some implementations fail with C<OpenCL::INVALID_BINARY> when the
+compilation state is successful but some later stage fails.
 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
 =item $ulong = $kernel->private_mem_size ($device)
 Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PRIVATE_MEM_SIZE> and returns the result.
 =for gengetinfo end kernel_work_group
+=item $kernel->setf ($format, ...)
+Sets the arguments of a kernel. Since OpenCL 1.1 doesn't have a generic
+way to set arguments (and with OpenCL 1.2 it might be rather slow), you
+need to specify a format argument, much as with C<printf>, to tell OpenCL
+what type of argument it is.
+The format arguments are single letters:
+   c   char
+   C   unsigned char
+   s   short
+   S   unsigned short
+   i   int
+   I   unsigned int
+   l   long
+   L   unsigned long
+   h   half float (0..65535)
+   f   float
+   d   double
+   z   local (octet size)
+   m   memory object (buffer or image)
+   a   sampler
+   e   event
+Space characters in the format string are ignored.
+Example: set the arguments for a kernel that expects an int, two floats, a buffer and an image.
+   $kernel->setf ("i ff mm", 5, 0.5, 3, $buffer, $image);
 =item $kernel->set_TYPE    ($index, $value)
 =item $kernel->set_char    ($index, $value)
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetUserEventStatus.html>
 =back
+=head2 THE OpenCL::Mapped CLASS
+This class represents objects mapped into host memory. They are
+represented by a blessed string scalar. The string data is the mapped
+memory area, that is, if you read or write it, then the mapped object is
+accessed directly.
+You must only ever use operations that modify the string in-place - for
+example, a C<substr> that doesn't change the length, or maybe a regex that
+doesn't change the length. Any other operation might cause the data to be
+copied.
+When the object is destroyed it will enqueue an implicit unmap operation
+on the queue that was used to create it.
+Example, replace the first two floats in the mapped buffer by 1 and 2.
+   my $mapped = $queue->map_buffer ($buf, ...
+   $mapped->event->wait; # make sure it's there
+   # now replace first 8 bytes by new data, which is exactly 8 bytes long
+   # we blindly assume device endianness to equal host endianness
+   # (and of course, we assume iee 754 single precision floats :)
+   substr $$mapped, 0, 8, pack "f*", 1, 2;
+=over 4
+=item $bool = $mapped->mapped
+Returns whether the object is still mapped - true before an C<unmap> is
+enqueued, false afterwards.
+=item $ev = $mapped->event
+Return the event object associated with the mapped object. Initially, this
+will be the event object created when mapping the object, and after an
+unmap, this will be the event object that the unmap operation created.
+=item $mapped->wait
+Same as C<< $mapped->event->wait >> - makes sure no operations on this
+mapped object are outstanding.
+=item $bytes = $mapped->size
+Returns the size of the mapped area, in bytes. Same as C<length $$mapped>.
+=item $ptr = $mapped->ptr
+Returns the raw memory address of the mapped area - same as C<$mapped+0>.
+=back
+=head2 THE OpenCL::MappedBuffer CLASS
+This is a subclass of OpenCL::Mapped, representing mapped buffers.
+=over 4
+=back
+=head2 THE OpenCL::MappedImage CLASS
+This is a subclass of OpenCL::Mapped, representing mapped images.
+=over 4
+=back
 =cut
 1;
 =head1 AUTHOR

Diff Legend

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

Comparing OpenCL/OpenCL.pm (file contents): Revision 1.58 by root, Wed Apr 25 20:29:03 2012 UTC vs. Revision 1.66 by root, Tue May 1 16:37:23 2012 UTC

Diff Legend

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.58 by root, Wed Apr 25 20:29:03 2012 UTC vs.
Revision 1.66 by root, Tue May 1 16:37:23 2012 UTC