cudaImageList< T > Struct Template Reference

Stores a large number of small images into a single large memory space. Optimizes GPU memory copies. It has no constructor/destructor, so it can be passed to CUDA kernels. The memory is allocated using cudaMallocPitch. It allows binding to a texture. More...

#include <cudaImageList.h>

Public Types
enum	{ MaxImageWidth = 8192 }

Public Member Functions
CUBOTH int	fullwidth ()
CUBOTH int	fullheight ()
CUBOTH int	capacity ()
CUBOTH int	numpixels ()
CUBOTH bool	isEmpty ()
template<int Flags>
void	allocateHostImageBuffer (pinned_array< T, Flags > &hostImgBuf)
CUBOTH T *	get (int i)
CUBOTH T	pixel_oobcheck (int x, int y, int imgIndex, T border=0.0f)
CUBOTH T &	pixel (int x, int y, int imgIndex)
CUBOTH T *	pixelAddress (int x, int y, int imgIndex)
CUBOTH bool	boundaryHit (float2 center, float radius)
void	free ()
void	copyImageToHost (int img, T *dst, bool async=false, cudaStream_t s=0)
void	copyImageToDevice (int img, T *src, bool async=false, cudaStream_t s=0)
void	copyToHost (T *dst, bool async=false, cudaStream_t s=0)
void	copyToDevice (T *src, bool async=false, cudaStream_t s=0)
void	copyToDevice (T *src, int numImages, bool async=false, cudaStream_t s=0)
void	clear ()
CUBOTH int	totalNumPixels ()
CUBOTH int	totalNumBytes ()
CUBOTH T	interpolate (float x, float y, int idx, bool &outside)
void	bind (texture< T, cudaTextureType2D, cudaReadModeElementType > &texref)
void	unbind (texture< T, cudaTextureType2D, cudaReadModeElementType > &texref)
CUBOTH void	computeImagePos (int &x, int &y, int idx)
__device__ T	interpolateFromTexture (texture< T, cudaTextureType2D, cudaReadModeElementType > texref, float x, float y, int idx, bool &outside)

Static Public Member Functions
static cudaImageList< T >	emptyList ()
static cudaImageList< T >	alloc (int w, int h, int amount)
static CUBOTH T	interp (T a, T b, float x)

Public Attributes
T *	data
size_t	pitch
int	w
int	h
int	count

Detailed Description

template<typename T>
struct cudaImageList< T >

Stores a large number of small images into a single large memory space. Optimizes GPU memory copies. It has no constructor/destructor, so it can be passed to CUDA kernels. The memory is allocated using cudaMallocPitch. It allows binding to a texture.

Todo:

Maybe this should be converted into a 3D cudaArray?

Definition at line 10 of file cudaImageList.h.

Member Enumeration Documentation

anonymous enum

template<typename T>

anonymous enum

Enumerator
MaxImageWidth

Definition at line 19 of file cudaImageList.h.

{ MaxImageWidth = 8192 };

Member Function Documentation

alloc()

template<typename T>

static cudaImageList<T> cudaImageList< T >::alloc ( int  w, int  h, int  amount  )

inlinestatic

Definition at line 35 of file cudaImageList.h.

                                                            {
        cudaImageList imgl;
        imgl.w = w; imgl.h = h;
        imgl.count = amount;

        if (cudaMallocPitch(&imgl.data, &imgl.pitch, sizeof(T)*imgl.fullwidth(), imgl.fullheight()) != cudaSuccess) {
            throw std::bad_alloc(SPrintf("cudaImageListf<%s> alloc %dx%dx%d failed", typeid(T).name(), w, h, amount).c_str());
        }
        return imgl;
    }

allocateHostImageBuffer()

template<typename T>

template<int Flags>

void cudaImageList< T >::allocateHostImageBuffer ( pinned_array< T, Flags > &  hostImgBuf )

inline

Definition at line 47 of file cudaImageList.h.

                                                                     {
        hostImgBuf.init( numpixels() );
    }

bind()

template<typename T>

void cudaImageList< T >::bind ( texture< T, cudaTextureType2D, cudaReadModeElementType > &  texref )

inline

Definition at line 166 of file cudaImageList.h.

                                                                              {
        cudaChannelFormatDesc desc = cudaCreateChannelDesc<T>();
        cudaBindTexture2D(NULL, &texref, data, &desc, w, h*count, pitch);
    }

boundaryHit()

template<typename T>

CUBOTH bool cudaImageList< T >::boundaryHit ( float2  center, float  radius  )

inline

Definition at line 79 of file cudaImageList.h.

    {
        return center.x + radius >= w ||
            center.x - radius < 0 ||
            center.y + radius >= h ||
            center.y - radius < 0;
    }

capacity()

template<typename T>

CUBOTH int cudaImageList< T >::capacity ( )

inline

Definition at line 21 of file cudaImageList.h.

{ return count; }

clear()

template<typename T>

void cudaImageList< T >::clear ( )

inline

Definition at line 136 of file cudaImageList.h.

                 {
        if(data) cudaMemset2D(data, pitch, 0, w*sizeof(T), count*h);
    }

computeImagePos()

template<typename T>

CUBOTH void cudaImageList< T >::computeImagePos ( int &  x, int &  y, int  idx  )

inline

Definition at line 174 of file cudaImageList.h.

    {
        y += idx * h;
    }

copyImageToDevice()

template<typename T>

void cudaImageList< T >::copyImageToDevice ( int  img, T *  src, bool  async = false, cudaStream_t  s = 0  )

inline

Definition at line 106 of file cudaImageList.h.

                                                                                {
        T* dst = pixelAddress (0,0, img); 

        if (async)
            cudaMemcpy2DAsync(dst, pitch, src, w*sizeof(T), w*sizeof(T), h, cudaMemcpyHostToDevice, s);
        else
            cudaMemcpy2D(dst, pitch, src, w*sizeof(T), w*sizeof(T), h, cudaMemcpyHostToDevice);
    }

copyImageToHost()

template<typename T>

void cudaImageList< T >::copyImageToHost ( int  img, T *  dst, bool  async = false, cudaStream_t  s = 0  )

inline

Definition at line 97 of file cudaImageList.h.

                                                                              {
        T* src = pixelAddress (0,0, img); 

        if (async)
            cudaMemcpy2DAsync(dst, sizeof(T)*w, src, pitch, w*sizeof(T), h, cudaMemcpyDeviceToHost, s);
        else
            cudaMemcpy2D(dst, sizeof(T)*w, src, pitch, w*sizeof(T), h, cudaMemcpyDeviceToHost);
    }

copyToDevice() [1/2]

template<typename T>

void cudaImageList< T >::copyToDevice ( T *  src, bool  async = false, cudaStream_t  s = 0  )

inline

Definition at line 122 of file cudaImageList.h.

                                                                  {
        if (async)
            cudaMemcpy2DAsync(data, pitch, src, w*sizeof(T), w*sizeof(T), count*h, cudaMemcpyHostToDevice, s);
        else
            cudaMemcpy2D(data, pitch, src, w*sizeof(T), w*sizeof(T), count*h, cudaMemcpyHostToDevice);
    }

copyToDevice() [2/2]

template<typename T>

void cudaImageList< T >::copyToDevice ( T *  src, int  numImages, bool  async = false, cudaStream_t  s = 0  )

inline

Definition at line 129 of file cudaImageList.h.

                                                                                 {
        if (async)
            cudaMemcpy2DAsync(data, pitch, src, w*sizeof(T), w*sizeof(T), numImages*h, cudaMemcpyHostToDevice, s);
        else
            cudaMemcpy2D(data, pitch, src, w*sizeof(T), w*sizeof(T), numImages*h, cudaMemcpyHostToDevice);
    }

copyToHost()

template<typename T>

void cudaImageList< T >::copyToHost ( T *  dst, bool  async = false, cudaStream_t  s = 0  )

inline

Definition at line 115 of file cudaImageList.h.

                                                                {
        if (async)
            cudaMemcpy2DAsync(dst, sizeof(T)*w, data, pitch, w*sizeof(T), count*h, cudaMemcpyDeviceToHost, s);
        else
            cudaMemcpy2D(dst, sizeof(T)*w, data, pitch, w*sizeof(T), count*h, cudaMemcpyDeviceToHost);
    }

emptyList()

template<typename T>

static cudaImageList<T> cudaImageList< T >::emptyList ( )

inlinestatic

Definition at line 24 of file cudaImageList.h.

                                        {
        cudaImageList imgl;
        imgl.data = 0;
        imgl.pitch = 0;
        imgl.w = imgl.h = 0;
        imgl.count = 0;
        return imgl;
    }

free()

template<typename T>

void cudaImageList< T >::free ( )

inline

Definition at line 88 of file cudaImageList.h.

    {
        if(data) {
            cudaFree(data);
            data=0;
        }
    }

fullheight()

template<typename T>

CUBOTH int cudaImageList< T >::fullheight ( )

inline

Definition at line 17 of file cudaImageList.h.

{ return h*count; }

fullwidth()

template<typename T>

CUBOTH int cudaImageList< T >::fullwidth ( )

inline

Definition at line 16 of file cudaImageList.h.

{ return w; }

get()

template<typename T>

CUBOTH T* cudaImageList< T >::get ( int  i )

inline

Definition at line 51 of file cudaImageList.h.

                         {
        return (T*)(((char*)data) + pitch*h*i);
    }

interp()

template<typename T>

static CUBOTH T cudaImageList< T >::interp ( T  a, T  b, float  x  )

inlinestatic

Definition at line 143 of file cudaImageList.h.

{ return a + (b-a)*x; }

interpolate()

template<typename T>

CUBOTH T cudaImageList< T >::interpolate ( float  x, float  y, int  idx, bool &  outside  )

inline

Definition at line 145 of file cudaImageList.h.

    {
        int rx=x, ry=y;

        if (rx < 0 || ry < 0 || rx >= w-1 || ry >= h-1) {
            outside=true;
            return 0.0f;
        }

        T v00 = pixel(rx, ry, idx);
        T v10 = pixel(rx+1, ry, idx);
        T v01 = pixel(rx, ry+1, idx);
        T v11 = pixel(rx+1, ry+1, idx);

        T v0 = interp (v00, v10, x-rx);
        T v1 = interp (v01, v11, x-rx);

        outside=false;
        return interp (v0, v1, y-ry);
    }

interpolateFromTexture()

template<typename T>

__device__ T cudaImageList< T >::interpolateFromTexture ( texture< T, cudaTextureType2D, cudaReadModeElementType >  texref, float  x, float  y, int  idx, bool &  outside  )

inline

Definition at line 180 of file cudaImageList.h.

    {
        int rx=x, ry=y;

        if (rx < 0 || ry < 0 || rx >= w-1 || ry >= h-1) {
            outside=true;
            return 0.0f;
        }

        computeImagePos(rx, ry, idx);

        float fx=x-floor(x), fy = y-floor(y);
        float u = rx + 0.5f;
        float v = ry + 0.5f;

        T v00 = tex2D(texref, u, v);
        T v10 = tex2D(texref, u+1, v);
        T v01 = tex2D(texref, u, v+1);
        T v11 = tex2D(texref, u+1, v+1);

        T v0 = interp (v00, v10, fx);
        T v1 = interp (v01, v11, fx);

        outside = false;
        return interp (v0, v1, fy);
    }

isEmpty()

template<typename T>

CUBOTH bool cudaImageList< T >::isEmpty ( )

inline

Definition at line 33 of file cudaImageList.h.

{ return data==0; }

numpixels()

template<typename T>

CUBOTH int cudaImageList< T >::numpixels ( )

inline

Definition at line 22 of file cudaImageList.h.

{ return w*h*count; }

pixel()

template<typename T>

CUBOTH T& cudaImageList< T >::pixel ( int  x, int  y, int  imgIndex  )

inline

Definition at line 64 of file cudaImageList.h.

                                               {
        computeImagePos(x,y,imgIndex);
        T* row = (T*) ( (char*)data + y*pitch );
        return row[x];
    }

pixel_oobcheck()

template<typename T>

CUBOTH T cudaImageList< T >::pixel_oobcheck ( int  x, int  y, int  imgIndex, T  border = 0.0f  )

inline

Definition at line 55 of file cudaImageList.h.

                                                                      {
        if (x < 0 || x >= w || y < 0 || y >= h)
            return border;

        computeImagePos(x,y,imgIndex);
        T* row = (T*) ( (char*)data + y*pitch );
        return row[x];
    }

pixelAddress()

template<typename T>

CUBOTH T* cudaImageList< T >::pixelAddress ( int  x, int  y, int  imgIndex  )

inline

Bug:

Possibly bugged, should be row + x * sizeof(T)? X is always 0 though, so never encountered.

Definition at line 71 of file cudaImageList.h.

                                                      {
        computeImagePos(x,y,imgIndex);
        T* row = (T*) ( (char*)data + y*pitch );
        return row + x;
    }

totalNumBytes()

template<typename T>

CUBOTH int cudaImageList< T >::totalNumBytes ( )

inline

Definition at line 141 of file cudaImageList.h.

{ return w*h*count*sizeof(T); }

totalNumPixels()

template<typename T>

CUBOTH int cudaImageList< T >::totalNumPixels ( )

inline

Definition at line 140 of file cudaImageList.h.

{ return w*h*count; }

unbind()

template<typename T>

void cudaImageList< T >::unbind ( texture< T, cudaTextureType2D, cudaReadModeElementType > &  texref )

inline

Definition at line 170 of file cudaImageList.h.

                                                                                {
        cudaUnbindTexture(&texref);
    }

Member Data Documentation

count

template<typename T>

int cudaImageList< T >::count

Definition at line 14 of file cudaImageList.h.

data

template<typename T>

T* cudaImageList< T >::data

Definition at line 11 of file cudaImageList.h.

h

template<typename T>

int cudaImageList< T >::h

Definition at line 13 of file cudaImageList.h.

pitch

template<typename T>

size_t cudaImageList< T >::pitch

Definition at line 12 of file cudaImageList.h.

w

template<typename T>

int cudaImageList< T >::w

Definition at line 13 of file cudaImageList.h.

---

The documentation for this struct was generated from the following file:

• cudatrack/cudaImageList.h