--- OpenCL/OpenCL.pm 2011/11/17 03:42:03 1.16 +++ OpenCL/OpenCL.pm 2012/04/24 14:24:42 1.52 @@ -29,7 +29,7 @@ kernel function in a program you can then create an OpenCL::Kernel object which represents basically a function call with argument values. -OpenCL::Memory objects of various flavours: OpenCL::Buffers objects (flat +OpenCL::Memory objects of various flavours: OpenCL::Buffer objects (flat memory areas, think arrays or structs) and OpenCL::Image objects (think 2d or 3d array) for bulk data and input and output for kernels. @@ -54,15 +54,19 @@ http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/ +If you are into UML class diagrams, the following diagram might help - if +not, it will be mildly cobfusing: + + http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/classDiagram.html + Here's a tutorial from AMD (very AMD-centric, too), not sure how useful it is, but at least it's free of charge: http://developer.amd.com/zones/OpenCLZone/courses/Documents/Introduction_to_OpenCL_Programming%20Training_Guide%20%28201005%29.pdf -If you are into UML class diagrams, the following diagram might help - if -not, it will be mildly cofusing: +And here's NVIDIA's OpenCL Best Practises Guide: - http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/classDiagram.html + http://developer.download.nvidia.com/compute/cuda/3_2/toolkit/docs/OpenCL_Best_Practices_Guide.pdf =head1 BASIC WORKFLOW @@ -99,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 } } @@ -147,20 +151,15 @@ 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]; } '; - 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); - + my $prog = $ctx->build_program ($src); my $kernel = $prog->kernel ("squareit"); =head2 Create some input and output float buffers, then call the @@ -209,6 +208,99 @@ # 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 = <build_program ($src); + my $kernel = $prog->kernel ("juliatunnel"); + + # program compiled, kernel ready, now draw and loop + + for (my $time; ; ++$time) { + # acquire objects from opengl + $queue->enqueue_acquire_gl_objects ([$tex]); + + # configure and run our kernel + $kernel->set_image2d (0, $tex); + $kernel->set_float (1, $time); + $queue->enqueue_nd_range_kernel ($kernel, undef, [256, 256], undef); + + # release objects to opengl again + $queue->enqueue_release_gl_objects ([$tex]); + + # wait + $queue->finish; + + # now draw the texture, the defaults should be all right + glTexParameterf GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST; + + glEnable GL_TEXTURE_2D; + glBegin GL_QUADS; + glTexCoord2f 0, 1; glVertex3i -1, -1, -1; + glTexCoord2f 0, 0; glVertex3i 1, -1, -1; + glTexCoord2f 1, 0; glVertex3i 1, 1, -1; + glTexCoord2f 1, 1; glVertex3i -1, 1, -1; + glEnd; + + glXSwapBuffers; + + select undef, undef, undef, 1/60; + } + =head1 DOCUMENTATION =head2 BASIC CONVENTIONS @@ -224,23 +316,24 @@ to free objects explicitly (C), the release function is called automatically once all Perl references to it go away. -=item * OpenCL uses CamelCase for function names (C), -while this module uses underscores as word separator and often leaves out -prefixes (C<< $platform->info >>). +=item * OpenCL uses CamelCase for function names +(e.g. C, C), while this module +uses underscores as word separator and often leaves out prefixes +(C, C<< $platform->info >>). =item * OpenCL often specifies fixed vector function arguments as short -arrays (C<$origin[3]>), while this module explicitly expects the -components as separate arguments- +arrays (C), while this module explicitly expects the +components as separate arguments (C<$orig_x, $orig_y, $orig_z>) in +function calls. -=item * Structures are often specified with their components, and returned -as arrayrefs. - -=item * Where possible, one of the pitch values is calculated from the -perl scalar length and need not be specified. +=item * Structures are often specified by flattening out their components +as with short vectors, and returned as arrayrefs. =item * When enqueuing commands, the wait list is specified by adding extra arguments to the function - anywhere a C<$wait_events...> argument -is documented this can be any number of event objects. +is documented this can be any number of event objects. As an extsnion +implemented by this module, C values will be ignored in the event +list. =item * When enqueuing commands, if the enqueue method is called in void context, no event is created. In all other contexts an event is returned @@ -270,6 +363,58 @@ 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 and C +specially: If either or both of these are C, then the OpenCL +module tries to dynamically resolve C and +C, 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 must be available and +capable of finding the function via C. + +=cut + +package OpenCL; + +use common::sense; + +BEGIN { + our $VERSION = '0.96'; + + require XSLoader; + XSLoader::load (__PACKAGE__, $VERSION); + + @OpenCL::Platform::ISA = + @OpenCL::Device::ISA = + @OpenCL::Context::ISA = + @OpenCL::Queue::ISA = + @OpenCL::Memory::ISA = + @OpenCL::Sampler::ISA = + @OpenCL::Program::ISA = + @OpenCL::Kernel::ISA = + @OpenCL::Event::ISA = OpenCL::Object::; + + @OpenCL::Buffer::ISA = + @OpenCL::Image::ISA = OpenCL::Memory::; + + @OpenCL::BufferObj::ISA = OpenCL::Buffer::; + + @OpenCL::Image2D::ISA = + @OpenCL::Image3D::ISA = + @OpenCL::Image2DArray::ISA = + @OpenCL::Image1D::ISA = + @OpenCL::Image1DArray::ISA = + @OpenCL::Image1DBuffer::ISA = OpenCL::Image::; + + @OpenCL::UserEvent::ISA = OpenCL::Event::; +} + =head2 THE OpenCL PACKAGE =over 4 @@ -285,9 +430,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. @@ -311,17 +456,32 @@ =back -=head2 THE OpenCL::Platform CLASS +=head2 THE OpenCL::Object CLASS + +This is the base class for all objects in the OpenCL module. The only +method it implements is the C method, which is only useful if you want +to interface to OpenCL on the C level. =over 4 -=item $packed_value = $platform->info ($name) +=item $iv = $obj->id -Calls C and returns the packed, raw value - for -strings, this will be the string, for other values you probably need to -use the correct C. This might get improved in the future. Hopefully. +OpenCL objects are represented by pointers or integers on the C level. If +you want to interface to an OpenCL object directly on the C level, then +you need this value, which is returned by this method. You should use an +C type in your code and cast that to the correct type. -L +=cut + +sub OpenCL::Object::id { + ${$_[0]} +} + +=back + +=head2 THE OpenCL::Platform CLASS + +=over 4 =item @devices = $platform->devices ($type = OpenCL::DEVICE_TYPE_ALL) @@ -329,17 +489,59 @@ =item $ctx = $platform->context_from_type ($properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $notify = undef) -Tries to create a context. Never worked for me, and you need devices explitly anyway. +Tries to create a context. Never worked for me, and you need devices explicitly anyway. L -=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. L +=item $packed_value = $platform->info ($name) + +Calls C and returns the packed, raw value - for +strings, this will be the string (possibly including terminating \0), for +other values you probably need to use the correct C. + +It's best to avoid this method and use one of the following convenience +wrappers. + +L + +=item $platform->unload_compiler + +Attempts to unload the compiler for this platform, for endless +profit. Does nothing on OpenCL 1.1. + +L + +=for gengetinfo begin platform + +=item $string = $platform->profile + +Calls C with C and returns the result. + +=item $string = $platform->version + +Calls C with C and returns the result. + +=item $string = $platform->name + +Calls C with C and returns the result. + +=item $string = $platform->vendor + +Calls C with C and returns the result. + +=item $string = $platform->extensions + +Calls C with C and returns the result. + +=for gengetinfo end platform + =back =head2 THE OpenCL::Device CLASS @@ -352,17 +554,306 @@ L +=for gengetinfo begin device + +=item $device_type = $device->type + +Calls C with C and returns the result. + +=item $uint = $device->vendor_id + +Calls C with C and returns the result. + +=item $uint = $device->max_compute_units + +Calls C with C and returns the result. + +=item $uint = $device->max_work_item_dimensions + +Calls C with C and returns the result. + +=item $int = $device->max_work_group_size + +Calls C with C and returns the result. + +=item @ints = $device->max_work_item_sizes + +Calls C with C and returns the result. + +=item $uint = $device->preferred_vector_width_char + +Calls C with C and returns the result. + +=item $uint = $device->preferred_vector_width_short + +Calls C with C and returns the result. + +=item $uint = $device->preferred_vector_width_int + +Calls C with C and returns the result. + +=item $uint = $device->preferred_vector_width_long + +Calls C with C and returns the result. + +=item $uint = $device->preferred_vector_width_float + +Calls C with C and returns the result. + +=item $uint = $device->preferred_vector_width_double + +Calls C with C and returns the result. + +=item $uint = $device->max_clock_frequency + +Calls C with C and returns the result. + +=item $bitfield = $device->address_bits + +Calls C with C and returns the result. + +=item $uint = $device->max_read_image_args + +Calls C with C and returns the result. + +=item $uint = $device->max_write_image_args + +Calls C with C and returns the result. + +=item $ulong = $device->max_mem_alloc_size + +Calls C with C and returns the result. + +=item $int = $device->image2d_max_width + +Calls C with C and returns the result. + +=item $int = $device->image2d_max_height + +Calls C with C and returns the result. + +=item $int = $device->image3d_max_width + +Calls C with C and returns the result. + +=item $int = $device->image3d_max_height + +Calls C with C and returns the result. + +=item $int = $device->image3d_max_depth + +Calls C with C and returns the result. + +=item $uint = $device->image_support + +Calls C with C and returns the result. + +=item $int = $device->max_parameter_size + +Calls C with C and returns the result. + +=item $uint = $device->max_samplers + +Calls C with C and returns the result. + +=item $uint = $device->mem_base_addr_align + +Calls C with C and returns the result. + +=item $uint = $device->min_data_type_align_size + +Calls C with C and returns the result. + +=item $device_fp_config = $device->single_fp_config + +Calls C with C and returns the result. + +=item $device_mem_cache_type = $device->global_mem_cache_type + +Calls C with C and returns the result. + +=item $uint = $device->global_mem_cacheline_size + +Calls C with C and returns the result. + +=item $ulong = $device->global_mem_cache_size + +Calls C with C and returns the result. + +=item $ulong = $device->global_mem_size + +Calls C with C and returns the result. + +=item $ulong = $device->max_constant_buffer_size + +Calls C with C and returns the result. + +=item $uint = $device->max_constant_args + +Calls C with C and returns the result. + +=item $device_local_mem_type = $device->local_mem_type + +Calls C with C and returns the result. + +=item $ulong = $device->local_mem_size + +Calls C with C and returns the result. + +=item $boolean = $device->error_correction_support + +Calls C with C and returns the result. + +=item $int = $device->profiling_timer_resolution + +Calls C with C and returns the result. + +=item $boolean = $device->endian_little + +Calls C with C and returns the result. + +=item $boolean = $device->available + +Calls C with C and returns the result. + +=item $boolean = $device->compiler_available + +Calls C with C and returns the result. + +=item $device_exec_capabilities = $device->execution_capabilities + +Calls C with C and returns the result. + +=item $command_queue_properties = $device->properties + +Calls C with C and returns the result. + +=item $ = $device->platform + +Calls C with C and returns the result. + +=item $string = $device->name + +Calls C with C and returns the result. + +=item $string = $device->vendor + +Calls C with C and returns the result. + +=item $string = $device->driver_version + +Calls C with C and returns the result. + +=item $string = $device->profile + +Calls C with C and returns the result. + +=item $string = $device->version + +Calls C with C and returns the result. + +=item $string = $device->extensions + +Calls C with C and returns the result. + +=item $uint = $device->preferred_vector_width_half + +Calls C with C and returns the result. + +=item $uint = $device->native_vector_width_char + +Calls C with C and returns the result. + +=item $uint = $device->native_vector_width_short + +Calls C with C and returns the result. + +=item $uint = $device->native_vector_width_int + +Calls C with C and returns the result. + +=item $uint = $device->native_vector_width_long + +Calls C with C and returns the result. + +=item $uint = $device->native_vector_width_float + +Calls C with C and returns the result. + +=item $uint = $device->native_vector_width_double + +Calls C with C and returns the result. + +=item $uint = $device->native_vector_width_half + +Calls C with C and returns the result. + +=item $device_fp_config = $device->double_fp_config + +Calls C with C and returns the result. + +=item $device_fp_config = $device->half_fp_config + +Calls C with C and returns the result. + +=item $boolean = $device->host_unified_memory + +Calls C with C and returns the result. + +=item $device = $device->parent_device_ext + +Calls C with C and returns the result. + +=item @device_partition_property_exts = $device->partition_types_ext + +Calls C with C and returns the result. + +=item @device_partition_property_exts = $device->affinity_domains_ext + +Calls C with C and returns the result. + +=item $uint = $device->reference_count_ext + +Calls C with C and returns the result. + +=item @device_partition_property_exts = $device->partition_style_ext + +Calls C with C and returns the result. + +=for gengetinfo end device + =back =head2 THE OpenCL::Context CLASS =over 4 -=item $packed_value = $ctx->info ($name) +=item $prog = $ctx->build_program ($program, $options = "") -See C<< $platform->info >> for details. +This convenience function tries to build the program on all devices in +the context. If the build fails, then the function will C with the +build log. Otherwise ti returns the program object. + +The C<$program> can either be a C object or a string +containing the program. In the latter case, a program objetc will be +created automatically. -L +=cut + +sub OpenCL::Context::build_program { + my ($self, $prog, $options) = @_; + + $prog = $self->program_with_source ($prog) + unless ref $prog; + + for my $dev ($self->devices) { + eval { $prog->build ($dev, $options); 1 } + or Carp::croak "Building OpenCL program for device '" . $dev->name . "' failed:\n" + . $prog->build_log ($dev); + } + + $prog +} =item $queue = $ctx->queue ($device, $properties) @@ -370,6 +861,10 @@ L +Example: create an out-of-order queue. + + $queue = $ctx->queue ($device, OpenCL::QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE); + =item $ev = $ctx->user_event Creates a new OpenCL::UserEvent object. @@ -378,26 +873,72 @@ =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 =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->image ($self, $flags, $channel_order, $channel_type, $type, $width, $height, $depth, $array_size = 0, $row_pitch = 0, $slice_pitch = 0, $num_mip_level = 0, $num_samples = 0, $*data = &PL_sv_undef) + +Creates a new OpenCL::Image object and optionally initialises it with +the given data values. + +L -=item $img = $ctx->image2d ($flags, $channel_order, $channel_type, $width, $height, $data) +=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 -=item $img = $ctx->image3d ($flags, $channel_order, $channel_type, $width, $height, $depth, $slice_pitch, $data) +=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 +=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 $img = $ctx->gl_texture ($flags, $target, $miplevel, $texture) + +Creates a new OpenCL::Image object that refers to the given OpenGL +texture object or buffer. + +http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateFromGLTexture.html + +=item $img = $ctx->gl_texture2d ($flags, $target, $miplevel, $texture) + +Creates a new OpenCL::Image2D object that refers to the given OpenGL +2D texture object. + +http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture2D.html + +=item $img = $ctx->gl_texture3d ($flags, $target, $miplevel, $texture) + +Creates a new OpenCL::Image3D object that refers to the given OpenGL +3D texture object. + +http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture3D.html + +=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 @@ -417,6 +958,32 @@ L +=item $packed_value = $ctx->info ($name) + +See C<< $platform->info >> for details. + +L + +=for gengetinfo begin context + +=item $uint = $context->reference_count + +Calls C with C and returns the result. + +=item @devices = $context->devices + +Calls C with C and returns the result. + +=item @property_ints = $context->properties + +Calls C with C and returns the result. + +=item $uint = $context->num_devices + +Calls C with C and returns the result. + +=for gengetinfo end context + =back =head2 THE OpenCL::Queue CLASS @@ -431,20 +998,20 @@ They also allow you to specify any number of other event objects that this request has to wait for before it starts executing, by simply passing the -event objects as extra parameters to the enqueue methods. +event objects as extra parameters to the enqueue methods. To simplify +program design, this module ignores any C values in the list of +events. This makes it possible to code operations such as this, without +having to put a valid event object into C<$event> first: + + $event = $queue->enqueue_xxx (..., $event); Queues execute in-order by default, without any parallelism, so in most cases (i.e. you use only one queue) it's not necessary to wait for or -create event objects. +create event objects, althoguh an our of order queue is often a bit +faster. =over 4 -=item $packed_value = $ctx->info ($name) - -See C<< $platform->info >> for details. - -L - =item $ev = $queue->enqueue_read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...) Reads data from buffer into the given string. @@ -461,23 +1028,25 @@ L -=item $ev = $queue->enqueue_read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...) - -L +=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, $data, $wait_events...) +http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBufferRect.html -L +=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, 4dst_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 =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. +L + +=item $ev = $queue->enqueue_write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...) + +L =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...) @@ -487,6 +1056,28 @@ L +=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. + +=item $ev = $queue->enqueue_fill_buffer ($mem, $pattern, $offset, $size, ...) + +Fills the given buffer object with repeated applications of C<$pattern>, +starting at C<$offset> for C<$size> octets. + +L + +=item $ev = $queue->enqueue_fill_image ($img, $r, $g, $b, $a, $x, $y, $z, $width, $height, $depth, ...) + +Fills the given image area with the given rgba colour components. The +components are normally floating point values between C<0> and C<1>, +except when the image channel data type is a signe dor unsigned +unnormalised format, in which case the range is determined by the format. + +L + =item $ev = $queue->enqueue_task ($kernel, $wait_events...) L @@ -509,17 +1100,31 @@ L -=item $ev = $queue->enqueue_marker +=item $ev = $queue->enqueue_acquire_gl_objects ([object, ...], $wait_events...) -L +Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired +for subsequent OpenCL usage. + +L + +=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 =item $ev = $queue->enqueue_wait_for_events ($wait_events...) L -=item $queue->enqueue_barrier +=item $ev = $queue->enqueue_marker ($wait_events...) -L +L + +=item $ev = $queue->enqueue_barrier ($wait_events...) + +L =item $queue->flush @@ -529,13 +1134,38 @@ L +=item $packed_value = $queue->info ($name) + +See C<< $platform->info >> for details. + +L + +=for gengetinfo begin command_queue + +=item $ctx = $command_queue->context + +Calls C with C and returns the result. + +=item $device = $command_queue->device + +Calls C with C and returns the result. + +=item $uint = $command_queue->reference_count + +Calls C with C and returns the result. + +=item $command_queue_properties = $command_queue->properties + +Calls C with C and returns the result. + +=for gengetinfo end command_queue + =back =head2 THE OpenCL::Memory CLASS This the superclass of all memory objects - OpenCL::Buffer, OpenCL::Image, -OpenCL::Image2D and OpenCL::Image3D. The subclasses of this class -currently only exist to allow type-checking. +OpenCL::Image2D and OpenCL::Image3D. =over 4 @@ -545,33 +1175,181 @@ L +=for gengetinfo begin mem + +=item $mem_object_type = $mem->type + +Calls C with C and returns the result. + +=item $mem_flags = $mem->flags + +Calls C with C and returns the result. + +=item $int = $mem->size + +Calls C with C and returns the result. + +=item $ptr_value = $mem->host_ptr + +Calls C with C and returns the result. + +=item $uint = $mem->map_count + +Calls C with C and returns the result. + +=item $uint = $mem->reference_count + +Calls C with C and returns the result. + +=item $ctx = $mem->context + +Calls C with C and returns the result. + +=item $mem = $mem->associated_memobject + +Calls C with C and returns the result. + +=item $int = $mem->offset + +Calls C with C 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 + =back -=head2 THE OpenCL::Sampler CLASS +=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 $packed_value = $sampler->info ($name) +=item $subbuf = $buf_obj->sub_buffer_region ($flags, $origin, $size) + +Creates an OpenCL::Buffer objects from this buffer and returns it. The +C is assumed to be C. + +L + +=back + +=head2 THE OpenCL::Image CLASS + +This is the superclass of all image objects - OpenCL::Image1D, +OpenCL::Image1DArray, OpenCL::Image1DBuffer, OpenCL::Image2D, +OpenCL::Image2DArray and OpenCL::Image3D. + +=over 4 + +=item $packed_value = $ev->image_info ($name) See C<< $platform->info >> for details. -L +The reason this method is not called C is that there already is an +C<< ->info >> method inherited from C. + +L + +=for gengetinfo begin image + +=item $int = $image->element_size + +Calls C with C and returns the result. + +=item $int = $image->row_pitch + +Calls C with C and returns the result. + +=item $int = $image->slice_pitch + +Calls C with C and returns the result. + +=item $int = $image->width + +Calls C with C and returns the result. + +=item $int = $image->height + +Calls C with C and returns the result. + +=item $int = $image->depth + +Calls C with C and returns the result. + +=for gengetinfo end image + +=for gengetinfo begin gl_texture + +=item $GLenum = $gl_texture->target + +Calls C with C and returns the result. + +=item $GLint = $gl_texture->gl_mipmap_level + +Calls C with C and returns the result. + +=for gengetinfo end gl_texture =back -=head2 THE OpenCL::Program CLASS +=head2 THE OpenCL::Sampler CLASS =over 4 -=item $packed_value = $program->info ($name) +=item $packed_value = $sampler->info ($name) See C<< $platform->info >> for details. -L +L + +=for gengetinfo begin sampler + +=item $uint = $sampler->reference_count + +Calls C with C and returns the result. + +=item $ctx = $sampler->context + +Calls C with C and returns the result. + +=item $addressing_mode = $sampler->normalized_coords + +Calls C with C and returns the result. + +=item $filter_mode = $sampler->addressing_mode + +Calls C with C and returns the result. + +=item $boolean = $sampler->filter_mode + +Calls C with C and returns the result. + +=for gengetinfo end sampler + +=back + +=head2 THE OpenCL::Program CLASS + +=over 4 =item $program->build ($device, $options = "") -Tries to build the program with the givne options. +Tries to build the program with the given options. See also the +C<$ctx->build> convenience function. L @@ -589,6 +1367,73 @@ L +=item @kernels = $program->kernels_in_program + +Returns all kernels successfully compiled for all devices in program. + +http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernelsInProgram.html + +=for gengetinfo begin program_build + +=item $build_status = $program->build_status ($device) + +Calls C with C and returns the result. + +=item $string = $program->build_options ($device) + +Calls C with C and returns the result. + +=item $string = $program->build_log ($device) + +Calls C with C and returns the result. + +=for gengetinfo end program_build + +=item $packed_value = $program->info ($name) + +See C<< $platform->info >> for details. + +L + +=for gengetinfo begin program + +=item $uint = $program->reference_count + +Calls C with C and returns the result. + +=item $ctx = $program->context + +Calls C with C and returns the result. + +=item $uint = $program->num_devices + +Calls C with C and returns the result. + +=item @devices = $program->devices + +Calls C with C and returns the result. + +=item $string = $program->source + +Calls C with C and returns the result. + +=item @ints = $program->binary_sizes + +Calls C with C and returns the result. + +=for gengetinfo end program + +=item @blobs = $program->binaries + +Returns a string for the compiled binary for every device associated with +the program, empty strings indicate missing programs, and an empty result +means no program binaries are available. + +These "binaries" are often, in fact, informative low-level assembly +sources. + +L + =back =head2 THE OpenCL::Kernel CLASS @@ -601,18 +1446,80 @@ L +=for gengetinfo begin kernel + +=item $string = $kernel->function_name + +Calls C with C and returns the result. + +=item $uint = $kernel->num_args + +Calls C with C and returns the result. + +=item $uint = $kernel->reference_count + +Calls C with C and returns the result. + +=item $ctx = $kernel->context + +Calls C with C and returns the result. + +=item $program = $kernel->program + +Calls C with C and returns the result. + +=for gengetinfo end kernel + +=item $packed_value = $kernel->work_group_info ($device, $name) + +See C<< $platform->info >> for details. + +The reason this method is not called C is that there already is an +C<< ->info >> method. + +L + +=for gengetinfo begin kernel_work_group + +=item $int = $kernel->work_group_size ($device) + +Calls C with C and returns the result. + +=item @ints = $kernel->compile_work_group_size ($device) + +Calls C with C and returns the result. + +=item $ulong = $kernel->local_mem_size ($device) + +Calls C with C and returns the result. + +=item $int = $kernel->preferred_work_group_size_multiple ($device) + +Calls C with C and returns the result. + +=item $ulong = $kernel->private_mem_size ($device) + +Calls C with C and returns the result. + +=for gengetinfo end kernel_work_group + =item $kernel->set_TYPE ($index, $value) This is a family of methods to set the kernel argument with the number C<$index> to the give C<$value>. TYPE is one of C, C, C, C, C, C, C, C, C, C, C, C, C, -C, C, C or C. +C, C, C, C or C. 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, and sampler and event must be -objects of that type. +must be an object of that type or C, local-memory arguments are +set by specifying the size, and sampler and event must be objects of that +type. + +Setting an argument for a kernel does NOT keep a reference to the object - +for example, if you set an argument to some image object, free the image, +and call the kernel, you will run into undefined behaviour. L @@ -625,17 +1532,70 @@ =over 4 +=item $ev->wait + +Waits for the event to complete. + +L + =item $packed_value = $ev->info ($name) See C<< $platform->info >> for details. L -=item $ev->wait +=for gengetinfo begin event -Waits for the event to complete. +=item $queue = $event->command_queue -L +Calls C with C and returns the result. + +=item $command_type = $event->command_type + +Calls C with C and returns the result. + +=item $uint = $event->reference_count + +Calls C with C and returns the result. + +=item $uint = $event->command_execution_status + +Calls C with C and returns the result. + +=item $ctx = $event->context + +Calls C with C and returns the result. + +=for gengetinfo end event + +=item $packed_value = $ev->profiling_info ($name) + +See C<< $platform->info >> for details. + +The reason this method is not called C is that there already is an +C<< ->info >> method. + +L + +=for gengetinfo begin profiling + +=item $ulong = $event->profiling_command_queued + +Calls C with C and returns the result. + +=item $ulong = $event->profiling_command_submit + +Calls C with C and returns the result. + +=item $ulong = $event->profiling_command_start + +Calls C with C and returns the result. + +=item $ulong = $event->profiling_command_end + +Calls C with C and returns the result. + +=for gengetinfo end profiling =back @@ -653,25 +1613,6 @@ =cut -package OpenCL; - -use common::sense; - -BEGIN { - our $VERSION = '0.14'; - - require XSLoader; - XSLoader::load (__PACKAGE__, $VERSION); - - @OpenCL::Buffer::ISA = - @OpenCL::Image::ISA = OpenCL::Memory::; - - @OpenCL::Image2D::ISA = - @OpenCL::Image3D::ISA = OpenCL::Image::; - - @OpenCL::UserEvent::ISA = OpenCL::Event::; -} - 1; =head1 AUTHOR