ViewVC Help
View File | Revision Log | Show Annotations | Download File
/cvs/OpenCL/OpenCL.pm
Revision: 1.88
Committed: Tue Oct 22 17:25:38 2013 UTC (10 years, 8 months ago) by root
Branch: MAIN
CVS Tags: HEAD
Changes since 1.87: +15 -8 lines
Log Message:
*** empty log message ***

File Contents

# Content
1 =head1 NAME
2
3 OpenCL - Open Computing Language Bindings
4
5 =head1 SYNOPSIS
6
7 use OpenCL;
8
9 =head1 DESCRIPTION
10
11 This is an early release which might be useful, but hasn't seen much testing.
12
13 =head2 OpenCL FROM 10000 FEET HEIGHT
14
15 Here is a high level overview of OpenCL:
16
17 First you need to find one or more OpenCL::Platforms (kind of like
18 vendors) - usually there is only one.
19
20 Each platform gives you access to a number of OpenCL::Device objects, e.g.
21 your graphics card.
22
23 From a platform and some device(s), you create an OpenCL::Context, which is
24 a very central object in OpenCL: Once you have a context you can create
25 most other objects:
26
27 OpenCL::Program objects, which store source code and, after building for a
28 specific device ("compiling and linking"), also binary programs. For each
29 kernel function in a program you can then create an OpenCL::Kernel object
30 which represents basically a function call with argument values.
31
32 OpenCL::Memory objects of various flavours: OpenCL::Buffer objects (flat
33 memory areas, think arrays or structs) and OpenCL::Image objects (think 2D
34 or 3D array) for bulk data and input and output for kernels.
35
36 OpenCL::Sampler objects, which are kind of like texture filter modes in
37 OpenGL.
38
39 OpenCL::Queue objects - command queues, which allow you to submit memory
40 reads, writes and copies, as well as kernel calls to your devices. They
41 also offer a variety of methods to synchronise request execution, for
42 example with barriers or OpenCL::Event objects.
43
44 OpenCL::Event objects are used to signal when something is complete.
45
46 =head2 HELPFUL RESOURCES
47
48 The OpenCL specs used to develop this module - download these and keept
49 hema round, they are required reference material:
50
51 http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
52 http://www.khronos.org/registry/cl/specs/opencl-1.2.pdf
53 http://www.khronos.org/registry/cl/specs/opencl-1.2-extensions.pdf
54
55 OpenCL manpages:
56
57 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/
58 http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/
59
60 If you are into UML class diagrams, the following diagram might help - if
61 not, it will be mildly confusing (also, the class hierarchy of this module
62 is much more fine-grained):
63
64 http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/classDiagram.html
65
66 Here's a tutorial from AMD (very AMD-centric, too), not sure how useful it
67 is, but at least it's free of charge:
68
69 http://developer.amd.com/zones/OpenCLZone/courses/Documents/Introduction_to_OpenCL_Programming%20Training_Guide%20%28201005%29.pdf
70
71 And here's NVIDIA's OpenCL Best Practises Guide:
72
73 http://developer.download.nvidia.com/compute/cuda/3_2/toolkit/docs/OpenCL_Best_Practices_Guide.pdf
74
75 =head1 BASIC WORKFLOW
76
77 To get something done, you basically have to do this once (refer to the
78 examples below for actual code, this is just a high-level description):
79
80 Find some platform (e.g. the first one) and some device(s) (e.g. the first
81 device of the platform), and create a context from those.
82
83 Create program objects from your OpenCL source code, then build (compile)
84 the programs for each device you want to run them on.
85
86 Create kernel objects for all kernels you want to use (surprisingly, these
87 are not device-specific).
88
89 Then, to execute stuff, you repeat these steps, possibly resuing or
90 sharing some buffers:
91
92 Create some input and output buffers from your context. Set these as
93 arguments to your kernel.
94
95 Enqueue buffer writes to initialise your input buffers (when not
96 initialised at creation time).
97
98 Enqueue the kernel execution.
99
100 Enqueue buffer reads for your output buffer to read results.
101
102 =head1 EXAMPLES
103
104 =head2 Enumerate all devices and get contexts for them.
105
106 Best run this once to get a feel for the platforms and devices in your
107 system.
108
109 for my $platform (OpenCL::platforms) {
110 printf "platform: %s\n", $platform->name;
111 printf "extensions: %s\n", $platform->extensions;
112 for my $device ($platform->devices) {
113 printf "+ device: %s\n", $device->name;
114 my $ctx = $platform->context (undef, [$device]);
115 # do stuff
116 }
117 }
118
119 =head2 Get a useful context and a command queue.
120
121 This is a useful boilerplate for any OpenCL program that only wants to use
122 one device,
123
124 my ($platform) = OpenCL::platforms; # find first platform
125 my ($dev) = $platform->devices; # find first device of platform
126 my $ctx = $platform->context (undef, [$dev]); # create context out of those
127 my $queue = $ctx->queue ($dev); # create a command queue for the device
128
129 =head2 Print all supported image formats of a context.
130
131 Best run this once for your context, to see whats available and how to
132 gather information.
133
134 for my $type (OpenCL::MEM_OBJECT_IMAGE2D, OpenCL::MEM_OBJECT_IMAGE3D) {
135 print "supported image formats for ", OpenCL::enum2str $type, "\n";
136
137 for my $f ($ctx->supported_image_formats (0, $type)) {
138 printf " %-10s %-20s\n", OpenCL::enum2str $f->[0], OpenCL::enum2str $f->[1];
139 }
140 }
141
142 =head2 Create a buffer with some predefined data, read it back synchronously,
143 then asynchronously.
144
145 my $buf = $ctx->buffer_sv (OpenCL::MEM_COPY_HOST_PTR, "helmut");
146
147 $queue->read_buffer ($buf, 1, 1, 3, my $data);
148 print "$data\n";
149
150 my $ev = $queue->read_buffer ($buf, 0, 1, 3, my $data);
151 $ev->wait;
152 print "$data\n"; # prints "elm"
153
154 =head2 Create and build a program, then create a kernel out of one of its
155 functions.
156
157 my $src = '
158 kernel void
159 squareit (global float *input, global float *output)
160 {
161 $id = get_global_id (0);
162 output [id] = input [id] * input [id];
163 }
164 ';
165
166 my $prog = $ctx->build_program ($src);
167 my $kernel = $prog->kernel ("squareit");
168
169 =head2 Create some input and output float buffers, then call the
170 'squareit' kernel on them.
171
172 my $input = $ctx->buffer_sv (OpenCL::MEM_COPY_HOST_PTR, pack "f*", 1, 2, 3, 4.5);
173 my $output = $ctx->buffer (0, OpenCL::SIZEOF_FLOAT * 5);
174
175 # set buffer
176 $kernel->set_buffer (0, $input);
177 $kernel->set_buffer (1, $output);
178
179 # execute it for all 4 numbers
180 $queue->nd_range_kernel ($kernel, undef, [4], undef);
181
182 # enqueue a synchronous read
183 $queue->read_buffer ($output, 1, 0, OpenCL::SIZEOF_FLOAT * 4, my $data);
184
185 # print the results:
186 printf "%s\n", join ", ", unpack "f*", $data;
187
188 =head2 The same enqueue operations as before, but assuming an out-of-order queue,
189 showing off barriers.
190
191 # execute it for all 4 numbers
192 $queue->nd_range_kernel ($kernel, undef, [4], undef);
193
194 # enqueue a barrier to ensure in-order execution
195 $queue->barrier;
196
197 # enqueue an async read
198 $queue->read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data);
199
200 # wait for all requests to finish
201 $queue->finish;
202
203 =head2 The same enqueue operations as before, but assuming an out-of-order queue,
204 showing off event objects and wait lists.
205
206 # execute it for all 4 numbers
207 my $ev = $queue->nd_range_kernel ($kernel, undef, [4], undef);
208
209 # enqueue an async read
210 $ev = $queue->read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data, $ev);
211
212 # wait for the last event to complete
213 $ev->wait;
214
215 =head2 Use the OpenGL module to share a texture between OpenCL and OpenGL and draw some julia
216 set flight effect.
217
218 This is quite a long example to get you going - you can also download it
219 from L<http://cvs.schmorp.de/OpenCL/examples/juliaflight>.
220
221 use OpenGL ":all";
222 use OpenCL;
223
224 my $S = $ARGV[0] || 256; # window/texture size, smaller is faster
225
226 # open a window and create a gl texture
227 OpenGL::glpOpenWindow width => $S, height => $S;
228 my $texid = glGenTextures_p 1;
229 glBindTexture GL_TEXTURE_2D, $texid;
230 glTexImage2D_c GL_TEXTURE_2D, 0, GL_RGBA8, $S, $S, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0;
231
232 # find and use the first opencl device that let's us get a shared opengl context
233 my $platform;
234 my $dev;
235 my $ctx;
236
237 sub get_context {
238 for (OpenCL::platforms) {
239 $platform = $_;
240 for ($platform->devices) {
241 $dev = $_;
242 $ctx = eval { $platform->context ([OpenCL::GLX_DISPLAY_KHR, undef, OpenCL::GL_CONTEXT_KHR, undef], [$dev]) }
243 and return;
244 }
245 }
246
247 die "cannot find suitable OpenCL device\n";
248 }
249
250 get_context;
251
252 my $queue = $ctx->queue ($dev);
253
254 # now attach an opencl image2d object to the opengl texture
255 my $tex = $ctx->gl_texture2d (OpenCL::MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, $texid);
256
257 # now the boring opencl code
258 my $src = <<EOF;
259 kernel void
260 juliatunnel (write_only image2d_t img, float time)
261 {
262 int2 xy = (int2)(get_global_id (0), get_global_id (1));
263 float2 p = convert_float2 (xy) / $S.f * 2.f - 1.f;
264
265 float2 m = (float2)(1.f, p.y) / fabs (p.x); // tunnel
266 m.x = fabs (fmod (m.x + time * 0.05f, 4.f) - 2.f);
267
268 float2 z = m;
269 float2 c = (float2)(sin (time * 0.01133f), cos (time * 0.02521f));
270
271 for (int i = 0; i < 25 && dot (z, z) < 4.f; ++i) // standard julia
272 z = (float2)(z.x * z.x - z.y * z.y, 2.f * z.x * z.y) + c;
273
274 float3 colour = (float3)(z.x, z.y, atan2 (z.y, z.x));
275 write_imagef (img, xy, (float4)(colour * p.x * p.x, 1.));
276 }
277 EOF
278
279 my $prog = $ctx->build_program ($src);
280 my $kernel = $prog->kernel ("juliatunnel");
281
282 # program compiled, kernel ready, now draw and loop
283
284 for (my $time; ; ++$time) {
285 # acquire objects from opengl
286 $queue->acquire_gl_objects ([$tex]);
287
288 # configure and run our kernel
289 $kernel->setf ("mf", $tex, $time*2); # mf = memory object, float
290 $queue->nd_range_kernel ($kernel, undef, [$S, $S], undef);
291
292 # release objects to opengl again
293 $queue->release_gl_objects ([$tex]);
294
295 # wait
296 $queue->finish;
297
298 # now draw the texture, the defaults should be all right
299 glTexParameterf GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST;
300
301 glEnable GL_TEXTURE_2D;
302 glBegin GL_QUADS;
303 glTexCoord2f 0, 1; glVertex3i -1, -1, -1;
304 glTexCoord2f 0, 0; glVertex3i 1, -1, -1;
305 glTexCoord2f 1, 0; glVertex3i 1, 1, -1;
306 glTexCoord2f 1, 1; glVertex3i -1, 1, -1;
307 glEnd;
308
309 glXSwapBuffers;
310
311 select undef, undef, undef, 1/60;
312 }
313
314 =head2 How to modify the previous example to not rely on GL sharing.
315
316 For those poor souls with only a sucky CPU OpenCL implementation, you
317 currently have to read the image into some perl scalar, and then modify a
318 texture or use glDrawPixels or so).
319
320 First, when you don't need gl sharing, you can create the context much simpler:
321
322 $ctx = $platform->context (undef, [$dev])
323
324 To use a texture, you would modify the above example by creating an
325 OpenCL::Image manually instead of deriving it from a texture:
326
327 my $tex = $ctx->image2d (OpenCL::MEM_WRITE_ONLY, OpenCL::RGBA, OpenCL::UNORM_INT8, $S, $S);
328
329 And in the draw loop, intead of acquire_gl_objects/release_gl_objects, you
330 would read the image2d after the kernel has written it:
331
332 $queue->read_image ($tex, 0, 0, 0, 0, $S, $S, 1, 0, 0, my $data);
333
334 And then you would upload the pixel data to the texture (or use glDrawPixels):
335
336 glTexSubImage2D_s GL_TEXTURE_2D, 0, 0, 0, $S, $S, GL_RGBA, GL_UNSIGNED_BYTE, $data;
337
338 The fully modified example can be found at
339 L<http://cvs.schmorp.de/OpenCL/examples/juliaflight-nosharing>.
340
341 =head2 Julia sets look soooo 80ies.
342
343 Then colour them differently, e.g. using orbit traps! Replace the loop and
344 colour calculation from the previous examples by this:
345
346 float2 dm = (float2)(1.f, 1.f);
347
348 for (int i = 0; i < 25; ++i)
349 {
350 z = (float2)(z.x * z.x - z.y * z.y, 2.f * z.x * z.y) + c;
351 dm = fmin (dm, (float2)(fabs (dot (z, z) - 1.f), fabs (z.x - 1.f)));
352 }
353
354 float3 colour = (float3)(dm.x * dm.y, dm.x * dm.y, dm.x);
355
356 Also try C<-10.f> instead of C<-1.f>.
357
358 =head1 DOCUMENTATION
359
360 =head2 BASIC CONVENTIONS
361
362 This is not a one-to-one C-style translation of OpenCL to Perl - instead
363 I attempted to make the interface as type-safe as possible by introducing
364 object syntax where it makes sense. There are a number of important
365 differences between the OpenCL C API and this module:
366
367 =over 4
368
369 =item * Object lifetime managament is automatic - there is no need
370 to free objects explicitly (C<clReleaseXXX>), the release function
371 is called automatically once all Perl references to it go away.
372
373 =item * OpenCL uses CamelCase for function names
374 (e.g. C<clGetPlatformIDs>, C<clGetPlatformInfo>), while this module
375 uses underscores as word separator and often leaves out prefixes
376 (C<OpenCL::platforms>, C<< $platform->info >>).
377
378 =item * OpenCL often specifies fixed vector function arguments as short
379 arrays (C<size_t origin[3]>), while this module explicitly expects the
380 components as separate arguments (C<$orig_x, $orig_y, $orig_z>) in
381 function calls.
382
383 =item * Structures are often specified by flattening out their components
384 as with short vectors, and returned as arrayrefs.
385
386 =item * When enqueuing commands, the wait list is specified by adding
387 extra arguments to the function - anywhere a C<$wait_events...> argument
388 is documented this can be any number of event objects. As an extsnion
389 implemented by this module, C<undef> values will be ignored in the event
390 list.
391
392 =item * When enqueuing commands, if the enqueue method is called in void
393 context, no event is created. In all other contexts an event is returned
394 by the method.
395
396 =item * This module expects all functions to return C<OpenCL::SUCCESS>. If any
397 other status is returned the function will throw an exception, so you
398 don't normally have to to any error checking.
399
400 =back
401
402 =head2 CONSTANTS
403
404 All C<CL_xxx> constants that this module supports are always available
405 in the C<OpenCL> namespace as C<OpenCL::xxx> (i.e. without the C<CL_>
406 prefix). Constants which are not defined in the header files used during
407 compilation, or otherwise are not available, will have the value C<0> (in
408 some cases, this will make them indistinguishable from real constants,
409 sorry).
410
411 The latest version of this module knows and exports the constants
412 listed in L<http://cvs.schmorp.de/OpenCL/constiv.h>.
413
414 =head2 OPENCL 1.1 VS. OPENCL 1.2
415
416 This module supports both OpenCL version 1.1 and 1.2, although the OpenCL
417 1.2 interface hasn't been tested much for lack of availability of an
418 actual implementation.
419
420 Every function or method in this manual page that interfaces to a
421 particular OpenCL function has a link to the its C manual page.
422
423 If the link contains a F<1.1>, then this function is an OpenCL 1.1
424 function. Most but not all also exist in OpenCL 1.2, and this module
425 tries to emulate the missing ones for you, when told to do so at
426 compiletime. You can check whether a function was removed in OpenCL 1.2 by
427 replacing the F<1.1> component in the URL by F<1.2>.
428
429 If the link contains a F<1.2>, then this is a OpenCL 1.2-only
430 function. Even if the module was compiled with OpenCL 1.2 header files
431 and has an 1.2 OpenCL library, calling such a function on a platform that
432 doesn't implement 1.2 causes undefined behaviour, usually a crash (But
433 this is not guaranteed).
434
435 You can find out whether this module was compiled to prefer 1.1
436 functionality by ooking at C<OpenCL::PREFER_1_1> - if it is true, then
437 1.1 functions generally are implemented using 1.1 OpenCL functions. If it
438 is false, then 1.1 functions missing from 1.2 are emulated by calling 1.2
439 fucntions.
440
441 This is a somewhat sorry state of affairs, but the Khronos group choose to
442 make every release of OpenCL source and binary incompatible with previous
443 releases.
444
445 =head2 PERL AND OPENCL TYPES
446
447 This handy(?) table lists OpenCL types and their perl, PDL and pack/unpack
448 format equivalents:
449
450 OpenCL perl PDL pack/unpack
451 char IV - c
452 uchar IV byte C
453 short IV short s
454 ushort IV ushort S
455 int IV long? l
456 uint IV - L
457 long IV longlong q
458 ulong IV - Q
459 float NV float f
460 half IV ushort S
461 double NV double d
462
463 =head2 GLX SUPPORT
464
465 Due to the sad state that OpenGL support is in in Perl (mostly the OpenGL
466 module, which has little to no documentation and has little to no support
467 for glX), this module, as a special extension, treats context creation
468 properties C<OpenCL::GLX_DISPLAY_KHR> and C<OpenCL::GL_CONTEXT_KHR>
469 specially: If either or both of these are C<undef>, then the OpenCL
470 module tries to dynamically resolve C<glXGetCurrentDisplay> and
471 C<glXGetCurrentContext>, call these functions and use their return values
472 instead.
473
474 For this to work, the OpenGL library must be loaded, a GLX context must
475 have been created and be made current, and C<dlsym> must be available and
476 capable of finding the function via C<RTLD_DEFAULT>.
477
478 =head2 EVENT SYSTEM
479
480 OpenCL can generate a number of (potentially) asynchronous events, for
481 example, after compiling a program, to signal a context-related error or,
482 perhaps most important, to signal completion of queued jobs (by setting
483 callbacks on OpenCL::Event objects).
484
485 The OpenCL module converts all these callbacks into events - you can
486 still register callbacks, but they are not executed when your OpenCL
487 implementation calls the actual callback, but only later. Therefore, none
488 of the limitations of OpenCL callbacks apply to the perl implementation:
489 it is perfectly safe to make blocking operations from event callbacks, and
490 enqueued operations don't need to be flushed.
491
492 To facilitate this, this module maintains an event queue - each
493 time an asynchronous event happens, it is queued, and perl will be
494 interrupted. This is implemented via the L<Async::Interrupt> module. In
495 addition, this module has L<AnyEvent> support, so it can seamlessly
496 integrate itself into many event loops.
497
498 Since L<Async::Interrupt> is a bit hard to understand, here are some case examples:
499
500 =head3 Don't use callbacks.
501
502 When your program never uses any callbacks, then there will never be any
503 notifications you need to take care of, and therefore no need to worry
504 about all this.
505
506 You can achieve a great deal by explicitly waiting for events, or using
507 barriers and flush calls. In many programs, there is no need at all to
508 tinker with asynchronous events.
509
510 =head3 Use AnyEvent
511
512 This module automatically registers a watcher that invokes all outstanding
513 event callbacks when AnyEvent is initialised (and block asynchronous
514 interruptions). Using this mode of operations is the safest and most
515 recommended one.
516
517 To use this, simply use AnyEvent and this module normally, make sure you
518 have an event loop running:
519
520 use Gtk2 -init;
521 use AnyEvent;
522
523 # initialise AnyEvent, by creating a watcher, or:
524 AnyEvent::detect;
525
526 my $e = $queue->marker;
527 $e->cb (sub {
528 warn "opencl is finished\n";
529 })
530
531 main Gtk2;
532
533 Note that this module will not initialise AnyEvent for you. Before
534 AnyEvent is initialised, the module will asynchronously interrupt perl
535 instead. To avoid any surprises, it's best to explicitly initialise
536 AnyEvent.
537
538 You can temporarily enable asynchronous interruptions (see next paragraph)
539 by calling C<$OpenCL::INTERRUPT->unblock> and disable them again by
540 calling C<$OpenCL::INTERRUPT->block>.
541
542 =head3 Let yourself be interrupted at any time
543
544 This mode is the default unless AnyEvent is loaded and initialised. In
545 this mode, OpenCL asynchronously interrupts a running perl program. The
546 emphasis is on both I<asynchronously> and I<running> here.
547
548 Asynchronously means that perl might execute your callbacks at any
549 time. For example, in the following code (I<THAT YOU SHOULD NOT COPY>),
550 the C<until> loop following the marker call will be interrupted by the
551 callback:
552
553 my $e = $queue->marker;
554 my $flag;
555 $e->cb (sub { $flag = 1 });
556 1 until $flag;
557 # $flag is now 1
558
559 The reason why you shouldn't blindly copy the above code is that
560 busy waiting is a really really bad thing, and really really bad for
561 performance.
562
563 While at first this asynchronous business might look exciting, it can be
564 really hard, because you need to be prepared for the callback code to be
565 executed at any time, which limits the amount of things the callback code
566 can do safely.
567
568 This can be mitigated somewhat by using C<<
569 $OpenCL::INTERRUPT->scope_block >> (see the L<Async::Interrupt>
570 documentation for details).
571
572 The other problem is that your program must be actively I<running> to be
573 interrupted. When you calculate stuff, your program is running. When you
574 hang in some C functions or other block execution (by calling C<sleep>,
575 C<select>, running an event loop and so on), your program is waiting, not
576 running.
577
578 One way around that would be to attach a read watcher to your event loop,
579 listening for events on C<< $OpenCL::INTERRUPT->pipe_fileno >>, using a
580 dummy callback (C<sub { }>) to temporarily execute some perl code.
581
582 That is then awfully close to using the built-in AnyEvent support above,
583 though, so consider that one instead.
584
585 =head3 Be creative
586
587 OpenCL exports the L<Async::Interrupt> object it uses in the global
588 variable C<$OpenCL::INTERRUPT>. You can configure it in any way you like.
589
590 So if you want to feel like a real pro, err, wait, if you feel no risk
591 menas no fun, you can experiment by implementing your own mode of
592 operations.
593
594 =cut
595
596 package OpenCL;
597
598 use common::sense;
599 use Carp ();
600 use Async::Interrupt ();
601
602 our $POLL_FUNC; # set by XS
603
604 BEGIN {
605 our $VERSION = '1.01';
606
607 require XSLoader;
608 XSLoader::load (__PACKAGE__, $VERSION);
609
610 @OpenCL::Platform::ISA =
611 @OpenCL::Device::ISA =
612 @OpenCL::Context::ISA =
613 @OpenCL::Queue::ISA =
614 @OpenCL::Memory::ISA =
615 @OpenCL::Sampler::ISA =
616 @OpenCL::Program::ISA =
617 @OpenCL::Kernel::ISA =
618 @OpenCL::Event::ISA = OpenCL::Object::;
619
620 @OpenCL::SubDevice::ISA = OpenCL::Device::;
621
622 @OpenCL::Buffer::ISA =
623 @OpenCL::Image::ISA = OpenCL::Memory::;
624
625 @OpenCL::BufferObj::ISA = OpenCL::Buffer::;
626
627 @OpenCL::Image2D::ISA =
628 @OpenCL::Image3D::ISA =
629 @OpenCL::Image2DArray::ISA =
630 @OpenCL::Image1D::ISA =
631 @OpenCL::Image1DArray::ISA =
632 @OpenCL::Image1DBuffer::ISA = OpenCL::Image::;
633
634 @OpenCL::UserEvent::ISA = OpenCL::Event::;
635
636 @OpenCL::MappedBuffer::ISA =
637 @OpenCL::MappedImage::ISA = OpenCL::Mapped::;
638 }
639
640 =head2 THE OpenCL PACKAGE
641
642 =over 4
643
644 =item $int = OpenCL::errno
645
646 The last error returned by a function - it's only valid after an error occured
647 and before calling another OpenCL function.
648
649 =item $str = OpenCL::err2str [$errval]
650
651 Converts an error value into a human readable string. If no error value is
652 given, then the last error will be used (as returned by OpenCL::errno).
653
654 The latest version of this module knows the error constants
655 listed in L<http://cvs.schmorp.de/OpenCL/errstr.h>.
656
657 =item $str = OpenCL::enum2str $enum
658
659 Converts most enum values (of parameter names, image format constants,
660 object types, addressing and filter modes, command types etc.) into a
661 human readable string. When confronted with some random integer it can be
662 very helpful to pass it through this function to maybe get some readable
663 string out of it.
664
665 The latest version of this module knows the enumaration constants
666 listed in L<http://cvs.schmorp.de/OpenCL/enumstr.h>.
667
668 =item @platforms = OpenCL::platforms
669
670 Returns all available OpenCL::Platform objects.
671
672 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformIDs.html>
673
674 =item $ctx = OpenCL::context_from_type $properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $callback->($err, $pvt) = $print_stderr
675
676 Tries to create a context from a default device and platform type - never worked for me.
677 Consider using C<< $platform->context_from_type >> instead.
678
679 type: OpenCL::DEVICE_TYPE_DEFAULT, OpenCL::DEVICE_TYPE_CPU, OpenCL::DEVICE_TYPE_GPU,
680 OpenCL::DEVICE_TYPE_ACCELERATOR, OpenCL::DEVICE_TYPE_CUSTOM, OpenCL::DEVICE_TYPE_ALL.
681
682 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html>
683
684 =item $ctx = OpenCL::context $properties, \@devices, $callback->($err, $pvt) = $print_stderr)
685
686 Create a new OpenCL::Context object using the given device object(s).
687 Consider using C<< $platform->context >> instead.
688
689 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContext.html>
690
691 =item OpenCL::wait_for_events $wait_events...
692
693 Waits for all events to complete.
694
695 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html>
696
697 =item OpenCL::poll
698
699 Checks if there are any outstanding events (see L<EVENT SYSTEM>) and
700 invokes their callbacks.
701
702 =item $OpenCL::INTERRUPT
703
704 The L<Async::Interrupt> object used to signal asynchronous events (see
705 L<EVENT SYSTEM>).
706
707 =cut
708
709 our $INTERRUPT = new Async::Interrupt c_cb => [$POLL_FUNC, 0];
710
711 &_eq_initialise ($INTERRUPT->signal_func);
712
713 =item $OpenCL::WATCHER
714
715 The L<AnyEvent> watcher object used to watch for asynchronous events (see
716 L<EVENT SYSTEM>). This variable is C<undef> until L<AnyEvent> has been
717 loaded I<and> initialised (e.g. by calling C<AnyEvent::detect>).
718
719 =cut
720
721 our $WATCHER;
722
723 sub _init_anyevent {
724 $INTERRUPT->block;
725 $WATCHER = AE::io ($INTERRUPT->pipe_fileno, 0, sub { $INTERRUPT->handle });
726 }
727
728 if (defined $AnyEvent::MODEL) {
729 _init_anyevent;
730 } else {
731 push @AnyEvent::post_detect, \&_init_anyevent;
732 }
733
734 =back
735
736 =head2 THE OpenCL::Object CLASS
737
738 This is the base class for all objects in the OpenCL module. The only
739 method it implements is the C<id> method, which is only useful if you want
740 to interface to OpenCL on the C level.
741
742 =over 4
743
744 =item $iv = $obj->id
745
746 OpenCL objects are represented by pointers or integers on the C level. If
747 you want to interface to an OpenCL object directly on the C level, then
748 you need this value, which is returned by this method. You should use an
749 C<IV> type in your code and cast that to the correct type.
750
751 =cut
752
753 sub OpenCL::Object::id {
754 ref $_[0] eq "SCALAR"
755 ? ${ $_[0] }
756 : $_[0][0]
757 }
758
759 =back
760
761 =head2 THE OpenCL::Platform CLASS
762
763 =over 4
764
765 =item @devices = $platform->devices ($type = OpenCL::DEVICE_TYPE_ALL)
766
767 Returns a list of matching OpenCL::Device objects.
768
769 =item $ctx = $platform->context_from_type ($properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $callback->($err, $pvt) = $print_stderr)
770
771 Tries to create a context. Never worked for me, and you need devices explicitly anyway.
772
773 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html>
774
775 =item $ctx = $platform->context ($properties, \@devices, $callback->($err, $pvt) = $print_stderr)
776
777 Create a new OpenCL::Context object using the given device object(s)- a
778 OpenCL::CONTEXT_PLATFORM property is supplied automatically.
779
780 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContext.html>
781
782 =item $packed_value = $platform->info ($name)
783
784 Calls C<clGetPlatformInfo> and returns the packed, raw value - for
785 strings, this will be the string (possibly including terminating \0), for
786 other values you probably need to use the correct C<unpack>.
787
788 It's best to avoid this method and use one of the following convenience
789 wrappers.
790
791 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformInfo.html>
792
793 =item $platform->unload_compiler
794
795 Attempts to unload the compiler for this platform, for endless
796 profit. Does nothing on OpenCL 1.1.
797
798 L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clUnloadPlatformCompiler.html>
799
800 =for gengetinfo begin platform
801
802 =item $string = $platform->profile
803
804 Calls C<clGetPlatformInfo> with C<OpenCL::PLATFORM_PROFILE> and returns the result.
805
806 =item $string = $platform->version
807
808 Calls C<clGetPlatformInfo> with C<OpenCL::PLATFORM_VERSION> and returns the result.
809
810 =item $string = $platform->name
811
812 Calls C<clGetPlatformInfo> with C<OpenCL::PLATFORM_NAME> and returns the result.
813
814 =item $string = $platform->vendor
815
816 Calls C<clGetPlatformInfo> with C<OpenCL::PLATFORM_VENDOR> and returns the result.
817
818 =item $string = $platform->extensions
819
820 Calls C<clGetPlatformInfo> with C<OpenCL::PLATFORM_EXTENSIONS> and returns the result.
821
822 =for gengetinfo end platform
823
824 =back
825
826 =head2 THE OpenCL::Device CLASS
827
828 =over 4
829
830 =item $packed_value = $device->info ($name)
831
832 See C<< $platform->info >> for details.
833
834 type: OpenCL::DEVICE_TYPE_DEFAULT, OpenCL::DEVICE_TYPE_CPU,
835 OpenCL::DEVICE_TYPE_GPU, OpenCL::DEVICE_TYPE_ACCELERATOR,
836 OpenCL::DEVICE_TYPE_CUSTOM, OpenCL::DEVICE_TYPE_ALL.
837
838 fp_config: OpenCL::FP_DENORM, OpenCL::FP_INF_NAN, OpenCL::FP_ROUND_TO_NEAREST,
839 OpenCL::FP_ROUND_TO_ZERO, OpenCL::FP_ROUND_TO_INF, OpenCL::FP_FMA,
840 OpenCL::FP_SOFT_FLOAT, OpenCL::FP_CORRECTLY_ROUNDED_DIVIDE_SQRT.
841
842 mem_cache_type: OpenCL::NONE, OpenCL::READ_ONLY_CACHE, OpenCL::READ_WRITE_CACHE.
843
844 local_mem_type: OpenCL::LOCAL, OpenCL::GLOBAL.
845
846 exec_capabilities: OpenCL::EXEC_KERNEL, OpenCL::EXEC_NATIVE_KERNEL.
847
848 command_queue_properties: OpenCL::QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,
849 OpenCL::QUEUE_PROFILING_ENABLE.
850
851 partition_properties: OpenCL::DEVICE_PARTITION_EQUALLY,
852 OpenCL::DEVICE_PARTITION_BY_COUNTS, OpenCL::DEVICE_PARTITION_BY_COUNTS_LIST_END,
853 OpenCL::DEVICE_PARTITION_BY_AFFINITY_DOMAIN.
854
855 affinity_domain: OpenCL::DEVICE_AFFINITY_DOMAIN_NUMA,
856 OpenCL::DEVICE_AFFINITY_DOMAIN_L4_CACHE, OpenCL::DEVICE_AFFINITY_DOMAIN_L3_CACHE,
857 OpenCL::DEVICE_AFFINITY_DOMAIN_L2_CACHE, OpenCL::DEVICE_AFFINITY_DOMAIN_L1_CACHE,
858 OpenCL::DEVICE_AFFINITY_DOMAIN_NEXT_PARTITIONABLE.
859
860 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetDeviceInfo.html>
861
862 =item @devices = $device->sub_devices (\@properties)
863
864 Creates OpencL::SubDevice objects by partitioning an existing device.
865
866 L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateSubDevices.html>
867
868 =for gengetinfo begin device
869
870 =item $device_type = $device->type
871
872 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_TYPE> and returns the result.
873
874 =item $uint = $device->vendor_id
875
876 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_VENDOR_ID> and returns the result.
877
878 =item $uint = $device->max_compute_units
879
880 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_MAX_COMPUTE_UNITS> and returns the result.
881
882 =item $uint = $device->max_work_item_dimensions
883
884 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_MAX_WORK_ITEM_DIMENSIONS> and returns the result.
885
886 =item $int = $device->max_work_group_size
887
888 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_MAX_WORK_GROUP_SIZE> and returns the result.
889
890 =item @ints = $device->max_work_item_sizes
891
892 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_MAX_WORK_ITEM_SIZES> and returns the result.
893
894 =item $uint = $device->preferred_vector_width_char
895
896 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_PREFERRED_VECTOR_WIDTH_CHAR> and returns the result.
897
898 =item $uint = $device->preferred_vector_width_short
899
900 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_PREFERRED_VECTOR_WIDTH_SHORT> and returns the result.
901
902 =item $uint = $device->preferred_vector_width_int
903
904 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_PREFERRED_VECTOR_WIDTH_INT> and returns the result.
905
906 =item $uint = $device->preferred_vector_width_long
907
908 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_PREFERRED_VECTOR_WIDTH_LONG> and returns the result.
909
910 =item $uint = $device->preferred_vector_width_float
911
912 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT> and returns the result.
913
914 =item $uint = $device->preferred_vector_width_double
915
916 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE> and returns the result.
917
918 =item $uint = $device->max_clock_frequency
919
920 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_MAX_CLOCK_FREQUENCY> and returns the result.
921
922 =item $bitfield = $device->address_bits
923
924 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_ADDRESS_BITS> and returns the result.
925
926 =item $uint = $device->max_read_image_args
927
928 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_MAX_READ_IMAGE_ARGS> and returns the result.
929
930 =item $uint = $device->max_write_image_args
931
932 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_MAX_WRITE_IMAGE_ARGS> and returns the result.
933
934 =item $ulong = $device->max_mem_alloc_size
935
936 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_MAX_MEM_ALLOC_SIZE> and returns the result.
937
938 =item $int = $device->image2d_max_width
939
940 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_IMAGE2D_MAX_WIDTH> and returns the result.
941
942 =item $int = $device->image2d_max_height
943
944 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_IMAGE2D_MAX_HEIGHT> and returns the result.
945
946 =item $int = $device->image3d_max_width
947
948 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_IMAGE3D_MAX_WIDTH> and returns the result.
949
950 =item $int = $device->image3d_max_height
951
952 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_IMAGE3D_MAX_HEIGHT> and returns the result.
953
954 =item $int = $device->image3d_max_depth
955
956 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_IMAGE3D_MAX_DEPTH> and returns the result.
957
958 =item $uint = $device->image_support
959
960 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_IMAGE_SUPPORT> and returns the result.
961
962 =item $int = $device->max_parameter_size
963
964 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_MAX_PARAMETER_SIZE> and returns the result.
965
966 =item $uint = $device->max_samplers
967
968 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_MAX_SAMPLERS> and returns the result.
969
970 =item $uint = $device->mem_base_addr_align
971
972 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_MEM_BASE_ADDR_ALIGN> and returns the result.
973
974 =item $uint = $device->min_data_type_align_size
975
976 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_MIN_DATA_TYPE_ALIGN_SIZE> and returns the result.
977
978 =item $device_fp_config = $device->single_fp_config
979
980 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_SINGLE_FP_CONFIG> and returns the result.
981
982 =item $device_mem_cache_type = $device->global_mem_cache_type
983
984 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_GLOBAL_MEM_CACHE_TYPE> and returns the result.
985
986 =item $uint = $device->global_mem_cacheline_size
987
988 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_GLOBAL_MEM_CACHELINE_SIZE> and returns the result.
989
990 =item $ulong = $device->global_mem_cache_size
991
992 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_GLOBAL_MEM_CACHE_SIZE> and returns the result.
993
994 =item $ulong = $device->global_mem_size
995
996 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_GLOBAL_MEM_SIZE> and returns the result.
997
998 =item $ulong = $device->max_constant_buffer_size
999
1000 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_MAX_CONSTANT_BUFFER_SIZE> and returns the result.
1001
1002 =item $uint = $device->max_constant_args
1003
1004 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_MAX_CONSTANT_ARGS> and returns the result.
1005
1006 =item $device_local_mem_type = $device->local_mem_type
1007
1008 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_LOCAL_MEM_TYPE> and returns the result.
1009
1010 =item $ulong = $device->local_mem_size
1011
1012 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_LOCAL_MEM_SIZE> and returns the result.
1013
1014 =item $boolean = $device->error_correction_support
1015
1016 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_ERROR_CORRECTION_SUPPORT> and returns the result.
1017
1018 =item $int = $device->profiling_timer_resolution
1019
1020 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_PROFILING_TIMER_RESOLUTION> and returns the result.
1021
1022 =item $boolean = $device->endian_little
1023
1024 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_ENDIAN_LITTLE> and returns the result.
1025
1026 =item $boolean = $device->available
1027
1028 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_AVAILABLE> and returns the result.
1029
1030 =item $boolean = $device->compiler_available
1031
1032 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_COMPILER_AVAILABLE> and returns the result.
1033
1034 =item $device_exec_capabilities = $device->execution_capabilities
1035
1036 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_EXECUTION_CAPABILITIES> and returns the result.
1037
1038 =item $command_queue_properties = $device->properties
1039
1040 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_QUEUE_PROPERTIES> and returns the result.
1041
1042 =item $ = $device->platform
1043
1044 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_PLATFORM> and returns the result.
1045
1046 =item $string = $device->name
1047
1048 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_NAME> and returns the result.
1049
1050 =item $string = $device->vendor
1051
1052 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_VENDOR> and returns the result.
1053
1054 =item $string = $device->driver_version
1055
1056 Calls C<clGetDeviceInfo> with C<OpenCL::DRIVER_VERSION> and returns the result.
1057
1058 =item $string = $device->profile
1059
1060 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_PROFILE> and returns the result.
1061
1062 =item $string = $device->version
1063
1064 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_VERSION> and returns the result.
1065
1066 =item $string = $device->extensions
1067
1068 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_EXTENSIONS> and returns the result.
1069
1070 =item $uint = $device->preferred_vector_width_half
1071
1072 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_PREFERRED_VECTOR_WIDTH_HALF> and returns the result.
1073
1074 =item $uint = $device->native_vector_width_char
1075
1076 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_NATIVE_VECTOR_WIDTH_CHAR> and returns the result.
1077
1078 =item $uint = $device->native_vector_width_short
1079
1080 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_NATIVE_VECTOR_WIDTH_SHORT> and returns the result.
1081
1082 =item $uint = $device->native_vector_width_int
1083
1084 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_NATIVE_VECTOR_WIDTH_INT> and returns the result.
1085
1086 =item $uint = $device->native_vector_width_long
1087
1088 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_NATIVE_VECTOR_WIDTH_LONG> and returns the result.
1089
1090 =item $uint = $device->native_vector_width_float
1091
1092 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_NATIVE_VECTOR_WIDTH_FLOAT> and returns the result.
1093
1094 =item $uint = $device->native_vector_width_double
1095
1096 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE> and returns the result.
1097
1098 =item $uint = $device->native_vector_width_half
1099
1100 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_NATIVE_VECTOR_WIDTH_HALF> and returns the result.
1101
1102 =item $device_fp_config = $device->double_fp_config
1103
1104 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_DOUBLE_FP_CONFIG> and returns the result.
1105
1106 =item $device_fp_config = $device->half_fp_config
1107
1108 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_HALF_FP_CONFIG> and returns the result.
1109
1110 =item $boolean = $device->host_unified_memory
1111
1112 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_HOST_UNIFIED_MEMORY> and returns the result.
1113
1114 =item $device = $device->parent_device_ext
1115
1116 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_PARENT_DEVICE_EXT> and returns the result.
1117
1118 =item @device_partition_property_exts = $device->partition_types_ext
1119
1120 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_PARTITION_TYPES_EXT> and returns the result.
1121
1122 =item @device_partition_property_exts = $device->affinity_domains_ext
1123
1124 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_AFFINITY_DOMAINS_EXT> and returns the result.
1125
1126 =item $uint = $device->reference_count_ext
1127
1128 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_REFERENCE_COUNT_EXT> and returns the result.
1129
1130 =item @device_partition_property_exts = $device->partition_style_ext
1131
1132 Calls C<clGetDeviceInfo> with C<OpenCL::DEVICE_PARTITION_STYLE_EXT> and returns the result.
1133
1134 =for gengetinfo end device
1135
1136 =back
1137
1138 =head2 THE OpenCL::Context CLASS
1139
1140 An OpenCL::Context is basically a container, or manager, for a number of
1141 devices of a platform. It is used to create all sorts of secondary objects
1142 such as buffers, queues, programs and so on.
1143
1144 All context creation functions and methods take a list of properties
1145 (type-value pairs). All property values can be specified as integers -
1146 some additionally support other types:
1147
1148 =over 4
1149
1150 =item OpenCL::CONTEXT_PLATFORM
1151
1152 Also accepts OpenCL::Platform objects.
1153
1154 =item OpenCL::GLX_DISPLAY_KHR
1155
1156 Also accepts C<undef>, in which case a deep and troubling hack is engaged
1157 to find the current glx display (see L<GLX SUPPORT>).
1158
1159 =item OpenCL::GL_CONTEXT_KHR
1160
1161 Also accepts C<undef>, in which case a deep and troubling hack is engaged
1162 to find the current glx context (see L<GLX SUPPORT>).
1163
1164 =back
1165
1166 =over 4
1167
1168 =item $prog = $ctx->build_program ($program, $options = "")
1169
1170 This convenience function tries to build the program on all devices in
1171 the context. If the build fails, then the function will C<croak> with the
1172 build log. Otherwise ti returns the program object.
1173
1174 The C<$program> can either be a C<OpenCL::Program> object or a string
1175 containing the program. In the latter case, a program objetc will be
1176 created automatically.
1177
1178 =cut
1179
1180 sub OpenCL::Context::build_program {
1181 my ($self, $prog, $options) = @_;
1182
1183 $prog = $self->program_with_source ($prog)
1184 unless ref $prog;
1185
1186 eval { $prog->build (undef, $options); 1 }
1187 or errno == BUILD_PROGRAM_FAILURE
1188 or errno == INVALID_BINARY # workaround nvidia bug
1189 or Carp::croak "OpenCL::Context->build_program: " . err2str;
1190
1191 # we check status for all devices
1192 for my $dev ($self->devices) {
1193 $prog->build_status ($dev) == BUILD_SUCCESS
1194 or Carp::croak "Building OpenCL program for device '" . $dev->name . "' failed:\n"
1195 . $prog->build_log ($dev);
1196 }
1197
1198 $prog
1199 }
1200
1201 =item $queue = $ctx->queue ($device, $properties)
1202
1203 Create a new OpenCL::Queue object from the context and the given device.
1204
1205 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateCommandQueue.html>
1206
1207 Example: create an out-of-order queue.
1208
1209 $queue = $ctx->queue ($device, OpenCL::QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE);
1210
1211 =item $ev = $ctx->user_event
1212
1213 Creates a new OpenCL::UserEvent object.
1214
1215 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateUserEvent.html>
1216
1217 =item $buf = $ctx->buffer ($flags, $len)
1218
1219 Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object with the
1220 given flags and octet-size.
1221
1222 flags: OpenCL::MEM_READ_WRITE, OpenCL::MEM_WRITE_ONLY, OpenCL::MEM_READ_ONLY,
1223 OpenCL::MEM_USE_HOST_PTR, OpenCL::MEM_ALLOC_HOST_PTR, OpenCL::MEM_COPY_HOST_PTR,
1224 OpenCL::MEM_HOST_WRITE_ONLY, OpenCL::MEM_HOST_READ_ONLY, OpenCL::MEM_HOST_NO_ACCESS.
1225
1226 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateBuffer.html>
1227
1228 A somewhat informative thread on the flags is:
1229
1230 L<http://www.khronos.org/message_boards/viewtopic.php?f=28&t=2440>
1231
1232 =item $buf = $ctx->buffer_sv ($flags, $data)
1233
1234 Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object and
1235 initialise it with the given data values.
1236
1237 =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)
1238
1239 Creates a new OpenCL::Image object and optionally initialises it with
1240 the given data values.
1241
1242 channel_order: OpenCL::R, OpenCL::A, OpenCL::RG, OpenCL::RA, OpenCL::RGB,
1243 OpenCL::RGBA, OpenCL::BGRA, OpenCL::ARGB, OpenCL::INTENSITY, OpenCL::LUMINANCE,
1244 OpenCL::Rx, OpenCL::RGx, OpenCL::RGBx.
1245
1246 channel_type: OpenCL::SNORM_INT8, OpenCL::SNORM_INT16, OpenCL::UNORM_INT8,
1247 OpenCL::UNORM_INT16, OpenCL::UNORM_SHORT_565, OpenCL::UNORM_SHORT_555,
1248 OpenCL::UNORM_INT_101010, OpenCL::SIGNED_INT8, OpenCL::SIGNED_INT16,
1249 OpenCL::SIGNED_INT32, OpenCL::UNSIGNED_INT8, OpenCL::UNSIGNED_INT16,
1250 OpenCL::UNSIGNED_INT32, OpenCL::HALF_FLOAT, OpenCL::FLOAT.
1251
1252 type: OpenCL::MEM_OBJECT_BUFFER, OpenCL::MEM_OBJECT_IMAGE2D,
1253 OpenCL::MEM_OBJECT_IMAGE3D, OpenCL::MEM_OBJECT_IMAGE2D_ARRAY,
1254 OpenCL::MEM_OBJECT_IMAGE1D, OpenCL::MEM_OBJECT_IMAGE1D_ARRAY,
1255 OpenCL::MEM_OBJECT_IMAGE1D_BUFFER.
1256
1257 L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateImage.html>
1258
1259 =item $img = $ctx->image2d ($flags, $channel_order, $channel_type, $width, $height, $row_pitch = 0, $data = undef)
1260
1261 Creates a new OpenCL::Image2D object and optionally initialises it with
1262 the given data values.
1263
1264 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateImage2D.html>
1265
1266 =item $img = $ctx->image3d ($flags, $channel_order, $channel_type, $width, $height, $depth, $row_pitch = 0, $slice_pitch = 0, $data = undef)
1267
1268 Creates a new OpenCL::Image3D object and optionally initialises it with
1269 the given data values.
1270
1271 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateImage3D.html>
1272
1273 =item $buffer = $ctx->gl_buffer ($flags, $bufobj)
1274
1275 Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object that refers to the given
1276 OpenGL buffer object.
1277
1278 flags: OpenCL::MEM_READ_WRITE, OpenCL::MEM_READ_ONLY, OpenCL::MEM_WRITE_ONLY.
1279
1280 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLBuffer.html
1281
1282 =item $img = $ctx->gl_texture ($flags, $target, $miplevel, $texture)
1283
1284 Creates a new OpenCL::Image object that refers to the given OpenGL
1285 texture object or buffer.
1286
1287 target: GL_TEXTURE_1D, GL_TEXTURE_1D_ARRAY, GL_TEXTURE_BUFFER,
1288 GL_TEXTURE_2D, GL_TEXTURE_2D_ARRAY, GL_TEXTURE_3D,
1289 GL_TEXTURE_CUBE_MAP_POSITIVE_X, GL_TEXTURE_CUBE_MAP_POSITIVE_Y,
1290 GL_TEXTURE_CUBE_MAP_POSITIVE_Z, GL_TEXTURE_CUBE_MAP_NEGATIVE_X,
1291 GL_TEXTURE_CUBE_MAP_NEGATIVE_Y, GL_TEXTURE_CUBE_MAP_NEGATIVE_Z,
1292 GL_TEXTURE_RECTANGLE/GL_TEXTURE_RECTANGLE_ARB.
1293
1294 http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateFromGLTexture.html
1295
1296 =item $img = $ctx->gl_texture2d ($flags, $target, $miplevel, $texture)
1297
1298 Creates a new OpenCL::Image2D object that refers to the given OpenGL
1299 2D texture object.
1300
1301 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture2D.html
1302
1303 =item $img = $ctx->gl_texture3d ($flags, $target, $miplevel, $texture)
1304
1305 Creates a new OpenCL::Image3D object that refers to the given OpenGL
1306 3D texture object.
1307
1308 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture3D.html
1309
1310 =item $ctx->gl_renderbuffer ($flags, $renderbuffer)
1311
1312 Creates a new OpenCL::Image2D object that refers to the given OpenGL
1313 render buffer.
1314
1315 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLRenderbuffer.html
1316
1317 =item @formats = $ctx->supported_image_formats ($flags, $image_type)
1318
1319 Returns a list of matching image formats - each format is an arrayref with
1320 two values, $channel_order and $channel_type, in it.
1321
1322 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetSupportedImageFormats.html>
1323
1324 =item $sampler = $ctx->sampler ($normalized_coords, $addressing_mode, $filter_mode)
1325
1326 Creates a new OpenCL::Sampler object.
1327
1328 addressing_mode: OpenCL::ADDRESS_NONE, OpenCL::ADDRESS_CLAMP_TO_EDGE,
1329 OpenCL::ADDRESS_CLAMP, OpenCL::ADDRESS_REPEAT, OpenCL::ADDRESS_MIRRORED_REPEAT.
1330
1331 filter_mode: OpenCL::FILTER_NEAREST, OpenCL::FILTER_LINEAR.
1332
1333 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateSampler.html>
1334
1335 =item $program = $ctx->program_with_source ($string)
1336
1337 Creates a new OpenCL::Program object from the given source code.
1338
1339 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateProgramWithSource.html>
1340
1341 =item ($program, \@status) = $ctx->program_with_binary (\@devices, \@binaries)
1342
1343 Creates a new OpenCL::Program object from the given binaries.
1344
1345 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateProgramWithBinary.html>
1346
1347 Example: clone an existing program object that contains a successfully
1348 compiled program, no matter how useless this is.
1349
1350 my $clone = $ctx->program_with_binary ([$prog->devices], [$prog->binaries]);
1351
1352 =item $program = $ctx->program_with_built_in_kernels (\@devices, $kernel_names)
1353
1354 Creates a new OpenCL::Program object from the given built-in kernel names.
1355
1356 L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateProgramWithBuiltInKernels.html>
1357
1358 =item $program = $ctx->link_program (\@devices, $options, \@programs, $cb->($program) = undef)
1359
1360 Links all (already compiled) program objects specified in C<@programs>
1361 together and returns a new OpenCL::Program object with the result.
1362
1363 L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clLinkProgram.html>
1364
1365 =item $packed_value = $ctx->info ($name)
1366
1367 See C<< $platform->info >> for details.
1368
1369 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetContextInfo.html>
1370
1371 =for gengetinfo begin context
1372
1373 =item $uint = $context->reference_count
1374
1375 Calls C<clGetContextInfo> with C<OpenCL::CONTEXT_REFERENCE_COUNT> and returns the result.
1376
1377 =item @devices = $context->devices
1378
1379 Calls C<clGetContextInfo> with C<OpenCL::CONTEXT_DEVICES> and returns the result.
1380
1381 =item @property_ints = $context->properties
1382
1383 Calls C<clGetContextInfo> with C<OpenCL::CONTEXT_PROPERTIES> and returns the result.
1384
1385 =item $uint = $context->num_devices
1386
1387 Calls C<clGetContextInfo> with C<OpenCL::CONTEXT_NUM_DEVICES> and returns the result.
1388
1389 =for gengetinfo end context
1390
1391 =back
1392
1393 =head2 THE OpenCL::Queue CLASS
1394
1395 An OpenCL::Queue represents an execution queue for OpenCL. You execute
1396 requests by calling their respective method and waiting for it to complete
1397 in some way.
1398
1399 Most methods that enqueue some request return an event object that can
1400 be used to wait for completion (optionally using a callback), unless
1401 the method is called in void context, in which case no event object is
1402 created.
1403
1404 They also allow you to specify any number of other event objects that this
1405 request has to wait for before it starts executing, by simply passing the
1406 event objects as extra parameters to the enqueue methods. To simplify
1407 program design, this module ignores any C<undef> values in the list of
1408 events. This makes it possible to code operations such as this, without
1409 having to put a valid event object into C<$event> first:
1410
1411 $event = $queue->xxx (..., $event);
1412
1413 Queues execute in-order by default, without any parallelism, so in most
1414 cases (i.e. you use only one queue) it's not necessary to wait for or
1415 create event objects, althoguh an our of order queue is often a bit
1416 faster.
1417
1418 =over 4
1419
1420 =item $ev = $queue->read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...)
1421
1422 Reads data from buffer into the given string.
1423
1424 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBuffer.html>
1425
1426 =item $ev = $queue->write_buffer ($buffer, $blocking, $offset, $data, $wait_events...)
1427
1428 Writes data to buffer from the given string.
1429
1430 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBuffer.html>
1431
1432 =item $ev = $queue->copy_buffer ($src, $dst, $src_offset, $dst_offset, $len, $wait_events...)
1433
1434 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBuffer.html>
1435
1436 $eue->read_buffer_rect ($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...)
1437
1438 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBufferRect.html
1439
1440 =item $ev = $queue->write_buffer_rect ($buf, $blocking, $buf_y, $host_x, $host_z, $height, $buf_row_pitch, $host_row_pitch, $data, $wait_events...)
1441
1442 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBufferRect.html
1443
1444 =item $ev = $queue->copy_buffer_to_image ($src_buffer, $dst_image, $src_offset, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $wait_events...)
1445
1446 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>
1447
1448 =item $ev = $queue->read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
1449
1450 C<$row_pitch> (and C<$slice_pitch>) can be C<0>, in which case the OpenCL
1451 module uses the image width (and height) to supply default values.
1452
1453 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadImage.html>
1454
1455 =item $ev = $queue->write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
1456
1457 C<$row_pitch> (and C<$slice_pitch>) can be C<0>, in which case the OpenCL
1458 module uses the image width (and height) to supply default values.
1459 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteImage.html>
1460
1461 =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...)
1462
1463 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html>
1464
1465 =item $ev = $queue->copy_image_to_buffer ($src_image, $dst_image, $src_x, $src_y, $src_z, $width, $height, $depth, $dst_offset, $wait_events...)
1466
1467 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html>
1468
1469 =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...)
1470
1471 Yeah.
1472
1473 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>.
1474
1475 =item $ev = $queue->fill_buffer ($mem, $pattern, $offset, $size, ...)
1476
1477 Fills the given buffer object with repeated applications of C<$pattern>,
1478 starting at C<$offset> for C<$size> octets.
1479
1480 L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillBuffer.html>
1481
1482 =item $ev = $queue->fill_image ($img, $r, $g, $b, $a, $x, $y, $z, $width, $height, $depth, ...)
1483
1484 Fills the given image area with the given rgba colour components. The
1485 components are normally floating point values between C<0> and C<1>,
1486 except when the image channel data type is a signe dor unsigned
1487 unnormalised format, in which case the range is determined by the format.
1488
1489 L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillImage.html>
1490
1491 =item $ev = $queue->task ($kernel, $wait_events...)
1492
1493 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html>
1494
1495 =item $ev = $queue->nd_range_kernel ($kernel, \@global_work_offset, \@global_work_size, \@local_work_size, $wait_events...)
1496
1497 Enqueues a kernel execution.
1498
1499 \@global_work_size must be specified as a reference to an array of
1500 integers specifying the work sizes (element counts).
1501
1502 \@global_work_offset must be either C<undef> (in which case all offsets
1503 are C<0>), or a reference to an array of work offsets, with the same number
1504 of elements as \@global_work_size.
1505
1506 \@local_work_size must be either C<undef> (in which case the
1507 implementation is supposed to choose good local work sizes), or a
1508 reference to an array of local work sizes, with the same number of
1509 elements as \@global_work_size.
1510
1511 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html>
1512
1513 =item $ev = $queue->migrate_mem_objects (\@mem_objects, $flags, $wait_events...)
1514
1515 Migrates a number of OpenCL::Memory objects to or from the device.
1516
1517 flags: OpenCL::MIGRATE_MEM_OBJECT_HOST, OpenCL::MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED
1518
1519 L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMigrateMemObjects.html>
1520
1521 =item $ev = $queue->acquire_gl_objects ([object, ...], $wait_events...)
1522
1523 Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired
1524 for subsequent OpenCL usage.
1525
1526 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueAcquireGLObjects.html>
1527
1528 =item $ev = $queue->release_gl_objects ([object, ...], $wait_events...)
1529
1530 Enqueues a list (an array-ref of OpenCL::Memory objects) to be released
1531 for subsequent OpenGL usage.
1532
1533 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReleaseGLObjects.html>
1534
1535 =item $ev = $queue->wait_for_events ($wait_events...)
1536
1537 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html>
1538
1539 =item $ev = $queue->marker ($wait_events...)
1540
1541 L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMarkerWithWaitList.html>
1542
1543 =item $ev = $queue->barrier ($wait_events...)
1544
1545 L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueBarrierWithWaitList.html>
1546
1547 =item $queue->flush
1548
1549 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clFlush.html>
1550
1551 =item $queue->finish
1552
1553 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clFinish.html>
1554
1555 =item $packed_value = $queue->info ($name)
1556
1557 See C<< $platform->info >> for details.
1558
1559 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetCommandQueueInfo.html>
1560
1561 =for gengetinfo begin command_queue
1562
1563 =item $ctx = $command_queue->context
1564
1565 Calls C<clGetCommandQueueInfo> with C<OpenCL::QUEUE_CONTEXT> and returns the result.
1566
1567 =item $device = $command_queue->device
1568
1569 Calls C<clGetCommandQueueInfo> with C<OpenCL::QUEUE_DEVICE> and returns the result.
1570
1571 =item $uint = $command_queue->reference_count
1572
1573 Calls C<clGetCommandQueueInfo> with C<OpenCL::QUEUE_REFERENCE_COUNT> and returns the result.
1574
1575 =item $command_queue_properties = $command_queue->properties
1576
1577 Calls C<clGetCommandQueueInfo> with C<OpenCL::QUEUE_PROPERTIES> and returns the result.
1578
1579 =for gengetinfo end command_queue
1580
1581 =back
1582
1583 =head3 MEMORY MAPPED BUFFERS
1584
1585 OpenCL allows you to map buffers and images to host memory (read: perl
1586 scalars). This is done much like reading or copying a buffer, by enqueuing
1587 a map or unmap operation on the command queue.
1588
1589 The map operations return an C<OpenCL::Mapped> object - see L<THE
1590 OpenCL::Mapped CLASS> section for details on what to do with these
1591 objects.
1592
1593 The object will be unmapped automatically when the mapped object is
1594 destroyed (you can use a barrier to make sure the unmap has finished,
1595 before using the buffer in a kernel), but you can also enqueue an unmap
1596 operation manually.
1597
1598 =over 4
1599
1600 =item $mapped_buffer = $queue->map_buffer ($buf, $blocking=1, $map_flags=OpenCL::MAP_READ|OpenCL::MAP_WRITE, $offset=0, $size=undef, $wait_events...)
1601
1602 Maps the given buffer into host memory and returns an
1603 C<OpenCL::MappedBuffer> object. If C<$size> is specified as undef, then
1604 the map will extend to the end of the buffer.
1605
1606 map_flags: OpenCL::MAP_READ, OpenCL::MAP_WRITE, OpenCL::MAP_WRITE_INVALIDATE_REGION.
1607
1608 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMapBuffer.html>
1609
1610 Example: map the buffer $buf fully and replace the first 4 bytes by "abcd", then unmap.
1611
1612 {
1613 my $mapped = $queue->map_buffer ($buf, 1, OpenCL::MAP_WRITE);
1614 substr $$mapped, 0, 4, "abcd";
1615 } # asynchronously unmap because $mapped is destroyed
1616
1617 =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...)
1618
1619 Maps the given image area into host memory and return an
1620 C<OpenCL::MappedImage> object.
1621
1622 If any of C<$width>, C<$height> and/or C<$depth> are C<undef> then they
1623 will be replaced by the maximum possible value.
1624
1625 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMapImage.html>
1626
1627 Example: map an image (with OpenCL::UNSIGNED_INT8 channel type) and set
1628 the first channel of the leftmost column to 5, then explicitly unmap
1629 it. You are not necessarily meant to do it this way, this example just
1630 shows you the accessors to use :)
1631
1632 my $mapped = $queue->map_image ($image, 1, OpenCL::MAP_WRITE);
1633
1634 $mapped->write ($_ * $mapped->row_pitch, pack "C", 5)
1635 for 0 .. $mapped->height - 1;
1636
1637 $mapped->unmap;.
1638 $mapped->wait; # only needed for out of order queues normally
1639
1640 =item $ev = $queue->unmap ($mapped, $wait_events...)
1641
1642 Unmaps the data from host memory. You must not call any methods that
1643 modify the data, or modify the data scalar directly, after calling this
1644 method.
1645
1646 The mapped event object will always be passed as part of the
1647 $wait_events. The mapped event object will be replaced by the new event
1648 object that this request creates.
1649
1650 =back
1651
1652 =head2 THE OpenCL::Memory CLASS
1653
1654 This the superclass of all memory objects - OpenCL::Buffer, OpenCL::Image,
1655 OpenCL::Image2D and OpenCL::Image3D.
1656
1657 =over 4
1658
1659 =item $packed_value = $memory->info ($name)
1660
1661 See C<< $platform->info >> for details.
1662
1663 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetMemObjectInfo.html>
1664
1665 =item $memory->destructor_callback ($cb->())
1666
1667 Sets a callback that will be invoked after the memory object is destructed.
1668
1669 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetMemObjectDestructorCallback.html>
1670
1671 =for gengetinfo begin mem
1672
1673 =item $mem_object_type = $mem->type
1674
1675 Calls C<clGetMemObjectInfo> with C<OpenCL::MEM_TYPE> and returns the result.
1676
1677 =item $mem_flags = $mem->flags
1678
1679 Calls C<clGetMemObjectInfo> with C<OpenCL::MEM_FLAGS> and returns the result.
1680
1681 =item $int = $mem->size
1682
1683 Calls C<clGetMemObjectInfo> with C<OpenCL::MEM_SIZE> and returns the result.
1684
1685 =item $ptr_value = $mem->host_ptr
1686
1687 Calls C<clGetMemObjectInfo> with C<OpenCL::MEM_HOST_PTR> and returns the result.
1688
1689 =item $uint = $mem->map_count
1690
1691 Calls C<clGetMemObjectInfo> with C<OpenCL::MEM_MAP_COUNT> and returns the result.
1692
1693 =item $uint = $mem->reference_count
1694
1695 Calls C<clGetMemObjectInfo> with C<OpenCL::MEM_REFERENCE_COUNT> and returns the result.
1696
1697 =item $ctx = $mem->context
1698
1699 Calls C<clGetMemObjectInfo> with C<OpenCL::MEM_CONTEXT> and returns the result.
1700
1701 =item $mem = $mem->associated_memobject
1702
1703 Calls C<clGetMemObjectInfo> with C<OpenCL::MEM_ASSOCIATED_MEMOBJECT> and returns the result.
1704
1705 =item $int = $mem->offset
1706
1707 Calls C<clGetMemObjectInfo> with C<OpenCL::MEM_OFFSET> and returns the result.
1708
1709 =for gengetinfo end mem
1710
1711 =item ($type, $name) = $mem->gl_object_info
1712
1713 Returns the OpenGL object type (e.g. OpenCL::GL_OBJECT_TEXTURE2D) and the
1714 object "name" (e.g. the texture name) used to create this memory object.
1715
1716 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetGLObjectInfo.html>
1717
1718 =back
1719
1720 =head2 THE OpenCL::Buffer CLASS
1721
1722 This is a subclass of OpenCL::Memory, and the superclass of
1723 OpenCL::BufferObj. Its purpose is simply to distinguish between buffers
1724 and sub-buffers.
1725
1726 =head2 THE OpenCL::BufferObj CLASS
1727
1728 This is a subclass of OpenCL::Buffer and thus OpenCL::Memory. It exists
1729 because one cna create sub buffers of OpenLC::BufferObj objects, but not
1730 sub buffers from these sub buffers.
1731
1732 =over 4
1733
1734 =item $subbuf = $buf_obj->sub_buffer_region ($flags, $origin, $size)
1735
1736 Creates an OpenCL::Buffer objects from this buffer and returns it. The
1737 C<buffer_create_type> is assumed to be C<OpenCL::BUFFER_CREATE_TYPE_REGION>.
1738
1739 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateSubBuffer.html>
1740
1741 =back
1742
1743 =head2 THE OpenCL::Image CLASS
1744
1745 This is the superclass of all image objects - OpenCL::Image1D,
1746 OpenCL::Image1DArray, OpenCL::Image1DBuffer, OpenCL::Image2D,
1747 OpenCL::Image2DArray and OpenCL::Image3D.
1748
1749 =over 4
1750
1751 =item $packed_value = $image->image_info ($name)
1752
1753 See C<< $platform->info >> for details.
1754
1755 The reason this method is not called C<info> is that there already is an
1756 C<< ->info >> method inherited from C<OpenCL::Memory>.
1757
1758 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetImageInfo.html>
1759
1760 =item ($channel_order, $channel_data_type) = $image->format
1761
1762 Returns the channel order and type used to create the image by calling
1763 C<clGetImageInfo> with C<OpenCL::IMAGE_FORMAT>.
1764
1765 =for gengetinfo begin image
1766
1767 =item $int = $image->element_size
1768
1769 Calls C<clGetImageInfo> with C<OpenCL::IMAGE_ELEMENT_SIZE> and returns the result.
1770
1771 =item $int = $image->row_pitch
1772
1773 Calls C<clGetImageInfo> with C<OpenCL::IMAGE_ROW_PITCH> and returns the result.
1774
1775 =item $int = $image->slice_pitch
1776
1777 Calls C<clGetImageInfo> with C<OpenCL::IMAGE_SLICE_PITCH> and returns the result.
1778
1779 =item $int = $image->width
1780
1781 Calls C<clGetImageInfo> with C<OpenCL::IMAGE_WIDTH> and returns the result.
1782
1783 =item $int = $image->height
1784
1785 Calls C<clGetImageInfo> with C<OpenCL::IMAGE_HEIGHT> and returns the result.
1786
1787 =item $int = $image->depth
1788
1789 Calls C<clGetImageInfo> with C<OpenCL::IMAGE_DEPTH> and returns the result.
1790
1791 =for gengetinfo end image
1792
1793 =for gengetinfo begin gl_texture
1794
1795 =item $GLenum = $gl_texture->target
1796
1797 Calls C<clGetGLTextureInfo> with C<OpenCL::GL_TEXTURE_TARGET> and returns the result.
1798
1799 =item $GLint = $gl_texture->gl_mipmap_level
1800
1801 Calls C<clGetGLTextureInfo> with C<OpenCL::GL_MIPMAP_LEVEL> and returns the result.
1802
1803 =for gengetinfo end gl_texture
1804
1805 =back
1806
1807 =head2 THE OpenCL::Sampler CLASS
1808
1809 =over 4
1810
1811 =item $packed_value = $sampler->info ($name)
1812
1813 See C<< $platform->info >> for details.
1814
1815 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetSamplerInfo.html>
1816
1817 =for gengetinfo begin sampler
1818
1819 =item $uint = $sampler->reference_count
1820
1821 Calls C<clGetSamplerInfo> with C<OpenCL::SAMPLER_REFERENCE_COUNT> and returns the result.
1822
1823 =item $ctx = $sampler->context
1824
1825 Calls C<clGetSamplerInfo> with C<OpenCL::SAMPLER_CONTEXT> and returns the result.
1826
1827 =item $addressing_mode = $sampler->normalized_coords
1828
1829 Calls C<clGetSamplerInfo> with C<OpenCL::SAMPLER_NORMALIZED_COORDS> and returns the result.
1830
1831 =item $filter_mode = $sampler->addressing_mode
1832
1833 Calls C<clGetSamplerInfo> with C<OpenCL::SAMPLER_ADDRESSING_MODE> and returns the result.
1834
1835 =item $boolean = $sampler->filter_mode
1836
1837 Calls C<clGetSamplerInfo> with C<OpenCL::SAMPLER_FILTER_MODE> and returns the result.
1838
1839 =for gengetinfo end sampler
1840
1841 =back
1842
1843 =head2 THE OpenCL::Program CLASS
1844
1845 =over 4
1846
1847 =item $program->build (\@devices = undef, $options = "", $cb->($program) = undef)
1848
1849 Tries to build the program with the given options. See also the
1850 C<$ctx->build> convenience function.
1851
1852 If a callback is specified, then it will be called when compilation is
1853 finished. Note that many OpenCL implementations block your program while
1854 compiling whether you use a callback or not. See C<build_async> if you
1855 want to make sure the build is done in the background.
1856
1857 Note that some OpenCL implementations act up badly, and don't call the
1858 callback in some error cases (but call it in others). This implementation
1859 assumes the callback will always be called, and leaks memory if this is
1860 not so. So best make sure you don't pass in invalid values.
1861
1862 Some implementations fail with C<OpenCL::INVALID_BINARY> when the
1863 compilation state is successful but some later stage fails.
1864
1865 options: C<-D name>, C<-D name=definition>, C<-I dir>,
1866 C<-cl-single-precision-constant>, C<-cl-denorms-are-zero>,
1867 C<-cl-fp32-correctly-rounded-divide-sqrt>, C<-cl-opt-disable>,
1868 C<-cl-mad-enable>, C<-cl-no-signed-zeros>, C<-cl-unsafe-math-optimizations>,
1869 C<-cl-finite-math-only>, C<-cl-fast-relaxed-math>,
1870 C<-w>, C<-Werror>, C<-cl-std=CL1.1/CL1.2>, C<-cl-kernel-arg-info>,
1871 C<-create-library>, C<-enable-link-options>.
1872
1873 build_status: OpenCL::BUILD_SUCCESS, OpenCL::BUILD_NONE,
1874 OpenCL::BUILD_ERROR, OpenCL::BUILD_IN_PROGRESS.
1875
1876 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clBuildProgram.html>
1877
1878 =item $program->build_async (\@devices = undef, $options = "", $cb->($program) = undef)
1879
1880 Similar to C<< ->build >>, except it starts a thread, and never fails (you
1881 need to check the compilation status form the callback, or by polling).
1882
1883 =item $program->compile (\@devices = undef, $options = "", \%headers = undef, $cb->($program) = undef)
1884
1885 Compiles the given program for the given devices (or all devices if
1886 undef). If C<$headers> is given, it must be a hashref with include name =>
1887 OpenCL::Program pairs.
1888
1889 L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCompileProgram.html>
1890
1891 =item $packed_value = $program->build_info ($device, $name)
1892
1893 Similar to C<< $platform->info >>, but returns build info for a previous
1894 build attempt for the given device.
1895
1896 binary_type: OpenCL::PROGRAM_BINARY_TYPE_NONE,
1897 OpenCL::PROGRAM_BINARY_TYPE_COMPILED_OBJECT,
1898 OpenCL::PROGRAM_BINARY_TYPE_LIBRARY,
1899 OpenCL::PROGRAM_BINARY_TYPE_EXECUTABLE.
1900
1901 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetBuildInfo.html>
1902
1903 =item $kernel = $program->kernel ($function_name)
1904
1905 Creates an OpenCL::Kernel object out of the named C<__kernel> function in
1906 the program.
1907
1908 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernel.html>
1909
1910 =item @kernels = $program->kernels_in_program
1911
1912 Returns all kernels successfully compiled for all devices in program.
1913
1914 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernelsInProgram.html
1915
1916 =for gengetinfo begin program_build
1917
1918 =item $build_status = $program->build_status ($device)
1919
1920 Calls C<clGetProgramBuildInfo> with C<OpenCL::PROGRAM_BUILD_STATUS> and returns the result.
1921
1922 =item $string = $program->build_options ($device)
1923
1924 Calls C<clGetProgramBuildInfo> with C<OpenCL::PROGRAM_BUILD_OPTIONS> and returns the result.
1925
1926 =item $string = $program->build_log ($device)
1927
1928 Calls C<clGetProgramBuildInfo> with C<OpenCL::PROGRAM_BUILD_LOG> and returns the result.
1929
1930 =item $binary_type = $program->binary_type ($device)
1931
1932 Calls C<clGetProgramBuildInfo> with C<OpenCL::PROGRAM_BINARY_TYPE> and returns the result.
1933
1934 =for gengetinfo end program_build
1935
1936 =item $packed_value = $program->info ($name)
1937
1938 See C<< $platform->info >> for details.
1939
1940 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetProgramInfo.html>
1941
1942 =for gengetinfo begin program
1943
1944 =item $uint = $program->reference_count
1945
1946 Calls C<clGetProgramInfo> with C<OpenCL::PROGRAM_REFERENCE_COUNT> and returns the result.
1947
1948 =item $ctx = $program->context
1949
1950 Calls C<clGetProgramInfo> with C<OpenCL::PROGRAM_CONTEXT> and returns the result.
1951
1952 =item $uint = $program->num_devices
1953
1954 Calls C<clGetProgramInfo> with C<OpenCL::PROGRAM_NUM_DEVICES> and returns the result.
1955
1956 =item @devices = $program->devices
1957
1958 Calls C<clGetProgramInfo> with C<OpenCL::PROGRAM_DEVICES> and returns the result.
1959
1960 =item $string = $program->source
1961
1962 Calls C<clGetProgramInfo> with C<OpenCL::PROGRAM_SOURCE> and returns the result.
1963
1964 =item @ints = $program->binary_sizes
1965
1966 Calls C<clGetProgramInfo> with C<OpenCL::PROGRAM_BINARY_SIZES> and returns the result.
1967
1968 =for gengetinfo end program
1969
1970 =item @blobs = $program->binaries
1971
1972 Returns a string for the compiled binary for every device associated with
1973 the program, empty strings indicate missing programs, and an empty result
1974 means no program binaries are available.
1975
1976 These "binaries" are often, in fact, informative low-level assembly
1977 sources.
1978
1979 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetProgramInfo.html>
1980
1981 =back
1982
1983 =head2 THE OpenCL::Kernel CLASS
1984
1985 =over 4
1986
1987 =item $packed_value = $kernel->info ($name)
1988
1989 See C<< $platform->info >> for details.
1990
1991 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetKernelInfo.html>
1992
1993 =for gengetinfo begin kernel
1994
1995 =item $string = $kernel->function_name
1996
1997 Calls C<clGetKernelInfo> with C<OpenCL::KERNEL_FUNCTION_NAME> and returns the result.
1998
1999 =item $uint = $kernel->num_args
2000
2001 Calls C<clGetKernelInfo> with C<OpenCL::KERNEL_NUM_ARGS> and returns the result.
2002
2003 =item $uint = $kernel->reference_count
2004
2005 Calls C<clGetKernelInfo> with C<OpenCL::KERNEL_REFERENCE_COUNT> and returns the result.
2006
2007 =item $ctx = $kernel->context
2008
2009 Calls C<clGetKernelInfo> with C<OpenCL::KERNEL_CONTEXT> and returns the result.
2010
2011 =item $program = $kernel->program
2012
2013 Calls C<clGetKernelInfo> with C<OpenCL::KERNEL_PROGRAM> and returns the result.
2014
2015 =for gengetinfo end kernel
2016
2017 =item $packed_value = $kernel->work_group_info ($device, $name)
2018
2019 See C<< $platform->info >> for details.
2020
2021 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetKernelWorkGroupInfo.html>
2022
2023 =for gengetinfo begin kernel_work_group
2024
2025 =item $int = $kernel->work_group_size ($device)
2026
2027 Calls C<clGetKernelWorkGroupInfo> with C<OpenCL::KERNEL_WORK_GROUP_SIZE> and returns the result.
2028
2029 =item @ints = $kernel->compile_work_group_size ($device)
2030
2031 Calls C<clGetKernelWorkGroupInfo> with C<OpenCL::KERNEL_COMPILE_WORK_GROUP_SIZE> and returns the result.
2032
2033 =item $ulong = $kernel->local_mem_size ($device)
2034
2035 Calls C<clGetKernelWorkGroupInfo> with C<OpenCL::KERNEL_LOCAL_MEM_SIZE> and returns the result.
2036
2037 =item $int = $kernel->preferred_work_group_size_multiple ($device)
2038
2039 Calls C<clGetKernelWorkGroupInfo> with C<OpenCL::KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE> and returns the result.
2040
2041 =item $ulong = $kernel->private_mem_size ($device)
2042
2043 Calls C<clGetKernelWorkGroupInfo> with C<OpenCL::KERNEL_PRIVATE_MEM_SIZE> and returns the result.
2044
2045 =for gengetinfo end kernel_work_group
2046
2047 =item $packed_value = $kernel->arg_info ($idx, $name)
2048
2049 See C<< $platform->info >> for details.
2050
2051 L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clGetKernelArgInfo.html>
2052
2053 =for gengetinfo begin kernel_arg
2054
2055 =item $kernel_arg_address_qualifier = $kernel->arg_address_qualifier ($idx)
2056
2057 Calls C<clGetKernelArgInfo> with C<OpenCL::KERNEL_ARG_ADDRESS_QUALIFIER> and returns the result.
2058
2059 =item $kernel_arg_access_qualifier = $kernel->arg_access_qualifier ($idx)
2060
2061 Calls C<clGetKernelArgInfo> with C<OpenCL::KERNEL_ARG_ACCESS_QUALIFIER> and returns the result.
2062
2063 =item $string = $kernel->arg_type_name ($idx)
2064
2065 Calls C<clGetKernelArgInfo> with C<OpenCL::KERNEL_ARG_TYPE_NAME> and returns the result.
2066
2067 =item $kernel_arg_type_qualifier = $kernel->arg_type_qualifier ($idx)
2068
2069 Calls C<clGetKernelArgInfo> with C<OpenCL::KERNEL_ARG_TYPE_QUALIFIER> and returns the result.
2070
2071 =item $string = $kernel->arg_name ($idx)
2072
2073 Calls C<clGetKernelArgInfo> with C<OpenCL::KERNEL_ARG_NAME> and returns the result.
2074
2075 =for gengetinfo end kernel_arg
2076
2077 =item $kernel->setf ($format, ...)
2078
2079 Sets the arguments of a kernel. Since OpenCL 1.1 doesn't have a generic
2080 way to set arguments (and with OpenCL 1.2 it might be rather slow), you
2081 need to specify a format argument, much as with C<printf>, to tell OpenCL
2082 what type of argument it is.
2083
2084 The format arguments are single letters:
2085
2086 c char
2087 C unsigned char
2088 s short
2089 S unsigned short
2090 i int
2091 I unsigned int
2092 l long
2093 L unsigned long
2094
2095 h half float (0..65535)
2096 f float
2097 d double
2098
2099 z local (octet size)
2100
2101 m memory object (buffer or image)
2102 a sampler
2103 e event
2104
2105 Space characters in the format string are ignored.
2106
2107 Example: set the arguments for a kernel that expects an int, two floats, a buffer and an image.
2108
2109 $kernel->setf ("i ff mm", 5, 0.5, 3, $buffer, $image);
2110
2111 =item $kernel->set_TYPE ($index, $value)
2112
2113 =item $kernel->set_char ($index, $value)
2114
2115 =item $kernel->set_uchar ($index, $value)
2116
2117 =item $kernel->set_short ($index, $value)
2118
2119 =item $kernel->set_ushort ($index, $value)
2120
2121 =item $kernel->set_int ($index, $value)
2122
2123 =item $kernel->set_uint ($index, $value)
2124
2125 =item $kernel->set_long ($index, $value)
2126
2127 =item $kernel->set_ulong ($index, $value)
2128
2129 =item $kernel->set_half ($index, $value)
2130
2131 =item $kernel->set_float ($index, $value)
2132
2133 =item $kernel->set_double ($index, $value)
2134
2135 =item $kernel->set_memory ($index, $value)
2136
2137 =item $kernel->set_buffer ($index, $value)
2138
2139 =item $kernel->set_image ($index, $value)
2140
2141 =item $kernel->set_sampler ($index, $value)
2142
2143 =item $kernel->set_local ($index, $value)
2144
2145 =item $kernel->set_event ($index, $value)
2146
2147 This is a family of methods to set the kernel argument with the number
2148 C<$index> to the give C<$value>.
2149
2150 Chars and integers (including the half type) are specified as integers,
2151 float and double as floating point values, memory/buffer/image must be
2152 an object of that type or C<undef>, local-memory arguments are set by
2153 specifying the size, and sampler and event must be objects of that type.
2154
2155 Note that C<set_memory> works for all memory objects (all types of buffers
2156 and images) - the main purpose of the more specific C<set_TYPE> functions
2157 is type checking.
2158
2159 Setting an argument for a kernel does NOT keep a reference to the object -
2160 for example, if you set an argument to some image object, free the image,
2161 and call the kernel, you will run into undefined behaviour.
2162
2163 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html>
2164
2165 =back
2166
2167 =head2 THE OpenCL::Event CLASS
2168
2169 This is the superclass for all event objects (including OpenCL::UserEvent
2170 objects).
2171
2172 =over 4
2173
2174 =item $ev->wait
2175
2176 Waits for the event to complete.
2177
2178 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html>
2179
2180 =item $ev->cb ($exec_callback_type, $callback->($event, $event_command_exec_status))
2181
2182 Adds a callback to the callback stack for the given event type. There is
2183 no way to remove a callback again.
2184
2185 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetEventCallback.html>
2186
2187 =item $packed_value = $ev->info ($name)
2188
2189 See C<< $platform->info >> for details.
2190
2191 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetEventInfo.html>
2192
2193 =for gengetinfo begin event
2194
2195 =item $queue = $event->command_queue
2196
2197 Calls C<clGetEventInfo> with C<OpenCL::EVENT_COMMAND_QUEUE> and returns the result.
2198
2199 =item $command_type = $event->command_type
2200
2201 Calls C<clGetEventInfo> with C<OpenCL::EVENT_COMMAND_TYPE> and returns the result.
2202
2203 =item $uint = $event->reference_count
2204
2205 Calls C<clGetEventInfo> with C<OpenCL::EVENT_REFERENCE_COUNT> and returns the result.
2206
2207 =item $uint = $event->command_execution_status
2208
2209 Calls C<clGetEventInfo> with C<OpenCL::EVENT_COMMAND_EXECUTION_STATUS> and returns the result.
2210
2211 =item $ctx = $event->context
2212
2213 Calls C<clGetEventInfo> with C<OpenCL::EVENT_CONTEXT> and returns the result.
2214
2215 =for gengetinfo end event
2216
2217 =item $packed_value = $ev->profiling_info ($name)
2218
2219 See C<< $platform->info >> for details.
2220
2221 The reason this method is not called C<info> is that there already is an
2222 C<< ->info >> method.
2223
2224 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetProfilingInfo.html>
2225
2226 =for gengetinfo begin profiling
2227
2228 =item $ulong = $event->profiling_command_queued
2229
2230 Calls C<clGetEventProfilingInfo> with C<OpenCL::PROFILING_COMMAND_QUEUED> and returns the result.
2231
2232 =item $ulong = $event->profiling_command_submit
2233
2234 Calls C<clGetEventProfilingInfo> with C<OpenCL::PROFILING_COMMAND_SUBMIT> and returns the result.
2235
2236 =item $ulong = $event->profiling_command_start
2237
2238 Calls C<clGetEventProfilingInfo> with C<OpenCL::PROFILING_COMMAND_START> and returns the result.
2239
2240 =item $ulong = $event->profiling_command_end
2241
2242 Calls C<clGetEventProfilingInfo> with C<OpenCL::PROFILING_COMMAND_END> and returns the result.
2243
2244 =for gengetinfo end profiling
2245
2246 =back
2247
2248 =head2 THE OpenCL::UserEvent CLASS
2249
2250 This is a subclass of OpenCL::Event.
2251
2252 =over 4
2253
2254 =item $ev->set_status ($execution_status)
2255
2256 Sets the execution status of the user event. Can only be called once,
2257 either with OpenCL::COMPLETE or a negative number as status.
2258
2259 execution_status: OpenCL::COMPLETE or a negative integer.
2260
2261 L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetUserEventStatus.html>
2262
2263 =back
2264
2265 =head2 THE OpenCL::Mapped CLASS
2266
2267 This class represents objects mapped into host memory. They are
2268 represented by a blessed string scalar. The string data is the mapped
2269 memory area, that is, if you read or write it, then the mapped object is
2270 accessed directly.
2271
2272 You must only ever use operations that modify the string in-place - for
2273 example, a C<substr> that doesn't change the length, or maybe a regex that
2274 doesn't change the length. Any other operation might cause the data to be
2275 copied.
2276
2277 When the object is destroyed it will enqueue an implicit unmap operation
2278 on the queue that was used to create it.
2279
2280 Keep in mind that you I<need> to unmap (or destroy) mapped objects before
2281 OpenCL sees the changes, even if some implementations don't need this
2282 sometimes.
2283
2284 Example, replace the first two floats in the mapped buffer by 1 and 2.
2285
2286 my $mapped = $queue->map_buffer ($buf, ...
2287 $mapped->event->wait; # make sure it's there
2288
2289 # now replace first 8 bytes by new data, which is exactly 8 bytes long
2290 # we blindly assume device endianness to equal host endianness
2291 # (and of course, we assume iee 754 single precision floats :)
2292 substr $$mapped, 0, 8, pack "f*", 1, 2;
2293
2294 =over 4
2295
2296 =item $ev = $mapped->unmap ($wait_events...)
2297
2298 Unmaps the mapped memory object, using the queue originally used to create
2299 it, quite similarly to C<< $queue->unmap ($mapped, ...) >>.
2300
2301 =item $bool = $mapped->mapped
2302
2303 Returns whether the object is still mapped - true before an C<unmap> is
2304 enqueued, false afterwards.
2305
2306 =item $ev = $mapped->event
2307
2308 Return the event object associated with the mapped object. Initially, this
2309 will be the event object created when mapping the object, and after an
2310 unmap, this will be the event object that the unmap operation created.
2311
2312 =item $mapped->wait
2313
2314 Same as C<< $mapped->event->wait >> - makes sure no operations on this
2315 mapped object are outstanding.
2316
2317 =item $bytes = $mapped->size
2318
2319 Returns the size of the mapped area, in bytes. Same as C<length $$mapped>.
2320
2321 =item $ptr = $mapped->ptr
2322
2323 Returns the raw memory address of the mapped area.
2324
2325 =item $mapped->set ($offset, $data)
2326
2327 Replaces the data at the given C<$offset> in the memory area by the new
2328 C<$data>. This method is safer than direct manipulation of C<$mapped>
2329 because it does bounds-checking, but also slower.
2330
2331 =item $data = $mapped->get ($offset, $length)
2332
2333 Returns (without copying) a scalar representing the data at the given
2334 C<$offset> and C<$length> in the mapped memory area. This is the same as
2335 the following substr, except much slower;
2336
2337 $data = substr $$mapped, $offset, $length
2338
2339 =cut
2340
2341 sub OpenCL::Mapped::get {
2342 substr ${$_[0]}, $_[1], $_[2]
2343 }
2344
2345 =back
2346
2347 =head2 THE OpenCL::MappedBuffer CLASS
2348
2349 This is a subclass of OpenCL::Mapped, representing mapped buffers.
2350
2351 =head2 THE OpenCL::MappedImage CLASS
2352
2353 This is a subclass of OpenCL::Mapped, representing mapped images.
2354
2355 =over 4
2356
2357 =item $pixels = $mapped->width
2358
2359 =item $pixels = $mapped->height
2360
2361 =item $pixels = $mapped->depth
2362
2363 Return the width/height/depth of the mapped image region, in pixels.
2364
2365 =item $bytes = $mapped->row_pitch
2366
2367 =item $bytes = $mapped->slice_pitch
2368
2369 Return the row or slice pitch of the image that has been mapped.
2370
2371 =item $bytes = $mapped->element_size
2372
2373 Return the size of a single pixel.
2374
2375 =item $data = $mapped->get_row ($count, $x=0, $y=0, $z=0)
2376
2377 Return C<$count> pixels from the given coordinates. The pixel data must
2378 be completely contained within a single row.
2379
2380 If C<$count> is C<undef>, then all the remaining pixels in that row are
2381 returned.
2382
2383 =item $mapped->set_row ($data, $x=0, $y=0, $z=0)
2384
2385 Write the given pixel data at the given coordinate. The pixel data must
2386 be completely contained within a single row.
2387
2388 =back
2389
2390 =cut
2391
2392 1;
2393
2394 =head1 AUTHOR
2395
2396 Marc Lehmann <schmorp@schmorp.de>
2397 http://home.schmorp.de/
2398
2399 =cut
2400