HDK
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
CE_Grid Class Reference

#include <CE_Grid.h>

+ Inheritance diagram for CE_Grid:

Public Member Functions

 CE_Grid ()
 
 CE_Grid (const CE_Grid &src)
 
virtual ~CE_Grid ()
 
const cl::Bufferbuffer () const
 
bool hasBuffer () const
 If the current OpenCL buffer is valid. More...
 
void size (int xres, int yres, int zres, int xghost=1, int yghost=1, int zghost=1, int xpad=1, int ypad=1, int zpad=1)
 
int getXRes () const
 
int getYRes () const
 
int getZRes () const
 
int getRes (int dim) const
 
UT_Vector3I getRes () const
 
UT_Vector3I getGhostRes () const
 
int getGhostRes (int dim) const
 
UT_Vector3I getPadding () const
 
int getPadding (int dim) const
 
int getOffset () const
 
int getGhostOffset () const
 
void setBorder (UT_VoxelBorderType type, fpreal32 t)
 
UT_VoxelBorderType getBorder () const
 
fpreal32 getBorderValue () const
 
void setBorderScale (fpreal32 scalex, fpreal32 scaley, fpreal32 scalez)
 
fpreal32 getBorderScale (int axis) const
 
void updateBorderCells () const
 
int64 numVoxels () const
 Returns the number of data voxels in the grid. More...
 
int64 numTotalVoxels () const
 
int64 totalVoxelMemory () const
 
void initFromVoxels (const UT_VoxelArrayF &src, int xghost=1, int yghost=1, int zghost=1, int xpad=1, int ypad=1, int zpad=1)
 
void matchAndCopyToVoxels (UT_VoxelArrayF &dest, bool includeGhostCells=false) const
 
void match (const CE_Grid &src)
 Match the src CE_Grid in terms of size and border conditions. More...
 
bool isMatching (const CE_Grid &src) const
 
bool isCongruent (const CE_Grid &src) const
 Returns true if this CE_Grid matches src in terms of offset and strides. More...
 
int getXStride () const
 
int getYStride () const
 
int getZStride () const
 
int getStride (int dim) const
 
UT_Vector3I getStrides () const
 
int getGhostStride (int axis) const
 Returns the stride along axis including ghost cells. More...
 
bool isAxis2D (int axis) const
 Returns whether the specified axis is 2-dimensional. More...
 
int getAxis2D () const
 Returns the 2-dimensional axis of this grid, or -1 if there is none. More...
 
int getXStride2D (int axis2d) const
 2D strides. More...
 
int getYStride2D (int axis2d) const
 
int getXStride2D3D () const
 
int getYStride2D3D () const
 
int getZStride2D3D () const
 Returns the true Z-stride if the grid is 3D, or 0 if 2D. More...
 
int getXRes2D3D () const
 
int getYRes2D3D () const
 
int getZRes2D3D () const
 Returns the true z resolution if the grid is 3D, or 1 if 2D. More...
 
void constant (fpreal32 v)
 
void zero ()
 
bool isConstant (fpreal32 *cval=0, bool checkBorders=false) const
 
void copyData (const CE_Grid &src)
 
CE_Gridoperator= (const CE_Grid &src)
 Assign to this, equivalent to match(src), followed by copyData(src). More...
 
void stealBuffer (CE_Grid &src)
 Steal the buffer from the other grid, leaving it unitialized. More...
 
CE_Gridoperator+= (const CE_Grid &src)
 
cl::NDRange getGlobalRange () const
 
cl::NDRange getGlobalRange2D (int axis2d) const
 
cl::NDRange getGlobalRange2D3D () const
 
cl::NDRange getLocalRange2D3D () const
 
cl::KernelFunctor bind (cl::Kernel k) const
 Bind a 3D kernel, with one work item per voxel. More...
 
cl::KernelFunctor bind (cl::Program prog, const char *kernelname) const
 
cl::KernelFunctor bind2D (int axis, cl::Kernel k) const
 
cl::KernelFunctor bind2D (int axis, cl::Program prog, const char *kernelname) const
 
cl::KernelFunctor bind2D3D (cl::Kernel k) const
 
cl::KernelFunctor bind2D3D (cl::Program prog, const char *kernelname) const
 
void linearCombination (fpreal32 c0, const CE_Grid &g0, fpreal32 d)
 
void linearCombination (fpreal32 c0, const CE_Grid &g0, fpreal32 c1, const CE_Grid &g1, fpreal32 d)
 
void linearCombination (fpreal32 c0, const CE_Grid &g0, fpreal32 c1, const CE_Grid &g1, fpreal32 c2, const CE_Grid &g2, fpreal32 d)
 
void divergence (const CE_Grid &x, const CE_Grid &y, const CE_Grid &z, fpreal32 scale, const UT_Vector3 voxelSize)
 
void applyGradient (const CE_Grid &p, fpreal32 scale, fpreal32 voxelSize, int axis)
 
fpreal64 sum () const
 Reductions of the grid to a single value. More...
 
fpreal64 sumAbs () const
 
fpreal64 sumSqr () const
 
fpreal64 min () const
 
fpreal64 minAbs () const
 
fpreal64 max () const
 
fpreal64 maxAbs () const
 
fpreal64 average () const
 
fpreal64 localAverage (UT_Vector3I &radius)
 
fpreal64 localSum (UT_Vector3I &radius)
 
fpreal64 localSumSqr (UT_Vector3I &radius)
 
fpreal64 localSumAbs (UT_Vector3I &radius)
 
fpreal64 localMin (UT_Vector3I &radius)
 
fpreal64 localMinAbs (UT_Vector3I &radius)
 
fpreal64 localMax (UT_Vector3I &radius)
 
fpreal64 localMaxAbs (UT_Vector3I &radius)
 
void computeNorms (fpreal64 &norminf, fpreal64 &norm2) const
 Compute the infinity-norm and 2-norm of the grid. More...
 

Static Public Member Functions

static int getXAxis2D (int axis2d)
 
static int getYAxis2D (int axis2d)
 
static cl::NDRange getLocalRange (const cl::NDRange &g)
 Create a local work item range for the supplied global range. More...
 

Protected Member Functions

const cl::BufferallocBuffer () const
 
void releaseBuffer ()
 
void setValue (fpreal32 cval) const
 
void getReductionRanges (const cl::Kernel &k, cl::NDRange &globalRange, cl::NDRange &localRange, uint &groupsize, uint &ngroups, size_t &accumsize) const
 
fpreal64 reduceFlat (cl::Buffer outgrid, uint groupsize, uint ngroups, size_t accumsize, const char *reduceFlags) const
 
fpreal64 doReduce (const char *reduceFlags) const
 
bool doLocalReduce (const char *options, UT_Vector3I &radius)
 

Protected Attributes

cl::Buffer myBuffer
 
bool myIsConstant
 
fpreal32 myConstantVal
 
UT_Vector3I myRes
 
UT_Vector3I myGhostCells
 
UT_Vector3I myPadding
 
int myAxis2d
 
UT_Vector3I myStrides
 
fpreal32 myBorderValue
 
UT_VoxelBorderType myBorderType
 
UT_Vector3F myBorderScale
 

Detailed Description

This class represents a 3-dimensional array of float values stored on an OpenCL device. It is roughly analagous to UT_VoxelArray, although it does not support tiling. There are various convenience functions for simplifying invocations of OpenCL kernels that operate on grids of values.

CE_Grid supports an arbitrary number of "ghost" cells, which are cells at the boundary edges of the grid that can be set according to different UT_VoxelBorderType values, as well as arbitrary number of "padding" cells. The padding cells can be used to ensure that several grids with slightly different resolutions can nonetheless have indentical x, y, and z stride values. This makes writing kernels that operate on grids that represent face-sampled (MAC) vector array much simpler, since the index into the different grids need only be calculated once.

Invoking an OpenCL kernel typically involves one of the bind() functions to bind the kernel to the range of work items in the grid, assuming one OpenCL work item per voxel. There are three different variants of bind(), the default assumes the grid is always 3-dimensional; bind2D() which assumes the grid has a 2-dimensional axis, and bind2D3D(), which will treat a 3-dimensional grid normally, but flatten a 2-d grid for optimal performance on OpenCL devices that prefer a large number of work items in the first dimension. This last function is the preferred method to call if the grid might represent 2D or 3D data.

A simple OpenCL kernel that doubles every value in a CE_Grid might look like:

void __kernel doubleit(__global float *grid,
uint offset, uint xstride, uint ystride, uint zstride)
{
size_t idx = offset + get_global_id(0) * xstride +
get_global_id(1) * ystride +
get_global_id(2) * zstride;
grid[idx] *= 2;
}

Because this kernel takes x, y, and z strides, it can be used for 2D and 3D data: in the 2D case zstride will always be 0. Invoking this kernel then typically looks like:

cl::Program prog = context->loadProgram("mykernels.cl");
cl::KernelFunctor doubleit = myGrid.bind2D3D(prog, "doubleit");
doubleit(myGrid.buffer(), myGrid.getOffset(),
myGrid.getXStride2D3D(),
myGrid.getYStride2D3D(),
myGrid.getZStride2D3D());

Definition at line 71 of file CE_Grid.h.

Constructor & Destructor Documentation

CE_Grid::CE_Grid ( )
CE_Grid::CE_Grid ( const CE_Grid src)
virtual CE_Grid::~CE_Grid ( )
virtual

Member Function Documentation

const cl::Buffer& CE_Grid::allocBuffer ( ) const
protected
void CE_Grid::applyGradient ( const CE_Grid p,
fpreal32  scale,
fpreal32  voxelSize,
int  axis 
)

Add scale * gradient of the supplied field along the supplied axis and voxelsize.

fpreal64 CE_Grid::average ( ) const
inline

Definition at line 381 of file CE_Grid.h.

cl::KernelFunctor CE_Grid::bind ( cl::Kernel  k) const

Bind a 3D kernel, with one work item per voxel.

cl::KernelFunctor CE_Grid::bind ( cl::Program  prog,
const char *  kernelname 
) const
cl::KernelFunctor CE_Grid::bind2D ( int  axis,
cl::Kernel  k 
) const

Bind a 2D kernel, treating the provided axis as the flat one, with one work item per voxel. This also allows calling 2d kernels along slices of a 3D grid.

cl::KernelFunctor CE_Grid::bind2D ( int  axis,
cl::Program  prog,
const char *  kernelname 
) const
cl::KernelFunctor CE_Grid::bind2D3D ( cl::Kernel  k) const

Bind a 2D-3D kernel, which should take x-, y-, and z-strides as as parameters, automatically flattening a 2D grid if necessary.

cl::KernelFunctor CE_Grid::bind2D3D ( cl::Program  prog,
const char *  kernelname 
) const
const cl::Buffer& CE_Grid::buffer ( ) const
inline

Return the underlying OpenCL buffer that can be used in kernel invocations. It allocates this buffer only on demand.

Definition at line 80 of file CE_Grid.h.

void CE_Grid::computeNorms ( fpreal64 norminf,
fpreal64 norm2 
) const

Compute the infinity-norm and 2-norm of the grid.

void CE_Grid::constant ( fpreal32  v)

Set this grid to the specified constant value. This will release the underlying OpenCL buffer and store the constant value in the CE_Grid object.

void CE_Grid::copyData ( const CE_Grid src)

Copy data from the source grid. This requires that the grids have the same data resolution, but ghost cells, padding, borders, etc. can differ.

void CE_Grid::divergence ( const CE_Grid x,
const CE_Grid y,
const CE_Grid z,
fpreal32  scale,
const UT_Vector3  voxelSize 
)

Compute scale * divergence of the vector field represented by the supplied grids and voxelsize.

bool CE_Grid::doLocalReduce ( const char *  options,
UT_Vector3I radius 
)
protected
fpreal64 CE_Grid::doReduce ( const char *  reduceFlags) const
protected
int CE_Grid::getAxis2D ( ) const
inline

Returns the 2-dimensional axis of this grid, or -1 if there is none.

Definition at line 196 of file CE_Grid.h.

UT_VoxelBorderType CE_Grid::getBorder ( ) const
inline

Definition at line 126 of file CE_Grid.h.

fpreal32 CE_Grid::getBorderScale ( int  axis) const
inline

Definition at line 131 of file CE_Grid.h.

fpreal32 CE_Grid::getBorderValue ( ) const
inline

Definition at line 127 of file CE_Grid.h.

int CE_Grid::getGhostOffset ( ) const

Returns the offset from the beginning of the buffer to the beginning of the data including ghost cells.

UT_Vector3I CE_Grid::getGhostRes ( ) const
inline

Definition at line 103 of file CE_Grid.h.

int CE_Grid::getGhostRes ( int  dim) const
inline

Definition at line 104 of file CE_Grid.h.

int CE_Grid::getGhostStride ( int  axis) const

Returns the stride along axis including ghost cells.

cl::NDRange CE_Grid::getGlobalRange ( ) const
inline

Return an OpenCL NDRange comprised of the entire grid resolution, implying one OpenCL work item per voxel.

Definition at line 295 of file CE_Grid.h.

cl::NDRange CE_Grid::getGlobalRange2D ( int  axis2d) const
inline

Return an OpenCL NDRange comprised of the flattened 2D grid resolution, implying one OpenCL work item per voxel.

Definition at line 304 of file CE_Grid.h.

cl::NDRange CE_Grid::getGlobalRange2D3D ( ) const
inline

Return an OpenCL NDRange comprised of the entire 3D grid resolution if the grid is 3D, else the flattened 2D grid resolution.

Definition at line 312 of file CE_Grid.h.

static cl::NDRange CE_Grid::getLocalRange ( const cl::NDRange g)
static

Create a local work item range for the supplied global range.

cl::NDRange CE_Grid::getLocalRange2D3D ( ) const
inline

Definition at line 320 of file CE_Grid.h.

int CE_Grid::getOffset ( ) const
inline

Returns the offset from the beginning of the buffer to the beginning of actual data.

Definition at line 111 of file CE_Grid.h.

UT_Vector3I CE_Grid::getPadding ( ) const
inline

Definition at line 106 of file CE_Grid.h.

int CE_Grid::getPadding ( int  dim) const
inline

Definition at line 107 of file CE_Grid.h.

void CE_Grid::getReductionRanges ( const cl::Kernel k,
cl::NDRange globalRange,
cl::NDRange localRange,
uint groupsize,
uint ngroups,
size_t &  accumsize 
) const
protected
int CE_Grid::getRes ( int  dim) const
inline

Definition at line 100 of file CE_Grid.h.

UT_Vector3I CE_Grid::getRes ( ) const
inline

Definition at line 101 of file CE_Grid.h.

int CE_Grid::getStride ( int  dim) const
inline

Definition at line 182 of file CE_Grid.h.

UT_Vector3I CE_Grid::getStrides ( ) const
inline

Definition at line 184 of file CE_Grid.h.

static int CE_Grid::getXAxis2D ( int  axis2d)
inlinestatic

When flattening a grid to 2-dimensions, this is the axis to treat as the x-axis, usually when using bind2D.

Definition at line 200 of file CE_Grid.h.

int CE_Grid::getXRes ( ) const
inline

Definition at line 96 of file CE_Grid.h.

int CE_Grid::getXRes2D3D ( ) const
inline

Returns the true x resolution if the grid is 3D, or the flattened x resolution if 2D.

Definition at line 244 of file CE_Grid.h.

int CE_Grid::getXStride ( ) const
inline

Definition at line 179 of file CE_Grid.h.

int CE_Grid::getXStride2D ( int  axis2d) const
inline

2D strides.

Definition at line 213 of file CE_Grid.h.

int CE_Grid::getXStride2D3D ( ) const
inline

Returns the true x-stride if the grid is 3D, or the flattened x-stride if 2D.

Definition at line 218 of file CE_Grid.h.

static int CE_Grid::getYAxis2D ( int  axis2d)
inlinestatic

When flattening a grid to 2-dimensions, this is the axis to treat as the y-axis, usually when using bind2D.

Definition at line 207 of file CE_Grid.h.

int CE_Grid::getYRes ( ) const
inline

Definition at line 97 of file CE_Grid.h.

int CE_Grid::getYRes2D3D ( ) const
inline

Returns the true y resolution if the grid is 3D, or the flattened y resolution if 2D.

Definition at line 253 of file CE_Grid.h.

int CE_Grid::getYStride ( ) const
inline

Definition at line 180 of file CE_Grid.h.

int CE_Grid::getYStride2D ( int  axis2d) const
inline

Definition at line 214 of file CE_Grid.h.

int CE_Grid::getYStride2D3D ( ) const
inline

Returns the true Y-stride if the grid is 3D, or the flattened y-stride if 2D.

Definition at line 227 of file CE_Grid.h.

int CE_Grid::getZRes ( ) const
inline

Definition at line 98 of file CE_Grid.h.

int CE_Grid::getZRes2D3D ( ) const
inline

Returns the true z resolution if the grid is 3D, or 1 if 2D.

Definition at line 261 of file CE_Grid.h.

int CE_Grid::getZStride ( ) const
inline

Definition at line 181 of file CE_Grid.h.

int CE_Grid::getZStride2D3D ( ) const
inline

Returns the true Z-stride if the grid is 3D, or 0 if 2D.

Definition at line 235 of file CE_Grid.h.

bool CE_Grid::hasBuffer ( ) const
inline

If the current OpenCL buffer is valid.

Definition at line 88 of file CE_Grid.h.

void CE_Grid::initFromVoxels ( const UT_VoxelArrayF src,
int  xghost = 1,
int  yghost = 1,
int  zghost = 1,
int  xpad = 1,
int  ypad = 1,
int  zpad = 1 
)

Initialize the CE_Grid from the supplied UT_VoxelArray. Note that in the case that UT_VoxelArray::isConstant(), this will be very fast and avoid allocating any actual GPU memory.

bool CE_Grid::isAxis2D ( int  axis) const
inline

Returns whether the specified axis is 2-dimensional.

Definition at line 190 of file CE_Grid.h.

bool CE_Grid::isCongruent ( const CE_Grid src) const

Returns true if this CE_Grid matches src in terms of offset and strides.

bool CE_Grid::isConstant ( fpreal32 cval = 0,
bool  checkBorders = false 
) const

Returns whether the grid is set to a constant value. If checkBorders is true, checks that the border type and value are equal as well.

bool CE_Grid::isMatching ( const CE_Grid src) const

Returns true if this CE_Grid matches src in terms of size and border conditions.

void CE_Grid::linearCombination ( fpreal32  c0,
const CE_Grid g0,
fpreal32  d 
)

Enqueue kernel that stores the linear combination c0 * g0 + d. Requires isCongruent(g0).

void CE_Grid::linearCombination ( fpreal32  c0,
const CE_Grid g0,
fpreal32  c1,
const CE_Grid g1,
fpreal32  d 
)

Enqueue kernel that stores the linear combination c0 * g0 + c1 * g1 + d. Requires isCongruent(g0) && isCongruent(g1).

void CE_Grid::linearCombination ( fpreal32  c0,
const CE_Grid g0,
fpreal32  c1,
const CE_Grid g1,
fpreal32  c2,
const CE_Grid g2,
fpreal32  d 
)

Enqueue kernel that stores the linear combination c0 * g0 + c1 * g1 + c2 * g2 + d Requires isCongruent(g0) && isCongruent(g1) && isCongruent(g2)

fpreal64 CE_Grid::localAverage ( UT_Vector3I radius)
fpreal64 CE_Grid::localMax ( UT_Vector3I radius)
fpreal64 CE_Grid::localMaxAbs ( UT_Vector3I radius)
fpreal64 CE_Grid::localMin ( UT_Vector3I radius)
fpreal64 CE_Grid::localMinAbs ( UT_Vector3I radius)
fpreal64 CE_Grid::localSum ( UT_Vector3I radius)
fpreal64 CE_Grid::localSumAbs ( UT_Vector3I radius)
fpreal64 CE_Grid::localSumSqr ( UT_Vector3I radius)
void CE_Grid::match ( const CE_Grid src)

Match the src CE_Grid in terms of size and border conditions.

void CE_Grid::matchAndCopyToVoxels ( UT_VoxelArrayF dest,
bool  includeGhostCells = false 
) const

Match destination UT_VoxelArray to this CE_Grid and copy data. Note if isConstant() is true, this is very fast. includeGhostCells can be set to include the ghost cell values in the sizing and copy operations.

fpreal64 CE_Grid::max ( ) const
fpreal64 CE_Grid::maxAbs ( ) const
fpreal64 CE_Grid::min ( ) const
fpreal64 CE_Grid::minAbs ( ) const
int64 CE_Grid::numTotalVoxels ( ) const
inline

Returns the number of total voxels in the grid, including ghost cells and padding.

Definition at line 144 of file CE_Grid.h.

int64 CE_Grid::numVoxels ( ) const
inline

Returns the number of data voxels in the grid.

Definition at line 139 of file CE_Grid.h.

CE_Grid& CE_Grid::operator+= ( const CE_Grid src)

Add the source array values to this. Calls linearCombination() and requires isCongruent(src).

CE_Grid& CE_Grid::operator= ( const CE_Grid src)

Assign to this, equivalent to match(src), followed by copyData(src).

fpreal64 CE_Grid::reduceFlat ( cl::Buffer  outgrid,
uint  groupsize,
uint  ngroups,
size_t  accumsize,
const char *  reduceFlags 
) const
protected
void CE_Grid::releaseBuffer ( )
protected
void CE_Grid::setBorder ( UT_VoxelBorderType  type,
fpreal32  t 
)

Like the identically-named function in UT_VoxelArray, set the values that determine grid border behaviour. Note that calling these functions does not actually set the ghost cell values; updateBorderCells() does that.

void CE_Grid::setBorderScale ( fpreal32  scalex,
fpreal32  scaley,
fpreal32  scalez 
)
void CE_Grid::setValue ( fpreal32  cval) const
protected
void CE_Grid::size ( int  xres,
int  yres,
int  zres,
int  xghost = 1,
int  yghost = 1,
int  zghost = 1,
int  xpad = 1,
int  ypad = 1,
int  zpad = 1 
)

Size the grid as specified. Note this does not actually allocate memory on the OpenCL device.

void CE_Grid::stealBuffer ( CE_Grid src)

Steal the buffer from the other grid, leaving it unitialized.

fpreal64 CE_Grid::sum ( ) const

Reductions of the grid to a single value.

fpreal64 CE_Grid::sumAbs ( ) const
fpreal64 CE_Grid::sumSqr ( ) const
int64 CE_Grid::totalVoxelMemory ( ) const
inline

Returns the memory required by the entire grid, including ghost cells and padding.

Definition at line 153 of file CE_Grid.h.

void CE_Grid::updateBorderCells ( ) const

Invoke a series of kernels to fill the ghost cells along each axis with the proper values according to the border value and type.

void CE_Grid::zero ( )
inline

Definition at line 273 of file CE_Grid.h.

Member Data Documentation

int CE_Grid::myAxis2d
protected

Definition at line 420 of file CE_Grid.h.

UT_Vector3F CE_Grid::myBorderScale
protected

Definition at line 425 of file CE_Grid.h.

UT_VoxelBorderType CE_Grid::myBorderType
protected

Definition at line 424 of file CE_Grid.h.

fpreal32 CE_Grid::myBorderValue
protected

Definition at line 423 of file CE_Grid.h.

cl::Buffer CE_Grid::myBuffer
mutableprotected

Definition at line 414 of file CE_Grid.h.

fpreal32 CE_Grid::myConstantVal
protected

Definition at line 416 of file CE_Grid.h.

UT_Vector3I CE_Grid::myGhostCells
protected

Definition at line 418 of file CE_Grid.h.

bool CE_Grid::myIsConstant
mutableprotected

Definition at line 415 of file CE_Grid.h.

UT_Vector3I CE_Grid::myPadding
protected

Definition at line 419 of file CE_Grid.h.

UT_Vector3I CE_Grid::myRes
protected

Definition at line 417 of file CE_Grid.h.

UT_Vector3I CE_Grid::myStrides
protected

Definition at line 422 of file CE_Grid.h.


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