Skip to content
GitLab
Commit 97764d64 authored by lvs's avatar lvs
Browse files

solved #30

parent 92f41f06
Expand all Hide whitespace changes
Inline Side-by-side
common/defines.h
View file @ 97764d64
......@@ -39,6 +39,7 @@ Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
#define NUM_OF_THREADS_PER_BLOCK 64
#endif
#define ATYPE_NUM 22
#define MAX_NUM_OF_ATOMS 256
#define MAX_NUM_OF_ATYPES 14
#define MAX_NUM_OF_ROTBONDS 32
......
device/calcenergy.cl
View file @ 97764d64
This diff is collapsed.
device/calcgradient.cl
View file @ 97764d64
......@@ -54,6 +54,9 @@ void map_priv_angle(float* angle)
}
}
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable
// Atomic operations used in gradients of intra contributors.
// Only atomic_cmpxchg() works on floats.
// So for atomic add on floats, this link was used:
......@@ -90,7 +93,6 @@ void atomicSub_g_f(volatile __local float *addr, float val)
} while( current.u32 != expected.u32 );
}
void gpu_calc_gradient(
int dockpars_rotbondlist_length,
char dockpars_num_of_atoms,
......@@ -107,6 +109,7 @@ void gpu_calc_gradient(
float dockpars_coeff_elec,
float dockpars_qasp,
float dockpars_coeff_desolv,
float dockpars_smooth,
// Some OpenCL compilers don't allow declaring
// local variables within non-kernel functions.
......@@ -120,32 +123,19 @@ void gpu_calc_gradient(
__local float* calc_coords_y,
__local float* calc_coords_z,
__constant float* atom_charges_const,
__constant char* atom_types_const,
__constant char* intraE_contributors_const,
float dockpars_smooth,
__constant float* reqm,
__constant float* reqm_hbond,
__constant uint* atom1_types_reqm,
__constant uint* atom2_types_reqm,
__constant float* VWpars_AC_const,
__constant float* VWpars_BD_const,
__constant float* dspars_S_const,
__constant float* dspars_V_const,
__constant int* rotlist_const,
__constant float* ref_coords_x_const,
__constant float* ref_coords_y_const,
__constant float* ref_coords_z_const,
__constant float* rotbonds_moving_vectors_const,
__constant float* rotbonds_unit_vectors_const,
__constant float* ref_orientation_quats_const,
__constant int* rotbonds_const,
__constant int* rotbonds_atoms_const,
__constant int* num_rotating_atoms_per_rotbond_const
__constant kernelconstant_interintra* kerconst_interintra,
__global const kernelconstant_intracontrib* kerconst_intracontrib,
__constant kernelconstant_intra* kerconst_intra,
__constant kernelconstant_rotlist* kerconst_rotlist,
__constant kernelconstant_conform* kerconst_conform
,
__constant int* rotbonds_const,
__global const int* rotbonds_atoms_const,
__constant int* num_rotating_atoms_per_rotbond_const
,
__constant float* angle_const,
__constant float* dependence_on_theta_const,
__constant float* dependence_on_rotangle_const
__global const float* angle_const,
__constant float* dependence_on_theta_const,
__constant float* dependence_on_rotangle_const
// Gradient-related arguments
// Calculate gradients (forces) for intermolecular energy
......@@ -209,7 +199,7 @@ void gpu_calc_gradient(
rotation_counter < dockpars_rotbondlist_length;
rotation_counter+=NUM_OF_THREADS_PER_BLOCK)
{
int rotation_list_element = rotlist_const[rotation_counter];
int rotation_list_element = kerconst_rotlist->rotlist_const[rotation_counter];
if ((rotation_list_element & RLIST_DUMMY_MASK) == 0) // If not dummy rotation
{
......@@ -220,9 +210,9 @@ void gpu_calc_gradient(
if ((rotation_list_element & RLIST_FIRSTROT_MASK) != 0) // If first rotation of this atom
{
atom_to_rotate[0] = ref_coords_x_const[atom_id];
atom_to_rotate[1] = ref_coords_y_const[atom_id];
atom_to_rotate[2] = ref_coords_z_const[atom_id];
atom_to_rotate[0] = kerconst_conform->ref_coords_x_const[atom_id];
atom_to_rotate[1] = kerconst_conform->ref_coords_y_const[atom_id];
atom_to_rotate[2] = kerconst_conform->ref_coords_z_const[atom_id];
}
else
{
......@@ -255,13 +245,13 @@ void gpu_calc_gradient(
{
uint rotbond_id = (rotation_list_element & RLIST_RBONDID_MASK) >> RLIST_RBONDID_SHIFT;
rotation_unitvec[0] = rotbonds_unit_vectors_const[3*rotbond_id];
rotation_unitvec[1] = rotbonds_unit_vectors_const[3*rotbond_id+1];
rotation_unitvec[2] = rotbonds_unit_vectors_const[3*rotbond_id+2];
rotation_unitvec[0] = kerconst_conform->rotbonds_unit_vectors_const[3*rotbond_id];
rotation_unitvec[1] = kerconst_conform->rotbonds_unit_vectors_const[3*rotbond_id+1];
rotation_unitvec[2] = kerconst_conform->rotbonds_unit_vectors_const[3*rotbond_id+2];
rotation_movingvec[0] = rotbonds_moving_vectors_const[3*rotbond_id];
rotation_movingvec[1] = rotbonds_moving_vectors_const[3*rotbond_id+1];
rotation_movingvec[2] = rotbonds_moving_vectors_const[3*rotbond_id+2];
rotation_movingvec[0] = kerconst_conform->rotbonds_moving_vectors_const[3*rotbond_id];
rotation_movingvec[1] = kerconst_conform->rotbonds_moving_vectors_const[3*rotbond_id+1];
rotation_movingvec[2] = kerconst_conform->rotbonds_moving_vectors_const[3*rotbond_id+2];
rotation_angle = genotype[6+rotbond_id]*DEG_TO_RAD;
......@@ -292,22 +282,22 @@ void gpu_calc_gradient(
quatrot_temp_y = quatrot_left_y;
quatrot_temp_z = quatrot_left_z;
quatrot_left_q = quatrot_temp_q*ref_orientation_quats_const[4*(*run_id)]-
quatrot_temp_x*ref_orientation_quats_const[4*(*run_id)+1]-
quatrot_temp_y*ref_orientation_quats_const[4*(*run_id)+2]-
quatrot_temp_z*ref_orientation_quats_const[4*(*run_id)+3];
quatrot_left_x = quatrot_temp_q*ref_orientation_quats_const[4*(*run_id)+1]+
ref_orientation_quats_const[4*(*run_id)]*quatrot_temp_x+
quatrot_temp_y*ref_orientation_quats_const[4*(*run_id)+3]-
ref_orientation_quats_const[4*(*run_id)+2]*quatrot_temp_z;
quatrot_left_y = quatrot_temp_q*ref_orientation_quats_const[4*(*run_id)+2]+
ref_orientation_quats_const[4*(*run_id)]*quatrot_temp_y+
ref_orientation_quats_const[4*(*run_id)+1]*quatrot_temp_z-
quatrot_temp_x*ref_orientation_quats_const[4*(*run_id)+3];
quatrot_left_z = quatrot_temp_q*ref_orientation_quats_const[4*(*run_id)+3]+
ref_orientation_quats_const[4*(*run_id)]*quatrot_temp_z+
quatrot_temp_x*ref_orientation_quats_const[4*(*run_id)+2]-
ref_orientation_quats_const[4*(*run_id)+1]*quatrot_temp_y;
quatrot_left_q = quatrot_temp_q*kerconst_conform->ref_orientation_quats_const[4*(*run_id)]-
quatrot_temp_x*kerconst_conform->ref_orientation_quats_const[4*(*run_id)+1]-
quatrot_temp_y*kerconst_conform->ref_orientation_quats_const[4*(*run_id)+2]-
quatrot_temp_z*kerconst_conform->ref_orientation_quats_const[4*(*run_id)+3];
quatrot_left_x = quatrot_temp_q*kerconst_conform->ref_orientation_quats_const[4*(*run_id)+1]+
kerconst_conform->ref_orientation_quats_const[4*(*run_id)]*quatrot_temp_x+
quatrot_temp_y*kerconst_conform->ref_orientation_quats_const[4*(*run_id)+3]-
kerconst_conform->ref_orientation_quats_const[4*(*run_id)+2]*quatrot_temp_z;
quatrot_left_y = quatrot_temp_q*kerconst_conform->ref_orientation_quats_const[4*(*run_id)+2]+
kerconst_conform->ref_orientation_quats_const[4*(*run_id)]*quatrot_temp_y+
kerconst_conform->ref_orientation_quats_const[4*(*run_id)+1]*quatrot_temp_z-
quatrot_temp_x*kerconst_conform->ref_orientation_quats_const[4*(*run_id)+3];
quatrot_left_z = quatrot_temp_q*kerconst_conform->ref_orientation_quats_const[4*(*run_id)+3]+
kerconst_conform->ref_orientation_quats_const[4*(*run_id)]*quatrot_temp_z+
quatrot_temp_x*kerconst_conform->ref_orientation_quats_const[4*(*run_id)+2]-
kerconst_conform->ref_orientation_quats_const[4*(*run_id)+1]*quatrot_temp_y;
}
quatrot_temp_q = 0 -
......@@ -358,11 +348,11 @@ void gpu_calc_gradient(
atom_id < dockpars_num_of_atoms;
atom_id+= NUM_OF_THREADS_PER_BLOCK)
{
uint atom_typeid = atom_types_const[atom_id];
uint atom_typeid = kerconst_interintra->atom_types_const[atom_id];
float x = calc_coords_x[atom_id];
float y = calc_coords_y[atom_id];
float z = calc_coords_z[atom_id];
float q = atom_charges_const[atom_id];
float q = kerconst_interintra->atom_charges_const[atom_id];
if ((x < 0) || (y < 0) || (z < 0) || (x >= dockpars_gridsize_x-1)
|| (y >= dockpars_gridsize_y-1)
......@@ -615,8 +605,8 @@ void gpu_calc_gradient(
float priv_gradient_per_intracontributor= 0.0f;
// Getting atom IDs
uint atom1_id = intraE_contributors_const[3*contributor_counter];
uint atom2_id = intraE_contributors_const[3*contributor_counter+1];
uint atom1_id = kerconst_intracontrib->intraE_contributors_const[3*contributor_counter];
uint atom2_id = kerconst_intracontrib->intraE_contributors_const[3*contributor_counter+1];
/*
printf ("%-5u %-5u %-5u\n", contributor_counter, atom1_id, atom2_id);
......@@ -633,11 +623,11 @@ void gpu_calc_gradient(
float atomic_distance = dist*dockpars_grid_spacing;
// Getting type IDs
uint atom1_typeid = atom_types_const[atom1_id];
uint atom2_typeid = atom_types_const[atom2_id];
uint atom1_typeid = kerconst_interintra->atom_types_const[atom1_id];
uint atom2_typeid = kerconst_interintra->atom_types_const[atom2_id];
uint atom1_type_vdw_hb = atom1_types_reqm [atom1_typeid];
uint atom2_type_vdw_hb = atom2_types_reqm [atom2_typeid];
uint atom1_type_vdw_hb = kerconst_intra->atom1_types_reqm_const [atom1_typeid];
uint atom2_type_vdw_hb = kerconst_intra->atom2_types_reqm_const [atom2_typeid];
//printf ("%-5u %-5u %-5u\n", contributor_counter, atom1_id, atom2_id);
// Getting optimum pair distance (opt_distance) from reqm and reqm_hbond
......@@ -647,13 +637,13 @@ void gpu_calc_gradient(
// (sum of the vdW radii of two like atoms (A)) in the case of hbond
float opt_distance;
if (intraE_contributors_const[3*contributor_counter+2] == 1) //H-bond
if (kerconst_intracontrib->intraE_contributors_const[3*contributor_counter+2] == 1) //H-bond
{
opt_distance = reqm_hbond [atom1_type_vdw_hb] + reqm_hbond [atom2_type_vdw_hb];
opt_distance = kerconst_intra->reqm_hbond_const [atom1_type_vdw_hb] + kerconst_intra->reqm_hbond_const [atom2_type_vdw_hb];
}
else //van der Waals
{
opt_distance = 0.5f*(reqm [atom1_type_vdw_hb] + reqm [atom2_type_vdw_hb]);
opt_distance = 0.5f*(kerconst_intra->reqm_const [atom1_type_vdw_hb] + kerconst_intra->reqm_const [atom2_type_vdw_hb]);
}
// Getting smoothed distance
......@@ -676,17 +666,17 @@ void gpu_calc_gradient(
if (atomic_distance < 8.0f)
{
// Calculating van der Waals / hydrogen bond term
priv_gradient_per_intracontributor += native_divide (-12*VWpars_AC_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],
priv_gradient_per_intracontributor += native_divide (-12*kerconst_intra->VWpars_AC_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],
native_powr(smoothed_distance/*atomic_distance*/, 13)
);
if (intraE_contributors_const[3*contributor_counter+2] == 1) { //H-bond
priv_gradient_per_intracontributor += native_divide (10*VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],
if (kerconst_intracontrib->intraE_contributors_const[3*contributor_counter+2] == 1) { //H-bond
priv_gradient_per_intracontributor += native_divide (10*kerconst_intra->VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],
native_powr(smoothed_distance/*atomic_distance*/, 11)
);
}
else { //van der Waals
priv_gradient_per_intracontributor += native_divide (6*VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],
priv_gradient_per_intracontributor += native_divide (6*kerconst_intra->VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],
native_powr(smoothed_distance/*atomic_distance*/, 7)
);
}
......@@ -702,12 +692,12 @@ void gpu_calc_gradient(
float lower = native_powr(atomic_distance, 2) * native_powr(DIEL_A * (native_exp(DIEL_B_TIMES_H*atomic_distance) + DIEL_K) + DIEL_B * native_exp(DIEL_B_TIMES_H*atomic_distance), 2);
priv_gradient_per_intracontributor += -dockpars_coeff_elec * atom_charges_const[atom1_id] * atom_charges_const[atom2_id] * native_divide (upper, lower);
priv_gradient_per_intracontributor += -dockpars_coeff_elec * kerconst_interintra->atom_charges_const[atom1_id] * kerconst_interintra->atom_charges_const[atom2_id] * native_divide (upper, lower);
// Calculating desolvation term
priv_gradient_per_intracontributor += (
(dspars_S_const[atom1_typeid] + dockpars_qasp*fabs(atom_charges_const[atom1_id])) * dspars_V_const[atom2_typeid] +
(dspars_S_const[atom2_typeid] + dockpars_qasp*fabs(atom_charges_const[atom2_id])) * dspars_V_const[atom1_typeid]
(kerconst_intra->dspars_S_const[atom1_typeid] + dockpars_qasp*fabs(kerconst_interintra->atom_charges_const[atom1_id])) * kerconst_intra->dspars_V_const[atom2_typeid] +
(kerconst_intra->dspars_S_const[atom2_typeid] + dockpars_qasp*fabs(kerconst_interintra->atom_charges_const[atom2_id])) * kerconst_intra->dspars_V_const[atom1_typeid]
) *
dockpars_coeff_desolv * /*-0.07716049382716049*/ -0.077160f * atomic_distance * native_exp(/*-0.038580246913580245*/ -0.038580f *native_powr(atomic_distance, 2));
} // if cuttoff2 - internuclear-distance at 20.48A
......@@ -749,8 +739,8 @@ void gpu_calc_gradient(
contributor_counter ++) {
// Getting atom IDs
uint atom1_id = intraE_contributors_const[3*contributor_counter];
uint atom2_id = intraE_contributors_const[3*contributor_counter+1];
uint atom1_id = kerconst_intracontrib->intraE_contributors_const[3*contributor_counter];
uint atom2_id = kerconst_intracontrib->intraE_contributors_const[3*contributor_counter+1];
// Calculating xyz distances in Angstroms of vector
// that goes from "atom1_id"-to-"atom2_id"
......@@ -1254,9 +1244,9 @@ void gpu_calc_gradient(
float3 rotation_unitvec;
/*
rotation_unitvec.x = rotbonds_unit_vectors_const[3*rotbond_id];
rotation_unitvec.y = rotbonds_unit_vectors_const[3*rotbond_id+1];
rotation_unitvec.z = rotbonds_unit_vectors_const[3*rotbond_id+2];
rotation_unitvec.x = kerconst_conform->rotbonds_unit_vectors_const[3*rotbond_id];
rotation_unitvec.y = kerconst_conform->rotbonds_unit_vectors_const[3*rotbond_id+1];
rotation_unitvec.z = kerconst_conform->rotbonds_unit_vectors_const[3*rotbond_id+2];
*/
rotation_unitvec.x = calc_coords_x[atom2_id] - calc_coords_x[atom1_id];
rotation_unitvec.y = calc_coords_y[atom2_id] - calc_coords_y[atom1_id];
......
device/kernel1.cl
View file @ 97764d64
......@@ -44,25 +44,13 @@ gpu_calc_initpop(
__global int* restrict dockpars_evals_of_new_entities,
int dockpars_pop_size,
float dockpars_qasp,
__constant float* atom_charges_const,
__constant char* atom_types_const,
__constant char* intraE_contributors_const,
float dockpars_smooth,
__constant float* reqm,
__constant float* reqm_hbond,
__constant uint* atom1_types_reqm,
__constant uint* atom2_types_reqm,
__constant float* VWpars_AC_const,
__constant float* VWpars_BD_const,
__constant float* dspars_S_const,
__constant float* dspars_V_const,
__constant int* rotlist_const,
__constant float* ref_coords_x_const,
__constant float* ref_coords_y_const,
__constant float* ref_coords_z_const,
__constant float* rotbonds_moving_vectors_const,
__constant float* rotbonds_unit_vectors_const,
__constant float* ref_orientation_quats_const
float dockpars_smooth,
__constant kernelconstant_interintra* kerconst_interintra,
__global const kernelconstant_intracontrib* kerconst_intracontrib,
__constant kernelconstant_intra* kerconst_intra,
__constant kernelconstant_rotlist* kerconst_rotlist,
__constant kernelconstant_conform* kerconst_conform
){
// Some OpenCL compilers don't allow declaring
// local variables within non-kernel functions.
......@@ -113,6 +101,8 @@ gpu_calc_initpop(
dockpars_coeff_elec,
dockpars_qasp,
dockpars_coeff_desolv,
dockpars_smooth,
genotype,
&energy,
&run_id,
......@@ -128,29 +118,14 @@ gpu_calc_initpop(
partial_interE,
partial_intraE,
#endif
atom_charges_const,
atom_types_const,
intraE_contributors_const,
#if 0
false,
#endif
dockpars_smooth,
reqm,
reqm_hbond,
atom1_types_reqm,
atom2_types_reqm,
VWpars_AC_const,
VWpars_BD_const,
dspars_S_const,
dspars_V_const,
rotlist_const,
ref_coords_x_const,
ref_coords_y_const,
ref_coords_z_const,
rotbonds_moving_vectors_const,
rotbonds_unit_vectors_const,
ref_orientation_quats_const
kerconst_interintra,
kerconst_intracontrib,
kerconst_intra,
kerconst_rotlist,
kerconst_conform
);
// =============================================================
......
device/kernel3.cl
View file @ 97764d64
......@@ -52,25 +52,13 @@ perform_LS(
uint dockpars_cons_limit,
uint dockpars_max_num_of_iters,
float dockpars_qasp,
__constant float* atom_charges_const,
__constant char* atom_types_const,
__constant char* intraE_contributors_const,
float dockpars_smooth,
__constant float* reqm,
__constant float* reqm_hbond,
__constant uint* atom1_types_reqm,
__constant uint* atom2_types_reqm,
__constant float* VWpars_AC_const,
__constant float* VWpars_BD_const,
__constant float* dspars_S_const,
__constant float* dspars_V_const,
__constant int* rotlist_const,
__constant float* ref_coords_x_const,
__constant float* ref_coords_y_const,
__constant float* ref_coords_z_const,
__constant float* rotbonds_moving_vectors_const,
__constant float* rotbonds_unit_vectors_const,
__constant float* ref_orientation_quats_const
float dockpars_smooth,
__constant kernelconstant_interintra* kerconst_interintra,
__global const kernelconstant_intracontrib* kerconst_intracontrib,
__constant kernelconstant_intra* kerconst_intra,
__constant kernelconstant_rotlist* kerconst_rotlist,
__constant kernelconstant_conform* kerconst_conform
)
//The GPU global function performs local search on the pre-defined entities of conformations_next.
//The number of blocks which should be started equals to num_of_lsentities*num_of_runs.
......@@ -203,6 +191,8 @@ perform_LS(
dockpars_coeff_elec,
dockpars_qasp,
dockpars_coeff_desolv,
dockpars_smooth,
genotype_candidate,
&candidate_energy,
&run_id,
......@@ -218,29 +208,14 @@ perform_LS(
partial_interE,
partial_intraE,
#endif
atom_charges_const,
atom_types_const,
intraE_contributors_const,
#if 0
false,
#endif
dockpars_smooth,
reqm,
reqm_hbond,
atom1_types_reqm,
atom2_types_reqm,
VWpars_AC_const,
VWpars_BD_const,
dspars_S_const,
dspars_V_const,
rotlist_const,
ref_coords_x_const,
ref_coords_y_const,
ref_coords_z_const,
rotbonds_moving_vectors_const,
rotbonds_unit_vectors_const,
ref_orientation_quats_const
kerconst_interintra,
kerconst_intracontrib,
kerconst_intra,
kerconst_rotlist,
kerconst_conform
);
// =================================================================
......@@ -308,6 +283,8 @@ perform_LS(
dockpars_coeff_elec,
dockpars_qasp,
dockpars_coeff_desolv,
dockpars_smooth,
genotype_candidate,
&candidate_energy,
&run_id,
......@@ -323,29 +300,14 @@ perform_LS(
partial_interE,
partial_intraE,
#endif
atom_charges_const,
atom_types_const,
intraE_contributors_const,
#if 0
false,
#endif
dockpars_smooth,
reqm,
reqm_hbond,
atom1_types_reqm,
atom2_types_reqm,
VWpars_AC_const,
VWpars_BD_const,
dspars_S_const,
dspars_V_const,
rotlist_const,
ref_coords_x_const,
ref_coords_y_const,
ref_coords_z_const,
rotbonds_moving_vectors_const,
rotbonds_unit_vectors_const,
ref_orientation_quats_const
kerconst_interintra,
kerconst_intracontrib,
kerconst_intra,
kerconst_rotlist,
kerconst_conform
);
// =================================================================
......
device/kernel4.cl
View file @ 97764d64
......@@ -53,25 +53,13 @@ gpu_gen_and_eval_newpops(
float dockpars_abs_max_dmov,
float dockpars_abs_max_dang,
float dockpars_qasp,
__constant float* atom_charges_const,
__constant char* atom_types_const,
__constant char* intraE_contributors_const,
float dockpars_smooth,
__constant float* reqm,
__constant float* reqm_hbond,
__constant uint* atom1_types_reqm,
__constant uint* atom2_types_reqm,
__constant float* VWpars_AC_const,
__constant float* VWpars_BD_const,
__constant float* dspars_S_const,
__constant float* dspars_V_const,
__constant int* rotlist_const,
__constant float* ref_coords_x_const,
__constant float* ref_coords_y_const,
__constant float* ref_coords_z_const,
__constant float* rotbonds_moving_vectors_const,
__constant float* rotbonds_unit_vectors_const,
__constant float* ref_orientation_quats_const
float dockpars_smooth,
__constant kernelconstant_interintra* kerconst_interintra,
__global const kernelconstant_intracontrib* kerconst_intracontrib,
__constant kernelconstant_intra* kerconst_intra,
__constant kernelconstant_rotlist* kerconst_rotlist,
__constant kernelconstant_conform* kerconst_conform
)
//The GPU global function
{
......@@ -265,6 +253,8 @@ gpu_gen_and_eval_newpops(
dockpars_coeff_elec,
dockpars_qasp,
dockpars_coeff_desolv,
dockpars_smooth,
offspring_genotype,
&energy,
&run_id,
......@@ -280,29 +270,14 @@ gpu_gen_and_eval_newpops(
partial_interE,
partial_intraE,
#endif
atom_charges_const,
atom_types_const,
intraE_contributors_const,
#if 0
false,
#endif
dockpars_smooth,
reqm,
reqm_hbond,
atom1_types_reqm,
atom2_types_reqm,
VWpars_AC_const,
VWpars_BD_const,
dspars_S_const,
dspars_V_const,
rotlist_const,
ref_coords_x_const,
ref_coords_y_const,
ref_coords_z_const,
rotbonds_moving_vectors_const,
rotbonds_unit_vectors_const,
ref_orientation_quats_const
kerconst_interintra,
kerconst_intracontrib,
kerconst_intra,
kerconst_rotlist,
kerconst_conform
);
// =============================================================
......
device/kernel_fire.cl
View file @ 97764d64
......@@ -16,56 +16,45 @@
__kernel void __attribute__ ((reqd_work_group_size(NUM_OF_THREADS_PER_BLOCK,1,1)))
gradient_minFire(
char dockpars_num_of_atoms,
char dockpars_num_of_atypes,
int dockpars_num_of_intraE_contributors,
char dockpars_gridsize_x,
char dockpars_gridsize_y,
char dockpars_gridsize_z,
char dockpars_num_of_atoms,
char dockpars_num_of_atypes,
int dockpars_num_of_intraE_contributors,
char dockpars_gridsize_x,
char dockpars_gridsize_y,
char dockpars_gridsize_z,
// g1 = gridsize_x
uint dockpars_gridsize_x_times_y, // g2 = gridsize_x * gridsize_y
uint dockpars_gridsize_x_times_y_times_z, // g3 = gridsize_x * gridsize_y * gridsize_z
float dockpars_grid_spacing,
__global const float* restrict dockpars_fgrids, // This is too large to be allocated in __constant
int dockpars_rotbondlist_length,
float dockpars_coeff_elec,
float dockpars_coeff_desolv,
__global float* restrict dockpars_conformations_next,
__global float* restrict dockpars_energies_next,
__global int* restrict dockpars_evals_of_new_entities,
__global uint* restrict dockpars_prng_states,
int dockpars_pop_size,
int dockpars_num_of_genes,
float dockpars_lsearch_rate,
uint dockpars_num_of_lsentities,
uint dockpars_max_num_of_iters,
float dockpars_qasp,
__constant float* atom_charges_const,
__constant char* atom_types_const,
__constant char* intraE_contributors_const,
float dockpars_smooth,
__constant float* reqm,
__constant float* reqm_hbond,
__constant uint* atom1_types_reqm,
__constant uint* atom2_types_reqm,
__constant float* VWpars_AC_const,
__constant float* VWpars_BD_const,
__constant float* dspars_S_const,
__constant float* dspars_V_const,
__constant int* rotlist_const,
__constant float* ref_coords_x_const,
__constant float* ref_coords_y_const,
__constant float* ref_coords_z_const,
__constant float* rotbonds_moving_vectors_const,
__constant float* rotbonds_unit_vectors_const,
__constant float* ref_orientation_quats_const,
__constant int* rotbonds_const,
__constant int* rotbonds_atoms_const,
__constant int* num_rotating_atoms_per_rotbond_const
uint dockpars_gridsize_x_times_y, // g2 = gridsize_x * gridsize_y
uint dockpars_gridsize_x_times_y_times_z, // g3 = gridsize_x * gridsize_y * gridsize_z
float dockpars_grid_spacing,
__global const float* restrict dockpars_fgrids, // This is too large to be allocated in __constant
int dockpars_rotbondlist_length,
float dockpars_coeff_elec,
float dockpars_coeff_desolv,
__global float* restrict dockpars_conformations_next,
__global float* restrict dockpars_energies_next,
__global int* restrict dockpars_evals_of_new_entities,
__global uint* restrict dockpars_prng_states,
int dockpars_pop_size,
int dockpars_num_of_genes,
float dockpars_lsearch_rate,
uint dockpars_num_of_lsentities,
uint dockpars_max_num_of_iters,
float dockpars_qasp,
float dockpars_smooth,
__constant kernelconstant_interintra* kerconst_interintra,
__global const kernelconstant_intracontrib* kerconst_intracontrib,
__constant kernelconstant_intra* kerconst_intra,
__constant kernelconstant_rotlist* kerconst_rotlist,