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

CUDAQuickSurf.cu File Reference

CUDA accelerated QuickSurf gaussian density and surface calculation. More...

#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <cuda.h>
#include "Inform.h"
#include "utilities.h"
#include "WKFThreads.h"
#include "WKFUtils.h"
#include "CUDAKernels.h"
#include "CUDASpatialSearch.h"
#include "CUDAMarchingCubes.h"
#include "CUDAQuickSurf.h"
#include "DispCmds.h"
#include "VMDDisplayList.h"

Go to the source code of this file.

Defines

#define CUERR
#define RESTRICT
#define GGRIDSZ   8.0f
#define GBLOCKSZX   8
#define GBLOCKSZY   8
#define GTEXBLOCKSZZ   2
#define GTEXUNROLL   4
#define GBLOCKSZZ   2
#define GUNROLL   4
#define MAXTHRDENS   ( GBLOCKSZX * GBLOCKSZY * GBLOCKSZZ )
#define MINBLOCKDENS   1

Functions

__host__ __device__ float3 operator * (float3 a, float b)
__host__ __device__ void operator+= (float3 &a, float3 b)
__device__ float3 make_float3 (float4 a)
__device__ void convert_density (float &df, float df2)
__device__ void convert_density (unsigned short &dh, float df2)
__device__ void convert_color (float3 &cf, float3 cf2, float invisovalue)
__device__ void convert_color (uchar4 &cu, float3 cf, float invisovalue)
__device__ void convert_color (float3 &cf, uchar4 cu)
template<class DENSITY, class VOLTEX> __global__ void __launch_bounds__ (MAXTHRDENS, MINBLOCKDENS) gaussdensity_fast_tex_norm(int natoms

Variables

__global__ void const float4
*RESTRICT 		sorted_xyzr
__global__ void const float4
*RESTRICT const float4 *RESTRICT 		sorted_color
__global__ void const float4
*RESTRICT const float4 *RESTRICT
int3 		volsz
__global__ void const float4
*RESTRICT const float4 *RESTRICT
int3 int3 		acncells
__global__ void const float4
*RESTRICT const float4 *RESTRICT
int3 int3 float 		acgridspacing
__global__ void const float4
*RESTRICT const float4 *RESTRICT
int3 int3 float float 		invacgridspacing
__global__ void const float4
*RESTRICT const float4 *RESTRICT
int3 int3 float float const
uint2 *RESTRICT 		cellStartEnd
__global__ void const float4
*RESTRICT const float4 *RESTRICT
int3 int3 float float const
uint2 *RESTRICT float 		gridspacing
__global__ void const float4
*RESTRICT const float4 *RESTRICT
int3 int3 float float const
uint2 *RESTRICT float unsigned
int 		z
__global__ void const float4
*RESTRICT const float4 *RESTRICT
int3 int3 float float const
uint2 *RESTRICT float unsigned
int DENSITY *RESTRICT 		densitygrid
__global__ void const float4
*RESTRICT const float4 *RESTRICT
int3 int3 float float const
uint2 *RESTRICT float unsigned
int DENSITY *RESTRICT VOLTEX
*RESTRICT 		voltexmap

Detailed Description

CUDA accelerated QuickSurf gaussian density and surface calculation.

This work is described in the following papers:

"GPU-Accelerated Molecular Visualization on Petascale Supercomputing Platforms" John E. Stone, Kirby L. Vandivort, and Klaus Schulten. UltraVis'13: Proceedings of the 8th International Workshop on Ultrascale Visualization, pp. 6:1-6:8, 2013. http://dx.doi.org/10.1145/2535571.2535595

"Early Experiences Scaling VMD Molecular Visualization and Analysis Jobs on Blue Waters" John E. Stone, Barry Isralewitz, and Klaus Schulten. Extreme Scaling Workshop (XSW), pp. 43-50, 2013. http://dx.doi.org/10.1109/XSW.2013.10

"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 CUDAQuickSurf.cu.

Define Documentation

#define CUERR
 

Value:

{ cudaError_t err; \
  if ((err = cudaGetLastError()) != cudaSuccess) { \
  printf("CUDA error: %s, %s line %d\n", cudaGetErrorString(err), __FILE__, __LINE__); \
  printf("Thread aborting...\n"); \
  return NULL; }}

Definition at line 67 of file CUDAQuickSurf.cu.

#define GBLOCKSZX   8
 

Definition at line 178 of file CUDAQuickSurf.cu.

Referenced by CUDAQuickSurf::calc_surf.

#define GBLOCKSZY   8
 

Definition at line 179 of file CUDAQuickSurf.cu.

Referenced by CUDAQuickSurf::calc_surf.

#define GBLOCKSZZ   2
 

Definition at line 184 of file CUDAQuickSurf.cu.

Referenced by CUDAQuickSurf::calc_surf.

#define GGRIDSZ   8.0f
 

Definition at line 177 of file CUDAQuickSurf.cu.

#define GTEXBLOCKSZZ   2
 

Definition at line 182 of file CUDAQuickSurf.cu.

Referenced by CUDAQuickSurf::calc_surf.

#define GTEXUNROLL   4
 

Definition at line 183 of file CUDAQuickSurf.cu.

Referenced by CUDAQuickSurf::calc_surf.

#define GUNROLL   4
 

Definition at line 185 of file CUDAQuickSurf.cu.

Referenced by CUDAQuickSurf::calc_surf.

#define MAXTHRDENS   ( GBLOCKSZX * GBLOCKSZY * GBLOCKSZZ )
 

Definition at line 193 of file CUDAQuickSurf.cu.

#define MINBLOCKDENS   1
 

Definition at line 201 of file CUDAQuickSurf.cu.

#define RESTRICT
 

Definition at line 171 of file CUDAQuickSurf.cu.

Function Documentation

template<class DENSITY, class VOLTEX>
__global__ void __launch_bounds__ (  MAXTHRDENS   ,
MINBLOCKDENS   
)  [static]
 

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

Definition at line 157 of file CUDAQuickSurf.cu.

Referenced by generateTriangleColorNormal, and sampleColors.

__device__ void convert_color (  uchar4 &    cu,
float3    cf,
float    invisovalue
)  [inline]
 

Definition at line 132 of file CUDAQuickSurf.cu.

__device__ void convert_color (  float3 &    cf,
float3    cf2,
float    invisovalue
)  [inline]
 

Definition at line 125 of file CUDAQuickSurf.cu.

__device__ void convert_density (  unsigned short &    dh,
float    df2
)  [inline]
 

Definition at line 114 of file CUDAQuickSurf.cu.

__device__ void convert_density (  float &    df,
float    df2
)  [inline]
 

Definition at line 108 of file CUDAQuickSurf.cu.

__device__ float3 make_float3 (  float4    a )  [inline]
 

Definition at line 94 of file CUDAQuickSurf.cu.

Referenced by any_hit_clip_sphere, any_hit_opaque, any_hit_transmission, CUDAQuickSurf::calc_surf, clip_ray_by_plane, CUDAMarchingCubes::computeIsosurface, cylinder_array_bounds, cylinder_array_color_bounds, exception, fmaf3, generateTriangleColorNormal, generateTriangleNormals, generateTriangleVerticesSMEM, operator *, operator+, operator-, ring_array_color_bounds, CUDAMarchingCubes::SetVolumeData, shade_ambient_occlusion, shade_light, shader_template, sphere_array_bounds, sphere_array_color_bounds, vmd_camera_cubemap_general, vmd_camera_dome_general, vmd_camera_equirectangular_general, vmd_camera_oculus_rift_general, vmd_camera_orthographic_general, vmd_camera_perspective_general, vmd_cuda_calc_density, and vmd_cuda_compare_sel_refmap.

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

Definition at line 82 of file CUDAQuickSurf.cu.

References make_float3.

__host__ __device__ void operator+= (  float3 &    a,
float3    b
)  [inline]
 

Definition at line 87 of file CUDAQuickSurf.cu.

Variable Documentation

__global__ void const float4* RESTRICT const float4* RESTRICT int3 int3 float acgridspacing
 

Definition at line 215 of file CUDAQuickSurf.cu.

Referenced by calc_ac_cell_ids, CUDAQuickSurf::calc_surf, gaussdensity_cc, gaussdensity_diff, gaussdensity_fast, gaussdensity_sumabsdiff, vmd_cuda_build_accel, vmd_cuda_cc_calc, and vmd_cuda_gaussdensity_calc.

__global__ void const float4* RESTRICT const float4* RESTRICT int3 int3 acncells
 

Definition at line 215 of file CUDAQuickSurf.cu.

Referenced by calc_ac_cell_ids, calc_densities, gaussdensity_cc, gaussdensity_diff, gaussdensity_fast, and gaussdensity_sumabsdiff.

__global__ void const float4* RESTRICT const float4* RESTRICT int3 int3 float float const uint2* RESTRICT cellStartEnd
 

Definition at line 215 of file CUDAQuickSurf.cu.

Referenced by calc_densities, gaussdensity_cc, gaussdensity_diff, gaussdensity_fast, and gaussdensity_sumabsdiff.

__global__ void const float4* RESTRICT const float4* RESTRICT int3 int3 float float const uint2* RESTRICT float unsigned int DENSITY* RESTRICT densitygrid
 

Definition at line 215 of file CUDAQuickSurf.cu.

Referenced by gaussdensity_fast.

__global__ void const float4* RESTRICT const float4* RESTRICT int3 int3 float float const uint2* RESTRICT float gridspacing
 

Definition at line 215 of file CUDAQuickSurf.cu.

Referenced by calc_ac_cell_ids, calc_densities, calc_density_bounds, calc_density_bounds_overlapping_map, QuickSurf::calc_density_map, QuickSurf::calc_surf, CUDAQuickSurf::calc_surf, cenergy, cudaenergythread, energythread, gaussdensity_cc, gaussdensity_diff, gaussdensity_fast, gaussdensity_sumabsdiff, openclenergythread, MolData::readData, MolData::readFile, vmd_cuda_build_accel, vmd_cuda_calc_density, vmd_cuda_cc_calc, vmd_cuda_compare_sel_refmap, vmd_cuda_device_props, vmd_cuda_gaussdensity_calc, vmd_cuda_vol_cpotential, vmd_gaussdensity_avx2, vmd_gaussdensity_opt, vmd_gaussdensity_threaded, vmd_opencl_vol_cpotential, vol_cpotential, and vol_cpotential_cpu.

__global__ void const float4* RESTRICT const float4* RESTRICT int3 int3 float float invacgridspacing
 

Definition at line 215 of file CUDAQuickSurf.cu.

Referenced by calc_ac_cell_ids, CUDAQuickSurf::calc_surf, gaussdensity_cc, gaussdensity_diff, gaussdensity_fast, and gaussdensity_sumabsdiff.

__global__ void const float4* RESTRICT const float4* RESTRICT sorted_color
 

Definition at line 215 of file CUDAQuickSurf.cu.

__global__ void const float4* RESTRICT sorted_xyzr
 

Definition at line 215 of file CUDAQuickSurf.cu.

Referenced by calc_densities, gaussdensity_cc, gaussdensity_diff, gaussdensity_fast, and gaussdensity_sumabsdiff.

__global__ void const float4* RESTRICT const float4* RESTRICT int3 volsz
 

Definition at line 215 of file CUDAQuickSurf.cu.

Referenced by calc_density_bounds, calc_density_bounds_overlapping_map, QuickSurf::calc_surf, CUDAQuickSurf::calc_surf, cc_threaded, CreateEmptyGrid, CreateProbGrid, gaussdensity_cc, gaussdensity_diff, gaussdensity_fast, gaussdensity_sumabsdiff, normalize_pmap, process_pmap, vmd_cuda_build_accel, vmd_cuda_build_density_atom_grid, vmd_cuda_calc_density, vmd_cuda_cc_calc, vmd_cuda_compare_sel_refmap, vmd_cuda_gaussdensity_calc, vmd_cuda_kernel_launch_dims, vmd_gaussdensity_threaded, VolIn_CleanGrid, volin_threaded, and volin_threaded_prob.

__global__ void const float4* RESTRICT const float4* RESTRICT int3 int3 float float const uint2* RESTRICT float unsigned int DENSITY* RESTRICT VOLTEX* RESTRICT voltexmap
 

Definition at line 215 of file CUDAQuickSurf.cu.

Referenced by vmd_gaussdensity_avx2, vmd_gaussdensity_opt, vmd_gaussdensity_threaded, volin_threaded, and volin_threaded_prob.

__global__ void const float4* RESTRICT const float4* RESTRICT int3 int3 float float const uint2* RESTRICT float unsigned int z
 

Definition at line 215 of file CUDAQuickSurf.cu.

Referenced by add, Displayable::add_cent_trans, Displayable::add_glob_trans, BaseMolecule::add_volume_data, colvarmodule::rvector::as_vector, average, bin_evaluation, OpenGLRenderer::build3Dmipmaps, calc_neighbors_kernel, CUDAQuickSurf::calc_surf, MoleculeList::center_from_top_molecule_reps, cephesfastexpf, Displayable::change_center, Surf::compute, IsoSurface::compute, IsoContour::compute, compute_allatoms, compute_occupancy_multiatom, convolution_x_kernel, convolution_x_kernel_s, convolution_y_kernel, convolution_y_kernel_s, convolution_z_kernel, convolution_z_kernel_s, MoleculeGraphics::cov, DrawMolecule::cov, cuda_shortrange, cuda_shortrange_nonperiodic, TachyonDisplayDevice::define_volume_texture, density_rotate, IsoContour::DoCell, IsoSurface::DoCellGeneral, IsoSurface::DoGridPosNorms, energythread, QMData::expand_atompos, fill_atom_bins, fmaf3, gaussdensity_cc, gaussdensity_diff, gaussdensity_fast, gaussdensity_sumabsdiff, VolumeTexture::generateChargeTexture, VolumeTexture::generateColorScaleTexture, VolumeTexture::generateContourLineTexture, VolumeTexture::generateHSVTexture, VolumeTexture::generateIndexTexture, VolumeTexture::generatePosTexture, get_intersection_normal, VRJugglerScene::getWandXYZ, M_VRJapp::getWandXYZ, hill_climb_kernel, interpolation, jacobi, linklist_evaluation, linklist_hashing, measure_center, measure_center_perresidue, measure_inertia, measure_move, measure_sasa, measure_sasa_perresidue, measure_sasalist, minmax_3fv_aligned, minmax_selected_3fv_aligned, molinfo_set, VMDTracker::moveto, Msmpot_compute_shortrng_bin_hashing, Quat::mult, multiply, colvarmodule::rvector::norm2, integrate_potential::nr_linbcg_sym, obj_transabout, obj_transvec, obj_transvecinv, colvarmodule::rvector::operator *=, colvarmodule::rvector::operator+=, colvarmodule::rvector::operator-=, colvarmodule::rvector::operator/=, colvarmodule::rvector::operator[], print_xyz, py_set_center, py_set_trans, py_translate, Quat::Quat, colvarmodule::quaternion::quaternion, RaycastGrid, DisplayDevice::right_eye_dir, ring_axes, rotate, Matrix4::scale, DrawMolecule::scale_factor, VMDApp::scene_translate_by, VMDApp::scene_translate_to, colvarmodule::rvector::set, Displayable::set_cent_trans, IsoSurface::set_color_voltex_rgb3fv, Displayable::set_glob_trans, setup_domain, OpenGLRenderer::setup_initial_opengl_state, subtract, text_cmd_translate, Matrix4::translate, transvec, Matrix4::transvec, transvecinv, Matrix4::transvecinv, colvarmodule::rvector::unit, VRJugglerTracker::update, vec_incr, vmd_cuda_cc_calc, vmd_cuda_gaussdensity_calc, vmd_gaussdensity_avx2, vmd_gaussdensity_opt, vmd_mattrans_d, vmd_measure_volinterior, volinthread, volinthread_prob, VolumetricData::voxel_coord, voxel_coord, VolumetricData::voxel_gradient_fast, VolumetricData::voxel_gradient_interpolate, VolumetricData::voxel_gradient_safe, VolumetricData::voxel_value, VolumetricData::voxel_value_interpolate, and VolumetricData::voxel_value_safe.

Generated on Wed Oct 9 02:43:57 2024 for VMD (current) by doxygen1.2.14 written by Dimitri van Heesch, © 1997-2002