Archives Discussions

[bubu]

Is this code legal?

How can I pass an image handle to a struct? How can I pass an array of images as kernel argument?

thx

typedef struct tMyStruct { float a, b, c, d; const __global float *buf; __read_only image2d_t img; int e, f, g; } MyStruct; __kernel void MyKernel ( const __constant MyStruct *c ) { .... read a value from c->img and c->buf... } __kernel void MyKernel2 ( __read_only image2d_t *imgs ) { .... read a value from imgs ... }

[nou]

read section 6.8

Variables that are declared to be of type image2d_t or image3d_t refer to image memory objects. These can only be specified as arguments to a function. Elements of an image cannot be directly accessed. Specific built-in functions are provided to read from and write to any location in the image.    Refer to section 6.11.13 for a list of image read and write functions.    The image2d_t or image3d_t data types cannot be declared in a struct.    The image2d_t or image3d_t data types cannot be used to declare local variables or as the return type of a function.    A image2d_t, or image3d_t argument to a function cannot be modified.

Arguments to __kernel functions that are declared to be a struct do not allow OpenCL objects to be passed as elements of the struct.

[LeeHowes]

bubu, remember what CL memory objects are. They're complex structures that have to be decoded by the runtime and turned into pointers for whichever device you're targeting. If you could put these references in structures all over the place then the runtime would not necessarily be able to find them to construct real parameters from. There has been discussion about how to do this correctly but no clear answer.

[bubu]

Well, imagine the following cases:

Imagine a programmer wants to write a renderer with multiple materials. Each material could have a reference to a texture. Without arrays of images or without having the chance to use an image handler in a struct... how is he supposed to implement that?

Imagine somebody want to write a 3d engine in OpenCL using their own rasterizer. Imagine each light has a reference to an image used as a bitmap projector....

Imagine all the people... lalalalal

It's a big problem ! You must allow us to use image handlers in structs or image arrays. Without that, none of the cases I described can be implemented in an easy way.

[MicahVillmow]

bubu,
This is something you will have to take up with Khronos, it is not something that AMD can decide as we don't define the specification ourselves.

[bubu]

Originally posted by: MicahVillmow bubu, This is something you will have to take up with Khronos, it is not something that AMD can decide as we don't define the specification ourselves.

So it cannot be done currently with OpenCL 1.0/1.1 then? Impossible?

Btw... cannot you propose these new functionality suggestions directly to khronos as implementor? Must be a 3rd party person using its forums?

[LeeHowes]

Btw... cannot you propose these new functionality suggestions directly to khronos as implementor? Must be a 3rd party person using its forums?

Yes, we can. We just can't really discuss what we have or haven't proposed for obvious reasons.

In this case you're not alone in wanting it.

[MicahVillmow]

The 1.0 and 1.1 specs have been ratified and finalized and are not subject to change. So it can not be done. A vendor can come out with vendor specific extensions, but that does not mean any other vendor must support it.

[bubu]

Ok, thx.

[nou]

well samples can't be part of array and all CLobj can't be part of struct. but where in specification is stated that you can't pack image2d_t into a array?

i can't find it.

and what about using 3D texture as texture array. HW support it as it is part of OpenGL/DX. and i can imagine exposing it with special filter mode like CLK_FILTER_LINEAR_ARRAY or something.

[bubu]

Originally posted by: nou you can't pack image2d_t into am array

I tried with:

__kernel void MyKernel ( _read_only image2d_t *imgs )

{

    const float4 pix = read_imagef ( ... imgs[get_global_id(0)]... );

}

and

__kernel void MyKernel ( _read_only image2d_t imgs[MAX_IMAGES] )

{

    const float4 pix = read_imagef ( ... imgs[get_global_id(0)]... );

}

but I got a compiler error for using ATI, NV and Intel SDK.

and what about using 3D texture as texture array.

Unfortunately, my images have different sizes and formats. On the other hand, I need 4k x 4k images and that's simply too much for a cell.

I've also tried to write the image's clmem to a global buffer and cast implicitly as:

__kernel void MyKernel ( __global float4*imgBuf )

{

    _read_only image2d_t img = ((_read_only image2d_t*)imgBuf)[10];

}

( I wrote the image clmem handles to a CL float4 buffer )

but I also get a a compiler error.

I think it's just and plain.. impossible.

[MicahVillmow]

bubu,
there currently is no such thing as a pointer to an image2d_t/image3d_t in OpenCL, it is a handle and cannot be casted. What you need to do is create a image3d_t where z = 0 is your first 2d image, z = 1 is your second 2d image, etc... The only constraint is that all images have to be the same size and format.

[bubu]

Originally posted by: MicahVillmow bubu, there currently is no such thing as a pointer to an image2d_t/image3d_t in OpenCL, it is a handle and cannot be casted. What you need to do is create a image3d_t where z = 0 is your first 2d image, z = 1 is your second 2d image, etc... The only constraint is that all images have to be the same size and format.

Yep, it's a good alternative but probably the artist would want to use textures of different sizes and formats 8( In my specific case, I wanted to use images for a cached database table... and I also need to use different sizes.

I just hope this could be included in the next CL version. I think somebody posted the suggestion at Khronos atm.