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

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.35 by root, Thu Apr 19 13:18:31 2012 UTC vs.
Revision 1.53 by root, Tue Apr 24 14:30:57 2012 UTC

         $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_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>.
+=cut
+package OpenCL;
+use common::sense;
+BEGIN {
+   our $VERSION = '0.96';
+   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
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html>
 =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 {
+   ${$_[0]}
+}
+=back
 =head2 THE OpenCL::Platform CLASS
 =over 4
 =item @devices = $platform->devices ($type = OpenCL::DEVICE_TYPE_ALL)
 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.
 =item @device_partition_property_exts = $device->affinity_domains_ext
 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.
+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.
 =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) = @_;
+   $prog = $self->program_with_source ($prog)
+      unless ref $prog;
+   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.
 =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)
+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::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 $ctx->gl_texture2d ($flags, $target, $miplevel, $texture)
+=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.
 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture3D.html
 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_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, ...)
+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>
 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 ($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
 for subsequent OpenCL usage.
 =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>
 =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.
 =for gengetinfo begin gl_texture
 =item $GLenum = $gl_texture->target
-Calls C<clGetGlTextureInfo> with C<CL_GL_TEXTURE_TARGET> and returns the result.
+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.
+Calls C<clGetGLTextureInfo> with C<CL_GL_MIPMAP_LEVEL> and returns the result.
 =for gengetinfo end gl_texture
 =back
 =over 4
 =item $program->build ($device, $options = "")
-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.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clBuildProgram.html>
 =item $packed_value = $program->build_info ($device, $name)
 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)
 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
 =back
 =cut
-package OpenCL;
-use common::sense;
-BEGIN {
-   our $VERSION = '0.92';
-   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
  Marc Lehmann <schmorp@schmorp.de>

Diff Legend

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

Comparing OpenCL/OpenCL.pm (file contents): Revision 1.35 by root, Thu Apr 19 13:18:31 2012 UTC vs. Revision 1.53 by root, Tue Apr 24 14:30:57 2012 UTC

Diff Legend

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.35 by root, Thu Apr 19 13:18:31 2012 UTC vs.
Revision 1.53 by root, Tue Apr 24 14:30:57 2012 UTC