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

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.24 by root, Sun Nov 20 22:31:48 2011 UTC vs.
Revision 1.70 by root, Thu May 3 23:32:47 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
105 for my $platform (OpenCL::platforms) { 108 for my $platform (OpenCL::platforms) {
106 printf "platform: %s\n", $platform->name; 109 printf "platform: %s\n", $platform->name;
107 printf "extensions: %s\n", $platform->extensions; 110 printf "extensions: %s\n", $platform->extensions;
108 for my $device ($platform->devices) { 111 for my $device ($platform->devices) {
109 printf "+ device: %s\n", $device->name; 112 printf "+ device: %s\n", $device->name;
110 my $ctx = $device->context; 113 my $ctx = $platform->context (undef, [$device]);
111 # do stuff 114 # do stuff
112 } 115 }
113 } 116 }
114 117
115=head2 Get a useful context and a command queue. 118=head2 Get a useful context and a command queue.
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.
152 155
153 my $src = ' 156 my $src = '
154 __kernel void 157 kernel void
155 squareit (__global float *input, __global float *output) 158 squareit (global float *input, global float *output)
156 { 159 {
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 - you can download it from
218L<http://cvs.schmorp.de/OpenCL/examples/juliaflight>.
219
220 use OpenGL ":all";
221 use OpenCL;
222
223 my $S = $ARGV[0] || 256; # window/texture size, smaller is faster
224
225 # open a window and create a gl texture
226 OpenGL::glpOpenWindow width => $S, height => $S;
227 my $texid = glGenTextures_p 1;
228 glBindTexture GL_TEXTURE_2D, $texid;
229 glTexImage2D_c GL_TEXTURE_2D, 0, GL_RGBA8, $S, $S, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0;
230
231 # find and use the first opencl device that let's us get a shared opengl context
232 my $platform;
233 my $dev;
234 my $ctx;
235
236 for (OpenCL::platforms) {
237 $platform = $_;
238 for ($platform->devices) {
239 $dev = $_;
240 $ctx = $platform->context ([OpenCL::GLX_DISPLAY_KHR, undef, OpenCL::GL_CONTEXT_KHR, undef], [$dev])
241 and last;
242 }
243 }
244
245 $ctx
246 or die "cannot find suitable OpenCL device\n";
247
248 my $queue = $ctx->queue ($dev);
249
250 # now attach an opencl image2d object to the opengl texture
251 my $tex = $ctx->gl_texture2d (OpenCL::MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, $texid);
252
253 # now the boring opencl code
254 my $src = <<EOF;
255 kernel void
256 juliatunnel (write_only image2d_t img, float time)
257 {
258 int2 xy = (int2)(get_global_id (0), get_global_id (1));
259 float2 p = convert_float2 (xy) / $S.f * 2.f - 1.f;
260
261 float2 m = (float2)(1.f, p.y) / fabs (p.x); // tunnel
262 m.x = fabs (fmod (m.x + time * 0.05f, 4.f) - 2.f);
263
264 float2 z = m;
265 float2 c = (float2)(sin (time * 0.01133f), cos (time * 0.02521f));
266
267 for (int i = 0; i < 25 && dot (z, z) < 4.f; ++i) // standard julia
268 z = (float2)(z.x * z.x - z.y * z.y, 2.f * z.x * z.y) + c;
269
270 float3 colour = (float3)(z.x, z.y, atan2 (z.y, z.x));
271 write_imagef (img, xy, (float4)(colour * p.x * p.x, 1.));
272 }
273 EOF
274
275 my $prog = $ctx->build_program ($src);
276 my $kernel = $prog->kernel ("juliatunnel");
277
278 # program compiled, kernel ready, now draw and loop
279
280 for (my $time; ; ++$time) {
281 # acquire objects from opengl
282 $queue->acquire_gl_objects ([$tex]);
283
284 # configure and run our kernel
285 $kernel->setf ("mf", $tex, $time*2); # mf = memory object, float
286 $queue->nd_range_kernel ($kernel, undef, [$S, $S], undef);
287
288 # release objects to opengl again
289 $queue->release_gl_objects ([$tex]);
290
291 # wait
292 $queue->finish;
293
294 # now draw the texture, the defaults should be all right
295 glTexParameterf GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST;
296
297 glEnable GL_TEXTURE_2D;
298 glBegin GL_QUADS;
299 glTexCoord2f 0, 1; glVertex3i -1, -1, -1;
300 glTexCoord2f 0, 0; glVertex3i 1, -1, -1;
301 glTexCoord2f 1, 0; glVertex3i 1, 1, -1;
302 glTexCoord2f 1, 1; glVertex3i -1, 1, -1;
303 glEnd;
304
305 glXSwapBuffers;
306
307 select undef, undef, undef, 1/60;
308 }
309
310=head2 How to modify the previous example to not rely on GL sharing.
311
312For those poor souls with only a sucky CPU OpenCL implementation, you
313currently have to read the image into some perl scalar, and then modify a
314texture or use glDrawPixels or so).
315
316First, when you don't need gl sharing, you can create the context much simpler:
317
318 $ctx = $platform->context (undef, [$dev])
319
320To use a texture, you would modify the above example by creating an
321OpenCL::Image manually instead of deriving it from a texture:
322
323 my $tex = $ctx->image2d (OpenCL::MEM_WRITE_ONLY, OpenCL::RGBA, OpenCL::UNORM_INT8, $S, $S);
324
325And in the darw loop, intead of acquire_gl_objects/release_gl_objects, you
326would read the image2d after the kernel has written it:
327
328 $queue->read_image ($tex, 0, 0, 0, 0, $S, $S, 1, 0, 0, my $data);
329
330And then you would upload the pixel data to the texture (or use glDrawPixels):
331
332 glTexSubImage2D_s GL_TEXTURE_2D, 0, 0, 0, $S, $S, GL_RGBA, GL_UNSIGNED_BYTE, $data;
333
334The fully modified example can be found at
335L<http://cvs.schmorp.de/OpenCL/examples/juliaflight-nosharing>.
215 336
216=head1 DOCUMENTATION 337=head1 DOCUMENTATION
217 338
218=head2 BASIC CONVENTIONS 339=head2 BASIC CONVENTIONS
219 340
241=item * Structures are often specified by flattening out their components 362=item * Structures are often specified by flattening out their components
242as with short vectors, and returned as arrayrefs. 363as with short vectors, and returned as arrayrefs.
243 364
244=item * When enqueuing commands, the wait list is specified by adding 365=item * When enqueuing commands, the wait list is specified by adding
245extra arguments to the function - anywhere a C<$wait_events...> argument 366extra arguments to the function - anywhere a C<$wait_events...> argument
246is documented this can be any number of event objects. 367is documented this can be any number of event objects. As an extsnion
368implemented by this module, C<undef> values will be ignored in the event
369list.
247 370
248=item * When enqueuing commands, if the enqueue method is called in void 371=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 372context, no event is created. In all other contexts an event is returned
250by the method. 373by the method.
251 374
271 ulong IV - Q 394 ulong IV - Q
272 float NV float f 395 float NV float f
273 half IV ushort S 396 half IV ushort S
274 double NV double d 397 double NV double d
275 398
399=head2 GLX SUPPORT
400
401Due to the sad state that OpenGL support is in in Perl (mostly the OpenGL
402module, which has little to no documentation and has little to no support
403for glX), this module, as a special extension, treats context creation
404properties C<OpenCL::GLX_DISPLAY_KHR> and C<OpenCL::GL_CONTEXT_KHR>
405specially: If either or both of these are C<undef>, then the OpenCL
406module tries to dynamically resolve C<glXGetCurrentDisplay> and
407C<glXGetCurrentContext>, call these functions and use their return values
408instead.
409
410For this to work, the OpenGL library must be loaded, a GLX context must
411have been created and be made current, and C<dlsym> must be available and
412capable of finding the function via C<RTLD_DEFAULT>.
413
414=head2 EVENT SYSTEM
415
416OpenCL can generate a number of (potentially) asynchronous events, for
417example, after compiling a program, to signal a context-related error or,
418perhaps most important, to signal completion of queued jobs (by setting
419callbacks on OpenCL::Event objects).
420
421To facilitate this, this module maintains an event queue - each
422time an asynchronous event happens, it is queued, and perl will be
423interrupted. This is implemented via the L<Async::Interrupt> module. In
424addition, this module has L<AnyEvent> support, so it can seamlessly
425integrate itself into many event loops.
426
427Since this module is a bit hard to understand, here are some case examples:
428
429=head3 Don't use callbacks.
430
431When your program never uses any callbacks, then there will never be any
432notifications you need to take care of, and therefore no need to worry
433about all this.
434
435You can achieve a great deal by explicitly waiting for events, or using
436barriers and flush calls. In many programs, there is no need at all to
437tinker with asynchronous events.
438
439=head3 Use AnyEvent
440
441This module automatically registers a watcher that invokes all outstanding
442event callbacks when AnyEvent is initialised (and block asynchronous
443interruptions). Using this mode of operations is the safest and most
444recommended one.
445
446To use this, simply use AnyEvent and this module normally, make sure you
447have an event loop running:
448
449 use Gtk2 -init;
450 use AnyEvent;
451
452 # initialise AnyEvent, by creating a watcher, or:
453 AnyEvent::detect;
454
455 my $e = $queue->marker;
456 $e->cb (sub {
457 warn "opencl is finished\n";
458 })
459
460 main Gtk2;
461
462Note that this module will not initialise AnyEvent for you. Before
463AnyEvent is initialised, the module will asynchronously interrupt perl
464instead. To avoid any surprises, it's best to explicitly initialise
465AnyEvent.
466
467You can temporarily enable asynchronous interruptions (see next paragraph)
468by calling C<$OpenCL::INTERRUPT->unblock> and disable them again by
469calling C<$OpenCL::INTERRUPT->block>.
470
471=head3 Let yourself be interrupted at any time
472
473This mode is the default unless AnyEvent is loaded and initialised. In
474this mode, OpenCL asynchronously interrupts a running perl program. The
475emphasis is on both I<asynchronously> and I<running> here.
476
477Asynchronously means that perl might execute your callbacks at any
478time. For example, in the following code (I<THAT YOU SHOULD NOT COPY>),
479the C<until> loop following the marker call will be interrupted by the
480callback:
481
482 my $e = $queue->marker;
483 my $flag;
484 $e->cb (sub { $flag = 1 });
485 1 until $flag;
486 # $flag is now 1
487
488The reason why you shouldn't blindly copy the above code is that
489busy waiting is a really really bad thing, and really really bad for
490performance.
491
492While at first this asynchronous business might look exciting, it can be
493really hard, because you need to be prepared for the callback code to be
494executed at any time, which limits the amount of things the callback code
495can do safely.
496
497This can be mitigated somewhat by using C<<
498$OpenCL::INTERRUPT->scope_block >> (see the L<Async::Interrupt>
499documentation for details).
500
501The other problem is that your program must be actively I<running> to be
502interrupted. When you calculate stuff, your program is running. When you
503hang in some C functions or other block execution (by calling C<sleep>,
504C<select>, running an event loop and so on), your program is waiting, not
505running.
506
507One way around that would be to attach a read watcher to your event loop,
508listening for events on C<< $OpenCL::INTERRUPT->pipe_fileno >>, using a
509dummy callback (C<sub { }>) to temporarily execute some perl code.
510
511That is then awfully close to using the built-in AnyEvent support above,
512though, so consider that one instead.
513
514=head3 Be creative
515
516OpenCL exports the L<Async::Interrupt> object it uses in the global
517variable C<$OpenCL::INTERRUPT>. You can configure it in any way you like.
518
519So if you want to feel like a real pro, err, wait, if you feel no risk
520menas no fun, you can experiment by implementing your own mode of
521operations.
522
523=cut
524
525package OpenCL;
526
527use common::sense;
528use Carp ();
529use Async::Interrupt ();
530
531our $POLL_FUNC; # set by XS
532
533BEGIN {
534 our $VERSION = '0.99';
535
536 require XSLoader;
537 XSLoader::load (__PACKAGE__, $VERSION);
538
539 @OpenCL::Platform::ISA =
540 @OpenCL::Device::ISA =
541 @OpenCL::Context::ISA =
542 @OpenCL::Queue::ISA =
543 @OpenCL::Memory::ISA =
544 @OpenCL::Sampler::ISA =
545 @OpenCL::Program::ISA =
546 @OpenCL::Kernel::ISA =
547 @OpenCL::Event::ISA = OpenCL::Object::;
548
549 @OpenCL::Buffer::ISA =
550 @OpenCL::Image::ISA = OpenCL::Memory::;
551
552 @OpenCL::BufferObj::ISA = OpenCL::Buffer::;
553
554 @OpenCL::Image2D::ISA =
555 @OpenCL::Image3D::ISA =
556 @OpenCL::Image2DArray::ISA =
557 @OpenCL::Image1D::ISA =
558 @OpenCL::Image1DArray::ISA =
559 @OpenCL::Image1DBuffer::ISA = OpenCL::Image::;
560
561 @OpenCL::UserEvent::ISA = OpenCL::Event::;
562
563 @OpenCL::MappedBuffer::ISA =
564 @OpenCL::MappedImage::ISA = OpenCL::Mapped::;
565}
566
276=head2 THE OpenCL PACKAGE 567=head2 THE OpenCL PACKAGE
277 568
278=over 4 569=over 4
279 570
280=item $int = OpenCL::errno 571=item $int = OpenCL::errno
281 572
282The last error returned by a function - it's only valid after an error occured 573The last error returned by a function - it's only valid after an error occured
283and before calling another OpenCL function. 574and before calling another OpenCL function.
284 575
285=item $str = OpenCL::err2str $errval 576=item $str = OpenCL::err2str [$errval]
286 577
287Comverts an error value into a human readable string. 578Converts an error value into a human readable string. IF no error value is
579given, then the last error will be used (as returned by OpenCL::errno).
288 580
289=item $str = OpenCL::enum2str $enum 581=item $str = OpenCL::enum2str $enum
290 582
291Converts most enum values (inof parameter names, image format constants, 583Converts most enum values (of parameter names, image format constants,
292object types, addressing and filter modes, command types etc.) into a 584object types, addressing and filter modes, command types etc.) into a
293human readbale string. When confronted with some random integer it can be 585human readable string. When confronted with some random integer it can be
294very helpful to pass it through this function to maybe get some readable 586very helpful to pass it through this function to maybe get some readable
295string out of it. 587string out of it.
296 588
297=item @platforms = OpenCL::platforms 589=item @platforms = OpenCL::platforms
298 590
299Returns all available OpenCL::Platform objects. 591Returns all available OpenCL::Platform objects.
300 592
301L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformIDs.html> 593L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformIDs.html>
302 594
303=item $ctx = OpenCL::context_from_type $properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $notify = undef 595=item $ctx = OpenCL::context_from_type $properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $callback->($err, $pvt) = $print_stderr
304 596
305Tries to create a context from a default device and platform - never worked for me. 597Tries to create a context from a default device and platform type - never worked for me.
306 598
307L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html> 599L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html>
308 600
601=item $ctx = OpenCL::context $properties, \@devices, $callback->($err, $pvt) = $print_stderr)
602
603Create a new OpenCL::Context object using the given device object(s). This
604function isn't implemented yet, use C<< $platform->context >> instead.
605
606L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContext.html>
607
309=item OpenCL::wait_for_events $wait_events... 608=item OpenCL::wait_for_events $wait_events...
310 609
311Waits for all events to complete. 610Waits for all events to complete.
312 611
313L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html> 612L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html>
314 613
614=item OpenCL::poll
615
616Checks if there are any outstanding events (see L<EVENT SYSTEM>) and
617invokes their callbacks.
618
619=item $OpenCL::INTERRUPT
620
621The L<Async::Interrupt> object used to signal asynchronous events (see
622L<EVENT SYSTEM>).
623
624=cut
625
626our $INTERRUPT = new Async::Interrupt c_cb => [$POLL_FUNC, 0];
627
628&_eq_initialise ($INTERRUPT->signal_func);
629
630=item $OpenCL::WATCHER
631
632The L<AnyEvent> watcher object used to watch for asynchronous events (see
633L<EVENT SYSTEM>). This variable is C<undef> until L<AnyEvent> has been
634loaded I<and> initialised (e.g. by calling C<AnyEvent::detect>).
635
636=cut
637
638our $WATCHER;
639
640sub _init_anyevent {
641 $INTERRUPT->block;
642 $WATCHER = AE::io ($INTERRUPT->pipe_fileno, 0, sub { $INTERRUPT->handle });
643}
644
645if (defined $AnyEvent::MODEL) {
646 _init_anyevent;
647} else {
648 push @AnyEvent::post_detect, \&_init_anyevent;
649}
650
315=back 651=back
316 652
653=head2 THE OpenCL::Object CLASS
654
655This is the base class for all objects in the OpenCL module. The only
656method it implements is the C<id> method, which is only useful if you want
657to interface to OpenCL on the C level.
658
659=over 4
660
661=item $iv = $obj->id
662
663OpenCL objects are represented by pointers or integers on the C level. If
664you want to interface to an OpenCL object directly on the C level, then
665you need this value, which is returned by this method. You should use an
666C<IV> type in your code and cast that to the correct type.
667
668=cut
669
670sub OpenCL::Object::id {
671 ref $_[0] eq "SCALAR"
672 ? ${ $_[0] }
673 : $_[0][0]
674}
675
676=back
677
317=head2 THE OpenCL::Platform CLASS 678=head2 THE OpenCL::Platform CLASS
318 679
319=over 4 680=over 4
320 681
321=item @devices = $platform->devices ($type = OpenCL::DEVICE_TYPE_ALL) 682=item @devices = $platform->devices ($type = OpenCL::DEVICE_TYPE_ALL)
322 683
323Returns a list of matching OpenCL::Device objects. 684Returns a list of matching OpenCL::Device objects.
324 685
325=item $ctx = $platform->context_from_type ($properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $notify = undef) 686=item $ctx = $platform->context_from_type ($properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $callback->($err, $pvt) = $print_stderr)
326 687
327Tries to create a context. Never worked for me, and you need devices explicitly anyway. 688Tries to create a context. Never worked for me, and you need devices explicitly anyway.
328 689
329L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html> 690L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html>
330 691
331=item $ctx = $device->context ($properties = undef, @$devices, $notify = undef) 692=item $ctx = $platform->context ($properties, \@devices, $callback->($err, $pvt) = $print_stderr)
332 693
333Create a new OpenCL::Context object using the given device object(s)- a 694Create a new OpenCL::Context object using the given device object(s)- a
334CL_CONTEXT_PLATFORM property is supplied automatically. 695CL_CONTEXT_PLATFORM property is supplied automatically.
335 696
336L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContext.html> 697L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContext.html>
344It's best to avoid this method and use one of the following convenience 705It's best to avoid this method and use one of the following convenience
345wrappers. 706wrappers.
346 707
347L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformInfo.html> 708L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformInfo.html>
348 709
710=item $platform->unload_compiler
711
712Attempts to unload the compiler for this platform, for endless
713profit. Does nothing on OpenCL 1.1.
714
715L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clUnloadPlatformCompiler.html>
716
349=for gengetinfo begin platform 717=for gengetinfo begin platform
350 718
351=item $string = $platform->profile 719=item $string = $platform->profile
352 720
353Calls C<clGetPlatformInfo> with C<CL_PLATFORM_PROFILE> and returns the result. 721Calls C<clGetPlatformInfo> with C<CL_PLATFORM_PROFILE> and returns the result.
638 1006
639=item @device_partition_property_exts = $device->affinity_domains_ext 1007=item @device_partition_property_exts = $device->affinity_domains_ext
640 1008
641Calls C<clGetDeviceInfo> with C<CL_DEVICE_AFFINITY_DOMAINS_EXT> and returns the result. 1009Calls C<clGetDeviceInfo> with C<CL_DEVICE_AFFINITY_DOMAINS_EXT> and returns the result.
642 1010
643=item $uint = $device->reference_count_ext 1011=item $uint = $device->reference_count_ext
644 1012
645Calls C<clGetDeviceInfo> with C<CL_DEVICE_REFERENCE_COUNT_EXT > and returns the result. 1013Calls C<clGetDeviceInfo> with C<CL_DEVICE_REFERENCE_COUNT_EXT> and returns the result.
646 1014
647=item @device_partition_property_exts = $device->partition_style_ext 1015=item @device_partition_property_exts = $device->partition_style_ext
648 1016
649Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARTITION_STYLE_EXT> and returns the result. 1017Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARTITION_STYLE_EXT> and returns the result.
650 1018
654 1022
655=head2 THE OpenCL::Context CLASS 1023=head2 THE OpenCL::Context CLASS
656 1024
657=over 4 1025=over 4
658 1026
1027=item $prog = $ctx->build_program ($program, $options = "")
1028
1029This convenience function tries to build the program on all devices in
1030the context. If the build fails, then the function will C<croak> with the
1031build log. Otherwise ti returns the program object.
1032
1033The C<$program> can either be a C<OpenCL::Program> object or a string
1034containing the program. In the latter case, a program objetc will be
1035created automatically.
1036
1037=cut
1038
1039sub OpenCL::Context::build_program {
1040 my ($self, $prog, $options) = @_;
1041
1042 $prog = $self->program_with_source ($prog)
1043 unless ref $prog;
1044
1045 eval { $prog->build (undef, $options); 1 }
1046 or errno == BUILD_PROGRAM_FAILURE
1047 or errno == INVALID_BINARY # workaround nvidia bug
1048 or Carp::croak "OpenCL::Context->build_program: " . err2str;
1049
1050 # we check status for all devices
1051 for my $dev ($self->devices) {
1052 $prog->build_status ($dev) == BUILD_SUCCESS
1053 or Carp::croak "Building OpenCL program for device '" . $dev->name . "' failed:\n"
1054 . $prog->build_log ($dev);
1055 }
1056
1057 $prog
1058}
1059
659=item $queue = $ctx->queue ($device, $properties) 1060=item $queue = $ctx->queue ($device, $properties)
660 1061
661Create a new OpenCL::Queue object from the context and the given device. 1062Create a new OpenCL::Queue object from the context and the given device.
662 1063
663L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateCommandQueue.html> 1064L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateCommandQueue.html>
664 1065
1066Example: create an out-of-order queue.
1067
1068 $queue = $ctx->queue ($device, OpenCL::QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE);
1069
665=item $ev = $ctx->user_event 1070=item $ev = $ctx->user_event
666 1071
667Creates a new OpenCL::UserEvent object. 1072Creates a new OpenCL::UserEvent object.
668 1073
669L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateUserEvent.html> 1074L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateUserEvent.html>
670 1075
671=item $buf = $ctx->buffer ($flags, $len) 1076=item $buf = $ctx->buffer ($flags, $len)
672 1077
673Creates a new OpenCL::Buffer object with the given flags and octet-size. 1078Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object with the
1079given flags and octet-size.
674 1080
675L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateBuffer.html> 1081L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateBuffer.html>
676 1082
677=item $buf = $ctx->buffer_sv ($flags, $data) 1083=item $buf = $ctx->buffer_sv ($flags, $data)
678 1084
679Creates a new OpenCL::Buffer object and initialise it with the given data values. 1085Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object and
1086initialise it with the given data values.
1087
1088=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)
1089
1090Creates a new OpenCL::Image object and optionally initialises it with
1091the given data values.
1092
1093L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateImage.html>
680 1094
681=item $img = $ctx->image2d ($flags, $channel_order, $channel_type, $width, $height, $row_pitch = 0, $data = undef) 1095=item $img = $ctx->image2d ($flags, $channel_order, $channel_type, $width, $height, $row_pitch = 0, $data = undef)
682 1096
683Creates a new OpenCL::Image2D object and optionally initialises it with the given data values. 1097Creates a new OpenCL::Image2D object and optionally initialises it with
1098the given data values.
684 1099
685L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateImage2D.html> 1100L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateImage2D.html>
686 1101
687=item $img = $ctx->image3d ($flags, $channel_order, $channel_type, $width, $height, $depth, $row_pitch = 0, $slice_pitch = 0, $data = undef) 1102=item $img = $ctx->image3d ($flags, $channel_order, $channel_type, $width, $height, $depth, $row_pitch = 0, $slice_pitch = 0, $data = undef)
688 1103
689Creates a new OpenCL::Image3D object and optionally initialises it with the given data values. 1104Creates a new OpenCL::Image3D object and optionally initialises it with
1105the given data values.
690 1106
691L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateImage3D.html> 1107L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateImage3D.html>
1108
1109=item $buffer = $ctx->gl_buffer ($flags, $bufobj)
1110
1111Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object that refers to the given
1112OpenGL buffer object.
1113
1114http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLBuffer.html
1115
1116=item $img = $ctx->gl_texture ($flags, $target, $miplevel, $texture)
1117
1118Creates a new OpenCL::Image object that refers to the given OpenGL
1119texture object or buffer.
1120
1121http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateFromGLTexture.html
1122
1123=item $img = $ctx->gl_texture2d ($flags, $target, $miplevel, $texture)
1124
1125Creates a new OpenCL::Image2D object that refers to the given OpenGL
11262D texture object.
1127
1128http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture2D.html
1129
1130=item $img = $ctx->gl_texture3d ($flags, $target, $miplevel, $texture)
1131
1132Creates a new OpenCL::Image3D object that refers to the given OpenGL
11333D texture object.
1134
1135http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture3D.html
1136
1137=item $ctx->gl_renderbuffer ($flags, $renderbuffer)
1138
1139Creates a new OpenCL::Image2D object that refers to the given OpenGL
1140render buffer.
1141
1142http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLRenderbuffer.html
692 1143
693=item @formats = $ctx->supported_image_formats ($flags, $image_type) 1144=item @formats = $ctx->supported_image_formats ($flags, $image_type)
694 1145
695Returns a list of matching image formats - each format is an arrayref with 1146Returns a list of matching image formats - each format is an arrayref with
696two values, $channel_order and $channel_type, in it. 1147two values, $channel_order and $channel_type, in it.
707 1158
708Creates a new OpenCL::Program object from the given source code. 1159Creates a new OpenCL::Program object from the given source code.
709 1160
710L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateProgramWithSource.html> 1161L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateProgramWithSource.html>
711 1162
1163=item ($program, \@status) = $ctx->program_with_binary (\@devices, \@binaries)
1164
1165Creates a new OpenCL::Program object from the given binaries.
1166
1167L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateProgramWithBinary.html>
1168
1169Example: clone an existing program object that contains a successfully
1170compiled program, no matter how useless this is.
1171
1172 my $clone = $ctx->program_with_binary ([$prog->devices], [$prog->binaries]);
1173
712=item $packed_value = $ctx->info ($name) 1174=item $packed_value = $ctx->info ($name)
713 1175
714See C<< $platform->info >> for details. 1176See C<< $platform->info >> for details.
715 1177
716L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetContextInfo.html> 1178L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetContextInfo.html>
738=back 1200=back
739 1201
740=head2 THE OpenCL::Queue CLASS 1202=head2 THE OpenCL::Queue CLASS
741 1203
742An OpenCL::Queue represents an execution queue for OpenCL. You execute 1204An OpenCL::Queue represents an execution queue for OpenCL. You execute
743requests by calling their respective C<enqueue_xxx> method and waitinf for 1205requests by calling their respective method and waiting for it to complete
744it to complete in some way. 1206in some way.
745 1207
746All the enqueue methods return an event object that can be used to wait 1208Most methods that enqueue some request return an event object that can
747for completion, unless the method is called in void context, in which case 1209be used to wait for completion (optionally using a callback), unless
748no event object is created. 1210the method is called in void context, in which case no event object is
1211created.
749 1212
750They also allow you to specify any number of other event objects that this 1213They also allow you to specify any number of other event objects that this
751request has to wait for before it starts executing, by simply passing the 1214request has to wait for before it starts executing, by simply passing the
752event objects as extra parameters to the enqueue methods. 1215event objects as extra parameters to the enqueue methods. To simplify
1216program design, this module ignores any C<undef> values in the list of
1217events. This makes it possible to code operations such as this, without
1218having to put a valid event object into C<$event> first:
1219
1220 $event = $queue->xxx (..., $event);
753 1221
754Queues execute in-order by default, without any parallelism, so in most 1222Queues execute in-order by default, without any parallelism, so in most
755cases (i.e. you use only one queue) it's not necessary to wait for or 1223cases (i.e. you use only one queue) it's not necessary to wait for or
756create event objects. 1224create event objects, althoguh an our of order queue is often a bit
1225faster.
757 1226
758=over 4 1227=over 4
759 1228
760=item $ev = $queue->enqueue_read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...) 1229=item $ev = $queue->read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...)
761 1230
762Reads data from buffer into the given string. 1231Reads data from buffer into the given string.
763 1232
764L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBuffer.html> 1233L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBuffer.html>
765 1234
766=item $ev = $queue->enqueue_write_buffer ($buffer, $blocking, $offset, $data, $wait_events...) 1235=item $ev = $queue->write_buffer ($buffer, $blocking, $offset, $data, $wait_events...)
767 1236
768Writes data to buffer from the given string. 1237Writes data to buffer from the given string.
769 1238
770L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBuffer.html> 1239L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBuffer.html>
771 1240
772=item $ev = $queue->enqueue_copy_buffer ($src, $dst, $src_offset, $dst_offset, $len, $wait_events...) 1241=item $ev = $queue->copy_buffer ($src, $dst, $src_offset, $dst_offset, $len, $wait_events...)
773 1242
774L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBuffer.html> 1243L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBuffer.html>
775 1244
1245=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...)
1246
1247http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBufferRect.html
1248
1249=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...)
1250
1251http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBufferRect.html
1252
1253=item $ev = $queue->copy_buffer_to_image ($src_buffer, $dst_image, $src_offset, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $wait_events...)
1254
1255L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>
1256
776=item $ev = $queue->enqueue_read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...) 1257=item $ev = $queue->read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
1258
1259C<$row_pitch> (and C<$slice_pitch>) can be C<0>, in which case the OpenCL
1260module uses the image width (and height) to supply default values.
777 1261
778L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadImage.html> 1262L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadImage.html>
779 1263
780=item $ev = $queue->enqueue_write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...) 1264=item $ev = $queue->write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
781 1265
1266C<$row_pitch> (and C<$slice_pitch>) can be C<0>, in which case the OpenCL
1267module uses the image width (and height) to supply default values.
782L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteImage.html> 1268L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteImage.html>
783 1269
1270=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...)
1271
1272L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html>
1273
1274=item $ev = $queue->copy_image_to_buffer ($src_image, $dst_image, $src_x, $src_y, $src_z, $width, $height, $depth, $dst_offset, $wait_events...)
1275
1276L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html>
1277
784=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...) 1278=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...)
785 1279
786Yeah. 1280Yeah.
787 1281
788L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferRect.html>
789
790=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...)
791
792L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>. 1282L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>.
793 1283
794=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...) 1284=item $ev = $queue->fill_buffer ($mem, $pattern, $offset, $size, ...)
795 1285
1286Fills the given buffer object with repeated applications of C<$pattern>,
1287starting at C<$offset> for C<$size> octets.
1288
1289L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillBuffer.html>
1290
1291=item $ev = $queue->fill_image ($img, $r, $g, $b, $a, $x, $y, $z, $width, $height, $depth, ...)
1292
1293Fills the given image area with the given rgba colour components. The
1294components are normally floating point values between C<0> and C<1>,
1295except when the image channel data type is a signe dor unsigned
1296unnormalised format, in which case the range is determined by the format.
1297
796L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html> 1298L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillImage.html>
797 1299
798=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...)
799
800L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html>
801
802=item $ev = $queue->enqueue_task ($kernel, $wait_events...) 1300=item $ev = $queue->task ($kernel, $wait_events...)
803 1301
804L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html> 1302L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html>
805 1303
806=item $ev = $queue->enqueue_nd_range_kernel ($kernel, @$global_work_offset, @$global_work_size, @$local_work_size, $wait_events...) 1304=item $ev = $queue->nd_range_kernel ($kernel, \@global_work_offset, \@global_work_size, \@local_work_size, $wait_events...)
807 1305
808Enqueues a kernel execution. 1306Enqueues a kernel execution.
809 1307
810@$global_work_size must be specified as a reference to an array of 1308\@global_work_size must be specified as a reference to an array of
811integers specifying the work sizes (element counts). 1309integers specifying the work sizes (element counts).
812 1310
813@$global_work_offset must be either C<undef> (in which case all offsets 1311\@global_work_offset must be either C<undef> (in which case all offsets
814are C<0>), or a reference to an array of work offsets, with the same number 1312are C<0>), or a reference to an array of work offsets, with the same number
815of elements as @$global_work_size. 1313of elements as \@global_work_size.
816 1314
817@$local_work_size must be either C<undef> (in which case the 1315\@local_work_size must be either C<undef> (in which case the
818implementation is supposed to choose good local work sizes), or a 1316implementation is supposed to choose good local work sizes), or a
819reference to an array of local work sizes, with the same number of 1317reference to an array of local work sizes, with the same number of
820elements as @$global_work_size. 1318elements as \@global_work_size.
821 1319
822L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html> 1320L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html>
823 1321
824=item $ev = $queue->enqueue_marker 1322=item $ev = $queue->acquire_gl_objects ([object, ...], $wait_events...)
825 1323
1324Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired
1325for subsequent OpenCL usage.
1326
826L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMarker.html> 1327L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueAcquireGLObjects.html>
827 1328
1329=item $ev = $queue->release_gl_objects ([object, ...], $wait_events...)
1330
1331Enqueues a list (an array-ref of OpenCL::Memory objects) to be released
1332for subsequent OpenGL usage.
1333
1334L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReleaseGLObjects.html>
1335
828=item $ev = $queue->enqueue_wait_for_events ($wait_events...) 1336=item $ev = $queue->wait_for_events ($wait_events...)
829 1337
830L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html> 1338L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html>
831 1339
832=item $queue->enqueue_barrier 1340=item $ev = $queue->marker ($wait_events...)
833 1341
1342L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMarkerWithWaitList.html>
1343
1344=item $ev = $queue->barrier ($wait_events...)
1345
834L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueBarrier.html> 1346L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueBarrierWithWaitList.html>
835 1347
836=item $queue->flush 1348=item $queue->flush
837 1349
838L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clFlush.html> 1350L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clFlush.html>
839 1351
864=item $command_queue_properties = $command_queue->properties 1376=item $command_queue_properties = $command_queue->properties
865 1377
866Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_PROPERTIES> and returns the result. 1378Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_PROPERTIES> and returns the result.
867 1379
868=for gengetinfo end command_queue 1380=for gengetinfo end command_queue
1381
1382=back
1383
1384=head3 MEMORY MAPPED BUFFERS
1385
1386OpenCL allows you to map buffers and images to host memory (read: perl
1387scalars). This is done much like reading or copying a buffer, by enqueuing
1388a map or unmap operation on the command queue.
1389
1390The map operations return an C<OpenCL::Mapped> object - see L<THE
1391OpenCL::Mapped CLASS> section for details on what to do with these
1392objects.
1393
1394The object will be unmapped automatically when the mapped object is
1395destroyed (you can use a barrier to make sure the unmap has finished,
1396before using the buffer in a kernel), but you can also enqueue an unmap
1397operation manually.
1398
1399=over 4
1400
1401=item $mapped_buffer = $queue->map_buffer ($buf, $blocking=1, $map_flags=OpenCL::MAP_READ|OpenCL::MAP_WRITE, $offset=0, $size=undef, $wait_events...)
1402
1403Maps the given buffer into host memory and returns an
1404C<OpenCL::MappedBuffer> object. If C<$size> is specified as undef, then
1405the map will extend to the end of the buffer.
1406
1407L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMapBuffer.html>
1408
1409Example: map the buffer $buf fully and replace the first 4 bytes by "abcd", then unmap.
1410
1411 {
1412 my $mapped = $queue->map_buffer ($buf, 1, OpenCL::MAP_WRITE);
1413 substr $$mapped, 0, 4, "abcd";
1414 } # asynchronously unmap because $mapped is destroyed
1415
1416=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...)
1417
1418Maps the given image area into host memory and return an
1419C<OpenCL::MappedImage> object.
1420
1421If any of C<$width>, C<$height> and/or C<$depth> are C<undef> then they
1422will be replaced by the maximum possible value.
1423
1424L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMapImage.html>
1425
1426Example: map an image (with OpenCL::UNSIGNED_INT8 channel type) and set
1427the first channel of the leftmost column to 5, then explicitly unmap
1428it. You are not necessarily meant to do it this way, this example just
1429shows you the accessors to use :)
1430
1431 my $mapped = $queue->map_image ($image, 1, OpenCL::MAP_WRITE);
1432
1433 $mapped->set ($_ * $mapped->row_pitch, pack "C", 5)
1434 for 0..$image->height;
1435
1436 $mapped->unmap;.
1437 $mapped->wait; # only needed for out of order queues normally
1438
1439=item $ev = $queue->unmap ($mapped, $wait_events...)
1440
1441Unmaps the data from host memory. You must not call any methods that
1442modify the data, or modify the data scalar directly, after calling this
1443method.
1444
1445The mapped event object will always be passed as part of the
1446$wait_events. The mapped event object will be replaced by the new event
1447object that this request creates.
869 1448
870=back 1449=back
871 1450
872=head2 THE OpenCL::Memory CLASS 1451=head2 THE OpenCL::Memory CLASS
873 1452
920 1499
921Calls C<clGetMemObjectInfo> with C<CL_MEM_OFFSET> and returns the result. 1500Calls C<clGetMemObjectInfo> with C<CL_MEM_OFFSET> and returns the result.
922 1501
923=for gengetinfo end mem 1502=for gengetinfo end mem
924 1503
1504=item ($type, $name) = $mem->gl_object_info
1505
1506Returns the OpenGL object type (e.g. OpenCL::GL_OBJECT_TEXTURE2D) and the
1507object "name" (e.g. the texture name) used to create this memory object.
1508
1509L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetGLObjectInfo.html>
1510
925=back 1511=back
926 1512
1513=head2 THE OpenCL::Buffer CLASS
1514
1515This is a subclass of OpenCL::Memory, and the superclass of
1516OpenCL::BufferObj. Its purpose is simply to distinguish between buffers
1517and sub-buffers.
1518
1519=head2 THE OpenCL::BufferObj CLASS
1520
1521This is a subclass of OpenCL::Buffer and thus OpenCL::Memory. It exists
1522because one cna create sub buffers of OpenLC::BufferObj objects, but not
1523sub buffers from these sub buffers.
1524
1525=over 4
1526
1527=item $subbuf = $buf_obj->sub_buffer_region ($flags, $origin, $size)
1528
1529Creates an OpenCL::Buffer objects from this buffer and returns it. The
1530C<buffer_create_type> is assumed to be C<CL_BUFFER_CREATE_TYPE_REGION>.
1531
1532L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateSubBuffer.html>
1533
1534=back
1535
927=head2 THE OpenCL::Image CLASS 1536=head2 THE OpenCL::Image CLASS
928 1537
929This is the superclass of all image objects - OpenCL::Image2D and OpenCL::Image3D. 1538This is the superclass of all image objects - OpenCL::Image1D,
1539OpenCL::Image1DArray, OpenCL::Image1DBuffer, OpenCL::Image2D,
1540OpenCL::Image2DArray and OpenCL::Image3D.
930 1541
931=over 4 1542=over 4
932 1543
933=item $packed_value = $ev->image_info ($name) 1544=item $packed_value = $image->image_info ($name)
934 1545
935See C<< $platform->info >> for details. 1546See C<< $platform->info >> for details.
936 1547
937The reason this method is not called C<info> is that there already is an 1548The reason this method is not called C<info> is that there already is an
938C<< ->info >> method inherited from C<OpenCL::Memory>. 1549C<< ->info >> method inherited from C<OpenCL::Memory>.
939 1550
940L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetImageInfo.html> 1551L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetImageInfo.html>
941 1552
1553=item ($channel_order, $channel_data_type) = $image->format
1554
1555Returns the channel order and type used to create the image by calling
1556C<clGetImageInfo> with C<CL_IMAGE_FORMAT>.
1557
942=for gengetinfo begin image 1558=for gengetinfo begin image
943 1559
944=item $int = $image->element_size 1560=item $int = $image->element_size
945 1561
946Calls C<clGetImageInfo> with C<CL_IMAGE_ELEMENT_SIZE> and returns the result. 1562Calls C<clGetImageInfo> with C<CL_IMAGE_ELEMENT_SIZE> and returns the result.
965 1581
966Calls C<clGetImageInfo> with C<CL_IMAGE_DEPTH> and returns the result. 1582Calls C<clGetImageInfo> with C<CL_IMAGE_DEPTH> and returns the result.
967 1583
968=for gengetinfo end image 1584=for gengetinfo end image
969 1585
1586=for gengetinfo begin gl_texture
1587
1588=item $GLenum = $gl_texture->target
1589
1590Calls C<clGetGLTextureInfo> with C<CL_GL_TEXTURE_TARGET> and returns the result.
1591
1592=item $GLint = $gl_texture->gl_mipmap_level
1593
1594Calls C<clGetGLTextureInfo> with C<CL_GL_MIPMAP_LEVEL> and returns the result.
1595
1596=for gengetinfo end gl_texture
1597
970=back 1598=back
971 1599
972=head2 THE OpenCL::Sampler CLASS 1600=head2 THE OpenCL::Sampler CLASS
973 1601
974=over 4 1602=over 4
1007 1635
1008=head2 THE OpenCL::Program CLASS 1636=head2 THE OpenCL::Program CLASS
1009 1637
1010=over 4 1638=over 4
1011 1639
1012=item $program->build ($device, $options = "") 1640=item $program->build (\@devices = undef, $options = "", $cb->($program) = undef)
1013 1641
1014Tries to build the program with the givne options. 1642Tries to build the program with the given options. See also the
1643C<$ctx->build> convenience function.
1644
1645If a callback is specified, then it will be called when compilation is
1646finished. Note that many OpenCL implementations block your program while
1647compiling whether you use a callback or not. See C<build_async> if you
1648want to make sure the build is done in the background.
1649
1650Note that some OpenCL implementations act up badly, and don't call the
1651callback in some error cases (but call it in others). This implementation
1652assumes the callback will always be called, and leaks memory if this is
1653not so. So best make sure you don't pass in invalid values.
1654
1655Some implementations fail with C<OpenCL::INVALID_BINARY> when the
1656compilation state is successful but some later stage fails.
1015 1657
1016L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clBuildProgram.html> 1658L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clBuildProgram.html>
1659
1660=item $program->build_async (\@devices = undef, $options = "", $cb->($program) = undef)
1661
1662Similar to C<< ->build >>, except it starts a thread, and never fails (you
1663need to check the compilation status form the callback, or by polling).
1017 1664
1018=item $packed_value = $program->build_info ($device, $name) 1665=item $packed_value = $program->build_info ($device, $name)
1019 1666
1020Similar to C<< $platform->info >>, but returns build info for a previous 1667Similar to C<< $platform->info >>, but returns build info for a previous
1021build attempt for the given device. 1668build attempt for the given device.
1026 1673
1027Creates an OpenCL::Kernel object out of the named C<__kernel> function in 1674Creates an OpenCL::Kernel object out of the named C<__kernel> function in
1028the program. 1675the program.
1029 1676
1030L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernel.html> 1677L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernel.html>
1678
1679=item @kernels = $program->kernels_in_program
1680
1681Returns all kernels successfully compiled for all devices in program.
1682
1683http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernelsInProgram.html
1031 1684
1032=for gengetinfo begin program_build 1685=for gengetinfo begin program_build
1033 1686
1034=item $build_status = $program->build_status ($device) 1687=item $build_status = $program->build_status ($device)
1035 1688
1157 1810
1158Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PRIVATE_MEM_SIZE> and returns the result. 1811Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PRIVATE_MEM_SIZE> and returns the result.
1159 1812
1160=for gengetinfo end kernel_work_group 1813=for gengetinfo end kernel_work_group
1161 1814
1815=item $kernel->setf ($format, ...)
1816
1817Sets the arguments of a kernel. Since OpenCL 1.1 doesn't have a generic
1818way to set arguments (and with OpenCL 1.2 it might be rather slow), you
1819need to specify a format argument, much as with C<printf>, to tell OpenCL
1820what type of argument it is.
1821
1822The format arguments are single letters:
1823
1824 c char
1825 C unsigned char
1826 s short
1827 S unsigned short
1828 i int
1829 I unsigned int
1830 l long
1831 L unsigned long
1832
1833 h half float (0..65535)
1834 f float
1835 d double
1836
1837 z local (octet size)
1838
1839 m memory object (buffer or image)
1840 a sampler
1841 e event
1842
1843Space characters in the format string are ignored.
1844
1845Example: set the arguments for a kernel that expects an int, two floats, a buffer and an image.
1846
1847 $kernel->setf ("i ff mm", 5, 0.5, 3, $buffer, $image);
1848
1162=item $kernel->set_TYPE ($index, $value) 1849=item $kernel->set_TYPE ($index, $value)
1163 1850
1851=item $kernel->set_char ($index, $value)
1852
1853=item $kernel->set_uchar ($index, $value)
1854
1855=item $kernel->set_short ($index, $value)
1856
1857=item $kernel->set_ushort ($index, $value)
1858
1859=item $kernel->set_int ($index, $value)
1860
1861=item $kernel->set_uint ($index, $value)
1862
1863=item $kernel->set_long ($index, $value)
1864
1865=item $kernel->set_ulong ($index, $value)
1866
1867=item $kernel->set_half ($index, $value)
1868
1869=item $kernel->set_float ($index, $value)
1870
1871=item $kernel->set_double ($index, $value)
1872
1873=item $kernel->set_memory ($index, $value)
1874
1875=item $kernel->set_buffer ($index, $value)
1876
1877=item $kernel->set_image ($index, $value)
1878
1879=item $kernel->set_sampler ($index, $value)
1880
1881=item $kernel->set_local ($index, $value)
1882
1883=item $kernel->set_event ($index, $value)
1884
1164This is a family of methods to set the kernel argument with the number C<$index> to the give C<$value>. 1885This is a family of methods to set the kernel argument with the number
1165 1886C<$index> to the give C<$value>.
1166TYPE is one of C<char>, C<uchar>, C<short>, C<ushort>, C<int>, C<uint>,
1167C<long>, C<ulong>, C<half>, C<float>, C<double>, C<memory>, C<buffer>,
1168C<image2d>, C<image3d>, C<sampler> or C<event>.
1169 1887
1170Chars and integers (including the half type) are specified as integers, 1888Chars and integers (including the half type) are specified as integers,
1171float and double as floating point values, memory/buffer/image2d/image3d 1889float and double as floating point values, memory/buffer/image must be
1172must be an object of that type or C<undef>, and sampler and event must be 1890an object of that type or C<undef>, local-memory arguments are set by
1173objects of that type. 1891specifying the size, and sampler and event must be objects of that type.
1892
1893Note that C<set_memory> works for all memory objects (all types of buffers
1894and images) - the main purpose of the more specific C<set_TYPE> functions
1895is type checking.
1896
1897Setting an argument for a kernel does NOT keep a reference to the object -
1898for example, if you set an argument to some image object, free the image,
1899and call the kernel, you will run into undefined behaviour.
1174 1900
1175L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html> 1901L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html>
1176 1902
1177=back 1903=back
1178 1904
1187 1913
1188Waits for the event to complete. 1914Waits for the event to complete.
1189 1915
1190L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html> 1916L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html>
1191 1917
1918=item $ev->cb ($exec_callback_type, $callback->($event, $event_command_exec_status))
1919
1920Adds a callback to the callback stack for the given event type. There is
1921no way to remove a callback again.
1922
1923L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetEventCallback.html>
1924
1192=item $packed_value = $ev->info ($name) 1925=item $packed_value = $ev->info ($name)
1193 1926
1194See C<< $platform->info >> for details. 1927See C<< $platform->info >> for details.
1195 1928
1196L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetEventInfo.html> 1929L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetEventInfo.html>
1256 1989
1257=over 4 1990=over 4
1258 1991
1259=item $ev->set_status ($execution_status) 1992=item $ev->set_status ($execution_status)
1260 1993
1994Sets the execution status of the user event. Can only be called once,
1995either with OpenCL::COMPLETE or a negative number as status.
1996
1261L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetUserEventStatus.html> 1997L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetUserEventStatus.html>
1262 1998
1263=back 1999=back
1264 2000
2001=head2 THE OpenCL::Mapped CLASS
2002
2003This class represents objects mapped into host memory. They are
2004represented by a blessed string scalar. The string data is the mapped
2005memory area, that is, if you read or write it, then the mapped object is
2006accessed directly.
2007
2008You must only ever use operations that modify the string in-place - for
2009example, a C<substr> that doesn't change the length, or maybe a regex that
2010doesn't change the length. Any other operation might cause the data to be
2011copied.
2012
2013When the object is destroyed it will enqueue an implicit unmap operation
2014on the queue that was used to create it.
2015
2016Keep in mind that you I<need> to unmap (or destroy) mapped objects before
2017OpenCL sees the changes, even if some implementations don't need this
2018sometimes.
2019
2020Example, replace the first two floats in the mapped buffer by 1 and 2.
2021
2022 my $mapped = $queue->map_buffer ($buf, ...
2023 $mapped->event->wait; # make sure it's there
2024
2025 # now replace first 8 bytes by new data, which is exactly 8 bytes long
2026 # we blindly assume device endianness to equal host endianness
2027 # (and of course, we assume iee 754 single precision floats :)
2028 substr $$mapped, 0, 8, pack "f*", 1, 2;
2029
2030=over 4
2031
2032=item $ev = $mapped->unmap ($wait_events...)
2033
2034Unmaps the mapped memory object, using the queue originally used to create
2035it, quite similarly to C<< $queue->unmap ($mapped, ...) >>.
2036
2037=item $bool = $mapped->mapped
2038
2039Returns whether the object is still mapped - true before an C<unmap> is
2040enqueued, false afterwards.
2041
2042=item $ev = $mapped->event
2043
2044Return the event object associated with the mapped object. Initially, this
2045will be the event object created when mapping the object, and after an
2046unmap, this will be the event object that the unmap operation created.
2047
2048=item $mapped->wait
2049
2050Same as C<< $mapped->event->wait >> - makes sure no operations on this
2051mapped object are outstanding.
2052
2053=item $bytes = $mapped->size
2054
2055Returns the size of the mapped area, in bytes. Same as C<length $$mapped>.
2056
2057=item $ptr = $mapped->ptr
2058
2059Returns the raw memory address of the mapped area.
2060
2061=item $mapped->set ($offset, $data)
2062
2063Replaces the data at the given C<$offset> in the memory area by the new
2064C<$data>. This method is safer than direct manipulation of C<$mapped>
2065because it does bounds-checking, but also slower.
2066
2067=item $data = $mapped->get ($offset, $length)
2068
2069Returns (without copying) a scalar representing the data at the given
2070C<$offset> and C<$length> in the mapped memory area. This is the same as
2071the following substr, except much slower;
2072
2073 $data = substr $$mapped, $offset, $length
2074
1265=cut 2075=cut
1266 2076
1267package OpenCL; 2077sub OpenCL::Mapped::get {
1268 2078 substr ${$_[0]}, $_[1], $_[2]
1269use common::sense;
1270
1271BEGIN {
1272 our $VERSION = '0.55';
1273
1274 require XSLoader;
1275 XSLoader::load (__PACKAGE__, $VERSION);
1276
1277 @OpenCL::Buffer::ISA =
1278 @OpenCL::Image::ISA = OpenCL::Memory::;
1279
1280 @OpenCL::Image2D::ISA =
1281 @OpenCL::Image3D::ISA = OpenCL::Image::;
1282
1283 @OpenCL::UserEvent::ISA = OpenCL::Event::;
1284} 2079}
2080
2081=back
2082
2083=head2 THE OpenCL::MappedBuffer CLASS
2084
2085This is a subclass of OpenCL::Mapped, representing mapped buffers.
2086
2087=head2 THE OpenCL::MappedImage CLASS
2088
2089This is a subclass of OpenCL::Mapped, representing mapped images.
2090
2091=over 4
2092
2093=item $bytes = $mapped->row_pitch
2094
2095=item $bytes = $mapped->slice_pitch
2096
2097Return the row or slice pitch of the image that has been mapped.
2098
2099=back
2100
2101
2102=cut
1285 2103
12861; 21041;
1287 2105
1288=head1 AUTHOR 2106=head1 AUTHOR
1289 2107

Diff Legend

Removed lines
+ Added lines
< Changed lines
> Changed lines