#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) |
__global__ void | cuorbitalcachedglobmem_packed (int numatoms, const float *RESTRICT wave_f, const float2 *RESTRICT basis_array, const flint *RESTRICT atominfo, const int2 *RESTRICT shellinfo, float voxelsize, float originx, float originy, float grid_z, int density, float *RESTRICT orbitalgrid) |
__host__ __device__ float2 | make_float2 (const float s) |
__host__ __device__ float2 | operator+ (const float2 &a, const float2 &b) |
__host__ __device__ void | operator+= (float2 &a, const float &b) |
__host__ __device__ void | operator+= (float2 &a, const float2 &b) |
__host__ __device__ float2 | operator- (const float2 &a, const float s) |
__host__ __device__ float2 | operator * (const float2 &a, const float2 &b) |
__host__ __device__ float2 | operator * (const float s, const float2 &a) |
__host__ __device__ float2 | operator * (const float2 &a, const float s) |
__global__ void | cuorbitalcachedglobmem_packed_2x (int numatoms, const float *RESTRICT wave_f, const float2 *RESTRICT basis_array, const flint *RESTRICT atominfo, const int2 *RESTRICT shellinfo, float voxelsize, float originx, float originy, float grid_z, int density, float2 *RESTRICT orbitalgrid) |
__host__ __device__ float4 | make_float4 (const float s) |
__host__ __device__ float4 | operator+ (const float4 &a, const float4 &b) |
__host__ __device__ float4 | make_float4 (const float3 &a, const float &b) |
__host__ __device__ void | operator+= (float4 &a, const float &b) |
__host__ __device__ void | operator+= (float4 &a, const float4 &b) |
__host__ __device__ float4 | operator- (const float4 &a, const float s) |
__host__ __device__ void | operator *= (float4 &a, const float &b) |
__host__ __device__ float4 | operator * (const float4 &a, const float4 &b) |
__host__ __device__ float4 | operator * (const float4 &a, const float s) |
__host__ __device__ float4 | operator * (const float s, const float4 &a) |
__global__ void | cuorbitalcachedglobmem_packed_4x_3D (int numatoms, const float *RESTRICT wave_f, const float2 *RESTRICT basis_array, const flint *RESTRICT atominfo, const int2 *RESTRICT shellinfo, float voxelsize, float originx, float originy, float grid_z, int density, float4 *RESTRICT orbitalgrid) |
__global__ void | cuorbitalcachedglobmem_packed_4x (int numatoms, const float *RESTRICT wave_f, const float2 *RESTRICT basis_array, const flint *RESTRICT atominfo, const int2 *RESTRICT shellinfo, float voxelsize, float originx, float originy, float grid_z, int density, float4 *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] |
This work is described in the following papers:
"NAMD goes quantum: an integrative suite for hybrid simulations" Marcelo C. R. Melo, Rafael C. Bernardi, Till Rudack, Maximilian Scheurer, Christoph Riplinger, James C. Phillips, Julio D. C. Maia, Gerd B. Rocha, Joćo V. Ribeiro, John E. Stone, Frank Neese, Klaus Schulten, and Zaida Luthey-Schulten. Nature Methods, 15: 351-354, 2018. https://www.nature.com/articles/nmeth.4638
"Early Experiences Porting the NAMD and VMD Molecular Simulation and Analysis Software to GPU-Accelerated OpenPOWER Platforms" John E. Stone, Antti-Pekka Hynninen, James C. Phillips, and Klaus Schulten. International Workshop on OpenPOWER for HPC (IWOPH'16), LNCS 9945, pp. 188-206, 2016. http://dx.doi.org/10.1007/978-3-319-46079-6_14
"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
"High Performance Computation and Interactive Display of Molecular Orbitals on GPUs and Multi-core CPUs" John E. Stone, Jan Saam, David J. Hardy, Kirby L. Vandivort, Wen-mei W. Hwu, and Klaus Schulten. In Proceedings of the 2nd Workshop on General-Purpose Processing on Graphics Processing Units, ACM International Conference Proceeding Series, volume 383, pp. 9-18, 2009. http://doi.acm.org/10.1145/1513895.1513897
Definition in file CUDAOrbital.cu.
|
Definition at line 89 of file CUDAOrbital.cu. Referenced by cuorbitalcachedglobmem, cuorbitalcachedglobmem_packed, cuorbitalcachedglobmem_packed_2x, cuorbitalcachedglobmem_packed_4x, cuorbitalcachedglobmem_packed_4x_3D, cuorbitalconstmem, and cuorbitaltiledshared. |
|
Definition at line 123 of file CUDAOrbital.cu. |
|
Definition at line 121 of file CUDAOrbital.cu. Referenced by cudaorbitalthread. |
|
Definition at line 122 of file CUDAOrbital.cu. Referenced by cudaorbitalthread. |
|
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 80 of file CUDAOrbital.cu. Referenced by cudaorbitalthread. |
|
Definition at line 136 of file CUDAOrbital.cu. Referenced by cuorbitalcachedglobmem, cuorbitalcachedglobmem_packed, cuorbitalcachedglobmem_packed_2x, cuorbitalcachedglobmem_packed_4x, cuorbitalcachedglobmem_packed_4x_3D, cuorbitalconstmem, and cuorbitaltiledshared. |
|
Definition at line 137 of file CUDAOrbital.cu. Referenced by cuorbitalcachedglobmem, cuorbitalcachedglobmem_packed, cuorbitalcachedglobmem_packed_2x, cuorbitalcachedglobmem_packed_4x, cuorbitalcachedglobmem_packed_4x_3D, cuorbitalconstmem, and cuorbitaltiledshared. |
|
Definition at line 138 of file CUDAOrbital.cu. Referenced by cuorbitalcachedglobmem, cuorbitalcachedglobmem_packed, cuorbitalcachedglobmem_packed_2x, cuorbitalcachedglobmem_packed_4x, cuorbitalcachedglobmem_packed_4x_3D, cuorbitalconstmem, and cuorbitaltiledshared. |
|
Definition at line 128 of file CUDAOrbital.cu. Referenced by cudaorbitalthread. |
|
Definition at line 129 of file CUDAOrbital.cu. Referenced by cudaorbitalthread. |
|
Definition at line 139 of file CUDAOrbital.cu. |
|
Definition at line 149 of file CUDAOrbital.cu. |
|
Definition at line 144 of file CUDAOrbital.cu. Referenced by cudaorbitalthread. |
|
Definition at line 146 of file CUDAOrbital.cu. |
|
Definition at line 152 of file CUDAOrbital.cu. Referenced by cudaorbitalthread. |
|
Definition at line 155 of file CUDAOrbital.cu. Referenced by cudaorbitalthread. |
|
Definition at line 158 of file CUDAOrbital.cu. Referenced by cudaorbitalthread. |
|
Definition at line 162 of file CUDAOrbital.cu. Referenced by cudaorbitalthread. |
|
Definition at line 332 of file CUDAOrbital.cu. Referenced by cuorbitaltiledshared. |
|
Definition at line 131 of file CUDAOrbital.cu. Referenced by cudaorbitalthread. |
|
Definition at line 338 of file CUDAOrbital.cu. Referenced by cuorbitaltiledshared. |
|
Definition at line 135 of file CUDAOrbital.cu. Referenced by cuorbitalcachedglobmem, cuorbitalcachedglobmem_packed, cuorbitalcachedglobmem_packed_2x, cuorbitalcachedglobmem_packed_4x, cuorbitalcachedglobmem_packed_4x_3D, cuorbitalconstmem, and cuorbitaltiledshared. |
|
Definition at line 76 of file CUDAOrbital.cu. Referenced by cuorbitalcachedglobmem, cuorbitalcachedglobmem_packed, cuorbitalcachedglobmem_packed_2x, cuorbitalcachedglobmem_packed_4x, and cuorbitalcachedglobmem_packed_4x_3D. |
|
Definition at line 134 of file CUDAOrbital.cu. Referenced by cuorbitalcachedglobmem, cuorbitalcachedglobmem_packed, cuorbitalcachedglobmem_packed_2x, cuorbitalcachedglobmem_packed_4x, cuorbitalcachedglobmem_packed_4x_3D, cuorbitalconstmem, and cuorbitaltiledshared. |
|
Definition at line 345 of file CUDAOrbital.cu. Referenced by cuorbitaltiledshared. |
|
Definition at line 126 of file CUDAOrbital.cu. |
|
Definition at line 127 of file CUDAOrbital.cu. |
|
Definition at line 119 of file CUDAOrbital.cu. Referenced by cudaorbitalthread, cuorbitalcachedglobmem, cuorbitalcachedglobmem_packed, cuorbitalcachedglobmem_packed_2x, cuorbitalcachedglobmem_packed_4x, and cuorbitalcachedglobmem_packed_4x_3D. |
|
Definition at line 120 of file CUDAOrbital.cu. Referenced by cudaorbitalthread. |
|
|
|
Definition at line 1375 of file CUDAOrbital.cu. |
|
|
Definition at line 550 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. |
|
Definition at line 691 of file CUDAOrbital.cu. References ANGS_TO_BOHR, D_SHELL, F_SHELL, G_SHELL, P_SHELL, RESTRICT, S_SHELL, and UNROLLX. |
|
Definition at line 868 of file CUDAOrbital.cu. References ANGS_TO_BOHR, D_SHELL, F_SHELL, G_SHELL, make_float2, P_SHELL, RESTRICT, S_SHELL, and UNROLLX. |
|
Definition at line 1222 of file CUDAOrbital.cu. References ANGS_TO_BOHR, D_SHELL, F_SHELL, G_SHELL, make_float4, P_SHELL, RESTRICT, S_SHELL, and UNROLLX. |
|
Definition at line 1063 of file CUDAOrbital.cu. References ANGS_TO_BOHR, D_SHELL, F_SHELL, G_SHELL, make_float4, P_SHELL, RESTRICT, S_SHELL, and UNROLLX. |
|
Definition at line 185 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. |
|
Definition at line 348 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. |
|
Definition at line 829 of file CUDAOrbital.cu. Referenced by cuorbitalcachedglobmem_packed_2x, gaussdensity_cc, operator *, operator+, operator-, vmd_camera_cubemap_general, vmd_camera_dome_general, vmd_camera_equirectangular_general, vmd_camera_oculus_rift_general, vmd_camera_orthographic_general, and vmd_camera_perspective_general. |
|
Definition at line 1025 of file CUDAOrbital.cu. Referenced by accumulate_color, cc_sumreduction, clear_accumulation_buffer, clip_ray_by_plane, cudaglobmembw, cuorbitalcachedglobmem_packed_4x, cuorbitalcachedglobmem_packed_4x_3D, gaussdensity_cc, make_float4, operator *, operator+, operator-, and vmd_cuda_gaussdensity_calc. |
|
Definition at line 1017 of file CUDAOrbital.cu. References make_float4. |
|
Definition at line 1053 of file CUDAOrbital.cu. References make_float4. |
|
Definition at line 1049 of file CUDAOrbital.cu. References make_float4. |
|
Definition at line 1045 of file CUDAOrbital.cu. References make_float4. |
|
Definition at line 857 of file CUDAOrbital.cu. References make_float2. |
|
Definition at line 853 of file CUDAOrbital.cu. References make_float2. |
|
Definition at line 849 of file CUDAOrbital.cu. References make_float2. |
|
Definition at line 1041 of file CUDAOrbital.cu. |
|
Definition at line 1021 of file CUDAOrbital.cu. References make_float4. |
|
Definition at line 833 of file CUDAOrbital.cu. References make_float2. |
|
Definition at line 1033 of file CUDAOrbital.cu. |
|
Definition at line 1029 of file CUDAOrbital.cu. |
|
Definition at line 841 of file CUDAOrbital.cu. |
|
Definition at line 837 of file CUDAOrbital.cu. |
|
Definition at line 1037 of file CUDAOrbital.cu. References make_float4. |
|
Definition at line 845 of file CUDAOrbital.cu. References make_float2. |
|
|
Definition at line 150 of file CUDAOrbital.cu. Referenced by cudaorbitalthread, and cuorbitalconstmem. |
|
Definition at line 147 of file CUDAOrbital.cu. Referenced by cudaorbitalthread, and cuorbitalconstmem. |
|
Definition at line 156 of file CUDAOrbital.cu. Referenced by cudaorbitalthread, and cuorbitalconstmem. |
|
Definition at line 159 of file CUDAOrbital.cu. Referenced by cudaorbitalthread, and cuorbitalconstmem. |
|
Definition at line 153 of file CUDAOrbital.cu. Referenced by cudaorbitalthread, and cuorbitalconstmem. |
|
Definition at line 160 of file CUDAOrbital.cu. Referenced by cudaorbitalthread, and cuorbitalconstmem. |
|
Definition at line 163 of file CUDAOrbital.cu. Referenced by cudaorbitalthread, and cuorbitalconstmem. |