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

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.51 by root, Tue Apr 24 13:45:58 2012 UTC vs.
Revision 1.68 by root, Tue May 1 22:25:13 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.
171 # set buffer 174 # set buffer
172 $kernel->set_buffer (0, $input); 175 $kernel->set_buffer (0, $input);
173 $kernel->set_buffer (1, $output); 176 $kernel->set_buffer (1, $output);
174 177
175 # execute it for all 4 numbers 178 # execute it for all 4 numbers
176 $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); 179 $queue->nd_range_kernel ($kernel, undef, [4], undef);
177 180
178 # enqueue a synchronous read 181 # enqueue a synchronous read
179 $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);
180 183
181 # print the results: 184 # print the results:
182 printf "%s\n", join ", ", unpack "f*", $data; 185 printf "%s\n", join ", ", unpack "f*", $data;
183 186
184=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,
185showing off barriers. 188showing off barriers.
186 189
187 # execute it for all 4 numbers 190 # execute it for all 4 numbers
188 $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); 191 $queue->nd_range_kernel ($kernel, undef, [4], undef);
189 192
190 # enqueue a barrier to ensure in-order execution 193 # enqueue a barrier to ensure in-order execution
191 $queue->enqueue_barrier; 194 $queue->barrier;
192 195
193 # enqueue an async read 196 # enqueue an async read
194 $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);
195 198
196 # wait for all requests to finish 199 # wait for all requests to finish
197 $queue->finish; 200 $queue->finish;
198 201
199=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,
200showing off event objects and wait lists. 203showing off event objects and wait lists.
201 204
202 # execute it for all 4 numbers 205 # execute it for all 4 numbers
203 my $ev = $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); 206 my $ev = $queue->nd_range_kernel ($kernel, undef, [4], undef);
204 207
205 # enqueue an async read 208 # enqueue an async read
206 $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);
207 210
208 # wait for the last event to complete 211 # wait for the last event to complete
209 $ev->wait; 212 $ev->wait;
210 213
211=head2 Use the OpenGL module to share a texture between OpenCL and OpenGL and draw some julia 214=head2 Use the OpenGL module to share a texture between OpenCL and OpenGL and draw some julia
212set tunnel effect. 215set tunnel effect.
213 216
214This is quite a long example to get you going. 217This is quite a long example to get you going - you can download it from
218L<http://cvs.schmorp.de/OpenCL/examples/juliaflight>.
215 219
216 use OpenGL ":all"; 220 use OpenGL ":all";
217 use OpenCL; 221 use OpenCL;
218 222
223 my $S = $ARGV[0] || 256; # window/texture size, smaller is faster
224
219 # open a window and create a gl texture 225 # open a window and create a gl texture
220 OpenGL::glpOpenWindow width => 256, height => 256; 226 OpenGL::glpOpenWindow width => $S, height => $S;
221 my $texid = glGenTextures_p 1; 227 my $texid = glGenTextures_p 1;
222 glBindTexture GL_TEXTURE_2D, $texid; 228 glBindTexture GL_TEXTURE_2D, $texid;
223 glTexImage2D_c GL_TEXTURE_2D, 0, GL_RGBA8, 256, 256, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0; 229 glTexImage2D_c GL_TEXTURE_2D, 0, GL_RGBA8, $S, $S, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0;
224 230
225 # find and use the first opencl device that let's us get a shared opengl context 231 # find and use the first opencl device that let's us get a shared opengl context
226 my $platform; 232 my $platform;
227 my $dev; 233 my $dev;
228 my $ctx; 234 my $ctx;
247 # now the boring opencl code 253 # now the boring opencl code
248 my $src = <<EOF; 254 my $src = <<EOF;
249 kernel void 255 kernel void
250 juliatunnel (write_only image2d_t img, float time) 256 juliatunnel (write_only image2d_t img, float time)
251 { 257 {
252 float2 p = (float2)(get_global_id (0), get_global_id (1)) / 256.f * 2.f - 1.f; 258 int2 xy = (int2)(get_global_id (0), get_global_id (1));
259 float2 p = convert_float2 (xy) / $S.f * 2.f - 1.f;
253 260
254 float2 m = (float2)(1.f, p.y) / fabs (p.x); 261 float2 m = (float2)(1.f, p.y) / fabs (p.x); // tunnel
255 m.x = fabs (fmod (m.x + time * 0.05f, 4.f)) - 2.f; 262 m.x = fabs (fmod (m.x + time * 0.05f, 4.f) - 2.f);
256 263
257 float2 z = m; 264 float2 z = m;
258 float2 c = (float2)(sin (time * 0.05005), cos (time * 0.06001)); 265 float2 c = (float2)(sin (time * 0.01133f), cos (time * 0.02521f));
259 266
260 for (int i = 0; i < 25 && dot (z, z) < 4.f; ++i) 267 for (int i = 0; i < 25 && dot (z, z) < 4.f; ++i) // standard julia
261 z = (float2)(z.x * z.x - z.y * z.y, 2.f * z.x * z.y) + c; 268 z = (float2)(z.x * z.x - z.y * z.y, 2.f * z.x * z.y) + c;
262 269
263 float3 colour = (float3)(z.x, z.y, z.x * z.y); 270 float3 colour = (float3)(z.x, z.y, atan2 (z.y, z.x));
264 write_imagef (img, (int2)(get_global_id (0), get_global_id (1)), (float4)(colour * p.x * p.x, 1.)); 271 write_imagef (img, xy, (float4)(colour * p.x * p.x, 1.));
265 } 272 }
266 EOF 273 EOF
267 274
268 my $prog = $ctx->build_program ($src); 275 my $prog = $ctx->build_program ($src);
269 my $kernel = $prog->kernel ("juliatunnel"); 276 my $kernel = $prog->kernel ("juliatunnel");
270 277
271 # program compiled, kernel ready, now draw and loop 278 # program compiled, kernel ready, now draw and loop
272 279
273 for (my $time; ; ++$time) { 280 for (my $time; ; ++$time) {
274 # acquire objects from opengl 281 # acquire objects from opengl
275 $queue->enqueue_acquire_gl_objects ([$tex]); 282 $queue->acquire_gl_objects ([$tex]);
276 283
277 # configure and run our kernel 284 # configure and run our kernel
278 $kernel->set_image2d (0, $tex); 285 $kernel->setf ("mf", $tex, $time*2); # mf = memory object, float
279 $kernel->set_float (1, $time);
280 $queue->enqueue_nd_range_kernel ($kernel, undef, [256, 256], undef); 286 $queue->nd_range_kernel ($kernel, undef, [$S, $S], undef);
281 287
282 # release objects to opengl again 288 # release objects to opengl again
283 $queue->enqueue_release_gl_objects ([$tex]); 289 $queue->release_gl_objects ([$tex]);
284 290
285 # wait 291 # wait
286 $queue->finish; 292 $queue->finish;
287 293
288 # now draw the texture, the defaults should be all right 294 # now draw the texture, the defaults should be all right
298 304
299 glXSwapBuffers; 305 glXSwapBuffers;
300 306
301 select undef, undef, undef, 1/60; 307 select undef, undef, undef, 1/60;
302 } 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>.
303 336
304=head1 DOCUMENTATION 337=head1 DOCUMENTATION
305 338
306=head2 BASIC CONVENTIONS 339=head2 BASIC CONVENTIONS
307 340
376 409
377For this to work, the OpenGL library must be loaded, a GLX context must 410For this to work, the OpenGL library must be loaded, a GLX context must
378have been created and be made current, and C<dlsym> must be available and 411have been created and be made current, and C<dlsym> must be available and
379capable of finding the function via C<RTLD_DEFAULT>. 412capable of finding the function via C<RTLD_DEFAULT>.
380 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.98';
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
381=head2 THE OpenCL PACKAGE 567=head2 THE OpenCL PACKAGE
382 568
383=over 4 569=over 4
384 570
385=item $int = OpenCL::errno 571=item $int = OpenCL::errno
386 572
387The 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
388and before calling another OpenCL function. 574and before calling another OpenCL function.
389 575
390=item $str = OpenCL::err2str $errval 576=item $str = OpenCL::err2str [$errval]
391 577
392Comverts 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).
393 580
394=item $str = OpenCL::enum2str $enum 581=item $str = OpenCL::enum2str $enum
395 582
396Converts most enum values (of parameter names, image format constants, 583Converts most enum values (of parameter names, image format constants,
397object types, addressing and filter modes, command types etc.) into a 584object types, addressing and filter modes, command types etc.) into a
403 590
404Returns all available OpenCL::Platform objects. 591Returns all available OpenCL::Platform objects.
405 592
406L<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>
407 594
408=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
409 596
410Tries 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.
411 598
412L<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>
413 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
414=item OpenCL::wait_for_events $wait_events... 608=item OpenCL::wait_for_events $wait_events...
415 609
416Waits for all events to complete. 610Waits for all events to complete.
417 611
418L<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>
419 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
420=back 651=back
421 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
422=head2 THE OpenCL::Platform CLASS 678=head2 THE OpenCL::Platform CLASS
423 679
424=over 4 680=over 4
425 681
426=item @devices = $platform->devices ($type = OpenCL::DEVICE_TYPE_ALL) 682=item @devices = $platform->devices ($type = OpenCL::DEVICE_TYPE_ALL)
427 683
428Returns a list of matching OpenCL::Device objects. 684Returns a list of matching OpenCL::Device objects.
429 685
430=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)
431 687
432Tries 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.
433 689
434L<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>
435 691
436=item $ctx = $platform->context ($properties = undef, @$devices, $notify = undef) 692=item $ctx = $platform->context ($properties, \@devices, $callback->($err, $pvt) = $print_stderr)
437 693
438Create 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
439CL_CONTEXT_PLATFORM property is supplied automatically. 695CL_CONTEXT_PLATFORM property is supplied automatically.
440 696
441L<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>
784 my ($self, $prog, $options) = @_; 1040 my ($self, $prog, $options) = @_;
785 1041
786 $prog = $self->program_with_source ($prog) 1042 $prog = $self->program_with_source ($prog)
787 unless ref $prog; 1043 unless ref $prog;
788 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
789 for my $dev ($self->devices) { 1051 for my $dev ($self->devices) {
790 eval { $prog->build ($dev, $options); 1 } 1052 $prog->build_status ($dev) == BUILD_SUCCESS
791 or Carp::croak "Building OpenCL program for device '" . $dev->name . "' failed:\n" 1053 or Carp::croak "Building OpenCL program for device '" . $dev->name . "' failed:\n"
792 . $prog->build_log ($dev); 1054 . $prog->build_log ($dev);
793 } 1055 }
794 1056
795 $prog 1057 $prog
796} 1058}
797 1059
821=item $buf = $ctx->buffer_sv ($flags, $data) 1083=item $buf = $ctx->buffer_sv ($flags, $data)
822 1084
823Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object and 1085Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object and
824initialise it with the given data values. 1086initialise it with the given data values.
825 1087
826=item $img = $ctx->image ($self, $flags, $channel_order, $channel_type, $type, $width, $height, $depth, $array_size = 0, $row_pitch = 0, $slice_pitch = 0, $num_mip_level = 0, $num_samples = 0, $*data = &PL_sv_undef) 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)
827 1089
828Creates a new OpenCL::Image object and optionally initialises it with 1090Creates a new OpenCL::Image object and optionally initialises it with
829the given data values. 1091the given data values.
830 1092
831L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateImage.html> 1093L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateImage.html>
927=back 1189=back
928 1190
929=head2 THE OpenCL::Queue CLASS 1191=head2 THE OpenCL::Queue CLASS
930 1192
931An OpenCL::Queue represents an execution queue for OpenCL. You execute 1193An OpenCL::Queue represents an execution queue for OpenCL. You execute
932requests by calling their respective C<enqueue_xxx> method and waitinf for 1194requests by calling their respective method and waiting for it to complete
933it to complete in some way. 1195in some way.
934 1196
935All the enqueue methods return an event object that can be used to wait 1197Most methods that enqueue some request return an event object that can
936for completion, unless the method is called in void context, in which case 1198be used to wait for completion (optionally using a callback), unless
937no event object is created. 1199the method is called in void context, in which case no event object is
1200created.
938 1201
939They also allow you to specify any number of other event objects that this 1202They also allow you to specify any number of other event objects that this
940request has to wait for before it starts executing, by simply passing the 1203request has to wait for before it starts executing, by simply passing the
941event objects as extra parameters to the enqueue methods. To simplify 1204event objects as extra parameters to the enqueue methods. To simplify
942program design, this module ignores any C<undef> values in the list of 1205program design, this module ignores any C<undef> values in the list of
943events. This makes it possible to code operations such as this, without 1206events. This makes it possible to code operations such as this, without
944having to put a valid event object into C<$event> first: 1207having to put a valid event object into C<$event> first:
945 1208
946 $event = $queue->enqueue_xxx (..., $event); 1209 $event = $queue->xxx (..., $event);
947 1210
948Queues execute in-order by default, without any parallelism, so in most 1211Queues execute in-order by default, without any parallelism, so in most
949cases (i.e. you use only one queue) it's not necessary to wait for or 1212cases (i.e. you use only one queue) it's not necessary to wait for or
950create event objects, althoguh an our of order queue is often a bit 1213create event objects, althoguh an our of order queue is often a bit
951faster. 1214faster.
952 1215
953=over 4 1216=over 4
954 1217
955=item $ev = $queue->enqueue_read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...) 1218=item $ev = $queue->read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...)
956 1219
957Reads data from buffer into the given string. 1220Reads data from buffer into the given string.
958 1221
959L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBuffer.html> 1222L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBuffer.html>
960 1223
961=item $ev = $queue->enqueue_write_buffer ($buffer, $blocking, $offset, $data, $wait_events...) 1224=item $ev = $queue->write_buffer ($buffer, $blocking, $offset, $data, $wait_events...)
962 1225
963Writes data to buffer from the given string. 1226Writes data to buffer from the given string.
964 1227
965L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBuffer.html> 1228L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBuffer.html>
966 1229
967=item $ev = $queue->enqueue_copy_buffer ($src, $dst, $src_offset, $dst_offset, $len, $wait_events...) 1230=item $ev = $queue->copy_buffer ($src, $dst, $src_offset, $dst_offset, $len, $wait_events...)
968 1231
969L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBuffer.html> 1232L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBuffer.html>
970 1233
971=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...) 1234=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...)
972 1235
973http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBufferRect.html 1236http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBufferRect.html
974 1237
975=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...) 1238=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...)
976 1239
977http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBufferRect.html 1240http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBufferRect.html
978 1241
979=item $ev = $queue->enqueue_read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
980
981L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferRect.html>
982
983=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...) 1242=item $ev = $queue->copy_buffer_to_image ($src_buffer, $dst_image, $src_offset, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $wait_events...)
1243
1244L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>
1245
1246=item $ev = $queue->read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
1247
1248C<$row_pitch> (and C<$slice_pitch>) can be C<0>, in which case the OpenCL
1249module uses the image width (and height) to supply default values.
984 1250
985L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadImage.html> 1251L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadImage.html>
986 1252
987=item $ev = $queue->enqueue_write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...) 1253=item $ev = $queue->write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
988 1254
1255C<$row_pitch> (and C<$slice_pitch>) can be C<0>, in which case the OpenCL
1256module uses the image width (and height) to supply default values.
989L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteImage.html> 1257L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteImage.html>
990 1258
991=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...) 1259=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...)
992 1260
993L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html> 1261L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html>
994 1262
995=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...) 1263=item $ev = $queue->copy_image_to_buffer ($src_image, $dst_image, $src_x, $src_y, $src_z, $width, $height, $depth, $dst_offset, $wait_events...)
996 1264
997L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html> 1265L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html>
998 1266
999=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...) 1267=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...)
1000 1268
1001Yeah. 1269Yeah.
1002 1270
1003L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>. 1271L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>.
1004 1272
1273=item $ev = $queue->fill_buffer ($mem, $pattern, $offset, $size, ...)
1274
1275Fills the given buffer object with repeated applications of C<$pattern>,
1276starting at C<$offset> for C<$size> octets.
1277
1278L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillBuffer.html>
1279
1280=item $ev = $queue->fill_image ($img, $r, $g, $b, $a, $x, $y, $z, $width, $height, $depth, ...)
1281
1282Fills the given image area with the given rgba colour components. The
1283components are normally floating point values between C<0> and C<1>,
1284except when the image channel data type is a signe dor unsigned
1285unnormalised format, in which case the range is determined by the format.
1286
1287L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillImage.html>
1288
1005=item $ev = $queue->enqueue_task ($kernel, $wait_events...) 1289=item $ev = $queue->task ($kernel, $wait_events...)
1006 1290
1007L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html> 1291L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html>
1008 1292
1009=item $ev = $queue->enqueue_nd_range_kernel ($kernel, @$global_work_offset, @$global_work_size, @$local_work_size, $wait_events...) 1293=item $ev = $queue->nd_range_kernel ($kernel, \@global_work_offset, \@global_work_size, \@local_work_size, $wait_events...)
1010 1294
1011Enqueues a kernel execution. 1295Enqueues a kernel execution.
1012 1296
1013@$global_work_size must be specified as a reference to an array of 1297\@global_work_size must be specified as a reference to an array of
1014integers specifying the work sizes (element counts). 1298integers specifying the work sizes (element counts).
1015 1299
1016@$global_work_offset must be either C<undef> (in which case all offsets 1300\@global_work_offset must be either C<undef> (in which case all offsets
1017are C<0>), or a reference to an array of work offsets, with the same number 1301are C<0>), or a reference to an array of work offsets, with the same number
1018of elements as @$global_work_size. 1302of elements as \@global_work_size.
1019 1303
1020@$local_work_size must be either C<undef> (in which case the 1304\@local_work_size must be either C<undef> (in which case the
1021implementation is supposed to choose good local work sizes), or a 1305implementation is supposed to choose good local work sizes), or a
1022reference to an array of local work sizes, with the same number of 1306reference to an array of local work sizes, with the same number of
1023elements as @$global_work_size. 1307elements as \@global_work_size.
1024 1308
1025L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html> 1309L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html>
1026 1310
1027=item $ev = $queue->enqueue_acquire_gl_objects ([object, ...], $wait_events...) 1311=item $ev = $queue->acquire_gl_objects ([object, ...], $wait_events...)
1028 1312
1029Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired 1313Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired
1030for subsequent OpenCL usage. 1314for subsequent OpenCL usage.
1031 1315
1032L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueAcquireGLObjects.html> 1316L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueAcquireGLObjects.html>
1033 1317
1034=item $ev = $queue->enqueue_release_gl_objects ([object, ...], $wait_events...) 1318=item $ev = $queue->release_gl_objects ([object, ...], $wait_events...)
1035 1319
1036Enqueues a list (an array-ref of OpenCL::Memory objects) to be released 1320Enqueues a list (an array-ref of OpenCL::Memory objects) to be released
1037for subsequent OpenGL usage. 1321for subsequent OpenGL usage.
1038 1322
1039L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReleaseGLObjects.html> 1323L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReleaseGLObjects.html>
1040 1324
1041=item $ev = $queue->enqueue_wait_for_events ($wait_events...) 1325=item $ev = $queue->wait_for_events ($wait_events...)
1042 1326
1043L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html> 1327L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html>
1044 1328
1045=item $ev = $queue->enqueue_marker ($wait_events...) 1329=item $ev = $queue->marker ($wait_events...)
1046 1330
1047L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMarkerWithWaitList.html> 1331L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMarkerWithWaitList.html>
1048 1332
1049=item $ev = $queue->enqueue_barrier ($wait_events...) 1333=item $ev = $queue->barrier ($wait_events...)
1050 1334
1051L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueBarrierWithWaitList.html> 1335L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueBarrierWithWaitList.html>
1052 1336
1053=item $queue->flush 1337=item $queue->flush
1054 1338
1081=item $command_queue_properties = $command_queue->properties 1365=item $command_queue_properties = $command_queue->properties
1082 1366
1083Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_PROPERTIES> and returns the result. 1367Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_PROPERTIES> and returns the result.
1084 1368
1085=for gengetinfo end command_queue 1369=for gengetinfo end command_queue
1370
1371=back
1372
1373=head3 MEMORY MAPPED BUFFERS
1374
1375OpenCL allows you to map buffers and images to host memory (read: perl
1376scalars). This is done much like reading or copying a buffer, by enqueuing
1377a map or unmap operation on the command queue.
1378
1379The map operations return a C<OpenCL::Mapped> object - see L<THE
1380OpenCL::Mapped CLASS> section for details on what to do with these
1381objects.
1382
1383The object will be unmapped automatically when the mapped object is
1384destroyed (you can use a barrier to make sure the unmap has finished,
1385before using the buffer in a kernel), but you can also enqueue an unmap
1386operation manually.
1387
1388=over 4
1389
1390=item $mapped_buffer = $queue->map_buffer ($buf, $data, $blocking=1, $map_flags=OpenCL::MAP_READ|OpenCL::MAP_WRITE, $offset=0, $size=0, $wait_events...)
1391
1392Maps the given buffer into host memory and returns a C<OpenCL::MappedBuffer> object.
1393
1394L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMapBuffer.html>
1395
1396=item $mapped_image = $queue->map_image ($img, $data, $blocking=1, $map_flags=OpenCL::MAP_READ|OpenCL::MAP_WRITE, $x=0, $y=0, $z=0, $width=0, $height=0, $depth=0, $wait_events...)
1397
1398Maps the given image area into host memory and return a
1399C<OpenCL::MappedImage> object. Although there are default values for most
1400arguments, you currently have to specify all arguments, otherwise the call
1401will fail.
1402
1403L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMapImage.html>
1404
1405=item $ev = $queue->unmap ($mapped, $wait_events...)
1406
1407Unmaps the data from host memory. You must not call any methods that
1408modify the data, or modify the data scalar directly, after calling this
1409method.
1410
1411The mapped event object will always be passed as part of the
1412$wait_events. The mapped event object will be replaced by the new event
1413object that this request creates.
1086 1414
1087=back 1415=back
1088 1416
1089=head2 THE OpenCL::Memory CLASS 1417=head2 THE OpenCL::Memory CLASS
1090 1418
1177OpenCL::Image1DArray, OpenCL::Image1DBuffer, OpenCL::Image2D, 1505OpenCL::Image1DArray, OpenCL::Image1DBuffer, OpenCL::Image2D,
1178OpenCL::Image2DArray and OpenCL::Image3D. 1506OpenCL::Image2DArray and OpenCL::Image3D.
1179 1507
1180=over 4 1508=over 4
1181 1509
1182=item $packed_value = $ev->image_info ($name) 1510=item $packed_value = $image->image_info ($name)
1183 1511
1184See C<< $platform->info >> for details. 1512See C<< $platform->info >> for details.
1185 1513
1186The reason this method is not called C<info> is that there already is an 1514The reason this method is not called C<info> is that there already is an
1187C<< ->info >> method inherited from C<OpenCL::Memory>. 1515C<< ->info >> method inherited from C<OpenCL::Memory>.
1188 1516
1189L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetImageInfo.html> 1517L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetImageInfo.html>
1190 1518
1519=item ($channel_order, $channel_data_type) = $image->format
1520
1521Returns the channel order and type used to create the image by calling
1522C<clGetImageInfo> with C<CL_IMAGE_FORMAT>.
1523
1191=for gengetinfo begin image 1524=for gengetinfo begin image
1192 1525
1193=item $int = $image->element_size 1526=item $int = $image->element_size
1194 1527
1195Calls C<clGetImageInfo> with C<CL_IMAGE_ELEMENT_SIZE> and returns the result. 1528Calls C<clGetImageInfo> with C<CL_IMAGE_ELEMENT_SIZE> and returns the result.
1268 1601
1269=head2 THE OpenCL::Program CLASS 1602=head2 THE OpenCL::Program CLASS
1270 1603
1271=over 4 1604=over 4
1272 1605
1273=item $program->build ($device, $options = "") 1606=item $program->build (\@devices = undef, $options = "", $cb->($program) = undef)
1274 1607
1275Tries to build the program with the given options. See also the 1608Tries to build the program with the given options. See also the
1276C<$ctx->build> convenience function. 1609C<$ctx->build> convenience function.
1277 1610
1611If a callback is specified, then it will be called when compilation is
1612finished. Note that many OpenCL implementations block your program while
1613compiling whether you use a callback or not. See C<build_async> if you
1614want to make sure the build is done in the background.
1615
1616Note that some OpenCL implementations act up badly, and don't call the
1617callback in some error cases (but call it in others). This implementation
1618assumes the callback will always be called, and leaks memory if this is
1619not so. So best make sure you don't pass in invalid values.
1620
1621Some implementations fail with C<OpenCL::INVALID_BINARY> when the
1622compilation state is successful but some later stage fails.
1623
1278L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clBuildProgram.html> 1624L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clBuildProgram.html>
1625
1626=item $program->build_async (\@devices = undef, $options = "", $cb->($program) = undef)
1627
1628Similar to C<< ->build >>, except it starts a thread, and never fails (you
1629need to check the compilation status form the callback, or by polling).
1279 1630
1280=item $packed_value = $program->build_info ($device, $name) 1631=item $packed_value = $program->build_info ($device, $name)
1281 1632
1282Similar to C<< $platform->info >>, but returns build info for a previous 1633Similar to C<< $platform->info >>, but returns build info for a previous
1283build attempt for the given device. 1634build attempt for the given device.
1425 1776
1426Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PRIVATE_MEM_SIZE> and returns the result. 1777Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PRIVATE_MEM_SIZE> and returns the result.
1427 1778
1428=for gengetinfo end kernel_work_group 1779=for gengetinfo end kernel_work_group
1429 1780
1781=item $kernel->setf ($format, ...)
1782
1783Sets the arguments of a kernel. Since OpenCL 1.1 doesn't have a generic
1784way to set arguments (and with OpenCL 1.2 it might be rather slow), you
1785need to specify a format argument, much as with C<printf>, to tell OpenCL
1786what type of argument it is.
1787
1788The format arguments are single letters:
1789
1790 c char
1791 C unsigned char
1792 s short
1793 S unsigned short
1794 i int
1795 I unsigned int
1796 l long
1797 L unsigned long
1798
1799 h half float (0..65535)
1800 f float
1801 d double
1802
1803 z local (octet size)
1804
1805 m memory object (buffer or image)
1806 a sampler
1807 e event
1808
1809Space characters in the format string are ignored.
1810
1811Example: set the arguments for a kernel that expects an int, two floats, a buffer and an image.
1812
1813 $kernel->setf ("i ff mm", 5, 0.5, 3, $buffer, $image);
1814
1430=item $kernel->set_TYPE ($index, $value) 1815=item $kernel->set_TYPE ($index, $value)
1431 1816
1817=item $kernel->set_char ($index, $value)
1818
1819=item $kernel->set_uchar ($index, $value)
1820
1821=item $kernel->set_short ($index, $value)
1822
1823=item $kernel->set_ushort ($index, $value)
1824
1825=item $kernel->set_int ($index, $value)
1826
1827=item $kernel->set_uint ($index, $value)
1828
1829=item $kernel->set_long ($index, $value)
1830
1831=item $kernel->set_ulong ($index, $value)
1832
1833=item $kernel->set_half ($index, $value)
1834
1835=item $kernel->set_float ($index, $value)
1836
1837=item $kernel->set_double ($index, $value)
1838
1839=item $kernel->set_memory ($index, $value)
1840
1841=item $kernel->set_buffer ($index, $value)
1842
1843=item $kernel->set_image ($index, $value)
1844
1845=item $kernel->set_sampler ($index, $value)
1846
1847=item $kernel->set_local ($index, $value)
1848
1849=item $kernel->set_event ($index, $value)
1850
1432This is a family of methods to set the kernel argument with the number C<$index> to the give C<$value>. 1851This is a family of methods to set the kernel argument with the number
1433 1852C<$index> to the give C<$value>.
1434TYPE is one of C<char>, C<uchar>, C<short>, C<ushort>, C<int>, C<uint>,
1435C<long>, C<ulong>, C<half>, C<float>, C<double>, C<memory>, C<buffer>,
1436C<image2d>, C<image3d>, C<sampler>, C<local> or C<event>.
1437 1853
1438Chars and integers (including the half type) are specified as integers, 1854Chars and integers (including the half type) are specified as integers,
1439float and double as floating point values, memory/buffer/image2d/image3d 1855float and double as floating point values, memory/buffer/image must be
1440must be an object of that type or C<undef>, local-memory arguments are 1856an object of that type or C<undef>, local-memory arguments are set by
1441set by specifying the size, and sampler and event must be objects of that 1857specifying the size, and sampler and event must be objects of that type.
1442type. 1858
1859Note that C<set_memory> works for all memory objects (all types of buffers
1860and images) - the main purpose of the more specific C<set_TYPE> functions
1861is type checking.
1443 1862
1444Setting an argument for a kernel does NOT keep a reference to the object - 1863Setting an argument for a kernel does NOT keep a reference to the object -
1445for example, if you set an argument to some image object, free the image, 1864for example, if you set an argument to some image object, free the image,
1446and call the kernel, you will run into undefined behaviour. 1865and call the kernel, you will run into undefined behaviour.
1447 1866
1460 1879
1461Waits for the event to complete. 1880Waits for the event to complete.
1462 1881
1463L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html> 1882L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html>
1464 1883
1884=item $ev->cb ($exec_callback_type, $callback->($event, $event_command_exec_status))
1885
1886Adds a callback to the callback stack for the given event type. There is
1887no way to remove a callback again.
1888
1889L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetEventCallback.html>
1890
1465=item $packed_value = $ev->info ($name) 1891=item $packed_value = $ev->info ($name)
1466 1892
1467See C<< $platform->info >> for details. 1893See C<< $platform->info >> for details.
1468 1894
1469L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetEventInfo.html> 1895L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetEventInfo.html>
1529 1955
1530=over 4 1956=over 4
1531 1957
1532=item $ev->set_status ($execution_status) 1958=item $ev->set_status ($execution_status)
1533 1959
1960Sets the execution status of the user event. Can only be called once,
1961either with OpenCL::COMPLETE or a negative number as status.
1962
1534L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetUserEventStatus.html> 1963L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetUserEventStatus.html>
1535 1964
1536=back 1965=back
1537 1966
1967=head2 THE OpenCL::Mapped CLASS
1968
1969This class represents objects mapped into host memory. They are
1970represented by a blessed string scalar. The string data is the mapped
1971memory area, that is, if you read or write it, then the mapped object is
1972accessed directly.
1973
1974You must only ever use operations that modify the string in-place - for
1975example, a C<substr> that doesn't change the length, or maybe a regex that
1976doesn't change the length. Any other operation might cause the data to be
1977copied.
1978
1979When the object is destroyed it will enqueue an implicit unmap operation
1980on the queue that was used to create it.
1981
1982Keep in mind that you I<need> to unmap (or destroy) mapped objects before
1983OpenCL sees the changes, even if some implementations don't need this
1984sometimes.
1985
1986Example, replace the first two floats in the mapped buffer by 1 and 2.
1987
1988 my $mapped = $queue->map_buffer ($buf, ...
1989 $mapped->event->wait; # make sure it's there
1990
1991 # now replace first 8 bytes by new data, which is exactly 8 bytes long
1992 # we blindly assume device endianness to equal host endianness
1993 # (and of course, we assume iee 754 single precision floats :)
1994 substr $$mapped, 0, 8, pack "f*", 1, 2;
1995
1996=over 4
1997
1998=item $ev = $mapped->unmap ($wait_events...)
1999
2000Unmaps the mapped memory object, using the queue originally used to create
2001it, quite similarly to C<< $queue->unmap ($mapped, ...) >>.
2002
2003=item $bool = $mapped->mapped
2004
2005Returns whether the object is still mapped - true before an C<unmap> is
2006enqueued, false afterwards.
2007
2008=item $ev = $mapped->event
2009
2010Return the event object associated with the mapped object. Initially, this
2011will be the event object created when mapping the object, and after an
2012unmap, this will be the event object that the unmap operation created.
2013
2014=item $mapped->wait
2015
2016Same as C<< $mapped->event->wait >> - makes sure no operations on this
2017mapped object are outstanding.
2018
2019=item $bytes = $mapped->size
2020
2021Returns the size of the mapped area, in bytes. Same as C<length $$mapped>.
2022
2023=item $ptr = $mapped->ptr
2024
2025Returns the raw memory address of the mapped area.
2026
2027=item $mapped->set ($offset, $data)
2028
2029Replaces the data at the given C<$offset> in the memory area by the new
2030C<$data>. This method is safer than direct manipulation of C<$mapped>
2031because it does bounds-checking, but also slower.
2032
2033=item $data = $mapped->get ($offset, $length)
2034
2035Returns (without copying) a scalar representing the data at the given
2036C<$offset> and C<$length> in the mapped memory area. This is the same as
2037the following substr, except much slower;
2038
2039 $data = substr $$mapped, $offset, $length
2040
1538=cut 2041=cut
1539 2042
1540package OpenCL; 2043sub OpenCL::Mapped::get {
1541 2044 substr ${$_[0]}, $_[1], $_[2]
1542use common::sense;
1543
1544BEGIN {
1545 our $VERSION = '0.96';
1546
1547 require XSLoader;
1548 XSLoader::load (__PACKAGE__, $VERSION);
1549
1550 @OpenCL::Buffer::ISA =
1551 @OpenCL::Image::ISA = OpenCL::Memory::;
1552
1553 @OpenCL::BufferObj::ISA = OpenCL::Buffer::;
1554
1555 @OpenCL::Image2D::ISA =
1556 @OpenCL::Image3D::ISA =
1557 @OpenCL::Image2DArray::ISA =
1558 @OpenCL::Image1D::ISA =
1559 @OpenCL::Image1DArray::ISA =
1560 @OpenCL::Image1DBuffer::ISA = OpenCL::Image::;
1561
1562 @OpenCL::UserEvent::ISA = OpenCL::Event::;
1563} 2045}
2046
2047=back
2048
2049=head2 THE OpenCL::MappedBuffer CLASS
2050
2051This is a subclass of OpenCL::Mapped, representing mapped buffers.
2052
2053=head2 THE OpenCL::MappedImage CLASS
2054
2055This is a subclass of OpenCL::Mapped, representing mapped images.
2056
2057=over 4
2058
2059=item $bytes = $mapped->row_pitch
2060
2061=item $bytes = $mapped->slice_pitch
2062
2063Return the row or slice pitch of the image that has been mapped.
2064
2065=back
2066
2067
2068=cut
1564 2069
15651; 20701;
1566 2071
1567=head1 AUTHOR 2072=head1 AUTHOR
1568 2073

Diff Legend

Removed lines
+ Added lines
< Changed lines
> Changed lines