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

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

Diff Legend

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