Main Page   Namespace List   Class Hierarchy   Alphabetical List   Compound List   File List   Namespace Members   Compound Members   File Members   Related Pages  

CUDAMarchingCubes.cu File Reference

CUDA implementation of Marching Cubes. More...

#include "CUDAKernels.h"
#include "CUDAMarchingCubes.h"
#include <thrust/device_ptr.h>
#include <thrust/scan.h>
#include <thrust/functional.h>
#include "CUDAParPrefixOps.h"

Go to the source code of this file.

Defines

#define CUDAMARCHINGCUBES_INTERNAL   1
#define RESTRICT
#define NTHREADS   48
#define KERNTHREADS   256

Functions

__host__ __device__ float3 operator+ (float3 a, float3 b)
__host__ __device__ uint3 operator+ (uint3 a, uint3 b)
__host__ __device__ uint2 operator+ (uint2 a, uint2 b)
__host__ __device__ float3 operator- (float3 a, float3 b)
__host__ __device__ uint3 operator- (uint3 a, uint3 b)
__host__ __device__ float3 operator * (float b, float3 a)
__host__ __device__ float dot (float3 a, float3 b)
__host__ __device__ float3 fmaf3 (float x, float3 y, float3 z)
__device__ __host__ float3 lerp (float3 a, float3 b, float t)
__host__ __device__ float length (float3 v)
__device__ void convert_color (float3 &cf, float3 cf2)
__device__ void convert_color (uchar4 &cu, float3 cf)
__device__ void convert_color (float3 &cf, uchar4 cu)
__device__ void convert_normal (float3 &nf, float3 nf2)
__device__ void convert_normal (char3 &cc, float3 cf)
__device__ float sampleVolume (const float *RESTRICT data, uint3 p, uint3 gridSize)
__device__ float3 sampleColors (const float3 *RESTRICT data, uint3 p, uint3 gridSize)
__device__ float3 sampleColors (const uchar4 *RESTRICT data, uint3 p, uint3 gridSize)
__device__ uint3 calcGridPos (unsigned int i, uint3 gridSize)
__device__ float3 vertexInterp (float isolevel, float3 p0, float3 p1, float f0, float f1)
template<int GRIDIS3D, int SUBGRID> __global__ void classifyVoxel (uint2 *RESTRICT voxelVerts, cudaTextureObject_t numVertsTexObj, const float *RESTRICT volume, uint3 gridSize, unsigned int numVoxels, float3 voxelSize, uint3 subGridStart, uint3 subGridEnd, float isoValue)
__global__ void compactVoxels (unsigned int *RESTRICT compactedVoxelArray, const uint2 *RESTRICT voxelOccupied, unsigned int lastVoxel, unsigned int numVoxels, unsigned int numVoxelsm1)
__global__ void generateTriangleVerticesSMEM (float3 *RESTRICT pos, const unsigned int *RESTRICT compactedVoxelArray, const uint2 *RESTRICT numVertsScanned, cudaTextureObject_t triTexObj, cudaTextureObject_t numVertsTexObj, const float *RESTRICT volume, uint3 gridSize, float3 voxelSize, float isoValue, unsigned int activeVoxels, unsigned int maxVertsM3)
__global__ void offsetTriangleVertices (float3 *RESTRICT pos, float3 origin, unsigned int numVertsM1)
template<class NORMAL> __global__ void generateTriangleNormals (const float3 *RESTRICT pos, NORMAL *norm, cudaTextureObject_t volTexObj, float3 gridSizeInv, float3 bBoxInv, unsigned int numVerts)
template<class VERTEXCOL, class VOLTEX, class NORMAL> __global__ void generateTriangleColorNormal (const float3 *RESTRICT pos, VERTEXCOL *RESTRICT col, NORMAL *RESTRICT norm, const VOLTEX *RESTRICT colors, cudaTextureObject_t volTexObj, uint3 gridSize, float3 gridSizeInv, float3 bBoxInv, unsigned int numVerts)
void ThrustScanWrapperUint2 (uint2 *output, uint2 *input, unsigned int numElements)
void ThrustScanWrapperArea (float *output, float *input, unsigned int numElements)
__global__ void computeTriangleAreas (const float3 *RESTRICT pos, float *RESTRICT area, unsigned int maxTria)


Detailed Description

CUDA implementation of Marching Cubes.

This class computes an isosurface for a given density grid using a CUDA Marching Cubes (MC) alorithm. The implementation is loosely modeled after the MC demo from the Nvidia GPU Computing SDK, but the design has been improved and extended in several ways. This implementation achieves higher performance by reducing the number of temporary memory buffers, reduces the number of scan calls by using vector integer types, and allows extraction of per-vertex normals and optionally computes per-vertex colors if a volumetric texture map is provided by the caller.

This work is described in the following papers:

"Evaluation of Emerging Energy-Efficient Heterogeneous Computing Platforms for Biomolecular and Cellular Simulation Workloads" John E. Stone, Michael J. Hallock, James C. Phillips, Joseph R. Peterson, Zaida Luthey-Schulten, and Klaus Schulten. 25th International Heterogeneity in Computing Workshop, 2016 IEEE International Parallel and Distributed Processing Symposium Workshops (IPDPSW), pp. 89-100, 2016. http://dx.doi.org/10.1109/IPDPSW.2016.130

"Fast Visualization of Gaussian Density Surfaces for Molecular Dynamics and Particle System Trajectories" Michael Krone, John E. Stone, Thomas Ertl, and Klaus Schulten. EuroVis - Short Papers, pp. 67-71, 2012. http://dx.doi.org/10.2312/PE/EuroVisShort/EuroVisShort2012/067-071

Revision:
1.44
Date:
2022/03/18 20:39:06

Author:
Michael Krone <michael.krone@visus.uni-stuttgart.de> , John Stone <johns@ks.uiuc.edu>
\copyright (C) Copyright 1995-2019 The Board of Trustees of the University of Illinois, All Rights Reserved
UIUC Open Source License: http://www.ks.uiuc.edu/Research/vmd/plugins/pluginlicense.html

Definition in file CUDAMarchingCubes.cu.


Define Documentation

#define CUDAMARCHINGCUBES_INTERNAL   1
 

Definition at line 61 of file CUDAMarchingCubes.cu.

#define KERNTHREADS   256
 

Definition at line 83 of file CUDAMarchingCubes.cu.

#define NTHREADS   48
 

Definition at line 80 of file CUDAMarchingCubes.cu.

Referenced by generateTriangleVerticesSMEM.

#define RESTRICT
 

Definition at line 75 of file CUDAMarchingCubes.cu.

Referenced by classifyVoxel, compactVoxels, computeTriangleAreas, generateTriangleColorNormal, generateTriangleNormals, generateTriangleVerticesSMEM, offsetTriangleVertices, sampleColors, and sampleVolume.


Function Documentation

__device__ uint3 calcGridPos unsigned int    i,
uint3    gridSize
 

Definition at line 213 of file CUDAMarchingCubes.cu.

Referenced by classifyVoxel, and generateTriangleVerticesSMEM.

template<int GRIDIS3D, int SUBGRID>
__global__ void classifyVoxel uint2 *RESTRICT    voxelVerts,
cudaTextureObject_t    numVertsTexObj,
const float *RESTRICT    volume,
uint3    gridSize,
unsigned int    numVoxels,
float3    voxelSize,
uint3    subGridStart,
uint3    subGridEnd,
float    isoValue
 

Definition at line 236 of file CUDAMarchingCubes.cu.

References calcGridPos, RESTRICT, and sampleVolume.

__global__ void compactVoxels unsigned int *RESTRICT    compactedVoxelArray,
const uint2 *RESTRICT    voxelOccupied,
unsigned int    lastVoxel,
unsigned int    numVoxels,
unsigned int    numVoxelsm1
 

Definition at line 333 of file CUDAMarchingCubes.cu.

References RESTRICT.

__global__ void computeTriangleAreas const float3 *RESTRICT    pos,
float *RESTRICT    area,
unsigned int    maxTria
 

Definition at line 649 of file CUDAMarchingCubes.cu.

References length, and RESTRICT.

__device__ void convert_color float3 &    cf,
uchar4    cu
[inline]
 

Definition at line 158 of file CUDAMarchingCubes.cu.

__device__ void convert_color uchar4 &    cu,
float3    cf
[inline]
 

Definition at line 149 of file CUDAMarchingCubes.cu.

__device__ void convert_color float3 &    cf,
float3    cf2
[inline]
 

Definition at line 145 of file CUDAMarchingCubes.cu.

__device__ void convert_normal char3 &    cc,
float3    cf
[inline]
 

Definition at line 173 of file CUDAMarchingCubes.cu.

References cc.

Referenced by generateTriangleColorNormal, and generateTriangleNormals.

__device__ void convert_normal float3 &    nf,
float3    nf2
[inline]
 

Definition at line 168 of file CUDAMarchingCubes.cu.

__host__ __device__ float dot float3    a,
float3    b
[inline]
 

Definition at line 114 of file CUDAMarchingCubes.cu.

Referenced by clip_ray_by_plane, cylinder_array_color_intersect, cylinder_array_intersect, fog_coord, jitter_sphere3f, length, main, miss_gradient_bg_sky_plane, miss_gradient_bg_sky_sphere, ray_sphere_clip_interval, ring_array_color_intersect, shade_ambient_occlusion, shade_light, shader_template, Wavefunction::sort_orbitals, and sphere_intersect_hearn_baker2.

__host__ __device__ float3 fmaf3 float    x,
float3    y,
float3    z
[inline]
 

Definition at line 119 of file CUDAMarchingCubes.cu.

References make_float3, and z.

Referenced by lerp.

template<class VERTEXCOL, class VOLTEX, class NORMAL>
__global__ void generateTriangleColorNormal const float3 *RESTRICT    pos,
VERTEXCOL *RESTRICT    col,
NORMAL *RESTRICT    norm,
const VOLTEX *RESTRICT    colors,
cudaTextureObject_t    volTexObj,
uint3    gridSize,
float3    gridSizeInv,
float3    bBoxInv,
unsigned int    numVerts
 

Definition at line 520 of file CUDAMarchingCubes.cu.

References convert_color, convert_normal, lerp, make_float3, n, RESTRICT, and sampleColors.

template<class NORMAL>
__global__ void generateTriangleNormals const float3 *RESTRICT    pos,
NORMAL *    norm,
cudaTextureObject_t    volTexObj,
float3    gridSizeInv,
float3    bBoxInv,
unsigned int    numVerts
 

Definition at line 479 of file CUDAMarchingCubes.cu.

References convert_normal, make_float3, n, and RESTRICT.

__global__ void generateTriangleVerticesSMEM float3 *RESTRICT    pos,
const unsigned int *RESTRICT    compactedVoxelArray,
const uint2 *RESTRICT    numVertsScanned,
cudaTextureObject_t    triTexObj,
cudaTextureObject_t    numVertsTexObj,
const float *RESTRICT    volume,
uint3    gridSize,
float3    voxelSize,
float    isoValue,
unsigned int    activeVoxels,
unsigned int    maxVertsM3
 

Definition at line 349 of file CUDAMarchingCubes.cu.

References calcGridPos, make_float3, NTHREADS, RESTRICT, sampleVolume, and vertexInterp.

__host__ __device__ float length float3    v [inline]
 

Definition at line 138 of file CUDAMarchingCubes.cu.

References dot.

Referenced by cmd_rawtimestep, R3dDisplayDevice::comment, compile_branch, compile_regex, computeTriangleAreas, cylinder_array_color_intersect, cylinder_array_intersect, TclTextInterp::doEvent, TclTextInterp::evalFile, fill_header, find_fixedlength, generic_tri_bounds, glwin_compile_shaders, imd_recv_header, imd_recv_header_nolengthswap, main, match, match_ref, obj_veclength, colvarmodule::matrix2d::row::operator vector1d, IMDSimThread::reader, ring_array_color_intersect, colvarmodule::matrix2d::row::row, JRegex::search, colvarmodule::matrix2d::row::set, VMDDisplayList::set_clip_normal, shader_template, sphere_fade_and_clip, transvec, Matrix4::transvec, transvecinv, Matrix4::transvecinv, IMDSimBlocking::update, vmdpcre_compile, and vmdpcre_exec.

__device__ __host__ float3 lerp float3    a,
float3    b,
float    t
[inline]
 

Definition at line 127 of file CUDAMarchingCubes.cu.

References fmaf3.

Referenced by generateTriangleColorNormal, and vertexInterp.

__global__ void offsetTriangleVertices float3 *RESTRICT    pos,
float3    origin,
unsigned int    numVertsM1
 

Definition at line 458 of file CUDAMarchingCubes.cu.

References RESTRICT.

__host__ __device__ float3 operator * float    b,
float3    a
[inline]
 

Definition at line 109 of file CUDAMarchingCubes.cu.

References make_float3.

__host__ __device__ uint2 operator+ uint2    a,
uint2    b
[inline]
 

Definition at line 96 of file CUDAMarchingCubes.cu.

__host__ __device__ uint3 operator+ uint3    a,
uint3    b
[inline]
 

Definition at line 93 of file CUDAMarchingCubes.cu.

__host__ __device__ float3 operator+ float3    a,
float3    b
[inline]
 

Definition at line 90 of file CUDAMarchingCubes.cu.

References make_float3.

__host__ __device__ uint3 operator- uint3    a,
uint3    b
[inline]
 

Definition at line 104 of file CUDAMarchingCubes.cu.

__host__ __device__ float3 operator- float3    a,
float3    b
[inline]
 

Definition at line 101 of file CUDAMarchingCubes.cu.

References make_float3.

__device__ float3 sampleColors const uchar4 *RESTRICT    data,
uint3    p,
uint3    gridSize
 

Definition at line 203 of file CUDAMarchingCubes.cu.

References convert_color, data, and RESTRICT.

Referenced by generateTriangleColorNormal.

__device__ float3 sampleColors const float3 *RESTRICT    data,
uint3    p,
uint3    gridSize
 

Definition at line 197 of file CUDAMarchingCubes.cu.

References data, and RESTRICT.

__device__ float sampleVolume const float *RESTRICT    data,
uint3    p,
uint3    gridSize
 

Definition at line 190 of file CUDAMarchingCubes.cu.

References data, and RESTRICT.

Referenced by classifyVoxel, and generateTriangleVerticesSMEM.

void ThrustScanWrapperArea float *    output,
float *    input,
unsigned int    numElements
 

Definition at line 640 of file CUDAMarchingCubes.cu.

Referenced by CUDAMarchingCubes::computeSurfaceArea.

void ThrustScanWrapperUint2 uint2 *    output,
uint2 *    input,
unsigned int    numElements
 

Definition at line 630 of file CUDAMarchingCubes.cu.

__device__ float3 vertexInterp float    isolevel,
float3    p0,
float3    p1,
float    f0,
float    f1
 

Definition at line 225 of file CUDAMarchingCubes.cu.

References lerp.

Referenced by generateTriangleVerticesSMEM.


Generated on Fri Mar 29 02:46:36 2024 for VMD (current) by doxygen1.2.14 written by Dimitri van Heesch, © 1997-2002