Commit f62ae0c7 authored by lvs's avatar lvs
Browse files

cleaned + better redefinitions

parent b3bb1d46
......@@ -206,7 +206,7 @@ NRUN := 100
POPSIZE := 150
TESTNAME := test
TESTLS := sd
NUM_LSIT := 100
NUM_LSIT := 10
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 -lsmet $(TESTLS) -lsit $(NUM_LSIT) -smooth 0.5
......
......@@ -42,9 +42,10 @@ Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
#define MAX_NUM_OF_ATOMS 256
#define MAX_NUM_OF_ATYPES 14
#define MAX_NUM_OF_ROTBONDS 32
#define MAX_INTRAE_CONTRIBUTORS 4096 // Ideally =MAX_NUM_OF_ATOMS * MAX_NUM_OF_ATOMS,
#define MAX_INTRAE_CONTRIBUTORS 8192 // Ideally =MAX_NUM_OF_ATOMS * MAX_NUM_OF_ATOMS,
// but too large for __local in gradient calculation.
// If calculated as #rotbons * #atoms = 32 * 256 = 8192,
// If calculated as #rotbons * #atoms, then = 32 * 256 = 8192.
// If implementing an array of MAX_INTRAE_CONTRIBUTORS size,
// then error CL_OUT_OF_RESOURCES.
#define MAX_NUM_OF_ROTATIONS (MAX_NUM_OF_ATOMS * MAX_NUM_OF_ROTBONDS)
#define MAX_POPSIZE 2048
......
......@@ -41,7 +41,10 @@ Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
//#define DEBUG_ENERGY_KERNEL5
///*
// Atomic operations used in gradients of intra contributors.
// Only atomic_cmpxchg() works on floats.
// So for atomic add on floats, this link was used:
// https://streamhpc.com/blog/2016-02-09/atomic-operations-for-floats-in-opencl-improved/
void atomicAdd_g_f(volatile __local float *addr, float val)
{
union{
......@@ -73,7 +76,7 @@ void atomicSub_g_f(volatile __local float *addr, float val)
current.u32 = atomic_cmpxchg( (volatile __local unsigned int *)addr, expected.u32, next.u32);
} while( current.u32 != expected.u32 );
}
//*/
void gpu_calc_gradient(
int dockpars_rotbondlist_length,
......@@ -142,12 +145,6 @@ void gpu_calc_gradient(
__local float* gradient_x,
__local float* gradient_y,
__local float* gradient_z,
//----------------------------------
// eliminate unnecessary local storage
//----------------------------------
/*
__local float* gradient_per_intracontributor,
*/
__local float* gradient_genotype
)
{
......@@ -166,18 +163,6 @@ void gpu_calc_gradient(
gradient_intra_z[atom_id] = 0.0f;
}
//----------------------------------
// eliminate unnecessary local storage
//----------------------------------
/*
// Initializing gradients per intramolecular contributor pairs
for (uint intracontrib_atompair_id = get_local_id(0);
intracontrib_atompair_id < dockpars_num_of_intraE_contributors;
intracontrib_atompair_id+= NUM_OF_THREADS_PER_BLOCK) {
gradient_per_intracontributor[intracontrib_atompair_id] = 0.0f;
}
*/
// Initializing gradient genotypes
for (uint gene_cnt = get_local_id(0);
gene_cnt < dockpars_num_of_genes;
......@@ -609,11 +594,9 @@ void gpu_calc_gradient(
contributor_counter < dockpars_num_of_intraE_contributors;
contributor_counter+= NUM_OF_THREADS_PER_BLOCK)
{
//----------------------------------
// eliminate unnecessary local storage
//----------------------------------
// Storing in a private variable
// the gradient contribution of each contributing atomic pair
float priv_gradient_per_intracontributor= 0.0f;
//----------------------------------
// Getting atom IDs
uint atom1_id = intraE_contributors_const[3*contributor_counter];
......@@ -687,25 +670,20 @@ void gpu_calc_gradient(
}
*/
//----------------------------------
// eliminate unnecessary local storage
//----------------------------------
// Calculating van der Waals / hydrogen bond term
/*gradient_per_intracontributor[contributor_counter]*/ priv_gradient_per_intracontributor += native_divide (-12*VWpars_AC_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],
native_powr(smoothed_distance/*atomic_distance*/, 13)
);
priv_gradient_per_intracontributor += native_divide (-12*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
/*gradient_per_intracontributor[contributor_counter]*/ priv_gradient_per_intracontributor += native_divide (10*VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],
native_powr(smoothed_distance/*atomic_distance*/, 11)
);
priv_gradient_per_intracontributor += native_divide (10*VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],
native_powr(smoothed_distance/*atomic_distance*/, 11)
);
}
else { //van der Waals
/*gradient_per_intracontributor[contributor_counter]*/ priv_gradient_per_intracontributor += native_divide (6*VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],
native_powr(smoothed_distance/*atomic_distance*/, 7)
);
priv_gradient_per_intracontributor += native_divide (6*VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],
native_powr(smoothed_distance/*atomic_distance*/, 7)
);
}
// Calculating electrostatic term
......@@ -714,23 +692,17 @@ 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);
/*gradient_per_intracontributor[contributor_counter]*/ 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 * atom_charges_const[atom1_id] * atom_charges_const[atom2_id] * native_divide (upper, lower);
// Calculating desolvation term
/*gradient_per_intracontributor[contributor_counter]*/ priv_gradient_per_intracontributor += (
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]
) *
dockpars_coeff_desolv * -0.07716049382716049 * atomic_distance * native_exp(-0.038580246913580245*native_powr(atomic_distance, 2));
//----------------------------------
//----------------------------------
// eliminate unnecessary local storage
//----------------------------------
// Accumulating gradients from "priv_gradient_per_intracontributor" for each each
// Decomposing "priv_gradient_per_intracontributor"
// into the contribution of each atom of the pair
float subx_div_dist = native_divide(subx, dist);
float suby_div_dist = native_divide(suby, dist);
float subz_div_dist = native_divide(subz, dist);
......@@ -738,66 +710,21 @@ void gpu_calc_gradient(
float priv_intra_gradient_x = priv_gradient_per_intracontributor * subx_div_dist;
float priv_intra_gradient_y = priv_gradient_per_intracontributor * suby_div_dist;
float priv_intra_gradient_z = priv_gradient_per_intracontributor * subz_div_dist;
/*
barrier(CLK_LOCAL_MEM_FENCE);
gradient_intra_x[atom1_id] -= priv_intra_gradient_x;
gradient_intra_y[atom1_id] -= priv_intra_gradient_y;
gradient_intra_z[atom1_id] -= priv_intra_gradient_z;
barrier(CLK_LOCAL_MEM_FENCE);
gradient_intra_x[atom2_id] += priv_intra_gradient_x;
gradient_intra_y[atom2_id] += priv_intra_gradient_y;
gradient_intra_z[atom2_id] += priv_intra_gradient_z;
*/
#if 1
// Calculating gradients in xyz components.
// Gradients for both atoms in a single contributor pair
// have the same magnitude, but opposite directions
/*
gradient_intra_x[atom1_id] -= priv_intra_gradient_x;
gradient_intra_y[atom1_id] -= priv_intra_gradient_y;
gradient_intra_z[atom1_id] -= priv_intra_gradient_z;
*/
atomicSub_g_f(&gradient_intra_x[atom1_id], priv_intra_gradient_x);
atomicSub_g_f(&gradient_intra_y[atom1_id], priv_intra_gradient_y);
atomicSub_g_f(&gradient_intra_z[atom1_id], priv_intra_gradient_z);
/*
gradient_intra_x[atom2_id] += priv_intra_gradient_x;
gradient_intra_y[atom2_id] += priv_intra_gradient_y;
gradient_intra_z[atom2_id] += priv_intra_gradient_z;
*/
//atomic_add(&gradient_intra_x[atom2_id], priv_intra_gradient_x);
//atomic_add(&gradient_intra_y[atom2_id], priv_intra_gradient_y);
//atomic_add(&gradient_intra_z[atom2_id], priv_intra_gradient_z);
atomicAdd_g_f(&gradient_intra_x[atom2_id], priv_intra_gradient_x);
atomicAdd_g_f(&gradient_intra_y[atom2_id], priv_intra_gradient_y);
atomicAdd_g_f(&gradient_intra_z[atom2_id], priv_intra_gradient_z);
//----------------------------------
#endif
}
} // End contributor_counter for-loop (INTRAMOLECULAR ENERGY)
//----------------------------------
// eliminate unnecessary local storage
//----------------------------------
......@@ -841,23 +768,6 @@ void gpu_calc_gradient(
}
*/
barrier(CLK_LOCAL_MEM_FENCE);
// Accumulating inter- and intramolecular gradients
......
......@@ -155,14 +155,6 @@ gradient_minimizer(
__local float gradient_intra_y[MAX_NUM_OF_ATOMS];
__local float gradient_intra_z[MAX_NUM_OF_ATOMS];
//----------------------------------
// eliminate unnecessary local storage
//----------------------------------
/*
__local float gradient_per_intracontributor[MAX_INTRAE_CONTRIBUTORS];
*/
// Accummulated gradient per each ligand atom
__local float gradient_x[MAX_NUM_OF_ATOMS];
__local float gradient_y[MAX_NUM_OF_ATOMS];
......@@ -265,12 +257,8 @@ gradient_minimizer(
__local float max_trans_gene, max_rota_gene, max_tors_gene;
__local float max_trans_stepsize, max_rota_stepsize, max_tors_stepsize;
__local float max_stepsize;
//----------------------------------
//----------------------------------
// fastergrad
//----------------------------------
// Storing torsion genotypes here
__local float torsions_genotype[ACTUAL_GENOTYPE_LENGTH];
//----------------------------------
......@@ -321,23 +309,11 @@ gradient_minimizer(
genotype[20] = 0.0f;
#endif
//----------------------------------
// fastergrad
//----------------------------------
/*
// Calculating maximum possible stepsize (alpha)
__local float max_trans_gene, max_rota_gene, max_tors_gene;
__local float max_trans_stepsize, max_rota_stepsize, max_tors_stepsize;
__local float max_stepsize;
*/
//----------------------------------
if (get_local_id(0) == 0) {
// Finding maximum of the absolute value
// for the three translation genes
max_trans_gene = max(fabs(genotype[0]), fabs(genotype[1]));
max_trans_gene = max(max_trans_gene, fabs(genotype[2]));
max_trans_gene = fmax/*max*/(fabs(genotype[0]), fabs(genotype[1]));
max_trans_gene = fmax/*max*/(max_trans_gene, fabs(genotype[2]));
// Note that MAX_DEV_TRANSLATION needs to be
// expressed in grid size first
......@@ -345,8 +321,8 @@ gradient_minimizer(
// Finding maximum of the absolute value
// for the three Shoemake rotation genes
max_rota_gene = max(fabs(genotype[3]), fabs(genotype[4]));//printf("max_rota_gene: %-10.7f\n", max_rota_gene);
max_rota_gene = max(max_rota_gene, fabs(genotype[5])); //printf("max_rota_gene: %-10.7f\n", max_rota_gene);
max_rota_gene = fmax/*max*/(fabs(genotype[3]), fabs(genotype[4]));//printf("max_rota_gene: %-10.7f\n", max_rota_gene);
max_rota_gene = fmax/*max*/(max_rota_gene, fabs(genotype[5])); //printf("max_rota_gene: %-10.7f\n", max_rota_gene);
// Note that MAX_DEV_ROTATION
// is already expressed within [0, 1]
......@@ -354,15 +330,6 @@ gradient_minimizer(
}
// Copying torsions genes
//----------------------------------
// fastergrad
//----------------------------------
/*
__local float torsions_genotype[ACTUAL_GENOTYPE_LENGTH];
*/
//----------------------------------
for(uint i = get_local_id(0);
i < dockpars_num_of_genes-6;
i+= NUM_OF_THREADS_PER_BLOCK) {
......@@ -394,11 +361,11 @@ gradient_minimizer(
if (get_local_id(0) == 0) {
// Calculating the maximum stepsize using previous three
max_stepsize = min(max_trans_stepsize, max_rota_stepsize);
max_stepsize = min(max_stepsize, max_tors_stepsize);
max_stepsize = fmin/*min*/(max_trans_stepsize, max_rota_stepsize);
max_stepsize = fmin/*min*/(max_stepsize, max_tors_stepsize);
// Capping the stepsize
stepsize = min(stepsize, max_stepsize);
stepsize = fmin/*min*/(stepsize, max_stepsize);
#if defined (DEBUG_MINIMIZER)
//printf("max_genes: %-0.7f %-10.7f %-10.7f %-10.7f\n", max_trans_gene, max_rota_gene, max_tors_gene, stepsize);
......@@ -472,12 +439,6 @@ gradient_minimizer(
gradient_x,
gradient_y,
gradient_z,
//----------------------------------
// eliminate unnecessary local storage
//----------------------------------
/*
gradient_per_intracontributor,
*/
gradient
);
// =============================================================
......@@ -507,8 +468,8 @@ gradient_minimizer(
#endif
// Putting genes back within bounds
candidate_genotype[i] = min(candidate_genotype[i], upper_bounds_genotype[i]);
candidate_genotype[i] = max(candidate_genotype[i], lower_bounds_genotype[i]);
candidate_genotype[i] = fmin/*min*/(candidate_genotype[i], upper_bounds_genotype[i]);
candidate_genotype[i] = fmax/*max*/(candidate_genotype[i], lower_bounds_genotype[i]);
}
......
Supports Markdown
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment