--- OpenCL/OpenCL.pm	2011/11/20 22:29:36	1.23
+++ OpenCL/OpenCL.pm	2012/04/20 20:24:20	1.43
@@ -103,11 +103,11 @@
 system.
 
    for my $platform (OpenCL::platforms) {
-      printf "platform: %s\n", $platform->info (OpenCL::PLATFORM_NAME);
-      printf "extensions: %s\n", $platform->info (OpenCL::PLATFORM_EXTENSIONS);
+      printf "platform: %s\n", $platform->name;
+      printf "extensions: %s\n", $platform->extensions;
       for my $device ($platform->devices) {
-         printf "+ device: %s\n", $device->info (OpenCL::DEVICE_NAME);
-         my $ctx = $device->context;
+         printf "+ device: %s\n", $device->name;
+         my $ctx = $platform->context (undef, [$device]);
          # do stuff
       }
    }
@@ -151,8 +151,8 @@
 functions.
 
    my $src = '
-      __kernel void
-      squareit (__global float *input, __global float *output)
+      kernel void
+      squareit (global float *input, global float *output)
       {
         $id = get_global_id (0);
         output [id] = input [id] * input [id];
@@ -162,8 +162,8 @@
    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_info ($dev, OpenCL::PROGRAM_BUILD_LOG);
+   eval { $prog->build ($dev, "-cl-fast-relaxed-math"); 1 }
+      or die $prog->build_log;
 
    my $kernel = $prog->kernel ("squareit");
 
@@ -213,6 +213,101 @@
    # 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
@@ -273,6 +368,21 @@
    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 THE OpenCL PACKAGE
 
 =over 4
@@ -288,9 +398,9 @@
 
 =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.
 
@@ -328,7 +438,7 @@
 
 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 = undef, @$devices, $notify = undef)
 
 Create a new OpenCL::Context object using the given device object(s)- a
 CL_CONTEXT_PLATFORM property is supplied automatically.
@@ -350,23 +460,23 @@
 
 =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
 
@@ -386,267 +496,267 @@
 
 =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
 
@@ -670,26 +780,58 @@
 
 =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->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 $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)
+
+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
@@ -719,19 +861,19 @@
 
 =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
 
@@ -773,23 +915,25 @@
 
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBuffer.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/clEnqueueReadImage.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->enqueue_write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
+http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBufferRect.html
 
-L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteImage.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->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...)
+http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBufferRect.html
 
-Yeah.
+=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/clEnqueueCopyBufferToImage.html>.
+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...)
 
@@ -799,6 +943,12 @@
 
 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/clEnqueueCopyBufferToImage.html>.
+
 =item $ev = $queue->enqueue_task ($kernel, $wait_events...)
 
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html>
@@ -821,10 +971,24 @@
 
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html>
 
-=item $ev = $queue->enqueue_marker
+=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.
+
+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>
@@ -851,19 +1015,19 @@
 
 =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
 
@@ -886,42 +1050,72 @@
 
 =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
@@ -943,30 +1137,42 @@
 
 =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
@@ -983,23 +1189,23 @@
 
 =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
 
@@ -1033,15 +1239,15 @@
 
 =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
 
@@ -1055,27 +1261,27 @@
 
 =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
 
@@ -1106,23 +1312,23 @@
 
 =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
 
@@ -1139,23 +1345,23 @@
 
 =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
 
@@ -1165,12 +1371,13 @@
 
 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.
 
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html>
 
@@ -1199,23 +1406,23 @@
 
 =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
 
@@ -1232,19 +1439,19 @@
 
 =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
 
@@ -1269,18 +1476,20 @@
 use common::sense;
 
 BEGIN {
-   our $VERSION = '0.55';
+   our $VERSION = '0.96';
 
    require XSLoader;
    XSLoader::load (__PACKAGE__, $VERSION);
 
-   @OpenCL::Buffer::ISA    =
-   @OpenCL::Image::ISA     = OpenCL::Memory::;
+   @OpenCL::Buffer::ISA =
+   @OpenCL::Image::ISA      = OpenCL::Memory::;
+
+   @OpenCL::BufferObj::ISA  = OpenCL::Buffer::;
 
-   @OpenCL::Image2D::ISA   =
-   @OpenCL::Image3D::ISA   = OpenCL::Image::;
+   @OpenCL::Image2D::ISA    =
+   @OpenCL::Image3D::ISA    = OpenCL::Image::;
 
-   @OpenCL::UserEvent::ISA = OpenCL::Event::;
+   @OpenCL::UserEvent::ISA  = OpenCL::Event::;
 }
 
 1;