Skip to content

Commit

Permalink
adapt to HOOMD-blue 2.5.0
Browse files Browse the repository at this point in the history
  • Loading branch information
Jens Glaser committed Apr 16, 2019
1 parent b0dd6c1 commit b6802f2
Show file tree
Hide file tree
Showing 6 changed files with 49 additions and 49 deletions.
24 changes: 12 additions & 12 deletions metadynamics/OrderParameterMesh.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -48,15 +48,15 @@ OrderParameterMesh::OrderParameterMesh(std::shared_ptr<SystemDefinition> sysdef,
throw std::runtime_error("Error setting up cv.mesh");
}

GPUArray<Scalar> mode_array(m_pdata->getNTypes(), m_exec_conf);
GlobalArray<Scalar> mode_array(m_pdata->getNTypes(), m_exec_conf);
m_mode.swap(mode_array);

m_mode_sq = 0.0;

ArrayHandle<Scalar> h_mode(m_mode, access_location::host, access_mode::overwrite);
std::copy(mode.begin(), mode.end(), h_mode.data);

GPUArray<int3> zero_modes_array(zero_modes.size(), m_exec_conf);
GlobalArray<int3> zero_modes_array(zero_modes.size(), m_exec_conf);
m_zero_modes.swap(zero_modes_array);

ArrayHandle<int3> h_zero_modes(m_zero_modes, access_location::host, access_mode::overwrite);
Expand Down Expand Up @@ -170,10 +170,10 @@ void OrderParameterMesh::setTable(const std::vector<Scalar> &K,
m_delta_k = (kmax - kmin) / Scalar(K.size() - 1);

// allocate the arrays
GPUArray<Scalar> table(K.size(), m_exec_conf);
GlobalArray<Scalar> table(K.size(), m_exec_conf);
m_table.swap(table);

GPUArray<Scalar> table_d(d_K.size(), m_exec_conf);
GlobalArray<Scalar> table_d(d_K.size(), m_exec_conf);
m_table_d.swap(table_d);

// access the array
Expand Down Expand Up @@ -212,17 +212,17 @@ void OrderParameterMesh::setupMesh()
m_n_inner_cells = m_mesh_points.x * m_mesh_points.y * m_mesh_points.z;

// allocate memory for influence function and k values
GPUArray<Scalar> inf_f(m_n_inner_cells, m_exec_conf);
GlobalArray<Scalar> inf_f(m_n_inner_cells, m_exec_conf);
m_inf_f.swap(inf_f);

// allocate memory for interpolationluence function and k values
GPUArray<Scalar> interpolation_f(m_n_inner_cells, m_exec_conf);
GlobalArray<Scalar> interpolation_f(m_n_inner_cells, m_exec_conf);
m_interpolation_f.swap(interpolation_f);

GPUArray<Scalar3> k(m_n_inner_cells, m_exec_conf);
GlobalArray<Scalar3> k(m_n_inner_cells, m_exec_conf);
m_k.swap(k);

GPUArray<Scalar> virial_mesh(6*m_n_inner_cells, m_exec_conf);
GlobalArray<Scalar> virial_mesh(6*m_n_inner_cells, m_exec_conf);
m_virial_mesh.swap(virial_mesh);

initializeFFT();
Expand Down Expand Up @@ -328,16 +328,16 @@ void OrderParameterMesh::initializeFFT()
}

// allocate mesh and transformed mesh
GPUArray<kiss_fft_cpx> mesh(m_n_cells,m_exec_conf);
GlobalArray<kiss_fft_cpx> mesh(m_n_cells,m_exec_conf);
m_mesh.swap(mesh);

GPUArray<kiss_fft_cpx> fourier_mesh(m_n_inner_cells, m_exec_conf);
GlobalArray<kiss_fft_cpx> fourier_mesh(m_n_inner_cells, m_exec_conf);
m_fourier_mesh.swap(fourier_mesh);

GPUArray<kiss_fft_cpx> fourier_mesh_G(m_n_inner_cells, m_exec_conf);
GlobalArray<kiss_fft_cpx> fourier_mesh_G(m_n_inner_cells, m_exec_conf);
m_fourier_mesh_G.swap(fourier_mesh_G);

GPUArray<kiss_fft_cpx> inv_fourier_mesh(m_n_cells, m_exec_conf);
GlobalArray<kiss_fft_cpx> inv_fourier_mesh(m_n_cells, m_exec_conf);
m_inv_fourier_mesh.swap(inv_fourier_mesh);
}

Expand Down
24 changes: 12 additions & 12 deletions metadynamics/OrderParameterMesh.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -80,30 +80,30 @@ class OrderParameterMesh : public CollectiveVariable
unsigned int m_n_cells; //!< Total number of inner cells
unsigned int m_radius; //!< Stencil radius (in units of mesh size)
unsigned int m_n_inner_cells; //!< Number of inner mesh points (without ghost cells)
GPUArray<Scalar> m_mode; //!< Per-type scalar multiplying density ("charges")
GlobalArray<Scalar> m_mode; //!< Per-type scalar multiplying density ("charges")
Scalar m_mode_sq; //!< Sum of squared mode amplitudes
GPUArray<Scalar> m_inf_f; //!< Fourier representation of the influence function (real part)
GPUArray<Scalar> m_interpolation_f; //!< Fourier representation of the interpolation function
GPUArray<Scalar3> m_k; //!< Mesh of k values
GlobalArray<Scalar> m_inf_f; //!< Fourier representation of the influence function (real part)
GlobalArray<Scalar> m_interpolation_f; //!< Fourier representation of the interpolation function
GlobalArray<Scalar3> m_k; //!< Mesh of k values
Scalar m_qstarsq; //!< Short wave length cut-off squared for density harmonics
bool m_is_first_step; //!< True if we have not yet computed the influence function
unsigned int m_cv_last_updated; //!< Timestep of last update of collective variable
bool m_box_changed; //!< True if box has changed since last compute
Scalar m_cv; //!< Current value of collective variable

GPUArray<Scalar> m_virial_mesh; //!< k-space mesh of virial tensor values
GlobalArray<Scalar> m_virial_mesh; //!< k-space mesh of virial tensor values

unsigned int m_q_max_last_computed; //!< Last time step at which q max was computed
Scalar3 m_q_max; //!< Current wave vector with maximum amplitude
Scalar m_sq_max; //!< Maximum structure factor

GPUArray<int3> m_zero_modes; //!< Fourier modes that should be zeroed
GlobalArray<int3> m_zero_modes; //!< Fourier modes that should be zeroed

Scalar m_k_min; //!< Minimum k of tabulated convolution kernel
Scalar m_k_max; //!< Maximum k of tabulated convolution kernel
Scalar m_delta_k; //!< Spacing between k values
GPUArray<Scalar> m_table; //!< Tabulated kernel
GPUArray<Scalar> m_table_d; //!< Tabulated kernel
GlobalArray<Scalar> m_table; //!< Tabulated kernel
GlobalArray<Scalar> m_table_d; //!< Tabulated kernel
bool m_use_table; //!< Whether to use the tabulated kernel

//! Helper function to be called when box changes
Expand Down Expand Up @@ -161,10 +161,10 @@ class OrderParameterMesh : public CollectiveVariable

bool m_kiss_fft_initialized; //!< True if a local KISS FFT has been set up

GPUArray<kiss_fft_cpx> m_mesh; //!< The particle density mesh
GPUArray<kiss_fft_cpx> m_fourier_mesh; //!< The fourier transformed mesh
GPUArray<kiss_fft_cpx> m_fourier_mesh_G; //!< Fourier transformed mesh times the influence function
GPUArray<kiss_fft_cpx> m_inv_fourier_mesh; //!< The inverse-Fourier transformed mesh
GlobalArray<kiss_fft_cpx> m_mesh; //!< The particle density mesh
GlobalArray<kiss_fft_cpx> m_fourier_mesh; //!< The fourier transformed mesh
GlobalArray<kiss_fft_cpx> m_fourier_mesh_G; //!< Fourier transformed mesh times the influence function
GlobalArray<kiss_fft_cpx> m_inv_fourier_mesh; //!< The inverse-Fourier transformed mesh

std::vector<std::string> m_log_names; //!< Name of the log quantity

Expand Down
24 changes: 12 additions & 12 deletions metadynamics/OrderParameterMeshGPU.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -25,16 +25,16 @@ OrderParameterMeshGPU::OrderParameterMeshGPU(std::shared_ptr<SystemDefinition> s
m_gpu_q_max(m_exec_conf)
{
unsigned int n_blocks = m_mesh_points.x*m_mesh_points.y*m_mesh_points.z/m_block_size+1;
GPUArray<Scalar> sum_partial(n_blocks,m_exec_conf);
GlobalArray<Scalar> sum_partial(n_blocks,m_exec_conf);
m_sum_partial.swap(sum_partial);

GPUArray<Scalar> sum_virial_partial(6*n_blocks,m_exec_conf);
GlobalArray<Scalar> sum_virial_partial(6*n_blocks,m_exec_conf);
m_sum_virial_partial.swap(sum_virial_partial);

GPUArray<Scalar> sum_virial(6,m_exec_conf);
GlobalArray<Scalar> sum_virial(6,m_exec_conf);
m_sum_virial.swap(sum_virial);

GPUArray<Scalar4> max_partial(n_blocks, m_exec_conf);
GlobalArray<Scalar4> max_partial(n_blocks, m_exec_conf);
m_max_partial.swap(max_partial);

// initial value of number of particles per bin
Expand Down Expand Up @@ -123,22 +123,22 @@ void OrderParameterMeshGPU::initializeFFT()
m_scratch_idx = Index2D(n_particle_bins,(2*m_radius+1)*(2*m_radius+1)*(2*m_radius+1));

// allocate mesh and transformed mesh
GPUArray<cufftComplex> mesh(m_n_cells,m_exec_conf);
GlobalArray<cufftComplex> mesh(m_n_cells,m_exec_conf);
m_mesh.swap(mesh);

GPUArray<cufftComplex> fourier_mesh(m_n_inner_cells, m_exec_conf);
GlobalArray<cufftComplex> fourier_mesh(m_n_inner_cells, m_exec_conf);
m_fourier_mesh.swap(fourier_mesh);

GPUArray<cufftComplex> fourier_mesh_G(m_n_inner_cells, m_exec_conf);
GlobalArray<cufftComplex> fourier_mesh_G(m_n_inner_cells, m_exec_conf);
m_fourier_mesh_G.swap(fourier_mesh_G);

GPUArray<cufftComplex> inv_fourier_mesh(m_n_cells, m_exec_conf);
GlobalArray<cufftComplex> inv_fourier_mesh(m_n_cells, m_exec_conf);
m_inv_fourier_mesh.swap(inv_fourier_mesh);

GPUArray<Scalar4> particle_bins(m_bin_idx.getNumElements(), m_exec_conf);
GlobalArray<Scalar4> particle_bins(m_bin_idx.getNumElements(), m_exec_conf);
m_particle_bins.swap(particle_bins);

GPUArray<unsigned int> n_cell(m_bin_idx.getW(), m_exec_conf);
GlobalArray<unsigned int> n_cell(m_bin_idx.getW(), m_exec_conf);
m_n_cell.swap(n_cell);

GPUFlags<unsigned int> cell_overflowed(m_exec_conf);
Expand All @@ -147,7 +147,7 @@ void OrderParameterMeshGPU::initializeFFT()
m_cell_overflowed.resetFlags(0);

// allocate scratch space for density reduction
GPUArray<Scalar> mesh_scratch(m_scratch_idx.getNumElements(), m_exec_conf);
GlobalArray<Scalar> mesh_scratch(m_scratch_idx.getNumElements(), m_exec_conf);
m_mesh_scratch.swap(mesh_scratch);
}

Expand Down Expand Up @@ -195,7 +195,7 @@ void OrderParameterMeshGPU::assignParticles()
m_cell_size = flags;

m_bin_idx = Index2D(m_bin_idx.getW(),m_cell_size);
GPUArray<Scalar4> particle_bins(m_bin_idx.getNumElements(),m_exec_conf);
GlobalArray<Scalar4> particle_bins(m_bin_idx.getNumElements(),m_exec_conf);
m_particle_bins.swap(particle_bins);
m_cell_overflowed.resetFlags(0);
}
Expand Down
2 changes: 1 addition & 1 deletion metadynamics/OrderParameterMeshGPU.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1344,7 +1344,7 @@ struct get_type : thrust::unary_function<Scalar4, unsigned int>
Scalar gpu_compute_mode_sq(unsigned int N,
const Scalar4 *d_postype,
const Scalar *d_mode,
const CachedAllocator& alloc)
CachedAllocator& alloc)
{
thrust::device_ptr<const Scalar> mode(d_mode);
thrust::device_ptr<const Scalar4> postype(d_postype);
Expand Down
2 changes: 1 addition & 1 deletion metadynamics/OrderParameterMeshGPU.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -106,4 +106,4 @@ void gpu_compute_q_max(unsigned int n_wave_vectors,
Scalar gpu_compute_mode_sq(unsigned int N,
const Scalar4 *d_postype,
const Scalar *d_mode,
const CachedAllocator& alloc);
CachedAllocator& alloc);
22 changes: 11 additions & 11 deletions metadynamics/OrderParameterMeshGPU.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -68,27 +68,27 @@ class OrderParameterMeshGPU : public OrderParameterMesh
dfft_plan m_dfft_plan_inverse; //!< Forward distributed FFT
#endif

GPUArray<cufftComplex> m_mesh; //!< The particle density mesh
GPUArray<cufftComplex> m_fourier_mesh; //!< The fourier transformed mesh
GPUArray<cufftComplex> m_fourier_mesh_G; //!< Fourier transformed mesh times the influence function
GPUArray<cufftComplex> m_inv_fourier_mesh; //!< The inverse-fourier transformed force mesh
GlobalArray<cufftComplex> m_mesh; //!< The particle density mesh
GlobalArray<cufftComplex> m_fourier_mesh; //!< The fourier transformed mesh
GlobalArray<cufftComplex> m_fourier_mesh_G; //!< Fourier transformed mesh times the influence function
GlobalArray<cufftComplex> m_inv_fourier_mesh; //!< The inverse-fourier transformed force mesh

Index2D m_bin_idx; //!< Total number of bins
GPUArray<Scalar4> m_particle_bins; //!< Cell list for particle positions and modes
GPUArray<Scalar> m_mesh_scratch; //!< Mesh with scratch space for density reduction
GlobalArray<Scalar4> m_particle_bins; //!< Cell list for particle positions and modes
GlobalArray<Scalar> m_mesh_scratch; //!< Mesh with scratch space for density reduction
Index2D m_scratch_idx; //!< Indexer for scratch space
GPUArray<unsigned int> m_n_cell; //!< Number of particles per cell
GlobalArray<unsigned int> m_n_cell; //!< Number of particles per cell
unsigned int m_cell_size; //!< Current max. number of particles per cell
GPUFlags<unsigned int> m_cell_overflowed; //!< Flag set to 1 if a cell overflows

GPUFlags<Scalar> m_sum; //!< Sum over fourier mesh values
GPUArray<Scalar> m_sum_partial; //!< Partial sums over fourier mesh values
GPUArray<Scalar> m_sum_virial_partial; //!< Partial sums over virial mesh values
GPUArray<Scalar> m_sum_virial; //!< Final sum over virial mesh values
GlobalArray<Scalar> m_sum_partial; //!< Partial sums over fourier mesh values
GlobalArray<Scalar> m_sum_virial_partial; //!< Partial sums over virial mesh values
GlobalArray<Scalar> m_sum_virial; //!< Final sum over virial mesh values
unsigned int m_block_size; //!< Block size for fourier mesh reduction

GPUFlags<Scalar4> m_gpu_q_max; //!< Return value for maximum Fourier mode reduction
GPUArray<Scalar4> m_max_partial; //!< Scratch space for reduction of maximum Fourier amplitude
GlobalArray<Scalar4> m_max_partial; //!< Scratch space for reduction of maximum Fourier amplitude
};

//! Define plus operator for complex data type (only need to compile by CommunicatorMesh base class)
Expand Down

0 comments on commit b6802f2

Please sign in to comment.