Mipmaps and Image Pyramid

Hi,

I would like to use image pyramids and mipmaps in OpenCL.

I need them for two purposes:

  1. To sample blocks in an arbitrary scale and rotation from a mipmap.
  2. To have an image integral for each scale of the mipmap’s pyramid separately.

I know that image2d arrays are not supported in OpenCL.
I wonder what is the best way to do that.

I thought of two options:

  1. To store all scales in one image (and use OpenCL’s image type) like this: http://en.wikipedia.org/wiki/File:MipMa … STS101.jpg
  2. To use OpenGL built-in mipmaps. Is it possible to sample blocks from OpenGL mipmaps in OpenCL?
    Any other ideas are welcome as well.

Thanks in advance,
Yoav

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.

Thanks for your quick and helpful reply!

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?

Thanks in advance,
Yoav

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!