|
| | CE_Grid () |
| |
| | CE_Grid (const CE_Grid &src) |
| |
| virtual | ~CE_Grid () |
| |
| const cl::Buffer & | buffer () 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 |
| |
| bool | indexToPos (int x, int y, int z, UT_Vector3F &pos) 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_Grid & | operator= (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_Grid & | operator+= (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 cl::NDRange *lrange=nullptr, int n=1) const |
| |
| 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 | scaledAddCornerFromCenter (fpreal32 c1, const CE_Grid &g1) |
| |
| 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, UT_Vector3 voxelSize) |
| |
| void | divergenceCenterToCorner (const CE_Grid &x, const CE_Grid &y, const CE_Grid &z, fpreal32 scale, UT_Vector3 voxelSize) |
| |
| void | applyGradient (const CE_Grid &p, fpreal32 scale, fpreal32 voxelSize, int axis) |
| |
| void | applyGradientCornerToCenter (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 | boxBlur (int passes, UT_Vector3 radius) |
| |
| void | computeNorms (fpreal64 &norminf, fpreal64 &norm2) const |
| | Compute the infinity-norm and 2-norm of the grid. More...
|
| |
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,
{
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:
doubleit(myGrid.buffer(), myGrid.getOffset(),
myGrid.getXStride2D3D(),
myGrid.getYStride2D3D(),
myGrid.getZStride2D3D());
Definition at line 71 of file CE_Grid.h.
Inheritance diagram for CE_Grid:
1.8.6 
