#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) |
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
Definition in file CUDAMarchingCubes.cu.
|
Definition at line 61 of file CUDAMarchingCubes.cu. |
|
Definition at line 83 of file CUDAMarchingCubes.cu. |
|
Definition at line 80 of file CUDAMarchingCubes.cu. Referenced by generateTriangleVerticesSMEM. |
|
Definition at line 75 of file CUDAMarchingCubes.cu. Referenced by classifyVoxel, compactVoxels, computeTriangleAreas, generateTriangleColorNormal, generateTriangleNormals, generateTriangleVerticesSMEM, offsetTriangleVertices, sampleColors, and sampleVolume. |
|
Definition at line 213 of file CUDAMarchingCubes.cu. Referenced by classifyVoxel, and generateTriangleVerticesSMEM. |
|
Definition at line 236 of file CUDAMarchingCubes.cu. References calcGridPos, RESTRICT, and sampleVolume. |
|
Definition at line 333 of file CUDAMarchingCubes.cu. References RESTRICT. |
|
Definition at line 649 of file CUDAMarchingCubes.cu. |
|
Definition at line 158 of file CUDAMarchingCubes.cu. |
|
Definition at line 149 of file CUDAMarchingCubes.cu. |
|
Definition at line 145 of file CUDAMarchingCubes.cu. |
|
Definition at line 173 of file CUDAMarchingCubes.cu. References cc. Referenced by generateTriangleColorNormal, and generateTriangleNormals. |
|
Definition at line 168 of file CUDAMarchingCubes.cu. |
|
|
Definition at line 119 of file CUDAMarchingCubes.cu. References make_float3, and z. Referenced by lerp. |
|
Definition at line 520 of file CUDAMarchingCubes.cu. References convert_color, convert_normal, lerp, make_float3, n, RESTRICT, and sampleColors. |
|
Definition at line 479 of file CUDAMarchingCubes.cu. References convert_normal, make_float3, n, and RESTRICT. |
|
Definition at line 349 of file CUDAMarchingCubes.cu. References calcGridPos, make_float3, NTHREADS, RESTRICT, sampleVolume, and vertexInterp. |
|
|
Definition at line 127 of file CUDAMarchingCubes.cu. References fmaf3. Referenced by generateTriangleColorNormal, and vertexInterp. |
|
Definition at line 458 of file CUDAMarchingCubes.cu. References RESTRICT. |
|
Definition at line 109 of file CUDAMarchingCubes.cu. References make_float3. |
|
Definition at line 96 of file CUDAMarchingCubes.cu. |
|
Definition at line 93 of file CUDAMarchingCubes.cu. |
|
Definition at line 90 of file CUDAMarchingCubes.cu. References make_float3. |
|
Definition at line 104 of file CUDAMarchingCubes.cu. |
|
Definition at line 101 of file CUDAMarchingCubes.cu. References make_float3. |
|
Definition at line 203 of file CUDAMarchingCubes.cu. References convert_color, data, and RESTRICT. Referenced by generateTriangleColorNormal. |
|
Definition at line 197 of file CUDAMarchingCubes.cu. |
|
Definition at line 190 of file CUDAMarchingCubes.cu. References data, and RESTRICT. Referenced by classifyVoxel, and generateTriangleVerticesSMEM. |
|
Definition at line 640 of file CUDAMarchingCubes.cu. Referenced by CUDAMarchingCubes::computeSurfaceArea. |
|
Definition at line 630 of file CUDAMarchingCubes.cu. |
|
Definition at line 225 of file CUDAMarchingCubes.cu. References lerp. Referenced by generateTriangleVerticesSMEM. |