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

CUDAOrbital.cu File Reference

#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <math.h>
#include <cuda.h>
#include "WKFThreads.h"
#include "CUDAKernels.h"
#include "OrbitalJIT.h"

Go to the source code of this file.

Compounds

union  flint_t
struct  orbthrparms

Defines

#define RESTRICT
#define CUERR
#define ANGS_TO_BOHR   1.8897259877218677f
#define UNROLLX   1
#define UNROLLY   1
#define BLOCKSIZEX   8
#define BLOCKSIZEY   8
#define BLOCKSIZE   BLOCKSIZEX * BLOCKSIZEY
#define TILESIZEX   BLOCKSIZEX*UNROLLX
#define TILESIZEY   BLOCKSIZEY*UNROLLY
#define GPU_X_ALIGNMASK   (TILESIZEX - 1)
#define GPU_Y_ALIGNMASK   (TILESIZEY - 1)
#define MEMCOALESCE   384
#define S_SHELL   0
#define P_SHELL   1
#define D_SHELL   2
#define F_SHELL   3
#define G_SHELL   4
#define H_SHELL   5
#define MAX_ATOM_SZ   256
#define MAX_ATOMPOS_SZ   (MAX_ATOM_SZ)
#define MAX_ATOM_BASIS_SZ   (MAX_ATOM_SZ)
#define MAX_ATOMSHELL_SZ   (MAX_ATOM_SZ)
#define MAX_BASIS_SZ   6144
#define MAX_SHELL_SZ   1024
#define MAX_WAVEF_SZ   6144
#define MAXSHELLCOUNT   15
#define MEMCOAMASK   (~15)
#define SHAREDSIZE   256

Typedefs

typedef flint_t flint

Functions

void * cudaorbitalthread (void *)
__global__ void cuorbitalconstmem (int numatoms, float voxelsize, float originx, float originy, float grid_z, int density, float *orbitalgrid)
__global__ void cuorbitaltiledshared (int numatoms, const float *wave_f, const float *basis_array, const flint *atominfo, const int *shellinfo, float voxelsize, float originx, float originy, float grid_z, int density, float *orbitalgrid)
__global__ void cuorbitalcachedglobmem (int numatoms, const float *RESTRICT wave_f, const float *RESTRICT basis_array, const flint *RESTRICT atominfo, const int *RESTRICT shellinfo, float voxelsize, float originx, float originy, float grid_z, int density, float *RESTRICT orbitalgrid)
int computepaddedsize (int orig, int tilesize)
int vmd_cuda_evaluate_orbital_grid (wkf_threadpool_t *devpool, int numatoms, const float *wave_f, int num_wave_f, const float *basis_array, int num_basis, const float *atompos, const int *atom_basis, const int *num_shells_per_atom, const int *num_prim_per_shell, const int *shell_types, int num_shells, const int *numvoxels, float voxelsize, const float *origin, int density, float *orbitalgrid)

Variables

__constant__ float const_atompos [MAX_ATOMPOS_SZ *3]
__constant__ int const_atom_basis [MAX_ATOM_BASIS_SZ]
__constant__ int const_num_shells_per_atom [MAX_ATOMSHELL_SZ]
__constant__ float const_basis_array [MAX_BASIS_SZ]
__constant__ int const_num_prim_per_shell [MAX_SHELL_SZ]
__constant__ int const_shell_types [MAX_SHELL_SZ]
__constant__ float const_wave_f [MAX_WAVEF_SZ]

Define Documentation

#define ANGS_TO_BOHR   1.8897259877218677f
 

Definition at line 54 of file CUDAOrbital.cu.

Referenced by cuorbitalcachedglobmem, cuorbitalconstmem, and cuorbitaltiledshared.

#define BLOCKSIZE   BLOCKSIZEX * BLOCKSIZEY
 

Definition at line 88 of file CUDAOrbital.cu.

#define BLOCKSIZEX   8
 

Definition at line 86 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread.

#define BLOCKSIZEY   8
 

Definition at line 87 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread.

#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 45 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread.

#define D_SHELL   2
 

Definition at line 101 of file CUDAOrbital.cu.

Referenced by cuorbitalcachedglobmem, cuorbitalconstmem, and cuorbitaltiledshared.

#define F_SHELL   3
 

Definition at line 102 of file CUDAOrbital.cu.

Referenced by cuorbitalcachedglobmem, cuorbitalconstmem, and cuorbitaltiledshared.

#define G_SHELL   4
 

Definition at line 103 of file CUDAOrbital.cu.

Referenced by cuorbitalcachedglobmem, cuorbitalconstmem, and cuorbitaltiledshared.

#define GPU_X_ALIGNMASK   (TILESIZEX - 1)
 

Definition at line 93 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread.

#define GPU_Y_ALIGNMASK   (TILESIZEY - 1)
 

Definition at line 94 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread.

#define H_SHELL   5
 

Definition at line 104 of file CUDAOrbital.cu.

#define MAX_ATOM_BASIS_SZ   (MAX_ATOM_SZ)
 

Definition at line 114 of file CUDAOrbital.cu.

#define MAX_ATOM_SZ   256
 

Definition at line 109 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread.

#define MAX_ATOMPOS_SZ   (MAX_ATOM_SZ)
 

Definition at line 111 of file CUDAOrbital.cu.

#define MAX_ATOMSHELL_SZ   (MAX_ATOM_SZ)
 

Definition at line 117 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread.

#define MAX_BASIS_SZ   6144
 

Definition at line 120 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread.

#define MAX_SHELL_SZ   1024
 

Definition at line 123 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread.

#define MAX_WAVEF_SZ   6144
 

Definition at line 127 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread.

#define MAXSHELLCOUNT   15
 

Definition at line 297 of file CUDAOrbital.cu.

Referenced by cuorbitaltiledshared.

#define MEMCOALESCE   384
 

Definition at line 96 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread.

#define MEMCOAMASK   (~15)
 

Definition at line 303 of file CUDAOrbital.cu.

Referenced by cuorbitaltiledshared.

#define P_SHELL   1
 

Definition at line 100 of file CUDAOrbital.cu.

Referenced by cuorbitalcachedglobmem, cuorbitalconstmem, and cuorbitaltiledshared.

#define RESTRICT
 

Definition at line 41 of file CUDAOrbital.cu.

Referenced by cuorbitalcachedglobmem.

#define S_SHELL   0
 

Definition at line 99 of file CUDAOrbital.cu.

Referenced by cuorbitalcachedglobmem, cuorbitalconstmem, and cuorbitaltiledshared.

#define SHAREDSIZE   256
 

Definition at line 310 of file CUDAOrbital.cu.

Referenced by cuorbitaltiledshared.

#define TILESIZEX   BLOCKSIZEX*UNROLLX
 

Definition at line 91 of file CUDAOrbital.cu.

#define TILESIZEY   BLOCKSIZEY*UNROLLY
 

Definition at line 92 of file CUDAOrbital.cu.

#define UNROLLX   1
 

Definition at line 84 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread, and cuorbitalcachedglobmem.

#define UNROLLY   1
 

Definition at line 85 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread.

Typedef Documentation

typedef union flint_t flint
 

Function Documentation

int computepaddedsize (  int    orig,
int    tilesize
)  [static]
 

Definition at line 654 of file CUDAOrbital.cu.

void * cudaorbitalthread (  void *    )  [static]
 

Definition at line 661 of file CUDAOrbital.cu.

References orbthrparms::atom_basis, orbthrparms::atompos, orbthrparms::basis_array, BLOCKSIZEX, BLOCKSIZEY, computepaddedsize, const_atom_basis, const_atompos, const_basis_array, const_num_prim_per_shell, const_num_shells_per_atom, const_shell_types, const_wave_f, CUERR, orbthrparms::density, wkf_tasktile_struct::end, flint_t::f, GPU_X_ALIGNMASK, GPU_Y_ALIGNMASK, flint_t::i, MAX_ATOM_SZ, MAX_ATOMSHELL_SZ, MAX_BASIS_SZ, MAX_SHELL_SZ, MAX_WAVEF_SZ, MEMCOALESCE, NULL, orbthrparms::num_basis, orbthrparms::num_prim_per_shell, orbthrparms::num_shells, orbthrparms::num_shells_per_atom, orbthrparms::num_wave_f, orbthrparms::numatoms, orbthrparms::numvoxels, orbthrparms::orbitalgrid, orbthrparms::origin, orbthrparms::shell_types, wkf_tasktile_struct::start, UNROLLX, UNROLLY, orbthrparms::voxelsize, orbthrparms::wave_f, WKF_SCHED_DONE, wkf_threadpool_next_tile, wkf_threadpool_worker_devscaletile, wkf_threadpool_worker_getdata, and wkf_threadpool_worker_getid.

Referenced by vmd_cuda_evaluate_orbital_grid.

__global__ void cuorbitalcachedglobmem (  int    numatoms,
const float *RESTRICT    wave_f,
const float *RESTRICT    basis_array,
const flint *RESTRICT    atominfo,
const int *RESTRICT    shellinfo,
float    voxelsize,
float    originx,
float    originy,
float    grid_z,
int    density,
float *RESTRICT    orbitalgrid
)  [static]
 

Definition at line 515 of file CUDAOrbital.cu.

References ANGS_TO_BOHR, D_SHELL, F_SHELL, G_SHELL, P_SHELL, RESTRICT, S_SHELL, and UNROLLX.

Referenced by vmd_cuda_evaluate_orbital_grid.

__global__ void cuorbitalconstmem (  int    numatoms,
float    voxelsize,
float    originx,
float    originy,
float    grid_z,
int    density,
float *    orbitalgrid
)  [static]
 

Definition at line 150 of file CUDAOrbital.cu.

References ANGS_TO_BOHR, const_atom_basis, const_atompos, const_basis_array, const_num_prim_per_shell, const_num_shells_per_atom, const_shell_types, const_wave_f, D_SHELL, F_SHELL, G_SHELL, P_SHELL, and S_SHELL.

__global__ void cuorbitaltiledshared (  int    numatoms,
const float *    wave_f,
const float *    basis_array,
const flint *    atominfo,
const int *    shellinfo,
float    voxelsize,
float    originx,
float    originy,
float    grid_z,
int    density,
float *    orbitalgrid
)  [static]
 

Definition at line 313 of file CUDAOrbital.cu.

References ANGS_TO_BOHR, D_SHELL, F_SHELL, G_SHELL, flint_t::i, MAXSHELLCOUNT, MEMCOAMASK, P_SHELL, S_SHELL, and SHAREDSIZE.

Referenced by vmd_cuda_evaluate_orbital_grid.

int vmd_cuda_evaluate_orbital_grid (  wkf_threadpool_t *    devpool,
int    numatoms,
const float *    wave_f,
int    num_wave_f,
const float *    basis_array,
int    num_basis,
const float *    atompos,
const int *    atom_basis,
const int *    num_shells_per_atom,
const int *    num_prim_per_shell,
const int *    shell_types,
int    num_shells,
const int *    numvoxels,
float    voxelsize,
const float *    origin,
int    density,
float *    orbitalgrid
) 
 

Definition at line 976 of file CUDAOrbital.cu.

References orbthrparms::atom_basis, orbthrparms::atompos, orbthrparms::basis_array, cudaorbitalthread, cuorbitalcachedglobmem, cuorbitaltiledshared, orbthrparms::density, wkf_tasktile_struct::end, NULL, orbthrparms::num_basis, orbthrparms::num_prim_per_shell, orbthrparms::num_shells, orbthrparms::num_shells_per_atom, orbthrparms::num_wave_f, orbthrparms::numatoms, orbthrparms::numvoxels, ORBITAL_JIT_CUDA, orbital_jit_generate, ORBITAL_JIT_OPENCL, orbthrparms::orbitalgrid, orbthrparms::origin, orbthrparms::shell_types, wkf_tasktile_struct::start, orbthrparms::voxelsize, orbthrparms::wave_f, wkf_threadpool_launch, and wkf_threadpool_sched_dynamic.

Variable Documentation

__constant__ int const_atom_basis[MAX_ATOM_BASIS_SZ] [static]
 

Definition at line 115 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread, and cuorbitalconstmem.

__constant__ float const_atompos[MAX_ATOMPOS_SZ * 3] [static]
 

Definition at line 112 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread, and cuorbitalconstmem.

__constant__ float const_basis_array[MAX_BASIS_SZ] [static]
 

Definition at line 121 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread, and cuorbitalconstmem.

__constant__ int const_num_prim_per_shell[MAX_SHELL_SZ] [static]
 

Definition at line 124 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread, and cuorbitalconstmem.

__constant__ int const_num_shells_per_atom[MAX_ATOMSHELL_SZ] [static]
 

Definition at line 118 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread, and cuorbitalconstmem.

__constant__ int const_shell_types[MAX_SHELL_SZ] [static]
 

Definition at line 125 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread, and cuorbitalconstmem.

__constant__ float const_wave_f[MAX_WAVEF_SZ] [static]
 

Definition at line 128 of file CUDAOrbital.cu.

Referenced by cudaorbitalthread, and cuorbitalconstmem.

Generated on Sun May 19 01:51:12 2013 for VMD (current) by doxygen1.2.14 written by Dimitri van Heesch, © 1997-2002