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

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.18 by root, Thu Nov 17 06:22:29 2011 UTC vs.
Revision 1.82 by root, Sat May 5 14:18:35 2012 UTC

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

Diff Legend

– Removed lines
+ Added lines
< Changed lines
> Changed lines