[ViewVC] Diff of: cvs/OpenCL/OpenCL.pm

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.34 by root, Thu Apr 19 13:06:55 2012 UTC vs.
Revision 1.43 by root, Fri Apr 20 20:24:20 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
		217	set tunnel effect.
		218
		219	This 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
		373	Due to the sad state that OpenGL support is in in Perl (mostly the OpenGL
		374	module, which has little to no documentation and has little to no support
		375	for glX), this module, as a special extension, treats context creation
		376	properties C<OpenCL::GLX_DISPLAY_KHR> and C<OpenCL::GL_CONTEXT_KHR>
		377	specially: If either or both of these are C<undef>, then the OpenCL
		378	module tries to dynamically resolve C<glXGetCurrentDisplay> and
		379	C<glXGetCurrentContext>, call these functions and use their return values
		380	instead.
		381
		382	For this to work, the OpenGL library must be loaded, a GLX context must
		383	have been created and be made current, and C<dlsym> must be available and
		384	capable 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
…		…
859	reference to an array of local work sizes, with the same number of	969	reference to an array of local work sizes, with the same number of
860	elements as @$global_work_size.	970	elements as @$global_work_size.
861		971
862	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html>	972	L<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
866	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMarker.html>	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...)
		979
		980	Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired
		981	for subsequent OpenCL usage.
		982
		983	L<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
		987	Enqueues a list (an array-ref of OpenCL::Memory objects) to be released
		988	for subsequent OpenGL usage.
		989
		990	L<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
870	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html>	994	L<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
1044	Calls C<clGetGlTextureInfo> with C<CL_GL_TEXTURE_TARGET> and returns the result.	1168	Calls 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
1048	Calls C<clGetGlTextureInfo> with C<CL_GL_MIPMAP_LEVEL> and returns the result.	1172	Calls 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
…		…
1245		1369
1246	This is a family of methods to set the kernel argument with the number C<$index> to the give C<$value>.	1370	This is a family of methods to set the kernel argument with the number C<$index> to the give C<$value>.
1247		1371
1248	TYPE is one of C<char>, C<uchar>, C<short>, C<ushort>, C<int>, C<uint>,	1372	TYPE is one of C<char>, C<uchar>, C<short>, C<ushort>, C<int>, C<uint>,
1249	C<long>, C<ulong>, C<half>, C<float>, C<double>, C<memory>, C<buffer>,	1373	C<long>, C<ulong>, C<half>, C<float>, C<double>, C<memory>, C<buffer>,
1250	C<image2d>, C<image3d>, C<sampler> or C<event>.	1374	C<image2d>, C<image3d>, C<sampler>, C<local> or C<event>.
1251		1375
1252	Chars and integers (including the half type) are specified as integers,	1376	Chars and integers (including the half type) are specified as integers,
1253	float and double as floating point values, memory/buffer/image2d/image3d	1377	float and double as floating point values, memory/buffer/image2d/image3d
1254	must be an object of that type or C<undef>, and sampler and event must be	1378	must be an object of that type or C<undef>, local-memory arguments are
1255	objects of that type.	1379	set by specifying the size, and sampler and event must be objects of that
		1380	type.
1256		1381
1257	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html>	1382	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html>
1258		1383
1259	=back	1384	=back
1260		1385
…		…
1349	package OpenCL;	1474	package OpenCL;
1350		1475
1351	use common::sense;	1476	use common::sense;
1352		1477
1353	BEGIN {	1478	BEGIN {
1354	our $VERSION = '0.92';	1479	our $VERSION = '0.96';
1355		1480
1356	require XSLoader;	1481	require XSLoader;
1357	XSLoader::load (__PACKAGE__, $VERSION);	1482	XSLoader::load (__PACKAGE__, $VERSION);
1358		1483
1359	@OpenCL::Buffer::ISA =	1484	@OpenCL::Buffer::ISA =

Diff Legend

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

Comparing OpenCL/OpenCL.pm (file contents): Revision 1.34 by root, Thu Apr 19 13:06:55 2012 UTC vs. Revision 1.43 by root, Fri Apr 20 20:24:20 2012 UTC

Diff Legend

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.34 by root, Thu Apr 19 13:06:55 2012 UTC vs.
Revision 1.43 by root, Fri Apr 20 20:24:20 2012 UTC