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

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.19 by root, Sat Nov 19 19:54:04 2011 UTC vs.
Revision 1.68 by root, Tue May 1 22:25:13 2012 UTC

… …
27OpenCL::Program objects, which store source code and, after building for a 27OpenCL::Program objects, which store source code and, after building for a
28specific device ("compiling and linking"), also binary programs. For each 28specific device ("compiling and linking"), also binary programs. For each
29kernel function in a program you can then create an OpenCL::Kernel object 29kernel function in a program you can then create an OpenCL::Kernel object
30which represents basically a function call with argument values. 30which represents basically a function call with argument values.
31 31
32OpenCL::Memory objects of various flavours: OpenCL::Buffers objects (flat 32OpenCL::Memory objects of various flavours: OpenCL::Buffer objects (flat
33memory areas, think arrays or structs) and OpenCL::Image objects (think 2d 33memory areas, think arrays or structs) and OpenCL::Image objects (think 2d
34or 3d array) for bulk data and input and output for kernels. 34or 3d array) for bulk data and input and output for kernels.
35 35
36OpenCL::Sampler objects, which are kind of like texture filter modes in 36OpenCL::Sampler objects, which are kind of like texture filter modes in
37OpenGL. 37OpenGL.
… …
43 43
44OpenCL::Event objects are used to signal when something is complete. 44OpenCL::Event objects are used to signal when something is complete.
45 45
46=head2 HELPFUL RESOURCES 46=head2 HELPFUL RESOURCES
47 47
48The OpenCL spec used to develop this module (1.2 spec was available, but 48The OpenCL specs used to develop this module:
49no implementation was available to me :).
50 49
51 http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf 50 http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
51 http://www.khronos.org/registry/cl/specs/opencl-1.2.pdf
52 http://www.khronos.org/registry/cl/specs/opencl-1.2-extensions.pdf
52 53
53OpenCL manpages: 54OpenCL manpages:
54 55
55 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/ 56 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/
57 http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/
56 58
57If you are into UML class diagrams, the following diagram might help - if 59If you are into UML class diagrams, the following diagram might help - if
58not, it will be mildly cobfusing: 60not, it will be mildly confusing (also, the class hierarchy of this module
61is much more fine-grained):
59 62
60 http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/classDiagram.html 63 http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/classDiagram.html
61 64
62Here's a tutorial from AMD (very AMD-centric, too), not sure how useful it 65Here's a tutorial from AMD (very AMD-centric, too), not sure how useful it
63is, but at least it's free of charge: 66is, but at least it's free of charge:
64 67
65 http://developer.amd.com/zones/OpenCLZone/courses/Documents/Introduction_to_OpenCL_Programming%20Training_Guide%20%28201005%29.pdf 68 http://developer.amd.com/zones/OpenCLZone/courses/Documents/Introduction_to_OpenCL_Programming%20Training_Guide%20%28201005%29.pdf
… …
101 104
102Best run this once to get a feel for the platforms and devices in your 105Best run this once to get a feel for the platforms and devices in your
103system. 106system.
104 107
105 for my $platform (OpenCL::platforms) { 108 for my $platform (OpenCL::platforms) {
106 printf "platform: %s\n", $platform->info (OpenCL::PLATFORM_NAME); 109 printf "platform: %s\n", $platform->name;
107 printf "extensions: %s\n", $platform->info (OpenCL::PLATFORM_EXTENSIONS); 110 printf "extensions: %s\n", $platform->extensions;
108 for my $device ($platform->devices) { 111 for my $device ($platform->devices) {
109 printf "+ device: %s\n", $device->info (OpenCL::DEVICE_NAME); 112 printf "+ device: %s\n", $device->name;
110 my $ctx = $device->context; 113 my $ctx = $platform->context (undef, [$device]);
111 # do stuff 114 # do stuff
112 } 115 }
113 } 116 }
114 117
115=head2 Get a useful context and a command queue. 118=head2 Get a useful context and a command queue.
… …
138=head2 Create a buffer with some predefined data, read it back synchronously, 141=head2 Create a buffer with some predefined data, read it back synchronously,
139then asynchronously. 142then asynchronously.
140 143
141 my $buf = $ctx->buffer_sv (OpenCL::MEM_COPY_HOST_PTR, "helmut"); 144 my $buf = $ctx->buffer_sv (OpenCL::MEM_COPY_HOST_PTR, "helmut");
142 145
143 $queue->enqueue_read_buffer ($buf, 1, 1, 3, my $data); 146 $queue->read_buffer ($buf, 1, 1, 3, my $data);
144 print "$data\n"; 147 print "$data\n";
145 148
146 my $ev = $queue->enqueue_read_buffer ($buf, 0, 1, 3, my $data); 149 my $ev = $queue->read_buffer ($buf, 0, 1, 3, my $data);
147 $ev->wait; 150 $ev->wait;
148 print "$data\n"; # prints "elm" 151 print "$data\n"; # prints "elm"
149 152
150=head2 Create and build a program, then create a kernel out of one of its 153=head2 Create and build a program, then create a kernel out of one of its
151functions. 154functions.
152 155
153 my $src = ' 156 my $src = '
154 __kernel void 157 kernel void
155 squareit (__global float *input, __global float *output) 158 squareit (global float *input, global float *output)
156 { 159 {
157 $id = get_global_id (0); 160 $id = get_global_id (0);
158 output [id] = input [id] * input [id]; 161 output [id] = input [id] * input [id];
159 } 162 }
160 '; 163 ';
161 164
162 my $prog = $ctx->program_with_source ($src); 165 my $prog = $ctx->build_program ($src);
163
164 # build croaks on compile errors, so catch it and print the compile errors
165 eval { $prog->build ($dev); 1 }
166 or die $prog->build_info ($dev, OpenCL::PROGRAM_BUILD_LOG);
167
168 my $kernel = $prog->kernel ("squareit"); 166 my $kernel = $prog->kernel ("squareit");
169 167
170=head2 Create some input and output float buffers, then call the 168=head2 Create some input and output float buffers, then call the
171'squareit' kernel on them. 169'squareit' kernel on them.
172 170
… …
176 # set buffer 174 # set buffer
177 $kernel->set_buffer (0, $input); 175 $kernel->set_buffer (0, $input);
178 $kernel->set_buffer (1, $output); 176 $kernel->set_buffer (1, $output);
179 177
180 # execute it for all 4 numbers 178 # execute it for all 4 numbers
181 $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); 179 $queue->nd_range_kernel ($kernel, undef, [4], undef);
182 180
183 # enqueue a synchronous read 181 # enqueue a synchronous read
184 $queue->enqueue_read_buffer ($output, 1, 0, OpenCL::SIZEOF_FLOAT * 4, my $data); 182 $queue->read_buffer ($output, 1, 0, OpenCL::SIZEOF_FLOAT * 4, my $data);
185 183
186 # print the results: 184 # print the results:
187 printf "%s\n", join ", ", unpack "f*", $data; 185 printf "%s\n", join ", ", unpack "f*", $data;
188 186
189=head2 The same enqueue operations as before, but assuming an out-of-order queue, 187=head2 The same enqueue operations as before, but assuming an out-of-order queue,
190showing off barriers. 188showing off barriers.
191 189
192 # execute it for all 4 numbers 190 # execute it for all 4 numbers
193 $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); 191 $queue->nd_range_kernel ($kernel, undef, [4], undef);
194 192
195 # enqueue a barrier to ensure in-order execution 193 # enqueue a barrier to ensure in-order execution
196 $queue->enqueue_barrier; 194 $queue->barrier;
197 195
198 # enqueue an async read 196 # enqueue an async read
199 $queue->enqueue_read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data); 197 $queue->read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data);
200 198
201 # wait for all requests to finish 199 # wait for all requests to finish
202 $queue->finish; 200 $queue->finish;
203 201
204=head2 The same enqueue operations as before, but assuming an out-of-order queue, 202=head2 The same enqueue operations as before, but assuming an out-of-order queue,
205showing off event objects and wait lists. 203showing off event objects and wait lists.
206 204
207 # execute it for all 4 numbers 205 # execute it for all 4 numbers
208 my $ev = $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); 206 my $ev = $queue->nd_range_kernel ($kernel, undef, [4], undef);
209 207
210 # enqueue an async read 208 # enqueue an async read
211 $ev = $queue->enqueue_read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data, $ev); 209 $ev = $queue->read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data, $ev);
212 210
213 # wait for the last event to complete 211 # wait for the last event to complete
214 $ev->wait; 212 $ev->wait;
213
214=head2 Use the OpenGL module to share a texture between OpenCL and OpenGL and draw some julia
215set tunnel effect.
216
217This is quite a long example to get you going - you can download it from
218L<http://cvs.schmorp.de/OpenCL/examples/juliaflight>.
219
220 use OpenGL ":all";
221 use OpenCL;
222
223 my $S = $ARGV[0] || 256; # window/texture size, smaller is faster
224
225 # open a window and create a gl texture
226 OpenGL::glpOpenWindow width => $S, height => $S;
227 my $texid = glGenTextures_p 1;
228 glBindTexture GL_TEXTURE_2D, $texid;
229 glTexImage2D_c GL_TEXTURE_2D, 0, GL_RGBA8, $S, $S, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0;
230
231 # find and use the first opencl device that let's us get a shared opengl context
232 my $platform;
233 my $dev;
234 my $ctx;
235
236 for (OpenCL::platforms) {
237 $platform = $_;
238 for ($platform->devices) {
239 $dev = $_;
240 $ctx = $platform->context ([OpenCL::GLX_DISPLAY_KHR, undef, OpenCL::GL_CONTEXT_KHR, undef], [$dev])
241 and last;
242 }
243 }
244
245 $ctx
246 or die "cannot find suitable OpenCL device\n";
247
248 my $queue = $ctx->queue ($dev);
249
250 # now attach an opencl image2d object to the opengl texture
251 my $tex = $ctx->gl_texture2d (OpenCL::MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, $texid);
252
253 # now the boring opencl code
254 my $src = <<EOF;
255 kernel void
256 juliatunnel (write_only image2d_t img, float time)
257 {
258 int2 xy = (int2)(get_global_id (0), get_global_id (1));
259 float2 p = convert_float2 (xy) / $S.f * 2.f - 1.f;
260
261 float2 m = (float2)(1.f, p.y) / fabs (p.x); // tunnel
262 m.x = fabs (fmod (m.x + time * 0.05f, 4.f) - 2.f);
263
264 float2 z = m;
265 float2 c = (float2)(sin (time * 0.01133f), cos (time * 0.02521f));
266
267 for (int i = 0; i < 25 && dot (z, z) < 4.f; ++i) // standard julia
268 z = (float2)(z.x * z.x - z.y * z.y, 2.f * z.x * z.y) + c;
269
270 float3 colour = (float3)(z.x, z.y, atan2 (z.y, z.x));
271 write_imagef (img, xy, (float4)(colour * p.x * p.x, 1.));
272 }
273 EOF
274
275 my $prog = $ctx->build_program ($src);
276 my $kernel = $prog->kernel ("juliatunnel");
277
278 # program compiled, kernel ready, now draw and loop
279
280 for (my $time; ; ++$time) {
281 # acquire objects from opengl
282 $queue->acquire_gl_objects ([$tex]);
283
284 # configure and run our kernel
285 $kernel->setf ("mf", $tex, $time*2); # mf = memory object, float
286 $queue->nd_range_kernel ($kernel, undef, [$S, $S], undef);
287
288 # release objects to opengl again
289 $queue->release_gl_objects ([$tex]);
290
291 # wait
292 $queue->finish;
293
294 # now draw the texture, the defaults should be all right
295 glTexParameterf GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST;
296
297 glEnable GL_TEXTURE_2D;
298 glBegin GL_QUADS;
299 glTexCoord2f 0, 1; glVertex3i -1, -1, -1;
300 glTexCoord2f 0, 0; glVertex3i 1, -1, -1;
301 glTexCoord2f 1, 0; glVertex3i 1, 1, -1;
302 glTexCoord2f 1, 1; glVertex3i -1, 1, -1;
303 glEnd;
304
305 glXSwapBuffers;
306
307 select undef, undef, undef, 1/60;
308 }
309
310=head2 How to modify the previous example to not rely on GL sharing.
311
312For those poor souls with only a sucky CPU OpenCL implementation, you
313currently have to read the image into some perl scalar, and then modify a
314texture or use glDrawPixels or so).
315
316First, when you don't need gl sharing, you can create the context much simpler:
317
318 $ctx = $platform->context (undef, [$dev])
319
320To use a texture, you would modify the above example by creating an
321OpenCL::Image manually instead of deriving it from a texture:
322
323 my $tex = $ctx->image2d (OpenCL::MEM_WRITE_ONLY, OpenCL::RGBA, OpenCL::UNORM_INT8, $S, $S);
324
325And in the darw loop, intead of acquire_gl_objects/release_gl_objects, you
326would read the image2d after the kernel has written it:
327
328 $queue->read_image ($tex, 0, 0, 0, 0, $S, $S, 1, 0, 0, my $data);
329
330And then you would upload the pixel data to the texture (or use glDrawPixels):
331
332 glTexSubImage2D_s GL_TEXTURE_2D, 0, 0, 0, $S, $S, GL_RGBA, GL_UNSIGNED_BYTE, $data;
333
334The fully modified example can be found at
335L<http://cvs.schmorp.de/OpenCL/examples/juliaflight-nosharing>.
215 336
216=head1 DOCUMENTATION 337=head1 DOCUMENTATION
217 338
218=head2 BASIC CONVENTIONS 339=head2 BASIC CONVENTIONS
219 340
… …
226 347
227=item * Object lifetime managament is automatic - there is no need 348=item * Object lifetime managament is automatic - there is no need
228to free objects explicitly (C<clReleaseXXX>), the release function 349to free objects explicitly (C<clReleaseXXX>), the release function
229is called automatically once all Perl references to it go away. 350is called automatically once all Perl references to it go away.
230 351
231=item * OpenCL uses CamelCase for function names (e.g. C<clGetPlatformIDs>, C<clGetPlatformInfo>), 352=item * OpenCL uses CamelCase for function names
353(e.g. C<clGetPlatformIDs>, C<clGetPlatformInfo>), while this module
232while this module uses underscores as word separator and often leaves out 354uses underscores as word separator and often leaves out prefixes
233prefixes (C<OpenCL::platforms>, C<< $platform->info >>). 355(C<OpenCL::platforms>, C<< $platform->info >>).
234 356
235=item * OpenCL often specifies fixed vector function arguments as short 357=item * OpenCL often specifies fixed vector function arguments as short
236arrays (C<size_t origin[3]>), while this module explicitly expects the 358arrays (C<size_t origin[3]>), while this module explicitly expects the
237components as separate arguments (C<$orig_x, $orig_y, $orig_z>) in 359components as separate arguments (C<$orig_x, $orig_y, $orig_z>) in
238function calls. 360function calls.
… …
240=item * Structures are often specified by flattening out their components 362=item * Structures are often specified by flattening out their components
241as with short vectors, and returned as arrayrefs. 363as with short vectors, and returned as arrayrefs.
242 364
243=item * When enqueuing commands, the wait list is specified by adding 365=item * When enqueuing commands, the wait list is specified by adding
244extra arguments to the function - anywhere a C<$wait_events...> argument 366extra arguments to the function - anywhere a C<$wait_events...> argument
245is documented this can be any number of event objects. 367is documented this can be any number of event objects. As an extsnion
368implemented by this module, C<undef> values will be ignored in the event
369list.
246 370
247=item * When enqueuing commands, if the enqueue method is called in void 371=item * When enqueuing commands, if the enqueue method is called in void
248context, no event is created. In all other contexts an event is returned 372context, no event is created. In all other contexts an event is returned
249by the method. 373by the method.
250 374
… …
270 ulong IV - Q 394 ulong IV - Q
271 float NV float f 395 float NV float f
272 half IV ushort S 396 half IV ushort S
273 double NV double d 397 double NV double d
274 398
399=head2 GLX SUPPORT
400
401Due to the sad state that OpenGL support is in in Perl (mostly the OpenGL
402module, which has little to no documentation and has little to no support
403for glX), this module, as a special extension, treats context creation
404properties C<OpenCL::GLX_DISPLAY_KHR> and C<OpenCL::GL_CONTEXT_KHR>
405specially: If either or both of these are C<undef>, then the OpenCL
406module tries to dynamically resolve C<glXGetCurrentDisplay> and
407C<glXGetCurrentContext>, call these functions and use their return values
408instead.
409
410For this to work, the OpenGL library must be loaded, a GLX context must
411have been created and be made current, and C<dlsym> must be available and
412capable of finding the function via C<RTLD_DEFAULT>.
413
414=head2 EVENT SYSTEM
415
416OpenCL can generate a number of (potentially) asynchronous events, for
417example, after compiling a program, to signal a context-related error or,
418perhaps most important, to signal completion of queued jobs (by setting
419callbacks on OpenCL::Event objects).
420
421To facilitate this, this module maintains an event queue - each
422time an asynchronous event happens, it is queued, and perl will be
423interrupted. This is implemented via the L<Async::Interrupt> module. In
424addition, this module has L<AnyEvent> support, so it can seamlessly
425integrate itself into many event loops.
426
427Since this module is a bit hard to understand, here are some case examples:
428
429=head3 Don't use callbacks.
430
431When your program never uses any callbacks, then there will never be any
432notifications you need to take care of, and therefore no need to worry
433about all this.
434
435You can achieve a great deal by explicitly waiting for events, or using
436barriers and flush calls. In many programs, there is no need at all to
437tinker with asynchronous events.
438
439=head3 Use AnyEvent
440
441This module automatically registers a watcher that invokes all outstanding
442event callbacks when AnyEvent is initialised (and block asynchronous
443interruptions). Using this mode of operations is the safest and most
444recommended one.
445
446To use this, simply use AnyEvent and this module normally, make sure you
447have an event loop running:
448
449 use Gtk2 -init;
450 use AnyEvent;
451
452 # initialise AnyEvent, by creating a watcher, or:
453 AnyEvent::detect;
454
455 my $e = $queue->marker;
456 $e->cb (sub {
457 warn "opencl is finished\n";
458 })
459
460 main Gtk2;
461
462Note that this module will not initialise AnyEvent for you. Before
463AnyEvent is initialised, the module will asynchronously interrupt perl
464instead. To avoid any surprises, it's best to explicitly initialise
465AnyEvent.
466
467You can temporarily enable asynchronous interruptions (see next paragraph)
468by calling C<$OpenCL::INTERRUPT->unblock> and disable them again by
469calling C<$OpenCL::INTERRUPT->block>.
470
471=head3 Let yourself be interrupted at any time
472
473This mode is the default unless AnyEvent is loaded and initialised. In
474this mode, OpenCL asynchronously interrupts a running perl program. The
475emphasis is on both I<asynchronously> and I<running> here.
476
477Asynchronously means that perl might execute your callbacks at any
478time. For example, in the following code (I<THAT YOU SHOULD NOT COPY>),
479the C<until> loop following the marker call will be interrupted by the
480callback:
481
482 my $e = $queue->marker;
483 my $flag;
484 $e->cb (sub { $flag = 1 });
485 1 until $flag;
486 # $flag is now 1
487
488The reason why you shouldn't blindly copy the above code is that
489busy waiting is a really really bad thing, and really really bad for
490performance.
491
492While at first this asynchronous business might look exciting, it can be
493really hard, because you need to be prepared for the callback code to be
494executed at any time, which limits the amount of things the callback code
495can do safely.
496
497This can be mitigated somewhat by using C<<
498$OpenCL::INTERRUPT->scope_block >> (see the L<Async::Interrupt>
499documentation for details).
500
501The other problem is that your program must be actively I<running> to be
502interrupted. When you calculate stuff, your program is running. When you
503hang in some C functions or other block execution (by calling C<sleep>,
504C<select>, running an event loop and so on), your program is waiting, not
505running.
506
507One way around that would be to attach a read watcher to your event loop,
508listening for events on C<< $OpenCL::INTERRUPT->pipe_fileno >>, using a
509dummy callback (C<sub { }>) to temporarily execute some perl code.
510
511That is then awfully close to using the built-in AnyEvent support above,
512though, so consider that one instead.
513
514=head3 Be creative
515
516OpenCL exports the L<Async::Interrupt> object it uses in the global
517variable C<$OpenCL::INTERRUPT>. You can configure it in any way you like.
518
519So if you want to feel like a real pro, err, wait, if you feel no risk
520menas no fun, you can experiment by implementing your own mode of
521operations.
522
523=cut
524
525package OpenCL;
526
527use common::sense;
528use Carp ();
529use Async::Interrupt ();
530
531our $POLL_FUNC; # set by XS
532
533BEGIN {
534 our $VERSION = '0.98';
535
536 require XSLoader;
537 XSLoader::load (__PACKAGE__, $VERSION);
538
539 @OpenCL::Platform::ISA =
540 @OpenCL::Device::ISA =
541 @OpenCL::Context::ISA =
542 @OpenCL::Queue::ISA =
543 @OpenCL::Memory::ISA =
544 @OpenCL::Sampler::ISA =
545 @OpenCL::Program::ISA =
546 @OpenCL::Kernel::ISA =
547 @OpenCL::Event::ISA = OpenCL::Object::;
548
549 @OpenCL::Buffer::ISA =
550 @OpenCL::Image::ISA = OpenCL::Memory::;
551
552 @OpenCL::BufferObj::ISA = OpenCL::Buffer::;
553
554 @OpenCL::Image2D::ISA =
555 @OpenCL::Image3D::ISA =
556 @OpenCL::Image2DArray::ISA =
557 @OpenCL::Image1D::ISA =
558 @OpenCL::Image1DArray::ISA =
559 @OpenCL::Image1DBuffer::ISA = OpenCL::Image::;
560
561 @OpenCL::UserEvent::ISA = OpenCL::Event::;
562
563 @OpenCL::MappedBuffer::ISA =
564 @OpenCL::MappedImage::ISA = OpenCL::Mapped::;
565}
566
275=head2 THE OpenCL PACKAGE 567=head2 THE OpenCL PACKAGE
276 568
277=over 4 569=over 4
278 570
279=item $int = OpenCL::errno 571=item $int = OpenCL::errno
280 572
281The last error returned by a function - it's only valid after an error occured 573The last error returned by a function - it's only valid after an error occured
282and before calling another OpenCL function. 574and before calling another OpenCL function.
283 575
284=item $str = OpenCL::err2str $errval 576=item $str = OpenCL::err2str [$errval]
285 577
286Comverts an error value into a human readable string. 578Converts an error value into a human readable string. IF no error value is
579given, then the last error will be used (as returned by OpenCL::errno).
287 580
288=item $str = OpenCL::enum2str $enum 581=item $str = OpenCL::enum2str $enum
289 582
290Converts most enum values (inof parameter names, image format constants, 583Converts most enum values (of parameter names, image format constants,
291object types, addressing and filter modes, command types etc.) into a 584object types, addressing and filter modes, command types etc.) into a
292human readbale string. When confronted with some random integer it can be 585human readable string. When confronted with some random integer it can be
293very helpful to pass it through this function to maybe get some readable 586very helpful to pass it through this function to maybe get some readable
294string out of it. 587string out of it.
295 588
296=item @platforms = OpenCL::platforms 589=item @platforms = OpenCL::platforms
297 590
298Returns all available OpenCL::Platform objects. 591Returns all available OpenCL::Platform objects.
299 592
300L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformIDs.html> 593L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformIDs.html>
301 594
302=item $ctx = OpenCL::context_from_type $properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $notify = undef 595=item $ctx = OpenCL::context_from_type $properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $callback->($err, $pvt) = $print_stderr
303 596
304Tries to create a context from a default device and platform - never worked for me. 597Tries to create a context from a default device and platform type - never worked for me.
305 598
306L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html> 599L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html>
307 600
601=item $ctx = OpenCL::context $properties, \@devices, $callback->($err, $pvt) = $print_stderr)
602
603Create a new OpenCL::Context object using the given device object(s). This
604function isn't implemented yet, use C<< $platform->context >> instead.
605
606L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContext.html>
607
308=item OpenCL::wait_for_events $wait_events... 608=item OpenCL::wait_for_events $wait_events...
309 609
310Waits for all events to complete. 610Waits for all events to complete.
311 611
312L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html> 612L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html>
313 613
614=item OpenCL::poll
615
616Checks if there are any outstanding events (see L<EVENT SYSTEM>) and
617invokes their callbacks.
618
619=item $OpenCL::INTERRUPT
620
621The L<Async::Interrupt> object used to signal asynchronous events (see
622L<EVENT SYSTEM>).
623
624=cut
625
626our $INTERRUPT = new Async::Interrupt c_cb => [$POLL_FUNC, 0];
627
628&_eq_initialise ($INTERRUPT->signal_func);
629
630=item $OpenCL::WATCHER
631
632The L<AnyEvent> watcher object used to watch for asynchronous events (see
633L<EVENT SYSTEM>). This variable is C<undef> until L<AnyEvent> has been
634loaded I<and> initialised (e.g. by calling C<AnyEvent::detect>).
635
636=cut
637
638our $WATCHER;
639
640sub _init_anyevent {
641 $INTERRUPT->block;
642 $WATCHER = AE::io ($INTERRUPT->pipe_fileno, 0, sub { $INTERRUPT->handle });
643}
644
645if (defined $AnyEvent::MODEL) {
646 _init_anyevent;
647} else {
648 push @AnyEvent::post_detect, \&_init_anyevent;
649}
650
314=back 651=back
315 652
653=head2 THE OpenCL::Object CLASS
654
655This is the base class for all objects in the OpenCL module. The only
656method it implements is the C<id> method, which is only useful if you want
657to interface to OpenCL on the C level.
658
659=over 4
660
661=item $iv = $obj->id
662
663OpenCL objects are represented by pointers or integers on the C level. If
664you want to interface to an OpenCL object directly on the C level, then
665you need this value, which is returned by this method. You should use an
666C<IV> type in your code and cast that to the correct type.
667
668=cut
669
670sub OpenCL::Object::id {
671 ref $_[0] eq "SCALAR"
672 ? ${ $_[0] }
673 : $_[0][0]
674}
675
676=back
677
316=head2 THE OpenCL::Platform CLASS 678=head2 THE OpenCL::Platform CLASS
317 679
318=over 4 680=over 4
319 681
320=item $packed_value = $platform->info ($name)
321
322Calls C<clGetPlatformInfo> and returns the packed, raw value - for
323strings, this will be the string, for other values you probably need to
324use the correct C<unpack>. This might get improved in the future. Hopefully.
325
326L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformInfo.html>
327
328=item @devices = $platform->devices ($type = OpenCL::DEVICE_TYPE_ALL) 682=item @devices = $platform->devices ($type = OpenCL::DEVICE_TYPE_ALL)
329 683
330Returns a list of matching OpenCL::Device objects. 684Returns a list of matching OpenCL::Device objects.
331 685
332=item $ctx = $platform->context_from_type ($properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $notify = undef) 686=item $ctx = $platform->context_from_type ($properties, $type = OpenCL::DEVICE_TYPE_DEFAULT, $callback->($err, $pvt) = $print_stderr)
333 687
334Tries to create a context. Never worked for me, and you need devices explitly anyway. 688Tries to create a context. Never worked for me, and you need devices explicitly anyway.
335 689
336L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html> 690L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContextFromType.html>
337 691
338=item $ctx = $device->context ($properties = undef, @$devices, $notify = undef) 692=item $ctx = $platform->context ($properties, \@devices, $callback->($err, $pvt) = $print_stderr)
339 693
340Create a new OpenCL::Context object using the given device object(s)- a 694Create a new OpenCL::Context object using the given device object(s)- a
341CL_CONTEXT_PLATFORM property is supplied automatically. 695CL_CONTEXT_PLATFORM property is supplied automatically.
342 696
343L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContext.html> 697L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateContext.html>
344 698
699=item $packed_value = $platform->info ($name)
700
701Calls C<clGetPlatformInfo> and returns the packed, raw value - for
702strings, this will be the string (possibly including terminating \0), for
703other values you probably need to use the correct C<unpack>.
704
705It's best to avoid this method and use one of the following convenience
706wrappers.
707
708L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformInfo.html>
709
710=item $platform->unload_compiler
711
712Attempts to unload the compiler for this platform, for endless
713profit. Does nothing on OpenCL 1.1.
714
715L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clUnloadPlatformCompiler.html>
716
717=for gengetinfo begin platform
718
719=item $string = $platform->profile
720
721Calls C<clGetPlatformInfo> with C<CL_PLATFORM_PROFILE> and returns the result.
722
723=item $string = $platform->version
724
725Calls C<clGetPlatformInfo> with C<CL_PLATFORM_VERSION> and returns the result.
726
727=item $string = $platform->name
728
729Calls C<clGetPlatformInfo> with C<CL_PLATFORM_NAME> and returns the result.
730
731=item $string = $platform->vendor
732
733Calls C<clGetPlatformInfo> with C<CL_PLATFORM_VENDOR> and returns the result.
734
735=item $string = $platform->extensions
736
737Calls C<clGetPlatformInfo> with C<CL_PLATFORM_EXTENSIONS> and returns the result.
738
739=for gengetinfo end platform
740
345=back 741=back
346 742
347=head2 THE OpenCL::Device CLASS 743=head2 THE OpenCL::Device CLASS
348 744
349=over 4 745=over 4
… …
352 748
353See C<< $platform->info >> for details. 749See C<< $platform->info >> for details.
354 750
355L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetDeviceInfo.html> 751L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetDeviceInfo.html>
356 752
753=for gengetinfo begin device
754
755=item $device_type = $device->type
756
757Calls C<clGetDeviceInfo> with C<CL_DEVICE_TYPE> and returns the result.
758
759=item $uint = $device->vendor_id
760
761Calls C<clGetDeviceInfo> with C<CL_DEVICE_VENDOR_ID> and returns the result.
762
763=item $uint = $device->max_compute_units
764
765Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_COMPUTE_UNITS> and returns the result.
766
767=item $uint = $device->max_work_item_dimensions
768
769Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS> and returns the result.
770
771=item $int = $device->max_work_group_size
772
773Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_WORK_GROUP_SIZE> and returns the result.
774
775=item @ints = $device->max_work_item_sizes
776
777Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_WORK_ITEM_SIZES> and returns the result.
778
779=item $uint = $device->preferred_vector_width_char
780
781Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR> and returns the result.
782
783=item $uint = $device->preferred_vector_width_short
784
785Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT> and returns the result.
786
787=item $uint = $device->preferred_vector_width_int
788
789Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT> and returns the result.
790
791=item $uint = $device->preferred_vector_width_long
792
793Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG> and returns the result.
794
795=item $uint = $device->preferred_vector_width_float
796
797Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT> and returns the result.
798
799=item $uint = $device->preferred_vector_width_double
800
801Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE> and returns the result.
802
803=item $uint = $device->max_clock_frequency
804
805Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_CLOCK_FREQUENCY> and returns the result.
806
807=item $bitfield = $device->address_bits
808
809Calls C<clGetDeviceInfo> with C<CL_DEVICE_ADDRESS_BITS> and returns the result.
810
811=item $uint = $device->max_read_image_args
812
813Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_READ_IMAGE_ARGS> and returns the result.
814
815=item $uint = $device->max_write_image_args
816
817Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_WRITE_IMAGE_ARGS> and returns the result.
818
819=item $ulong = $device->max_mem_alloc_size
820
821Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_MEM_ALLOC_SIZE> and returns the result.
822
823=item $int = $device->image2d_max_width
824
825Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE2D_MAX_WIDTH> and returns the result.
826
827=item $int = $device->image2d_max_height
828
829Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE2D_MAX_HEIGHT> and returns the result.
830
831=item $int = $device->image3d_max_width
832
833Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE3D_MAX_WIDTH> and returns the result.
834
835=item $int = $device->image3d_max_height
836
837Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE3D_MAX_HEIGHT> and returns the result.
838
839=item $int = $device->image3d_max_depth
840
841Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE3D_MAX_DEPTH> and returns the result.
842
843=item $uint = $device->image_support
844
845Calls C<clGetDeviceInfo> with C<CL_DEVICE_IMAGE_SUPPORT> and returns the result.
846
847=item $int = $device->max_parameter_size
848
849Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_PARAMETER_SIZE> and returns the result.
850
851=item $uint = $device->max_samplers
852
853Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_SAMPLERS> and returns the result.
854
855=item $uint = $device->mem_base_addr_align
856
857Calls C<clGetDeviceInfo> with C<CL_DEVICE_MEM_BASE_ADDR_ALIGN> and returns the result.
858
859=item $uint = $device->min_data_type_align_size
860
861Calls C<clGetDeviceInfo> with C<CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE> and returns the result.
862
863=item $device_fp_config = $device->single_fp_config
864
865Calls C<clGetDeviceInfo> with C<CL_DEVICE_SINGLE_FP_CONFIG> and returns the result.
866
867=item $device_mem_cache_type = $device->global_mem_cache_type
868
869Calls C<clGetDeviceInfo> with C<CL_DEVICE_GLOBAL_MEM_CACHE_TYPE> and returns the result.
870
871=item $uint = $device->global_mem_cacheline_size
872
873Calls C<clGetDeviceInfo> with C<CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE> and returns the result.
874
875=item $ulong = $device->global_mem_cache_size
876
877Calls C<clGetDeviceInfo> with C<CL_DEVICE_GLOBAL_MEM_CACHE_SIZE> and returns the result.
878
879=item $ulong = $device->global_mem_size
880
881Calls C<clGetDeviceInfo> with C<CL_DEVICE_GLOBAL_MEM_SIZE> and returns the result.
882
883=item $ulong = $device->max_constant_buffer_size
884
885Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE> and returns the result.
886
887=item $uint = $device->max_constant_args
888
889Calls C<clGetDeviceInfo> with C<CL_DEVICE_MAX_CONSTANT_ARGS> and returns the result.
890
891=item $device_local_mem_type = $device->local_mem_type
892
893Calls C<clGetDeviceInfo> with C<CL_DEVICE_LOCAL_MEM_TYPE> and returns the result.
894
895=item $ulong = $device->local_mem_size
896
897Calls C<clGetDeviceInfo> with C<CL_DEVICE_LOCAL_MEM_SIZE> and returns the result.
898
899=item $boolean = $device->error_correction_support
900
901Calls C<clGetDeviceInfo> with C<CL_DEVICE_ERROR_CORRECTION_SUPPORT> and returns the result.
902
903=item $int = $device->profiling_timer_resolution
904
905Calls C<clGetDeviceInfo> with C<CL_DEVICE_PROFILING_TIMER_RESOLUTION> and returns the result.
906
907=item $boolean = $device->endian_little
908
909Calls C<clGetDeviceInfo> with C<CL_DEVICE_ENDIAN_LITTLE> and returns the result.
910
911=item $boolean = $device->available
912
913Calls C<clGetDeviceInfo> with C<CL_DEVICE_AVAILABLE> and returns the result.
914
915=item $boolean = $device->compiler_available
916
917Calls C<clGetDeviceInfo> with C<CL_DEVICE_COMPILER_AVAILABLE> and returns the result.
918
919=item $device_exec_capabilities = $device->execution_capabilities
920
921Calls C<clGetDeviceInfo> with C<CL_DEVICE_EXECUTION_CAPABILITIES> and returns the result.
922
923=item $command_queue_properties = $device->properties
924
925Calls C<clGetDeviceInfo> with C<CL_DEVICE_QUEUE_PROPERTIES> and returns the result.
926
927=item $ = $device->platform
928
929Calls C<clGetDeviceInfo> with C<CL_DEVICE_PLATFORM> and returns the result.
930
931=item $string = $device->name
932
933Calls C<clGetDeviceInfo> with C<CL_DEVICE_NAME> and returns the result.
934
935=item $string = $device->vendor
936
937Calls C<clGetDeviceInfo> with C<CL_DEVICE_VENDOR> and returns the result.
938
939=item $string = $device->driver_version
940
941Calls C<clGetDeviceInfo> with C<CL_DRIVER_VERSION> and returns the result.
942
943=item $string = $device->profile
944
945Calls C<clGetDeviceInfo> with C<CL_DEVICE_PROFILE> and returns the result.
946
947=item $string = $device->version
948
949Calls C<clGetDeviceInfo> with C<CL_DEVICE_VERSION> and returns the result.
950
951=item $string = $device->extensions
952
953Calls C<clGetDeviceInfo> with C<CL_DEVICE_EXTENSIONS> and returns the result.
954
955=item $uint = $device->preferred_vector_width_half
956
957Calls C<clGetDeviceInfo> with C<CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF> and returns the result.
958
959=item $uint = $device->native_vector_width_char
960
961Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR> and returns the result.
962
963=item $uint = $device->native_vector_width_short
964
965Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT> and returns the result.
966
967=item $uint = $device->native_vector_width_int
968
969Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_INT> and returns the result.
970
971=item $uint = $device->native_vector_width_long
972
973Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG> and returns the result.
974
975=item $uint = $device->native_vector_width_float
976
977Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT> and returns the result.
978
979=item $uint = $device->native_vector_width_double
980
981Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE> and returns the result.
982
983=item $uint = $device->native_vector_width_half
984
985Calls C<clGetDeviceInfo> with C<CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF> and returns the result.
986
987=item $device_fp_config = $device->double_fp_config
988
989Calls C<clGetDeviceInfo> with C<CL_DEVICE_DOUBLE_FP_CONFIG> and returns the result.
990
991=item $device_fp_config = $device->half_fp_config
992
993Calls C<clGetDeviceInfo> with C<CL_DEVICE_HALF_FP_CONFIG> and returns the result.
994
995=item $boolean = $device->host_unified_memory
996
997Calls C<clGetDeviceInfo> with C<CL_DEVICE_HOST_UNIFIED_MEMORY> and returns the result.
998
999=item $device = $device->parent_device_ext
1000
1001Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARENT_DEVICE_EXT> and returns the result.
1002
1003=item @device_partition_property_exts = $device->partition_types_ext
1004
1005Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARTITION_TYPES_EXT> and returns the result.
1006
1007=item @device_partition_property_exts = $device->affinity_domains_ext
1008
1009Calls C<clGetDeviceInfo> with C<CL_DEVICE_AFFINITY_DOMAINS_EXT> and returns the result.
1010
1011=item $uint = $device->reference_count_ext
1012
1013Calls C<clGetDeviceInfo> with C<CL_DEVICE_REFERENCE_COUNT_EXT> and returns the result.
1014
1015=item @device_partition_property_exts = $device->partition_style_ext
1016
1017Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARTITION_STYLE_EXT> and returns the result.
1018
1019=for gengetinfo end device
1020
357=back 1021=back
358 1022
359=head2 THE OpenCL::Context CLASS 1023=head2 THE OpenCL::Context CLASS
360 1024
361=over 4 1025=over 4
362 1026
363=item $packed_value = $ctx->info ($name) 1027=item $prog = $ctx->build_program ($program, $options = "")
364 1028
365See C<< $platform->info >> for details. 1029This convenience function tries to build the program on all devices in
1030the context. If the build fails, then the function will C<croak> with the
1031build log. Otherwise ti returns the program object.
366 1032
367L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetContextInfo.html> 1033The C<$program> can either be a C<OpenCL::Program> object or a string
1034containing the program. In the latter case, a program objetc will be
1035created automatically.
1036
1037=cut
1038
1039sub OpenCL::Context::build_program {
1040 my ($self, $prog, $options) = @_;
1041
1042 $prog = $self->program_with_source ($prog)
1043 unless ref $prog;
1044
1045 eval { $prog->build (undef, $options); 1 }
1046 or errno == BUILD_PROGRAM_FAILURE
1047 or errno == INVALID_BINARY # workaround nvidia bug
1048 or Carp::croak "OpenCL::Context->build_program: " . err2str;
1049
1050 # we check status for all devices
1051 for my $dev ($self->devices) {
1052 $prog->build_status ($dev) == BUILD_SUCCESS
1053 or Carp::croak "Building OpenCL program for device '" . $dev->name . "' failed:\n"
1054 . $prog->build_log ($dev);
1055 }
1056
1057 $prog
1058}
368 1059
369=item $queue = $ctx->queue ($device, $properties) 1060=item $queue = $ctx->queue ($device, $properties)
370 1061
371Create a new OpenCL::Queue object from the context and the given device. 1062Create a new OpenCL::Queue object from the context and the given device.
372 1063
373L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateCommandQueue.html> 1064L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateCommandQueue.html>
374 1065
1066Example: create an out-of-order queue.
1067
1068 $queue = $ctx->queue ($device, OpenCL::QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE);
1069
375=item $ev = $ctx->user_event 1070=item $ev = $ctx->user_event
376 1071
377Creates a new OpenCL::UserEvent object. 1072Creates a new OpenCL::UserEvent object.
378 1073
379L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateUserEvent.html> 1074L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateUserEvent.html>
380 1075
381=item $buf = $ctx->buffer ($flags, $len) 1076=item $buf = $ctx->buffer ($flags, $len)
382 1077
383Creates a new OpenCL::Buffer object with the given flags and octet-size. 1078Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object with the
1079given flags and octet-size.
384 1080
385L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateBuffer.html> 1081L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateBuffer.html>
386 1082
387=item $buf = $ctx->buffer_sv ($flags, $data) 1083=item $buf = $ctx->buffer_sv ($flags, $data)
388 1084
389Creates a new OpenCL::Buffer object and initialise it with the given data values. 1085Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object and
1086initialise it with the given data values.
1087
1088=item $img = $ctx->image ($self, $flags, $channel_order, $channel_type, $type, $width, $height, $depth = 0, $array_size = 0, $row_pitch = 0, $slice_pitch = 0, $num_mip_level = 0, $num_samples = 0, $*data = &PL_sv_undef)
1089
1090Creates a new OpenCL::Image object and optionally initialises it with
1091the given data values.
1092
1093L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateImage.html>
390 1094
391=item $img = $ctx->image2d ($flags, $channel_order, $channel_type, $width, $height, $row_pitch = 0, $data = undef) 1095=item $img = $ctx->image2d ($flags, $channel_order, $channel_type, $width, $height, $row_pitch = 0, $data = undef)
392 1096
393Creates a new OpenCL::Image2D object and optionally initialises it with the given data values. 1097Creates a new OpenCL::Image2D object and optionally initialises it with
1098the given data values.
394 1099
395L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateImage2D.html> 1100L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateImage2D.html>
396 1101
397=item $img = $ctx->image3d ($flags, $channel_order, $channel_type, $width, $height, $depth, $row_pitch = 0, $slice_pitch = 0, $data = undef) 1102=item $img = $ctx->image3d ($flags, $channel_order, $channel_type, $width, $height, $depth, $row_pitch = 0, $slice_pitch = 0, $data = undef)
398 1103
399Creates a new OpenCL::Image3D object and optionally initialises it with the given data values. 1104Creates a new OpenCL::Image3D object and optionally initialises it with
1105the given data values.
400 1106
401L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateImage3D.html> 1107L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateImage3D.html>
1108
1109=item $buffer = $ctx->gl_buffer ($flags, $bufobj)
1110
1111Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object that refers to the given
1112OpenGL buffer object.
1113
1114http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLBuffer.html
1115
1116=item $img = $ctx->gl_texture ($flags, $target, $miplevel, $texture)
1117
1118Creates a new OpenCL::Image object that refers to the given OpenGL
1119texture object or buffer.
1120
1121http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateFromGLTexture.html
1122
1123=item $img = $ctx->gl_texture2d ($flags, $target, $miplevel, $texture)
1124
1125Creates a new OpenCL::Image2D object that refers to the given OpenGL
11262D texture object.
1127
1128http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture2D.html
1129
1130=item $img = $ctx->gl_texture3d ($flags, $target, $miplevel, $texture)
1131
1132Creates a new OpenCL::Image3D object that refers to the given OpenGL
11333D texture object.
1134
1135http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture3D.html
1136
1137=item $ctx->gl_renderbuffer ($flags, $renderbuffer)
1138
1139Creates a new OpenCL::Image2D object that refers to the given OpenGL
1140render buffer.
1141
1142http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLRenderbuffer.html
402 1143
403=item @formats = $ctx->supported_image_formats ($flags, $image_type) 1144=item @formats = $ctx->supported_image_formats ($flags, $image_type)
404 1145
405Returns a list of matching image formats - each format is an arrayref with 1146Returns a list of matching image formats - each format is an arrayref with
406two values, $channel_order and $channel_type, in it. 1147two values, $channel_order and $channel_type, in it.
… …
417 1158
418Creates a new OpenCL::Program object from the given source code. 1159Creates a new OpenCL::Program object from the given source code.
419 1160
420L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateProgramWithSource.html> 1161L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateProgramWithSource.html>
421 1162
1163=item $packed_value = $ctx->info ($name)
1164
1165See C<< $platform->info >> for details.
1166
1167L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetContextInfo.html>
1168
1169=for gengetinfo begin context
1170
1171=item $uint = $context->reference_count
1172
1173Calls C<clGetContextInfo> with C<CL_CONTEXT_REFERENCE_COUNT> and returns the result.
1174
1175=item @devices = $context->devices
1176
1177Calls C<clGetContextInfo> with C<CL_CONTEXT_DEVICES> and returns the result.
1178
1179=item @property_ints = $context->properties
1180
1181Calls C<clGetContextInfo> with C<CL_CONTEXT_PROPERTIES> and returns the result.
1182
1183=item $uint = $context->num_devices
1184
1185Calls C<clGetContextInfo> with C<CL_CONTEXT_NUM_DEVICES> and returns the result.
1186
1187=for gengetinfo end context
1188
422=back 1189=back
423 1190
424=head2 THE OpenCL::Queue CLASS 1191=head2 THE OpenCL::Queue CLASS
425 1192
426An OpenCL::Queue represents an execution queue for OpenCL. You execute 1193An OpenCL::Queue represents an execution queue for OpenCL. You execute
427requests by calling their respective C<enqueue_xxx> method and waitinf for 1194requests by calling their respective method and waiting for it to complete
428it to complete in some way. 1195in some way.
429 1196
430All the enqueue methods return an event object that can be used to wait 1197Most methods that enqueue some request return an event object that can
431for completion, unless the method is called in void context, in which case 1198be used to wait for completion (optionally using a callback), unless
432no event object is created. 1199the method is called in void context, in which case no event object is
1200created.
433 1201
434They also allow you to specify any number of other event objects that this 1202They also allow you to specify any number of other event objects that this
435request has to wait for before it starts executing, by simply passing the 1203request has to wait for before it starts executing, by simply passing the
436event objects as extra parameters to the enqueue methods. 1204event objects as extra parameters to the enqueue methods. To simplify
1205program design, this module ignores any C<undef> values in the list of
1206events. This makes it possible to code operations such as this, without
1207having to put a valid event object into C<$event> first:
1208
1209 $event = $queue->xxx (..., $event);
437 1210
438Queues execute in-order by default, without any parallelism, so in most 1211Queues execute in-order by default, without any parallelism, so in most
439cases (i.e. you use only one queue) it's not necessary to wait for or 1212cases (i.e. you use only one queue) it's not necessary to wait for or
440create event objects. 1213create event objects, althoguh an our of order queue is often a bit
1214faster.
441 1215
442=over 4 1216=over 4
443 1217
444=item $packed_value = $ctx->info ($name)
445
446See C<< $platform->info >> for details.
447
448L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetCommandQueueInfo.html>
449
450=item $ev = $queue->enqueue_read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...) 1218=item $ev = $queue->read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...)
451 1219
452Reads data from buffer into the given string. 1220Reads data from buffer into the given string.
453 1221
454L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBuffer.html> 1222L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBuffer.html>
455 1223
456=item $ev = $queue->enqueue_write_buffer ($buffer, $blocking, $offset, $data, $wait_events...) 1224=item $ev = $queue->write_buffer ($buffer, $blocking, $offset, $data, $wait_events...)
457 1225
458Writes data to buffer from the given string. 1226Writes data to buffer from the given string.
459 1227
460L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBuffer.html> 1228L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBuffer.html>
461 1229
462=item $ev = $queue->enqueue_copy_buffer ($src, $dst, $src_offset, $dst_offset, $len, $wait_events...) 1230=item $ev = $queue->copy_buffer ($src, $dst, $src_offset, $dst_offset, $len, $wait_events...)
463 1231
464L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBuffer.html> 1232L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBuffer.html>
465 1233
1234=item $ev = $queue->read_buffer_rect (OpenCL::Memory buf, cl_bool blocking, $buf_x, $buf_y, $buf_z, $host_x, $host_y, $host_z, $width, $height, $depth, $buf_row_pitch, $buf_slice_pitch, $host_row_pitch, $host_slice_pitch, $data, $wait_events...)
1235
1236http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBufferRect.html
1237
1238=item $ev = $queue->write_buffer_rect (OpenCL::Memory buf, cl_bool blocking, $buf_x, $buf_y, $buf_z, $host_x, $host_y, $host_z, $width, $height, $depth, $buf_row_pitch, $buf_slice_pitch, $host_row_pitch, $host_slice_pitch, $data, $wait_events...)
1239
1240http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBufferRect.html
1241
1242=item $ev = $queue->copy_buffer_to_image ($src_buffer, $dst_image, $src_offset, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $wait_events...)
1243
1244L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>
1245
466=item $ev = $queue->enqueue_read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...) 1246=item $ev = $queue->read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
1247
1248C<$row_pitch> (and C<$slice_pitch>) can be C<0>, in which case the OpenCL
1249module uses the image width (and height) to supply default values.
467 1250
468L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadImage.html> 1251L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadImage.html>
469 1252
470=item $ev = $queue->enqueue_write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...) 1253=item $ev = $queue->write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
471 1254
1255C<$row_pitch> (and C<$slice_pitch>) can be C<0>, in which case the OpenCL
1256module uses the image width (and height) to supply default values.
472L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteImage.html> 1257L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteImage.html>
473 1258
1259=item $ev = $queue->copy_image ($src_image, $dst_image, $src_x, $src_y, $src_z, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $wait_events...)
1260
1261L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html>
1262
1263=item $ev = $queue->copy_image_to_buffer ($src_image, $dst_image, $src_x, $src_y, $src_z, $width, $height, $depth, $dst_offset, $wait_events...)
1264
1265L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html>
1266
474=item $ev = $queue->enqueue_copy_buffer_rect ($src, $dst, $src_x, $src_y, $src_z, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $src_row_pitch, $src_slice_pitch, $dst_row_pitch, $dst_slice_pitch, $wait_event...) 1267=item $ev = $queue->copy_buffer_rect ($src, $dst, $src_x, $src_y, $src_z, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $src_row_pitch, $src_slice_pitch, $dst_row_pitch, $dst_slice_pitch, $wait_event...)
475 1268
476Yeah. 1269Yeah.
477 1270
478L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferRect.html>
479
480=item $ev = $queue->enqueue_copy_buffer_to_image ($src_buffer, $dst_image, $src_offset, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $wait_events...)
481
482L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>. 1271L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>.
483 1272
484=item $ev = $queue->enqueue_copy_image ($src_image, $dst_image, $src_x, $src_y, $src_z, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $wait_events...) 1273=item $ev = $queue->fill_buffer ($mem, $pattern, $offset, $size, ...)
485 1274
1275Fills the given buffer object with repeated applications of C<$pattern>,
1276starting at C<$offset> for C<$size> octets.
1277
1278L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillBuffer.html>
1279
1280=item $ev = $queue->fill_image ($img, $r, $g, $b, $a, $x, $y, $z, $width, $height, $depth, ...)
1281
1282Fills the given image area with the given rgba colour components. The
1283components are normally floating point values between C<0> and C<1>,
1284except when the image channel data type is a signe dor unsigned
1285unnormalised format, in which case the range is determined by the format.
1286
486L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html> 1287L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillImage.html>
487 1288
488=item $ev = $queue->enqueue_copy_image_to_buffer ($src_image, $dst_image, $src_x, $src_y, $src_z, $width, $height, $depth, $dst_offset, $wait_events...)
489
490L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html>
491
492=item $ev = $queue->enqueue_task ($kernel, $wait_events...) 1289=item $ev = $queue->task ($kernel, $wait_events...)
493 1290
494L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html> 1291L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html>
495 1292
496=item $ev = $queue->enqueue_nd_range_kernel ($kernel, @$global_work_offset, @$global_work_size, @$local_work_size, $wait_events...) 1293=item $ev = $queue->nd_range_kernel ($kernel, \@global_work_offset, \@global_work_size, \@local_work_size, $wait_events...)
497 1294
498Enqueues a kernel execution. 1295Enqueues a kernel execution.
499 1296
500@$global_work_size must be specified as a reference to an array of 1297\@global_work_size must be specified as a reference to an array of
501integers specifying the work sizes (element counts). 1298integers specifying the work sizes (element counts).
502 1299
503@$global_work_offset must be either C<undef> (in which case all offsets 1300\@global_work_offset must be either C<undef> (in which case all offsets
504are C<0>), or a reference to an array of work offsets, with the same number 1301are C<0>), or a reference to an array of work offsets, with the same number
505of elements as @$global_work_size. 1302of elements as \@global_work_size.
506 1303
507@$local_work_size must be either C<undef> (in which case the 1304\@local_work_size must be either C<undef> (in which case the
508implementation is supposed to choose good local work sizes), or a 1305implementation is supposed to choose good local work sizes), or a
509reference to an array of local work sizes, with the same number of 1306reference to an array of local work sizes, with the same number of
510elements as @$global_work_size. 1307elements as \@global_work_size.
511 1308
512L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html> 1309L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html>
513 1310
514=item $ev = $queue->enqueue_marker 1311=item $ev = $queue->acquire_gl_objects ([object, ...], $wait_events...)
515 1312
1313Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired
1314for subsequent OpenCL usage.
1315
516L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMarker.html> 1316L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueAcquireGLObjects.html>
517 1317
1318=item $ev = $queue->release_gl_objects ([object, ...], $wait_events...)
1319
1320Enqueues a list (an array-ref of OpenCL::Memory objects) to be released
1321for subsequent OpenGL usage.
1322
1323L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReleaseGLObjects.html>
1324
518=item $ev = $queue->enqueue_wait_for_events ($wait_events...) 1325=item $ev = $queue->wait_for_events ($wait_events...)
519 1326
520L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html> 1327L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html>
521 1328
522=item $queue->enqueue_barrier 1329=item $ev = $queue->marker ($wait_events...)
523 1330
1331L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMarkerWithWaitList.html>
1332
1333=item $ev = $queue->barrier ($wait_events...)
1334
524L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueBarrier.html> 1335L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueBarrierWithWaitList.html>
525 1336
526=item $queue->flush 1337=item $queue->flush
527 1338
528L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clFlush.html> 1339L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clFlush.html>
529 1340
530=item $queue->finish 1341=item $queue->finish
531 1342
532L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clFinish.html> 1343L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clFinish.html>
533 1344
1345=item $packed_value = $queue->info ($name)
1346
1347See C<< $platform->info >> for details.
1348
1349L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetCommandQueueInfo.html>
1350
1351=for gengetinfo begin command_queue
1352
1353=item $ctx = $command_queue->context
1354
1355Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_CONTEXT> and returns the result.
1356
1357=item $device = $command_queue->device
1358
1359Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_DEVICE> and returns the result.
1360
1361=item $uint = $command_queue->reference_count
1362
1363Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_REFERENCE_COUNT> and returns the result.
1364
1365=item $command_queue_properties = $command_queue->properties
1366
1367Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_PROPERTIES> and returns the result.
1368
1369=for gengetinfo end command_queue
1370
534=back 1371=back
535 1372
1373=head3 MEMORY MAPPED BUFFERS
1374
1375OpenCL allows you to map buffers and images to host memory (read: perl
1376scalars). This is done much like reading or copying a buffer, by enqueuing
1377a map or unmap operation on the command queue.
1378
1379The map operations return a C<OpenCL::Mapped> object - see L<THE
1380OpenCL::Mapped CLASS> section for details on what to do with these
1381objects.
1382
1383The object will be unmapped automatically when the mapped object is
1384destroyed (you can use a barrier to make sure the unmap has finished,
1385before using the buffer in a kernel), but you can also enqueue an unmap
1386operation manually.
1387
1388=over 4
1389
1390=item $mapped_buffer = $queue->map_buffer ($buf, $data, $blocking=1, $map_flags=OpenCL::MAP_READ|OpenCL::MAP_WRITE, $offset=0, $size=0, $wait_events...)
1391
1392Maps the given buffer into host memory and returns a C<OpenCL::MappedBuffer> object.
1393
1394L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMapBuffer.html>
1395
1396=item $mapped_image = $queue->map_image ($img, $data, $blocking=1, $map_flags=OpenCL::MAP_READ|OpenCL::MAP_WRITE, $x=0, $y=0, $z=0, $width=0, $height=0, $depth=0, $wait_events...)
1397
1398Maps the given image area into host memory and return a
1399C<OpenCL::MappedImage> object. Although there are default values for most
1400arguments, you currently have to specify all arguments, otherwise the call
1401will fail.
1402
1403L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMapImage.html>
1404
1405=item $ev = $queue->unmap ($mapped, $wait_events...)
1406
1407Unmaps the data from host memory. You must not call any methods that
1408modify the data, or modify the data scalar directly, after calling this
1409method.
1410
1411The mapped event object will always be passed as part of the
1412$wait_events. The mapped event object will be replaced by the new event
1413object that this request creates.
1414
1415=back
1416
536=head2 THE OpenCL::Memory CLASS 1417=head2 THE OpenCL::Memory CLASS
537 1418
538This the superclass of all memory objects - OpenCL::Buffer, OpenCL::Image, 1419This the superclass of all memory objects - OpenCL::Buffer, OpenCL::Image,
539OpenCL::Image2D and OpenCL::Image3D. The subclasses of this class 1420OpenCL::Image2D and OpenCL::Image3D.
540currently only exist to allow type-checking.
541 1421
542=over 4 1422=over 4
543 1423
544=item $packed_value = $memory->info ($name) 1424=item $packed_value = $memory->info ($name)
545 1425
546See C<< $platform->info >> for details. 1426See C<< $platform->info >> for details.
547 1427
548L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetMemObjectInfo.html> 1428L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetMemObjectInfo.html>
549 1429
1430=for gengetinfo begin mem
1431
1432=item $mem_object_type = $mem->type
1433
1434Calls C<clGetMemObjectInfo> with C<CL_MEM_TYPE> and returns the result.
1435
1436=item $mem_flags = $mem->flags
1437
1438Calls C<clGetMemObjectInfo> with C<CL_MEM_FLAGS> and returns the result.
1439
1440=item $int = $mem->size
1441
1442Calls C<clGetMemObjectInfo> with C<CL_MEM_SIZE> and returns the result.
1443
1444=item $ptr_value = $mem->host_ptr
1445
1446Calls C<clGetMemObjectInfo> with C<CL_MEM_HOST_PTR> and returns the result.
1447
1448=item $uint = $mem->map_count
1449
1450Calls C<clGetMemObjectInfo> with C<CL_MEM_MAP_COUNT> and returns the result.
1451
1452=item $uint = $mem->reference_count
1453
1454Calls C<clGetMemObjectInfo> with C<CL_MEM_REFERENCE_COUNT> and returns the result.
1455
1456=item $ctx = $mem->context
1457
1458Calls C<clGetMemObjectInfo> with C<CL_MEM_CONTEXT> and returns the result.
1459
1460=item $mem = $mem->associated_memobject
1461
1462Calls C<clGetMemObjectInfo> with C<CL_MEM_ASSOCIATED_MEMOBJECT> and returns the result.
1463
1464=item $int = $mem->offset
1465
1466Calls C<clGetMemObjectInfo> with C<CL_MEM_OFFSET> and returns the result.
1467
1468=for gengetinfo end mem
1469
1470=item ($type, $name) = $mem->gl_object_info
1471
1472Returns the OpenGL object type (e.g. OpenCL::GL_OBJECT_TEXTURE2D) and the
1473object "name" (e.g. the texture name) used to create this memory object.
1474
1475L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetGLObjectInfo.html>
1476
550=back 1477=back
551 1478
1479=head2 THE OpenCL::Buffer CLASS
1480
1481This is a subclass of OpenCL::Memory, and the superclass of
1482OpenCL::BufferObj. Its purpose is simply to distinguish between buffers
1483and sub-buffers.
1484
1485=head2 THE OpenCL::BufferObj CLASS
1486
1487This is a subclass of OpenCL::Buffer and thus OpenCL::Memory. It exists
1488because one cna create sub buffers of OpenLC::BufferObj objects, but not
1489sub buffers from these sub buffers.
1490
1491=over 4
1492
1493=item $subbuf = $buf_obj->sub_buffer_region ($flags, $origin, $size)
1494
1495Creates an OpenCL::Buffer objects from this buffer and returns it. The
1496C<buffer_create_type> is assumed to be C<CL_BUFFER_CREATE_TYPE_REGION>.
1497
1498L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateSubBuffer.html>
1499
1500=back
1501
1502=head2 THE OpenCL::Image CLASS
1503
1504This is the superclass of all image objects - OpenCL::Image1D,
1505OpenCL::Image1DArray, OpenCL::Image1DBuffer, OpenCL::Image2D,
1506OpenCL::Image2DArray and OpenCL::Image3D.
1507
1508=over 4
1509
1510=item $packed_value = $image->image_info ($name)
1511
1512See C<< $platform->info >> for details.
1513
1514The reason this method is not called C<info> is that there already is an
1515C<< ->info >> method inherited from C<OpenCL::Memory>.
1516
1517L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetImageInfo.html>
1518
1519=item ($channel_order, $channel_data_type) = $image->format
1520
1521Returns the channel order and type used to create the image by calling
1522C<clGetImageInfo> with C<CL_IMAGE_FORMAT>.
1523
1524=for gengetinfo begin image
1525
1526=item $int = $image->element_size
1527
1528Calls C<clGetImageInfo> with C<CL_IMAGE_ELEMENT_SIZE> and returns the result.
1529
1530=item $int = $image->row_pitch
1531
1532Calls C<clGetImageInfo> with C<CL_IMAGE_ROW_PITCH> and returns the result.
1533
1534=item $int = $image->slice_pitch
1535
1536Calls C<clGetImageInfo> with C<CL_IMAGE_SLICE_PITCH> and returns the result.
1537
1538=item $int = $image->width
1539
1540Calls C<clGetImageInfo> with C<CL_IMAGE_WIDTH> and returns the result.
1541
1542=item $int = $image->height
1543
1544Calls C<clGetImageInfo> with C<CL_IMAGE_HEIGHT> and returns the result.
1545
1546=item $int = $image->depth
1547
1548Calls C<clGetImageInfo> with C<CL_IMAGE_DEPTH> and returns the result.
1549
1550=for gengetinfo end image
1551
1552=for gengetinfo begin gl_texture
1553
1554=item $GLenum = $gl_texture->target
1555
1556Calls C<clGetGLTextureInfo> with C<CL_GL_TEXTURE_TARGET> and returns the result.
1557
1558=item $GLint = $gl_texture->gl_mipmap_level
1559
1560Calls C<clGetGLTextureInfo> with C<CL_GL_MIPMAP_LEVEL> and returns the result.
1561
1562=for gengetinfo end gl_texture
1563
1564=back
1565
552=head2 THE OpenCL::Sampler CLASS 1566=head2 THE OpenCL::Sampler CLASS
553 1567
554=over 4 1568=over 4
555 1569
556=item $packed_value = $sampler->info ($name) 1570=item $packed_value = $sampler->info ($name)
557 1571
558See C<< $platform->info >> for details. 1572See C<< $platform->info >> for details.
559 1573
560L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetSamplerInfo.html> 1574L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetSamplerInfo.html>
561 1575
1576=for gengetinfo begin sampler
1577
1578=item $uint = $sampler->reference_count
1579
1580Calls C<clGetSamplerInfo> with C<CL_SAMPLER_REFERENCE_COUNT> and returns the result.
1581
1582=item $ctx = $sampler->context
1583
1584Calls C<clGetSamplerInfo> with C<CL_SAMPLER_CONTEXT> and returns the result.
1585
1586=item $addressing_mode = $sampler->normalized_coords
1587
1588Calls C<clGetSamplerInfo> with C<CL_SAMPLER_NORMALIZED_COORDS> and returns the result.
1589
1590=item $filter_mode = $sampler->addressing_mode
1591
1592Calls C<clGetSamplerInfo> with C<CL_SAMPLER_ADDRESSING_MODE> and returns the result.
1593
1594=item $boolean = $sampler->filter_mode
1595
1596Calls C<clGetSamplerInfo> with C<CL_SAMPLER_FILTER_MODE> and returns the result.
1597
1598=for gengetinfo end sampler
1599
562=back 1600=back
563 1601
564=head2 THE OpenCL::Program CLASS 1602=head2 THE OpenCL::Program CLASS
565 1603
566=over 4 1604=over 4
567 1605
568=item $packed_value = $program->info ($name) 1606=item $program->build (\@devices = undef, $options = "", $cb->($program) = undef)
569 1607
570See C<< $platform->info >> for details.
571
572L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetProgramInfo.html>
573
574=item $program->build ($device, $options = "")
575
576Tries to build the program with the givne options. 1608Tries to build the program with the given options. See also the
1609C<$ctx->build> convenience function.
1610
1611If a callback is specified, then it will be called when compilation is
1612finished. Note that many OpenCL implementations block your program while
1613compiling whether you use a callback or not. See C<build_async> if you
1614want to make sure the build is done in the background.
1615
1616Note that some OpenCL implementations act up badly, and don't call the
1617callback in some error cases (but call it in others). This implementation
1618assumes the callback will always be called, and leaks memory if this is
1619not so. So best make sure you don't pass in invalid values.
1620
1621Some implementations fail with C<OpenCL::INVALID_BINARY> when the
1622compilation state is successful but some later stage fails.
577 1623
578L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clBuildProgram.html> 1624L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clBuildProgram.html>
1625
1626=item $program->build_async (\@devices = undef, $options = "", $cb->($program) = undef)
1627
1628Similar to C<< ->build >>, except it starts a thread, and never fails (you
1629need to check the compilation status form the callback, or by polling).
579 1630
580=item $packed_value = $program->build_info ($device, $name) 1631=item $packed_value = $program->build_info ($device, $name)
581 1632
582Similar to C<< $platform->info >>, but returns build info for a previous 1633Similar to C<< $platform->info >>, but returns build info for a previous
583build attempt for the given device. 1634build attempt for the given device.
… …
589Creates an OpenCL::Kernel object out of the named C<__kernel> function in 1640Creates an OpenCL::Kernel object out of the named C<__kernel> function in
590the program. 1641the program.
591 1642
592L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernel.html> 1643L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernel.html>
593 1644
1645=item @kernels = $program->kernels_in_program
1646
1647Returns all kernels successfully compiled for all devices in program.
1648
1649http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernelsInProgram.html
1650
1651=for gengetinfo begin program_build
1652
1653=item $build_status = $program->build_status ($device)
1654
1655Calls C<clGetProgramBuildInfo> with C<CL_PROGRAM_BUILD_STATUS> and returns the result.
1656
1657=item $string = $program->build_options ($device)
1658
1659Calls C<clGetProgramBuildInfo> with C<CL_PROGRAM_BUILD_OPTIONS> and returns the result.
1660
1661=item $string = $program->build_log ($device)
1662
1663Calls C<clGetProgramBuildInfo> with C<CL_PROGRAM_BUILD_LOG> and returns the result.
1664
1665=for gengetinfo end program_build
1666
1667=item $packed_value = $program->info ($name)
1668
1669See C<< $platform->info >> for details.
1670
1671L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetProgramInfo.html>
1672
1673=for gengetinfo begin program
1674
1675=item $uint = $program->reference_count
1676
1677Calls C<clGetProgramInfo> with C<CL_PROGRAM_REFERENCE_COUNT> and returns the result.
1678
1679=item $ctx = $program->context
1680
1681Calls C<clGetProgramInfo> with C<CL_PROGRAM_CONTEXT> and returns the result.
1682
1683=item $uint = $program->num_devices
1684
1685Calls C<clGetProgramInfo> with C<CL_PROGRAM_NUM_DEVICES> and returns the result.
1686
1687=item @devices = $program->devices
1688
1689Calls C<clGetProgramInfo> with C<CL_PROGRAM_DEVICES> and returns the result.
1690
1691=item $string = $program->source
1692
1693Calls C<clGetProgramInfo> with C<CL_PROGRAM_SOURCE> and returns the result.
1694
1695=item @ints = $program->binary_sizes
1696
1697Calls C<clGetProgramInfo> with C<CL_PROGRAM_BINARY_SIZES> and returns the result.
1698
1699=for gengetinfo end program
1700
1701=item @blobs = $program->binaries
1702
1703Returns a string for the compiled binary for every device associated with
1704the program, empty strings indicate missing programs, and an empty result
1705means no program binaries are available.
1706
1707These "binaries" are often, in fact, informative low-level assembly
1708sources.
1709
1710L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetProgramInfo.html>
1711
594=back 1712=back
595 1713
596=head2 THE OpenCL::Kernel CLASS 1714=head2 THE OpenCL::Kernel CLASS
597 1715
598=over 4 1716=over 4
… …
601 1719
602See C<< $platform->info >> for details. 1720See C<< $platform->info >> for details.
603 1721
604L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetKernelInfo.html> 1722L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetKernelInfo.html>
605 1723
1724=for gengetinfo begin kernel
1725
1726=item $string = $kernel->function_name
1727
1728Calls C<clGetKernelInfo> with C<CL_KERNEL_FUNCTION_NAME> and returns the result.
1729
1730=item $uint = $kernel->num_args
1731
1732Calls C<clGetKernelInfo> with C<CL_KERNEL_NUM_ARGS> and returns the result.
1733
1734=item $uint = $kernel->reference_count
1735
1736Calls C<clGetKernelInfo> with C<CL_KERNEL_REFERENCE_COUNT> and returns the result.
1737
1738=item $ctx = $kernel->context
1739
1740Calls C<clGetKernelInfo> with C<CL_KERNEL_CONTEXT> and returns the result.
1741
1742=item $program = $kernel->program
1743
1744Calls C<clGetKernelInfo> with C<CL_KERNEL_PROGRAM> and returns the result.
1745
1746=for gengetinfo end kernel
1747
1748=item $packed_value = $kernel->work_group_info ($device, $name)
1749
1750See C<< $platform->info >> for details.
1751
1752The reason this method is not called C<info> is that there already is an
1753C<< ->info >> method.
1754
1755L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetKernelWorkGroupInfo.html>
1756
1757=for gengetinfo begin kernel_work_group
1758
1759=item $int = $kernel->work_group_size ($device)
1760
1761Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_WORK_GROUP_SIZE> and returns the result.
1762
1763=item @ints = $kernel->compile_work_group_size ($device)
1764
1765Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_COMPILE_WORK_GROUP_SIZE> and returns the result.
1766
1767=item $ulong = $kernel->local_mem_size ($device)
1768
1769Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_LOCAL_MEM_SIZE> and returns the result.
1770
1771=item $int = $kernel->preferred_work_group_size_multiple ($device)
1772
1773Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE> and returns the result.
1774
1775=item $ulong = $kernel->private_mem_size ($device)
1776
1777Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PRIVATE_MEM_SIZE> and returns the result.
1778
1779=for gengetinfo end kernel_work_group
1780
1781=item $kernel->setf ($format, ...)
1782
1783Sets the arguments of a kernel. Since OpenCL 1.1 doesn't have a generic
1784way to set arguments (and with OpenCL 1.2 it might be rather slow), you
1785need to specify a format argument, much as with C<printf>, to tell OpenCL
1786what type of argument it is.
1787
1788The format arguments are single letters:
1789
1790 c char
1791 C unsigned char
1792 s short
1793 S unsigned short
1794 i int
1795 I unsigned int
1796 l long
1797 L unsigned long
1798
1799 h half float (0..65535)
1800 f float
1801 d double
1802
1803 z local (octet size)
1804
1805 m memory object (buffer or image)
1806 a sampler
1807 e event
1808
1809Space characters in the format string are ignored.
1810
1811Example: set the arguments for a kernel that expects an int, two floats, a buffer and an image.
1812
1813 $kernel->setf ("i ff mm", 5, 0.5, 3, $buffer, $image);
1814
606=item $kernel->set_TYPE ($index, $value) 1815=item $kernel->set_TYPE ($index, $value)
607 1816
1817=item $kernel->set_char ($index, $value)
1818
1819=item $kernel->set_uchar ($index, $value)
1820
1821=item $kernel->set_short ($index, $value)
1822
1823=item $kernel->set_ushort ($index, $value)
1824
1825=item $kernel->set_int ($index, $value)
1826
1827=item $kernel->set_uint ($index, $value)
1828
1829=item $kernel->set_long ($index, $value)
1830
1831=item $kernel->set_ulong ($index, $value)
1832
1833=item $kernel->set_half ($index, $value)
1834
1835=item $kernel->set_float ($index, $value)
1836
1837=item $kernel->set_double ($index, $value)
1838
1839=item $kernel->set_memory ($index, $value)
1840
1841=item $kernel->set_buffer ($index, $value)
1842
1843=item $kernel->set_image ($index, $value)
1844
1845=item $kernel->set_sampler ($index, $value)
1846
1847=item $kernel->set_local ($index, $value)
1848
1849=item $kernel->set_event ($index, $value)
1850
608This is a family of methods to set the kernel argument with the number C<$index> to the give C<$value>. 1851This is a family of methods to set the kernel argument with the number
609 1852C<$index> to the give C<$value>.
610TYPE is one of C<char>, C<uchar>, C<short>, C<ushort>, C<int>, C<uint>,
611C<long>, C<ulong>, C<half>, C<float>, C<double>, C<memory>, C<buffer>,
612C<image2d>, C<image3d>, C<sampler> or C<event>.
613 1853
614Chars and integers (including the half type) are specified as integers, 1854Chars and integers (including the half type) are specified as integers,
615float and double as floating point values, memory/buffer/image2d/image3d 1855float and double as floating point values, memory/buffer/image must be
616must be an object of that type or C<undef>, and sampler and event must be 1856an object of that type or C<undef>, local-memory arguments are set by
617objects of that type. 1857specifying the size, and sampler and event must be objects of that type.
1858
1859Note that C<set_memory> works for all memory objects (all types of buffers
1860and images) - the main purpose of the more specific C<set_TYPE> functions
1861is type checking.
1862
1863Setting an argument for a kernel does NOT keep a reference to the object -
1864for example, if you set an argument to some image object, free the image,
1865and call the kernel, you will run into undefined behaviour.
618 1866
619L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html> 1867L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html>
620 1868
621=back 1869=back
622 1870
… …
625This is the superclass for all event objects (including OpenCL::UserEvent 1873This is the superclass for all event objects (including OpenCL::UserEvent
626objects). 1874objects).
627 1875
628=over 4 1876=over 4
629 1877
1878=item $ev->wait
1879
1880Waits for the event to complete.
1881
1882L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html>
1883
1884=item $ev->cb ($exec_callback_type, $callback->($event, $event_command_exec_status))
1885
1886Adds a callback to the callback stack for the given event type. There is
1887no way to remove a callback again.
1888
1889L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetEventCallback.html>
1890
630=item $packed_value = $ev->info ($name) 1891=item $packed_value = $ev->info ($name)
631 1892
632See C<< $platform->info >> for details. 1893See C<< $platform->info >> for details.
633 1894
634L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetEventInfo.html> 1895L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetEventInfo.html>
635 1896
636=item $ev->wait 1897=for gengetinfo begin event
637 1898
638Waits for the event to complete. 1899=item $queue = $event->command_queue
639 1900
1901Calls C<clGetEventInfo> with C<CL_EVENT_COMMAND_QUEUE> and returns the result.
1902
1903=item $command_type = $event->command_type
1904
1905Calls C<clGetEventInfo> with C<CL_EVENT_COMMAND_TYPE> and returns the result.
1906
1907=item $uint = $event->reference_count
1908
1909Calls C<clGetEventInfo> with C<CL_EVENT_REFERENCE_COUNT> and returns the result.
1910
1911=item $uint = $event->command_execution_status
1912
1913Calls C<clGetEventInfo> with C<CL_EVENT_COMMAND_EXECUTION_STATUS> and returns the result.
1914
1915=item $ctx = $event->context
1916
1917Calls C<clGetEventInfo> with C<CL_EVENT_CONTEXT> and returns the result.
1918
1919=for gengetinfo end event
1920
1921=item $packed_value = $ev->profiling_info ($name)
1922
1923See C<< $platform->info >> for details.
1924
1925The reason this method is not called C<info> is that there already is an
1926C<< ->info >> method.
1927
640L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clWaitForEvents.html> 1928L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetProfilingInfo.html>
1929
1930=for gengetinfo begin profiling
1931
1932=item $ulong = $event->profiling_command_queued
1933
1934Calls C<clGetEventProfilingInfo> with C<CL_PROFILING_COMMAND_QUEUED> and returns the result.
1935
1936=item $ulong = $event->profiling_command_submit
1937
1938Calls C<clGetEventProfilingInfo> with C<CL_PROFILING_COMMAND_SUBMIT> and returns the result.
1939
1940=item $ulong = $event->profiling_command_start
1941
1942Calls C<clGetEventProfilingInfo> with C<CL_PROFILING_COMMAND_START> and returns the result.
1943
1944=item $ulong = $event->profiling_command_end
1945
1946Calls C<clGetEventProfilingInfo> with C<CL_PROFILING_COMMAND_END> and returns the result.
1947
1948=for gengetinfo end profiling
641 1949
642=back 1950=back
643 1951
644=head2 THE OpenCL::UserEvent CLASS 1952=head2 THE OpenCL::UserEvent CLASS
645 1953
… …
647 1955
648=over 4 1956=over 4
649 1957
650=item $ev->set_status ($execution_status) 1958=item $ev->set_status ($execution_status)
651 1959
1960Sets the execution status of the user event. Can only be called once,
1961either with OpenCL::COMPLETE or a negative number as status.
1962
652L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetUserEventStatus.html> 1963L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetUserEventStatus.html>
653 1964
654=back 1965=back
655 1966
1967=head2 THE OpenCL::Mapped CLASS
1968
1969This class represents objects mapped into host memory. They are
1970represented by a blessed string scalar. The string data is the mapped
1971memory area, that is, if you read or write it, then the mapped object is
1972accessed directly.
1973
1974You must only ever use operations that modify the string in-place - for
1975example, a C<substr> that doesn't change the length, or maybe a regex that
1976doesn't change the length. Any other operation might cause the data to be
1977copied.
1978
1979When the object is destroyed it will enqueue an implicit unmap operation
1980on the queue that was used to create it.
1981
1982Keep in mind that you I<need> to unmap (or destroy) mapped objects before
1983OpenCL sees the changes, even if some implementations don't need this
1984sometimes.
1985
1986Example, replace the first two floats in the mapped buffer by 1 and 2.
1987
1988 my $mapped = $queue->map_buffer ($buf, ...
1989 $mapped->event->wait; # make sure it's there
1990
1991 # now replace first 8 bytes by new data, which is exactly 8 bytes long
1992 # we blindly assume device endianness to equal host endianness
1993 # (and of course, we assume iee 754 single precision floats :)
1994 substr $$mapped, 0, 8, pack "f*", 1, 2;
1995
1996=over 4
1997
1998=item $ev = $mapped->unmap ($wait_events...)
1999
2000Unmaps the mapped memory object, using the queue originally used to create
2001it, quite similarly to C<< $queue->unmap ($mapped, ...) >>.
2002
2003=item $bool = $mapped->mapped
2004
2005Returns whether the object is still mapped - true before an C<unmap> is
2006enqueued, false afterwards.
2007
2008=item $ev = $mapped->event
2009
2010Return the event object associated with the mapped object. Initially, this
2011will be the event object created when mapping the object, and after an
2012unmap, this will be the event object that the unmap operation created.
2013
2014=item $mapped->wait
2015
2016Same as C<< $mapped->event->wait >> - makes sure no operations on this
2017mapped object are outstanding.
2018
2019=item $bytes = $mapped->size
2020
2021Returns the size of the mapped area, in bytes. Same as C<length $$mapped>.
2022
2023=item $ptr = $mapped->ptr
2024
2025Returns the raw memory address of the mapped area.
2026
2027=item $mapped->set ($offset, $data)
2028
2029Replaces the data at the given C<$offset> in the memory area by the new
2030C<$data>. This method is safer than direct manipulation of C<$mapped>
2031because it does bounds-checking, but also slower.
2032
2033=item $data = $mapped->get ($offset, $length)
2034
2035Returns (without copying) a scalar representing the data at the given
2036C<$offset> and C<$length> in the mapped memory area. This is the same as
2037the following substr, except much slower;
2038
2039 $data = substr $$mapped, $offset, $length
2040
656=cut 2041=cut
657 2042
658package OpenCL; 2043sub OpenCL::Mapped::get {
659 2044 substr ${$_[0]}, $_[1], $_[2]
660use common::sense;
661
662BEGIN {
663 our $VERSION = '0.15';
664
665 require XSLoader;
666 XSLoader::load (__PACKAGE__, $VERSION);
667
668 @OpenCL::Buffer::ISA =
669 @OpenCL::Image::ISA = OpenCL::Memory::;
670
671 @OpenCL::Image2D::ISA =
672 @OpenCL::Image3D::ISA = OpenCL::Image::;
673
674 @OpenCL::UserEvent::ISA = OpenCL::Event::;
675} 2045}
2046
2047=back
2048
2049=head2 THE OpenCL::MappedBuffer CLASS
2050
2051This is a subclass of OpenCL::Mapped, representing mapped buffers.
2052
2053=head2 THE OpenCL::MappedImage CLASS
2054
2055This is a subclass of OpenCL::Mapped, representing mapped images.
2056
2057=over 4
2058
2059=item $bytes = $mapped->row_pitch
2060
2061=item $bytes = $mapped->slice_pitch
2062
2063Return the row or slice pitch of the image that has been mapped.
2064
2065=back
2066
2067
2068=cut
676 2069
6771; 20701;
678 2071
679=head1 AUTHOR 2072=head1 AUTHOR
680 2073

Diff Legend

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