ViewVC Help
View File | Revision Log | Show Annotations | Download File
/cvs/OpenCL/OpenCL.pm
(Generate patch)

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.32 by root, Thu Apr 19 12:12:03 2012 UTC vs.
Revision 1.60 by root, Wed Apr 25 21:48:44 2012 UTC

43 43
44OpenCL::Event objects are used to signal when something is complete. 44OpenCL::Event objects are used to signal when something is complete.
45 45
46=head2 HELPFUL RESOURCES 46=head2 HELPFUL RESOURCES
47 47
48The OpenCL spec used to develop this module (1.2 spec was available, but 48The OpenCL specs used to develop this module:
49no implementation was available to me :).
50 49
51 http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf 50 http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
51 http://www.khronos.org/registry/cl/specs/opencl-1.2.pdf
52 http://www.khronos.org/registry/cl/specs/opencl-1.2-extensions.pdf
52 53
53OpenCL manpages: 54OpenCL manpages:
54 55
55 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/ 56 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/
57 http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/
56 58
57If you are into UML class diagrams, the following diagram might help - if 59If you are into UML class diagrams, the following diagram might help - if
58not, it will be mildly cobfusing: 60not, it will be mildly confusing (also, the class hierarchy of this module
61is much more fine-grained):
59 62
60 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/classDiagram.html 63 http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/classDiagram.html
61 64
62Here's a tutorial from AMD (very AMD-centric, too), not sure how useful it 65Here's a tutorial from AMD (very AMD-centric, too), not sure how useful it
63is, but at least it's free of charge: 66is, but at least it's free of charge:
64 67
65 http://developer.amd.com/zones/OpenCLZone/courses/Documents/Introduction_to_OpenCL_Programming%20Training_Guide%20%28201005%29.pdf 68 http://developer.amd.com/zones/OpenCLZone/courses/Documents/Introduction_to_OpenCL_Programming%20Training_Guide%20%28201005%29.pdf
138=head2 Create a buffer with some predefined data, read it back synchronously, 141=head2 Create a buffer with some predefined data, read it back synchronously,
139then asynchronously. 142then asynchronously.
140 143
141 my $buf = $ctx->buffer_sv (OpenCL::MEM_COPY_HOST_PTR, "helmut"); 144 my $buf = $ctx->buffer_sv (OpenCL::MEM_COPY_HOST_PTR, "helmut");
142 145
143 $queue->enqueue_read_buffer ($buf, 1, 1, 3, my $data); 146 $queue->read_buffer ($buf, 1, 1, 3, my $data);
144 print "$data\n"; 147 print "$data\n";
145 148
146 my $ev = $queue->enqueue_read_buffer ($buf, 0, 1, 3, my $data); 149 my $ev = $queue->read_buffer ($buf, 0, 1, 3, my $data);
147 $ev->wait; 150 $ev->wait;
148 print "$data\n"; # prints "elm" 151 print "$data\n"; # prints "elm"
149 152
150=head2 Create and build a program, then create a kernel out of one of its 153=head2 Create and build a program, then create a kernel out of one of its
151functions. 154functions.
157 $id = get_global_id (0); 160 $id = get_global_id (0);
158 output [id] = input [id] * input [id]; 161 output [id] = input [id] * input [id];
159 } 162 }
160 '; 163 ';
161 164
162 my $prog = $ctx->program_with_source ($src); 165 my $prog = $ctx->build_program ($src);
163
164 # build croaks on compile errors, so catch it and print the compile errors
165 eval { $prog->build ($dev); 1 }
166 or die $prog->build_log;
167
168 my $kernel = $prog->kernel ("squareit"); 166 my $kernel = $prog->kernel ("squareit");
169 167
170=head2 Create some input and output float buffers, then call the 168=head2 Create some input and output float buffers, then call the
171'squareit' kernel on them. 169'squareit' kernel on them.
172 170
176 # set buffer 174 # set buffer
177 $kernel->set_buffer (0, $input); 175 $kernel->set_buffer (0, $input);
178 $kernel->set_buffer (1, $output); 176 $kernel->set_buffer (1, $output);
179 177
180 # execute it for all 4 numbers 178 # execute it for all 4 numbers
181 $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); 179 $queue->nd_range_kernel ($kernel, undef, [4], undef);
182 180
183 # enqueue a synchronous read 181 # enqueue a synchronous read
184 $queue->enqueue_read_buffer ($output, 1, 0, OpenCL::SIZEOF_FLOAT * 4, my $data); 182 $queue->read_buffer ($output, 1, 0, OpenCL::SIZEOF_FLOAT * 4, my $data);
185 183
186 # print the results: 184 # print the results:
187 printf "%s\n", join ", ", unpack "f*", $data; 185 printf "%s\n", join ", ", unpack "f*", $data;
188 186
189=head2 The same enqueue operations as before, but assuming an out-of-order queue, 187=head2 The same enqueue operations as before, but assuming an out-of-order queue,
190showing off barriers. 188showing off barriers.
191 189
192 # execute it for all 4 numbers 190 # execute it for all 4 numbers
193 $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); 191 $queue->nd_range_kernel ($kernel, undef, [4], undef);
194 192
195 # enqueue a barrier to ensure in-order execution 193 # enqueue a barrier to ensure in-order execution
196 $queue->enqueue_barrier; 194 $queue->barrier;
197 195
198 # enqueue an async read 196 # enqueue an async read
199 $queue->enqueue_read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data); 197 $queue->read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data);
200 198
201 # wait for all requests to finish 199 # wait for all requests to finish
202 $queue->finish; 200 $queue->finish;
203 201
204=head2 The same enqueue operations as before, but assuming an out-of-order queue, 202=head2 The same enqueue operations as before, but assuming an out-of-order queue,
205showing off event objects and wait lists. 203showing off event objects and wait lists.
206 204
207 # execute it for all 4 numbers 205 # execute it for all 4 numbers
208 my $ev = $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); 206 my $ev = $queue->nd_range_kernel ($kernel, undef, [4], undef);
209 207
210 # enqueue an async read 208 # enqueue an async read
211 $ev = $queue->enqueue_read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data, $ev); 209 $ev = $queue->read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data, $ev);
212 210
213 # wait for the last event to complete 211 # wait for the last event to complete
214 $ev->wait; 212 $ev->wait;
213
214=head2 Use the OpenGL module to share a texture between OpenCL and OpenGL and draw some julia
215set tunnel effect.
216
217This is quite a long example to get you going.
218
219 use OpenGL ":all";
220 use OpenCL;
221
222 # open a window and create a gl texture
223 OpenGL::glpOpenWindow width => 256, height => 256;
224 my $texid = glGenTextures_p 1;
225 glBindTexture GL_TEXTURE_2D, $texid;
226 glTexImage2D_c GL_TEXTURE_2D, 0, GL_RGBA8, 256, 256, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0;
227
228 # find and use the first opencl device that let's us get a shared opengl context
229 my $platform;
230 my $dev;
231 my $ctx;
232
233 for (OpenCL::platforms) {
234 $platform = $_;
235 for ($platform->devices) {
236 $dev = $_;
237 $ctx = $platform->context ([OpenCL::GLX_DISPLAY_KHR, undef, OpenCL::GL_CONTEXT_KHR, undef], [$dev])
238 and last;
239 }
240 }
241
242 $ctx
243 or die "cannot find suitable OpenCL device\n";
244
245 my $queue = $ctx->queue ($dev);
246
247 # now attach an opencl image2d object to the opengl texture
248 my $tex = $ctx->gl_texture2d (OpenCL::MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, $texid);
249
250 # now the boring opencl code
251 my $src = <<EOF;
252 kernel void
253 juliatunnel (write_only image2d_t img, float time)
254 {
255 float2 p = (float2)(get_global_id (0), get_global_id (1)) / 256.f * 2.f - 1.f;
256
257 float2 m = (float2)(1.f, p.y) / fabs (p.x);
258 m.x = fabs (fmod (m.x + time * 0.05f, 4.f)) - 2.f;
259
260 float2 z = m;
261 float2 c = (float2)(sin (time * 0.05005), cos (time * 0.06001));
262
263 for (int i = 0; i < 25 && dot (z, z) < 4.f; ++i)
264 z = (float2)(z.x * z.x - z.y * z.y, 2.f * z.x * z.y) + c;
265
266 float3 colour = (float3)(z.x, z.y, z.x * z.y);
267 write_imagef (img, (int2)(get_global_id (0), get_global_id (1)), (float4)(colour * p.x * p.x, 1.));
268 }
269 EOF
270
271 my $prog = $ctx->build_program ($src);
272 my $kernel = $prog->kernel ("juliatunnel");
273
274 # program compiled, kernel ready, now draw and loop
275
276 for (my $time; ; ++$time) {
277 # acquire objects from opengl
278 $queue->acquire_gl_objects ([$tex]);
279
280 # configure and run our kernel
281 $kernel->set_image2d (0, $tex);
282 $kernel->set_float (1, $time);
283 $queue->nd_range_kernel ($kernel, undef, [256, 256], undef);
284
285 # release objects to opengl again
286 $queue->release_gl_objects ([$tex]);
287
288 # wait
289 $queue->finish;
290
291 # now draw the texture, the defaults should be all right
292 glTexParameterf GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST;
293
294 glEnable GL_TEXTURE_2D;
295 glBegin GL_QUADS;
296 glTexCoord2f 0, 1; glVertex3i -1, -1, -1;
297 glTexCoord2f 0, 0; glVertex3i 1, -1, -1;
298 glTexCoord2f 1, 0; glVertex3i 1, 1, -1;
299 glTexCoord2f 1, 1; glVertex3i -1, 1, -1;
300 glEnd;
301
302 glXSwapBuffers;
303
304 select undef, undef, undef, 1/60;
305 }
215 306
216=head1 DOCUMENTATION 307=head1 DOCUMENTATION
217 308
218=head2 BASIC CONVENTIONS 309=head2 BASIC CONVENTIONS
219 310
241=item * Structures are often specified by flattening out their components 332=item * Structures are often specified by flattening out their components
242as with short vectors, and returned as arrayrefs. 333as with short vectors, and returned as arrayrefs.
243 334
244=item * When enqueuing commands, the wait list is specified by adding 335=item * When enqueuing commands, the wait list is specified by adding
245extra arguments to the function - anywhere a C<$wait_events...> argument 336extra arguments to the function - anywhere a C<$wait_events...> argument
246is documented this can be any number of event objects. 337is documented this can be any number of event objects. As an extsnion
338implemented by this module, C<undef> values will be ignored in the event
339list.
247 340
248=item * When enqueuing commands, if the enqueue method is called in void 341=item * When enqueuing commands, if the enqueue method is called in void
249context, no event is created. In all other contexts an event is returned 342context, no event is created. In all other contexts an event is returned
250by the method. 343by the method.
251 344
271 ulong IV - Q 364 ulong IV - Q
272 float NV float f 365 float NV float f
273 half IV ushort S 366 half IV ushort S
274 double NV double d 367 double NV double d
275 368
276=head2 OpenGL sharing 369=head2 GLX SUPPORT
277 370
278This module can be optionally compiled with support for 371Due to the sad state that OpenGL support is in in Perl (mostly the OpenGL
279OpenGL sharing. The sharing functions are only available when 372module, which has little to no documentation and has little to no support
280C<OpenCL::HAVE_OPENGL> returns true, otherwise they are absent and cannot 373for glX), this module, as a special extension, treats context creation
281be called. 374properties C<OpenCL::GLX_DISPLAY_KHR> and C<OpenCL::GL_CONTEXT_KHR>
375specially: If either or both of these are C<undef>, then the OpenCL
376module tries to dynamically resolve C<glXGetCurrentDisplay> and
377C<glXGetCurrentContext>, call these functions and use their return values
378instead.
379
380For this to work, the OpenGL library must be loaded, a GLX context must
381have been created and be made current, and C<dlsym> must be available and
382capable of finding the function via C<RTLD_DEFAULT>.
383
384=head2 EVENT SYSTEM
385
386OpenCL can generate a number of (potentially) asynchronous events, for
387example, after compiling a program, to signal a context-related error or,
388perhaps most important, to signal completion of queued jobs (by setting
389callbacks on OpenCL::Event objects).
390
391To facilitate this, this module maintains an event queue - each
392time an asynchronous event happens, it is queued, and perl will be
393interrupted. This is implemented via the L<Async::Interrupt> module. In
394addition, this module has L<AnyEvent> support, so it can seamlessly
395integrate itself into many event loops.
396
397Since this module is a bit hard to understand, here are some case examples:
398
399=head3 Don't use callbacks.
400
401When your program never uses any callbacks, then there will never be any
402notifications you need to take care of, and therefore no need to worry
403about all this.
404
405You can achieve a great deal by explicitly waiting for events, or using
406barriers and flush calls. In many programs, there is no need at all to
407tinker with asynchronous events.
408
409=head3 Use AnyEvent
410
411This module automatically registers a watcher that invokes all outstanding
412event callbacks when AnyEvent is initialised (and block asynchronous
413interruptions). Using this mode of operations is the safest and most
414recommended one.
415
416To use this, simply use AnyEvent and this module normally, make sure you
417have an event loop running:
418
419 use Gtk2 -init;
420 use AnyEvent;
421
422 # initialise AnyEvent, by creating a watcher, or:
423 AnyEvent::detect;
424
425 my $e = $queue->marker;
426 $e->cb (sub {
427 warn "opencl is finished\n";
428 })
429
430 main Gtk2;
431
432Note that this module will not initialise AnyEvent for you. Before
433AnyEvent is initialised, the module will asynchronously interrupt perl
434instead. To avoid any surprises, it's best to explicitly initialise
435AnyEvent.
436
437You can temporarily enable asynchronous interruptions (see next paragraph)
438by calling C<$OpenCL::INTERRUPT->unblock> and disable them again by
439calling C<$OpenCL::INTERRUPT->block>.
440
441=head3 Let yourself be interrupted at any time
442
443This mode is the default unless AnyEvent is loaded and initialised. In
444this mode, OpenCL asynchronously interrupts a running perl program. The
445emphasis is on both I<asynchronously> and I<running> here.
446
447Asynchronously means that perl might execute your callbacks at any
448time. For example, in the following code (I<THAT YOU SHOULD NOT COPY>),
449the C<until> loop following the marker call will be interrupted by the
450callback:
451
452 my $e = $queue->marker;
453 my $flag;
454 $e->cb (sub { $flag = 1 });
455 1 until $flag;
456 # $flag is now 1
457
458The reason why you shouldn't blindly copy the above code is that
459busy waiting is a really really bad thing, and really really bad for
460performance.
461
462While at first this asynchronous business might look exciting, it can be
463really hard, because you need to be prepared for the callback code to be
464executed at any time, which limits the amount of things the callback code
465can do safely.
466
467This can be mitigated somewhat by using C<<
468$OpenCL::INTERRUPT->scope_block >> (see the L<Async::Interrupt>
469documentation for details).
470
471The other problem is that your program must be actively I<running> to be
472interrupted. When you calculate stuff, your program is running. When you
473hang in some C functions or other block execution (by calling C<sleep>,
474C<select>, running an event loop and so on), your program is waiting, not
475running.
476
477One way around that would be to attach a read watcher to your event loop,
478listening for events on C<< $OpenCL::INTERRUPT->pipe_fileno >>, using a
479dummy callback (C<sub { }>) to temporarily execute some perl code.
480
481That is then awfully close to using the built-in AnyEvent support above,
482though, so consider that one instead.
483
484=head3 Be creative
485
486OpenCL exports the L<Async::Interrupt> object it uses in the global
487variable C<$OpenCL::INTERRUPT>. You can configure it in any way you like.
488
489So if you want to feel like a real pro, err, wait, if you feel no risk
490menas no fun, you can experiment by implementing your own mode of
491operations.
492
493=cut
494
495package OpenCL;
496
497use common::sense;
498use Async::Interrupt ();
499
500our $POLL_FUNC; # set by XS
501
502BEGIN {
503 our $VERSION = '0.97';
504
505 require XSLoader;
506 XSLoader::load (__PACKAGE__, $VERSION);
507
508 @OpenCL::Platform::ISA =
509 @OpenCL::Device::ISA =
510 @OpenCL::Context::ISA =
511 @OpenCL::Queue::ISA =
512 @OpenCL::Memory::ISA =
513 @OpenCL::Sampler::ISA =
514 @OpenCL::Program::ISA =
515 @OpenCL::Kernel::ISA =
516 @OpenCL::Event::ISA = OpenCL::Object::;
517
518 @OpenCL::Buffer::ISA =
519 @OpenCL::Image::ISA = OpenCL::Memory::;
520
521 @OpenCL::BufferObj::ISA = OpenCL::Buffer::;
522
523 @OpenCL::Image2D::ISA =
524 @OpenCL::Image3D::ISA =
525 @OpenCL::Image2DArray::ISA =
526 @OpenCL::Image1D::ISA =
527 @OpenCL::Image1DArray::ISA =
528 @OpenCL::Image1DBuffer::ISA = OpenCL::Image::;
529
530 @OpenCL::UserEvent::ISA = OpenCL::Event::;
531}
282 532
283=head2 THE OpenCL PACKAGE 533=head2 THE OpenCL PACKAGE
284 534
285=over 4 535=over 4
286 536
305 555
306Returns all available OpenCL::Platform objects. 556Returns all available OpenCL::Platform objects.
307 557
308L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformIDs.html> 558L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformIDs.html>
309 559
310=item $ctx = OpenCL::context_from_type $properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $notify = undef 560=item $ctx = OpenCL::context_from_type $properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $callback->($err, $pvt) = $print_stderr
311 561
312Tries to create a context from a default device and platform - never worked for me. 562Tries to create a context from a default device and platform type - never worked for me.
313 563
314L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html> 564L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html>
315 565
566=item $ctx = OpenCL::context $properties, \@devices, $callback->($err, $pvt) = $print_stderr)
567
568Create a new OpenCL::Context object using the given device object(s). This
569function isn't implemented yet, use C<< $platform->context >> instead.
570
571L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContext.html>
572
316=item OpenCL::wait_for_events $wait_events... 573=item OpenCL::wait_for_events $wait_events...
317 574
318Waits for all events to complete. 575Waits for all events to complete.
319 576
320L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html> 577L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html>
321 578
579=item OpenCL::poll
580
581Checks if there are any outstanding events (see L<EVENT SYSTEM>) and
582invokes their callbacks.
583
584=item $OpenCL::INTERRUPT
585
586The L<Async::Interrupt> object used to signal asynchronous events (see
587L<EVENT SYSTEM>).
588
589=cut
590
591our $INTERRUPT = new Async::Interrupt c_cb => [$POLL_FUNC, 0];
592
593&_eq_initialise ($INTERRUPT->signal_func);
594
595=item $OpenCL::WATCHER
596
597The L<AnyEvent> watcher object used to watch for asynchronous events (see
598L<EVENT SYSTEM>). This variable is C<undef> until L<AnyEvent> has been
599loaded I<and> initialised (e.g. by calling C<AnyEvent::detect>).
600
601=cut
602
603our $WATCHER;
604
605sub _init_anyevent {
606 $INTERRUPT->block;
607 $WATCHER = AE::io ($INTERRUPT->pipe_fileno, 0, sub { $INTERRUPT->handle });
608}
609
610if (defined $AnyEvent::MODEL) {
611 _init_anyevent;
612} else {
613 push @AnyEvent::post_detect, \&_init_anyevent;
614}
615
322=back 616=back
323 617
618=head2 THE OpenCL::Object CLASS
619
620This is the base class for all objects in the OpenCL module. The only
621method it implements is the C<id> method, which is only useful if you want
622to interface to OpenCL on the C level.
623
624=over 4
625
626=item $iv = $obj->id
627
628OpenCL objects are represented by pointers or integers on the C level. If
629you want to interface to an OpenCL object directly on the C level, then
630you need this value, which is returned by this method. You should use an
631C<IV> type in your code and cast that to the correct type.
632
633=cut
634
635sub OpenCL::Object::id {
636 ref $_[0] eq "SCALAR"
637 ? ${ $_[0] }
638 : $_[0][0]
639}
640
641=back
642
324=head2 THE OpenCL::Platform CLASS 643=head2 THE OpenCL::Platform CLASS
325 644
326=over 4 645=over 4
327 646
328=item @devices = $platform->devices ($type = OpenCL::DEVICE_TYPE_ALL) 647=item @devices = $platform->devices ($type = OpenCL::DEVICE_TYPE_ALL)
329 648
330Returns a list of matching OpenCL::Device objects. 649Returns a list of matching OpenCL::Device objects.
331 650
332=item $ctx = $platform->context_from_type ($properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $notify = undef) 651=item $ctx = $platform->context_from_type ($properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $callback->($err, $pvt) = $print_stderr)
333 652
334Tries to create a context. Never worked for me, and you need devices explicitly anyway. 653Tries to create a context. Never worked for me, and you need devices explicitly anyway.
335 654
336L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html> 655L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html>
337 656
338=item $ctx = $platform->context ($properties = undef, @$devices, $notify = undef) 657=item $ctx = $platform->context ($properties, \@devices, $callback->($err, $pvt) = $print_stderr)
339 658
340Create a new OpenCL::Context object using the given device object(s)- a 659Create a new OpenCL::Context object using the given device object(s)- a
341CL_CONTEXT_PLATFORM property is supplied automatically. 660CL_CONTEXT_PLATFORM property is supplied automatically.
342 661
343L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContext.html> 662L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContext.html>
351It's best to avoid this method and use one of the following convenience 670It's best to avoid this method and use one of the following convenience
352wrappers. 671wrappers.
353 672
354L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformInfo.html> 673L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformInfo.html>
355 674
675=item $platform->unload_compiler
676
677Attempts to unload the compiler for this platform, for endless
678profit. Does nothing on OpenCL 1.1.
679
680L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clUnloadPlatformCompiler.html>
681
356=for gengetinfo begin platform 682=for gengetinfo begin platform
357 683
358=item $string = $platform->profile 684=item $string = $platform->profile
359 685
360Calls C<clGetPlatformInfo> with C<CL_PLATFORM_PROFILE> and returns the result. 686Calls C<clGetPlatformInfo> with C<CL_PLATFORM_PROFILE> and returns the result.
645 971
646=item @device_partition_property_exts = $device->affinity_domains_ext 972=item @device_partition_property_exts = $device->affinity_domains_ext
647 973
648Calls C<clGetDeviceInfo> with C<CL_DEVICE_AFFINITY_DOMAINS_EXT> and returns the result. 974Calls C<clGetDeviceInfo> with C<CL_DEVICE_AFFINITY_DOMAINS_EXT> and returns the result.
649 975
650=item $uint = $device->reference_count_ext 976=item $uint = $device->reference_count_ext
651 977
652Calls C<clGetDeviceInfo> with C<CL_DEVICE_REFERENCE_COUNT_EXT > and returns the result. 978Calls C<clGetDeviceInfo> with C<CL_DEVICE_REFERENCE_COUNT_EXT> and returns the result.
653 979
654=item @device_partition_property_exts = $device->partition_style_ext 980=item @device_partition_property_exts = $device->partition_style_ext
655 981
656Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARTITION_STYLE_EXT> and returns the result. 982Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARTITION_STYLE_EXT> and returns the result.
657 983
661 987
662=head2 THE OpenCL::Context CLASS 988=head2 THE OpenCL::Context CLASS
663 989
664=over 4 990=over 4
665 991
992=item $prog = $ctx->build_program ($program, $options = "")
993
994This convenience function tries to build the program on all devices in
995the context. If the build fails, then the function will C<croak> with the
996build log. Otherwise ti returns the program object.
997
998The C<$program> can either be a C<OpenCL::Program> object or a string
999containing the program. In the latter case, a program objetc will be
1000created automatically.
1001
1002=cut
1003
1004sub OpenCL::Context::build_program {
1005 my ($self, $prog, $options) = @_;
1006
1007 require Carp;
1008
1009 $prog = $self->program_with_source ($prog)
1010 unless ref $prog;
1011
1012 # we build separately per device so we instantly know which one failed
1013 for my $dev ($self->devices) {
1014 eval { $prog->build ([$dev], $options); 1 }
1015 or Carp::croak ("Building OpenCL program for device '" . $dev->name . "' failed:\n"
1016 . $prog->build_log ($dev));
1017 }
1018
1019 $prog
1020}
1021
666=item $queue = $ctx->queue ($device, $properties) 1022=item $queue = $ctx->queue ($device, $properties)
667 1023
668Create a new OpenCL::Queue object from the context and the given device. 1024Create a new OpenCL::Queue object from the context and the given device.
669 1025
670L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateCommandQueue.html> 1026L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateCommandQueue.html>
1027
1028Example: create an out-of-order queue.
1029
1030 $queue = $ctx->queue ($device, OpenCL::QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE);
671 1031
672=item $ev = $ctx->user_event 1032=item $ev = $ctx->user_event
673 1033
674Creates a new OpenCL::UserEvent object. 1034Creates a new OpenCL::UserEvent object.
675 1035
685=item $buf = $ctx->buffer_sv ($flags, $data) 1045=item $buf = $ctx->buffer_sv ($flags, $data)
686 1046
687Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object and 1047Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object and
688initialise it with the given data values. 1048initialise it with the given data values.
689 1049
1050=item $img = $ctx->image ($self, $flags, $channel_order, $channel_type, $type, $width, $height, $depth = 0, $array_size = 0, $row_pitch = 0, $slice_pitch = 0, $num_mip_level = 0, $num_samples = 0, $*data = &PL_sv_undef)
1051
1052Creates a new OpenCL::Image object and optionally initialises it with
1053the given data values.
1054
1055L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateImage.html>
1056
690=item $img = $ctx->image2d ($flags, $channel_order, $channel_type, $width, $height, $row_pitch = 0, $data = undef) 1057=item $img = $ctx->image2d ($flags, $channel_order, $channel_type, $width, $height, $row_pitch = 0, $data = undef)
691 1058
692Creates a new OpenCL::Image2D object and optionally initialises it with 1059Creates a new OpenCL::Image2D object and optionally initialises it with
693the given data values. 1060the given data values.
694 1061
699Creates a new OpenCL::Image3D object and optionally initialises it with 1066Creates a new OpenCL::Image3D object and optionally initialises it with
700the given data values. 1067the given data values.
701 1068
702L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateImage3D.html> 1069L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateImage3D.html>
703 1070
1071=item $buffer = $ctx->gl_buffer ($flags, $bufobj)
1072
1073Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object that refers to the given
1074OpenGL buffer object.
1075
1076http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLBuffer.html
1077
1078=item $img = $ctx->gl_texture ($flags, $target, $miplevel, $texture)
1079
1080Creates a new OpenCL::Image object that refers to the given OpenGL
1081texture object or buffer.
1082
1083http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateFromGLTexture.html
1084
1085=item $img = $ctx->gl_texture2d ($flags, $target, $miplevel, $texture)
1086
1087Creates a new OpenCL::Image2D object that refers to the given OpenGL
10882D texture object.
1089
1090http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture2D.html
1091
1092=item $img = $ctx->gl_texture3d ($flags, $target, $miplevel, $texture)
1093
1094Creates a new OpenCL::Image3D object that refers to the given OpenGL
10953D texture object.
1096
1097http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture3D.html
1098
1099=item $ctx->gl_renderbuffer ($flags, $renderbuffer)
1100
1101Creates a new OpenCL::Image2D object that refers to the given OpenGL
1102render buffer.
1103
1104http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLRenderbuffer.html
1105
704=item @formats = $ctx->supported_image_formats ($flags, $image_type) 1106=item @formats = $ctx->supported_image_formats ($flags, $image_type)
705 1107
706Returns a list of matching image formats - each format is an arrayref with 1108Returns a list of matching image formats - each format is an arrayref with
707two values, $channel_order and $channel_type, in it. 1109two values, $channel_order and $channel_type, in it.
708 1110
749=back 1151=back
750 1152
751=head2 THE OpenCL::Queue CLASS 1153=head2 THE OpenCL::Queue CLASS
752 1154
753An OpenCL::Queue represents an execution queue for OpenCL. You execute 1155An OpenCL::Queue represents an execution queue for OpenCL. You execute
754requests by calling their respective C<enqueue_xxx> method and waitinf for 1156requests by calling their respective method and waiting for it to complete
755it to complete in some way. 1157in some way.
756 1158
757All the enqueue methods return an event object that can be used to wait 1159Most methods that enqueue some request return an event object that can
758for completion, unless the method is called in void context, in which case 1160be used to wait for completion (optionally using a callback), unless
759no event object is created. 1161the method is called in void context, in which case no event object is
1162created.
760 1163
761They also allow you to specify any number of other event objects that this 1164They also allow you to specify any number of other event objects that this
762request has to wait for before it starts executing, by simply passing the 1165request has to wait for before it starts executing, by simply passing the
763event objects as extra parameters to the enqueue methods. 1166event objects as extra parameters to the enqueue methods. To simplify
1167program design, this module ignores any C<undef> values in the list of
1168events. This makes it possible to code operations such as this, without
1169having to put a valid event object into C<$event> first:
1170
1171 $event = $queue->xxx (..., $event);
764 1172
765Queues execute in-order by default, without any parallelism, so in most 1173Queues execute in-order by default, without any parallelism, so in most
766cases (i.e. you use only one queue) it's not necessary to wait for or 1174cases (i.e. you use only one queue) it's not necessary to wait for or
767create event objects. 1175create event objects, althoguh an our of order queue is often a bit
1176faster.
768 1177
769=over 4 1178=over 4
770 1179
771=item $ev = $queue->enqueue_read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...) 1180=item $ev = $queue->read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...)
772 1181
773Reads data from buffer into the given string. 1182Reads data from buffer into the given string.
774 1183
775L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBuffer.html> 1184L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBuffer.html>
776 1185
777=item $ev = $queue->enqueue_write_buffer ($buffer, $blocking, $offset, $data, $wait_events...) 1186=item $ev = $queue->write_buffer ($buffer, $blocking, $offset, $data, $wait_events...)
778 1187
779Writes data to buffer from the given string. 1188Writes data to buffer from the given string.
780 1189
781L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBuffer.html> 1190L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBuffer.html>
782 1191
783=item $ev = $queue->enqueue_copy_buffer ($src, $dst, $src_offset, $dst_offset, $len, $wait_events...) 1192=item $ev = $queue->copy_buffer ($src, $dst, $src_offset, $dst_offset, $len, $wait_events...)
784 1193
785L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBuffer.html> 1194L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBuffer.html>
786 1195
787=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...) 1196=item $ev = $queue->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...)
788 1197
789http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBufferRect.html 1198http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBufferRect.html
790 1199
791=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...) 1200=item $ev = $queue->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...)
792 1201
793http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBufferRect.html 1202http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBufferRect.html
794 1203
795=item $ev = $queue->enqueue_read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...) 1204=item $ev = $queue->read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
796 1205
797L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferRect.html> 1206L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferRect.html>
798 1207
799=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...) 1208=item $ev = $queue->copy_buffer_to_image ($src_buffer, $dst_image, $src_offset, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $wait_events...)
800 1209
801L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadImage.html> 1210L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadImage.html>
802 1211
803=item $ev = $queue->enqueue_write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...) 1212=item $ev = $queue->write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
804 1213
805L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteImage.html> 1214L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteImage.html>
806 1215
807=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...) 1216=item $ev = $queue->copy_image ($src_image, $dst_image, $src_x, $src_y, $src_z, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $wait_events...)
808 1217
809L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html> 1218L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html>
810 1219
811=item $ev = $queue->enqueue_copy_image_to_buffer ($src_image, $dst_image, $src_x, $src_y, $src_z, $width, $height, $depth, $dst_offset, $wait_events...) 1220=item $ev = $queue->copy_image_to_buffer ($src_image, $dst_image, $src_x, $src_y, $src_z, $width, $height, $depth, $dst_offset, $wait_events...)
812 1221
813L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html> 1222L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html>
814 1223
815=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...) 1224=item $ev = $queue->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...)
816 1225
817Yeah. 1226Yeah.
818 1227
819L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>. 1228L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>.
820 1229
1230=item $ev = $queue->fill_buffer ($mem, $pattern, $offset, $size, ...)
1231
1232Fills the given buffer object with repeated applications of C<$pattern>,
1233starting at C<$offset> for C<$size> octets.
1234
1235L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillBuffer.html>
1236
1237=item $ev = $queue->fill_image ($img, $r, $g, $b, $a, $x, $y, $z, $width, $height, $depth, ...)
1238
1239Fills the given image area with the given rgba colour components. The
1240components are normally floating point values between C<0> and C<1>,
1241except when the image channel data type is a signe dor unsigned
1242unnormalised format, in which case the range is determined by the format.
1243
1244L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillImage.html>
1245
821=item $ev = $queue->enqueue_task ($kernel, $wait_events...) 1246=item $ev = $queue->task ($kernel, $wait_events...)
822 1247
823L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html> 1248L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html>
824 1249
825=item $ev = $queue->enqueue_nd_range_kernel ($kernel, @$global_work_offset, @$global_work_size, @$local_work_size, $wait_events...) 1250=item $ev = $queue->nd_range_kernel ($kernel, \@global_work_offset, \@global_work_size, \@local_work_size, $wait_events...)
826 1251
827Enqueues a kernel execution. 1252Enqueues a kernel execution.
828 1253
829@$global_work_size must be specified as a reference to an array of 1254\@global_work_size must be specified as a reference to an array of
830integers specifying the work sizes (element counts). 1255integers specifying the work sizes (element counts).
831 1256
832@$global_work_offset must be either C<undef> (in which case all offsets 1257\@global_work_offset must be either C<undef> (in which case all offsets
833are C<0>), or a reference to an array of work offsets, with the same number 1258are C<0>), or a reference to an array of work offsets, with the same number
834of elements as @$global_work_size. 1259of elements as \@global_work_size.
835 1260
836@$local_work_size must be either C<undef> (in which case the 1261\@local_work_size must be either C<undef> (in which case the
837implementation is supposed to choose good local work sizes), or a 1262implementation is supposed to choose good local work sizes), or a
838reference to an array of local work sizes, with the same number of 1263reference to an array of local work sizes, with the same number of
839elements as @$global_work_size. 1264elements as \@global_work_size.
840 1265
841L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html> 1266L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html>
842 1267
843=item $ev = $queue->enqueue_marker 1268=item $ev = $queue->acquire_gl_objects ([object, ...], $wait_events...)
844 1269
1270Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired
1271for subsequent OpenCL usage.
1272
845L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMarker.html> 1273L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueAcquireGLObjects.html>
846 1274
1275=item $ev = $queue->release_gl_objects ([object, ...], $wait_events...)
1276
1277Enqueues a list (an array-ref of OpenCL::Memory objects) to be released
1278for subsequent OpenGL usage.
1279
1280L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReleaseGLObjects.html>
1281
847=item $ev = $queue->enqueue_wait_for_events ($wait_events...) 1282=item $ev = $queue->wait_for_events ($wait_events...)
848 1283
849L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html> 1284L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html>
850 1285
851=item $queue->enqueue_barrier 1286=item $ev = $queue->marker ($wait_events...)
852 1287
1288L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMarkerWithWaitList.html>
1289
1290=item $ev = $queue->barrier ($wait_events...)
1291
853L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueBarrier.html> 1292L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueBarrierWithWaitList.html>
854 1293
855=item $queue->flush 1294=item $queue->flush
856 1295
857L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clFlush.html> 1296L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clFlush.html>
858 1297
938=item $int = $mem->offset 1377=item $int = $mem->offset
939 1378
940Calls C<clGetMemObjectInfo> with C<CL_MEM_OFFSET> and returns the result. 1379Calls C<clGetMemObjectInfo> with C<CL_MEM_OFFSET> and returns the result.
941 1380
942=for gengetinfo end mem 1381=for gengetinfo end mem
1382
1383=item ($type, $name) = $mem->gl_object_info
1384
1385Returns the OpenGL object type (e.g. OpenCL::GL_OBJECT_TEXTURE2D) and the
1386object "name" (e.g. the texture name) used to create this memory object.
1387
1388L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetGLObjectInfo.html>
943 1389
944=back 1390=back
945 1391
946=head2 THE OpenCL::Buffer CLASS 1392=head2 THE OpenCL::Buffer CLASS
947 1393
966 1412
967=back 1413=back
968 1414
969=head2 THE OpenCL::Image CLASS 1415=head2 THE OpenCL::Image CLASS
970 1416
971This is the superclass of all image objects - OpenCL::Image2D and OpenCL::Image3D. 1417This is the superclass of all image objects - OpenCL::Image1D,
1418OpenCL::Image1DArray, OpenCL::Image1DBuffer, OpenCL::Image2D,
1419OpenCL::Image2DArray and OpenCL::Image3D.
972 1420
973=over 4 1421=over 4
974 1422
975=item $packed_value = $ev->image_info ($name) 1423=item $packed_value = $image->image_info ($name)
976 1424
977See C<< $platform->info >> for details. 1425See C<< $platform->info >> for details.
978 1426
979The reason this method is not called C<info> is that there already is an 1427The reason this method is not called C<info> is that there already is an
980C<< ->info >> method inherited from C<OpenCL::Memory>. 1428C<< ->info >> method inherited from C<OpenCL::Memory>.
981 1429
982L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetImageInfo.html> 1430L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetImageInfo.html>
983 1431
1432=item ($channel_order, $channel_data_type) = $image->format
1433
1434Returns the channel order and type used to create the image by calling
1435C<clGetImageInfo> with C<CL_IMAGE_FORMAT>.
1436
984=for gengetinfo begin image 1437=for gengetinfo begin image
985 1438
986=item $int = $image->element_size 1439=item $int = $image->element_size
987 1440
988Calls C<clGetImageInfo> with C<CL_IMAGE_ELEMENT_SIZE> and returns the result. 1441Calls C<clGetImageInfo> with C<CL_IMAGE_ELEMENT_SIZE> and returns the result.
1007 1460
1008Calls C<clGetImageInfo> with C<CL_IMAGE_DEPTH> and returns the result. 1461Calls C<clGetImageInfo> with C<CL_IMAGE_DEPTH> and returns the result.
1009 1462
1010=for gengetinfo end image 1463=for gengetinfo end image
1011 1464
1465=for gengetinfo begin gl_texture
1466
1467=item $GLenum = $gl_texture->target
1468
1469Calls C<clGetGLTextureInfo> with C<CL_GL_TEXTURE_TARGET> and returns the result.
1470
1471=item $GLint = $gl_texture->gl_mipmap_level
1472
1473Calls C<clGetGLTextureInfo> with C<CL_GL_MIPMAP_LEVEL> and returns the result.
1474
1475=for gengetinfo end gl_texture
1476
1012=back 1477=back
1013 1478
1014=head2 THE OpenCL::Sampler CLASS 1479=head2 THE OpenCL::Sampler CLASS
1015 1480
1016=over 4 1481=over 4
1049 1514
1050=head2 THE OpenCL::Program CLASS 1515=head2 THE OpenCL::Program CLASS
1051 1516
1052=over 4 1517=over 4
1053 1518
1054=item $program->build ($device, $options = "") 1519=item $program->build (\@devices = undef, $options = "", $cb->($program) = undef)
1055 1520
1056Tries to build the program with the givne options. 1521Tries to build the program with the given options. See also the
1522C<$ctx->build> convenience function.
1523
1524If a callback is specified, then it will be called when compilation is
1525finished. Note that many OpenCL implementations block your program while
1526compiling whether you use a callback or not. See C<build_async> if you
1527want to make sure the build is done in the background.
1528
1529Note that some OpenCL implementations atc up badly, and don't call the
1530callback in some error cases (but call it in others). This implementation
1531assumes the callback will always be called, and leaks memory if this is
1532not so. So best make sure you don't pass in invalid values.
1057 1533
1058L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clBuildProgram.html> 1534L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clBuildProgram.html>
1535
1536=item $program->build_async (\@devices = undef, $options = "", $cb->($program) = undef)
1537
1538Similar to C<< ->build >>, except it starts a thread, and never fails (you
1539need to check the compilation status form the callback, or by polling).
1059 1540
1060=item $packed_value = $program->build_info ($device, $name) 1541=item $packed_value = $program->build_info ($device, $name)
1061 1542
1062Similar to C<< $platform->info >>, but returns build info for a previous 1543Similar to C<< $platform->info >>, but returns build info for a previous
1063build attempt for the given device. 1544build attempt for the given device.
1068 1549
1069Creates an OpenCL::Kernel object out of the named C<__kernel> function in 1550Creates an OpenCL::Kernel object out of the named C<__kernel> function in
1070the program. 1551the program.
1071 1552
1072L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernel.html> 1553L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernel.html>
1554
1555=item @kernels = $program->kernels_in_program
1556
1557Returns all kernels successfully compiled for all devices in program.
1558
1559http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernelsInProgram.html
1073 1560
1074=for gengetinfo begin program_build 1561=for gengetinfo begin program_build
1075 1562
1076=item $build_status = $program->build_status ($device) 1563=item $build_status = $program->build_status ($device)
1077 1564
1199 1686
1200Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PRIVATE_MEM_SIZE> and returns the result. 1687Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PRIVATE_MEM_SIZE> and returns the result.
1201 1688
1202=for gengetinfo end kernel_work_group 1689=for gengetinfo end kernel_work_group
1203 1690
1691=item $kernel->setf ($format, ...)
1692
1693Sets the arguments of a kernel. Since OpenCL 1.1 doesn't have a generic
1694way to set arguments (and with OpenCL 1.2 it might be rather slow), you
1695need to specify a format argument, much as with C<printf>, to tell OpenCL
1696what type of argument it is.
1697
1698The format arguments are single letters:
1699
1700 c char
1701 C unsigned char
1702 s short
1703 S unsigned short
1704 i int
1705 I unsigned int
1706 l long
1707 L unsigned long
1708
1709 h half float (0..65535)
1710 f float
1711 d double
1712
1713 z local (octet size)
1714
1715 m memory object (buffer or image)
1716 a sampler
1717 e event
1718
1719Space characters in the format string are ignored.
1720
1721Example: set the arguments for a kernel that expects an int, two floats, a buffer and an image.
1722
1723 $kernel->setf ("i ff mm", 5, 0.5, 3, $buffer, $image);
1724
1204=item $kernel->set_TYPE ($index, $value) 1725=item $kernel->set_TYPE ($index, $value)
1205 1726
1727=item $kernel->set_char ($index, $value)
1728
1729=item $kernel->set_uchar ($index, $value)
1730
1731=item $kernel->set_short ($index, $value)
1732
1733=item $kernel->set_ushort ($index, $value)
1734
1735=item $kernel->set_int ($index, $value)
1736
1737=item $kernel->set_uint ($index, $value)
1738
1739=item $kernel->set_long ($index, $value)
1740
1741=item $kernel->set_ulong ($index, $value)
1742
1743=item $kernel->set_half ($index, $value)
1744
1745=item $kernel->set_float ($index, $value)
1746
1747=item $kernel->set_double ($index, $value)
1748
1749=item $kernel->set_memory ($index, $value)
1750
1751=item $kernel->set_buffer ($index, $value)
1752
1753=item $kernel->set_image ($index, $value)
1754
1755=item $kernel->set_sampler ($index, $value)
1756
1757=item $kernel->set_local ($index, $value)
1758
1759=item $kernel->set_event ($index, $value)
1760
1206This is a family of methods to set the kernel argument with the number C<$index> to the give C<$value>. 1761This is a family of methods to set the kernel argument with the number
1207 1762C<$index> to the give C<$value>.
1208TYPE is one of C<char>, C<uchar>, C<short>, C<ushort>, C<int>, C<uint>,
1209C<long>, C<ulong>, C<half>, C<float>, C<double>, C<memory>, C<buffer>,
1210C<image2d>, C<image3d>, C<sampler> or C<event>.
1211 1763
1212Chars and integers (including the half type) are specified as integers, 1764Chars and integers (including the half type) are specified as integers,
1213float and double as floating point values, memory/buffer/image2d/image3d 1765float and double as floating point values, memory/buffer/image must be
1214must be an object of that type or C<undef>, and sampler and event must be 1766an object of that type or C<undef>, local-memory arguments are set by
1215objects of that type. 1767specifying the size, and sampler and event must be objects of that type.
1768
1769Note that C<set_memory> works for all memory objects (all types of buffers
1770and images) - the main purpose of the more specific C<set_TYPE> functions
1771is type checking.
1772
1773Setting an argument for a kernel does NOT keep a reference to the object -
1774for example, if you set an argument to some image object, free the image,
1775and call the kernel, you will run into undefined behaviour.
1216 1776
1217L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html> 1777L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html>
1218 1778
1219=back 1779=back
1220 1780
1229 1789
1230Waits for the event to complete. 1790Waits for the event to complete.
1231 1791
1232L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html> 1792L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html>
1233 1793
1794=item $ev->cb ($exec_callback_type, $callback->($event, $event_command_exec_status))
1795
1796Adds a callback to the callback stack for the given event type. There is
1797no way to remove a callback again.
1798
1799L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetEventCallback.html>
1800
1234=item $packed_value = $ev->info ($name) 1801=item $packed_value = $ev->info ($name)
1235 1802
1236See C<< $platform->info >> for details. 1803See C<< $platform->info >> for details.
1237 1804
1238L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetEventInfo.html> 1805L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetEventInfo.html>
1298 1865
1299=over 4 1866=over 4
1300 1867
1301=item $ev->set_status ($execution_status) 1868=item $ev->set_status ($execution_status)
1302 1869
1870Sets the execution status of the user event. Can only be called once,
1871either with OpenCL::COMPLETE or a negative number as status.
1872
1303L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetUserEventStatus.html> 1873L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetUserEventStatus.html>
1304 1874
1305=back 1875=back
1306 1876
1307=cut 1877=cut
1308
1309package OpenCL;
1310
1311use common::sense;
1312
1313BEGIN {
1314 our $VERSION = '0.92';
1315
1316 require XSLoader;
1317 XSLoader::load (__PACKAGE__, $VERSION);
1318
1319 @OpenCL::Buffer::ISA =
1320 @OpenCL::Image::ISA = OpenCL::Memory::;
1321
1322 @OpenCL::BufferObj::ISA = OpenCL::Buffer::;
1323
1324 @OpenCL::Image2D::ISA =
1325 @OpenCL::Image3D::ISA = OpenCL::Image::;
1326
1327 @OpenCL::UserEvent::ISA = OpenCL::Event::;
1328}
1329 1878
13301; 18791;
1331 1880
1332=head1 AUTHOR 1881=head1 AUTHOR
1333 1882

Diff Legend

Removed lines
+ Added lines
< Changed lines
> Changed lines