… | |
… | |
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 |
… | |
… | |
260 | m.x = fabs (fmod (m.x + time * 0.05f, 4.f)) - 2.f; |
260 | m.x = fabs (fmod (m.x + time * 0.05f, 4.f)) - 2.f; |
261 | |
261 | |
262 | float2 z = m; |
262 | float2 z = m; |
263 | float2 c = (float2)(sin (time * 0.05005), cos (time * 0.06001)); |
263 | float2 c = (float2)(sin (time * 0.05005), cos (time * 0.06001)); |
264 | |
264 | |
265 | for (int i = 0; i < 100 && dot (z, z) < 4.f; ++i) |
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; |
266 | z = (float2)(z.x * z.x - z.y * z.y, 2.f * z.x * z.y) + c; |
267 | |
267 | |
268 | float3 colour = (float3)(z.x, z.y, z.x * z.y); |
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.)); |
269 | write_imagef (img, (int2)(get_global_id (0), get_global_id (1)), (float4)(colour * p.x * p.x, 1.)); |
270 | } |
270 | } |
… | |
… | |
288 | |
288 | |
289 | # release objects to opengl again |
289 | # release objects to opengl again |
290 | $queue->enqueue_release_gl_objects ([$tex]); |
290 | $queue->enqueue_release_gl_objects ([$tex]); |
291 | |
291 | |
292 | # wait |
292 | # wait |
293 | $queue->flush; |
293 | $queue->finish; |
294 | |
294 | |
295 | # now draw the texture, the defaults should be all right |
295 | # now draw the texture, the defaults should be all right |
296 | glTexParameterf GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST; |
296 | glTexParameterf GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST; |
297 | |
297 | |
298 | glEnable GL_TEXTURE_2D; |
298 | glEnable GL_TEXTURE_2D; |
… | |
… | |
336 | =item * Structures are often specified by flattening out their components |
336 | =item * Structures are often specified by flattening out their components |
337 | as with short vectors, and returned as arrayrefs. |
337 | as with short vectors, and returned as arrayrefs. |
338 | |
338 | |
339 | =item * When enqueuing commands, the wait list is specified by adding |
339 | =item * When enqueuing commands, the wait list is specified by adding |
340 | extra arguments to the function - anywhere a C<$wait_events...> argument |
340 | extra arguments to the function - anywhere a C<$wait_events...> argument |
341 | is documented this can be any number of event objects. |
341 | is documented this can be any number of event objects. As an extsnion |
|
|
342 | implemented by this module, C<undef> values will be ignored in the event |
|
|
343 | list. |
342 | |
344 | |
343 | =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 |
344 | context, no event is created. In all other contexts an event is returned |
346 | context, no event is created. In all other contexts an event is returned |
345 | by the method. |
347 | by the method. |
346 | |
348 | |
… | |
… | |
454 | It's best to avoid this method and use one of the following convenience |
456 | It's best to avoid this method and use one of the following convenience |
455 | wrappers. |
457 | wrappers. |
456 | |
458 | |
457 | L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformInfo.html> |
459 | L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformInfo.html> |
458 | |
460 | |
|
|
461 | =item $platform->unload_compiler |
|
|
462 | |
|
|
463 | Attempts to unload the compiler for this platform, for endless |
|
|
464 | profit. Does nothing on OpenCL 1.1. |
|
|
465 | |
|
|
466 | L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clUnloadPlatformCompiler.html> |
|
|
467 | |
459 | =for gengetinfo begin platform |
468 | =for gengetinfo begin platform |
460 | |
469 | |
461 | =item $string = $platform->profile |
470 | =item $string = $platform->profile |
462 | |
471 | |
463 | Calls C<clGetPlatformInfo> with C<CL_PLATFORM_PROFILE> and returns the result. |
472 | Calls C<clGetPlatformInfo> with C<CL_PLATFORM_PROFILE> and returns the result. |
… | |
… | |
748 | |
757 | |
749 | =item @device_partition_property_exts = $device->affinity_domains_ext |
758 | =item @device_partition_property_exts = $device->affinity_domains_ext |
750 | |
759 | |
751 | Calls C<clGetDeviceInfo> with C<CL_DEVICE_AFFINITY_DOMAINS_EXT> and returns the result. |
760 | Calls C<clGetDeviceInfo> with C<CL_DEVICE_AFFINITY_DOMAINS_EXT> and returns the result. |
752 | |
761 | |
753 | =item $uint = $device->reference_count_ext |
762 | =item $uint = $device->reference_count_ext |
754 | |
763 | |
755 | Calls C<clGetDeviceInfo> with C<CL_DEVICE_REFERENCE_COUNT_EXT > and returns the result. |
764 | Calls C<clGetDeviceInfo> with C<CL_DEVICE_REFERENCE_COUNT_EXT> and returns the result. |
756 | |
765 | |
757 | =item @device_partition_property_exts = $device->partition_style_ext |
766 | =item @device_partition_property_exts = $device->partition_style_ext |
758 | |
767 | |
759 | Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARTITION_STYLE_EXT> and returns the result. |
768 | Calls C<clGetDeviceInfo> with C<CL_DEVICE_PARTITION_STYLE_EXT> and returns the result. |
760 | |
769 | |
… | |
… | |
769 | =item $queue = $ctx->queue ($device, $properties) |
778 | =item $queue = $ctx->queue ($device, $properties) |
770 | |
779 | |
771 | Create a new OpenCL::Queue object from the context and the given device. |
780 | Create a new OpenCL::Queue object from the context and the given device. |
772 | |
781 | |
773 | L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateCommandQueue.html> |
782 | L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateCommandQueue.html> |
|
|
783 | |
|
|
784 | Example: create an out-of-order queue. |
|
|
785 | |
|
|
786 | $queue = $ctx->queue ($device, OpenCL::QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE); |
774 | |
787 | |
775 | =item $ev = $ctx->user_event |
788 | =item $ev = $ctx->user_event |
776 | |
789 | |
777 | Creates a new OpenCL::UserEvent object. |
790 | Creates a new OpenCL::UserEvent object. |
778 | |
791 | |
… | |
… | |
788 | =item $buf = $ctx->buffer_sv ($flags, $data) |
801 | =item $buf = $ctx->buffer_sv ($flags, $data) |
789 | |
802 | |
790 | Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object and |
803 | Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object and |
791 | initialise it with the given data values. |
804 | initialise it with the given data values. |
792 | |
805 | |
|
|
806 | =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) |
|
|
807 | |
|
|
808 | Creates a new OpenCL::Image object and optionally initialises it with |
|
|
809 | the given data values. |
|
|
810 | |
|
|
811 | L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateImage.html> |
|
|
812 | |
793 | =item $img = $ctx->image2d ($flags, $channel_order, $channel_type, $width, $height, $row_pitch = 0, $data = undef) |
813 | =item $img = $ctx->image2d ($flags, $channel_order, $channel_type, $width, $height, $row_pitch = 0, $data = undef) |
794 | |
814 | |
795 | Creates a new OpenCL::Image2D object and optionally initialises it with |
815 | Creates a new OpenCL::Image2D object and optionally initialises it with |
796 | the given data values. |
816 | the given data values. |
797 | |
817 | |
… | |
… | |
809 | Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object that refers to the given |
829 | Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object that refers to the given |
810 | OpenGL buffer object. |
830 | OpenGL buffer object. |
811 | |
831 | |
812 | http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLBuffer.html |
832 | http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLBuffer.html |
813 | |
833 | |
|
|
834 | =item $img = $ctx->gl_texture ($flags, $target, $miplevel, $texture) |
|
|
835 | |
|
|
836 | Creates a new OpenCL::Image object that refers to the given OpenGL |
|
|
837 | texture object or buffer. |
|
|
838 | |
|
|
839 | http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateFromGLTexture.html |
|
|
840 | |
814 | =item $ctx->gl_texture2d ($flags, $target, $miplevel, $texture) |
841 | =item $img = $ctx->gl_texture2d ($flags, $target, $miplevel, $texture) |
815 | |
842 | |
816 | Creates a new OpenCL::Image2D object that refers to the given OpenGL |
843 | Creates a new OpenCL::Image2D object that refers to the given OpenGL |
817 | 2D texture object. |
844 | 2D texture object. |
818 | |
845 | |
819 | http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture2D.html |
846 | http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture2D.html |
820 | |
847 | |
821 | =item $ctx->gl_texture3d ($flags, $target, $miplevel, $texture) |
848 | =item $img = $ctx->gl_texture3d ($flags, $target, $miplevel, $texture) |
822 | |
849 | |
823 | Creates a new OpenCL::Image3D object that refers to the given OpenGL |
850 | Creates a new OpenCL::Image3D object that refers to the given OpenGL |
824 | 3D texture object. |
851 | 3D texture object. |
825 | |
852 | |
826 | http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture3D.html |
853 | http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture3D.html |
… | |
… | |
889 | for completion, unless the method is called in void context, in which case |
916 | for completion, unless the method is called in void context, in which case |
890 | no event object is created. |
917 | no event object is created. |
891 | |
918 | |
892 | They also allow you to specify any number of other event objects that this |
919 | They also allow you to specify any number of other event objects that this |
893 | request has to wait for before it starts executing, by simply passing the |
920 | request has to wait for before it starts executing, by simply passing the |
894 | event objects as extra parameters to the enqueue methods. |
921 | event objects as extra parameters to the enqueue methods. To simplify |
|
|
922 | program design, this module ignores any C<undef> values in the list of |
|
|
923 | events. This makes it possible to code operations such as this, without |
|
|
924 | having to put a valid event object into C<$event> first: |
|
|
925 | |
|
|
926 | $event = $queue->enqueue_xxx (..., $event); |
895 | |
927 | |
896 | Queues execute in-order by default, without any parallelism, so in most |
928 | Queues execute in-order by default, without any parallelism, so in most |
897 | cases (i.e. you use only one queue) it's not necessary to wait for or |
929 | cases (i.e. you use only one queue) it's not necessary to wait for or |
898 | create event objects. |
930 | create event objects, althoguh an our of order queue is often a bit |
|
|
931 | faster. |
899 | |
932 | |
900 | =over 4 |
933 | =over 4 |
901 | |
934 | |
902 | =item $ev = $queue->enqueue_read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...) |
935 | =item $ev = $queue->enqueue_read_buffer ($buffer, $blocking, $offset, $len, $data, $wait_events...) |
903 | |
936 | |
… | |
… | |
969 | reference to an array of local work sizes, with the same number of |
1002 | reference to an array of local work sizes, with the same number of |
970 | elements as @$global_work_size. |
1003 | elements as @$global_work_size. |
971 | |
1004 | |
972 | L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html> |
1005 | L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html> |
973 | |
1006 | |
974 | =item $ev = $queue->enqueue_marker ($wait_events...) |
|
|
975 | |
|
|
976 | L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMarker.html> |
|
|
977 | |
|
|
978 | =item $ev = $queue->enqueue_acquire_gl_objects ([object, ...], $wait_events...) |
1007 | =item $ev = $queue->enqueue_acquire_gl_objects ([object, ...], $wait_events...) |
979 | |
1008 | |
980 | Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired |
1009 | Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired |
981 | for subsequent OpenCL usage. |
1010 | for subsequent OpenCL usage. |
982 | |
1011 | |
… | |
… | |
991 | |
1020 | |
992 | =item $ev = $queue->enqueue_wait_for_events ($wait_events...) |
1021 | =item $ev = $queue->enqueue_wait_for_events ($wait_events...) |
993 | |
1022 | |
994 | L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html> |
1023 | L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html> |
995 | |
1024 | |
996 | =item $queue->enqueue_barrier |
1025 | =item $ev = $queue->enqueue_marker ($wait_events...) |
997 | |
1026 | |
|
|
1027 | L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueMarkerWithWaitList.html> |
|
|
1028 | |
|
|
1029 | =item $ev = $queue->enqueue_barrier ($wait_events...) |
|
|
1030 | |
998 | L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueBarrier.html> |
1031 | L<http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueBarrierWithWaitList.html> |
999 | |
1032 | |
1000 | =item $queue->flush |
1033 | =item $queue->flush |
1001 | |
1034 | |
1002 | L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clFlush.html> |
1035 | L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clFlush.html> |
1003 | |
1036 | |
… | |
… | |
1118 | |
1151 | |
1119 | =back |
1152 | =back |
1120 | |
1153 | |
1121 | =head2 THE OpenCL::Image CLASS |
1154 | =head2 THE OpenCL::Image CLASS |
1122 | |
1155 | |
1123 | This is the superclass of all image objects - OpenCL::Image2D and OpenCL::Image3D. |
1156 | This is the superclass of all image objects - OpenCL::Image1D, |
|
|
1157 | OpenCL::Image1DArray, OpenCL::Image1DBuffer, OpenCL::Image2D, |
|
|
1158 | OpenCL::Image2DArray and OpenCL::Image3D. |
1124 | |
1159 | |
1125 | =over 4 |
1160 | =over 4 |
1126 | |
1161 | |
1127 | =item $packed_value = $ev->image_info ($name) |
1162 | =item $packed_value = $ev->image_info ($name) |
1128 | |
1163 | |
… | |
… | |
1232 | |
1267 | |
1233 | Creates an OpenCL::Kernel object out of the named C<__kernel> function in |
1268 | Creates an OpenCL::Kernel object out of the named C<__kernel> function in |
1234 | the program. |
1269 | the program. |
1235 | |
1270 | |
1236 | L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernel.html> |
1271 | L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernel.html> |
|
|
1272 | |
|
|
1273 | =item @kernels = $program->kernels_in_program |
|
|
1274 | |
|
|
1275 | Returns all kernels successfully compiled for all devices in program. |
|
|
1276 | |
|
|
1277 | http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateKernelsInProgram.html |
1237 | |
1278 | |
1238 | =for gengetinfo begin program_build |
1279 | =for gengetinfo begin program_build |
1239 | |
1280 | |
1240 | =item $build_status = $program->build_status ($device) |
1281 | =item $build_status = $program->build_status ($device) |
1241 | |
1282 | |
… | |
… | |
1369 | |
1410 | |
1370 | This is a family of methods to set the kernel argument with the number C<$index> to the give C<$value>. |
1411 | This is a family of methods to set the kernel argument with the number C<$index> to the give C<$value>. |
1371 | |
1412 | |
1372 | TYPE is one of C<char>, C<uchar>, C<short>, C<ushort>, C<int>, C<uint>, |
1413 | TYPE is one of C<char>, C<uchar>, C<short>, C<ushort>, C<int>, C<uint>, |
1373 | C<long>, C<ulong>, C<half>, C<float>, C<double>, C<memory>, C<buffer>, |
1414 | C<long>, C<ulong>, C<half>, C<float>, C<double>, C<memory>, C<buffer>, |
1374 | C<image2d>, C<image3d>, C<sampler> or C<event>. |
1415 | C<image2d>, C<image3d>, C<sampler>, C<local> or C<event>. |
1375 | |
1416 | |
1376 | Chars and integers (including the half type) are specified as integers, |
1417 | Chars and integers (including the half type) are specified as integers, |
1377 | float and double as floating point values, memory/buffer/image2d/image3d |
1418 | float and double as floating point values, memory/buffer/image2d/image3d |
1378 | must be an object of that type or C<undef>, and sampler and event must be |
1419 | must be an object of that type or C<undef>, local-memory arguments are |
1379 | objects of that type. |
1420 | set by specifying the size, and sampler and event must be objects of that |
|
|
1421 | type. |
|
|
1422 | |
|
|
1423 | Setting an argument for a kernel does NOT keep a reference to the object - |
|
|
1424 | for example, if you set an argument to some image object, free the image, |
|
|
1425 | and call the kernel, you will run into undefined behaviour. |
1380 | |
1426 | |
1381 | L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html> |
1427 | L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html> |
1382 | |
1428 | |
1383 | =back |
1429 | =back |
1384 | |
1430 | |
… | |
… | |
1473 | package OpenCL; |
1519 | package OpenCL; |
1474 | |
1520 | |
1475 | use common::sense; |
1521 | use common::sense; |
1476 | |
1522 | |
1477 | BEGIN { |
1523 | BEGIN { |
1478 | our $VERSION = '0.95'; |
1524 | our $VERSION = '0.96'; |
1479 | |
1525 | |
1480 | require XSLoader; |
1526 | require XSLoader; |
1481 | XSLoader::load (__PACKAGE__, $VERSION); |
1527 | XSLoader::load (__PACKAGE__, $VERSION); |
1482 | |
1528 | |
1483 | @OpenCL::Buffer::ISA = |
1529 | @OpenCL::Buffer::ISA = |
1484 | @OpenCL::Image::ISA = OpenCL::Memory::; |
1530 | @OpenCL::Image::ISA = OpenCL::Memory::; |
1485 | |
1531 | |
1486 | @OpenCL::BufferObj::ISA = OpenCL::Buffer::; |
1532 | @OpenCL::BufferObj::ISA = OpenCL::Buffer::; |
1487 | |
1533 | |
1488 | @OpenCL::Image2D::ISA = |
1534 | @OpenCL::Image2D::ISA = |
|
|
1535 | @OpenCL::Image3D::ISA = |
|
|
1536 | @OpenCL::Image2DArray::ISA = |
|
|
1537 | @OpenCL::Image1D::ISA = |
|
|
1538 | @OpenCL::Image1DArray::ISA = |
1489 | @OpenCL::Image3D::ISA = OpenCL::Image::; |
1539 | @OpenCL::Image1DBuffer::ISA = OpenCL::Image::; |
1490 | |
1540 | |
1491 | @OpenCL::UserEvent::ISA = OpenCL::Event::; |
1541 | @OpenCL::UserEvent::ISA = OpenCL::Event::; |
1492 | } |
1542 | } |
1493 | |
1543 | |
1494 | 1; |
1544 | 1; |
1495 | |
1545 | |
1496 | =head1 AUTHOR |
1546 | =head1 AUTHOR |