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

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

Diff Legend

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