Commit 43feab9e authored by Leonardo Solis's avatar Leonardo Solis
Browse files

finalized version that exclude out-of-bound genomes in kergrad

parent 2a05277c
......@@ -201,16 +201,16 @@ odock: check-env-all stringify $(SRC)
# 7cpa: for testing gradients of torsion genes (15 torsions)
# 3tmn: for testing gradients of torsion genes (1 torsion)
PDB := 1stp
PDB := 3ce3
NRUN := 100
POPSIZE := 150
POPSIZE := 500
TESTNAME:= test
test: odock
$(BIN_DIR)/$(TARGET) -ffile ./input/$(PDB)/derived/$(PDB)_protein.maps.fld -lfile ./input/$(PDB)/derived/$(PDB)_ligand.pdbqt -nrun $(NRUN) -psize $(POPSIZE) -resnam $(TESTNAME) -gfpop 1
ASTEX_PDB := 2bsm
ASTEX_NRUN:= 100
ASTEX_NRUN:= 50
ASTEX_POPSIZE := 500
ASTEX_TESTNAME := test_astex
......
......@@ -65,12 +65,35 @@ Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
#define DIEL_LAMBDA 0.003627f
#define DIEL_H DIEL_LAMBDA
#define DIEL_K 7.7839f
#define DIEL_B_TIMES_H (DIEL_B*DIEL_H)
#define DIEL_B_TIMES_H_TIMES_K (DIEL_B_TIMES_H*DIEL_K)
#define DIEL_B_TIMES_H (DIEL_B * DIEL_H)
#define DIEL_B_TIMES_H_TIMES_K (DIEL_B_TIMES_H * DIEL_K)
// Used for Shoemake to quternion transformation
// Used for Shoemake to quaternion transformation
#define PI_TIMES_2 (float)(2.0f*M_PI)
// -------------------------------------------
// Gradient-related defines
// -------------------------------------------
#define INFINITESIMAL_RADIAN 1E-5
#define HALF_INFINITESIMAL_RADIAN (0.5f * INFINITESIMAL_RADIAN)
#define INV_INFINITESIMAL_RADIAN (1/INFINITESIMAL_RADIAN)
#define COS_HALF_INFINITESIMAL_RADIAN cos(HALF_INFINITESIMAL_RADIAN)
#define SIN_HALF_INFINITESIMAL_RADIAN sin(HALF_INFINITESIMAL_RADIAN)
// Enable printf statements for debugging
// the gradient-based minimizer
//#define DEBUG_MINIMIZER
#define TRANGENE_ALPHA 1E-3
#define ROTAGENE_ALPHA 1E-8
#define TORSGENE_ALPHA 1E-13
#endif /* CALCENERGY_BASIC_H_ */
......@@ -55,9 +55,17 @@ Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
#define LS_CONT_FACTOR 0.5f
// Improvements over Pechan's implementation
//#define NATIVE_PRECISION
#define MAPPED_COPY
// TODO: convert this into a program arg
#define GRADIENT_ENABLED
#endif /* DEFINES_H_ */
......@@ -35,6 +35,9 @@ void gpu_calc_energy(
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
__global const float* restrict dockpars_fgrids, // This is too large to be allocated in __constant
char dockpars_num_of_atypes,
int dockpars_num_of_intraE_contributors,
......@@ -90,9 +93,8 @@ void gpu_calc_energy(
#endif
uchar g1 = dockpars_gridsize_x;
uint g2 = dockpars_gridsize_x * dockpars_gridsize_y;
uint g3 = dockpars_gridsize_x * dockpars_gridsize_y * dockpars_gridsize_z;
uint g2 = dockpars_gridsize_x_times_y /*dockpars_gridsize_x * dockpars_gridsize_y*/;
uint g3 = dockpars_gridsize_x_times_y_times_z /*dockpars_gridsize_x * dockpars_gridsize_y * dockpars_gridsize_z*/;
// ================================================
// CALCULATING ATOMIC POSITIONS AFTER ROTATIONS
......@@ -462,18 +464,19 @@ void gpu_calc_energy(
#endif
// Calculating desolvation term
// 1/25.92 = 0.038580246913580245
partial_energies[get_local_id(0)] += ((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]) *
dockpars_coeff_desolv*native_exp(-atomic_distance*native_divide(atomic_distance,25.92f));
dockpars_coeff_desolv*native_exp(-0.03858025f*native_powr(atomic_distance, 2));
#if defined (DEBUG_ENERGY_KERNEL1) || defined (DEBUG_ENERGY_KERNEL4) || defined (DEBUG_ENERGY_KERNEL3)
partial_intraE[get_local_id(0)] += ((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]) *
dockpars_coeff_desolv*native_exp(-atomic_distance*native_divide(atomic_distance,25.92f));
dockpars_coeff_desolv*native_exp(-0.03858025f*native_powr(atomic_distance, 2));
#endif
......
......@@ -46,6 +46,9 @@ void gpu_calc_gradient(
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
__global const float* restrict dockpars_fgrids, // This is too large to be allocated in __constant
char dockpars_num_of_atypes,
int dockpars_num_of_intraE_contributors,
......@@ -133,13 +136,11 @@ void gpu_calc_gradient(
gene_cnt+= NUM_OF_THREADS_PER_BLOCK) {
gradient_genotype[gene_cnt] = 0.0f;
}
barrier(CLK_LOCAL_MEM_FENCE);
uchar g1 = dockpars_gridsize_x;
uint g2 = dockpars_gridsize_x * dockpars_gridsize_y;
uint g3 = dockpars_gridsize_x * dockpars_gridsize_y * dockpars_gridsize_z;
uint g2 = dockpars_gridsize_x_times_y /*dockpars_gridsize_x * dockpars_gridsize_y*/;
uint g3 = dockpars_gridsize_x_times_y_times_z /*dockpars_gridsize_x * dockpars_gridsize_y * dockpars_gridsize_z*/;
// ================================================
// CALCULATING ATOMIC POSITIONS AFTER ROTATIONS
......@@ -781,16 +782,25 @@ void gpu_calc_gradient(
printf("%-20s %-10.5f\n", "torque length: ", torque_length);
#endif
/*
// Infinitesimal rotation in radians
const float infinitesimal_radian = 1E-5;
*/
// Finding the quaternion that performs
// the infinitesimal rotation around torque axis
float4 quat_torque;
quat_torque.w = native_cos(infinitesimal_radian*0.5f);
quat_torque.x = fast_normalize(torque_rot).x * native_sin(infinitesimal_radian*0.5f);
quat_torque.y = fast_normalize(torque_rot).y * native_sin(infinitesimal_radian*0.5f);
quat_torque.z = fast_normalize(torque_rot).z * native_sin(infinitesimal_radian*0.5f);
#if 0
quat_torque.w = native_cos(HALF_INFINITESIMAL_RADIAN/*infinitesimal_radian*0.5f*/);
quat_torque.x = fast_normalize(torque_rot).x * native_sin(HALF_INFINITESIMAL_RADIAN/*infinitesimal_radian*0.5f*/);
quat_torque.y = fast_normalize(torque_rot).y * native_sin(HALF_INFINITESIMAL_RADIAN/*infinitesimal_radian*0.5f*/);
quat_torque.z = fast_normalize(torque_rot).z * native_sin(HALF_INFINITESIMAL_RADIAN/*infinitesimal_radian*0.5f*/);
#endif
quat_torque.w = COS_HALF_INFINITESIMAL_RADIAN;
quat_torque.x = fast_normalize(torque_rot).x * SIN_HALF_INFINITESIMAL_RADIAN;
quat_torque.y = fast_normalize(torque_rot).y * SIN_HALF_INFINITESIMAL_RADIAN;
quat_torque.z = fast_normalize(torque_rot).z * SIN_HALF_INFINITESIMAL_RADIAN;
#if defined (DEBUG_GRAD_ROTATION_GENES)
printf("%-20s %-10.5f %-10.5f %-10.5f %-10.5f\n", "quat_torque (w,x,y,z): ", quat_torque.w, quat_torque.x, quat_torque.y, quat_torque.z);
......@@ -859,7 +869,8 @@ void gpu_calc_gradient(
// the displacement in shoemake space is not distorted.
// The correct amount of displacement in shoemake space is obtained
// by multiplying the infinitesimal displacement by shoemake_scaling:
float shoemake_scaling = torque_length / infinitesimal_radian;
//float shoemake_scaling = native_divide(torque_length, INFINITESIMAL_RADIAN/*infinitesimal_radian*/);
float shoemake_scaling = torque_length * INV_INFINITESIMAL_RADIAN;
// Derivates in cube3
// "current_u2" and "current_u3" are mapped into
......@@ -877,7 +888,7 @@ void gpu_calc_gradient(
float temp_u1 = genotype[3];
if ((0.0f < temp_u1) && (temp_u1 < 1.0f)){
grad_u1 *= ((1.0f/temp_u1) + (1.0f/(1.0f-temp_u1)));
grad_u1 *= (native_divide(1.0f, temp_u1) + native_divide(1.0f, (1.0f-temp_u1)));
}
grad_u2 *= 4.0f * (1.0f-temp_u1);
grad_u3 *= 4.0f * temp_u1;
......@@ -988,7 +999,7 @@ void gpu_calc_gradient(
float torque_on_axis = dot(rotation_unitvec, torque_tor);
// Assignment of gene-based gradient
gradient_genotype[rotbond_id+6] = torque_on_axis * (M_PI / 180.0f);
gradient_genotype[rotbond_id+6] = torque_on_axis * DEG_TO_RAD /*(M_PI / 180.0f)*/;
#if defined (DEBUG_GRAD_TORSION_GENES)
printf("gradient_torsion [%u] :%f\n", rotbond_id+6, gradient_genotype [rotbond_id+6]);
......
......@@ -28,6 +28,9 @@ gpu_calc_initpop(
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,
......@@ -89,6 +92,9 @@ gpu_calc_initpop(
dockpars_gridsize_x,
dockpars_gridsize_y,
dockpars_gridsize_z,
// g1 = gridsize_x
dockpars_gridsize_x_times_y, // g2 = gridsize_x * gridsize_y
dockpars_gridsize_x_times_y_times_z, // g3 = gridsize_x * gridsize_y * gridsize_z
dockpars_fgrids,
dockpars_num_of_atypes,
dockpars_num_of_intraE_contributors,
......
......@@ -30,6 +30,9 @@ perform_LS(
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,
......@@ -191,6 +194,9 @@ perform_LS(
dockpars_gridsize_x,
dockpars_gridsize_y,
dockpars_gridsize_z,
// g1 = gridsize_x
dockpars_gridsize_x_times_y, // g2 = gridsize_x * gridsize_y
dockpars_gridsize_x_times_y_times_z, // g3 = gridsize_x * gridsize_y * gridsize_z
dockpars_fgrids,
dockpars_num_of_atypes,
dockpars_num_of_intraE_contributors,
......@@ -293,6 +299,9 @@ perform_LS(
dockpars_gridsize_x,
dockpars_gridsize_y,
dockpars_gridsize_z,
// g1 = gridsize_x
dockpars_gridsize_x_times_y, // g2 = gridsize_x * gridsize_y
dockpars_gridsize_x_times_y_times_z, // g3 = gridsize_x * gridsize_y * gridsize_z
dockpars_fgrids,
dockpars_num_of_atypes,
dockpars_num_of_intraE_contributors,
......
......@@ -30,6 +30,9 @@ gpu_gen_and_eval_newpops(
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,
......@@ -241,6 +244,9 @@ gpu_gen_and_eval_newpops(
dockpars_gridsize_x,
dockpars_gridsize_y,
dockpars_gridsize_z,
// g1 = gridsize_x
dockpars_gridsize_x_times_y, // g2 = gridsize_x * gridsize_y
dockpars_gridsize_x_times_y_times_z, // g3 = gridsize_x * gridsize_y * gridsize_z
dockpars_fgrids,
dockpars_num_of_atypes,
dockpars_num_of_intraE_contributors,
......
// Gradient-based steepest descent minimizer
// Alternative to Solis-Wetts
//#define DEBUG_MINIMIZER
#define TRANGENE_ALPHA 1E-3//1E-8
#define ROTAGENE_ALPHA 1E-8//1E-15
#define TORSGENE_ALPHA 1E-13//1E-4
__kernel void __attribute__ ((reqd_work_group_size(NUM_OF_THREADS_PER_BLOCK,1,1)))
gradient_minimizer(
char dockpars_num_of_atoms,
......@@ -15,6 +9,9 @@ gradient_minimizer(
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,
......@@ -48,9 +45,12 @@ gradient_minimizer(
__constant int* num_rotating_atoms_per_rotbond_const,
// Specific gradient-minimizer args
uint gradMin_maxiters,
uint gradMin_maxfails,
uint gradMin_maxfails
/*
,
float gradMin_alpha,
__constant float* gradMin_conformation_min_perturbation // minimal values for gene perturbation, originally as the scalar "dxmin"
*/
)
//The GPU global function performs gradient-based minimization on (some) entities of conformations_next.
//The number of OpenCL compute units (CU) which should be started equals to num_of_minEntities*num_of_runs.
......@@ -87,6 +87,16 @@ gradient_minimizer(
///*
run_id = get_group_id(0) / dockpars_num_of_lsentities;
entity_id = get_group_id(0) % dockpars_num_of_lsentities;
// Since entity 0 is the best one due to elitism,
// it should be subjected to random selection
if (entity_id == 0) {
// If entity 0 is not selected according to LS-rate,
// choosing an other entity
if (100.0f*gpu_randf(dockpars_prng_states) > dockpars_lsearch_rate) {
entity_id = dockpars_num_of_lsentities;
}
}
//*/
energy = dockpars_energies_next[run_id*dockpars_pop_size+entity_id];
......@@ -247,6 +257,9 @@ gradient_minimizer(
dockpars_gridsize_x,
dockpars_gridsize_y,
dockpars_gridsize_z,
// g1 = gridsize_x
dockpars_gridsize_x_times_y, // g2 = gridsize_x * gridsize_y
dockpars_gridsize_x_times_y_times_z, // g3 = gridsize_x * gridsize_y * gridsize_z
dockpars_fgrids,
dockpars_num_of_atypes,
dockpars_num_of_intraE_contributors,
......@@ -370,6 +383,9 @@ gradient_minimizer(
dockpars_gridsize_x,
dockpars_gridsize_y,
dockpars_gridsize_z,
// g1 = gridsize_x
dockpars_gridsize_x_times_y, // g2 = gridsize_x * gridsize_y
dockpars_gridsize_x_times_y_times_z, // g3 = gridsize_x * gridsize_y * gridsize_z
dockpars_fgrids,
dockpars_num_of_atypes,
dockpars_num_of_intraE_contributors,
......@@ -431,7 +447,7 @@ gradient_minimizer(
genotype [i] = candidate_genotype[i];
// Up-scaling alpha by one order magnitud
alpha = alpha*/*10*/(5/(failure_cnt == 0?(failure_cnt+1):(failure_cnt)));
alpha = alpha*/*10*/(5/(failure_cnt == 0?1:(failure_cnt)));
#if defined (DEBUG_MINIMIZER)
//printf("(%-3u) %-15.15f %-10.10f %-10.10f %-10.10f\n", i, alpha, genotype[i], gradient[i], candidate_genotype[i]);
......
......@@ -59,8 +59,10 @@ Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
typedef struct {
unsigned int max_num_of_iters;
unsigned int max_num_of_consec_fails;
/*
float alpha;
float conformation_min_perturbation [ACTUAL_GENOTYPE_LENGTH];
*/
} Gradientparameters;
int docking_with_gpu(const Gridinfo* mygrid,
......
......@@ -398,6 +398,9 @@ filled with clock() */
dockpars.max_num_of_iters = (unsigned int) mypars->max_num_of_iters;
dockpars.qasp = mypars->qasp;
unsigned int g2 = dockpars.gridsize_x * dockpars.gridsize_y;
unsigned int g3 = dockpars.gridsize_x * dockpars.gridsize_y * dockpars.gridsize_z;
blocksPerGridForEachLSEntity = dockpars.num_of_lsentities*mypars->num_of_runs;
/*
......@@ -411,7 +414,9 @@ filled with clock() */
Gradientparameters gradientpars;
gradientpars.max_num_of_iters = 30;
gradientpars.max_num_of_consec_fails = (unsigned int)(0.05 * gradientpars.max_num_of_iters);
gradientpars.alpha = 0.000001f; //0.001f; // TODO: find out why 0.001f, 0.0001f (100 runs, 500 popsize) throws segmentation fault
/*
gradientpars.alpha = 0.000001f;
// Set minimum values for input conformation (translation genes as x, y, z)
for (unsigned int gene_cnt=0; gene_cnt<3; gene_cnt++) {
......@@ -433,6 +438,8 @@ filled with clock() */
mallocBufferObject(context, CL_MEM_READ_ONLY, size_conformation_min_perturbation, &mem_gradpars_conformation_min_perturbation);
memcopyBufferObjectToDevice(command_queue, mem_gradpars_conformation_min_perturbation, gradientpars.conformation_min_perturbation, size_conformation_min_perturbation);
*/
///*
// Initially, the number of entities that undergo gradient-minimization,
......@@ -489,30 +496,32 @@ filled with clock() */
setKernelArg(kernel1,3, sizeof(dockpars.gridsize_x), &dockpars.gridsize_x);
setKernelArg(kernel1,4, sizeof(dockpars.gridsize_y), &dockpars.gridsize_y);
setKernelArg(kernel1,5, sizeof(dockpars.gridsize_z), &dockpars.gridsize_z);
setKernelArg(kernel1,6, sizeof(dockpars.grid_spacing), &dockpars.grid_spacing);
setKernelArg(kernel1,7, sizeof(mem_dockpars_fgrids), &mem_dockpars_fgrids);
setKernelArg(kernel1,8, sizeof(dockpars.rotbondlist_length), &dockpars.rotbondlist_length);
setKernelArg(kernel1,9, sizeof(dockpars.coeff_elec), &dockpars.coeff_elec);
setKernelArg(kernel1,10,sizeof(dockpars.coeff_desolv), &dockpars.coeff_desolv);
setKernelArg(kernel1,11,sizeof(mem_dockpars_conformations_current), &mem_dockpars_conformations_current);
setKernelArg(kernel1,12,sizeof(mem_dockpars_energies_current), &mem_dockpars_energies_current);
setKernelArg(kernel1,13,sizeof(mem_dockpars_evals_of_new_entities), &mem_dockpars_evals_of_new_entities);
setKernelArg(kernel1,14,sizeof(dockpars.pop_size), &dockpars.pop_size);
setKernelArg(kernel1,15,sizeof(dockpars.qasp), &dockpars.qasp);
setKernelArg(kernel1,16,sizeof(mem_atom_charges_const), &mem_atom_charges_const);
setKernelArg(kernel1,17,sizeof(mem_atom_types_const), &mem_atom_types_const);
setKernelArg(kernel1,18,sizeof(mem_intraE_contributors_const), &mem_intraE_contributors_const);
setKernelArg(kernel1,19,sizeof(mem_VWpars_AC_const), &mem_VWpars_AC_const);
setKernelArg(kernel1,20,sizeof(mem_VWpars_BD_const), &mem_VWpars_BD_const);
setKernelArg(kernel1,21,sizeof(mem_dspars_S_const), &mem_dspars_S_const);
setKernelArg(kernel1,22,sizeof(mem_dspars_V_const), &mem_dspars_V_const);
setKernelArg(kernel1,23,sizeof(mem_rotlist_const), &mem_rotlist_const);
setKernelArg(kernel1,24,sizeof(mem_ref_coords_x_const), &mem_ref_coords_x_const);
setKernelArg(kernel1,25,sizeof(mem_ref_coords_y_const), &mem_ref_coords_y_const);
setKernelArg(kernel1,26,sizeof(mem_ref_coords_z_const), &mem_ref_coords_z_const);
setKernelArg(kernel1,27,sizeof(mem_rotbonds_moving_vectors_const), &mem_rotbonds_moving_vectors_const);
setKernelArg(kernel1,28,sizeof(mem_rotbonds_unit_vectors_const), &mem_rotbonds_unit_vectors_const);
setKernelArg(kernel1,29,sizeof(mem_ref_orientation_quats_const), &mem_ref_orientation_quats_const);
setKernelArg(kernel1,6, sizeof(g2), &g2);
setKernelArg(kernel1,7, sizeof(g3), &g3);
setKernelArg(kernel1,8, sizeof(dockpars.grid_spacing), &dockpars.grid_spacing);
setKernelArg(kernel1,9, sizeof(mem_dockpars_fgrids), &mem_dockpars_fgrids);
setKernelArg(kernel1,10,sizeof(dockpars.rotbondlist_length), &dockpars.rotbondlist_length);
setKernelArg(kernel1,11,sizeof(dockpars.coeff_elec), &dockpars.coeff_elec);
setKernelArg(kernel1,12,sizeof(dockpars.coeff_desolv), &dockpars.coeff_desolv);
setKernelArg(kernel1,13,sizeof(mem_dockpars_conformations_current), &mem_dockpars_conformations_current);
setKernelArg(kernel1,14,sizeof(mem_dockpars_energies_current), &mem_dockpars_energies_current);
setKernelArg(kernel1,15,sizeof(mem_dockpars_evals_of_new_entities), &mem_dockpars_evals_of_new_entities);
setKernelArg(kernel1,16,sizeof(dockpars.pop_size), &dockpars.pop_size);
setKernelArg(kernel1,17,sizeof(dockpars.qasp), &dockpars.qasp);
setKernelArg(kernel1,18,sizeof(mem_atom_charges_const), &mem_atom_charges_const);
setKernelArg(kernel1,19,sizeof(mem_atom_types_const), &mem_atom_types_const);
setKernelArg(kernel1,20,sizeof(mem_intraE_contributors_const), &mem_intraE_contributors_const);
setKernelArg(kernel1,21,sizeof(mem_VWpars_AC_const), &mem_VWpars_AC_const);
setKernelArg(kernel1,22,sizeof(mem_VWpars_BD_const), &mem_VWpars_BD_const);
setKernelArg(kernel1,23,sizeof(mem_dspars_S_const), &mem_dspars_S_const);
setKernelArg(kernel1,24,sizeof(mem_dspars_V_const), &mem_dspars_V_const);
setKernelArg(kernel1,25,sizeof(mem_rotlist_const), &mem_rotlist_const);
setKernelArg(kernel1,26,sizeof(mem_ref_coords_x_const), &mem_ref_coords_x_const);
setKernelArg(kernel1,27,sizeof(mem_ref_coords_y_const), &mem_ref_coords_y_const);
setKernelArg(kernel1,28,sizeof(mem_ref_coords_z_const), &mem_ref_coords_z_const);
setKernelArg(kernel1,29,sizeof(mem_rotbonds_moving_vectors_const), &mem_rotbonds_moving_vectors_const);
setKernelArg(kernel1,30,sizeof(mem_rotbonds_unit_vectors_const), &mem_rotbonds_unit_vectors_const);
setKernelArg(kernel1,31,sizeof(mem_ref_orientation_quats_const), &mem_ref_orientation_quats_const);
kernel1_gxsize = blocksPerGridForEachEntity * threadsPerBlock;
kernel1_lxsize = threadsPerBlock;
#ifdef DOCK_DEBUG
......@@ -539,40 +548,41 @@ filled with clock() */
setKernelArg(kernel4,3, sizeof(dockpars.gridsize_x), &dockpars.gridsize_x);
setKernelArg(kernel4,4, sizeof(dockpars.gridsize_y), &dockpars.gridsize_y);
setKernelArg(kernel4,5, sizeof(dockpars.gridsize_z), &dockpars.gridsize_z);
setKernelArg(kernel4,6, sizeof(dockpars.grid_spacing), &dockpars.grid_spacing);
setKernelArg(kernel4,7, sizeof(mem_dockpars_fgrids), &mem_dockpars_fgrids);
setKernelArg(kernel4,8, sizeof(dockpars.rotbondlist_length), &dockpars.rotbondlist_length);
setKernelArg(kernel4,9, sizeof(dockpars.coeff_elec), &dockpars.coeff_elec);
setKernelArg(kernel4,10,sizeof(dockpars.coeff_desolv), &dockpars.coeff_desolv);
setKernelArg(kernel4,11,sizeof(mem_dockpars_conformations_current), &mem_dockpars_conformations_current);
setKernelArg(kernel4,12,sizeof(mem_dockpars_energies_current), &mem_dockpars_energies_current);
setKernelArg(kernel4,13,sizeof(mem_dockpars_conformations_next), &mem_dockpars_conformations_next);
setKernelArg(kernel4,14,sizeof(mem_dockpars_energies_next), &mem_dockpars_energies_next);
setKernelArg(kernel4,15,sizeof(mem_dockpars_evals_of_new_entities), &mem_dockpars_evals_of_new_entities);
setKernelArg(kernel4,16,sizeof(mem_dockpars_prng_states), &mem_dockpars_prng_states);
setKernelArg(kernel4,17,sizeof(dockpars.pop_size), &dockpars.pop_size);
setKernelArg(kernel4,18,sizeof(dockpars.num_of_genes), &dockpars.num_of_genes);
setKernelArg(kernel4,19,sizeof(dockpars.tournament_rate), &dockpars.tournament_rate);
setKernelArg(kernel4,20,sizeof(dockpars.crossover_rate), &dockpars.crossover_rate);
setKernelArg(kernel4,21,sizeof(dockpars.mutation_rate), &dockpars.mutation_rate);
setKernelArg(kernel4,22,sizeof(dockpars.abs_max_dmov), &dockpars.abs_max_dmov);
setKernelArg(kernel4,23,sizeof(dockpars.abs_max_dang), &dockpars.abs_max_dang);
setKernelArg(kernel4,24,sizeof(dockpars.qasp), &dockpars.qasp);
setKernelArg(kernel4,25,sizeof(mem_atom_charges_const), &mem_atom_charges_const);
setKernelArg(kernel4,26,sizeof(mem_atom_types_const), &mem_atom_types_const);
setKernelArg(kernel4,27,sizeof(mem_intraE_contributors_const), &mem_intraE_contributors_const);
setKernelArg(kernel4,28,sizeof(mem_VWpars_AC_const), &mem_VWpars_AC_const);
setKernelArg(kernel4,29,sizeof(mem_VWpars_BD_const), &mem_VWpars_BD_const);
setKernelArg(kernel4,30,sizeof(mem_dspars_S_const), &mem_dspars_S_const);
setKernelArg(kernel4,31,sizeof(mem_dspars_V_const), &mem_dspars_V_const);
setKernelArg(kernel4,32,sizeof(mem_rotlist_const), &mem_rotlist_const);
setKernelArg(kernel4,33,sizeof(mem_ref_coords_x_const), &mem_ref_coords_x_const);
setKernelArg(kernel4,34,sizeof(mem_ref_coords_y_const), &mem_ref_coords_y_const);
setKernelArg(kernel4,35,sizeof(mem_ref_coords_z_const), &mem_ref_coords_z_const);
setKernelArg(kernel4,36,sizeof(mem_rotbonds_moving_vectors_const), &mem_rotbonds_moving_vectors_const);
setKernelArg(kernel4,37,sizeof(mem_rotbonds_unit_vectors_const), &mem_rotbonds_unit_vectors_const);
setKernelArg(kernel4,38,sizeof(mem_ref_orientation_quats_const), &mem_ref_orientation_quats_const);
setKernelArg(kernel4,6, sizeof(g2), &g2);
setKernelArg(kernel4,7, sizeof(g3), &g3);
setKernelArg(kernel4,8, sizeof(dockpars.grid_spacing), &dockpars.grid_spacing);
setKernelArg(kernel4,9, sizeof(mem_dockpars_fgrids), &mem_dockpars_fgrids);
setKernelArg(kernel4,10,sizeof(dockpars.rotbondlist_length), &dockpars.rotbondlist_length);
setKernelArg(kernel4,11,sizeof(dockpars.coeff_elec), &dockpars.coeff_elec);
setKernelArg(kernel4,12,sizeof(dockpars.coeff_desolv), &dockpars.coeff_desolv);
setKernelArg(kernel4,13,sizeof(mem_dockpars_conformations_current), &mem_dockpars_conformations_current);
setKernelArg(kernel4,14,sizeof(mem_dockpars_energies_current), &mem_dockpars_energies_current);
setKernelArg(kernel4,15,sizeof(mem_dockpars_conformations_next), &mem_dockpars_conformations_next);
setKernelArg(kernel4,16,sizeof(mem_dockpars_energies_next), &mem_dockpars_energies_next);
setKernelArg(kernel4,17,sizeof(mem_dockpars_evals_of_new_entities), &mem_dockpars_evals_of_new_entities);
setKernelArg(kernel4,18,sizeof(mem_dockpars_prng_states), &mem_dockpars_prng_states);
setKernelArg(kernel4,19,sizeof(dockpars.pop_size), &dockpars.pop_size);
setKernelArg(kernel4,20,sizeof(dockpars.num_of_genes), &dockpars.num_of_genes);
setKernelArg(kernel4,21,sizeof(dockpars.tournament_rate), &dockpars.tournament_rate);
setKernelArg(kernel4,22,sizeof(dockpars.crossover_rate), &dockpars.crossover_rate);
setKernelArg(kernel4,23,sizeof(dockpars.mutation_rate), &dockpars.mutation_rate);
setKernelArg(kernel4,24,sizeof(dockpars.abs_max_dmov), &dockpars.abs_max_dmov);
setKernelArg(kernel4,25,sizeof(dockpars.abs_max_dang), &dockpars.abs_max_dang);
setKernelArg(kernel4,26,sizeof(dockpars.qasp), &dockpars.qasp);
setKernelArg(kernel4,27,sizeof(mem_atom_charges_const), &mem_atom_charges_const);
setKernelArg(kernel4,28,sizeof(mem_atom_types_const), &mem_atom_types_const);
setKernelArg(kernel4,29,sizeof(mem_intraE_contributors_const), &mem_intraE_contributors_const);
setKernelArg(kernel4,30,sizeof(mem_VWpars_AC_const), &mem_VWpars_AC_const);
setKernelArg(kernel4,31,sizeof(mem_VWpars_BD_const), &mem_VWpars_BD_const);
setKernelArg(kernel4,32,sizeof(mem_dspars_S_const), &mem_dspars_S_const);
setKernelArg(kernel4,33,sizeof(mem_dspars_V_const), &mem_dspars_V_const);
setKernelArg(kernel4,34,sizeof(mem_rotlist_const), &mem_rotlist_const);
setKernelArg(kernel4,35,sizeof(mem_ref_coords_x_const), &mem_ref_coords_x_const);
setKernelArg(kernel4,36,sizeof(mem_ref_coords_y_const), &mem_ref_coords_y_const);
setKernelArg(kernel4,37,sizeof(mem_ref_coords_z_const), &mem_ref_coords_z_const);
setKernelArg(kernel4,38,sizeof(mem_rotbonds_moving_vectors_const), &mem_rotbonds_moving_vectors_const);
setKernelArg(kernel4,39,sizeof(mem_rotbonds_unit_vectors_const), &mem_rotbonds_unit_vectors_const);
setKernelArg(kernel4,40,sizeof(mem_ref_orientation_quats_const), &mem_ref_orientation_quats_const);
kernel4_gxsize = blocksPerGridForEachEntity * threadsPerBlock;
kernel4_lxsize = threadsPerBlock;
#ifdef DOCK_DEBUG
......@@ -583,45 +593,47 @@ filled with clock() */
#if !defined (GRADIENT_ENABLED)
// Kernel3
setKernelArg(kernel3,0,sizeof(dockpars.num_of_atoms), &dockpars.num_of_atoms);
setKernelArg(kernel3,1,sizeof(dockpars.num_of_atypes), &dockpars.num_of_atypes);
setKernelArg(kernel3,2,sizeof(dockpars.num_of_intraE_contributors), &dockpars.num_of_intraE_contributors);
setKernelArg(kernel3,3,sizeof(dockpars.gridsize_x), &dockpars.gridsize_x);
setKernelArg(kernel3,4,sizeof(dockpars.gridsize_y), &dockpars.gridsize_y);
setKernelArg(kernel3,5,sizeof(dockpars.gridsize_z), &dockpars.gridsize_z);
setKernelArg(kernel3,6,sizeof(dockpars.grid_spacing), &dockpars.grid_spacing);
setKernelArg(kernel3,7,sizeof(mem_dockpars_fgrids), &mem_dockpars_fgrids);
setKernelArg(kernel3,8,sizeof(dockpars.rotbondlist_length), &dockpars.rotbondlist_length);
setKernelArg(kernel3,9,sizeof(dockpars.coeff_elec), &dockpars.coeff_elec);
setKernelArg(kernel3,10,sizeof(dockpars.coeff_desolv), &dockpars.coeff_desolv);
setKernelArg(kernel3,11,sizeof(mem_dockpars_conformations_next), &mem_dockpars_conformations_next);
setKernelArg(kernel3,12,sizeof(mem_dockpars_energies_next), &mem_dockpars_energies_next);
setKernelArg(kernel3,13,sizeof(mem_dockpars_evals_of_new_entities), &mem_dockpars_evals_of_new_entities);
setKernelArg(kernel3,14,sizeof(mem_dockpars_prng_states), &mem_dockpars_prng_states);
setKernelArg(kernel3,15,sizeof(dockpars.pop_size), &dockpars.pop_size);
setKernelArg(kernel3,16,sizeof(dockpars.num_of_genes), &dockpars.num_of_genes);
setKernelArg(kernel3,17,sizeof(dockpars.lsearch_rate), &dockpars.lsearch_rate);
setKernelArg(kernel3,18,sizeof(dockpars.num_of_lsentities), &dockpars.num_of_lsentities);
setKernelArg(kernel3,19,sizeof(dockpars.rho_lower_bound), &dockpars.rho_lower_bound);
setKernelArg(kernel3,20,sizeof(dockpars.base_dmov_mul_sqrt3), &dockpars.base_dmov_mul_sqrt3);
setKernelArg(kernel3,21,sizeof(dockpars.base_dang_mul_sqrt3), &dockpars.base_dang_mul_sqrt3);
setKernelArg(kernel3,22,sizeof(dockpars.cons_limit), &dockpars.cons_limit);
setKernelArg(kernel3,23,sizeof(dockpars.max_num_of_iters), &dockpars.max_num_of_iters);
setKernelArg(kernel3,24,sizeof(dockpars.qasp), &dockpars.qasp);
setKernelArg(kernel3,25,sizeof(mem_atom_charges_const), &mem_atom_charges_const);
setKernelArg(kernel3,26,sizeof(mem_atom_types_const), &mem_atom_types_const);
setKernelArg(kernel3,27,sizeof(mem_intraE_contributors_const), &mem_intraE_contributors_const);
setKernelArg(kernel3,28,sizeof(mem_VWpars_AC_const), &mem_VWpars_AC_const);
setKernelArg(kernel3,29,sizeof(mem_VWpars_BD_const), &mem_VWpars_BD_const);
setKernelArg(kernel3,30,sizeof(mem_dspars_S_const), &mem_dspars_S_const);
setKernelArg(kernel3,31,sizeof(mem_dspars_V_const), &mem_dspars_V_const);
setKernelArg(kernel3,32,sizeof(mem_rotlist_const), &mem_rotlist_const);
setKernelArg(kernel3,33,sizeof(mem_ref_coords_x_const), &mem_ref_coords_x_const);
setKernelArg(kernel3,34,sizeof(mem_ref_coords_y_const), &mem_ref_coords_y_const);
setKernelArg(kernel3,35,sizeof(mem_ref_coords_z_const), &mem_ref_coords_z_const);
setKernelArg(kernel3,36,sizeof(mem_rotbonds_moving_vectors_const), &mem_rotbonds_moving_vectors_const);
setKernelArg(kernel3,37,sizeof(mem_rotbonds_unit_vectors_const), &mem_rotbonds_unit_vectors_const);
setKernelArg(kernel3,38,sizeof(mem_ref_orientation_quats_const), &mem_ref_orientation_quats_const);
setKernelArg(kernel3,0, sizeof(dockpars.num_of_atoms), &dockpars.num_of_atoms);
setKernelArg(kernel3,1, sizeof(dockpars.num_of_atypes), &dockpars.num_of_atypes);
setKernelArg(kernel3,2, sizeof(dockpars.num_of_intraE_contributors), &dockpars.num_of_intraE_contributors);
setKernelArg(kernel3,3, sizeof(dockpars.gridsize_x), &dockpars.gridsize_x);
setKernelArg(kernel3,4, sizeof(dockpars.gridsize_y), &dockpars.gridsize_y);
setKernelArg(kernel3,5, sizeof(dockpars.gridsize_z), &dockpars.gridsize_z);
setKernelArg(kernel3,6, sizeof(g2), &g2);
setKernelArg(kernel3,7, sizeof(g3), &g3);
setKernelArg(kernel3,8, sizeof(dockpars.grid_spacing), &dockpars.grid_spacing);
setKernelArg(kernel3,9, sizeof(mem_dockpars_fgrids), &mem_dockpars_fgrids);
setKernelArg(kernel3,10,sizeof(dockpars.rotbondlist_length), &dockpars.rotbondlist_length);
setKernelArg(kernel3,11,sizeof(dockpars.coeff_elec), &dockpars.coeff_elec);
setKernelArg(kernel3,12,sizeof(dockpars.coeff_desolv), &dockpars.coeff_desolv);
setKernelArg(kernel3,13,sizeof(mem_dockpars_conformations_next), &mem_dockpars_conformations_next);
setKernelArg(kernel3,14,sizeof(mem_dockpars_energies_next), &mem_dockpars_energies_next);