Image4DMemory< T > Class Template Reference

#include <cudaImageList.h>

Classes
struct	KernelParams

Public Member Functions
	Image4DMemory (int w, int h, int d, int L)
	~Image4DMemory ()
void	free ()
void	copyToDevice (T *src, bool async=false, cudaStream_t s=0)
void	copyToHost (T *dst, bool async=false, cudaStream_t s=0)
void	clear ()
float *	getImgAddr (int2 imgpos)
void	copyImageToHost (int z, int l, T *dst, bool async=false, cudaStream_t s=0)
void	copyImageToDevice (int z, int l, T *src, bool async=false, cudaStream_t s=0)
KernelParams	bind ()
void	unbind ()

Static Public Member Functions
static __device__ T	read (const KernelParams &kp, int x, int y, int2 imgpos)
static __device__ void	write (T value, const KernelParams &kp, int x, int y, int2 imgpos)

Public Attributes
KernelParams	kp
int	layers
int	totalImg
int	rows

Detailed Description

template<typename T>
class Image4DMemory< T >

Definition at line 404 of file cudaImageList.h.

Constructor & Destructor Documentation

Image4DMemory()

template<typename T >

Image4DMemory< T >::Image4DMemory ( int  w, int  h, int  d, int  L  )

inline

Definition at line 424 of file cudaImageList.h.

                                              {
        kp.imgh = h;
        kp.imgw = w;
        kp.depth = d;
        layers = L;
        totalImg = d*L;

        rows = 2048 / kp.imgh;
        kp.cols = (totalImg + rows - 1) / rows;

        CheckCUDAError( cudaMallocPitch (&kp.d_data, &kp.pitch, sizeof(T) * kp.cols * kp.imgw, rows * kp.imgh) );
    }

~Image4DMemory()

template<typename T >

Image4DMemory< T >::~Image4DMemory ( )

inline

Definition at line 437 of file cudaImageList.h.

                     {
        free();
    }

Member Function Documentation

bind()

template<typename T >

KernelParams Image4DMemory< T >::bind ( )

inline

Definition at line 494 of file cudaImageList.h.

{ return kp; }

clear()

template<typename T >

void Image4DMemory< T >::clear ( )

inline

Definition at line 462 of file cudaImageList.h.

    {
        cudaMemset2D(kp.d_data, kp.pitch, 0, sizeof(T)*(kp.cols*kp.imgw), rows*kp.imgh);
    }

copyImageToDevice()

template<typename T >

void Image4DMemory< T >::copyImageToDevice ( int  z, int  l, T *  src, bool  async = false, cudaStream_t  s = 0  )

inline

Definition at line 484 of file cudaImageList.h.

    {
        int2 imgpos = kp.GetImagePos(z, l);
        if (async)
            cudaMemcpy2DAsync(getImgAddr(imgpos), kp.pitch, src, sizeof(T)*kp.imgw, sizeof(T)*kp.imgw, kp.imgh, cudaMemcpyHostToDevice, s);
        else
            cudaMemcpy2D(getImgAddr(imgpos), kp.pitch, src, sizeof(T)*kp.imgw, sizeof(T)*kp.imgw, kp.imgh, cudaMemcpyHostToDevice);
    }

copyImageToHost()

template<typename T >

void Image4DMemory< T >::copyImageToHost ( int  z, int  l, T *  dst, bool  async = false, cudaStream_t  s = 0  )

inline

Definition at line 475 of file cudaImageList.h.

    {
        int2 imgpos = kp.GetImagePos(z, l);
        if (async)
            cudaMemcpy2DAsync(dst, sizeof(T)*kp.imgw, getImgAddr(imgpos), kp.pitch, kp.imgw * sizeof(T), kp.imgh, cudaMemcpyDeviceToHost, s);
        else
            cudaMemcpy2D(dst, sizeof(T)*kp.imgw, getImgAddr(imgpos), kp.pitch, kp.imgw * sizeof(T), kp.imgh, cudaMemcpyDeviceToHost);
    }

copyToDevice()

template<typename T >

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

inline

Definition at line 445 of file cudaImageList.h.

    {
        for (int L=0;L<layers;L++) {
            for (int i=0;i<kp.depth;i++)
                copyImageToDevice(i, L, &src[ kp.imgw * kp.imgh * ( kp.depth * L + i ) ], async, s);
        }
    }

copyToHost()

template<typename T >

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

inline

Definition at line 453 of file cudaImageList.h.

    {
        for (int L=0;L<layers;L++) {
            for (int i=0;i<kp.depth;i++)
                copyImageToHost(i, L, &dst[ kp.imgw * kp.imgh * ( kp.depth * L + i ) ], async, s);
        }
    }

free()

template<typename T >

void Image4DMemory< T >::free ( )

inline

Definition at line 440 of file cudaImageList.h.

               {
        if(kp.d_data) cudaFree(kp.d_data);
        kp.d_data=0;
    }

getImgAddr()

template<typename T >

float* Image4DMemory< T >::getImgAddr ( int2  imgpos )

inline

Definition at line 467 of file cudaImageList.h.

    {
        char* d = (char*)kp.d_data;
        d += imgpos.y * kp.pitch;
        return &((float*)d)[imgpos.x];
    }

read()

template<typename T >

static __device__ T Image4DMemory< T >::read ( const KernelParams &  kp, int  x, int  y, int2  imgpos  )

inlinestatic

Definition at line 497 of file cudaImageList.h.

                                                                                {
        return ((T*)( (char*)kp.d_data + (y + imgpos.y) * kp.pitch))[ x + imgpos.x ];
    }

unbind()

template<typename T >

void Image4DMemory< T >::unbind ( )

inline

Definition at line 495 of file cudaImageList.h.

{}

write()

template<typename T >

static __device__ void Image4DMemory< T >::write ( T  value, const KernelParams &  kp, int  x, int  y, int2  imgpos  )

inlinestatic

Definition at line 500 of file cudaImageList.h.

                                                                                             {
        ((T*)( (char*)kp.d_data + (y + imgpos.y) * kp.pitch)) [ x + imgpos.x ] = value;
    }

Member Data Documentation

kp

template<typename T >

KernelParams Image4DMemory< T >::kp

Definition at line 420 of file cudaImageList.h.

layers

template<typename T >

int Image4DMemory< T >::layers

Definition at line 421 of file cudaImageList.h.

rows

template<typename T >

int Image4DMemory< T >::rows

Definition at line 422 of file cudaImageList.h.

totalImg

template<typename T >

int Image4DMemory< T >::totalImg

Definition at line 421 of file cudaImageList.h.

---

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

• cudatrack/cudaImageList.h