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

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.34 by root, Thu Apr 19 13:06:55 2012 UTC vs.
Revision 1.40 by root, Thu Apr 19 22:33:27 2012 UTC

… …
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
… …
271 ulong IV - Q 366 ulong IV - Q
272 float NV float f 367 float NV float f
273 half IV ushort S 368 half IV ushort S
274 double NV double d 369 double NV double d
275 370
371=head2 GLX SUPPORT
372
373Due to the sad state that OpenGL support is in in Perl (mostly the OpenGL
374module, which has little to no documentation and has little to no support
375for glX), this module, as a special extension, treats context creation
376properties C<OpenCL::GLX_DISPLAY_KHR> and C<OpenCL::GL_CONTEXT_KHR>
377specially: If either or both of these are C<undef>, then the OpenCL
378module tries to dynamically resolve C<glXGetCurrentDisplay> and
379C<glXGetCurrentContext>, call these functions and use their return values
380instead.
381
382For this to work, the OpenGL library must be loaded, a GLX context must
383have been created and be made current, and C<dlsym> must be available and
384capable of finding the function via C<RTLD_DEFAULT>.
385
276=head2 THE OpenCL PACKAGE 386=head2 THE OpenCL PACKAGE
277 387
278=over 4 388=over 4
279 389
280=item $int = OpenCL::errno 390=item $int = OpenCL::errno
… …
859reference to an array of local work sizes, with the same number of 969reference to an array of local work sizes, with the same number of
860elements as @$global_work_size. 970elements as @$global_work_size.
861 971
862L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html> 972L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html>
863 973
864=item $ev = $queue->enqueue_marker 974=item $ev = $queue->enqueue_marker ($wait_events...)
865 975
866L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMarker.html> 976L<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...)
979
980Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired
981for subsequent OpenCL usage.
982
983L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueAcquireGLObjects.html>
984
985=item $ev = $queue->enqueue_release_gl_objects ([object, ...], $wait_events...)
986
987Enqueues a list (an array-ref of OpenCL::Memory objects) to be released
988for subsequent OpenGL usage.
989
990L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReleaseGLObjects.html>
867 991
868=item $ev = $queue->enqueue_wait_for_events ($wait_events...) 992=item $ev = $queue->enqueue_wait_for_events ($wait_events...)
869 993
870L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html> 994L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html>
871 995
… …
1039 1163
1040=for gengetinfo begin gl_texture 1164=for gengetinfo begin gl_texture
1041 1165
1042=item $GLenum = $gl_texture->target 1166=item $GLenum = $gl_texture->target
1043 1167
1044Calls C<clGetGlTextureInfo> with C<CL_GL_TEXTURE_TARGET> and returns the result. 1168Calls C<clGetGLTextureInfo> with C<CL_GL_TEXTURE_TARGET> and returns the result.
1045 1169
1046=item $GLint = $gl_texture->gl_mipmap_level 1170=item $GLint = $gl_texture->gl_mipmap_level
1047 1171
1048Calls C<clGetGlTextureInfo> with C<CL_GL_MIPMAP_LEVEL> and returns the result. 1172Calls C<clGetGLTextureInfo> with C<CL_GL_MIPMAP_LEVEL> and returns the result.
1049 1173
1050=for gengetinfo end gl_texture 1174=for gengetinfo end gl_texture
1051 1175
1052=back 1176=back
1053 1177
… …
1349package OpenCL; 1473package OpenCL;
1350 1474
1351use common::sense; 1475use common::sense;
1352 1476
1353BEGIN { 1477BEGIN {
1354 our $VERSION = '0.92'; 1478 our $VERSION = '0.95';
1355 1479
1356 require XSLoader; 1480 require XSLoader;
1357 XSLoader::load (__PACKAGE__, $VERSION); 1481 XSLoader::load (__PACKAGE__, $VERSION);
1358 1482
1359 @OpenCL::Buffer::ISA = 1483 @OpenCL::Buffer::ISA =

Diff Legend

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