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

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.37 by root, Thu Apr 19 14:36:46 2012 UTC vs.
Revision 1.48 by root, Sat Apr 21 19:53:29 2012 UTC

… …
160 '; 160 ';
161 161
162 my $prog = $ctx->program_with_source ($src); 162 my $prog = $ctx->program_with_source ($src);
163 163
164 # build croaks on compile errors, so catch it and print the compile errors 164 # build croaks on compile errors, so catch it and print the compile errors
165 eval { $prog->build ($dev); 1 } 165 eval { $prog->build ($dev, "-cl-fast-relaxed-math"); 1 }
166 or die $prog->build_log; 166 or die $prog->build_log;
167 167
168 my $kernel = $prog->kernel ("squareit"); 168 my $kernel = $prog->kernel ("squareit");
169 169
170=head2 Create some input and output float buffers, then call the 170=head2 Create some input and output float buffers, then call the
… …
210 # enqueue an async read 210 # enqueue an async read
211 $ev = $queue->enqueue_read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data, $ev); 211 $ev = $queue->enqueue_read_buffer ($output, 0, 0, OpenCL::SIZEOF_FLOAT * 4, my $data, $ev);
212 212
213 # wait for the last event to complete 213 # wait for the last event to complete
214 $ev->wait; 214 $ev->wait;
215
216=head2 Use the OpenGL module to share a texture between OpenCL and OpenGL and draw some julia
217set tunnel effect.
218
219This is quite a long example to get you going.
220
221 use OpenGL ":all";
222 use OpenCL;
223
224 # open a window and create a gl texture
225 OpenGL::glpOpenWindow width => 256, height => 256;
226 my $texid = glGenTextures_p 1;
227 glBindTexture GL_TEXTURE_2D, $texid;
228 glTexImage2D_c GL_TEXTURE_2D, 0, GL_RGBA8, 256, 256, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0;
229
230 # find and use the first opencl device that let's us get a shared opengl context
231 my $platform;
232 my $dev;
233 my $ctx;
234
235 for (OpenCL::platforms) {
236 $platform = $_;
237 for ($platform->devices) {
238 $dev = $_;
239 $ctx = $platform->context ([OpenCL::GLX_DISPLAY_KHR, undef, OpenCL::GL_CONTEXT_KHR, undef], [$dev])
240 and last;
241 }
242 }
243
244 $ctx
245 or die "cannot find suitable OpenCL device\n";
246
247 my $queue = $ctx->queue ($dev);
248
249 # now attach an opencl image2d object to the opengl texture
250 my $tex = $ctx->gl_texture2d (OpenCL::MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, $texid);
251
252 # now the boring opencl code
253 my $src = <<EOF;
254 kernel void
255 juliatunnel (write_only image2d_t img, float time)
256 {
257 float2 p = (float2)(get_global_id (0), get_global_id (1)) / 256.f * 2.f - 1.f;
258
259 float2 m = (float2)(1.f, p.y) / fabs (p.x);
260 m.x = fabs (fmod (m.x + time * 0.05f, 4.f)) - 2.f;
261
262 float2 z = m;
263 float2 c = (float2)(sin (time * 0.05005), cos (time * 0.06001));
264
265 for (int i = 0; i < 25 && dot (z, z) < 4.f; ++i)
266 z = (float2)(z.x * z.x - z.y * z.y, 2.f * z.x * z.y) + c;
267
268 float3 colour = (float3)(z.x, z.y, z.x * z.y);
269 write_imagef (img, (int2)(get_global_id (0), get_global_id (1)), (float4)(colour * p.x * p.x, 1.));
270 }
271 EOF
272 my $prog = $ctx->program_with_source ($src);
273 eval { $prog->build ($dev); 1 }
274 or die $prog->build_log ($dev);
275
276 my $kernel = $prog->kernel ("juliatunnel");
277
278 # program compiled, kernel ready, now draw and loop
279
280 for (my $time; ; ++$time) {
281 # acquire objects from opengl
282 $queue->enqueue_acquire_gl_objects ([$tex]);
283
284 # configure and run our kernel
285 $kernel->set_image2d (0, $tex);
286 $kernel->set_float (1, $time);
287 $queue->enqueue_nd_range_kernel ($kernel, undef, [256, 256], undef);
288
289 # release objects to opengl again
290 $queue->enqueue_release_gl_objects ([$tex]);
291
292 # wait
293 $queue->finish;
294
295 # now draw the texture, the defaults should be all right
296 glTexParameterf GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST;
297
298 glEnable GL_TEXTURE_2D;
299 glBegin GL_QUADS;
300 glTexCoord2f 0, 1; glVertex3i -1, -1, -1;
301 glTexCoord2f 0, 0; glVertex3i 1, -1, -1;
302 glTexCoord2f 1, 0; glVertex3i 1, 1, -1;
303 glTexCoord2f 1, 1; glVertex3i -1, 1, -1;
304 glEnd;
305
306 glXSwapBuffers;
307
308 select undef, undef, undef, 1/60;
309 }
215 310
216=head1 DOCUMENTATION 311=head1 DOCUMENTATION
217 312
218=head2 BASIC CONVENTIONS 313=head2 BASIC CONVENTIONS
219 314
… …
241=item * Structures are often specified by flattening out their components 336=item * Structures are often specified by flattening out their components
242as with short vectors, and returned as arrayrefs. 337as with short vectors, and returned as arrayrefs.
243 338
244=item * When enqueuing commands, the wait list is specified by adding 339=item * When enqueuing commands, the wait list is specified by adding
245extra arguments to the function - anywhere a C<$wait_events...> argument 340extra arguments to the function - anywhere a C<$wait_events...> argument
246is documented this can be any number of event objects. 341is documented this can be any number of event objects. As an extsnion
342implemented by this module, C<undef> values will be ignored in the event
343list.
247 344
248=item * When enqueuing commands, if the enqueue method is called in void 345=item * When enqueuing commands, if the enqueue method is called in void
249context, no event is created. In all other contexts an event is returned 346context, no event is created. In all other contexts an event is returned
250by the method. 347by the method.
251 348
… …
275 372
276=head2 GLX SUPPORT 373=head2 GLX SUPPORT
277 374
278Due to the sad state that OpenGL support is in in Perl (mostly the OpenGL 375Due to the sad state that OpenGL support is in in Perl (mostly the OpenGL
279module, which has little to no documentation and has little to no support 376module, which has little to no documentation and has little to no support
280for glx), this module, as a special extension, treats context creation 377for glX), this module, as a special extension, treats context creation
281properties C<OpenCL::GLX_DISPLAY_KHR> and C<OpenCL::GL_CONTEXT_KHR> 378properties C<OpenCL::GLX_DISPLAY_KHR> and C<OpenCL::GL_CONTEXT_KHR>
282specially: If either or both of these are C<undef>, then the OpenCL 379specially: If either or both of these are C<undef>, then the OpenCL
283module tries to dynamically resolve C<glxGetCurrentDisplay> and 380module tries to dynamically resolve C<glXGetCurrentDisplay> and
284C<glxGetCurrentContext>, call these functions and use their return values 381C<glXGetCurrentContext>, call these functions and use their return values
285instead. 382instead.
286 383
287For this to work, the OpenGL library must be loaded, a GLX context must 384For this to work, the OpenGL library must be loaded, a GLX context must
288have been created and be made current, and C<dlsym> must be available and 385have been created and be made current, and C<dlsym> must be available and
289capable of finding the function via C<RTLD_DEFAULT>. 386capable of finding the function via C<RTLD_DEFAULT>.
… …
653 750
654=item @device_partition_property_exts = $device->affinity_domains_ext 751=item @device_partition_property_exts = $device->affinity_domains_ext
655 752
656Calls C<clGetDeviceInfo> with C<CL_DEVICE_AFFINITY_DOMAINS_EXT> and returns the result. 753Calls C<clGetDeviceInfo> with C<CL_DEVICE_AFFINITY_DOMAINS_EXT> and returns the result.
657 754
658=item $uint = $device->reference_count_ext 755=item $uint = $device->reference_count_ext
659 756
660Calls C<clGetDeviceInfo> with C<CL_DEVICE_REFERENCE_COUNT_EXT > and returns the result. 757Calls C<clGetDeviceInfo> with C<CL_DEVICE_REFERENCE_COUNT_EXT> and returns the result.
661 758
662=item @device_partition_property_exts = $device->partition_style_ext 759=item @device_partition_property_exts = $device->partition_style_ext
663 760
664Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARTITION_STYLE_EXT> and returns the result. 761Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARTITION_STYLE_EXT> and returns the result.
665 762
… …
674=item $queue = $ctx->queue ($device, $properties) 771=item $queue = $ctx->queue ($device, $properties)
675 772
676Create a new OpenCL::Queue object from the context and the given device. 773Create a new OpenCL::Queue object from the context and the given device.
677 774
678L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateCommandQueue.html> 775L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateCommandQueue.html>
776
777Example: create an out-of-order queue.
778
779 $queue = $ctx->queue ($device, OpenCL::QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE);
679 780
680=item $ev = $ctx->user_event 781=item $ev = $ctx->user_event
681 782
682Creates a new OpenCL::UserEvent object. 783Creates a new OpenCL::UserEvent object.
683 784
… …
714Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object that refers to the given 815Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object that refers to the given
715OpenGL buffer object. 816OpenGL buffer object.
716 817
717http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLBuffer.html 818http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLBuffer.html
718 819
820=item $img = $ctx->gl_texture ($flags, $target, $miplevel, $texture)
821
822Creates a new OpenCL::Image object that refers to the given OpenGL
823texture object or buffer.
824
825http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateFromGLTexture.html
826
719=item $ctx->gl_texture2d ($flags, $target, $miplevel, $texture) 827=item $img = $ctx->gl_texture2d ($flags, $target, $miplevel, $texture)
720 828
721Creates a new OpenCL::Image2D object that refers to the given OpenGL 829Creates a new OpenCL::Image2D object that refers to the given OpenGL
7222D texture object. 8302D texture object.
723 831
724http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture2D.html 832http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture2D.html
725 833
726=item $ctx->gl_texture3d ($flags, $target, $miplevel, $texture) 834=item $img = $ctx->gl_texture3d ($flags, $target, $miplevel, $texture)
727 835
728Creates a new OpenCL::Image3D object that refers to the given OpenGL 836Creates a new OpenCL::Image3D object that refers to the given OpenGL
7293D texture object. 8373D texture object.
730 838
731http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture3D.html 839http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture3D.html
… …
794for completion, unless the method is called in void context, in which case 902for completion, unless the method is called in void context, in which case
795no event object is created. 903no event object is created.
796 904
797They also allow you to specify any number of other event objects that this 905They also allow you to specify any number of other event objects that this
798request has to wait for before it starts executing, by simply passing the 906request has to wait for before it starts executing, by simply passing the
799event objects as extra parameters to the enqueue methods. 907event objects as extra parameters to the enqueue methods. To simplify
908program design, this module ignores any C<undef> values in the list of
909events. This makes it possible to code operations such as this, without
910having to put a valid event object into C<$event> first:
911
912 $event = $queue->enqueue_xxx (..., $event);
800 913
801Queues execute in-order by default, without any parallelism, so in most 914Queues execute in-order by default, without any parallelism, so in most
802cases (i.e. you use only one queue) it's not necessary to wait for or 915cases (i.e. you use only one queue) it's not necessary to wait for or
803create event objects. 916create event objects, althoguh an our of order queue is often a bit
917faster.
804 918
805=over 4 919=over 4
806 920
807=item $ev = $queue->enqueue_read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...) 921=item $ev = $queue->enqueue_read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...)
808 922
… …
874reference to an array of local work sizes, with the same number of 988reference to an array of local work sizes, with the same number of
875elements as @$global_work_size. 989elements as @$global_work_size.
876 990
877L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html> 991L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html>
878 992
879=item $ev = $queue->enqueue_marker ($wait_events...)
880
881L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMarker.html>
882
883=item $ev = $queue->enqueue_acquire_gl_objects ([object, ...], $wait_events...) 993=item $ev = $queue->enqueue_acquire_gl_objects ([object, ...], $wait_events...)
884 994
885Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired 995Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired
886for subsequent OpenCL usage. 996for subsequent OpenCL usage.
887 997
… …
896 1006
897=item $ev = $queue->enqueue_wait_for_events ($wait_events...) 1007=item $ev = $queue->enqueue_wait_for_events ($wait_events...)
898 1008
899L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html> 1009L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html>
900 1010
901=item $queue->enqueue_barrier 1011=item $ev = $queue->enqueue_marker ($wait_events...)
902 1012
1013L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMarkerWithWaitList.html>
1014
1015=item $ev = $queue->enqueue_barrier ($wait_events...)
1016
903L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueBarrier.html> 1017L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueBarrierWithWaitList.html>
904 1018
905=item $queue->flush 1019=item $queue->flush
906 1020
907L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clFlush.html> 1021L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clFlush.html>
908 1022
… …
1023 1137
1024=back 1138=back
1025 1139
1026=head2 THE OpenCL::Image CLASS 1140=head2 THE OpenCL::Image CLASS
1027 1141
1028This is the superclass of all image objects - OpenCL::Image2D and OpenCL::Image3D. 1142This is the superclass of all image objects - OpenCL::Image1D,
1143OpenCL::Image1DArray, OpenCL::Image1DBuffer, OpenCL::Image2D,
1144OpenCL::Image2DArray and OpenCL::Image3D.
1029 1145
1030=over 4 1146=over 4
1031 1147
1032=item $packed_value = $ev->image_info ($name) 1148=item $packed_value = $ev->image_info ($name)
1033 1149
… …
1274 1390
1275This is a family of methods to set the kernel argument with the number C<$index> to the give C<$value>. 1391This is a family of methods to set the kernel argument with the number C<$index> to the give C<$value>.
1276 1392
1277TYPE is one of C<char>, C<uchar>, C<short>, C<ushort>, C<int>, C<uint>, 1393TYPE is one of C<char>, C<uchar>, C<short>, C<ushort>, C<int>, C<uint>,
1278C<long>, C<ulong>, C<half>, C<float>, C<double>, C<memory>, C<buffer>, 1394C<long>, C<ulong>, C<half>, C<float>, C<double>, C<memory>, C<buffer>,
1279C<image2d>, C<image3d>, C<sampler> or C<event>. 1395C<image2d>, C<image3d>, C<sampler>, C<local> or C<event>.
1280 1396
1281Chars and integers (including the half type) are specified as integers, 1397Chars and integers (including the half type) are specified as integers,
1282float and double as floating point values, memory/buffer/image2d/image3d 1398float and double as floating point values, memory/buffer/image2d/image3d
1283must be an object of that type or C<undef>, and sampler and event must be 1399must be an object of that type or C<undef>, local-memory arguments are
1284objects of that type. 1400set by specifying the size, and sampler and event must be objects of that
1401type.
1285 1402
1286L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html> 1403L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html>
1287 1404
1288=back 1405=back
1289 1406
… …
1378package OpenCL; 1495package OpenCL;
1379 1496
1380use common::sense; 1497use common::sense;
1381 1498
1382BEGIN { 1499BEGIN {
1383 our $VERSION = '0.92'; 1500 our $VERSION = '0.96';
1384 1501
1385 require XSLoader; 1502 require XSLoader;
1386 XSLoader::load (__PACKAGE__, $VERSION); 1503 XSLoader::load (__PACKAGE__, $VERSION);
1387 1504
1388 @OpenCL::Buffer::ISA = 1505 @OpenCL::Buffer::ISA =
1389 @OpenCL::Image::ISA = OpenCL::Memory::; 1506 @OpenCL::Image::ISA = OpenCL::Memory::;
1390 1507
1391 @OpenCL::BufferObj::ISA = OpenCL::Buffer::; 1508 @OpenCL::BufferObj::ISA = OpenCL::Buffer::;
1392 1509
1393 @OpenCL::Image2D::ISA = 1510 @OpenCL::Image2D::ISA =
1511 @OpenCL::Image3D::ISA =
1512 @OpenCL::Image2DArray::ISA =
1513 @OpenCL::Image1D::ISA =
1514 @OpenCL::Image1DArray::ISA =
1394 @OpenCL::Image3D::ISA = OpenCL::Image::; 1515 @OpenCL::Image1DBuffer::ISA = OpenCL::Image::;
1395 1516
1396 @OpenCL::UserEvent::ISA = OpenCL::Event::; 1517 @OpenCL::UserEvent::ISA = OpenCL::Event::;
1397} 1518}
1398 1519
13991; 15201;
1400 1521
1401=head1 AUTHOR 1522=head1 AUTHOR

Diff Legend

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