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

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.36 by root, Thu Apr 19 14:34:56 2012 UTC vs.
Revision 1.46 by root, Sat Apr 21 19:17:09 2012 UTC

    ';
    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 }
+   eval { $prog->build ($dev, "-cl-fast-relaxed-math"); 1 }
       or die $prog->build_log;
    my $kernel = $prog->kernel ("squareit");
 =head2 Create some input and output float buffers, then call the
    # 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->program_with_source ($src);
+   eval { $prog->build ($dev); 1 }
+      or die $prog->build_log ($dev);
+   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.
 =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
+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
+module tries to dynamically resolve C<glXGetCurrentDisplay> and
-C<glxGetCurrentContext>, call these functions and use their return values
+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>.
 =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.
 =item @device_partition_property_exts = $device->partition_style_ext
 =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.
 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...)
 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>
 =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
 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.
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html>
 =back
 package OpenCL;
 use common::sense;
 BEGIN {
-   our $VERSION = '0.92';
+   our $VERSION = '0.96';
    require XSLoader;
    XSLoader::load (__PACKAGE__, $VERSION);
    @OpenCL::Buffer::ISA =

Diff Legend

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

Comparing OpenCL/OpenCL.pm (file contents): Revision 1.36 by root, Thu Apr 19 14:34:56 2012 UTC vs. Revision 1.46 by root, Sat Apr 21 19:17:09 2012 UTC

Diff Legend

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.36 by root, Thu Apr 19 14:34:56 2012 UTC vs.
Revision 1.46 by root, Sat Apr 21 19:17:09 2012 UTC