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

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.56 by root, Tue Apr 24 23:53:12 2012 UTC vs.
Revision 1.68 by root, Tue May 1 22:25:13 2012 UTC

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
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.
171 # set buffer 174 # set buffer
172 $kernel->set_buffer (0, $input); 175 $kernel->set_buffer (0, $input);
173 $kernel->set_buffer (1, $output); 176 $kernel->set_buffer (1, $output);
174 177
175 # execute it for all 4 numbers 178 # execute it for all 4 numbers
176 $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); 179 $queue->nd_range_kernel ($kernel, undef, [4], undef);
177 180
178 # enqueue a synchronous read 181 # enqueue a synchronous read
179 $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);
180 183
181 # print the results: 184 # print the results:
182 printf "%s\n", join ", ", unpack "f*", $data; 185 printf "%s\n", join ", ", unpack "f*", $data;
183 186
184=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,
185showing off barriers. 188showing off barriers.
186 189
187 # execute it for all 4 numbers 190 # execute it for all 4 numbers
188 $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); 191 $queue->nd_range_kernel ($kernel, undef, [4], undef);
189 192
190 # enqueue a barrier to ensure in-order execution 193 # enqueue a barrier to ensure in-order execution
191 $queue->enqueue_barrier; 194 $queue->barrier;
192 195
193 # enqueue an async read 196 # enqueue an async read
194 $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);
195 198
196 # wait for all requests to finish 199 # wait for all requests to finish
197 $queue->finish; 200 $queue->finish;
198 201
199=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,
200showing off event objects and wait lists. 203showing off event objects and wait lists.
201 204
202 # execute it for all 4 numbers 205 # execute it for all 4 numbers
203 my $ev = $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); 206 my $ev = $queue->nd_range_kernel ($kernel, undef, [4], undef);
204 207
205 # enqueue an async read 208 # enqueue an async read
206 $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);
207 210
208 # wait for the last event to complete 211 # wait for the last event to complete
209 $ev->wait; 212 $ev->wait;
210 213
211=head2 Use the OpenGL module to share a texture between OpenCL and OpenGL and draw some julia 214=head2 Use the OpenGL module to share a texture between OpenCL and OpenGL and draw some julia
212set tunnel effect. 215set tunnel effect.
213 216
214This is quite a long example to get you going. 217This is quite a long example to get you going - you can download it from
218L<http://cvs.schmorp.de/OpenCL/examples/juliaflight>.
215 219
216 use OpenGL ":all"; 220 use OpenGL ":all";
217 use OpenCL; 221 use OpenCL;
218 222
223 my $S = $ARGV[0] || 256; # window/texture size, smaller is faster
224
219 # open a window and create a gl texture 225 # open a window and create a gl texture
220 OpenGL::glpOpenWindow width => 256, height => 256; 226 OpenGL::glpOpenWindow width => $S, height => $S;
221 my $texid = glGenTextures_p 1; 227 my $texid = glGenTextures_p 1;
222 glBindTexture GL_TEXTURE_2D, $texid; 228 glBindTexture GL_TEXTURE_2D, $texid;
223 glTexImage2D_c GL_TEXTURE_2D, 0, GL_RGBA8, 256, 256, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0; 229 glTexImage2D_c GL_TEXTURE_2D, 0, GL_RGBA8, $S, $S, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0;
224 230
225 # find and use the first opencl device that let's us get a shared opengl context 231 # find and use the first opencl device that let's us get a shared opengl context
226 my $platform; 232 my $platform;
227 my $dev; 233 my $dev;
228 my $ctx; 234 my $ctx;
247 # now the boring opencl code 253 # now the boring opencl code
248 my $src = <<EOF; 254 my $src = <<EOF;
249 kernel void 255 kernel void
250 juliatunnel (write_only image2d_t img, float time) 256 juliatunnel (write_only image2d_t img, float time)
251 { 257 {
252 float2 p = (float2)(get_global_id (0), get_global_id (1)) / 256.f * 2.f - 1.f; 258 int2 xy = (int2)(get_global_id (0), get_global_id (1));
259 float2 p = convert_float2 (xy) / $S.f * 2.f - 1.f;
253 260
254 float2 m = (float2)(1.f, p.y) / fabs (p.x); 261 float2 m = (float2)(1.f, p.y) / fabs (p.x); // tunnel
255 m.x = fabs (fmod (m.x + time * 0.05f, 4.f)) - 2.f; 262 m.x = fabs (fmod (m.x + time * 0.05f, 4.f) - 2.f);
256 263
257 float2 z = m; 264 float2 z = m;
258 float2 c = (float2)(sin (time * 0.05005), cos (time * 0.06001)); 265 float2 c = (float2)(sin (time * 0.01133f), cos (time * 0.02521f));
259 266
260 for (int i = 0; i < 25 && dot (z, z) < 4.f; ++i) 267 for (int i = 0; i < 25 && dot (z, z) < 4.f; ++i) // standard julia
261 z = (float2)(z.x * z.x - z.y * z.y, 2.f * z.x * z.y) + c; 268 z = (float2)(z.x * z.x - z.y * z.y, 2.f * z.x * z.y) + c;
262 269
263 float3 colour = (float3)(z.x, z.y, z.x * z.y); 270 float3 colour = (float3)(z.x, z.y, atan2 (z.y, z.x));
264 write_imagef (img, (int2)(get_global_id (0), get_global_id (1)), (float4)(colour * p.x * p.x, 1.)); 271 write_imagef (img, xy, (float4)(colour * p.x * p.x, 1.));
265 } 272 }
266 EOF 273 EOF
267 274
268 my $prog = $ctx->build_program ($src); 275 my $prog = $ctx->build_program ($src);
269 my $kernel = $prog->kernel ("juliatunnel"); 276 my $kernel = $prog->kernel ("juliatunnel");
270 277
271 # program compiled, kernel ready, now draw and loop 278 # program compiled, kernel ready, now draw and loop
272 279
273 for (my $time; ; ++$time) { 280 for (my $time; ; ++$time) {
274 # acquire objects from opengl 281 # acquire objects from opengl
275 $queue->enqueue_acquire_gl_objects ([$tex]); 282 $queue->acquire_gl_objects ([$tex]);
276 283
277 # configure and run our kernel 284 # configure and run our kernel
278 $kernel->set_image2d (0, $tex); 285 $kernel->setf ("mf", $tex, $time*2); # mf = memory object, float
279 $kernel->set_float (1, $time);
280 $queue->enqueue_nd_range_kernel ($kernel, undef, [256, 256], undef); 286 $queue->nd_range_kernel ($kernel, undef, [$S, $S], undef);
281 287
282 # release objects to opengl again 288 # release objects to opengl again
283 $queue->enqueue_release_gl_objects ([$tex]); 289 $queue->release_gl_objects ([$tex]);
284 290
285 # wait 291 # wait
286 $queue->finish; 292 $queue->finish;
287 293
288 # now draw the texture, the defaults should be all right 294 # now draw the texture, the defaults should be all right
298 304
299 glXSwapBuffers; 305 glXSwapBuffers;
300 306
301 select undef, undef, undef, 1/60; 307 select undef, undef, undef, 1/60;
302 } 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>.
303 336
304=head1 DOCUMENTATION 337=head1 DOCUMENTATION
305 338
306=head2 BASIC CONVENTIONS 339=head2 BASIC CONVENTIONS
307 340
417 use AnyEvent; 450 use AnyEvent;
418 451
419 # initialise AnyEvent, by creating a watcher, or: 452 # initialise AnyEvent, by creating a watcher, or:
420 AnyEvent::detect; 453 AnyEvent::detect;
421 454
422 my $e = $queue->enqueue_marker; 455 my $e = $queue->marker;
423 $e->cb (sub { 456 $e->cb (sub {
424 warn "opencl is finished\n"; 457 warn "opencl is finished\n";
425 }) 458 })
426 459
427 main Gtk2; 460 main Gtk2;
444Asynchronously means that perl might execute your callbacks at any 477Asynchronously means that perl might execute your callbacks at any
445time. For example, in the following code (I<THAT YOU SHOULD NOT COPY>), 478time. For example, in the following code (I<THAT YOU SHOULD NOT COPY>),
446the C<until> loop following the marker call will be interrupted by the 479the C<until> loop following the marker call will be interrupted by the
447callback: 480callback:
448 481
449 my $e = $queue->enqueue_marker; 482 my $e = $queue->marker;
450 my $flag; 483 my $flag;
451 $e->cb (sub { $flag = 1 }); 484 $e->cb (sub { $flag = 1 });
452 1 until $flag; 485 1 until $flag;
453 # $flag is now 1 486 # $flag is now 1
454 487
490=cut 523=cut
491 524
492package OpenCL; 525package OpenCL;
493 526
494use common::sense; 527use common::sense;
528use Carp ();
495use Async::Interrupt (); 529use Async::Interrupt ();
496 530
497our $POLL_FUNC; # set by XS 531our $POLL_FUNC; # set by XS
498 532
499BEGIN { 533BEGIN {
500 our $VERSION = '0.97'; 534 our $VERSION = '0.98';
501 535
502 require XSLoader; 536 require XSLoader;
503 XSLoader::load (__PACKAGE__, $VERSION); 537 XSLoader::load (__PACKAGE__, $VERSION);
504 538
505 @OpenCL::Platform::ISA = 539 @OpenCL::Platform::ISA =
523 @OpenCL::Image1D::ISA = 557 @OpenCL::Image1D::ISA =
524 @OpenCL::Image1DArray::ISA = 558 @OpenCL::Image1DArray::ISA =
525 @OpenCL::Image1DBuffer::ISA = OpenCL::Image::; 559 @OpenCL::Image1DBuffer::ISA = OpenCL::Image::;
526 560
527 @OpenCL::UserEvent::ISA = OpenCL::Event::; 561 @OpenCL::UserEvent::ISA = OpenCL::Event::;
562
563 @OpenCL::MappedBuffer::ISA =
564 @OpenCL::MappedImage::ISA = OpenCL::Mapped::;
528} 565}
529 566
530=head2 THE OpenCL PACKAGE 567=head2 THE OpenCL PACKAGE
531 568
532=over 4 569=over 4
534=item $int = OpenCL::errno 571=item $int = OpenCL::errno
535 572
536The 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
537and before calling another OpenCL function. 574and before calling another OpenCL function.
538 575
539=item $str = OpenCL::err2str $errval 576=item $str = OpenCL::err2str [$errval]
540 577
541Comverts 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).
542 580
543=item $str = OpenCL::enum2str $enum 581=item $str = OpenCL::enum2str $enum
544 582
545Converts most enum values (of parameter names, image format constants, 583Converts most enum values (of parameter names, image format constants,
546object types, addressing and filter modes, command types etc.) into a 584object types, addressing and filter modes, command types etc.) into a
999=cut 1037=cut
1000 1038
1001sub OpenCL::Context::build_program { 1039sub OpenCL::Context::build_program {
1002 my ($self, $prog, $options) = @_; 1040 my ($self, $prog, $options) = @_;
1003 1041
1004 require Carp;
1005
1006 $prog = $self->program_with_source ($prog) 1042 $prog = $self->program_with_source ($prog)
1007 unless ref $prog; 1043 unless ref $prog;
1008 1044
1009 # we build separately per device so we instantly know which one failed 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
1010 for my $dev ($self->devices) { 1051 for my $dev ($self->devices) {
1011 eval { $prog->build ([$dev], $options); 1 } 1052 $prog->build_status ($dev) == BUILD_SUCCESS
1012 or Carp::croak ("Building OpenCL program for device '" . $dev->name . "' failed:\n" 1053 or Carp::croak "Building OpenCL program for device '" . $dev->name . "' failed:\n"
1013 . $prog->build_log ($dev)); 1054 . $prog->build_log ($dev);
1014 } 1055 }
1015 1056
1016 $prog 1057 $prog
1017} 1058}
1018 1059
1042=item $buf = $ctx->buffer_sv ($flags, $data) 1083=item $buf = $ctx->buffer_sv ($flags, $data)
1043 1084
1044Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object and 1085Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object and
1045initialise it with the given data values. 1086initialise it with the given data values.
1046 1087
1047=item $img = $ctx->image ($self, $flags, $channel_order, $channel_type, $type, $width, $height, $depth, $array_size = 0, $row_pitch = 0, $slice_pitch = 0, $num_mip_level = 0, $num_samples = 0, $*data = &PL_sv_undef) 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)
1048 1089
1049Creates a new OpenCL::Image object and optionally initialises it with 1090Creates a new OpenCL::Image object and optionally initialises it with
1050the given data values. 1091the given data values.
1051 1092
1052L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateImage.html> 1093L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateImage.html>
1148=back 1189=back
1149 1190
1150=head2 THE OpenCL::Queue CLASS 1191=head2 THE OpenCL::Queue CLASS
1151 1192
1152An OpenCL::Queue represents an execution queue for OpenCL. You execute 1193An OpenCL::Queue represents an execution queue for OpenCL. You execute
1153requests by calling their respective C<enqueue_xxx> method and waitinf for 1194requests by calling their respective method and waiting for it to complete
1154it to complete in some way. 1195in some way.
1155 1196
1156All 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
1157for completion, unless the method is called in void context, in which case 1198be used to wait for completion (optionally using a callback), unless
1158no event object is created. 1199the method is called in void context, in which case no event object is
1200created.
1159 1201
1160They 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
1161request 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
1162event objects as extra parameters to the enqueue methods. To simplify 1204event objects as extra parameters to the enqueue methods. To simplify
1163program design, this module ignores any C<undef> values in the list of 1205program design, this module ignores any C<undef> values in the list of
1164events. This makes it possible to code operations such as this, without 1206events. This makes it possible to code operations such as this, without
1165having to put a valid event object into C<$event> first: 1207having to put a valid event object into C<$event> first:
1166 1208
1167 $event = $queue->enqueue_xxx (..., $event); 1209 $event = $queue->xxx (..., $event);
1168 1210
1169Queues execute in-order by default, without any parallelism, so in most 1211Queues execute in-order by default, without any parallelism, so in most
1170cases (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
1171create event objects, althoguh an our of order queue is often a bit 1213create event objects, althoguh an our of order queue is often a bit
1172faster. 1214faster.
1173 1215
1174=over 4 1216=over 4
1175 1217
1176=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...)
1177 1219
1178Reads data from buffer into the given string. 1220Reads data from buffer into the given string.
1179 1221
1180L<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>
1181 1223
1182=item $ev = $queue->enqueue_write_buffer ($buffer, $blocking, $offset, $data, $wait_events...) 1224=item $ev = $queue->write_buffer ($buffer, $blocking, $offset, $data, $wait_events...)
1183 1225
1184Writes data to buffer from the given string. 1226Writes data to buffer from the given string.
1185 1227
1186L<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>
1187 1229
1188=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...)
1189 1231
1190L<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>
1191 1233
1192=item $ev = $queue->enqueue_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...) 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...)
1193 1235
1194http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBufferRect.html 1236http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBufferRect.html
1195 1237
1196=item $ev = $queue->enqueue_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...) 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...)
1197 1239
1198http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBufferRect.html 1240http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBufferRect.html
1199 1241
1200=item $ev = $queue->enqueue_read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
1201
1202L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferRect.html>
1203
1204=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...) 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
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.
1205 1250
1206L<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>
1207 1252
1208=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...)
1209 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.
1210L<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>
1211 1258
1212=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...) 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...)
1213 1260
1214L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html> 1261L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html>
1215 1262
1216=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...) 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...)
1217 1264
1218L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html> 1265L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html>
1219 1266
1220=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...)
1221 1268
1222Yeah. 1269Yeah.
1223 1270
1224L<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>.
1225 1272
1226=item $ev = $queue->enqueue_fill_buffer ($mem, $pattern, $offset, $size, ...) 1273=item $ev = $queue->fill_buffer ($mem, $pattern, $offset, $size, ...)
1227 1274
1228Fills the given buffer object with repeated applications of C<$pattern>, 1275Fills the given buffer object with repeated applications of C<$pattern>,
1229starting at C<$offset> for C<$size> octets. 1276starting at C<$offset> for C<$size> octets.
1230 1277
1231L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillBuffer.html> 1278L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillBuffer.html>
1232 1279
1233=item $ev = $queue->enqueue_fill_image ($img, $r, $g, $b, $a, $x, $y, $z, $width, $height, $depth, ...) 1280=item $ev = $queue->fill_image ($img, $r, $g, $b, $a, $x, $y, $z, $width, $height, $depth, ...)
1234 1281
1235Fills the given image area with the given rgba colour components. The 1282Fills the given image area with the given rgba colour components. The
1236components are normally floating point values between C<0> and C<1>, 1283components are normally floating point values between C<0> and C<1>,
1237except when the image channel data type is a signe dor unsigned 1284except when the image channel data type is a signe dor unsigned
1238unnormalised format, in which case the range is determined by the format. 1285unnormalised format, in which case the range is determined by the format.
1239 1286
1240L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillImage.html> 1287L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillImage.html>
1241 1288
1242=item $ev = $queue->enqueue_task ($kernel, $wait_events...) 1289=item $ev = $queue->task ($kernel, $wait_events...)
1243 1290
1244L<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>
1245 1292
1246=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...)
1247 1294
1248Enqueues a kernel execution. 1295Enqueues a kernel execution.
1249 1296
1250@$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
1251integers specifying the work sizes (element counts). 1298integers specifying the work sizes (element counts).
1252 1299
1253@$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
1254are 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
1255of elements as @$global_work_size. 1302of elements as \@global_work_size.
1256 1303
1257@$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
1258implementation is supposed to choose good local work sizes), or a 1305implementation is supposed to choose good local work sizes), or a
1259reference 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
1260elements as @$global_work_size. 1307elements as \@global_work_size.
1261 1308
1262L<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>
1263 1310
1264=item $ev = $queue->enqueue_acquire_gl_objects ([object, ...], $wait_events...) 1311=item $ev = $queue->acquire_gl_objects ([object, ...], $wait_events...)
1265 1312
1266Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired 1313Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired
1267for subsequent OpenCL usage. 1314for subsequent OpenCL usage.
1268 1315
1269L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueAcquireGLObjects.html> 1316L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueAcquireGLObjects.html>
1270 1317
1271=item $ev = $queue->enqueue_release_gl_objects ([object, ...], $wait_events...) 1318=item $ev = $queue->release_gl_objects ([object, ...], $wait_events...)
1272 1319
1273Enqueues a list (an array-ref of OpenCL::Memory objects) to be released 1320Enqueues a list (an array-ref of OpenCL::Memory objects) to be released
1274for subsequent OpenGL usage. 1321for subsequent OpenGL usage.
1275 1322
1276L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReleaseGLObjects.html> 1323L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReleaseGLObjects.html>
1277 1324
1278=item $ev = $queue->enqueue_wait_for_events ($wait_events...) 1325=item $ev = $queue->wait_for_events ($wait_events...)
1279 1326
1280L<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>
1281 1328
1282=item $ev = $queue->enqueue_marker ($wait_events...) 1329=item $ev = $queue->marker ($wait_events...)
1283 1330
1284L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMarkerWithWaitList.html> 1331L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMarkerWithWaitList.html>
1285 1332
1286=item $ev = $queue->enqueue_barrier ($wait_events...) 1333=item $ev = $queue->barrier ($wait_events...)
1287 1334
1288L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueBarrierWithWaitList.html> 1335L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueBarrierWithWaitList.html>
1289 1336
1290=item $queue->flush 1337=item $queue->flush
1291 1338
1318=item $command_queue_properties = $command_queue->properties 1365=item $command_queue_properties = $command_queue->properties
1319 1366
1320Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_PROPERTIES> and returns the result. 1367Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_PROPERTIES> and returns the result.
1321 1368
1322=for gengetinfo end command_queue 1369=for gengetinfo end command_queue
1370
1371=back
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.
1323 1414
1324=back 1415=back
1325 1416
1326=head2 THE OpenCL::Memory CLASS 1417=head2 THE OpenCL::Memory CLASS
1327 1418
1520If a callback is specified, then it will be called when compilation is 1611If a callback is specified, then it will be called when compilation is
1521finished. Note that many OpenCL implementations block your program while 1612finished. Note that many OpenCL implementations block your program while
1522compiling whether you use a callback or not. See C<build_async> if you 1613compiling whether you use a callback or not. See C<build_async> if you
1523want to make sure the build is done in the background. 1614want to make sure the build is done in the background.
1524 1615
1525Note that some OpenCL implementations atc up badly, and don't call the 1616Note that some OpenCL implementations act up badly, and don't call the
1526callback in some error cases (but call it in others). This implementation 1617callback in some error cases (but call it in others). This implementation
1527assumes the callback will always be called, and leaks memory if this is 1618assumes the callback will always be called, and leaks memory if this is
1528not so. So best make sure you don't pass in invalid values. 1619not so. So best make sure you don't pass in invalid values.
1529 1620
1621Some implementations fail with C<OpenCL::INVALID_BINARY> when the
1622compilation state is successful but some later stage fails.
1623
1530L<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>
1531 1625
1532=item $program->build_async (\@devices = undef, $options = "", $cb->($program) = undef) 1626=item $program->build_async (\@devices = undef, $options = "", $cb->($program) = undef)
1533 1627
1534Similar to C<< ->build >>, except it starts a thread, and never fails (you 1628Similar to C<< ->build >>, except it starts a thread, and never fails (you
1682 1776
1683Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PRIVATE_MEM_SIZE> and returns the result. 1777Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PRIVATE_MEM_SIZE> and returns the result.
1684 1778
1685=for gengetinfo end kernel_work_group 1779=for gengetinfo end kernel_work_group
1686 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
1687=item $kernel->set_TYPE ($index, $value) 1815=item $kernel->set_TYPE ($index, $value)
1688 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
1689This 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
1690 1852C<$index> to the give C<$value>.
1691TYPE is one of C<char>, C<uchar>, C<short>, C<ushort>, C<int>, C<uint>,
1692C<long>, C<ulong>, C<half>, C<float>, C<double>, C<memory>, C<buffer>,
1693C<image2d>, C<image3d>, C<sampler>, C<local> or C<event>.
1694 1853
1695Chars and integers (including the half type) are specified as integers, 1854Chars and integers (including the half type) are specified as integers,
1696float and double as floating point values, memory/buffer/image2d/image3d 1855float and double as floating point values, memory/buffer/image must be
1697must be an object of that type or C<undef>, local-memory arguments are 1856an object of that type or C<undef>, local-memory arguments are set by
1698set by specifying the size, and sampler and event must be objects of that 1857specifying the size, and sampler and event must be objects of that type.
1699type. 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.
1700 1862
1701Setting an argument for a kernel does NOT keep a reference to the object - 1863Setting an argument for a kernel does NOT keep a reference to the object -
1702for example, if you set an argument to some image object, free the image, 1864for example, if you set an argument to some image object, free the image,
1703and call the kernel, you will run into undefined behaviour. 1865and call the kernel, you will run into undefined behaviour.
1704 1866
1800 1962
1801L<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>
1802 1964
1803=back 1965=back
1804 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
2041=cut
2042
2043sub OpenCL::Mapped::get {
2044 substr ${$_[0]}, $_[1], $_[2]
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
1805=cut 2068=cut
1806 2069
18071; 20701;
1808 2071
1809=head1 AUTHOR 2072=head1 AUTHOR

Diff Legend

Removed lines
+ Added lines
< Changed lines
> Changed lines