Multiple Image Inputs to OpenCL Kernel

How do I get multiple image inputs to an OpenCL kernel program?
I want to convert pixels of one image into a vertex array, and sample color information from another image to create an array of colour float4s. Both images will have the same dimensions.

I currently have the following code:

__kernel void read_rgba8888(__rd image2d_t GeometryImg, float displacement, float Aspect, __global float4 *mesh, __global float4 *colors)
{
	int	tid_x = get_global_id(0),
		tid_y = get_global_id(1),
		indx = tid_y * get_global_size(0) + tid_x;
		
	float4 color, vertex;
	int2   pos = (int2)(tid_x, tid_y);
	sampler_t smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
	
	color = read_imagef(GeometryImg, smp, (int2)(tid_x, tid_y));
	colors[indx] = color;
	
	// Vertex Position
	float4 vPos;
	vPos.x = (float)pos.x / get_global_size(0);
	vPos.x = (fmod((float)pos.y, (float)2.0) < 1.0) ? 1.0 - vPos.x : vPos.x;
	vPos.x -= 0.5;
	vPos.x *= Aspect;
	
	vPos.y = 1.0 - (float)pos.y / get_global_size(1) - 0.5;
	// Displacement is determined by the displacement amount * the "length" of the color vector (ie brightness)
	vPos.z = displacement * length(color.xyzw) - 0.5;
	vPos.w = 1.0;
		
	vertex = vPos;
	mesh[indx] = vertex;
}

Any tips on optimising the above much appreciated, too.

Cheers guys,

a|x
http://machinesdontcare.wordpress.com

You can have multiple images as arguments to your kernel. I believe the minimum number required to be OpenCL compliant is 8, but you can query how many your device supports with clGetDeviceInfo. The only real restriction is that you can not read-from and write-to the same image in one kernel.

I see, thanks for getting back to me.

Sorry if this is a really basic question, but how would I declare the second image input, and sample from it?

I’m having problems finding documentation on writing OpenCL programs- plenty on backend setup etc. but not much in terms of the basics of the OpenCL language itself.

Cheers,

a|x

Here’s an example of using 2 images.

__kernel void
foo(read_only image2d_t imgA, read_only image2d_t imgB, sampler_t sampler, …)
{
int x = (int)get_global_id(0);
int y = (int)get_global_id(1);

float4 clrA = read_imagef(imgA, sampler, (float2)(x, y));
float4 clrB = read_imagef(imgB, sampler, (float2)(x, y));
.....

}

Ah, great, thank you.
Sampler1 and 2 can be defined in the kernel itself, rather than passed in, presumably…

Thanks a lot, I will experiment.

a|x

Samplers can also be defined in the kernel. You can reuse a sampler for multiple images or you can have separate sampler / image.

Great, thanks, got it working now.

Cheers,

a|x
http://machinesdontcare.wordpress.com