--- OpenCL/OpenCL.pm	2012/05/01 22:04:56	1.67
+++ OpenCL/OpenCL.pm	2012/05/05 12:54:03	1.77
@@ -45,7 +45,8 @@
 
 =head2 HELPFUL RESOURCES
 
-The OpenCL specs used to develop this module:
+The OpenCL specs used to develop this module - download these and keept
+hema round, they are required reference material:
 
    http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
    http://www.khronos.org/registry/cl/specs/opencl-1.2.pdf
@@ -98,6 +99,37 @@
 
 Enqueue buffer reads for your output buffer to read results.
 
+=head1 OPENCL 1.1 VS. OPENCL 1.2
+
+This module supports both OpenCL version 1.1 and 1.2, although the OpenCL
+1.2 interface hasn't been tested much for lack of availability of an
+actual implementation.
+
+Every function or method in this manual page that interfaces to a
+particular OpenCL function has a link to the its C manual page.
+
+If the link contains a F<1.1>, then this function is an OpenCL 1.1
+function. Most but not all also exist in OpenCL 1.2, and this module
+tries to emulate the missing ones for you, when told to do so at
+compiletime. You cna check whether a function was removed in OpenCL 1.2 by
+replacing the F<1.1> component in the URL by F<1.2>.
+
+If the link contains a F<1.2>, then this is a OpenCL 1.2-only
+function. Even if the module was compiled with OpenCL 1.2 header files
+and has an 1.2 OpenCL library, calling such a function on a platform that
+doesn't implement 1.2 causes undefined behaviour, usually a crash (But
+this is not guaranteed).
+
+You can find out whether this module was compiled to prefer 1.1
+functionality by ooking at C<OpenCL::PREFER_1_1> - if it is true, then
+1.1 functions generally are implemented using 1.1 OpenCL functions. If it
+is false, then 1.1 functions missing from 1.2 are emulated by calling 1.2
+fucntions.
+
+This is a somewhat sorry state of affairs, but the Khronos group choose to
+make every release of OpenCL source and binary incompatible with previous
+releases.
+
 =head1 EXAMPLES
 
 =head2 Enumerate all devices and get contexts for them.
@@ -212,7 +244,7 @@
    $ev->wait;
 
 =head2 Use the OpenGL module to share a texture between OpenCL and OpenGL and draw some julia
-set tunnel effect.
+set flight effect.
 
 This is quite a long example to get you going - you can download it from
 L<http://cvs.schmorp.de/OpenCL/examples/juliaflight>.
@@ -322,7 +354,7 @@
 
    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
+And in the draw 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);
@@ -334,6 +366,23 @@
 The fully modified example can be found at
 L<http://cvs.schmorp.de/OpenCL/examples/juliaflight-nosharing>.
 
+=head2 Julia sets look soooo 80ies.
+
+Then colour them differently, e.g. using orbit traps! Replace the loop and
+colour calculation from the previous examples by this:
+
+  float2 dm = (float2)(1.f, 1.f);
+
+  for (int i = 0; i < 25; ++i)
+    {
+      z = (float2)(z.x * z.x - z.y * z.y, 2.f * z.x * z.y) + c;
+      dm = fmin (dm, (float2)(fabs (dot (z, z) - 1.f), fabs (z.x - 1.f)));
+    }
+
+  float3 colour = (float3)(dm.x * dm.y, dm.x * dm.y, dm.x);
+
+Also try C<-10.f> instead of C<-1.f>.
+
 =head1 DOCUMENTATION
 
 =head2 BASIC CONVENTIONS
@@ -418,13 +467,20 @@
 perhaps most important, to signal completion of queued jobs (by setting
 callbacks on OpenCL::Event objects).
 
+The OpenCL module converts all these callbacks into events - you can
+still register callbacks, but they are not executed when your OpenCL
+implementation calls the actual callback, but only later. Therefore, none
+of the limitations of OpenCL callbacks apply to the perl implementation:
+it is perfectly safe to make blocking operations from event callbacks, and
+enqueued operations don't need to be flushed.
+
 To facilitate this, this module maintains an event queue - each
 time an asynchronous event happens, it is queued, and perl will be
 interrupted. This is implemented via the L<Async::Interrupt> module. In
 addition, this module has L<AnyEvent> support, so it can seamlessly
 integrate itself into many event loops.
 
-Since this module is a bit hard to understand, here are some case examples:
+Since L<Async::Interrupt> is a bit hard to understand, here are some case examples:
 
 =head3 Don't use callbacks.
 
@@ -531,7 +587,7 @@
 our $POLL_FUNC; # set by XS
 
 BEGIN {
-   our $VERSION = '0.98';
+   our $VERSION = '0.99';
 
    require XSLoader;
    XSLoader::load (__PACKAGE__, $VERSION);
@@ -546,6 +602,8 @@
    @OpenCL::Kernel::ISA        =
    @OpenCL::Event::ISA         = OpenCL::Object::;
 
+   @OpenCL::SubDevice::ISA     = OpenCL::Device::;
+
    @OpenCL::Buffer::ISA        =
    @OpenCL::Image::ISA         = OpenCL::Memory::;
 
@@ -595,13 +653,17 @@
 =item $ctx = OpenCL::context_from_type $properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $callback->($err, $pvt) = $print_stderr
 
 Tries to create a context from a default device and platform type - never worked for me.
+Consider using C<< $platform->context_from_type >> instead.
+
+type: OpenCL::DEVICE_TYPE_DEFAULT, OpenCL::DEVICE_TYPE_CPU, OpenCL::DEVICE_TYPE_GPU,
+OpenCL::DEVICE_TYPE_ACCELERATOR, OpenCL::DEVICE_TYPE_CUSTOM, OpenCL::DEVICE_TYPE_ALL.
 
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html>
 
 =item $ctx = OpenCL::context $properties, \@devices, $callback->($err, $pvt) = $print_stderr)
 
-Create a new OpenCL::Context object using the given device object(s). This
-function isn't implemented yet, use C<< $platform->context >> instead.
+Create a new OpenCL::Context object using the given device object(s).
+Consider using C<< $platform->context >> instead.
 
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContext.html>
 
@@ -718,23 +780,23 @@
 
 =item $string = $platform->profile
 
-Calls C<clGetPlatformInfo> with C<CL_PLATFORM_PROFILE> and returns the result.
+Calls C<clGetPlatformInfo> with C<OpenCL::PLATFORM_PROFILE> and returns the result.
 
 =item $string = $platform->version
 
-Calls C<clGetPlatformInfo> with C<CL_PLATFORM_VERSION> and returns the result.
+Calls C<clGetPlatformInfo> with C<OpenCL::PLATFORM_VERSION> and returns the result.
 
 =item $string = $platform->name
 
-Calls C<clGetPlatformInfo> with C<CL_PLATFORM_NAME> and returns the result.
+Calls C<clGetPlatformInfo> with C<OpenCL::PLATFORM_NAME> and returns the result.
 
 =item $string = $platform->vendor
 
-Calls C<clGetPlatformInfo> with C<CL_PLATFORM_VENDOR> and returns the result.
+Calls C<clGetPlatformInfo> with C<OpenCL::PLATFORM_VENDOR> and returns the result.
 
 =item $string = $platform->extensions
 
-Calls C<clGetPlatformInfo> with C<CL_PLATFORM_EXTENSIONS> and returns the result.
+Calls C<clGetPlatformInfo> with C<OpenCL::PLATFORM_EXTENSIONS> and returns the result.
 
 =for gengetinfo end platform
 
@@ -750,271 +812,277 @@
 
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetDeviceInfo.html>
 
+=item @devices = $device->sub_devices (\@properties)
+
+Creates OpencL::SubDevice objects by partitioning an existing device.
+
+L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateSubDevices.html>
+
 =for gengetinfo begin device
 
 =item $device_type = $device->type
 
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_TYPE> and returns the result.
+Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_TYPE> and returns the result.
 
 =item $uint = $device->vendor_id
 
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_VENDOR_ID> and returns the result.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_ENDIAN_LITTLE> and returns the result.
 
 =item $boolean = $device->available
 
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_AVAILABLE> and returns the result.
+Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_AVAILABLE> and returns the result.
 
 =item $boolean = $device->compiler_available
 
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_COMPILER_AVAILABLE> and returns the result.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_QUEUE_PROPERTIES> and returns the result.
 
 =item $ = $device->platform
 
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_PLATFORM> and returns the result.
+Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_PLATFORM> and returns the result.
 
 =item $string = $device->name
 
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_NAME> and returns the result.
+Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_NAME> and returns the result.
 
 =item $string = $device->vendor
 
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_VENDOR> and returns the result.
+Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_VENDOR> and returns the result.
 
 =item $string = $device->driver_version
 
-Calls C<clGetDeviceInfo> with C<CL_DRIVER_VERSION> and returns the result.
+Calls C<clGetDeviceInfo> with C<OpenCL::DRIVER_VERSION> and returns the result.
 
 =item $string = $device->profile
 
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_PROFILE> and returns the result.
+Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_PROFILE> and returns the result.
 
 =item $string = $device->version
 
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_VERSION> and returns the result.
+Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_VERSION> and returns the result.
 
 =item $string = $device->extensions
 
-Calls C<clGetDeviceInfo> with C<CL_DEVICE_EXTENSIONS> and returns the result.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::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<OpenCL::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.
+Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_PARTITION_STYLE_EXT> and returns the result.
 
 =for gengetinfo end device
 
@@ -1022,6 +1090,32 @@
 
 =head2 THE OpenCL::Context CLASS
 
+An OpenCL::Context is basically a container, or manager, for a number of
+devices of a platform. It is used to create all sorts of secondary objects
+such as buffers, queues, programs and so on.
+
+All context creation functions and methods take a list of properties
+(type-value pairs). All property values can be specified as integers -
+some additionally support other types:
+
+=over 4
+
+=item OpenCL::CONTEXT_PLATFORM
+
+Also accepts OpenCL::Platform objects.
+
+=item OpenCL::GLX_DISPLAY_KHR
+
+Also accepts C<undef>, in which case a deep and troubling hack is engaged
+to find the current glx display (see L<GLX SUPPORT>).
+
+=item OpenCL::GL_CONTEXT_KHR
+
+Also accepts C<undef>, in which case a deep and troubling hack is engaged
+to find the current glx context (see L<GLX SUPPORT>).
+
+=back
+
 =over 4
 
 =item $prog = $ctx->build_program ($program, $options = "")
@@ -1078,6 +1172,10 @@
 Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object with the
 given flags and octet-size.
 
+flags: OpenCL::MEM_READ_WRITE, OpenCL::MEM_WRITE_ONLY, OpenCL::MEM_READ_ONLY,
+OpenCL::MEM_USE_HOST_PTR, OpenCL::MEM_ALLOC_HOST_PTR, OpenCL::MEM_COPY_HOST_PTR,
+OpenCL::MEM_HOST_WRITE_ONLY, OpenCL::MEM_HOST_READ_ONLY, OpenCL::MEM_HOST_NO_ACCESS.
+
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateBuffer.html>
 
 =item $buf = $ctx->buffer_sv ($flags, $data)
@@ -1090,6 +1188,22 @@
 Creates a new OpenCL::Image object and optionally initialises it with
 the given data values.
 
+channel_order: OpenCL::R, OpenCL::A, OpenCL::RG, OpenCL::RA, OpenCL::RGB,
+OpenCL::RGBA, OpenCL::BGRA, OpenCL::ARGB, OpenCL::INTENSITY, OpenCL::LUMINANCE,
+OpenCL::Rx, OpenCL::RGx, OpenCL::RGBx.
+
+channel_type: OpenCL::SNORM_INT8, OpenCL::SNORM_INT16, OpenCL::UNORM_INT8,
+OpenCL::UNORM_INT16, OpenCL::UNORM_SHORT_565, OpenCL::UNORM_SHORT_555,
+OpenCL::UNORM_INT_101010, OpenCL::SIGNED_INT8, OpenCL::SIGNED_INT16,
+OpenCL::SIGNED_INT32, OpenCL::UNSIGNED_INT8, OpenCL::UNSIGNED_INT16,
+OpenCL::UNSIGNED_INT32, OpenCL::HALF_FLOAT, OpenCL::FLOAT.
+
+
+type: OpenCL::MEM_OBJECT_BUFFER, OpenCL::MEM_OBJECT_IMAGE2D,
+OpenCL::MEM_OBJECT_IMAGE3D, OpenCL::MEM_OBJECT_IMAGE2D_ARRAY,
+OpenCL::MEM_OBJECT_IMAGE1D, OpenCL::MEM_OBJECT_IMAGE1D_ARRAY,
+OpenCL::MEM_OBJECT_IMAGE1D_BUFFER.
+
 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)
@@ -1111,6 +1225,8 @@
 Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object that refers to the given
 OpenGL buffer object.
 
+flags: OpenCL::MEM_READ_WRITE, OpenCL::MEM_READ_ONLY, OpenCL::MEM_WRITE_ONLY.
+
 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLBuffer.html
 
 =item $img = $ctx->gl_texture ($flags, $target, $miplevel, $texture)
@@ -1118,6 +1234,13 @@
 Creates a new OpenCL::Image object that refers to the given OpenGL
 texture object or buffer.
 
+target: GL_TEXTURE_1D, GL_TEXTURE_1D_ARRAY, GL_TEXTURE_BUFFER,
+GL_TEXTURE_2D, GL_TEXTURE_2D_ARRAY, GL_TEXTURE_3D,
+GL_TEXTURE_CUBE_MAP_POSITIVE_X, GL_TEXTURE_CUBE_MAP_POSITIVE_Y,
+GL_TEXTURE_CUBE_MAP_POSITIVE_Z, GL_TEXTURE_CUBE_MAP_NEGATIVE_X,
+GL_TEXTURE_CUBE_MAP_NEGATIVE_Y, GL_TEXTURE_CUBE_MAP_NEGATIVE_Z,
+GL_TEXTURE_RECTANGLE/GL_TEXTURE_RECTANGLE_ARB.
+
 http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateFromGLTexture.html
 
 =item $img = $ctx->gl_texture2d ($flags, $target, $miplevel, $texture)
@@ -1152,6 +1275,11 @@
 
 Creates a new OpenCL::Sampler object.
 
+addressing_mode: OpenCL::ADDRESS_NONE, OpenCL::ADDRESS_CLAMP_TO_EDGE,
+OpenCL::ADDRESS_CLAMP, OpenCL::ADDRESS_REPEAT, OpenCL::ADDRESS_MIRRORED_REPEAT.
+
+filter_mode: OpenCL::FILTER_NEAREST, OpenCL::FILTER_LINEAR.
+
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateSampler.html>
 
 =item $program = $ctx->program_with_source ($string)
@@ -1160,6 +1288,30 @@
 
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateProgramWithSource.html>
 
+=item ($program, \@status) = $ctx->program_with_binary (\@devices, \@binaries)
+
+Creates a new OpenCL::Program object from the given binaries.
+
+L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateProgramWithBinary.html>
+
+Example: clone an existing program object that contains a successfully
+compiled program, no matter how useless this is.
+
+   my $clone = $ctx->program_with_binary ([$prog->devices], [$prog->binaries]);
+
+=item $program = $ctx->program_with_built_in_kernels (\@devices, $kernel_names)
+
+Creates a new OpenCL::Program object from the given built-in kernel names.
+
+L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateProgramWithBuiltInKernels.html>
+
+=item $program = $ctx->link_program (\@devices, $options, \@programs, $cb->($program) = undef)
+
+Links all (already compiled) program objects specified in C<@programs>
+together and returns a new OpenCL::Program object with the result.
+
+L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clLinkProgram.html>
+
 =item $packed_value = $ctx->info ($name)
 
 See C<< $platform->info >> for details.
@@ -1170,19 +1322,19 @@
 
 =item $uint = $context->reference_count
 
-Calls C<clGetContextInfo> with C<CL_CONTEXT_REFERENCE_COUNT> and returns the result.
+Calls C<clGetContextInfo> with C<OpenCL::CONTEXT_REFERENCE_COUNT> and returns the result.
 
 =item @devices = $context->devices
 
-Calls C<clGetContextInfo> with C<CL_CONTEXT_DEVICES> and returns the result.
+Calls C<clGetContextInfo> with C<OpenCL::CONTEXT_DEVICES> and returns the result.
 
 =item @property_ints = $context->properties
 
-Calls C<clGetContextInfo> with C<CL_CONTEXT_PROPERTIES> and returns the result.
+Calls C<clGetContextInfo> with C<OpenCL::CONTEXT_PROPERTIES> and returns the result.
 
 =item $uint = $context->num_devices
 
-Calls C<clGetContextInfo> with C<CL_CONTEXT_NUM_DEVICES> and returns the result.
+Calls C<clGetContextInfo> with C<OpenCL::CONTEXT_NUM_DEVICES> and returns the result.
 
 =for gengetinfo end context
 
@@ -1308,6 +1460,14 @@
 
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html>
 
+=item $ev = $queue->migrate_mem_objects (\@mem_objects, $flags, $wait_events...)
+
+Migrates a number of OpenCL::Memory objects to or from the device.
+
+flags: OpenCL::MIGRATE_MEM_OBJECT_HOST, OpenCL::MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED
+
+L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMigrateMemObjects.html>
+
 =item $ev = $queue->acquire_gl_objects ([object, ...], $wait_events...)
 
 Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired
@@ -1352,19 +1512,19 @@
 
 =item $ctx = $command_queue->context
 
-Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_CONTEXT> and returns the result.
+Calls C<clGetCommandQueueInfo> with C<OpenCL::QUEUE_CONTEXT> and returns the result.
 
 =item $device = $command_queue->device
 
-Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_DEVICE> and returns the result.
+Calls C<clGetCommandQueueInfo> with C<OpenCL::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.
+Calls C<clGetCommandQueueInfo> with C<OpenCL::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.
+Calls C<clGetCommandQueueInfo> with C<OpenCL::QUEUE_PROPERTIES> and returns the result.
 
 =for gengetinfo end command_queue
 
@@ -1376,7 +1536,7 @@
 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
+The map operations return an C<OpenCL::Mapped> object - see L<THE
 OpenCL::Mapped CLASS> section for details on what to do with these
 objects.
 
@@ -1387,21 +1547,46 @@
 
 =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...)
+=item $mapped_buffer = $queue->map_buffer ($buf, $blocking=1, $map_flags=OpenCL::MAP_READ|OpenCL::MAP_WRITE, $offset=0, $size=undef, $wait_events...)
 
-Maps the given buffer into host memory and returns a C<OpenCL::MappedBuffer> object.
+Maps the given buffer into host memory and returns an
+C<OpenCL::MappedBuffer> object. If C<$size> is specified as undef, then
+the map will extend to the end of the buffer.
+
+map_flags: OpenCL::MAP_READ, OpenCL::MAP_WRITE, OpenCL::MAP_WRITE_INVALIDATE_REGION.
 
 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...)
+Example: map the buffer $buf fully and replace the first 4 bytes by "abcd", then unmap.
+
+   {
+     my $mapped = $queue->map_buffer ($buf, 1, OpenCL::MAP_WRITE);
+     substr $$mapped, 0, 4, "abcd";
+   } # asynchronously unmap because $mapped is destroyed
+
+=item $mapped_image = $queue->map_image ($img, $blocking=1, $map_flags=OpenCL::MAP_READ|OpenCL::MAP_WRITE, $x=0, $y=0, $z=0, $width=undef, $height=undef, $depth=undef, $wait_events...)
+
+Maps the given image area into host memory and return an
+C<OpenCL::MappedImage> object.
 
-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.
+If any of C<$width>, C<$height> and/or C<$depth> are C<undef> then they
+will be replaced by the maximum possible value.
 
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMapImage.html>
 
+Example: map an image (with OpenCL::UNSIGNED_INT8 channel type) and set
+the first channel of the leftmost column to 5, then explicitly unmap
+it. You are not necessarily meant to do it this way, this example just
+shows you the accessors to use :)
+
+   my $mapped = $queue->map_image ($image, 1, OpenCL::MAP_WRITE);
+
+   $mapped->set ($_ * $mapped->row_pitch, pack "C", 5)
+      for 0..$image->height;
+
+   $mapped->unmap;.
+   $mapped->wait; # only needed for out of order queues normally
+
 =item $ev = $queue->unmap ($mapped, $wait_events...)
 
 Unmaps the data from host memory. You must not call any methods that
@@ -1431,39 +1616,39 @@
 
 =item $mem_object_type = $mem->type
 
-Calls C<clGetMemObjectInfo> with C<CL_MEM_TYPE> and returns the result.
+Calls C<clGetMemObjectInfo> with C<OpenCL::MEM_TYPE> and returns the result.
 
 =item $mem_flags = $mem->flags
 
-Calls C<clGetMemObjectInfo> with C<CL_MEM_FLAGS> and returns the result.
+Calls C<clGetMemObjectInfo> with C<OpenCL::MEM_FLAGS> and returns the result.
 
 =item $int = $mem->size
 
-Calls C<clGetMemObjectInfo> with C<CL_MEM_SIZE> and returns the result.
+Calls C<clGetMemObjectInfo> with C<OpenCL::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.
+Calls C<clGetMemObjectInfo> with C<OpenCL::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.
+Calls C<clGetMemObjectInfo> with C<OpenCL::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.
+Calls C<clGetMemObjectInfo> with C<OpenCL::MEM_REFERENCE_COUNT> and returns the result.
 
 =item $ctx = $mem->context
 
-Calls C<clGetMemObjectInfo> with C<CL_MEM_CONTEXT> and returns the result.
+Calls C<clGetMemObjectInfo> with C<OpenCL::MEM_CONTEXT> and returns the result.
 
 =item $mem = $mem->associated_memobject
 
-Calls C<clGetMemObjectInfo> with C<CL_MEM_ASSOCIATED_MEMOBJECT> and returns the result.
+Calls C<clGetMemObjectInfo> with C<OpenCL::MEM_ASSOCIATED_MEMOBJECT> and returns the result.
 
 =item $int = $mem->offset
 
-Calls C<clGetMemObjectInfo> with C<CL_MEM_OFFSET> and returns the result.
+Calls C<clGetMemObjectInfo> with C<OpenCL::MEM_OFFSET> and returns the result.
 
 =for gengetinfo end mem
 
@@ -1525,27 +1710,27 @@
 
 =item $int = $image->element_size
 
-Calls C<clGetImageInfo> with C<CL_IMAGE_ELEMENT_SIZE> and returns the result.
+Calls C<clGetImageInfo> with C<OpenCL::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.
+Calls C<clGetImageInfo> with C<OpenCL::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.
+Calls C<clGetImageInfo> with C<OpenCL::IMAGE_SLICE_PITCH> and returns the result.
 
 =item $int = $image->width
 
-Calls C<clGetImageInfo> with C<CL_IMAGE_WIDTH> and returns the result.
+Calls C<clGetImageInfo> with C<OpenCL::IMAGE_WIDTH> and returns the result.
 
 =item $int = $image->height
 
-Calls C<clGetImageInfo> with C<CL_IMAGE_HEIGHT> and returns the result.
+Calls C<clGetImageInfo> with C<OpenCL::IMAGE_HEIGHT> and returns the result.
 
 =item $int = $image->depth
 
-Calls C<clGetImageInfo> with C<CL_IMAGE_DEPTH> and returns the result.
+Calls C<clGetImageInfo> with C<OpenCL::IMAGE_DEPTH> and returns the result.
 
 =for gengetinfo end image
 
@@ -1553,11 +1738,11 @@
 
 =item $GLenum = $gl_texture->target
 
-Calls C<clGetGLTextureInfo> with C<CL_GL_TEXTURE_TARGET> and returns the result.
+Calls C<clGetGLTextureInfo> with C<OpenCL::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<OpenCL::GL_MIPMAP_LEVEL> and returns the result.
 
 =for gengetinfo end gl_texture
 
@@ -1577,23 +1762,23 @@
 
 =item $uint = $sampler->reference_count
 
-Calls C<clGetSamplerInfo> with C<CL_SAMPLER_REFERENCE_COUNT> and returns the result.
+Calls C<clGetSamplerInfo> with C<OpenCL::SAMPLER_REFERENCE_COUNT> and returns the result.
 
 =item $ctx = $sampler->context
 
-Calls C<clGetSamplerInfo> with C<CL_SAMPLER_CONTEXT> and returns the result.
+Calls C<clGetSamplerInfo> with C<OpenCL::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.
+Calls C<clGetSamplerInfo> with C<OpenCL::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.
+Calls C<clGetSamplerInfo> with C<OpenCL::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.
+Calls C<clGetSamplerInfo> with C<OpenCL::SAMPLER_FILTER_MODE> and returns the result.
 
 =for gengetinfo end sampler
 
@@ -1621,6 +1806,17 @@
 Some implementations fail with C<OpenCL::INVALID_BINARY> when the
 compilation state is successful but some later stage fails.
 
+options: C<-D name>, C<-D name=definition>, C<-I dir>,
+C<-cl-single-precision-constant>, C<-cl-denorms-are-zero>,
+C<-cl-fp32-correctly-rounded-divide-sqrt>, C<-cl-opt-disable>,
+C<-cl-mad-enable>, C<-cl-no-signed-zeros>, C<-cl-unsafe-math-optimizations>,
+C<-cl-finite-math-only>, C<-cl-fast-relaxed-math>,
+C<-w>, C<-Werror>, C<-cl-std=CL1.1/CL1.2>, C<-cl-kernel-arg-info>,
+C<-create-library>, C<-enable-link-options>.
+
+build_status: OpenCL::BUILD_SUCCESS, OpenCL::BUILD_NONE,
+OpenCL::BUILD_ERROR, OpenCL::BUILD_IN_PROGRESS.
+
 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)
@@ -1628,6 +1824,14 @@
 Similar to C<< ->build >>, except it starts a thread, and never fails (you
 need to check the compilation status form the callback, or by polling).
 
+=item $program->compile (\@devices = undef, $options = "", \%headers = undef, $cb->($program) = undef)
+
+Compiles the given program for the given devices (or all devices if
+undef). If C<$headers> is given, it must be a hashref with include name =>
+OpenCL::Program pairs.
+
+L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCompileProgram.html>
+
 =item $packed_value = $program->build_info ($device, $name)
 
 Similar to C<< $platform->info >>, but returns build info for a previous
@@ -1652,15 +1856,15 @@
 
 =item $build_status = $program->build_status ($device)
 
-Calls C<clGetProgramBuildInfo> with C<CL_PROGRAM_BUILD_STATUS> and returns the result.
+Calls C<clGetProgramBuildInfo> with C<OpenCL::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.
+Calls C<clGetProgramBuildInfo> with C<OpenCL::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.
+Calls C<clGetProgramBuildInfo> with C<OpenCL::PROGRAM_BUILD_LOG> and returns the result.
 
 =for gengetinfo end program_build
 
@@ -1674,27 +1878,27 @@
 
 =item $uint = $program->reference_count
 
-Calls C<clGetProgramInfo> with C<CL_PROGRAM_REFERENCE_COUNT> and returns the result.
+Calls C<clGetProgramInfo> with C<OpenCL::PROGRAM_REFERENCE_COUNT> and returns the result.
 
 =item $ctx = $program->context
 
-Calls C<clGetProgramInfo> with C<CL_PROGRAM_CONTEXT> and returns the result.
+Calls C<clGetProgramInfo> with C<OpenCL::PROGRAM_CONTEXT> and returns the result.
 
 =item $uint = $program->num_devices
 
-Calls C<clGetProgramInfo> with C<CL_PROGRAM_NUM_DEVICES> and returns the result.
+Calls C<clGetProgramInfo> with C<OpenCL::PROGRAM_NUM_DEVICES> and returns the result.
 
 =item @devices = $program->devices
 
-Calls C<clGetProgramInfo> with C<CL_PROGRAM_DEVICES> and returns the result.
+Calls C<clGetProgramInfo> with C<OpenCL::PROGRAM_DEVICES> and returns the result.
 
 =item $string = $program->source
 
-Calls C<clGetProgramInfo> with C<CL_PROGRAM_SOURCE> and returns the result.
+Calls C<clGetProgramInfo> with C<OpenCL::PROGRAM_SOURCE> and returns the result.
 
 =item @ints = $program->binary_sizes
 
-Calls C<clGetProgramInfo> with C<CL_PROGRAM_BINARY_SIZES> and returns the result.
+Calls C<clGetProgramInfo> with C<OpenCL::PROGRAM_BINARY_SIZES> and returns the result.
 
 =for gengetinfo end program
 
@@ -1725,23 +1929,23 @@
 
 =item $string = $kernel->function_name
 
-Calls C<clGetKernelInfo> with C<CL_KERNEL_FUNCTION_NAME> and returns the result.
+Calls C<clGetKernelInfo> with C<OpenCL::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.
+Calls C<clGetKernelInfo> with C<OpenCL::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.
+Calls C<clGetKernelInfo> with C<OpenCL::KERNEL_REFERENCE_COUNT> and returns the result.
 
 =item $ctx = $kernel->context
 
-Calls C<clGetKernelInfo> with C<CL_KERNEL_CONTEXT> and returns the result.
+Calls C<clGetKernelInfo> with C<OpenCL::KERNEL_CONTEXT> and returns the result.
 
 =item $program = $kernel->program
 
-Calls C<clGetKernelInfo> with C<CL_KERNEL_PROGRAM> and returns the result.
+Calls C<clGetKernelInfo> with C<OpenCL::KERNEL_PROGRAM> and returns the result.
 
 =for gengetinfo end kernel
 
@@ -1749,35 +1953,62 @@
 
 See C<< $platform->info >> for details.
 
-The reason this method is not called C<info> is that there already is an
-C<< ->info >> method.
-
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetKernelWorkGroupInfo.html>
 
 =for gengetinfo begin kernel_work_group
 
 =item $int = $kernel->work_group_size ($device)
 
-Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_WORK_GROUP_SIZE> and returns the result.
+Calls C<clGetKernelWorkGroupInfo> with C<OpenCL::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.
+Calls C<clGetKernelWorkGroupInfo> with C<OpenCL::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.
+Calls C<clGetKernelWorkGroupInfo> with C<OpenCL::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.
+Calls C<clGetKernelWorkGroupInfo> with C<OpenCL::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.
+Calls C<clGetKernelWorkGroupInfo> with C<OpenCL::KERNEL_PRIVATE_MEM_SIZE> and returns the result.
 
 =for gengetinfo end kernel_work_group
 
+=item $packed_value = $kernel->arg_info ($idx, $name)
+
+See C<< $platform->info >> for details.
+
+L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clGetKernelArgInfo.html>
+
+=for gengetinfo begin kernel_arg
+
+=item $kernel_arg_address_qualifier = $kernel->arg_address_qualifier ($idx)
+
+Calls C<clGetKernelArgInfo> with C<OpenCL::KERNEL_ARG_ADDRESS_QUALIFIER> and returns the result.
+
+=item $kernel_arg_access_qualifier = $kernel->arg_access_qualifier ($idx)
+
+Calls C<clGetKernelArgInfo> with C<OpenCL::KERNEL_ARG_ACCESS_QUALIFIER> and returns the result.
+
+=item $string = $kernel->arg_type_name ($idx)
+
+Calls C<clGetKernelArgInfo> with C<OpenCL::KERNEL_ARG_TYPE_NAME> and returns the result.
+
+=item $kernel_arg_type_qualifier = $kernel->arg_type_qualifier ($idx)
+
+Calls C<clGetKernelArgInfo> with C<OpenCL::KERNEL_ARG_TYPE_QUALIFIER> and returns the result.
+
+=item $string = $kernel->arg_name ($idx)
+
+Calls C<clGetKernelArgInfo> with C<OpenCL::KERNEL_ARG_NAME> and returns the result.
+
+=for gengetinfo end kernel_arg
+
 =item $kernel->setf ($format, ...)
 
 Sets the arguments of a kernel. Since OpenCL 1.1 doesn't have a generic
@@ -1898,23 +2129,23 @@
 
 =item $queue = $event->command_queue
 
-Calls C<clGetEventInfo> with C<CL_EVENT_COMMAND_QUEUE> and returns the result.
+Calls C<clGetEventInfo> with C<OpenCL::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.
+Calls C<clGetEventInfo> with C<OpenCL::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.
+Calls C<clGetEventInfo> with C<OpenCL::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.
+Calls C<clGetEventInfo> with C<OpenCL::EVENT_COMMAND_EXECUTION_STATUS> and returns the result.
 
 =item $ctx = $event->context
 
-Calls C<clGetEventInfo> with C<CL_EVENT_CONTEXT> and returns the result.
+Calls C<clGetEventInfo> with C<OpenCL::EVENT_CONTEXT> and returns the result.
 
 =for gengetinfo end event
 
@@ -1931,19 +2162,19 @@
 
 =item $ulong = $event->profiling_command_queued
 
-Calls C<clGetEventProfilingInfo> with C<CL_PROFILING_COMMAND_QUEUED> and returns the result.
+Calls C<clGetEventProfilingInfo> with C<OpenCL::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.
+Calls C<clGetEventProfilingInfo> with C<OpenCL::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.
+Calls C<clGetEventProfilingInfo> with C<OpenCL::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.
+Calls C<clGetEventProfilingInfo> with C<OpenCL::PROFILING_COMMAND_END> and returns the result.
 
 =for gengetinfo end profiling
 
@@ -1960,6 +2191,8 @@
 Sets the execution status of the user event. Can only be called once,
 either with OpenCL::COMPLETE or a negative number as status.
 
+execution_status: OpenCL::COMPLETE or a negative integer.
+
 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetUserEventStatus.html>
 
 =back
@@ -2022,27 +2255,25 @@
 
 =item $ptr = $mapped->ptr
 
-Returns the raw memory address of the mapped area - same as C<$mapped+0>.
+Returns the raw memory address of the mapped area.
 
 =item $mapped->set ($offset, $data)
 
 Replaces the data at the given C<$offset> in the memory area by the new
-C<$data>. This method is safer but slower than direct manipulation of
-C<$$mapped> with substr.
+C<$data>. This method is safer than direct manipulation of C<$mapped>
+because it does bounds-checking, but also slower.
 
 =item $data = $mapped->get ($offset, $length)
 
 Returns (without copying) a scalar representing the data at the given
 C<$offset> and C<$length> in the mapped memory area. This is the same as
-the following substr, except much slower:
+the following substr, except much slower;
 
    $data = substr $$mapped, $offset, $length
 
-#TODO: really?
-
 =cut
 
-sub get {
+sub OpenCL::Mapped::get {
    substr ${$_[0]}, $_[1], $_[2]
 }