array of images?

I need an array of images in my kernel (so I can pass multiple images without runtime-generating the kernel code), but I have failed to make this work (and have not found anything in spec that would forbid it)

as this would require something like “__global __read_only image2d_t * textures”, which doesn’t compile (due to syntax errors on multiple modifiers), I have “typedef __read_only image2d_t i2; … __global i2 * textures”, which partly works.

for “packing” I have a simple runtime-generated kernel, that copies the image2d_t objects to a buffer with size sizeof(cl_mem)*texture_count, but when I try to read these images I get a linker error (on nvidia implementation, don’t have access to ATi hardware)

the error can be simply replicated by simple kernel with typedef (included further down), which brings me to think that the problem is not in having arrays of (or pointers to) image objects, but rather in incorrect handling of typedef.

typedef __read_only image2d_t i2;
__kernel void test_kernel(i2 texture, __global float4 * out)
{
int2 coords=(get_global_id(0),get_global_id(1));
const sampler_t texsampler = CLK_NORMALIZED_COORDS_FALSE|CLK_ADDRESS_REPEAT|CLK_FILTER_LINEAR;
*out = read_imagef(texture,texsampler,coords);
}

this kernel compiles fine if the parameter type is specified without typedefing, but that is not possible with the array case

anoyone has any thoughts? perhaps a quote from the spec saying I can’t do this? or a way I can specify array of images without running into syntax or linking errors?

Arrays of images are not supported in OpenCL 1.0 and 1.1. The only way for you to pass an array of images is to declare each image as a separate argument.

Affie, I know you are right. However, I think we have a bug in the specs because I don’t see this limitation explained anywhere either. I expected to see a comment in section 6.8, such as the one we have one for samplers: Samplers cannot be declared as arrays, but I don’t see the equivalent for images.

OK, thanks for the info, maybe I can put it into OpenCL 1.2+ wishlist or is there some sort of hard technical reason?

however, the aforementioned typedef example should IMHO work, as it is not an array or pointer or anything, just another name for the type

and just to be clear, I declare each image in a separate argument, but in a much simpler kernel, which is runtime generated and looks like this (for 3 textures) and which compiles fine

#include "types.h"
__kernel void pack_textures_2D(__global i2 * images, __read_only image2d_t tex0, __read_only image2d_t tex1, __read_only image2d_t tex2)
{
 images[0] = tex0;
 images[1] = tex1;
 images[2] = tex2;
}

where “types.h” contains the typedef for i2 (and other struct I have in my kernels), and the first argument then gets passed to other kernels.

I understand that while images are cl_mem (pointer to device memory) on the host side, they may get some special treatement on the device side (loading into texturing units or what not)

one more thing, is having a image2d_t inside a struct supported? eg. for defining material properties (color, transparency,texture). this could also indirectly solve the image array problem (if I could have an array of these structs)

PS: sorry for spam, can’t edit the previous post

is having a image2d_t inside a struct supported

No. See section 6.8: The image2d_t or image3d_t data types cannot be declared in a struct