Commit 1354f05c authored by Leonardo Solis's avatar Leonardo Solis
Browse files

corrected grid forces for ligand with rotbond

parent 1948124e
...@@ -4,6 +4,7 @@ initpop.txt ...@@ -4,6 +4,7 @@ initpop.txt
*.xml *.xml
input/albumin_dock/ input/albumin_dock/
ocladock.wiki/ ocladock.wiki/
bin/
final_population_run* final_population_run*
device/stringify_tmp device/stringify_tmp
KernelProgramBuildInfo.txt KernelProgramBuildInfo.txt
......
...@@ -197,7 +197,11 @@ odock: check-env-all stringify $(SRC) ...@@ -197,7 +197,11 @@ odock: check-env-all stringify $(SRC)
# Example # Example
PDB := 1ac8 # 1ac8: for testing gradients of translation and rotation genes
# 7cpa: for testing gradients of torsion genes (15 torsions)
# 3tmn: for testing gradients of torsion genes (1 torsion)
PDB := 3tmn
NRUN := 1 NRUN := 1
POPSIZE := 10 POPSIZE := 10
TESTNAME:= test TESTNAME:= test
......
...@@ -22,6 +22,9 @@ Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA. ...@@ -22,6 +22,9 @@ Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
*/ */
#define DEBUG_ENERGY
#include "calcenergy_basic.h" #include "calcenergy_basic.h"
// All related pragmas are in defines.h (accesible by host and device code) // All related pragmas are in defines.h (accesible by host and device code)
...@@ -53,6 +56,11 @@ void gpu_calc_energy( ...@@ -53,6 +56,11 @@ void gpu_calc_energy(
__local float* calc_coords_z, __local float* calc_coords_z,
__local float* partial_energies, __local float* partial_energies,
#if defined (DEBUG_ENERGY)
__local float* partial_interE,
__local float* partial_intraE,
#endif
__constant float* atom_charges_const, __constant float* atom_charges_const,
__constant char* atom_types_const, __constant char* atom_types_const,
__constant char* intraE_contributors_const, __constant char* intraE_contributors_const,
...@@ -76,6 +84,12 @@ void gpu_calc_energy( ...@@ -76,6 +84,12 @@ void gpu_calc_energy(
{ {
partial_energies[get_local_id(0)] = 0.0f; partial_energies[get_local_id(0)] = 0.0f;
#if defined (DEBUG_ENERGY)
partial_interE[get_local_id(0)] = 0.0f;
partial_intraE[get_local_id(0)] = 0.0f;
#endif
#if defined (IMPROVE_GRID) #if defined (IMPROVE_GRID)
// INTERMOLECULAR for-loop (intermediate results) // INTERMOLECULAR for-loop (intermediate results)
// It stores a product of two chars // It stores a product of two chars
...@@ -276,6 +290,10 @@ void gpu_calc_energy( ...@@ -276,6 +290,10 @@ void gpu_calc_energy(
|| (y >= dockpars_gridsize_y-1) || (y >= dockpars_gridsize_y-1)
|| (z >= dockpars_gridsize_z-1)){ || (z >= dockpars_gridsize_z-1)){
partial_energies[get_local_id(0)] += 16777216.0f; //100000.0f; partial_energies[get_local_id(0)] += 16777216.0f; //100000.0f;
#if defined (DEBUG_ENERGY)
partial_interE[get_local_id(0)] += 16777216.0f;
#endif
} }
else else
{ {
...@@ -380,6 +398,10 @@ void gpu_calc_energy( ...@@ -380,6 +398,10 @@ void gpu_calc_energy(
// Calculating affinity energy // Calculating affinity energy
partial_energies[get_local_id(0)] += TRILININTERPOL(cube, weights); partial_energies[get_local_id(0)] += TRILININTERPOL(cube, weights);
#if defined (DEBUG_ENERGY)
partial_interE[get_local_id(0)] += TRILININTERPOL(cube, weights);
#endif
// Capturing electrostatic values // Capturing electrostatic values
atom_typeid = dockpars_num_of_atypes; atom_typeid = dockpars_num_of_atypes;
...@@ -444,6 +466,10 @@ void gpu_calc_energy( ...@@ -444,6 +466,10 @@ void gpu_calc_energy(
// Calculating electrosatic energy // Calculating electrosatic energy
partial_energies[get_local_id(0)] += q * TRILININTERPOL(cube, weights); partial_energies[get_local_id(0)] += q * TRILININTERPOL(cube, weights);
#if defined (DEBUG_ENERGY)
partial_interE[get_local_id(0)] += q * TRILININTERPOL(cube, weights);
#endif
// Capturing desolvation values // Capturing desolvation values
atom_typeid = dockpars_num_of_atypes+1; atom_typeid = dockpars_num_of_atypes+1;
...@@ -507,10 +533,36 @@ void gpu_calc_energy( ...@@ -507,10 +533,36 @@ void gpu_calc_energy(
// Calculating desolvation energy // Calculating desolvation energy
partial_energies[get_local_id(0)] += fabs(q) * TRILININTERPOL(cube, weights); partial_energies[get_local_id(0)] += fabs(q) * TRILININTERPOL(cube, weights);
#if defined (DEBUG_ENERGY)
partial_interE[get_local_id(0)] += fabs(q) * TRILININTERPOL(cube, weights);
#endif
} }
} // End atom_id for-loop (INTERMOLECULAR ENERGY) } // End atom_id for-loop (INTERMOLECULAR ENERGY)
#if defined (DEBUG_ENERGY)
barrier(CLK_LOCAL_MEM_FENCE);
if (get_local_id(0) == 0)
{
float energy_interE = partial_interE[0];
for (uint contributor_counter=1;
contributor_counter<NUM_OF_THREADS_PER_BLOCK;
contributor_counter++)
{
energy_interE += partial_interE[contributor_counter];
}
partial_interE[0] = energy_interE;
//printf("%-20s %-10.8f\n", "energy_interE: ", energy_interE);
}
barrier(CLK_LOCAL_MEM_FENCE);
#endif
// In paper: intermolecular and internal energy calculation // In paper: intermolecular and internal energy calculation
// are independent from each other, -> NO BARRIER NEEDED // are independent from each other, -> NO BARRIER NEEDED
// but require different operations, // but require different operations,
...@@ -559,33 +611,49 @@ void gpu_calc_energy( ...@@ -559,33 +611,49 @@ void gpu_calc_energy(
// Calculating van der Waals / hydrogen bond term // Calculating van der Waals / hydrogen bond term
#if defined (NATIVE_PRECISION) #if defined (NATIVE_PRECISION)
partial_energies[get_local_id(0)] += native_divide(VWpars_AC_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],native_powr(atomic_distance,12)); partial_energies[get_local_id(0)] += native_divide(VWpars_AC_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],native_powr(atomic_distance,12));
#if defined (DEBUG_ENERGY)
partial_intraE[get_local_id(0)] += native_divide(VWpars_AC_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],native_powr(atomic_distance,12));
#endif
#elif defined (HALF_PRECISION) #elif defined (HALF_PRECISION)
pfor (uint rotation_counter = get_local_id(0); for (uint rotation_counter = get_local_id(0);
rotation_counter < dockpars_rotbondlist_length; rotation_counter < dockpars_rotbondlist_length;
rotation_counter+=NUM_OF_THREADS_PER_BLOCK) rotation_counter+=NUM_OF_THREADS_PER_BLOCK)
{artial_energies[get_local_id(0)] += half_divide(VWpars_AC_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],half_powr(atomic_distance,12)); {partial_energies[get_local_id(0)] += half_divide(VWpars_AC_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],half_powr(atomic_distance,12));
#else // Full precision #else // Full precision
partial_energies[get_local_id(0)] += VWpars_AC_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid]/powr(atomic_distance,12); partial_energies[get_local_id(0)] += VWpars_AC_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid]/powr(atomic_distance,12);
#endif #endif
if (intraE_contributors_const[3*contributor_counter+2] == 1) //H-bond if (intraE_contributors_const[3*contributor_counter+2] == 1) { //H-bond
#if defined (NATIVE_PRECISION) #if defined (NATIVE_PRECISION)
partial_energies[get_local_id(0)] -= native_divide(VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],native_powr(atomic_distance,10)); partial_energies[get_local_id(0)] -= native_divide(VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],native_powr(atomic_distance,10));
#if defined (DEBUG_ENERGY)
partial_intraE[get_local_id(0)] -= native_divide(VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],native_powr(atomic_distance,10));
#endif
#elif defined (HALF_PRECISION) #elif defined (HALF_PRECISION)
partial_energies[get_local_id(0)] -= half_divide(VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],half_powr(atomic_distance,10)); partial_energies[get_local_id(0)] -= half_divide(VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],half_powr(atomic_distance,10));
#else // Full precision #else // Full precision
partial_energies[get_local_id(0)] -= VWpars_BD_const[atom1_typeid*dockpars_num_of_atypes+atom2_typeid]/powr(atomic_distance,10); partial_energies[get_local_id(0)] -= VWpars_BD_const[atom1_typeid*dockpars_num_of_atypes+atom2_typeid]/powr(atomic_distance,10);
#endif #endif
}
else //van der Waals else { //van der Waals
#if defined (NATIVE_PRECISION) #if defined (NATIVE_PRECISION)
partial_energies[get_local_id(0)] -= native_divide(VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],native_powr(atomic_distance,6)); partial_energies[get_local_id(0)] -= native_divide(VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],native_powr(atomic_distance,6));
#if defined (DEBUG_ENERGY)
partial_intraE[get_local_id(0)] -= native_divide(VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],native_powr(atomic_distance,6));
#endif
#elif defined (HALF_PRECISION) #elif defined (HALF_PRECISION)
partial_energies[get_local_id(0)] -= half_divide(VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],half_powr(atomic_distance,6)); partial_energies[get_local_id(0)] -= half_divide(VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],half_powr(atomic_distance,6));
#else // Full precision #else // Full precision
partial_energies[get_local_id(0)] -= VWpars_BD_const[atom1_typeid*dockpars_num_of_atypes+atom2_typeid]/powr(atomic_distance,6); partial_energies[get_local_id(0)] -= VWpars_BD_const[atom1_typeid*dockpars_num_of_atypes+atom2_typeid]/powr(atomic_distance,6);
#endif #endif
}
// Calculating electrostatic term // Calculating electrostatic term
/* /*
...@@ -600,7 +668,24 @@ void gpu_calc_energy( ...@@ -600,7 +668,24 @@ void gpu_calc_energy(
atomic_distance * (-8.5525f + half_divide(86.9525f,(1.0f + 7.7839f*half_exp(-0.3154f*atomic_distance)))) atomic_distance * (-8.5525f + half_divide(86.9525f,(1.0f + 7.7839f*half_exp(-0.3154f*atomic_distance))))
); );
#else // Full precision #else // Full precision
partial_energies[get_local_id(0)] += dockpars_coeff_elec*atom_charges_const[atom1_id]*atom_charges_const[atom2_id]/ partial_energies[get_local_ #if defined (DEBUG_ENERGY)
barrier(CLK_LOCAL_MEM_FENCE);
if (get_local_id(0) == 0)
{
float energy_interE = partial_interE[0];
for (uint contributor_counter=1;
contributor_counter<NUM_OF_THREADS_PER_BLOCK;
contributor_counter++)
{
energy_interE += partial_interE[contributor_counter];
}
}
printf("%-20s %-10.8f\n", "energy_interE: ", energy_interE);
barrier(CLK_LOCAL_MEM_FENCE);
#endifid(0)] += dockpars_coeff_elec*atom_charges_const[atom1_id]*atom_charges_const[atom2_id]/
(atomic_distance*(-8.5525f + 86.9525f/(1.0f + 7.7839f*exp(-0.3154f*atomic_distance)))); (atomic_distance*(-8.5525f + 86.9525f/(1.0f + 7.7839f*exp(-0.3154f*atomic_distance))));
#endif #endif
*/ */
...@@ -611,6 +696,15 @@ void gpu_calc_energy( ...@@ -611,6 +696,15 @@ void gpu_calc_energy(
dockpars_coeff_elec * atom_charges_const[atom1_id] * atom_charges_const[atom2_id], dockpars_coeff_elec * atom_charges_const[atom1_id] * atom_charges_const[atom2_id],
atomic_distance * (DIEL_A + native_divide(DIEL_B,(1.0f + DIEL_K*native_exp(-DIEL_B_TIMES_H*atomic_distance)))) atomic_distance * (DIEL_A + native_divide(DIEL_B,(1.0f + DIEL_K*native_exp(-DIEL_B_TIMES_H*atomic_distance))))
); );
#if defined (DEBUG_ENERGY)
partial_intraE[get_local_id(0)] += native_divide (
dockpars_coeff_elec * atom_charges_const[atom1_id] * atom_charges_const[atom2_id],
atomic_distance * (DIEL_A + native_divide(DIEL_B,(1.0f + DIEL_K*native_exp(-DIEL_B_TIMES_H*atomic_distance))))
);
#endif
#elif defined (HALF_PRECISION) #elif defined (HALF_PRECISION)
partial_energies[get_local_id(0)] += half_divide ( partial_energies[get_local_id(0)] += half_divide (
dockpars_coeff_elec * atom_charges_const[atom1_id] * atom_charges_const[atom2_id], dockpars_coeff_elec * atom_charges_const[atom1_id] * atom_charges_const[atom2_id],
...@@ -635,6 +729,15 @@ void gpu_calc_energy( ...@@ -635,6 +729,15 @@ void gpu_calc_energy(
(dspars_S_const[atom2_typeid] + (dspars_S_const[atom2_typeid] +
dockpars_qasp*fabs(atom_charges_const[atom2_id]))*dspars_V_const[atom1_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(-atomic_distance*native_divide(atomic_distance,25.92f));
#if defined (DEBUG_ENERGY)
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));
#endif
#elif defined (HALF_PRECISION) #elif defined (HALF_PRECISION)
partial_energies[get_local_id(0)] += ((dspars_S_const[atom1_typeid] + partial_energies[get_local_id(0)] += ((dspars_S_const[atom1_typeid] +
dockpars_qasp*fabs(atom_charges_const[atom1_id]))*dspars_V_const[atom2_typeid] + dockpars_qasp*fabs(atom_charges_const[atom1_id]))*dspars_V_const[atom2_typeid] +
...@@ -666,6 +769,26 @@ void gpu_calc_energy( ...@@ -666,6 +769,26 @@ void gpu_calc_energy(
} }
} }
#if defined (DEBUG_ENERGY)
barrier(CLK_LOCAL_MEM_FENCE);
if (get_local_id(0) == 0)
{
float energy_intraE = partial_intraE[0];
for (uint contributor_counter=1;
contributor_counter<NUM_OF_THREADS_PER_BLOCK;
contributor_counter++)
{
energy_intraE += partial_intraE[contributor_counter];
}
partial_intraE[0] = energy_intraE;
//printf("%-20s %-10.8f\n", "energy_intraE: ", energy_intraE);
}
barrier(CLK_LOCAL_MEM_FENCE);
#endif
} }
#include "kernel1.cl" #include "kernel1.cl"
......
...@@ -34,6 +34,12 @@ Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA. ...@@ -34,6 +34,12 @@ Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
// whose population includes the current entity (which can be determined with get_group_id(0)), // whose population includes the current entity (which can be determined with get_group_id(0)),
// since this determines which reference orientation should be used. // since this determines which reference orientation should be used.
//#define DEBUG_GRAD_TRANSLATION_GENES
//#define DEBUG_GRAD_ROTATION_GENES
#define DEBUG_GRAD_TORSION_GENES
//#define DEBUG_GRAD
void gpu_calc_gradient( void gpu_calc_gradient(
int dockpars_rotbondlist_length, int dockpars_rotbondlist_length,
char dockpars_num_of_atoms, char dockpars_num_of_atoms,
...@@ -71,12 +77,14 @@ void gpu_calc_gradient( ...@@ -71,12 +77,14 @@ void gpu_calc_gradient(
__constant float* ref_coords_z_const, __constant float* ref_coords_z_const,
__constant float* rotbonds_moving_vectors_const, __constant float* rotbonds_moving_vectors_const,
__constant float* rotbonds_unit_vectors_const, __constant float* rotbonds_unit_vectors_const,
__constant float* ref_orientation_quats_const __constant float* ref_orientation_quats_const,
__constant int* rotbonds_const,
__constant int* rotbonds_atoms_const,
__constant int* num_rotating_atoms_per_rotbond_const
// Gradient-related arguments // Gradient-related arguments
// Calculate gradients (forces) for intermolecular energy // Calculate gradients (forces) for intermolecular energy
// Derived from autodockdev/maps.py // Derived from autodockdev/maps.py
// "is_enabled_gradient_calc": enables gradient calculation. // "is_enabled_gradient_calc": enables gradient calculation.
// In Genetic-Generation: no need for gradients // In Genetic-Generation: no need for gradients
// In Gradient-Minimizer: must calculate gradients // In Gradient-Minimizer: must calculate gradients
...@@ -321,6 +329,8 @@ void gpu_calc_gradient( ...@@ -321,6 +329,8 @@ void gpu_calc_gradient(
float dy = y - y_low; float dy = y - y_low;
float dz = z - z_low; float dz = z - z_low;
//printf("%-15s %-5u %-10.8f %-10.8f %-10.8f\n", "dx,dy,dz", atom_id, dx, dy, dz);
// Capturing affinity values // Capturing affinity values
uint ylow_times_g1 = y_low*g1; uint ylow_times_g1 = y_low*g1;
uint yhigh_times_g1 = y_high*g1; uint yhigh_times_g1 = y_high*g1;
...@@ -447,6 +457,8 @@ void gpu_calc_gradient( ...@@ -447,6 +457,8 @@ void gpu_calc_gradient(
vz_y1 = (1 - dx) * z62 + dx * z75; // y = 1 vz_y1 = (1 - dx) * z62 + dx * z75; // y = 1
gradient_inter_z[atom_id] += (1 - dy) * vz_y0 + dy * vz_y1; gradient_inter_z[atom_id] += (1 - dy) * vz_y0 + dy * vz_y1;
//printf("%-15s %-3u %-10.8f %-10.8f %-10.8f %-10.8f %-10.8f %-10.8f\n", "atom aff", atom_id, vx_z0, vx_z1, vy_z0, vy_z1, vz_y0, vz_y1);
// ------------------------------------------------------------------- // -------------------------------------------------------------------
// Calculating gradients (forces) corresponding to // Calculating gradients (forces) corresponding to
// "elec" intermolecular energy // "elec" intermolecular energy
...@@ -473,7 +485,7 @@ void gpu_calc_gradient( ...@@ -473,7 +485,7 @@ void gpu_calc_gradient(
x76 = cube [1][1][1] - cube [0][1][1]; // z = 1 x76 = cube [1][1][1] - cube [0][1][1]; // z = 1
vx_z0 = (1 - dy) * x10 + dy * x52; // z = 0 vx_z0 = (1 - dy) * x10 + dy * x52; // z = 0
vx_z1 = (1 - dy) * x43 + dy * x76; // z = 1 vx_z1 = (1 - dy) * x43 + dy * x76; // z = 1
gradient_inter_x[atom_id] += (1 - dz) * vx_z0 + dz * vx_z1; gradient_inter_x[atom_id] += q * ((1 - dz) * vx_z0 + dz * vx_z1);
// Vector in y-direction // Vector in y-direction
y20 = cube[0][1][0] - cube [0][0][0]; // z = 0 y20 = cube[0][1][0] - cube [0][0][0]; // z = 0
...@@ -482,7 +494,7 @@ void gpu_calc_gradient( ...@@ -482,7 +494,7 @@ void gpu_calc_gradient(
y74 = cube[1][1][1] - cube [1][0][1]; // z = 1 y74 = cube[1][1][1] - cube [1][0][1]; // z = 1
vy_z0 = (1 - dx) * y20 + dx * y51; // z = 0 vy_z0 = (1 - dx) * y20 + dx * y51; // z = 0
vy_z1 = (1 - dx) * y63 + dx * y74; // z = 1 vy_z1 = (1 - dx) * y63 + dx * y74; // z = 1
gradient_inter_y[atom_id] += (1 - dz) * vy_z0 + dz * vy_z1; gradient_inter_y[atom_id] += q *((1 - dz) * vy_z0 + dz * vy_z1);
// Vectors in z-direction // Vectors in z-direction
z30 = cube [0][0][1] - cube [0][0][0]; // y = 0 z30 = cube [0][0][1] - cube [0][0][0]; // y = 0
...@@ -491,7 +503,9 @@ void gpu_calc_gradient( ...@@ -491,7 +503,9 @@ void gpu_calc_gradient(
z75 = cube [1][1][1] - cube [1][1][0]; // y = 1 z75 = cube [1][1][1] - cube [1][1][0]; // y = 1
vz_y0 = (1 - dx) * z30 + dx * z41; // y = 0 vz_y0 = (1 - dx) * z30 + dx * z41; // y = 0
vz_y1 = (1 - dx) * z62 + dx * z75; // y = 1 vz_y1 = (1 - dx) * z62 + dx * z75; // y = 1
gradient_inter_z[atom_id] += (1 - dy) * vz_y0 + dy * vz_y1; gradient_inter_z[atom_id] += q *((1 - dy) * vz_y0 + dy * vz_y1);
//printf("%-15s %-3u %-10.8f %-10.8f %-10.8f %-10.8f %-10.8f %-10.8f\n", "elec", atom_id, vx_z0, vx_z1, vy_z0, vy_z1, vz_y0, vz_y1);
// ------------------------------------------------------------------- // -------------------------------------------------------------------
// Calculating gradients (forces) corresponding to // Calculating gradients (forces) corresponding to
...@@ -519,7 +533,7 @@ void gpu_calc_gradient( ...@@ -519,7 +533,7 @@ void gpu_calc_gradient(
x76 = cube [1][1][1] - cube [0][1][1]; // z = 1 x76 = cube [1][1][1] - cube [0][1][1]; // z = 1
vx_z0 = (1 - dy) * x10 + dy * x52; // z = 0 vx_z0 = (1 - dy) * x10 + dy * x52; // z = 0
vx_z1 = (1 - dy) * x43 + dy * x76; // z = 1 vx_z1 = (1 - dy) * x43 + dy * x76; // z = 1
gradient_inter_x[atom_id] += (1 - dz) * vx_z0 + dz * vx_z1; gradient_inter_x[atom_id] += fabs(q) * ((1 - dz) * vx_z0 + dz * vx_z1);
// Vector in y-direction // Vector in y-direction
y20 = cube[0][1][0] - cube [0][0][0]; // z = 0 y20 = cube[0][1][0] - cube [0][0][0]; // z = 0
...@@ -528,7 +542,7 @@ void gpu_calc_gradient( ...@@ -528,7 +542,7 @@ void gpu_calc_gradient(
y74 = cube[1][1][1] - cube [1][0][1]; // z = 1 y74 = cube[1][1][1] - cube [1][0][1]; // z = 1
vy_z0 = (1 - dx) * y20 + dx * y51; // z = 0 vy_z0 = (1 - dx) * y20 + dx * y51; // z = 0
vy_z1 = (1 - dx) * y63 + dx * y74; // z = 1 vy_z1 = (1 - dx) * y63 + dx * y74; // z = 1
gradient_inter_y[atom_id] += (1 - dz) * vy_z0 + dz * vy_z1; gradient_inter_y[atom_id] += fabs(q) *((1 - dz) * vy_z0 + dz * vy_z1);
// Vectors in z-direction // Vectors in z-direction
z30 = cube [0][0][1] - cube [0][0][0]; // y = 0 z30 = cube [0][0][1] - cube [0][0][0]; // y = 0
...@@ -537,7 +551,10 @@ void gpu_calc_gradient( ...@@ -537,7 +551,10 @@ void gpu_calc_gradient(
z75 = cube [1][1][1] - cube [1][1][0]; // y = 1 z75 = cube [1][1][1] - cube [1][1][0]; // y = 1
vz_y0 = (1 - dx) * z30 + dx * z41; // y = 0 vz_y0 = (1 - dx) * z30 + dx * z41; // y = 0
vz_y1 = (1 - dx) * z62 + dx * z75; // y = 1 vz_y1 = (1 - dx) * z62 + dx * z75; // y = 1
gradient_inter_z[atom_id] += (1 - dy) * vz_y0 + dy * vz_y1; gradient_inter_z[atom_id] += fabs(q) *((1 - dy) * vz_y0 + dy * vz_y1);
//printf("%-15s %-3u %-10.8f %-10.8f %-10.8f %-10.8f %-10.8f %-10.8f\n", "desol", atom_id, vx_z0, vx_z1, vy_z0, vy_z1, vz_y0, vz_y1);
// ------------------------------------------------------------------- // -------------------------------------------------------------------
} }
...@@ -559,6 +576,9 @@ void gpu_calc_gradient( ...@@ -559,6 +576,9 @@ void gpu_calc_gradient(
// Getting atom IDs // Getting atom IDs
uint atom1_id = intraE_contributors_const[3*contributor_counter]; uint atom1_id = intraE_contributors_const[3*contributor_counter];
uint atom2_id = intraE_contributors_const[3*contributor_counter+1]; uint atom2_id = intraE_contributors_const[3*contributor_counter+1];
/*
printf ("%-5u %-5u %-5u\n", contributor_counter, atom1_id, atom2_id);
*/
// Calculating vector components of vector going // Calculating vector components of vector going
// from first atom's to second atom's coordinates // from first atom's to second atom's coordinates
...@@ -587,7 +607,8 @@ void gpu_calc_gradient( ...@@ -587,7 +607,8 @@ void gpu_calc_gradient(
if (intraE_contributors_const[3*contributor_counter+2] == 1) { //H-bond if (intraE_contributors_const[3*contributor_counter+2] == 1) { //H-bond
gradient_per_intracontributor[contributor_counter] += native_divide (10*VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid], gradient_per_intracontributor[contributor_counter] += native_divide (10*VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],
native_powr(atomic_distance, 11) native_powr(atomic_distance, 11)
); );
} }
else { //van der Waals else { //van der Waals
gradient_per_intracontributor[contributor_counter] += native_divide (6*VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid], gradient_per_intracontributor[contributor_counter] += native_divide (6*VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],
...@@ -615,7 +636,7 @@ void gpu_calc_gradient( ...@@ -615,7 +636,7 @@ void gpu_calc_gradient(
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
// Accumulating gradients of each atom from "gradient_per_intracontributor" // Accumulating gradients from "gradient_per_intracontributor" for each each
if (get_local_id(0) == 0) { if (get_local_id(0) == 0) {
for (uint contributor_counter = 0; for (uint contributor_counter = 0;
contributor_counter < dockpars_num_of_intraE_contributors; contributor_counter < dockpars_num_of_intraE_contributors;
...@@ -625,22 +646,23 @@ void gpu_calc_gradient( ...@@ -625,22 +646,23 @@ void gpu_calc_gradient(
uint atom1_id = intraE_contributors_const[3*contributor_counter]; uint atom1_id = intraE_contributors_const[3*contributor_counter];
uint atom2_id = intraE_contributors_const[3*contributor_counter+1]; uint atom2_id = intraE_contributors_const[3*contributor_counter+1];
// Calculating xyz distances in Angstroms // Calculating xyz distances in Angstroms of vector
// between"atom1_id"-to-"atom2_id" // that goes from "atom1_id"-to-"atom2_id"
float subx = (calc_coords_x[atom1_id] - calc_coords_x[atom2_id]) * dockpars_grid_spacing; float subx = (calc_coords_x[atom2_id] - calc_coords_x[atom1_id]);
float suby = (calc_coords_y[atom1_id] - calc_coords_y[atom2_id]) * dockpars_grid_spacing; float suby = (calc_coords_y[atom2_id] - calc_coords_y[atom1_id]);
float subz = (calc_coords_z[atom1_id] - calc_coords_z[atom2_id]) * dockpars_grid_spacing; float subz = (calc_coords_z[atom2_id] - calc_coords_z[atom1_id]);
float dist = native_sqrt(subx*subx + suby*suby + subz*subz);
// Calculating gradients in xyz components. // Calculating gradients in xyz components.
// Gradients for both atoms in a single contributor pair // Gradients for both atoms in a single contributor pair
// have the same magnitude, but opposite directions // have the same magnitude, but opposite directions
gradient_intra_x[atom1_id] += gradient_per_intracontributor[contributor_counter] * subx; gradient_intra_x[atom1_id] -= gradient_per_intracontributor[contributor_counter] * subx / dist;
gradient_intra_y[atom1_id] += gradient_per_intracontributor[contributor_counter] * suby; gradient_intra_y[atom1_id] -= gradient_per_intracontributor[contributor_counter] * suby / dist;
gradient_intra_z[atom1_id] += gradient_per_intracontributor[contributor_counter] * subz; gradient_intra_z[atom1_id] -= gradient_per_intracontributor[contributor_counter] * subz / dist;
gradient_intra_x[atom2_id] -= gradient_per_intracontributor[contributor_counter] * subx; gradient_intra_x[atom2_id] += gradient_per_intracontributor[contributor_counter] * subx / dist;
gradient_intra_y[atom2_id] -= gradient_per_intracontributor[contributor_counter] * suby; gradient_intra_y[atom2_id] += gradient_per_intracontributor[contributor_counter] * suby / dist;
gradient_intra_z[atom2_id] -= gradient_per_intracontributor[contributor_counter] * subz; gradient_intra_z[atom2_id] += gradient_per_intracontributor[contributor_counter] * subz / dist;
} }
} }
...@@ -651,9 +673,26 @@ void gpu_calc_gradient( ...@@ -651,9 +673,26 @@ void gpu_calc_gradient(
for (uint atom_cnt = get_local_id(0); for (uint atom_cnt = get_local_id(0);
atom_cnt < dockpars_num_of_atoms; atom_cnt < dockpars_num_of_atoms;
atom_cnt+= NUM_OF_THREADS_PER_BLOCK) { atom_cnt+= NUM_OF_THREADS_PER_BLOCK) {
gradient_x[atom_cnt] = gradient_inter_x[atom_cnt] + gradient_intra_x[atom_cnt];
gradient_y[atom_cnt] = gradient_inter_y[atom_cnt] + gradient_intra_y[atom_cnt]; // Grid gradients were calculated in the grid space,
gradient_z[atom_cnt] = gradient_inter_z[atom_cnt] + gradient_intra_z[atom_cnt]; // so they have to be put back in Angstrom.
// Intramolecular gradients were already in Angstrom,
// so no scaling for them is required.
gradient_inter_x[atom_cnt] = gradient_inter_x[atom_cnt] / dockpars_grid_spacing;
gradient_inter_y[atom_cnt] = gradient_inter_y[atom_cnt] / dockpars_grid_spacing;
gradient_inter_z[atom_cnt] = gradient_inter_z[atom_cnt] / dockpars_grid_spacing;
gradient_x[atom_cnt] = gradient_inter_x[atom_cnt] ;//+ gradient_intra_x[atom_cnt];