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

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.31 by root, Mon Apr 16 06:39:54 2012 UTC vs.
Revision 1.44 by root, Sat Apr 21 17:47:23 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
…		…
241	=item * Structures are often specified by flattening out their components	336	=item * Structures are often specified by flattening out their components
242	as with short vectors, and returned as arrayrefs.	337	as 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
245	extra arguments to the function - anywhere a C<$wait_events...> argument	340	extra arguments to the function - anywhere a C<$wait_events...> argument
246	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.
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
249	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
250	by the method.	347	by the method.
251		348
…		…
271	ulong IV - Q	368	ulong IV - Q
272	float NV float f	369	float NV float f
273	half IV ushort S	370	half IV ushort S
274	double NV double d	371	double NV double d
275		372
		373	=head2 GLX SUPPORT
		374
		375	Due to the sad state that OpenGL support is in in Perl (mostly the OpenGL
		376	module, which has little to no documentation and has little to no support
		377	for glX), this module, as a special extension, treats context creation
		378	properties C<OpenCL::GLX_DISPLAY_KHR> and C<OpenCL::GL_CONTEXT_KHR>
		379	specially: If either or both of these are C<undef>, then the OpenCL
		380	module tries to dynamically resolve C<glXGetCurrentDisplay> and
		381	C<glXGetCurrentContext>, call these functions and use their return values
		382	instead.
		383
		384	For this to work, the OpenGL library must be loaded, a GLX context must
		385	have been created and be made current, and C<dlsym> must be available and
		386	capable of finding the function via C<RTLD_DEFAULT>.
		387
276	=head2 THE OpenCL PACKAGE	388	=head2 THE OpenCL PACKAGE
277		389
278	=over 4	390	=over 4
279		391
280	=item $int = OpenCL::errno	392	=item $int = OpenCL::errno
…		…
692	Creates a new OpenCL::Image3D object and optionally initialises it with	804	Creates a new OpenCL::Image3D object and optionally initialises it with
693	the given data values.	805	the given data values.
694		806
695	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateImage3D.html>	807	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateImage3D.html>
696		808
		809	=item $buffer = $ctx->gl_buffer ($flags, $bufobj)
		810
		811	Creates a new OpenCL::Buffer (actually OpenCL::BufferObj) object that refers to the given
		812	OpenGL buffer object.
		813
		814	http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLBuffer.html
		815
		816	=item $ctx->gl_texture2d ($flags, $target, $miplevel, $texture)
		817
		818	Creates a new OpenCL::Image2D object that refers to the given OpenGL
		819	2D texture object.
		820
		821	http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture2D.html
		822
		823	=item $ctx->gl_texture3d ($flags, $target, $miplevel, $texture)
		824
		825	Creates a new OpenCL::Image3D object that refers to the given OpenGL
		826	3D texture object.
		827
		828	http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLTexture3D.html
		829
		830	=item $ctx->gl_renderbuffer ($flags, $renderbuffer)
		831
		832	Creates a new OpenCL::Image2D object that refers to the given OpenGL
		833	render buffer.
		834
		835	http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clCreateFromGLRenderbuffer.html
		836
697	=item @formats = $ctx->supported_image_formats ($flags, $image_type)	837	=item @formats = $ctx->supported_image_formats ($flags, $image_type)
698		838
699	Returns a list of matching image formats - each format is an arrayref with	839	Returns a list of matching image formats - each format is an arrayref with
700	two values, $channel_order and $channel_type, in it.	840	two values, $channel_order and $channel_type, in it.
701		841
…		…
831	reference to an array of local work sizes, with the same number of	971	reference to an array of local work sizes, with the same number of
832	elements as @$global_work_size.	972	elements as @$global_work_size.
833		973
834	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html>	974	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueNDRangeKernel.html>
835		975
836	=item $ev = $queue->enqueue_marker	976	=item $ev = $queue->enqueue_marker ($wait_events...)
837		977
838	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMarker.html>	978	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueMarker.html>
		979
		980	=item $ev = $queue->enqueue_acquire_gl_objects ([object, ...], $wait_events...)
		981
		982	Enqueues a list (an array-ref of OpenCL::Memory objects) to be acquired
		983	for subsequent OpenCL usage.
		984
		985	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueAcquireGLObjects.html>
		986
		987	=item $ev = $queue->enqueue_release_gl_objects ([object, ...], $wait_events...)
		988
		989	Enqueues a list (an array-ref of OpenCL::Memory objects) to be released
		990	for subsequent OpenGL usage.
		991
		992	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueReleaseGLObjects.html>
839		993
840	=item $ev = $queue->enqueue_wait_for_events ($wait_events...)	994	=item $ev = $queue->enqueue_wait_for_events ($wait_events...)
841		995
842	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html>	996	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clEnqueueWaitForEvents.html>
843		997
…		…
931	=item $int = $mem->offset	1085	=item $int = $mem->offset
932		1086
933	Calls C<clGetMemObjectInfo> with C<CL_MEM_OFFSET> and returns the result.	1087	Calls C<clGetMemObjectInfo> with C<CL_MEM_OFFSET> and returns the result.
934		1088
935	=for gengetinfo end mem	1089	=for gengetinfo end mem
		1090
		1091	=item ($type, $name) = $mem->gl_object_info
		1092
		1093	Returns the OpenGL object type (e.g. OpenCL::GL_OBJECT_TEXTURE2D) and the
		1094	object "name" (e.g. the texture name) used to create this memory object.
		1095
		1096	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetGLObjectInfo.html>
936		1097
937	=back	1098	=back
938		1099
939	=head2 THE OpenCL::Buffer CLASS	1100	=head2 THE OpenCL::Buffer CLASS
940		1101
…		…
1000		1161
1001	Calls C<clGetImageInfo> with C<CL_IMAGE_DEPTH> and returns the result.	1162	Calls C<clGetImageInfo> with C<CL_IMAGE_DEPTH> and returns the result.
1002		1163
1003	=for gengetinfo end image	1164	=for gengetinfo end image
1004		1165
		1166	=for gengetinfo begin gl_texture
		1167
		1168	=item $GLenum = $gl_texture->target
		1169
		1170	Calls C<clGetGLTextureInfo> with C<CL_GL_TEXTURE_TARGET> and returns the result.
		1171
		1172	=item $GLint = $gl_texture->gl_mipmap_level
		1173
		1174	Calls C<clGetGLTextureInfo> with C<CL_GL_MIPMAP_LEVEL> and returns the result.
		1175
		1176	=for gengetinfo end gl_texture
		1177
1005	=back	1178	=back
1006		1179
1007	=head2 THE OpenCL::Sampler CLASS	1180	=head2 THE OpenCL::Sampler CLASS
1008		1181
1009	=over 4	1182	=over 4
…		…
1198		1371
1199	This is a family of methods to set the kernel argument with the number C<$index> to the give C<$value>.	1372	This is a family of methods to set the kernel argument with the number C<$index> to the give C<$value>.
1200		1373
1201	TYPE is one of C<char>, C<uchar>, C<short>, C<ushort>, C<int>, C<uint>,	1374	TYPE is one of C<char>, C<uchar>, C<short>, C<ushort>, C<int>, C<uint>,
1202	C<long>, C<ulong>, C<half>, C<float>, C<double>, C<memory>, C<buffer>,	1375	C<long>, C<ulong>, C<half>, C<float>, C<double>, C<memory>, C<buffer>,
1203	C<image2d>, C<image3d>, C<sampler> or C<event>.	1376	C<image2d>, C<image3d>, C<sampler>, C<local> or C<event>.
1204		1377
1205	Chars and integers (including the half type) are specified as integers,	1378	Chars and integers (including the half type) are specified as integers,
1206	float and double as floating point values, memory/buffer/image2d/image3d	1379	float and double as floating point values, memory/buffer/image2d/image3d
1207	must be an object of that type or C<undef>, and sampler and event must be	1380	must be an object of that type or C<undef>, local-memory arguments are
1208	objects of that type.	1381	set by specifying the size, and sampler and event must be objects of that
		1382	type.
1209		1383
1210	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html>	1384	L<http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetKernelArg.html>
1211		1385
1212	=back	1386	=back
1213		1387
…		…
1302	package OpenCL;	1476	package OpenCL;
1303		1477
1304	use common::sense;	1478	use common::sense;
1305		1479
1306	BEGIN {	1480	BEGIN {
1307	our $VERSION = '0.92';	1481	our $VERSION = '0.96';
1308		1482
1309	require XSLoader;	1483	require XSLoader;
1310	XSLoader::load (__PACKAGE__, $VERSION);	1484	XSLoader::load (__PACKAGE__, $VERSION);
1311		1485
1312	@OpenCL::Buffer::ISA =	1486	@OpenCL::Buffer::ISA =

Diff Legend

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

Comparing OpenCL/OpenCL.pm (file contents): Revision 1.31 by root, Mon Apr 16 06:39:54 2012 UTC vs. Revision 1.44 by root, Sat Apr 21 17:47:23 2012 UTC

Diff Legend

Comparing OpenCL/OpenCL.pm (file contents):
Revision 1.31 by root, Mon Apr 16 06:39:54 2012 UTC vs.
Revision 1.44 by root, Sat Apr 21 17:47:23 2012 UTC