PDA

View Full Version : Multiple Image Inputs to OpenCL Kernel



toneburst
11-16-2009, 02:20 PM
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

dbs2
11-17-2009, 11:59 AM
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.

toneburst
11-17-2009, 01:00 PM
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

affie
11-17-2009, 02:41 PM
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));
.....
}

toneburst
11-17-2009, 03:44 PM
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

affie
11-17-2009, 07:03 PM
Samplers can also be defined in the kernel. You can reuse a sampler for multiple images or you can have separate sampler / image.

toneburst
11-18-2009, 10:46 AM
Great, thanks, got it working now.

Cheers,

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