Commit 187afb5c authored by lvs's avatar lvs
Browse files

added smoothing in intraE with SolisWets

parent 389c51af
......@@ -203,11 +203,12 @@ odock: check-env-all stringify $(SRC)
PDB := 3ce3
NRUN := 100
POPSIZE := 500
POPSIZE := 150
TESTNAME:= test
TESTLS := sw
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
$(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)
ASTEX_PDB := 2bsm
ASTEX_NRUN:= 10
......
......@@ -67,6 +67,15 @@ void gpu_calc_energy(
__constant float* atom_charges_const,
__constant char* atom_types_const,
__constant char* intraE_contributors_const,
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
float dockpars_smooth,
__constant float* reqm,
__constant float* reqm_hbond,
// -------------------------------------------
__constant float* VWpars_AC_const,
__constant float* VWpars_BD_const,
__constant float* dspars_S_const,
......@@ -419,36 +428,97 @@ void gpu_calc_energy(
// Calculating atomic_distance
float atomic_distance = native_sqrt(subx*subx + suby*suby + subz*subz)*dockpars_grid_spacing;
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
/*
if (atomic_distance < 1.0f)
atomic_distance = 1.0f;
*/
// Calculating energy contributions
/*
if ((atomic_distance < 8.0f) && (atomic_distance < 20.48f))
*/
if (atomic_distance < 8.0f)
{
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
// Getting type IDs
uint atom1_typeid = atom_types_const[atom1_id];
uint atom2_typeid = atom_types_const[atom2_id];
// Getting optimum pair distance (opt_distance) from reqm and reqm_hbond
// reqm: equilibrium internuclear separation
// (sum of the vdW radii of two like atoms (A)) in the case of vdW
// reqm_hbond: equilibrium internuclear separation
// (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
{
opt_distance = reqm_hbond [atom1_typeid] + reqm_hbond [atom2_typeid];
}
else //van der Waals
{
opt_distance = 0.5f*(reqm [atom1_typeid] + reqm [atom2_typeid]);
}
// Getting smoothed distance
// smoothed_distance = function(atomic_distance, opt_distance)
float smoothed_distance;
float delta_distance = 0.5f*dockpars_smooth;
if (atomic_distance <= (opt_distance - delta_distance)) {
smoothed_distance = atomic_distance + delta_distance;
}
else if (atomic_distance < (opt_distance + delta_distance)) {
smoothed_distance = opt_distance;
}
else { // else if (atomic_distance >= (opt_distance + delta_distance))
smoothed_distance = atomic_distance - delta_distance;
}
/*
if (get_local_id (0) == 0) {
if (intraE_contributors_const[3*contributor_counter+2] == 1) //H-bond
{
printf("%-5s %u %u %f %f %f %f %f %f\n", "hbond", atom1_typeid, atom2_typeid, reqm_hbond [atom1_typeid], reqm_hbond [atom2_typeid], opt_distance, delta_distance, atomic_distance, smoothed_distance);
}
else //van der Waals
{
printf("%-5s %u %u %f %f %f %f %f %f\n", "vdw", atom1_typeid, atom2_typeid, reqm [atom1_typeid], reqm [atom2_typeid], opt_distance, delta_distance, atomic_distance, smoothed_distance);
}
}
*/
// Calculating van der Waals / hydrogen bond term
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(smoothed_distance/*atomic_distance*/,12));
#if defined (DEBUG_ENERGY_KERNEL1) || defined (DEBUG_ENERGY_KERNEL4) || defined (DEBUG_ENERGY_KERNEL3)
partial_intraE[get_local_id(0)] += native_divide(VWpars_AC_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],native_powr(atomic_distance,12));
partial_intraE[get_local_id(0)] += native_divide(VWpars_AC_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],native_powr(smoothed_distance/*atomic_distance*/,12));
#endif
if (intraE_contributors_const[3*contributor_counter+2] == 1) { //H-bond
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(smoothed_distance/*atomic_distance*/,10));
#if defined (DEBUG_ENERGY_KERNEL1) || defined (DEBUG_ENERGY_KERNEL4) || defined (DEBUG_ENERGY_KERNEL3)
partial_intraE[get_local_id(0)] -= native_divide(VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],native_powr(atomic_distance,10));
partial_intraE[get_local_id(0)] -= native_divide(VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],native_powr(smoothed_distance/*atomic_distance*/,10));
#endif
}
else { //van der Waals
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(smoothed_distance/*atomic_distance*/,6));
#if defined (DEBUG_ENERGY_KERNEL1) || defined (DEBUG_ENERGY_KERNEL4) || defined (DEBUG_ENERGY_KERNEL3)
partial_intraE[get_local_id(0)] -= native_divide(VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],native_powr(atomic_distance,6));
partial_intraE[get_local_id(0)] -= native_divide(VWpars_BD_const[atom1_typeid * dockpars_num_of_atypes+atom2_typeid],native_powr(smoothed_distance/*atomic_distance*/,6));
#endif
}
// -------------------------------------------
// Calculating electrostatic term
partial_energies[get_local_id(0)] += native_divide (
......
......@@ -44,6 +44,15 @@ gpu_calc_initpop(
__constant float* atom_charges_const,
__constant char* atom_types_const,
__constant char* intraE_contributors_const,
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
float dockpars_smooth,
__constant float* reqm,
__constant float* reqm_hbond,
// -------------------------------------------
__constant float* VWpars_AC_const,
__constant float* VWpars_BD_const,
__constant float* dspars_S_const,
......@@ -121,6 +130,15 @@ gpu_calc_initpop(
atom_charges_const,
atom_types_const,
intraE_contributors_const,
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
dockpars_smooth,
reqm,
reqm_hbond,
// -------------------------------------------
VWpars_AC_const,
VWpars_BD_const,
dspars_S_const,
......
......@@ -55,6 +55,15 @@ perform_LS(
__constant float* atom_charges_const,
__constant char* atom_types_const,
__constant char* intraE_contributors_const,
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
float dockpars_smooth,
__constant float* reqm,
__constant float* reqm_hbond,
// -------------------------------------------
__constant float* VWpars_AC_const,
__constant float* VWpars_BD_const,
__constant float* dspars_S_const,
......@@ -223,6 +232,15 @@ perform_LS(
atom_charges_const,
atom_types_const,
intraE_contributors_const,
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
dockpars_smooth,
reqm,
reqm_hbond,
// -------------------------------------------
VWpars_AC_const,
VWpars_BD_const,
dspars_S_const,
......@@ -328,6 +346,15 @@ perform_LS(
atom_charges_const,
atom_types_const,
intraE_contributors_const,
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
dockpars_smooth,
reqm,
reqm_hbond,
// -------------------------------------------
VWpars_AC_const,
VWpars_BD_const,
dspars_S_const,
......
......@@ -55,6 +55,15 @@ gpu_gen_and_eval_newpops(
__constant float* atom_charges_const,
__constant char* atom_types_const,
__constant char* intraE_contributors_const,
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
float dockpars_smooth,
__constant float* reqm,
__constant float* reqm_hbond,
// -------------------------------------------
__constant float* VWpars_AC_const,
__constant float* VWpars_BD_const,
__constant float* dspars_S_const,
......@@ -273,6 +282,15 @@ gpu_gen_and_eval_newpops(
atom_charges_const,
atom_types_const,
intraE_contributors_const,
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
dockpars_smooth,
reqm,
reqm_hbond,
// -------------------------------------------
VWpars_AC_const,
VWpars_BD_const,
dspars_S_const,
......
......@@ -29,6 +29,15 @@ gradient_minimizer(
__constant float* atom_charges_const,
__constant char* atom_types_const,
__constant char* intraE_contributors_const,
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
float dockpars_smooth,
__constant float* reqm,
__constant float* reqm_hbond,
// -------------------------------------------
__constant float* VWpars_AC_const,
__constant float* VWpars_BD_const,
__constant float* dspars_S_const,
......@@ -488,6 +497,15 @@ gradient_minimizer(
atom_charges_const,
atom_types_const,
intraE_contributors_const,
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
dockpars_smooth,
reqm,
reqm_hbond,
// -------------------------------------------
VWpars_AC_const,
VWpars_BD_const,
dspars_S_const,
......
......@@ -66,6 +66,13 @@ typedef struct
float abs_max_dmov;
float abs_max_dang;
float lsearch_rate;
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
float smooth;
// -------------------------------------------
unsigned int num_of_lsentities;
float rho_lower_bound;
float base_dmov_mul_sqrt3;
......@@ -89,6 +96,14 @@ typedef struct
float atom_charges_const[MAX_NUM_OF_ATOMS];
char atom_types_const [MAX_NUM_OF_ATOMS];
char intraE_contributors_const[3*MAX_INTRAE_CONTRIBUTORS];
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
float reqm_const [ATYPE_NUM];
float reqm_hbond_const [ATYPE_NUM];
// -------------------------------------------
float VWpars_AC_const [MAX_NUM_OF_ATYPES*MAX_NUM_OF_ATYPES];
float VWpars_BD_const [MAX_NUM_OF_ATYPES*MAX_NUM_OF_ATYPES];
float dspars_S_const [MAX_NUM_OF_ATYPES];
......
......@@ -53,6 +53,13 @@ typedef struct
float mutation_rate;
float crossover_rate;
float lsearch_rate;
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
float smooth;
// -------------------------------------------
unsigned long num_of_ls;
char ls_method[128];
float tournament_rate;
......
......@@ -92,6 +92,17 @@ typedef struct
int atom_rigid_structures [MAX_NUM_OF_ATOMS];
char bonds [MAX_NUM_OF_ATOMS][MAX_NUM_OF_ATOMS];
char intraE_contributors [MAX_NUM_OF_ATOMS][MAX_NUM_OF_ATOMS];
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
// Sizes are hardcoded, ATYPE_NUM=22 float elements as in
// https://git.esa.informatik.tu-darmstadt.de/docking/ocladock/blob/master/host/src/processligand.cpp#L456
// See "User Guide AutoDock 4.2" (page 34)
double reqm [ATYPE_NUM];
double reqm_hbond [ATYPE_NUM];
// -------------------------------------------
double VWpars_A [MAX_NUM_OF_ATYPES][MAX_NUM_OF_ATYPES];
double VWpars_B [MAX_NUM_OF_ATYPES][MAX_NUM_OF_ATYPES];
double VWpars_C [MAX_NUM_OF_ATYPES][MAX_NUM_OF_ATYPES];
......
......@@ -99,6 +99,14 @@ int prepare_const_fields_for_gpu(Liganddata* myligand_reference,
float atom_charges[MAX_NUM_OF_ATOMS];
char atom_types[MAX_NUM_OF_ATOMS];
char intraE_contributors[3*MAX_INTRAE_CONTRIBUTORS];
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
float reqm [ATYPE_NUM];
float reqm_hbond [ATYPE_NUM];
// -------------------------------------------
float VWpars_AC[MAX_NUM_OF_ATYPES*MAX_NUM_OF_ATYPES];
float VWpars_BD[MAX_NUM_OF_ATYPES*MAX_NUM_OF_ATYPES];
float dspars_S[MAX_NUM_OF_ATYPES];
......@@ -184,6 +192,16 @@ int prepare_const_fields_for_gpu(Liganddata* myligand_reference,
}
}
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
// reqm, reqm_hbond: equilibrium internuclear separation for vdW and hbond
for (i= 0; i<ATYPE_NUM/*myligand_reference->num_of_atypes*/; i++) {
reqm[i] = myligand_reference->reqm[i];
reqm_hbond[i] = myligand_reference->reqm_hbond[i];
}
// -------------------------------------------
//van der Waals parameters
for (i=0; i<myligand_reference->num_of_atypes; i++)
for (j=0; j<myligand_reference->num_of_atypes; j++)
......@@ -341,6 +359,14 @@ int prepare_const_fields_for_gpu(Liganddata* myligand_reference,
for (m=0;m<MAX_NUM_OF_ATOMS;m++){ KerConst->atom_charges_const[m] = atom_charges[m]; }
for (m=0;m<MAX_NUM_OF_ATOMS;m++){ KerConst->atom_types_const[m] = atom_types[m]; }
for (m=0;m<3*MAX_INTRAE_CONTRIBUTORS;m++){ KerConst->intraE_contributors_const[m] = intraE_contributors[m]; }
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
for (m=0;m<ATYPE_NUM;m++){ KerConst->reqm_const[m] = reqm[m]; }
for (m=0;m<ATYPE_NUM;m++){ KerConst->reqm_hbond_const[m] = reqm_hbond[m]; }
// -------------------------------------------
for (m=0;m<MAX_NUM_OF_ATYPES*MAX_NUM_OF_ATYPES;m++){ KerConst->VWpars_AC_const[m] = VWpars_AC[m]; }
for (m=0;m<MAX_NUM_OF_ATYPES*MAX_NUM_OF_ATYPES;m++){ KerConst->VWpars_BD_const[m] = VWpars_BD[m]; }
for (m=0;m<MAX_NUM_OF_ATYPES;m++) { KerConst->dspars_S_const[m] = dspars_S[m]; }
......
......@@ -149,6 +149,13 @@ void get_commandpars(const int* argc,
strcpy(mypars->ls_method, "sw"); // "sw": Solis-Wetts,
// "sd": Steepest-Descent",
// "bfgs": Broyden-Fletcher-Goldfarb-Shanno (to be implemented)
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
mypars->smooth = 0.5f;
// -------------------------------------------
mypars->tournament_rate = 60; // 60%
mypars->rho_lower_bound = 0.01; // 0.01
mypars->base_dmov_mul_sqrt3 = 2.0/(*spacing)*sqrt(3.0); // 2 A
......@@ -272,6 +279,23 @@ void get_commandpars(const int* argc,
// MISSING: unsigned long num_of_ls
// ---------------------------------
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
if (strcmp("-smooth", argv [i]) == 0)
{
arg_recognized = 1;
sscanf(argv [i+1], "%f", &tempfloat);
// smooth is measured in Angstrom
if ((tempfloat >= 0.0f) && (tempfloat <= 0.5f))
mypars->smooth = tempfloat;
else
printf("Warning: value of -smooth argument ignored. Value must be a float between 0 and 0.5.\n");
}
// -------------------------------------------
//Argument: local search method:
// "sw": Solis-Wetts
......
......@@ -281,6 +281,14 @@ filled with clock() */
cl_mem mem_atom_charges_const;
cl_mem mem_atom_types_const;
cl_mem mem_intraE_contributors_const;
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
cl_mem mem_reqm_const;
cl_mem mem_reqm_hbond_const;
// -------------------------------------------
cl_mem mem_VWpars_AC_const;
cl_mem mem_VWpars_BD_const;
cl_mem mem_dspars_S_const;
......@@ -302,6 +310,14 @@ filled with clock() */
mallocBufferObject(context,CL_MEM_READ_ONLY,MAX_NUM_OF_ATOMS*sizeof(float), &mem_atom_charges_const);
mallocBufferObject(context,CL_MEM_READ_ONLY,MAX_NUM_OF_ATOMS*sizeof(char), &mem_atom_types_const);
mallocBufferObject(context,CL_MEM_READ_ONLY,3*MAX_INTRAE_CONTRIBUTORS*sizeof(char), &mem_intraE_contributors_const);
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
mallocBufferObject(context,CL_MEM_READ_ONLY,ATYPE_NUM*sizeof(float),&mem_reqm_const);
mallocBufferObject(context,CL_MEM_READ_ONLY,ATYPE_NUM*sizeof(float),&mem_reqm_hbond_const);
// -------------------------------------------
mallocBufferObject(context,CL_MEM_READ_ONLY,MAX_NUM_OF_ATYPES*MAX_NUM_OF_ATYPES*sizeof(float), &mem_VWpars_AC_const);
mallocBufferObject(context,CL_MEM_READ_ONLY,MAX_NUM_OF_ATYPES*MAX_NUM_OF_ATYPES*sizeof(float), &mem_VWpars_BD_const);
mallocBufferObject(context,CL_MEM_READ_ONLY,MAX_NUM_OF_ATYPES*sizeof(float), &mem_dspars_S_const);
......@@ -320,6 +336,14 @@ filled with clock() */
memcopyBufferObjectToDevice(command_queue,mem_atom_charges_const, &KerConst.atom_charges_const, MAX_NUM_OF_ATOMS*sizeof(float));
memcopyBufferObjectToDevice(command_queue,mem_atom_types_const, &KerConst.atom_types_const, MAX_NUM_OF_ATOMS*sizeof(char));
memcopyBufferObjectToDevice(command_queue,mem_intraE_contributors_const, &KerConst.intraE_contributors_const, 3*MAX_INTRAE_CONTRIBUTORS*sizeof(char));
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
memcopyBufferObjectToDevice(command_queue,mem_reqm_const, &KerConst.reqm_const, ATYPE_NUM*sizeof(float));
memcopyBufferObjectToDevice(command_queue,mem_reqm_hbond_const, &KerConst.reqm_hbond_const, ATYPE_NUM*sizeof(float));
// -------------------------------------------
memcopyBufferObjectToDevice(command_queue,mem_VWpars_AC_const, &KerConst.VWpars_AC_const, MAX_NUM_OF_ATYPES*MAX_NUM_OF_ATYPES*sizeof(float));
memcopyBufferObjectToDevice(command_queue,mem_VWpars_BD_const, &KerConst.VWpars_BD_const, MAX_NUM_OF_ATYPES*MAX_NUM_OF_ATYPES*sizeof(float));
memcopyBufferObjectToDevice(command_queue,mem_dspars_S_const, &KerConst.dspars_S_const, MAX_NUM_OF_ATYPES*sizeof(float));
......@@ -398,10 +422,19 @@ filled with clock() */
dockpars.max_num_of_iters = (unsigned int) mypars->max_num_of_iters;
dockpars.qasp = mypars->qasp;
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
//float smooth = 0.5f;
dockpars.smooth = mypars->smooth;
//printf("dockpars.smooth: %f\n", dockpars.smooth);
// -------------------------------------------
unsigned int g2 = dockpars.gridsize_x * dockpars.gridsize_y;
unsigned int g3 = dockpars.gridsize_x * dockpars.gridsize_y * dockpars.gridsize_z;
printf("Local-search chosen method is: %s\n", (strcmp(mypars->ls_method, "sw") == 0)?"Solis-Wetts":
printf("Local-search chosen method is: %s\n", (strcmp(mypars->ls_method, "sw") == 0)?"Solis-Wets":
(strcmp(mypars->ls_method, "sd") == 0)?"Steepest descent": "Unknown");
......@@ -519,17 +552,26 @@ filled with clock() */
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);
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
setKernelArg(kernel1,21,sizeof(dockpars.smooth), &dockpars.smooth);
setKernelArg(kernel1,22,sizeof(mem_reqm_const), &mem_reqm_const);
setKernelArg(kernel1,23,sizeof(mem_reqm_hbond_const), &mem_reqm_hbond_const);
// -------------------------------------------
setKernelArg(kernel1,24/*21*/,sizeof(mem_VWpars_AC_const), &mem_VWpars_AC_const);
setKernelArg(kernel1,25/*22*/,sizeof(mem_VWpars_BD_const), &mem_VWpars_BD_const);
setKernelArg(kernel1,26/*23*/,sizeof(mem_dspars_S_const), &mem_dspars_S_const);
setKernelArg(kernel1,27/*24*/,sizeof(mem_dspars_V_const), &mem_dspars_V_const);
setKernelArg(kernel1,28/*25*/,sizeof(mem_rotlist_const), &mem_rotlist_const);
setKernelArg(kernel1,29/*26*/,sizeof(mem_ref_coords_x_const), &mem_ref_coords_x_const);
setKernelArg(kernel1,30/*27*/,sizeof(mem_ref_coords_y_const), &mem_ref_coords_y_const);
setKernelArg(kernel1,31/*28*/,sizeof(mem_ref_coords_z_const), &mem_ref_coords_z_const);
setKernelArg(kernel1,32/*29*/,sizeof(mem_rotbonds_moving_vectors_const), &mem_rotbonds_moving_vectors_const);
setKernelArg(kernel1,33/*30*/,sizeof(mem_rotbonds_unit_vectors_const), &mem_rotbonds_unit_vectors_const);
setKernelArg(kernel1,34/*31*/,sizeof(mem_ref_orientation_quats_const), &mem_ref_orientation_quats_const);
kernel1_gxsize = blocksPerGridForEachEntity * threadsPerBlock;
kernel1_lxsize = threadsPerBlock;
#ifdef DOCK_DEBUG
......@@ -580,17 +622,26 @@ filled with clock() */
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);
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
setKernelArg(kernel4,30,sizeof(dockpars.smooth), &dockpars.smooth);
setKernelArg(kernel4,31,sizeof(mem_reqm_const), &mem_reqm_const);
setKernelArg(kernel4,32,sizeof(mem_reqm_hbond_const), &mem_reqm_hbond_const);
// -------------------------------------------
setKernelArg(kernel4,33/*30*/,sizeof(mem_VWpars_AC_const), &mem_VWpars_AC_const);
setKernelArg(kernel4,34/*31*/,sizeof(mem_VWpars_BD_const), &mem_VWpars_BD_const);
setKernelArg(kernel4,35/*32*/,sizeof(mem_dspars_S_const), &mem_dspars_S_const);
setKernelArg(kernel4,36/*33*/,sizeof(mem_dspars_V_const), &mem_dspars_V_const);
setKernelArg(kernel4,37/*34*/,sizeof(mem_rotlist_const), &mem_rotlist_const);
setKernelArg(kernel4,38/*35*/,sizeof(mem_ref_coords_x_const), &mem_ref_coords_x_const);
setKernelArg(kernel4,39/*36*/,sizeof(mem_ref_coords_y_const), &mem_ref_coords_y_const);
setKernelArg(kernel4,40/*37*/,sizeof(mem_ref_coords_z_const), &mem_ref_coords_z_const);
setKernelArg(kernel4,41/*38*/,sizeof(mem_rotbonds_moving_vectors_const), &mem_rotbonds_moving_vectors_const);
setKernelArg(kernel4,42/*39*/,sizeof(mem_rotbonds_unit_vectors_const), &mem_rotbonds_unit_vectors_const);
setKernelArg(kernel4,43/*40*/,sizeof(mem_ref_orientation_quats_const), &mem_ref_orientation_quats_const);
kernel4_gxsize = blocksPerGridForEachEntity * threadsPerBlock;
kernel4_lxsize = threadsPerBlock;
#ifdef DOCK_DEBUG
......@@ -634,17 +685,26 @@ if (strcmp(mypars->ls_method, "sw") == 0) {
setKernelArg(kernel3,27,sizeof(mem_atom_charges_const), &mem_atom_charges_const);
setKernelArg(kernel3,28,sizeof(mem_atom_types_const), &mem_atom_types_const);
setKernelArg(kernel3,29,sizeof(mem_intraE_contributors_const), &mem_intraE_contributors_const);
setKernelArg(kernel3,30,sizeof(mem_VWpars_AC_const), &mem_VWpars_AC_const);
setKernelArg(kernel3,31,sizeof(mem_VWpars_BD_const), &mem_VWpars_BD_const);
setKernelArg(kernel3,32,sizeof(mem_dspars_S_const), &mem_dspars_S_const);
setKernelArg(kernel3,33,sizeof(mem_dspars_V_const), &mem_dspars_V_const);
setKernelArg(kernel3,34,sizeof(mem_rotlist_const), &mem_rotlist_const);
setKernelArg(kernel3,35,sizeof(mem_ref_coords_x_const), &mem_ref_coords_x_const);
setKernelArg(kernel3,36,sizeof(mem_ref_coords_y_const), &mem_ref_coords_y_const);
setKernelArg(kernel3,37,sizeof(mem_ref_coords_z_const), &mem_ref_coords_z_const);
setKernelArg(kernel3,38,sizeof(mem_rotbonds_moving_vectors_const), &mem_rotbonds_moving_vectors_const);
setKernelArg(kernel3,39,sizeof(mem_rotbonds_unit_vectors_const), &mem_rotbonds_unit_vectors_const);
setKernelArg(kernel3,40,sizeof(mem_ref_orientation_quats_const), &mem_ref_orientation_quats_const);
// -------------------------------------------
// Smoothed pairwise potentials
// -------------------------------------------
setKernelArg(kernel3,30,sizeof(dockpars.smooth), &dockpars.smooth);
setKernelArg(kernel3,31,sizeof(mem_reqm_const), &mem_reqm_const);
setKernelArg(kernel3,32,sizeof(mem_reqm_hbond_const), &mem_reqm_hbond_const);
// -------------------------------------------
setKernelArg(kernel3,33/*30*/,sizeof(mem_VWpars_AC_const), &mem_VWpars_AC_const);
setKernelArg(kernel3,34/*31*/,sizeof(mem_VWpars_BD_const), &mem_VWpars_BD_const);
setKernelArg(kernel3,35/*32*/,sizeof(mem_dspars_S_const), &mem_dspars_S_const);
setKernelArg(kernel3,36/*33*/,sizeof(mem_dspars_V_const), &mem_dspars_V_const);
setKernelArg(kernel3,37/*34*/,sizeof(mem_rotlist_const), &mem_rotlist_const);
setKernelArg(kernel3,38/*35*/,sizeof(mem_ref_coords_x_const), &mem_ref_coords_x_const);
setKernelArg(kernel3,39/*36*/,sizeof(mem_ref_coords_y_const), &mem_ref_coords_y_const);
setKernelArg(kernel3,40/*37*/,sizeof(mem_ref_coords_z_const), &mem_ref_coords_z_const);
setKernelArg(kernel3,41/*38*/,sizeof(mem_rotbonds_moving_vectors_const), &mem_rotbonds_moving_vectors_const);
setKernelArg(kernel3,42/*39*/,sizeof(mem_rotbonds_unit_vectors_const), &mem_rotbonds_unit_vectors_const);
setKernelArg(kernel3,43/*40*/,sizeof(mem_ref_orientation_quats_const), &mem_ref_orientation_quats_const);
kernel3_gxsize = blocksPerGridForEachLSEntity * threadsPerBlock;
kernel3_lxsize = threadsPerBlock;
#ifdef DOCK_DEBUG
......@@ -682,22 +742,31 @@ if (strcmp(mypars->ls_method, "sw") == 0) {
setKernelArg(kernel5,22,sizeof(mem_atom_charges_const), &mem_atom_charges_const);
setKernelArg(kernel5,23,sizeof(mem_atom_types_const), &mem_atom_types_const);
setKernelArg(kernel5,24,sizeof(mem_intraE_contributors_const), &mem_intraE_contributors_const);