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

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.57 by root, Tue Apr 24 23:58:34 2012 UTC vs.
Revision 1.69 by root, Thu May 3 23:30:08 2012 UTC

141=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,
142then asynchronously. 142then asynchronously.
143 143
144 my $buf = $ctx->buffer_sv (OpenCL::MEM_COPY_HOST_PTR, "helmut"); 144 my $buf = $ctx->buffer_sv (OpenCL::MEM_COPY_HOST_PTR, "helmut");
145 145
146 $queue->enqueue_read_buffer ($buf, 1, 1, 3, my $data); 146 $queue->read_buffer ($buf, 1, 1, 3, my $data);
147 print "$data\n"; 147 print "$data\n";
148 148
149 my $ev = $queue->enqueue_read_buffer ($buf, 0, 1, 3, my $data); 149 my $ev = $queue->read_buffer ($buf, 0, 1, 3, my $data);
150 $ev->wait; 150 $ev->wait;
151 print "$data\n"; # prints "elm" 151 print "$data\n"; # prints "elm"
152 152
153=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
154functions. 154functions.
174 # set buffer 174 # set buffer
175 $kernel->set_buffer (0, $input); 175 $kernel->set_buffer (0, $input);
176 $kernel->set_buffer (1, $output); 176 $kernel->set_buffer (1, $output);
177 177
178 # execute it for all 4 numbers 178 # execute it for all 4 numbers
179 $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); 179 $queue->nd_range_kernel ($kernel, undef, [4], undef);
180 180
181 # enqueue a synchronous read 181 # enqueue a synchronous read
182 $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);
183 183
184 # print the results: 184 # print the results:
185 printf "%s\n", join ", ", unpack "f*", $data; 185 printf "%s\n", join ", ", unpack "f*", $data;
186 186
187=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,
188showing off barriers. 188showing off barriers.
189 189
190 # execute it for all 4 numbers 190 # execute it for all 4 numbers
191 $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); 191 $queue->nd_range_kernel ($kernel, undef, [4], undef);
192 192
193 # enqueue a barrier to ensure in-order execution 193 # enqueue a barrier to ensure in-order execution
194 $queue->enqueue_barrier; 194 $queue->barrier;
195 195
196 # enqueue an async read 196 # enqueue an async read
197 $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);
198 198
199 # wait for all requests to finish 199 # wait for all requests to finish
200 $queue->finish; 200 $queue->finish;
201 201
202=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,
203showing off event objects and wait lists. 203showing off event objects and wait lists.
204 204
205 # execute it for all 4 numbers 205 # execute it for all 4 numbers
206 my $ev = $queue->enqueue_nd_range_kernel ($kernel, undef, [4], undef); 206 my $ev = $queue->nd_range_kernel ($kernel, undef, [4], undef);
207 207
208 # enqueue an async read 208 # enqueue an async read
209 $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);
210 210
211 # wait for the last event to complete 211 # wait for the last event to complete
212 $ev->wait; 212 $ev->wait;
213 213
214=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
215set tunnel effect. 215set tunnel effect.
216 216
217This 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>.
218 219
219 use OpenGL ":all"; 220 use OpenGL ":all";
220 use OpenCL; 221 use OpenCL;
221 222
223 my $S = $ARGV[0] || 256; # window/texture size, smaller is faster
224
222 # open a window and create a gl texture 225 # open a window and create a gl texture
223 OpenGL::glpOpenWindow width => 256, height => 256; 226 OpenGL::glpOpenWindow width => $S, height => $S;
224 my $texid = glGenTextures_p 1; 227 my $texid = glGenTextures_p 1;
225 glBindTexture GL_TEXTURE_2D, $texid; 228 glBindTexture GL_TEXTURE_2D, $texid;
226 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;
227 230
228 # 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
229 my $platform; 232 my $platform;
230 my $dev; 233 my $dev;
231 my $ctx; 234 my $ctx;
250 # now the boring opencl code 253 # now the boring opencl code
251 my $src = <<EOF; 254 my $src = <<EOF;
252 kernel void 255 kernel void
253 juliatunnel (write_only image2d_t img, float time) 256 juliatunnel (write_only image2d_t img, float time)
254 { 257 {
255 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;
256 260
257 float2 m = (float2)(1.f, p.y) / fabs (p.x); 261 float2 m = (float2)(1.f, p.y) / fabs (p.x); // tunnel
258 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);
259 263
260 float2 z = m; 264 float2 z = m;
261 float2 c = (float2)(sin (time * 0.05005), cos (time * 0.06001)); 265 float2 c = (float2)(sin (time * 0.01133f), cos (time * 0.02521f));
262 266
263 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
264 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;
265 269
266 float3 colour = (float3)(z.x, z.y, z.x * z.y); 270 float3 colour = (float3)(z.x, z.y, atan2 (z.y, z.x));
267 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.));
268 } 272 }
269 EOF 273 EOF
270 274
271 my $prog = $ctx->build_program ($src); 275 my $prog = $ctx->build_program ($src);
272 my $kernel = $prog->kernel ("juliatunnel"); 276 my $kernel = $prog->kernel ("juliatunnel");
273 277
274 # program compiled, kernel ready, now draw and loop 278 # program compiled, kernel ready, now draw and loop
275 279
276 for (my $time; ; ++$time) { 280 for (my $time; ; ++$time) {
277 # acquire objects from opengl 281 # acquire objects from opengl
278 $queue->enqueue_acquire_gl_objects ([$tex]); 282 $queue->acquire_gl_objects ([$tex]);
279 283
280 # configure and run our kernel 284 # configure and run our kernel
281 $kernel->set_image2d (0, $tex); 285 $kernel->setf ("mf", $tex, $time*2); # mf = memory object, float
282 $kernel->set_float (1, $time);
283 $queue->enqueue_nd_range_kernel ($kernel, undef, [256, 256], undef); 286 $queue->nd_range_kernel ($kernel, undef, [$S, $S], undef);
284 287
285 # release objects to opengl again 288 # release objects to opengl again
286 $queue->enqueue_release_gl_objects ([$tex]); 289 $queue->release_gl_objects ([$tex]);
287 290
288 # wait 291 # wait
289 $queue->finish; 292 $queue->finish;
290 293
291 # now draw the texture, the defaults should be all right 294 # now draw the texture, the defaults should be all right
301 304
302 glXSwapBuffers; 305 glXSwapBuffers;
303 306
304 select undef, undef, undef, 1/60; 307 select undef, undef, undef, 1/60;
305 } 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>.
306 336
307=head1 DOCUMENTATION 337=head1 DOCUMENTATION
308 338
309=head2 BASIC CONVENTIONS 339=head2 BASIC CONVENTIONS
310 340
420 use AnyEvent; 450 use AnyEvent;
421 451
422 # initialise AnyEvent, by creating a watcher, or: 452 # initialise AnyEvent, by creating a watcher, or:
423 AnyEvent::detect; 453 AnyEvent::detect;
424 454
425 my $e = $queue->enqueue_marker; 455 my $e = $queue->marker;
426 $e->cb (sub { 456 $e->cb (sub {
427 warn "opencl is finished\n"; 457 warn "opencl is finished\n";
428 }) 458 })
429 459
430 main Gtk2; 460 main Gtk2;
447Asynchronously means that perl might execute your callbacks at any 477Asynchronously means that perl might execute your callbacks at any
448time. 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>),
449the 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
450callback: 480callback:
451 481
452 my $e = $queue->enqueue_marker; 482 my $e = $queue->marker;
453 my $flag; 483 my $flag;
454 $e->cb (sub { $flag = 1 }); 484 $e->cb (sub { $flag = 1 });
455 1 until $flag; 485 1 until $flag;
456 # $flag is now 1 486 # $flag is now 1
457 487
493=cut 523=cut
494 524
495package OpenCL; 525package OpenCL;
496 526
497use common::sense; 527use common::sense;
528use Carp ();
498use Async::Interrupt (); 529use Async::Interrupt ();
499 530
500our $POLL_FUNC; # set by XS 531our $POLL_FUNC; # set by XS
501 532
502BEGIN { 533BEGIN {
503 our $VERSION = '0.97'; 534 our $VERSION = '0.98';
504 535
505 require XSLoader; 536 require XSLoader;
506 XSLoader::load (__PACKAGE__, $VERSION); 537 XSLoader::load (__PACKAGE__, $VERSION);
507 538
508 @OpenCL::Platform::ISA = 539 @OpenCL::Platform::ISA =
526 @OpenCL::Image1D::ISA = 557 @OpenCL::Image1D::ISA =
527 @OpenCL::Image1DArray::ISA = 558 @OpenCL::Image1DArray::ISA =
528 @OpenCL::Image1DBuffer::ISA = OpenCL::Image::; 559 @OpenCL::Image1DBuffer::ISA = OpenCL::Image::;
529 560
530 @OpenCL::UserEvent::ISA = OpenCL::Event::; 561 @OpenCL::UserEvent::ISA = OpenCL::Event::;
562
563 @OpenCL::MappedBuffer::ISA =
564 @OpenCL::MappedImage::ISA = OpenCL::Mapped::;
531} 565}
532 566
533=head2 THE OpenCL PACKAGE 567=head2 THE OpenCL PACKAGE
534 568
535=over 4 569=over 4
537=item $int = OpenCL::errno 571=item $int = OpenCL::errno
538 572
539The 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
540and before calling another OpenCL function. 574and before calling another OpenCL function.
541 575
542=item $str = OpenCL::err2str $errval 576=item $str = OpenCL::err2str [$errval]
543 577
544Comverts 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).
545 580
546=item $str = OpenCL::enum2str $enum 581=item $str = OpenCL::enum2str $enum
547 582
548Converts most enum values (of parameter names, image format constants, 583Converts most enum values (of parameter names, image format constants,
549object types, addressing and filter modes, command types etc.) into a 584object types, addressing and filter modes, command types etc.) into a
1002=cut 1037=cut
1003 1038
1004sub OpenCL::Context::build_program { 1039sub OpenCL::Context::build_program {
1005 my ($self, $prog, $options) = @_; 1040 my ($self, $prog, $options) = @_;
1006 1041
1007 require Carp;
1008
1009 $prog = $self->program_with_source ($prog) 1042 $prog = $self->program_with_source ($prog)
1010 unless ref $prog; 1043 unless ref $prog;
1011 1044
1012 # 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
1013 for my $dev ($self->devices) { 1051 for my $dev ($self->devices) {
1014 eval { $prog->build ([$dev], $options); 1 } 1052 $prog->build_status ($dev) == BUILD_SUCCESS
1015 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"
1016 . $prog->build_log ($dev)); 1054 . $prog->build_log ($dev);
1017 } 1055 }
1018 1056
1019 $prog 1057 $prog
1020} 1058}
1021 1059
1045=item $buf = $ctx->buffer_sv ($flags, $data) 1083=item $buf = $ctx->buffer_sv ($flags, $data)
1046 1084
1047Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object and 1085Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object and
1048initialise it with the given data values. 1086initialise it with the given data values.
1049 1087
1050=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)
1051 1089
1052Creates a new OpenCL::Image object and optionally initialises it with 1090Creates a new OpenCL::Image object and optionally initialises it with
1053the given data values. 1091the given data values.
1054 1092
1055L<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>
1120 1158
1121Creates a new OpenCL::Program object from the given source code. 1159Creates a new OpenCL::Program object from the given source code.
1122 1160
1123L<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>
1124 1162
1163=item ($program, \@status) = $ctx->program_with_binary (\@devices, \@binaries)
1164
1165Creates a new OpenCL::Program object from the given binaries.
1166
1167L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateProgramWithBinary.html>
1168
1169Example: clone an existing program object that contains a successfully
1170compiled program, no matter how useless this is.
1171
1172 my $clone = $ctx->program_with_binary ([$prog->devices], [$prog->binaries]);
1173
1125=item $packed_value = $ctx->info ($name) 1174=item $packed_value = $ctx->info ($name)
1126 1175
1127See C<< $platform->info >> for details. 1176See C<< $platform->info >> for details.
1128 1177
1129L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetContextInfo.html> 1178L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetContextInfo.html>
1151=back 1200=back
1152 1201
1153=head2 THE OpenCL::Queue CLASS 1202=head2 THE OpenCL::Queue CLASS
1154 1203
1155An OpenCL::Queue represents an execution queue for OpenCL. You execute 1204An OpenCL::Queue represents an execution queue for OpenCL. You execute
1156requests by calling their respective C<enqueue_xxx> method and waitinf for 1205requests by calling their respective method and waiting for it to complete
1157it to complete in some way. 1206in some way.
1158 1207
1159All the enqueue methods return an event object that can be used to wait 1208Most methods that enqueue some request return an event object that can
1160for completion, unless the method is called in void context, in which case 1209be used to wait for completion (optionally using a callback), unless
1161no event object is created. 1210the method is called in void context, in which case no event object is
1211created.
1162 1212
1163They also allow you to specify any number of other event objects that this 1213They also allow you to specify any number of other event objects that this
1164request has to wait for before it starts executing, by simply passing the 1214request has to wait for before it starts executing, by simply passing the
1165event objects as extra parameters to the enqueue methods. To simplify 1215event objects as extra parameters to the enqueue methods. To simplify
1166program design, this module ignores any C<undef> values in the list of 1216program design, this module ignores any C<undef> values in the list of
1167events. This makes it possible to code operations such as this, without 1217events. This makes it possible to code operations such as this, without
1168having to put a valid event object into C<$event> first: 1218having to put a valid event object into C<$event> first:
1169 1219
1170 $event = $queue->enqueue_xxx (..., $event); 1220 $event = $queue->xxx (..., $event);
1171 1221
1172Queues execute in-order by default, without any parallelism, so in most 1222Queues execute in-order by default, without any parallelism, so in most
1173cases (i.e. you use only one queue) it's not necessary to wait for or 1223cases (i.e. you use only one queue) it's not necessary to wait for or
1174create event objects, althoguh an our of order queue is often a bit 1224create event objects, althoguh an our of order queue is often a bit
1175faster. 1225faster.
1176 1226
1177=over 4 1227=over 4
1178 1228
1179=item $ev = $queue->enqueue_read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...) 1229=item $ev = $queue->read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...)
1180 1230
1181Reads data from buffer into the given string. 1231Reads data from buffer into the given string.
1182 1232
1183L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBuffer.html> 1233L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBuffer.html>
1184 1234
1185=item $ev = $queue->enqueue_write_buffer ($buffer, $blocking, $offset, $data, $wait_events...) 1235=item $ev = $queue->write_buffer ($buffer, $blocking, $offset, $data, $wait_events...)
1186 1236
1187Writes data to buffer from the given string. 1237Writes data to buffer from the given string.
1188 1238
1189L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBuffer.html> 1239L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBuffer.html>
1190 1240
1191=item $ev = $queue->enqueue_copy_buffer ($src, $dst, $src_offset, $dst_offset, $len, $wait_events...) 1241=item $ev = $queue->copy_buffer ($src, $dst, $src_offset, $dst_offset, $len, $wait_events...)
1192 1242
1193L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBuffer.html> 1243L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBuffer.html>
1194 1244
1195=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...) 1245=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...)
1196 1246
1197http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBufferRect.html 1247http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBufferRect.html
1198 1248
1199=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...) 1249=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...)
1200 1250
1201http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBufferRect.html 1251http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBufferRect.html
1202 1252
1203=item $ev = $queue->enqueue_read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
1204
1205L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferRect.html>
1206
1207=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...) 1253=item $ev = $queue->copy_buffer_to_image ($src_buffer, $dst_image, $src_offset, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $wait_events...)
1254
1255L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>
1256
1257=item $ev = $queue->read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
1258
1259C<$row_pitch> (and C<$slice_pitch>) can be C<0>, in which case the OpenCL
1260module uses the image width (and height) to supply default values.
1208 1261
1209L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadImage.html> 1262L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadImage.html>
1210 1263
1211=item $ev = $queue->enqueue_write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...) 1264=item $ev = $queue->write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
1212 1265
1266C<$row_pitch> (and C<$slice_pitch>) can be C<0>, in which case the OpenCL
1267module uses the image width (and height) to supply default values.
1213L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteImage.html> 1268L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteImage.html>
1214 1269
1215=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...) 1270=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...)
1216 1271
1217L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html> 1272L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html>
1218 1273
1219=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...) 1274=item $ev = $queue->copy_image_to_buffer ($src_image, $dst_image, $src_x, $src_y, $src_z, $width, $height, $depth, $dst_offset, $wait_events...)
1220 1275
1221L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html> 1276L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html>
1222 1277
1223=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...) 1278=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...)
1224 1279
1225Yeah. 1280Yeah.
1226 1281
1227L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>. 1282L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>.
1228 1283
1229=item $ev = $queue->enqueue_fill_buffer ($mem, $pattern, $offset, $size, ...) 1284=item $ev = $queue->fill_buffer ($mem, $pattern, $offset, $size, ...)
1230 1285
1231Fills the given buffer object with repeated applications of C<$pattern>, 1286Fills the given buffer object with repeated applications of C<$pattern>,
1232starting at C<$offset> for C<$size> octets. 1287starting at C<$offset> for C<$size> octets.
1233 1288
1234L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillBuffer.html> 1289L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillBuffer.html>
1235 1290
1236=item $ev = $queue->enqueue_fill_image ($img, $r, $g, $b, $a, $x, $y, $z, $width, $height, $depth, ...) 1291=item $ev = $queue->fill_image ($img, $r, $g, $b, $a, $x, $y, $z, $width, $height, $depth, ...)
1237 1292
1238Fills the given image area with the given rgba colour components. The 1293Fills the given image area with the given rgba colour components. The
1239components are normally floating point values between C<0> and C<1>, 1294components are normally floating point values between C<0> and C<1>,
1240except when the image channel data type is a signe dor unsigned 1295except when the image channel data type is a signe dor unsigned
1241unnormalised format, in which case the range is determined by the format. 1296unnormalised format, in which case the range is determined by the format.
1242 1297
1243L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillImage.html> 1298L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillImage.html>
1244 1299
1245=item $ev = $queue->enqueue_task ($kernel, $wait_events...) 1300=item $ev = $queue->task ($kernel, $wait_events...)
1246 1301
1247L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html> 1302L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html>
1248 1303
1249=item $ev = $queue->enqueue_nd_range_kernel ($kernel, \@global_work_offset, \@global_work_size, \@local_work_size, $wait_events...) 1304=item $ev = $queue->nd_range_kernel ($kernel, \@global_work_offset, \@global_work_size, \@local_work_size, $wait_events...)
1250 1305
1251Enqueues a kernel execution. 1306Enqueues a kernel execution.
1252 1307
1253\@global_work_size must be specified as a reference to an array of 1308\@global_work_size must be specified as a reference to an array of
1254integers specifying the work sizes (element counts). 1309integers specifying the work sizes (element counts).
1262reference to an array of local work sizes, with the same number of 1317reference to an array of local work sizes, with the same number of
1263elements as \@global_work_size. 1318elements as \@global_work_size.
1264 1319
1265L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html> 1320L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html>
1266 1321
1267=item $ev = $queue->enqueue_acquire_gl_objects ([object, ...], $wait_events...) 1322=item $ev = $queue->acquire_gl_objects ([object, ...], $wait_events...)
1268 1323
1269Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired 1324Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired
1270for subsequent OpenCL usage. 1325for subsequent OpenCL usage.
1271 1326
1272L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueAcquireGLObjects.html> 1327L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueAcquireGLObjects.html>
1273 1328
1274=item $ev = $queue->enqueue_release_gl_objects ([object, ...], $wait_events...) 1329=item $ev = $queue->release_gl_objects ([object, ...], $wait_events...)
1275 1330
1276Enqueues a list (an array-ref of OpenCL::Memory objects) to be released 1331Enqueues a list (an array-ref of OpenCL::Memory objects) to be released
1277for subsequent OpenGL usage. 1332for subsequent OpenGL usage.
1278 1333
1279L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReleaseGLObjects.html> 1334L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReleaseGLObjects.html>
1280 1335
1281=item $ev = $queue->enqueue_wait_for_events ($wait_events...) 1336=item $ev = $queue->wait_for_events ($wait_events...)
1282 1337
1283L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html> 1338L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html>
1284 1339
1285=item $ev = $queue->enqueue_marker ($wait_events...) 1340=item $ev = $queue->marker ($wait_events...)
1286 1341
1287L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMarkerWithWaitList.html> 1342L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMarkerWithWaitList.html>
1288 1343
1289=item $ev = $queue->enqueue_barrier ($wait_events...) 1344=item $ev = $queue->barrier ($wait_events...)
1290 1345
1291L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueBarrierWithWaitList.html> 1346L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueBarrierWithWaitList.html>
1292 1347
1293=item $queue->flush 1348=item $queue->flush
1294 1349
1321=item $command_queue_properties = $command_queue->properties 1376=item $command_queue_properties = $command_queue->properties
1322 1377
1323Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_PROPERTIES> and returns the result. 1378Calls C<clGetCommandQueueInfo> with C<CL_QUEUE_PROPERTIES> and returns the result.
1324 1379
1325=for gengetinfo end command_queue 1380=for gengetinfo end command_queue
1381
1382=back
1383
1384=head3 MEMORY MAPPED BUFFERS
1385
1386OpenCL allows you to map buffers and images to host memory (read: perl
1387scalars). This is done much like reading or copying a buffer, by enqueuing
1388a map or unmap operation on the command queue.
1389
1390The map operations return an C<OpenCL::Mapped> object - see L<THE
1391OpenCL::Mapped CLASS> section for details on what to do with these
1392objects.
1393
1394The object will be unmapped automatically when the mapped object is
1395destroyed (you can use a barrier to make sure the unmap has finished,
1396before using the buffer in a kernel), but you can also enqueue an unmap
1397operation manually.
1398
1399=over 4
1400
1401=item $mapped_buffer = $queue->map_buffer ($buf, $blocking=1, $map_flags=OpenCL::MAP_READ|OpenCL::MAP_WRITE, $offset=0, $size=undef, $wait_events...)
1402
1403Maps the given buffer into host memory and returns an
1404C<OpenCL::MappedBuffer> object. If C<$size> is specified as undef, then
1405the map will extend to the end of the buffer.
1406
1407L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMapBuffer.html>
1408
1409Example: map the buffer $buf fully and replace the first 4 bytes by "abcd", then unmap.
1410
1411 {
1412 my $mapped = $queue->map_buffer ($buf, 1, OpenCL::MAP_WRITE);
1413 substr $$mapped, 0, 4, "abcd";
1414 } # asynchronously unmap because $mapped is destroyed
1415
1416=item $mapped_image = $queue->map_image ($img, $blocking=1, $map_flags=OpenCL::MAP_READ|OpenCL::MAP_WRITE, $x=0, $y=0, $z=0, $width=undef, $height=undef, $depth=undef, $wait_events...)
1417
1418Maps the given image area into host memory and return an
1419C<OpenCL::MappedImage> object.
1420
1421If any of C<$width>, C<$height> and/or C<$depth> are C<undef> then they
1422will be replaced by the maximum possible value.
1423
1424L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMapImage.html>
1425
1426Example: map an image (with OpenCL::UNSIGNED_INT8 channel type) and set
1427the first channel of the leftmost column to 5, then explicitly unmap
1428it. You are not necessarily meant to do it this way, this example just
1429shows you the accessors to use :)
1430
1431 my $mapped = $queue->map_image ($image, 1, OpenCL::MAP_WRITE);
1432
1433 $mapped->set ($_ * $mapped->row_pitch, pack "C", 5)
1434 for 0..$image->height;
1435
1436 $mapped->unmap;.
1437 $mapped->wait; # only needed for out of order queues normally
1438
1439=item $ev = $queue->unmap ($mapped, $wait_events...)
1440
1441Unmaps the data from host memory. You must not call any methods that
1442modify the data, or modify the data scalar directly, after calling this
1443method.
1444
1445The mapped event object will always be passed as part of the
1446$wait_events. The mapped event object will be replaced by the new event
1447object that this request creates.
1326 1448
1327=back 1449=back
1328 1450
1329=head2 THE OpenCL::Memory CLASS 1451=head2 THE OpenCL::Memory CLASS
1330 1452
1523If a callback is specified, then it will be called when compilation is 1645If a callback is specified, then it will be called when compilation is
1524finished. Note that many OpenCL implementations block your program while 1646finished. Note that many OpenCL implementations block your program while
1525compiling whether you use a callback or not. See C<build_async> if you 1647compiling whether you use a callback or not. See C<build_async> if you
1526want to make sure the build is done in the background. 1648want to make sure the build is done in the background.
1527 1649
1528Note that some OpenCL implementations atc up badly, and don't call the 1650Note that some OpenCL implementations act up badly, and don't call the
1529callback in some error cases (but call it in others). This implementation 1651callback in some error cases (but call it in others). This implementation
1530assumes the callback will always be called, and leaks memory if this is 1652assumes the callback will always be called, and leaks memory if this is
1531not so. So best make sure you don't pass in invalid values. 1653not so. So best make sure you don't pass in invalid values.
1532 1654
1655Some implementations fail with C<OpenCL::INVALID_BINARY> when the
1656compilation state is successful but some later stage fails.
1657
1533L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clBuildProgram.html> 1658L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clBuildProgram.html>
1534 1659
1535=item $program->build_async (\@devices = undef, $options = "", $cb->($program) = undef) 1660=item $program->build_async (\@devices = undef, $options = "", $cb->($program) = undef)
1536 1661
1537Similar to C<< ->build >>, except it starts a thread, and never fails (you 1662Similar to C<< ->build >>, except it starts a thread, and never fails (you
1685 1810
1686Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PRIVATE_MEM_SIZE> and returns the result. 1811Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PRIVATE_MEM_SIZE> and returns the result.
1687 1812
1688=for gengetinfo end kernel_work_group 1813=for gengetinfo end kernel_work_group
1689 1814
1815=item $kernel->setf ($format, ...)
1816
1817Sets the arguments of a kernel. Since OpenCL 1.1 doesn't have a generic
1818way to set arguments (and with OpenCL 1.2 it might be rather slow), you
1819need to specify a format argument, much as with C<printf>, to tell OpenCL
1820what type of argument it is.
1821
1822The format arguments are single letters:
1823
1824 c char
1825 C unsigned char
1826 s short
1827 S unsigned short
1828 i int
1829 I unsigned int
1830 l long
1831 L unsigned long
1832
1833 h half float (0..65535)
1834 f float
1835 d double
1836
1837 z local (octet size)
1838
1839 m memory object (buffer or image)
1840 a sampler
1841 e event
1842
1843Space characters in the format string are ignored.
1844
1845Example: set the arguments for a kernel that expects an int, two floats, a buffer and an image.
1846
1847 $kernel->setf ("i ff mm", 5, 0.5, 3, $buffer, $image);
1848
1690=item $kernel->set_TYPE ($index, $value) 1849=item $kernel->set_TYPE ($index, $value)
1691 1850
1851=item $kernel->set_char ($index, $value)
1852
1853=item $kernel->set_uchar ($index, $value)
1854
1855=item $kernel->set_short ($index, $value)
1856
1857=item $kernel->set_ushort ($index, $value)
1858
1859=item $kernel->set_int ($index, $value)
1860
1861=item $kernel->set_uint ($index, $value)
1862
1863=item $kernel->set_long ($index, $value)
1864
1865=item $kernel->set_ulong ($index, $value)
1866
1867=item $kernel->set_half ($index, $value)
1868
1869=item $kernel->set_float ($index, $value)
1870
1871=item $kernel->set_double ($index, $value)
1872
1873=item $kernel->set_memory ($index, $value)
1874
1875=item $kernel->set_buffer ($index, $value)
1876
1877=item $kernel->set_image ($index, $value)
1878
1879=item $kernel->set_sampler ($index, $value)
1880
1881=item $kernel->set_local ($index, $value)
1882
1883=item $kernel->set_event ($index, $value)
1884
1692This is a family of methods to set the kernel argument with the number C<$index> to the give C<$value>. 1885This is a family of methods to set the kernel argument with the number
1693 1886C<$index> to the give C<$value>.
1694TYPE is one of C<char>, C<uchar>, C<short>, C<ushort>, C<int>, C<uint>,
1695C<long>, C<ulong>, C<half>, C<float>, C<double>, C<memory>, C<buffer>,
1696C<image2d>, C<image3d>, C<sampler>, C<local> or C<event>.
1697 1887
1698Chars and integers (including the half type) are specified as integers, 1888Chars and integers (including the half type) are specified as integers,
1699float and double as floating point values, memory/buffer/image2d/image3d 1889float and double as floating point values, memory/buffer/image must be
1700must be an object of that type or C<undef>, local-memory arguments are 1890an object of that type or C<undef>, local-memory arguments are set by
1701set by specifying the size, and sampler and event must be objects of that 1891specifying the size, and sampler and event must be objects of that type.
1702type. 1892
1893Note that C<set_memory> works for all memory objects (all types of buffers
1894and images) - the main purpose of the more specific C<set_TYPE> functions
1895is type checking.
1703 1896
1704Setting an argument for a kernel does NOT keep a reference to the object - 1897Setting an argument for a kernel does NOT keep a reference to the object -
1705for example, if you set an argument to some image object, free the image, 1898for example, if you set an argument to some image object, free the image,
1706and call the kernel, you will run into undefined behaviour. 1899and call the kernel, you will run into undefined behaviour.
1707 1900
1803 1996
1804L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetUserEventStatus.html> 1997L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetUserEventStatus.html>
1805 1998
1806=back 1999=back
1807 2000
2001=head2 THE OpenCL::Mapped CLASS
2002
2003This class represents objects mapped into host memory. They are
2004represented by a blessed string scalar. The string data is the mapped
2005memory area, that is, if you read or write it, then the mapped object is
2006accessed directly.
2007
2008You must only ever use operations that modify the string in-place - for
2009example, a C<substr> that doesn't change the length, or maybe a regex that
2010doesn't change the length. Any other operation might cause the data to be
2011copied.
2012
2013When the object is destroyed it will enqueue an implicit unmap operation
2014on the queue that was used to create it.
2015
2016Keep in mind that you I<need> to unmap (or destroy) mapped objects before
2017OpenCL sees the changes, even if some implementations don't need this
2018sometimes.
2019
2020Example, replace the first two floats in the mapped buffer by 1 and 2.
2021
2022 my $mapped = $queue->map_buffer ($buf, ...
2023 $mapped->event->wait; # make sure it's there
2024
2025 # now replace first 8 bytes by new data, which is exactly 8 bytes long
2026 # we blindly assume device endianness to equal host endianness
2027 # (and of course, we assume iee 754 single precision floats :)
2028 substr $$mapped, 0, 8, pack "f*", 1, 2;
2029
2030=over 4
2031
2032=item $ev = $mapped->unmap ($wait_events...)
2033
2034Unmaps the mapped memory object, using the queue originally used to create
2035it, quite similarly to C<< $queue->unmap ($mapped, ...) >>.
2036
2037=item $bool = $mapped->mapped
2038
2039Returns whether the object is still mapped - true before an C<unmap> is
2040enqueued, false afterwards.
2041
2042=item $ev = $mapped->event
2043
2044Return the event object associated with the mapped object. Initially, this
2045will be the event object created when mapping the object, and after an
2046unmap, this will be the event object that the unmap operation created.
2047
2048=item $mapped->wait
2049
2050Same as C<< $mapped->event->wait >> - makes sure no operations on this
2051mapped object are outstanding.
2052
2053=item $bytes = $mapped->size
2054
2055Returns the size of the mapped area, in bytes. Same as C<length $$mapped>.
2056
2057=item $ptr = $mapped->ptr
2058
2059Returns the raw memory address of the mapped area.
2060
2061=item $mapped->set ($offset, $data)
2062
2063Replaces the data at the given C<$offset> in the memory area by the new
2064C<$data>. This method is safer than direct manipulation of C<$mapped>
2065because it does bounds-checking, but also slower.
2066
2067=item $data = $mapped->get ($offset, $length)
2068
2069Returns (without copying) a scalar representing the data at the given
2070C<$offset> and C<$length> in the mapped memory area. This is the same as
2071the following substr, except much slower;
2072
2073 $data = substr $$mapped, $offset, $length
2074
2075=cut
2076
2077sub OpenCL::Mapped::get {
2078 substr ${$_[0]}, $_[1], $_[2]
2079}
2080
2081=back
2082
2083=head2 THE OpenCL::MappedBuffer CLASS
2084
2085This is a subclass of OpenCL::Mapped, representing mapped buffers.
2086
2087=head2 THE OpenCL::MappedImage CLASS
2088
2089This is a subclass of OpenCL::Mapped, representing mapped images.
2090
2091=over 4
2092
2093=item $bytes = $mapped->row_pitch
2094
2095=item $bytes = $mapped->slice_pitch
2096
2097Return the row or slice pitch of the image that has been mapped.
2098
2099=back
2100
2101
1808=cut 2102=cut
1809 2103
18101; 21041;
1811 2105
1812=head1 AUTHOR 2106=head1 AUTHOR

Diff Legend

Removed lines
+ Added lines
< Changed lines
> Changed lines