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

CUDAMDFF.cu File Reference

CUDA kernels for computing simulated cryo-EM density maps, cross correlations, and other MDFF-related calculations. More...

#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <cuda.h>
#include <float.h>
#include "Inform.h"
#include "utilities.h"
#include "WKFThreads.h"
#include "WKFUtils.h"
#include "CUDAKernels.h"
#include "CUDASpatialSearch.h"
#include "AtomSel.h"
#include "VMDApp.h"
#include "MoleculeList.h"
#include "VolumetricData.h"
#include "VolMapCreate.h"
#include "CUDAMDFF.h"
#include <tcl.h>
#include "TclCommands.h"

Go to the source code of this file.

Defines

#define MIN(X, Y)   (((X)<(Y))? (X) : (Y))
#define MAX(X, Y)   (((X)>(Y))? (X) : (Y))
#define CUERR
#define RESTRICT
#define GGRIDSZ   8.0f
#define GBLOCKSZX   8
#define GBLOCKSZY   8
#define GBLOCKSZZ   2
#define GUNROLL   4
#define TOTALBLOCKSZ   (GBLOCKSZX * GBLOCKSZY * GBLOCKSZZ)
#define MAINDENSITYLOOP
#define CKACCUM(a, b)   a+=b;

Functions

int check_gpu_compute20 (cudaError_t &err)
__host__ __device__ float2 operator+ (float2 a, float2 b)
__host__ __device__ void operator+= (float2 &a, float2 b)
__host__ __device__ float4 operator+ (float4 a, float4 b)
__host__ __device__ void operator+= (float4 &a, float4 b)
__device__ void convert_density (float &df, float df2)
__device__ void convert_density (unsigned short &dh, float df2)
__device__ void reset_atomic_counter (unsigned int *counter)
__device__ void calc_ac_cell_ids (int3 &abmin, int3 &abmax, int3 acncells, float3 acoriginoffset, float gridspacing, float acgridspacing, float invacgridspacing)
__device__ void calc_densities (int xindex, int yindex, int zindex, const float4 *RESTRICT sorted_xyzr, int3 abmin, int3 abmax, int3 acncells, float3 acoriginoffset, const uint2 *RESTRICT cellStartEnd, float gridspacing, float &densityval1,float &densityval2,float &densityval3,float &densityval4)
__global__ void gaussdensity_fast (int natoms, const float4 *RESTRICT sorted_xyzr, int3 volsz, int3 acncells, float3 acoriginoffset, float acgridspacing, float invacgridspacing, const uint2 *RESTRICT cellStartEnd, float gridspacing, unsigned int z, float *RESTRICT densitygrid)
__global__ void gaussdensity_diff (int natoms, const float4 *RESTRICT sorted_xyzr, int3 volsz, int3 acncells, float3 acoriginoffset, float acgridspacing, float invacgridspacing, const uint2 *RESTRICT cellStartEnd, float gridspacing, unsigned int z, int3 refvolsz, int3 refoffset, const float *RESTRICT refmap, float *RESTRICT diffmap)
__device__ float sumabsdiff_sumreduction (int tid, int totaltb, float *sumabsdiffs_s, float *sumabsdiffs)
__global__ void gaussdensity_sumabsdiff (int totaltb, int natoms, const float4 *RESTRICT sorted_xyzr, int3 volsz, int3 acncells, float3 acoriginoffset, float acgridspacing, float invacgridspacing, const uint2 *RESTRICT cellStartEnd, float gridspacing, unsigned int z, int3 refvolsz, int3 refoffset, const float *RESTRICT refmap, float *RESTRICT sumabsdiff)
__device__ float calc_cc (float sums_ref, float sums_synth, float squares_ref, float squares_synth, int lsize, float lcc)
__device__ void cc_sumreduction (int tid, int totaltb, float4 &total_cc_sums, float &total_lcc, int &total_lsize, float4 *RESTRICT tb_cc_sums, float *RESTRICT tb_lcc, int *RESTRICT tb_lsize)
__device__ void thread_cc_sum (float ref, float density, float2 &thread_cc_means, float2 &thread_cc_squares, float &thread_lcc, int &thread_lsize)
__global__ void gaussdensity_cc (int totaltb, int natoms, const float4 *sorted_xyzr, int3 volsz, int3 acncells, float3 acoriginoffset, float acgridspacing, float invacgridspacing, const uint2 *RESTRICT cellStartEnd, float gridspacing, unsigned int z, float threshold, int3 refvolsz, int3 refoffset, const float *refmap, float4 *RESTRICT tb_cc_sums, float *RESTRICT tb_lcc, int *RESTRICT tb_lsize, float *RESTRICT tb_CC)
int vmd_cuda_build_accel (cudaError_t &err, const long int natoms, const int3 volsz, const float gausslim, const float radscale, const float maxrad, const float gridspacing, const float *xyzr_f, float4 *&xyzr_d, float4 *&sorted_xyzr_d, uint2 *&cellStartEnd_d, float &acgridspacing, int3 &accelcells)
int vmd_cuda_destroy_accel (float4 *&xyzr_d, float4 *&sorted_xyzr_d, uint2 *&cellStartEnd_d)
void vmd_cuda_kernel_launch_dims (int3 volsz, dim3 &Bsz, dim3 &Gsz)
int vmd_cuda_gaussdensity_calc (long int natoms, float3 acoriginoffset, int3 acvolsz, const float *xyzr_f, float *volmap, int3 volsz, float maxrad, float radscale, float gridspacing, float gausslim, int3 refvolsz, int3 refoffset, const float *refmap, float *diffmap, int verbose)
int vmd_cuda_cc_calc (long int natoms, float3 acoriginoffset, int3 acvolsz, const float *xyzr_f, int3 volsz, float maxrad, float radscale, float gridspacing, float gausslim, int3 refvolsz, int3 refoffset, const float *devrefmap, float ccthreshdensity, float &result_cc, const float *origin, VolumetricData **spatialccvol, int verbose)
float gausslim_quality (int quality)
int calc_density_bounds (const AtomSel *sel, MoleculeList *mlist, int verbose, int quality, float radscale, float gridspacing, float &maxrad, float *origin, float *xaxis, float *yaxis, float *zaxis, int3 &volsz)
int map_uniform_spacing (double xax, double yax, double zax, int szx, int szy, int szz)
int calc_density_bounds_overlapping_map (int verbose, float &gridspacing, float *origin, float *xaxis, float *yaxis, float *zaxis, int3 &volsz, int3 &refoffset, const VolumetricData *refmap)
float * build_xyzr_from_coords (const AtomSel *sel, const float *atompos, const float *atomradii, const float *origin)
float * build_xyzr_from_sel (const AtomSel *sel, MoleculeList *mlist, const float *origin)
int vmd_cuda_calc_density (const AtomSel *sel, MoleculeList *mlist, int quality, float radscale, float gridspacing, VolumetricData **synthvol, const VolumetricData *refmap, VolumetricData **diffvol, int verbose)
int vmd_cuda_compare_sel_refmap (const AtomSel *sel, MoleculeList *mlist, int quality, float radscale, float gridspacing, const VolumetricData *refmap, VolumetricData **synthvol, VolumetricData **diffvol, VolumetricData **spatialccvol, float *CC, float ccthreshdensity, int verbose)

Variables

__device__ unsigned int tbcatomic [3] = {0, 0, 0}

Detailed Description

CUDA kernels for computing simulated cryo-EM density maps, cross correlations, and other MDFF-related calculations.

"GPU-Accelerated Analysis and Visualization of Large Structures Solved by Molecular Dynamics Flexible Fitting" John E. Stone, Ryan McGreevy, Barry Isralewitz, and Klaus Schulten. Faraday Discussions, 169:265-283, 2014. http://dx.doi.org/10.1039/C4FD00005F

Definition in file CUDAMDFF.cu.

Define Documentation

#define CKACCUM (  a,
b    )     a+=b;
 

Definition at line 708 of file CUDAMDFF.cu.

#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 64 of file CUDAMDFF.cu.

#define GBLOCKSZX   8
 

Definition at line 156 of file CUDAMDFF.cu.

Referenced by vmd_cuda_kernel_launch_dims.

#define GBLOCKSZY   8
 

Definition at line 157 of file CUDAMDFF.cu.

Referenced by vmd_cuda_kernel_launch_dims.

#define GBLOCKSZZ   2
 

Definition at line 159 of file CUDAMDFF.cu.

Referenced by vmd_cuda_kernel_launch_dims.

#define GGRIDSZ   8.0f
 

Definition at line 155 of file CUDAMDFF.cu.

#define GUNROLL   4
 

Definition at line 160 of file CUDAMDFF.cu.

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

#define MAINDENSITYLOOP
 

Value:

float densityval1=0.0f; \
  float densityval2=0.0f; \
  float densityval3=0.0f; \
  float densityval4=0.0f; \
  calc_densities(xindex, yindex, zindex, sorted_xyzr, abmin, abmax, acncells, \
                 acoriginoffset, cellStartEnd, gridspacing,                   \
                 densityval1, densityval2, densityval3, densityval4);

Definition at line 288 of file CUDAMDFF.cu.

Referenced by gaussdensity_cc, gaussdensity_diff, gaussdensity_fast, and gaussdensity_sumabsdiff.

#define MAX (  X,
Y    )     (((X)>(Y))? (X) : (Y))
 

Definition at line 61 of file CUDAMDFF.cu.

Referenced by calc_density_bounds, QuickSurf::calc_surf, VolMapCreate::combo_addframe, VolMapCreateOccupancy::compute_frame, VolMapCreateDensity::compute_frame, VolMapCreateMask::compute_frame, VolMapCreateDensity::compute_init, VolMapCreate::compute_init, init_from_intersection, init_from_union, overlap_f, VolumetricData::pad, IsoSurface::set_color_voltex_rgb3fv, vmd_gaussdensity_avx2, and vmd_gaussdensity_opt.

#define MIN (  X,
Y    )     (((X)<(Y))? (X) : (Y))
 

Definition at line 60 of file CUDAMDFF.cu.

Referenced by VolMapCreate::combo_addframe, VolMapCreateOccupancy::compute_frame, VolMapCreateDensity::compute_frame, VolMapCreateMask::compute_frame, init_from_intersection, init_from_union, overlap_f, VolumetricData::pad, rdf_cpu, IsoSurface::set_color_voltex_rgb3fv, vmd_gaussdensity_avx2, and vmd_gaussdensity_opt.

#define RESTRICT
 

Definition at line 148 of file CUDAMDFF.cu.

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

#define TOTALBLOCKSZ   (GBLOCKSZX * GBLOCKSZY * GBLOCKSZZ)
 

Definition at line 162 of file CUDAMDFF.cu.

Referenced by gaussdensity_cc, and gaussdensity_sumabsdiff.

Function Documentation

float* build_xyzr_from_coords (  const AtomSel *    sel,
const float *    atompos,
const float *    atomradii,
const float *    origin
)  [static]
 

Definition at line 2075 of file CUDAMDFF.cu.

References AtomSel::firstsel, AtomSel::lastsel, AtomSel::on, and AtomSel::selected.

Referenced by build_xyzr_from_sel.

float* build_xyzr_from_sel (  const AtomSel *    sel,
MoleculeList *    mlist,
const float *    origin
)  [static]
 

Definition at line 2102 of file CUDAMDFF.cu.

References build_xyzr_from_coords, AtomSel::coordinates, NameList< float * >::data, BaseMolecule::extraflt, MoleculeList::mol_from_id, and AtomSel::molid.

Referenced by vmd_cuda_calc_density, and vmd_cuda_compare_sel_refmap.

__device__ void calc_ac_cell_ids (  int3 &    abmin,
int3 &    abmax,
int3    acncells,
float3    acoriginoffset,
float    gridspacing,
float    acgridspacing,
float    invacgridspacing
)  [inline]
 

Definition at line 186 of file CUDAMDFF.cu.

References acgridspacing, acncells, gridspacing, GUNROLL, and invacgridspacing.

Referenced by gaussdensity_cc, gaussdensity_diff, gaussdensity_fast, and gaussdensity_sumabsdiff.

__device__ float calc_cc (  float    sums_ref,
float    sums_synth,
float    squares_ref,
float    squares_synth,
int    lsize,
float    lcc
)  [inline]
 

Definition at line 555 of file CUDAMDFF.cu.

References cc.

__device__ void calc_densities (  int    xindex,
int    yindex,
int    zindex,
const float4 *RESTRICT    sorted_xyzr,
int3    abmin,
int3    abmax,
int3    acncells,
float3    acoriginoffset,
const uint2 *RESTRICT    cellStartEnd,
float    gridspacing,
float &    densityval1,
float &    densityval2,
float &    densityval3,
float &    densityval4
)  [inline]
 

Definition at line 212 of file CUDAMDFF.cu.

References acncells, cellStartEnd, gridspacing, GUNROLL, RESTRICT, and sorted_xyzr.

int calc_density_bounds (  const AtomSel *    sel,
MoleculeList *    mlist,
int    verbose,
int    quality,
float    radscale,
float    gridspacing,
float &    maxrad,
float *    origin,
float *    xaxis,
float *    yaxis,
float *    zaxis,
int3 &    volsz
)  [static]
 

Definition at line 1896 of file CUDAMDFF.cu.

References AtomSel::coordinates, NameList< float * >::data, BaseMolecule::extraflt, AtomSel::firstsel, BaseMolecule::get_radii_minmax, gridspacing, AtomSel::lastsel, MAX, minmax_selected_3fv_aligned, MoleculeList::mol_from_id, AtomSel::molid, AtomSel::num_atoms, AtomSel::on, vec_copy, vec_zero, VMD_PI, and volsz.

Referenced by vmd_cuda_calc_density, and vmd_cuda_compare_sel_refmap.

int calc_density_bounds_overlapping_map (  int    verbose,
float &    gridspacing,
float *    origin,
float *    xaxis,
float *    yaxis,
float *    zaxis,
int3 &    volsz,
int3 &    refoffset,
const VolumetricData *    refmap
)  [static]
 

Definition at line 2002 of file CUDAMDFF.cu.

References gridspacing, VolumetricData::origin, volsz, VolumetricData::xaxis, VolumetricData::xsize, VolumetricData::yaxis, VolumetricData::ysize, VolumetricData::zaxis, and VolumetricData::zsize.

Referenced by vmd_cuda_calc_density, and vmd_cuda_compare_sel_refmap.

__device__ void cc_sumreduction (  int    tid,
int    totaltb,
float4 &    total_cc_sums,
float &    total_lcc,
int &    total_lsize,
float4 *RESTRICT    tb_cc_sums,
float *RESTRICT    tb_lcc,
int *RESTRICT    tb_lsize
)  [inline]
 

Definition at line 633 of file CUDAMDFF.cu.

References make_float4, and RESTRICT.

Referenced by gaussdensity_cc.

int check_gpu_compute20 (  cudaError_t &    err )  [static]
 

Definition at line 75 of file CUDAMDFF.cu.

Referenced by vmd_cuda_cc_calc, and vmd_cuda_gaussdensity_calc.

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

Definition at line 137 of file CUDAMDFF.cu.

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

Definition at line 131 of file CUDAMDFF.cu.

__global__ void gaussdensity_cc (  int    totaltb,
int    natoms,
const float4 *    sorted_xyzr,
int3    volsz,
int3    acncells,
float3    acoriginoffset,
float    acgridspacing,
float    invacgridspacing,
const uint2 *RESTRICT    cellStartEnd,
float    gridspacing,
unsigned int    z,
float    threshold,
int3    refvolsz,
int3    refoffset,
const float *    refmap,
float4 *RESTRICT    tb_cc_sums,
float *RESTRICT    tb_lcc,
int *RESTRICT    tb_lsize,
float *RESTRICT    tb_CC
)  [static]
 

Definition at line 713 of file CUDAMDFF.cu.

References acgridspacing, acncells, calc_ac_cell_ids, calc_cc, cc, cc_sumreduction, cellStartEnd, gridspacing, GUNROLL, invacgridspacing, MAINDENSITYLOOP, make_float2, make_float4, mask, NULL, reset_atomic_counter, RESTRICT, sorted_xyzr, tbcatomic, thread_cc_sum, threshold, TOTALBLOCKSZ, volsz, and z.

__global__ void gaussdensity_diff (  int    natoms,
const float4 *RESTRICT    sorted_xyzr,
int3    volsz,
int3    acncells,
float3    acoriginoffset,
float    acgridspacing,
float    invacgridspacing,
const uint2 *RESTRICT    cellStartEnd,
float    gridspacing,
unsigned int    z,
int3    refvolsz,
int3    refoffset,
const float *RESTRICT    refmap,
float *RESTRICT    diffmap
)  [static]
 

Definition at line 351 of file CUDAMDFF.cu.

References acgridspacing, acncells, calc_ac_cell_ids, cellStartEnd, gridspacing, GUNROLL, invacgridspacing, MAINDENSITYLOOP, RESTRICT, sorted_xyzr, volsz, and z.

__global__ void gaussdensity_fast (  int    natoms,
const float4 *RESTRICT    sorted_xyzr,
int3    volsz,
int3    acncells,
float3    acoriginoffset,
float    acgridspacing,
float    invacgridspacing,
const uint2 *RESTRICT    cellStartEnd,
float    gridspacing,
unsigned int    z,
float *RESTRICT    densitygrid
)  [static]
 

Definition at line 302 of file CUDAMDFF.cu.

References acgridspacing, acncells, calc_ac_cell_ids, cellStartEnd, densitygrid, gridspacing, GUNROLL, invacgridspacing, MAINDENSITYLOOP, RESTRICT, sorted_xyzr, volsz, and z.

__global__ void gaussdensity_sumabsdiff (  int    totaltb,
int    natoms,
const float4 *RESTRICT    sorted_xyzr,
int3    volsz,
int3    acncells,
float3    acoriginoffset,
float    acgridspacing,
float    invacgridspacing,
const uint2 *RESTRICT    cellStartEnd,
float    gridspacing,
unsigned int    z,
int3    refvolsz,
int3    refoffset,
const float *RESTRICT    refmap,
float *RESTRICT    sumabsdiff
)  [static]
 

Definition at line 436 of file CUDAMDFF.cu.

References acgridspacing, acncells, calc_ac_cell_ids, cellStartEnd, gridspacing, GUNROLL, invacgridspacing, MAINDENSITYLOOP, reset_atomic_counter, RESTRICT, sorted_xyzr, sumabsdiff_sumreduction, tbcatomic, TOTALBLOCKSZ, volsz, and z.

float gausslim_quality (  int    quality )  [static]
 

Definition at line 1876 of file CUDAMDFF.cu.

Referenced by vmd_cuda_calc_density, and vmd_cuda_compare_sel_refmap.

int map_uniform_spacing (  double    xax,
double    yax,
double    zax,
int    szx,
int    szy,
int    szz
)  [static]
 

Definition at line 1977 of file CUDAMDFF.cu.

Referenced by vmd_cuda_compare_sel_refmap.

__host__ __device__ float4 operator+ (  float4    a,
float4    b
)  [inline]
 

Definition at line 112 of file CUDAMDFF.cu.

References make_float4.

__host__ __device__ float2 operator+ (  float2    a,
float2    b
)  [inline]
 

Definition at line 102 of file CUDAMDFF.cu.

References make_float2.

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

Definition at line 117 of file CUDAMDFF.cu.

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

Definition at line 107 of file CUDAMDFF.cu.

__device__ void reset_atomic_counter (  unsigned int *    counter ) 
 

Definition at line 175 of file CUDAMDFF.cu.

Referenced by gaussdensity_cc, and gaussdensity_sumabsdiff.

__device__ float sumabsdiff_sumreduction (  int    tid,
int    totaltb,
float *    sumabsdiffs_s,
float *    sumabsdiffs
) 
 

Definition at line 407 of file CUDAMDFF.cu.

Referenced by gaussdensity_sumabsdiff.

__device__ void thread_cc_sum (  float    ref,
float    density,
float2 &    thread_cc_means,
float2 &    thread_cc_squares,
float &    thread_lcc,
int &    thread_lsize
)  [inline]
 

Definition at line 677 of file CUDAMDFF.cu.

Referenced by gaussdensity_cc.

int vmd_cuda_build_accel (  cudaError_t &    err,
const long int    natoms,
const int3    volsz,
const float    gausslim,
const float    radscale,
const float    maxrad,
const float    gridspacing,
const float *    xyzr_f,
float4 *&    xyzr_d,
float4 *&    sorted_xyzr_d,
uint2 *&    cellStartEnd_d,
float &    acgridspacing,
int3 &    accelcells
)  [static]
 

Definition at line 1182 of file CUDAMDFF.cu.

References acgridspacing, gridspacing, NULL, vmd_cuda_build_density_atom_grid, and volsz.

Referenced by vmd_cuda_cc_calc, and vmd_cuda_gaussdensity_calc.

int vmd_cuda_calc_density (  const AtomSel *    sel,
MoleculeList *    mlist,
int    quality,
float    radscale,
float    gridspacing,
VolumetricData **    synthvol,
const VolumetricData *    refmap,
VolumetricData **    diffvol,
int    verbose
) 
 

Definition at line 2118 of file CUDAMDFF.cu.

References build_xyzr_from_sel, calc_density_bounds, calc_density_bounds_overlapping_map, VolumetricData::data, gausslim_quality, gridspacing, make_float3, NULL, VolumetricData::origin, AtomSel::selected, vec_copy, vmd_cuda_gaussdensity_calc, volmap, volsz, wkf_timer_create, wkf_timer_destroy, wkf_timer_start, wkf_timer_timenow, wkf_timerhandle, VolumetricData::xaxis, VolumetricData::xsize, VolumetricData::yaxis, VolumetricData::ysize, VolumetricData::zaxis, and VolumetricData::zsize.

int vmd_cuda_cc_calc (  long int    natoms,
float3    acoriginoffset,
int3    acvolsz,
const float *    xyzr_f,
int3    volsz,
float    maxrad,
float    radscale,
float    gridspacing,
float    gausslim,
int3    refvolsz,
int3    refoffset,
const float *    devrefmap,
float    ccthreshdensity,
float &    result_cc,
const float *    origin,
VolumetricData **    spatialccvol,
int    verbose
)  [static]
 

Definition at line 1598 of file CUDAMDFF.cu.

References acgridspacing, cc, check_gpu_compute20, gridspacing, GUNROLL, NULL, vmd_cuda_build_accel, vmd_cuda_destroy_accel, vmd_cuda_kernel_launch_dims, volsz, wkf_timer_create, wkf_timer_destroy, wkf_timer_start, wkf_timer_stop, wkf_timer_time, wkf_timer_timenow, wkf_timerhandle, and z.

Referenced by vmd_cuda_compare_sel_refmap.

int vmd_cuda_compare_sel_refmap (  const AtomSel *    sel,
MoleculeList *    mlist,
int    quality,
float    radscale,
float    gridspacing,
const VolumetricData *    refmap,
VolumetricData **    synthvol,
VolumetricData **    diffvol,
VolumetricData **    spatialccvol,
float *    CC,
float    ccthreshdensity,
int    verbose
) 
 

Definition at line 2324 of file CUDAMDFF.cu.

References build_xyzr_from_sel, calc_density_bounds, calc_density_bounds_overlapping_map, VolumetricData::data, gausslim_quality, gridspacing, make_float3, map_uniform_spacing, NULL, AtomSel::selected, vec_copy, vmd_cuda_cc_calc, vmd_cuda_gaussdensity_calc, volsz, wkf_timer_create, wkf_timer_destroy, wkf_timer_start, wkf_timer_timenow, wkf_timerhandle, VolumetricData::xaxis, VolumetricData::xsize, VolumetricData::yaxis, VolumetricData::ysize, VolumetricData::zaxis, and VolumetricData::zsize.

int vmd_cuda_destroy_accel (  float4 *&    xyzr_d,
float4 *&    sorted_xyzr_d,
uint2 *&    cellStartEnd_d
)  [static]
 

Definition at line 1245 of file CUDAMDFF.cu.

References NULL.

Referenced by vmd_cuda_cc_calc, and vmd_cuda_gaussdensity_calc.

int vmd_cuda_gaussdensity_calc (  long int    natoms,
float3    acoriginoffset,
int3    acvolsz,
const float *    xyzr_f,
float *    volmap,
int3    volsz,
float    maxrad,
float    radscale,
float    gridspacing,
float    gausslim,
int3    refvolsz,
int3    refoffset,
const float *    refmap,
float *    diffmap,
int    verbose
)  [static]
 

Definition at line 1279 of file CUDAMDFF.cu.

References acgridspacing, cc, check_gpu_compute20, gridspacing, make_float4, NULL, vmd_cuda_build_accel, vmd_cuda_destroy_accel, vmd_cuda_kernel_launch_dims, volmap, volsz, wkf_timer_create, wkf_timer_destroy, wkf_timer_start, wkf_timer_stop, wkf_timer_time, wkf_timer_timenow, wkf_timerhandle, and z.

Referenced by vmd_cuda_calc_density, and vmd_cuda_compare_sel_refmap.

void vmd_cuda_kernel_launch_dims (  int3    volsz,
dim3 &    Bsz,
dim3 &    Gsz
)  [static]
 

Definition at line 1266 of file CUDAMDFF.cu.

References GBLOCKSZX, GBLOCKSZY, GBLOCKSZZ, GUNROLL, and volsz.

Referenced by vmd_cuda_cc_calc, and vmd_cuda_gaussdensity_calc.

Variable Documentation

__device__ unsigned int tbcatomic[3] = {0, 0, 0}
 

Definition at line 173 of file CUDAMDFF.cu.

Referenced by gaussdensity_cc, and gaussdensity_sumabsdiff.

Generated on Fri Nov 8 02:45:57 2024 for VMD (current) by doxygen1.2.14 written by Dimitri van Heesch, © 1997-2002