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

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.58 by root, Wed Apr 25 20:29:03 2012 UTC vs.
Revision 1.64 by root, Mon Apr 30 23:12:43 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
… …
420 use AnyEvent; 423 use AnyEvent;
421 424
422 # initialise AnyEvent, by creating a watcher, or: 425 # initialise AnyEvent, by creating a watcher, or:
423 AnyEvent::detect; 426 AnyEvent::detect;
424 427
425 my $e = $queue->enqueue_marker; 428 my $e = $queue->marker;
426 $e->cb (sub { 429 $e->cb (sub {
427 warn "opencl is finished\n"; 430 warn "opencl is finished\n";
428 }) 431 })
429 432
430 main Gtk2; 433 main Gtk2;
… …
447Asynchronously means that perl might execute your callbacks at any 450Asynchronously means that perl might execute your callbacks at any
448time. For example, in the following code (I<THAT YOU SHOULD NOT COPY>), 451time. 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 452the C<until> loop following the marker call will be interrupted by the
450callback: 453callback:
451 454
452 my $e = $queue->enqueue_marker; 455 my $e = $queue->marker;
453 my $flag; 456 my $flag;
454 $e->cb (sub { $flag = 1 }); 457 $e->cb (sub { $flag = 1 });
455 1 until $flag; 458 1 until $flag;
456 # $flag is now 1 459 # $flag is now 1
457 460
… …
493=cut 496=cut
494 497
495package OpenCL; 498package OpenCL;
496 499
497use common::sense; 500use common::sense;
501use Carp ();
498use Async::Interrupt (); 502use Async::Interrupt ();
499 503
500our $POLL_FUNC; # set by XS 504our $POLL_FUNC; # set by XS
501 505
502BEGIN { 506BEGIN {
503 our $VERSION = '0.97'; 507 our $VERSION = '0.98';
504 508
505 require XSLoader; 509 require XSLoader;
506 XSLoader::load (__PACKAGE__, $VERSION); 510 XSLoader::load (__PACKAGE__, $VERSION);
507 511
508 @OpenCL::Platform::ISA = 512 @OpenCL::Platform::ISA =
… …
537=item $int = OpenCL::errno 541=item $int = OpenCL::errno
538 542
539The last error returned by a function - it's only valid after an error occured 543The last error returned by a function - it's only valid after an error occured
540and before calling another OpenCL function. 544and before calling another OpenCL function.
541 545
542=item $str = OpenCL::err2str $errval 546=item $str = OpenCL::err2str [$errval]
543 547
544Comverts an error value into a human readable string. 548Converts an error value into a human readable string. IF no error value is
549given, then the last error will be used (as returned by OpenCL::errno).
545 550
546=item $str = OpenCL::enum2str $enum 551=item $str = OpenCL::enum2str $enum
547 552
548Converts most enum values (of parameter names, image format constants, 553Converts most enum values (of parameter names, image format constants,
549object types, addressing and filter modes, command types etc.) into a 554object types, addressing and filter modes, command types etc.) into a
… …
1002=cut 1007=cut
1003 1008
1004sub OpenCL::Context::build_program { 1009sub OpenCL::Context::build_program {
1005 my ($self, $prog, $options) = @_; 1010 my ($self, $prog, $options) = @_;
1006 1011
1007 require Carp;
1008
1009 $prog = $self->program_with_source ($prog) 1012 $prog = $self->program_with_source ($prog)
1010 unless ref $prog; 1013 unless ref $prog;
1011 1014
1012 # we build separately per device so we instantly know which one failed 1015 eval { $prog->build (undef, $options); 1 }
1016 or errno == BUILD_PROGRAM_FAILURE
1017 or errno == INVALID_BINARY # workaround nvidia bug
1018 or Carp::croak "OpenCL::Context->build_program: " . err2str;
1019
1020 # we check status for all devices
1013 for my $dev ($self->devices) { 1021 for my $dev ($self->devices) {
1014 eval { $prog->build ([$dev], $options); 1 } 1022 $prog->build_status ($dev) == BUILD_SUCCESS
1015 or Carp::croak ("Building OpenCL program for device '" . $dev->name . "' failed:\n" 1023 or Carp::croak "Building OpenCL program for device '" . $dev->name . "' failed:\n"
1016 . $prog->build_log ($dev)); 1024 . $prog->build_log ($dev);
1017 } 1025 }
1018 1026
1019 $prog 1027 $prog
1020} 1028}
1021 1029
… …
1045=item $buf = $ctx->buffer_sv ($flags, $data) 1053=item $buf = $ctx->buffer_sv ($flags, $data)
1046 1054
1047Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object and 1055Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object and
1048initialise it with the given data values. 1056initialise it with the given data values.
1049 1057
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) 1058=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 1059
1052Creates a new OpenCL::Image object and optionally initialises it with 1060Creates a new OpenCL::Image object and optionally initialises it with
1053the given data values. 1061the given data values.
1054 1062
1055L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateImage.html> 1063L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateImage.html>
… …
1151=back 1159=back
1152 1160
1153=head2 THE OpenCL::Queue CLASS 1161=head2 THE OpenCL::Queue CLASS
1154 1162
1155An OpenCL::Queue represents an execution queue for OpenCL. You execute 1163An OpenCL::Queue represents an execution queue for OpenCL. You execute
1156requests by calling their respective C<enqueue_xxx> method and waitinf for 1164requests by calling their respective method and waiting for it to complete
1157it to complete in some way. 1165in some way.
1158 1166
1159All the enqueue methods return an event object that can be used to wait 1167Most methods that enqueue some request return an event object that can
1160for completion, unless the method is called in void context, in which case 1168be used to wait for completion (optionally using a callback), unless
1161no event object is created. 1169the method is called in void context, in which case no event object is
1170created.
1162 1171
1163They also allow you to specify any number of other event objects that this 1172They 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 1173request has to wait for before it starts executing, by simply passing the
1165event objects as extra parameters to the enqueue methods. To simplify 1174event objects as extra parameters to the enqueue methods. To simplify
1166program design, this module ignores any C<undef> values in the list of 1175program design, this module ignores any C<undef> values in the list of
1167events. This makes it possible to code operations such as this, without 1176events. This makes it possible to code operations such as this, without
1168having to put a valid event object into C<$event> first: 1177having to put a valid event object into C<$event> first:
1169 1178
1170 $event = $queue->enqueue_xxx (..., $event); 1179 $event = $queue->xxx (..., $event);
1171 1180
1172Queues execute in-order by default, without any parallelism, so in most 1181Queues 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 1182cases (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 1183create event objects, althoguh an our of order queue is often a bit
1175faster. 1184faster.
1176 1185
1177=over 4 1186=over 4
1178 1187
1179=item $ev = $queue->enqueue_read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...) 1188=item $ev = $queue->read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...)
1180 1189
1181Reads data from buffer into the given string. 1190Reads data from buffer into the given string.
1182 1191
1183L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBuffer.html> 1192L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBuffer.html>
1184 1193
1185=item $ev = $queue->enqueue_write_buffer ($buffer, $blocking, $offset, $data, $wait_events...) 1194=item $ev = $queue->write_buffer ($buffer, $blocking, $offset, $data, $wait_events...)
1186 1195
1187Writes data to buffer from the given string. 1196Writes data to buffer from the given string.
1188 1197
1189L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBuffer.html> 1198L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBuffer.html>
1190 1199
1191=item $ev = $queue->enqueue_copy_buffer ($src, $dst, $src_offset, $dst_offset, $len, $wait_events...) 1200=item $ev = $queue->copy_buffer ($src, $dst, $src_offset, $dst_offset, $len, $wait_events...)
1192 1201
1193L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBuffer.html> 1202L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBuffer.html>
1194 1203
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...) 1204=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 1205
1197http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBufferRect.html 1206http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadBufferRect.html
1198 1207
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...) 1208=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 1209
1201http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBufferRect.html 1210http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteBufferRect.html
1202 1211
1203=item $ev = $queue->enqueue_read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...) 1212=item $ev = $queue->read_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
1204 1213
1205L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferRect.html> 1214L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferRect.html>
1206 1215
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...) 1216=item $ev = $queue->copy_buffer_to_image ($src_buffer, $dst_image, $src_offset, $dst_x, $dst_y, $dst_z, $width, $height, $depth, $wait_events...)
1208 1217
1209L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadImage.html> 1218L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReadImage.html>
1210 1219
1211=item $ev = $queue->enqueue_write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...) 1220=item $ev = $queue->write_image ($src, $blocking, $x, $y, $z, $width, $height, $depth, $row_pitch, $slice_pitch, $data, $wait_events...)
1212 1221
1213L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteImage.html> 1222L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWriteImage.html>
1214 1223
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...) 1224=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 1225
1217L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html> 1226L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImage.html>
1218 1227
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...) 1228=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 1229
1221L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html> 1230L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyImageToBuffer.html>
1222 1231
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...) 1232=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 1233
1225Yeah. 1234Yeah.
1226 1235
1227L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>. 1236L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueCopyBufferToImage.html>.
1228 1237
1229=item $ev = $queue->enqueue_fill_buffer ($mem, $pattern, $offset, $size, ...) 1238=item $ev = $queue->fill_buffer ($mem, $pattern, $offset, $size, ...)
1230 1239
1231Fills the given buffer object with repeated applications of C<$pattern>, 1240Fills the given buffer object with repeated applications of C<$pattern>,
1232starting at C<$offset> for C<$size> octets. 1241starting at C<$offset> for C<$size> octets.
1233 1242
1234L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillBuffer.html> 1243L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillBuffer.html>
1235 1244
1236=item $ev = $queue->enqueue_fill_image ($img, $r, $g, $b, $a, $x, $y, $z, $width, $height, $depth, ...) 1245=item $ev = $queue->fill_image ($img, $r, $g, $b, $a, $x, $y, $z, $width, $height, $depth, ...)
1237 1246
1238Fills the given image area with the given rgba colour components. The 1247Fills the given image area with the given rgba colour components. The
1239components are normally floating point values between C<0> and C<1>, 1248components are normally floating point values between C<0> and C<1>,
1240except when the image channel data type is a signe dor unsigned 1249except when the image channel data type is a signe dor unsigned
1241unnormalised format, in which case the range is determined by the format. 1250unnormalised format, in which case the range is determined by the format.
1242 1251
1243L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillImage.html> 1252L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueFillImage.html>
1244 1253
1245=item $ev = $queue->enqueue_task ($kernel, $wait_events...) 1254=item $ev = $queue->task ($kernel, $wait_events...)
1246 1255
1247L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html> 1256L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueTask.html>
1248 1257
1249=item $ev = $queue->enqueue_nd_range_kernel ($kernel, \@global_work_offset, \@global_work_size, \@local_work_size, $wait_events...) 1258=item $ev = $queue->nd_range_kernel ($kernel, \@global_work_offset, \@global_work_size, \@local_work_size, $wait_events...)
1250 1259
1251Enqueues a kernel execution. 1260Enqueues a kernel execution.
1252 1261
1253\@global_work_size must be specified as a reference to an array of 1262\@global_work_size must be specified as a reference to an array of
1254integers specifying the work sizes (element counts). 1263integers specifying the work sizes (element counts).
… …
1262reference to an array of local work sizes, with the same number of 1271reference to an array of local work sizes, with the same number of
1263elements as \@global_work_size. 1272elements as \@global_work_size.
1264 1273
1265L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html> 1274L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html>
1266 1275
1267=item $ev = $queue->enqueue_acquire_gl_objects ([object, ...], $wait_events...) 1276=item $ev = $queue->acquire_gl_objects ([object, ...], $wait_events...)
1268 1277
1269Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired 1278Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired
1270for subsequent OpenCL usage. 1279for subsequent OpenCL usage.
1271 1280
1272L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueAcquireGLObjects.html> 1281L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueAcquireGLObjects.html>
1273 1282
1274=item $ev = $queue->enqueue_release_gl_objects ([object, ...], $wait_events...) 1283=item $ev = $queue->release_gl_objects ([object, ...], $wait_events...)
1275 1284
1276Enqueues a list (an array-ref of OpenCL::Memory objects) to be released 1285Enqueues a list (an array-ref of OpenCL::Memory objects) to be released
1277for subsequent OpenGL usage. 1286for subsequent OpenGL usage.
1278 1287
1279L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReleaseGLObjects.html> 1288L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReleaseGLObjects.html>
1280 1289
1281=item $ev = $queue->enqueue_wait_for_events ($wait_events...) 1290=item $ev = $queue->wait_for_events ($wait_events...)
1282 1291
1283L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html> 1292L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html>
1284 1293
1285=item $ev = $queue->enqueue_marker ($wait_events...) 1294=item $ev = $queue->marker ($wait_events...)
1286 1295
1287L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMarkerWithWaitList.html> 1296L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMarkerWithWaitList.html>
1288 1297
1289=item $ev = $queue->enqueue_barrier ($wait_events...) 1298=item $ev = $queue->barrier ($wait_events...)
1290 1299
1291L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueBarrierWithWaitList.html> 1300L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueBarrierWithWaitList.html>
1292 1301
1293=item $queue->flush 1302=item $queue->flush
1294 1303
… …
1523If a callback is specified, then it will be called when compilation is 1532If a callback is specified, then it will be called when compilation is
1524finished. Note that many OpenCL implementations block your program while 1533finished. Note that many OpenCL implementations block your program while
1525compiling whether you use a callback or not. See C<build_async> if you 1534compiling whether you use a callback or not. See C<build_async> if you
1526want to make sure the build is done in the background. 1535want to make sure the build is done in the background.
1527 1536
1528Note that some OpenCL implementations atc up badly, and don't call the 1537Note that some OpenCL implementations act up badly, and don't call the
1529callback in some error cases (but call it in others). This implementation 1538callback in some error cases (but call it in others). This implementation
1530assumes the callback will always be called, and leaks memory if this is 1539assumes 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. 1540not so. So best make sure you don't pass in invalid values.
1532 1541
1542Some implementations fail with C<OpenCL::INVALID_BINARY> when the
1543compilation state is successful but some later stage fails.
1544
1533L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clBuildProgram.html> 1545L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clBuildProgram.html>
1534 1546
1535=item $program->build_async (\@devices = undef, $options = "", $cb->($program) = undef) 1547=item $program->build_async (\@devices = undef, $options = "", $cb->($program) = undef)
1536 1548
1537Similar to C<< ->build >>, except it starts a thread, and never fails (you 1549Similar to C<< ->build >>, except it starts a thread, and never fails (you
… …
1684=item $ulong = $kernel->private_mem_size ($device) 1696=item $ulong = $kernel->private_mem_size ($device)
1685 1697
1686Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PRIVATE_MEM_SIZE> and returns the result. 1698Calls C<clGetKernelWorkGroupInfo> with C<CL_KERNEL_PRIVATE_MEM_SIZE> and returns the result.
1687 1699
1688=for gengetinfo end kernel_work_group 1700=for gengetinfo end kernel_work_group
1701
1702=item $kernel->setf ($format, ...)
1703
1704Sets the arguments of a kernel. Since OpenCL 1.1 doesn't have a generic
1705way to set arguments (and with OpenCL 1.2 it might be rather slow), you
1706need to specify a format argument, much as with C<printf>, to tell OpenCL
1707what type of argument it is.
1708
1709The format arguments are single letters:
1710
1711 c char
1712 C unsigned char
1713 s short
1714 S unsigned short
1715 i int
1716 I unsigned int
1717 l long
1718 L unsigned long
1719
1720 h half float (0..65535)
1721 f float
1722 d double
1723
1724 z local (octet size)
1725
1726 m memory object (buffer or image)
1727 a sampler
1728 e event
1729
1730Space characters in the format string are ignored.
1731
1732Example: set the arguments for a kernel that expects an int, two floats, a buffer and an image.
1733
1734 $kernel->setf ("i ff mm", 5, 0.5, 3, $buffer, $image);
1689 1735
1690=item $kernel->set_TYPE ($index, $value) 1736=item $kernel->set_TYPE ($index, $value)
1691 1737
1692=item $kernel->set_char ($index, $value) 1738=item $kernel->set_char ($index, $value)
1693 1739

Diff Legend

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