You can use the gl sharing extension to load a specific mipmap level as described in the clCreateFromGLTexture2D Documentation, so you could iterate over all mip levels and create a cl_image for each.
If your kernel needs to be able to access multiple different miplevels, however, then you might have to set a static number of miplevels as your parameters in your kernel, something like:
__kernel
void mipmapit( __read_only image2d_t img_mip0,
__read_only image2d_t img_mip1,
__read_only image2d_t img_mip2,
__read_only image2d_t img_mip3,
const unsigned int numLevels )
{
if(numLevels>0)
{
// do something with img_mip0
}
if(numLevels>1)
{
// do something with img_mip1
}
if(numLevels>2)
{
// do something with img_mip2
}
if(numLevels>3)
{
// do something with img_mip3
}
}
but this is quite ugly and limited.
Storing all scales in one image sounds like a much better route if you have a good way of indexing the image locations dynamically.
These solutions are good for passing the pyramid to the OpenCL kernel.
However, I would like to access any arbitrary (non-integer) scale of the mipmap from within the kernel, using trilinear interpolation (i.e. linearly interpolate the bilinear interpolation results from the two nearest integer scales of the pyramid).
The kernel should be able to access many different scales in a loop (the order of them is unknown outside the kernel).
As far as I understand, OpenGL (or maybe the GPU hardware?) has built-in support in rendering mipmaps using trilinear interpolations.
I would like to use those trilinear interpolations to access any non-integer coordinate in the space-scale volume from within my kernel. Is it possible to do so, or should I implement a trilinear interpolation from a mipmap by myself inside the kernel?
OpenCL supports linear filtering (see the sampler_t documentation ), in 2 and 3 dimensional images, but not miplevels…
Hmm, thinking about it now, i have a weird alternative–if you could create a 3d image where the 3rd dimension is the miplevel (a “stack” of images at different mip levels), you could use linear filtering in 3d (which is hardware-supported for most GPUs) for a dynamic number of levels and avoid manual implementation of the filtering mechanism. Still, creating the “mip-stack” could run into problems, especially since the images are of different resolutions.
In the end, I think doing multiple scales within one image will be the best, but you will have to interpolate manually between miplevels ( not too rough, you have the smoothstep() function ), as well as index the locations and dimensions of each miplevel manually.
Miplevel access could be a good feature to post on suggestions for future releases also!