Skip to content

Commit

Permalink
Make use of how for bricks, vdel = 0. Also removes unused arguments…
Browse files Browse the repository at this point in the history 
… and kernels.

That is, gvdb->vdel[0] is always equal to (1, 1, 1) after the change removing
`voxelsize` (see issue #74).

- The CUDA code's `getNodeAtPoint` now takes four arguments and no longer
returns `vdel` (previously, it always returned `(1, 1, 1)`)

- Removes `pstep` from the `gvdbBrickFunc` signature, as this depends only on `dir`.

- Removes unused functions `rayTricubic`, `rayTrilinear`, and `raySurfaceDepthBrick`.
  • Loading branch information
@NBickford-NV
NBickford-NV committed Jul 9, 2020
1 parent 98e7d1b commit fce5f5b
Show file tree
Hide file tree
Showing 6 changed files with 74 additions and 211 deletions.
8 changes: 4 additions & 4 deletions source/gvdb_library/kernels/cuda_gvdb_dda.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -70,10 +70,10 @@ struct HDDAState {
}

// Prepare a DDA for stepping through a brick, which omits the addition by t.x from Prepare.
// TODO: Is vdel always (1, 1, 1) here since 1.1.1?
__device__ void PrepareLeaf(float3 vmin, float3 vdel){
tDel = fabs3(vdel / dir);
float3 pFlt = (pos + t.x * dir - vmin) / vdel;
// For bricks, the size of a child node in voxels (vdel) is always 1.
__device__ void PrepareLeaf(float3 vmin){
tDel = fabs3(1.0f / dir);
float3 pFlt = pos + t.x * dir - vmin;
tSide = ((floor3(pFlt) - pFlt + 0.5f) * make_float3(pStep) + 0.5) * tDel;
p = make_int3(floor3(pFlt));
}
Expand Down
14 changes: 7 additions & 7 deletions source/gvdb_library/kernels/cuda_gvdb_module.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -253,16 +253,16 @@ extern "C" __global__ void gvdbSection3D ( VDBInfo* gvdb, uchar chan, uchar4* ou
if ( t > 0 ) { // yes..
wpos += t*rdir; // get point of surface

float3 offs, vmin, vdel; uint64 nid;
VDBNode* node = getNodeAtPoint ( gvdb, wpos, &offs, &vmin, &vdel, &nid ); // find vdb node at point
float3 offs, vmin; uint64 nid;
VDBNode* node = getNodeAtPoint ( gvdb, wpos, &offs, &vmin, &nid ); // find vdb node at point
if ( node != 0x0 ) {
//---- debugging: show apron
// float3 p = offs + (wpos-vmin)*(34.0/16.0) - make_float3(gvdb.atlas_apron);
// clr = transfer ( tex3D ( volTexIn, p.x, p.y, p.z ) );
t = getTrilinear ( gvdb, chan, wpos, offs, vmin, vdel ); // t <= voxel value
t = getTrilinear ( gvdb, chan, wpos, offs, vmin ); // t <= voxel value
clr = transfer ( gvdb, t ); // clr at point on surface
if ( gvdb->clr_chan != CHAN_UNDEF ) {
float3 p = offs + (wpos - vmin)/vdel;
float3 p = offs + (wpos - vmin);
clr *= make_float4( make_float3( getColor(gvdb, gvdb->clr_chan, p) ), 1.0 );
}
} else {
Expand Down Expand Up @@ -300,13 +300,13 @@ extern "C" __global__ void gvdbSection2D ( VDBInfo* gvdb, uchar chan, uchar4* ou
wpos = SCN_SLICE_PNT + spnt * SCN_SLICE_NORM;

// get leaf node at hit point
float3 offs, vmin, vdel;
float3 offs, vmin;
uint64 nid;
VDBNode* node = getNodeAtPoint ( gvdb, wpos, &offs, &vmin, &vdel, &nid );
VDBNode* node = getNodeAtPoint ( gvdb, wpos, &offs, &vmin, &nid );
if ( node == 0x0 ) { outBuf [ y*scn.width + x ] = make_uchar4(bgclr.x*255, bgclr.y*255, bgclr.z*255, 255); return; }

// get tricubic data value
clr = transfer ( gvdb, getTrilinear ( gvdb, chan, wpos, offs, vmin, vdel ) );
clr = transfer ( gvdb, getTrilinear ( gvdb, chan, wpos, offs, vmin ) );
bgclr = lerp3 ( bgclr, make_float3(clr.x,clr.y,clr.z), clr.w );

outBuf [ y*scn.width + x ] = make_uchar4( bgclr.x*255, bgclr.y*255, bgclr.z*255, 255 );
Expand Down
10 changes: 5 additions & 5 deletions source/gvdb_library/kernels/cuda_gvdb_nodes.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -249,19 +249,19 @@ inline __device__ VDBNode* getNodeIdAtPoint ( VDBInfo* gvdb, float3 pos, uint64*

// Gets the node pointer, atlas-space AABB minimum (not including apron), index-space AABB minimum, and voxel size of
// the brick (leaf node) containing the index-space position `pos`.
// `offs` is the brick's mValue (i.e. atlas-space AABB minimum not including apron)
// `offs` is the brick's atlas-space AABB minimum (i.e. mValue. This does not include apron.)
// `vmin` is the brick's index-space AABB minimum
// `vdel` is the world-space size of each voxel in index-space (which should be 1 since 1.1.1)
inline __device__ VDBNode* getNodeAtPoint ( VDBInfo* gvdb, float3 pos, float3* offs, float3* vmin, float3* vdel, uint64* node_id )
// Note that this function used to return `vdel`, the size of a brick's children in voxels. Since issue #74 on
// 2019-07-14, voxelsize has been deprecated, so this was always (1, 1, 1).
inline __device__ VDBNode* getNodeAtPoint ( VDBInfo* gvdb, float3 pos, float3* offs, float3* vmin, uint64* node_id )
{
// iteratively get node at world point
VDBNode* node = getNode ( gvdb, gvdb->top_lev, 0, pos, node_id );
if ( node == 0x0 ) return 0x0;

// compute node bounding box
*vmin = make_float3(node->mPos);
*vdel = gvdb->vdel[ node->mLev ];
*offs = make_float3( node->mValue );
*offs = make_float3(node->mValue);
return node;
}

Expand Down
6 changes: 3 additions & 3 deletions source/gvdb_library/kernels/cuda_gvdb_operators.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -121,15 +121,15 @@ __device__ void UpdateApron(VDBInfo* gvdb, const uchar channel, const int brickC
{
float3 worldPos;
if (!getAtlasToWorld(gvdb, atlasVoxel, worldPos)) return;
float3 offs, vmin, vdel; uint64 nodeID;
VDBNode* node = getNodeAtPoint(gvdb, worldPos, &offs, &vmin, &vdel, &nodeID);
float3 offs, vmin; uint64 nodeID;
VDBNode* node = getNodeAtPoint(gvdb, worldPos, &offs, &vmin, &nodeID);

if (node == 0x0) {
// Out of range, use the boundary value
value = boundaryValue;
}
else {
offs += (worldPos - vmin) / vdel; // Get the atlas position
offs += (worldPos - vmin); // Get the atlas position
value = surf3Dread<T>(gvdb->volOut[channel], uint(offs.x) * sizeof(T), uint(offs.y), uint(offs.z));
}
}
Expand Down
18 changes: 8 additions & 10 deletions source/gvdb_library/kernels/cuda_gvdb_particles.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -38,9 +38,9 @@ extern "C" __global__ void gvdbInsertPoints ( VDBInfo* gvdb, int num_pnts, char*
float3 wpos = (*(float3*) (ppos + i*pos_stride + pos_off)); // NOTE: +ptrans is below. Allows check for wpos.z==NOHIT

if ( wpos.z == NOHIT ) { pnode[i] = ID_UNDEFL; return; } // If position invalid, return.
float3 offs, vmin, vdel; // Get GVDB node at the particle point
float3 offs, vmin; // Get GVDB node at the particle point
uint64 nid;
VDBNode* node = getNodeAtPoint ( gvdb, wpos + ptrans, &offs, &vmin, &vdel, &nid );
VDBNode* node = getNodeAtPoint ( gvdb, wpos + ptrans, &offs, &vmin, &nid );
if ( node == 0x0 ) { pnode[i] = ID_UNDEFL; return; } // If no brick at location, return.

__syncthreads();
Expand All @@ -59,9 +59,9 @@ extern "C" __global__ void gvdbInsertSupportPoints ( VDBInfo* gvdb, int num_pnts
float3 wdir = (*(float3*) (pdir + i*dir_stride + dir_off));

if ( wpos.z == NOHIT ) { pnode[i] = ID_UNDEFL; return; } // If position invalid, return.
float3 offs, vmin, vdel; // Get GVDB node at the particle point
float3 offs, vmin; // Get GVDB node at the particle point
uint64 nid;
VDBNode* node = getNodeAtPoint ( gvdb, wpos + ptrans + wdir * offset, &offs, &vmin, &vdel, &nid );
VDBNode* node = getNodeAtPoint ( gvdb, wpos + ptrans + wdir * offset, &offs, &vmin, &nid );
if ( node == 0x0 ) { pnode[i] = ID_UNDEFL; return; } // If no brick at location, return.

__syncthreads();
Expand Down Expand Up @@ -547,7 +547,7 @@ extern "C" __global__ void gvdbScatterPointDensity (VDBInfo* gvdb, int num_pnts,
float3 vmin;
float w;
VDBNode* node = getNode ( gvdb, 0, pnode[i], &vmin ); // Get node
float3 p = (wpos-vmin)/gvdb->vdel[0];
float3 p = wpos-vmin;
float3 pi = make_float3(int(p.x), int(p.y), int(p.z));

// range of pi.x,pi.y,pi.z = [0, gvdb->res0-1]
Expand Down Expand Up @@ -604,7 +604,7 @@ extern "C" __global__ void gvdbAddSupportVoxel (VDBInfo* gvdb, int num_pnts, fl
float3 vmin;
float w;
VDBNode* node = getNode ( gvdb, 0, pnode[i], &vmin ); // Get node
float3 p = (wpos-vmin)/gvdb->vdel[0];
float3 p = wpos-vmin;
float3 pi = make_float3(int(p.x), int(p.y), int(p.z));

// -- should be ok that pi.x,pi.y,pi.z = 0
Expand Down Expand Up @@ -720,8 +720,7 @@ extern "C" __global__ void gvdbGatherDensity (VDBInfo* gvdb, int num_pnts, int n

VDBNode* node = getNode(gvdb, 0, sc_nid[sc_id]);
float3 vmin = make_float3(node->mPos);
float3 vdel = gvdb->vdel[0];
int3 vox = node->mValue + make_int3((wpos.x - vmin.x) / vdel.x, (wpos.y - vmin.y) / vdel.y, (wpos.z - vmin.z) / vdel.z);
int3 vox = node->mValue + make_int3(wpos.x - vmin.x, wpos.y - vmin.y, wpos.z - vmin.z);

float3 jpos;
float4 clr = make_float4(0,0,0,1);
Expand Down Expand Up @@ -761,8 +760,7 @@ extern "C" __global__ void gvdbGatherLevelSet (VDBInfo* gvdb, int num_pnts, int

VDBNode* node = getNode(gvdb, 0, sc_nid[sc_id]);
float3 vmin = make_float3(node->mPos);
float3 vdel = gvdb->vdel[0];
int3 vox = node->mValue + make_int3((wpos.x - vmin.x) / vdel.x, (wpos.y - vmin.y) / vdel.y, (wpos.z - vmin.z) / vdel.z);
int3 vox = node->mValue + make_int3(wpos.x - vmin.x, wpos.y - vmin.y, wpos.z - vmin.z);

float3 jpos;
float4 clr = make_float4(0,0,0,1);
Expand Down
Loading

0 comments on commit fce5f5b

Please sign in to comment.