Skip to content
GitLab
Projects
Groups
Snippets
/
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in
Toggle navigation
Menu
Open sidebar
docking
ocladock
Commits
6a2748bf
Commit
6a2748bf
authored
Apr 20, 2018
by
Leonardo Solis
Browse files
gradient kernel ready to debug
parent
43aa19f1
Changes
12
Expand all
Hide whitespace changes
Inline
Side-by-side
Makefile
View file @
6a2748bf
...
...
@@ -202,8 +202,8 @@ odock: check-env-all stringify $(SRC)
# 3tmn: for testing gradients of torsion genes (1 torsion)
PDB
:=
7cpa
NRUN
:=
1
POPSIZE
:=
1
NRUN
:=
1
00
POPSIZE
:=
500
TESTNAME
:=
test
test
:
odock
...
...
common/defines.h
View file @
6a2748bf
...
...
@@ -55,8 +55,9 @@ 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 IMPROVE_GRID
//#define NATIVE_PRECISION
#define MAPPED_COPY
#define GRADIENT_ENABLED
#endif
/* DEFINES_H_ */
device/auxiliary_genetic.cl
View file @
6a2748bf
...
...
@@ -70,13 +70,7 @@ float gpu_randf(
#
if
defined
(
REPRO
)
state
=
0.55f
; //0.55f;
#
else
#
if
defined
(
NATIVE_PRECISION
)
state
=
native_divide
(
gpu_rand
(
prng_states
)
,
MAX_UINT
)
*0.999999f
;
#
elif
defined
(
HALF_PRECISION
)
state
=
half_divide
(
gpu_rand
(
prng_states
)
,
MAX_UINT
)
*0.999999f
;
#
else
//
Full
precision
state
=
(((
float
)
gpu_rand
(
prng_states
))
/MAX_UINT
)
*0.999999f
;
#
endif
#
endif
return
state
;
...
...
device/auxiliary_gradient.cl
View file @
6a2748bf
#
define
DEBUG_GRADDESC_ENABLED
//
#define
DEBUG_GRADDESC_ENABLED
//
Implementation
of
auxiliary
functions
//
for
the
gradient-based
minimizer
void
is_gradDescent_enabled
(
__local
bool*
is_perturb_gt_gene_min,
__local
float*
local_gNorm,
float
gradMin_tol,
__local
bool*
is_genotype_valid,
__local
uint*
local_nIter,
uint
gradMin_maxiter,
__local
float*
local_perturbation,
__local
float*
local_genotype,
__constant
float*
gradMin_conformation_min_perturbation,
__local
float*
local_candidate_genotype,
uint
gradMin_numElements,
__local
bool*
is_gradDescentEn
)
{
bool
is_gNorm_gt_gMin
;
bool
is_nIter_lt_maxIter
;
bool
is_
perturb_gt_genotype
;
bool
is_
valid
;
if
(
get_local_id
(
0
)
==
0
)
{
//is_gNorm_gt_gMin
=
(
local_gNorm[0]
>=
gradMin_tol
)
;
is_nIter_lt_maxIter
=
(
local_nIter[0]
<
gradMin_maxiter
)
;
//is_perturb_gt_genotype
=
true
;
is_valid
=
true
;
}
/*
//
Verifying
that
Shoemake
genes
do
not
get
out
of
valid
range.
//
If
they
do
so,
then
set
them
to
0.0f
if
(
get_local_id
(
0
)
<
3
)
{
if
((
local_genotype[get_local_id
(
0
)
]
<
0.0f
)
&&
(
local_genotype[get_local_id
(
0
)
]
>
1.0f
))
{
local_genotype[get_local_id
(
0
)
]
=
0.0f
;
if
((
local_
candidate_
genotype[get_local_id
(
0
)
]
<
0.0f
)
&&
(
local_
candidate_
genotype[get_local_id
(
0
)
]
>
1.0f
))
{
local_
candidate_
genotype[get_local_id
(
0
)
]
=
0.0f
;
}
}
*/
#
if
0
//
For
every
gene,
let
's
determine
//
if
perturbation
is
greater
than
min
conformation
//
Using
every
of
its
genes,
let
's
determine
//
if
candidate
genotype
is
valid
for
(
uint
i
=
get_local_id
(
0
)
;
i
<
gradMin_numElements
;
i+=
NUM_OF_THREADS_PER_BLOCK
)
{
is_
perturb_gt_gene_min[i]
=
(
local_perturbation[i]
>=
gradMin_conformation_min_perturbation
[i]
)
;
is_
genotype_valid[i]
=
!isnan
(
local_candidate_genotype[i]
)
&&
!isinf
(
local_candidate_genotype
[i]
)
;
}
#
endif
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
if
(
get_local_id
(
0
)
==
0
)
{
#
if
0
//
Reduce
all
is_perturb_gt_gene_min
's
//
Reduce
all
"valid"
values
of
genes
//
into
their
corresponding
genotype
for
(
uint
i
=
0
;
i
<
gradMin_numElements
;
i++
)
{
is_
perturb_gt_genotype
=
is_perturb_gt_genotype
&&
is_perturb_gt_gene_min
[i]
;
is_
valid
=
is_valid
&&
is_genotype_valid
[i]
;
#
if
defined
(
DEBUG_GRADDESC_ENABLED
)
printf
(
"is_
perturb_gt_gene_min
[%u]?: %s\n"
,
i,
(
is_
perturb_gt_gene_min
[i]
==
true
)
?
"yes"
:
"no"
)
;
//
printf
(
"is_
genotype_valid
[%u]?: %s\n"
,
i,
(
is_
genotype_valid
[i]
==
true
)
?
"yes"
:
"no"
)
;
#
endif
}
#
endif
//
Reduce
all
three
previous
//
partial
evaluations
(
gNorm,
nIter,
perturb
)
into
a
final
one
//*is_gradDescentEn
=
is_gNorm_gt_gMin
&&
is_nIter_lt_maxIter
&&
is_perturb_gt_genotype
;
*is_gradDescentEn
=
is_nIter_lt_maxIter
;
//
Reduce
all
previous
partial
evaluations
(
nIter,
valid
)
into
a
final
one
*is_gradDescentEn
=
is_nIter_lt_maxIter
&&
is_valid
;
#
if
1
#
if
defined
(
DEBUG_GRADDESC_ENABLED
)
if
(
get_local_id
(
0
)
==
0
)
{
//printf
(
"is_gNorm_gt_gMin?: %s\n"
,
(
is_gNorm_gt_gMin
==
true
)
?
"yes"
:
"no"
)
;
//printf
(
"is_nIter_lt_maxIter?: %s\n"
,
(
is_nIter_lt_maxIter
==
true
)
?
"yes"
:
"no"
)
;
//printf
(
"is_perturb_gt_genotype?: %s\n"
,
(
is_perturb_gt_genotype
==
true
)
?
"yes"
:
"no"
)
;
printf
(
"Continue gradient iteration?: %s\n"
,
(
*is_gradDescentEn
==
true
)
?
"yes"
:
"no"
)
;
}
//printf
(
"is_gNorm_gt_gMin?: %s\n"
,
(
is_gNorm_gt_gMin
==
true
)
?
"yes"
:
"no"
)
;
//printf
(
"is_nIter_lt_maxIter?: %s\n"
,
(
is_nIter_lt_maxIter
==
true
)
?
"yes"
:
"no"
)
;
//printf
(
"is_valid?: %s\n"
,
(
is_valid
==
true
)
?
"yes"
:
"no"
)
;
printf
(
"Continue grad-mini?: %s\n"
,
(
*is_gradDescentEn
==
true
)
?
"yes"
:
"no"
)
;
#
endif
#
endif
}
}
void
gradient_norm
(
__local
float*
vector1,
uint
inputSize,
__local
float*
init,
__local
float*
inner_product
)
{
float
temp
=
0.0f
;
//
Element-wise
multiplication
for
(
uint
i
=
get_local_id
(
0
)
;
i
<
inputSize
;
i+=
NUM_OF_THREADS_PER_BLOCK
)
{
init[i]
=
vector1[i]
*
vector1[i]
;
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
//
Accumulating
dot
product,
//
and
then
getting
the
norm
if
(
get_local_id
(
0
)
==
0
)
{
for
(
uint
i
=
0
;
i
<
inputSize
;
i
++
)
{
temp
+=
init[i]
;
}
*inner_product
=
native_sqrt
(
temp
)
;
}
}
device/calcenergy.cl
View file @
6a2748bf
This diff is collapsed.
Click to expand it.
device/calcgradient.cl
View file @
6a2748bf
...
...
@@ -38,7 +38,7 @@ Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
//#define
DEBUG_GRAD_TRANSLATION_GENES
//#define
DEBUG_GRAD_ROTATION_GENES
//#define
DEBUG_GRAD_TORSION_GENES
//#define
DEBUG_
GRAD
//#define
DEBUG_
ENERGY_KERNEL5
void
gpu_calc_gradient
(
int
dockpars_rotbondlist_length,
...
...
@@ -53,11 +53,13 @@ void gpu_calc_gradient(
float
dockpars_coeff_elec,
float
dockpars_qasp,
float
dockpars_coeff_desolv,
//
Some
OpenCL
compilers
don
't
allow
declaring
//
local
variables
within
non-kernel
functions.
//
These
local
variables
must
be
declared
in
a
kernel,
//
and
then
passed
to
non-kernel
functions.
__local
float*
genotype,
__local
float*
energy,
__local
int*
run_id,
__local
float*
calc_coords_x,
...
...
@@ -212,8 +214,6 @@ void gpu_calc_gradient(
atom_to_rotate[2]
-=
rotation_movingvec[2]
;
//
Transforming
torsion
angles
into
quaternions
//
FIXME:
add
precision
choices
with
preprocessor
directives:
//
NATIVE_PRECISION,
HALF_PRECISION,
Full
precision
rotation_angle
=
native_divide
(
rotation_angle,
2.0f
)
;
float
sin_angle
=
native_sin
(
rotation_angle
)
;
quatrot_left_q
=
native_cos
(
rotation_angle
)
;
...
...
@@ -331,6 +331,17 @@ void gpu_calc_gradient(
//printf("%-15s %-5u %-10.8f %-10.8f %-10.8f\n", "dx,dy,dz", atom_id, dx, dy, dz);
// Calculating interpolation weights
float weights[2][2][2];
weights [0][0][0] = (1-dx)*(1-dy)*(1-dz);
weights [1][0][0] = dx*(1-dy)*(1-dz);
weights [0][1][0] = (1-dx)*dy*(1-dz);
weights [1][1][0] = dx*dy*(1-dz);
weights [0][0][1] = (1-dx)*(1-dy)*dz;
weights [1][0][1] = dx*(1-dy)*dz;
weights [0][1][1] = (1-dx)*dy*dz;
weights [1][1][1] = dx*dy*dz;
// Capturing affinity values
uint ylow_times_g1 = y_low*g1;
uint yhigh_times_g1 = y_high*g1;
...
...
@@ -401,16 +412,6 @@ void gpu_calc_gradient(
//
-------------------------------------------------------------------
//
Vector
in
x-direction
/*
x10
=
grid[int
(
vertices[1]
)
]
-
grid[int
(
vertices[0]
)
]
#
z
=
0
x52
=
grid[int
(
vertices[5]
)
]
-
grid[int
(
vertices[2]
)
]
#
z
=
0
x43
=
grid[int
(
vertices[4]
)
]
-
grid[int
(
vertices[3]
)
]
#
z
=
1
x76
=
grid[int
(
vertices[7]
)
]
-
grid[int
(
vertices[6]
)
]
#
z
=
1
vx_z0
=
(
1-yd
)
*
x10
+
yd
*
x52
#
z
=
0
vx_z1
=
(
1-yd
)
*
x43
+
yd
*
x76
#
z
=
1
gradient[0]
=
(
1-zd
)
*
vx_z0
+
zd
*
vx_z1
*/
x10
=
cube
[1][0][0]
-
cube
[0][0][0]
; // z = 0
x52
=
cube
[1][1][0]
-
cube
[0][1][0]
; // z = 0
x43
=
cube
[1][0][1]
-
cube
[0][0][1]
; // z = 1
...
...
@@ -420,16 +421,6 @@ void gpu_calc_gradient(
gradient_inter_x[atom_id]
+=
(
1
-
dz
)
*
vx_z0
+
dz
*
vx_z1
;
//
Vector
in
y-direction
/*
y20
=
grid[int
(
vertices[2]
)
]
-
grid[int
(
vertices[0]
)
]
#
z
=
0
y51
=
grid[int
(
vertices[5]
)
]
-
grid[int
(
vertices[1]
)
]
#
z
=
0
y63
=
grid[int
(
vertices[6]
)
]
-
grid[int
(
vertices[3]
)
]
#
z
=
1
y74
=
grid[int
(
vertices[7]
)
]
-
grid[int
(
vertices[4]
)
]
#
z
=
1
vy_z0
=
(
1-xd
)
*
y20
+
xd
*
y51
#
z
=
0
y_z1
=
(
1-xd
)
*
y63
+
xd
*
y74
#
z
=
1
gradient[1]
=
(
1-zd
)
*
vy_z0
+
zd
*
vy_z1
*/
y20
=
cube[0][1][0]
-
cube
[0][0][0]
; // z = 0
y51
=
cube[1][1][0]
-
cube
[1][0][0]
; // z = 0
y63
=
cube[0][1][1]
-
cube
[0][0][1]
; // z = 1
...
...
@@ -439,16 +430,6 @@ void gpu_calc_gradient(
gradient_inter_y[atom_id]
+=
(
1
-
dz
)
*
vy_z0
+
dz
*
vy_z1
;
//
Vectors
in
z-direction
/*
z30
=
grid[int
(
vertices[3]
)
]
-
grid[int
(
vertices[0]
)
]
#
y
=
0
z41
=
grid[int
(
vertices[4]
)
]
-
grid[int
(
vertices[1]
)
]
#
y
=
0
z62
=
grid[int
(
vertices[6]
)
]
-
grid[int
(
vertices[2]
)
]
#
y
=
1
z75
=
grid[int
(
vertices[7]
)
]
-
grid[int
(
vertices[5]
)
]
#
y
=
1
vz_y0
=
(
1-xd
)
*
z30
+
xd
*
z41
#
y
=
0
vz_y1
=
(
1-xd
)
*
z62
+
xd
*
z75
#
y
=
1
gradient[2]
=
(
1-yd
)
*
vz_y0
+
yd
*
vz_y1
*/
z30
=
cube
[0][0][1]
-
cube
[0][0][0]
; // y = 0
z41
=
cube
[1][0][1]
-
cube
[1][0][0]
; // y = 0
z62
=
cube
[0][1][1]
-
cube
[0][1][0]
; // y = 1
...
...
@@ -605,7 +586,6 @@ void gpu_calc_gradient(
gradient_per_intracontributor[contributor_counter]
+=
native_divide
(
10*VWpars_BD_const[atom1_typeid
*
dockpars_num_of_atypes+atom2_typeid],
native_powr
(
atomic_distance,
11
)
)
;
}
else
{
//van
der
Waals
gradient_per_intracontributor[contributor_counter]
+=
native_divide
(
6*VWpars_BD_const[atom1_typeid
*
dockpars_num_of_atypes+atom2_typeid],
...
...
@@ -651,16 +631,20 @@ void gpu_calc_gradient(
float
subz
=
(
calc_coords_z[atom2_id]
-
calc_coords_z[atom1_id]
)
;
float
dist
=
native_sqrt
(
subx*subx
+
suby*suby
+
subz*subz
)
;
float
subx_div_dist
=
native_divide
(
subx,
dist
)
;
float
suby_div_dist
=
native_divide
(
suby,
dist
)
;
float
subz_div_dist
=
native_divide
(
subz,
dist
)
;
//
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]
-=
gradient_per_intracontributor[contributor_counter]
*
subx
/
dist
;
gradient_intra_y[atom1_id]
-=
gradient_per_intracontributor[contributor_counter]
*
suby
/
dist
;
gradient_intra_z[atom1_id]
-=
gradient_per_intracontributor[contributor_counter]
*
subz
/
dist
;
gradient_intra_x[atom1_id]
-=
gradient_per_intracontributor[contributor_counter]
*
subx
_div_
dist
;
gradient_intra_y[atom1_id]
-=
gradient_per_intracontributor[contributor_counter]
*
suby
_div_
dist
;
gradient_intra_z[atom1_id]
-=
gradient_per_intracontributor[contributor_counter]
*
subz
_div_
dist
;
gradient_intra_x[atom2_id]
+=
gradient_per_intracontributor[contributor_counter]
*
subx
/
dist
;
gradient_intra_y[atom2_id]
+=
gradient_per_intracontributor[contributor_counter]
*
suby
/
dist
;
gradient_intra_z[atom2_id]
+=
gradient_per_intracontributor[contributor_counter]
*
subz
/
dist
;
gradient_intra_x[atom2_id]
+=
gradient_per_intracontributor[contributor_counter]
*
subx
_div_
dist
;
gradient_intra_y[atom2_id]
+=
gradient_per_intracontributor[contributor_counter]
*
suby
_div_
dist
;
gradient_intra_z[atom2_id]
+=
gradient_per_intracontributor[contributor_counter]
*
subz
_div_
dist
;
//printf
(
"%-20s %-10u %-5u %-5u %-10.8f\n"
,
"grad_intracontrib"
,
contributor_counter,
atom1_id,
atom2_id,
gradient_per_intracontributor[contributor_counter]
)
;
}
...
...
@@ -679,9 +663,9 @@ void gpu_calc_gradient(
//
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_inter_x[atom_cnt]
=
native_divide
(
gradient_inter_x[atom_cnt]
,
dockpars_grid_spacing
)
;
gradient_inter_y[atom_cnt]
=
native_divide
(
gradient_inter_y[atom_cnt]
,
dockpars_grid_spacing
)
;
gradient_inter_z[atom_cnt]
=
native_divide
(
gradient_inter_z[atom_cnt]
,
dockpars_grid_spacing
)
;
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]
;
...
...
@@ -722,7 +706,6 @@ void gpu_calc_gradient(
printf
(
"gradient_y:%f\n"
,
gradient_genotype
[1]
)
;
printf
(
"gradient_z:%f\n"
,
gradient_genotype
[2]
)
;
#
endif
}
//
------------------------------------------
...
...
@@ -748,13 +731,13 @@ void gpu_calc_gradient(
printf
(
"%-20s %-10.5f %-10.5f %-10.5f\n"
,
"initial torque: "
,
torque_rot.x,
torque_rot.y,
torque_rot.z
)
;
#
endif
//
C
enter
of
rotation
//
Declaring
a
variable
to
hold
the
c
enter
of
rotation
//
In
getparameters.cpp,
it
indicates
//
translation
genes
are
in
grid
spacing
(
instead
of
Angstroms
)
float3
about
;
about.x
=
/*30*/
genotype[0]
;
about.y
=
/*30*/
genotype[1]
;
about.z
=
/*30*/
genotype[2]
;
about.x
=
genotype[0]
;
about.y
=
genotype[1]
;
about.z
=
genotype[2]
;
//
Temporal
variable
to
calculate
translation
differences.
//
They
are
converted
back
to
Angstroms
here
...
...
@@ -1004,9 +987,9 @@ void gpu_calc_gradient(
//
Assignment
of
gene-based
gradient
gradient_genotype[rotbond_id+6]
=
torque_on_axis
*
(
M_PI
/
180.0f
)
;
//
#if
defined
(
DEBUG_GRAD_TORSION_GENES
)
//
printf
(
"gradient_torsion [%u] :%f\n"
,
rotbond_id+6,
gradient_genotype
[rotbond_id+6]
)
;
//
#endif
#
if
defined
(
DEBUG_GRAD_TORSION_GENES
)
printf
(
"gradient_torsion [%u] :%f\n"
,
rotbond_id+6,
gradient_genotype
[rotbond_id+6]
)
;
#
endif
}
//
End
of
iterations
over
rotatable
bonds
}
...
...
device/kernel1.cl
View file @
6a2748bf
...
...
@@ -20,8 +20,6 @@ along with this program; if not, write to the Free Software
Foundation,
Inc.,
51
Franklin
Street,
Fifth
Floor,
Boston,
MA
02110-1301,
USA.
*/
__kernel
void
__attribute__
((
reqd_work_group_size
(
NUM_OF_THREADS_PER_BLOCK,1,1
)))
gpu_calc_initpop
(
char
dockpars_num_of_atoms,
...
...
@@ -67,9 +65,9 @@ gpu_calc_initpop(
__local
float
calc_coords_y[MAX_NUM_OF_ATOMS]
;
__local
float
calc_coords_z[MAX_NUM_OF_ATOMS]
;
__local
float
partial_energies[NUM_OF_THREADS_PER_BLOCK]
;
#
if
defined
(
DEBUG_ENERGY
)
__local
float
partial_interE
[NUM_OF_THREADS_PER_BLOCK]
;
__local
float
partial_intraE
[NUM_OF_THREADS_PER_BLOCK]
;
#
if
defined
(
DEBUG_ENERGY
_KERNEL1
)
__local
float
partial_interE[NUM_OF_THREADS_PER_BLOCK]
;
__local
float
partial_intraE[NUM_OF_THREADS_PER_BLOCK]
;
#
endif
//
Copying
genotype
from
global
memory
...
...
@@ -109,7 +107,7 @@ gpu_calc_initpop(
calc_coords_y,
calc_coords_z,
partial_energies,
#
if
defined
(
DEBUG_ENERGY
)
#
if
defined
(
DEBUG_ENERGY
_KERNEL1
)
partial_interE,
partial_intraE,
#
endif
...
...
@@ -134,5 +132,9 @@ gpu_calc_initpop(
if
(
get_local_id
(
0
)
==
0
)
{
dockpars_energies_current[get_group_id
(
0
)
]
=
energy
;
dockpars_evals_of_new_entities[get_group_id
(
0
)
]
=
1
;
#
if
defined
(
DEBUG_ENERGY_KERNEL1
)
printf
(
"%-18s [%-5s]---{%-5s} [%-10.8f]---{%-10.8f}\n"
,
"-ENERGY-KERNEL1-"
,
"GRIDS"
,
"INTRA"
,
partial_interE[0],
partial_intraE[0]
)
;
#
endif
}
}
device/kernel3.cl
View file @
6a2748bf
...
...
@@ -97,7 +97,7 @@ perform_LS(
__local
float
calc_coords_z[MAX_NUM_OF_ATOMS]
;
__local
float
partial_energies[NUM_OF_THREADS_PER_BLOCK]
;
#
if
defined
(
DEBUG_ENERGY
)
#
if
defined
(
DEBUG_ENERGY
_KERNEL3
)
__local
float
partial_interE
[NUM_OF_THREADS_PER_BLOCK]
;
__local
float
partial_intraE
[NUM_OF_THREADS_PER_BLOCK]
;
#
endif
...
...
@@ -209,7 +209,7 @@ perform_LS(
calc_coords_y,
calc_coords_z,
partial_energies,
#
if
defined
(
DEBUG_ENERGY
)
#
if
defined
(
DEBUG_ENERGY
_KERNEL3
)
partial_interE,
partial_intraE,
#
endif
...
...
@@ -233,6 +233,10 @@ perform_LS(
if
(
get_local_id
(
0
)
==
0
)
{
evaluation_cnt++
;
#
if
defined
(
DEBUG_ENERGY_KERNEL3
)
printf
(
"%-18s [%-5s]---{%-5s} [%-10.8f]---{%-10.8f}\n"
,
"-ENERGY-KERNEL3-"
,
"GRIDS"
,
"INTRA"
,
partial_interE[0],
partial_intraE[0]
)
;
#
endif
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
...
...
@@ -307,7 +311,7 @@ perform_LS(
calc_coords_y,
calc_coords_z,
partial_energies,
#
if
defined
(
DEBUG_ENERGY
)
#
if
defined
(
DEBUG_ENERGY
_KERNEL3
)
partial_interE,
partial_intraE,
#
endif
...
...
@@ -331,6 +335,10 @@ perform_LS(
if
(
get_local_id
(
0
)
==
0
)
{
evaluation_cnt++
;
#
if
defined
(
DEBUG_ENERGY_KERNEL3
)
printf
(
"%-18s [%-5s]---{%-5s} [%-10.8f]---{%-10.8f}\n"
,
"-ENERGY-KERNEL3-"
,
"GRIDS"
,
"INTRA"
,
partial_interE[0],
partial_intraE[0]
)
;
#
endif
}
barrier
(
CLK_LOCAL_MEM_FENCE
)
;
...
...
device/kernel4.cl
View file @
6a2748bf
...
...
@@ -89,7 +89,7 @@ gpu_gen_and_eval_newpops(
__local
float
calc_coords_y[MAX_NUM_OF_ATOMS]
;
__local
float
calc_coords_z[MAX_NUM_OF_ATOMS]
;
__local
float
partial_energies[NUM_OF_THREADS_PER_BLOCK]
;
#
if
defined
(
DEBUG_ENERGY
)
#
if
defined
(
DEBUG_ENERGY
_KERNEL4
)
__local
float
partial_interE
[NUM_OF_THREADS_PER_BLOCK]
;
__local
float
partial_intraE
[NUM_OF_THREADS_PER_BLOCK]
;
#
endif
...
...
@@ -259,7 +259,7 @@ gpu_gen_and_eval_newpops(
calc_coords_y,
calc_coords_z,
partial_energies,
#
if
defined
(
DEBUG_ENERGY
)
#
if
defined
(
DEBUG_ENERGY
_KERNEL4
)
partial_interE,
partial_intraE,
#
endif
...
...
@@ -284,6 +284,10 @@ gpu_gen_and_eval_newpops(
if
(
get_local_id
(
0
)
==
0
)
{
dockpars_evals_of_new_entities[get_group_id
(
0
)
]
=
1
;
dockpars_energies_next[get_group_id
(
0
)
]
=
energy
;
#
if
defined
(
DEBUG_ENERGY_KERNEL4
)
printf
(
"%-18s [%-5s]---{%-5s} [%-10.8f]---{%-10.8f}\n"
,
"-ENERGY-KERNEL4-"
,
"GRIDS"
,
"INTRA"
,
partial_interE[0],
partial_intraE[0]
)
;
#
endif
}
//
Copying
new
offspring
to
next
generation
...
...
device/kernel_gradient.cl
View file @
6a2748bf
This diff is collapsed.
Click to expand it.
host/inc/performdocking.h
View file @
6a2748bf
...
...
@@ -57,10 +57,8 @@ Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
// This should be ultimately configurable by the user as program exec. flags.
typedef
struct
{
float
tolerance
;
unsigned
int
max_num_of_iters
;
float
alpha
;
float
h
;
float
conformation_min_perturbation
[
ACTUAL_GENOTYPE_LENGTH
];
}
Gradientparameters
;
...
...
host/src/performdocking.cpp
View file @
6a2748bf
...
...
@@ -408,10 +408,8 @@ filled with clock() */
// Setup here (temporarily?) the gradient and associated parameters.
// This should be ultimately configurable by the user as program exec. flags.
Gradientparameters
gradientpars
;
gradientpars
.
tolerance
=
1.e-6
;
gradientpars
.
max_num_of_iters
=
4
;
// Same as Solis-Wetts local search
gradientpars
.
max_num_of_iters
=
40
;
// Same as Solis-Wetts local search
gradientpars
.
alpha
=
0.000001
f
;
//0.001f; // TODO: find out why 0.001f, 0.0001f (100 runs, 500 popsize) throws segmentation fault
gradientpars
.
h
=
0.0001
f
;
// Set minimum values for input conformation (translation genes as x, y, z)
for
(
unsigned
int
gene_cnt
=
0
;
gene_cnt
<
3
;
gene_cnt
++
)
{
...
...
@@ -440,8 +438,10 @@ filled with clock() */
blocksPerGridForEachGradMinimizerEntity = dockpars.num_of_lsentities*mypars->num_of_runs;
*/
///*
// test, only one entity per reach run, undergoes gradient minimization
blocksPerGridForEachGradMinimizerEntity
=
mypars
->
num_of_runs
;
//*/
clock_start_docking
=
clock
();
...
...
@@ -558,6 +558,8 @@ filled with clock() */
#endif
// End of Kernel4
#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
);
...
...
@@ -600,12 +602,11 @@ filled with clock() */
setKernelArg
(
kernel3
,
38
,
sizeof
(
mem_ref_orientation_quats_const
),
&
mem_ref_orientation_quats_const
);
kernel3_gxsize
=
blocksPerGridForEachLSEntity
*
threadsPerBlock
;
kernel3_lxsize
=
threadsPerBlock
;
#ifdef DOCK_DEBUG
printf
(
"%-25s %10s %8u %10s %4u
\n
"
,
"K_LOCAL_SEARCH: "
,
"gSize: "
,
kernel3_gxsize
,
"lSize: "
,
kernel3_lxsize
);
fflush
(
stdout
);
#endif
#ifdef DOCK_DEBUG
printf
(
"%-25s %10s %8u %10s %4u
\n
"
,
"K_LOCAL_SEARCH: "
,
"gSize: "
,
kernel3_gxsize
,
"lSize: "
,
kernel3_lxsize
);
fflush
(
stdout
);
#endif
// End of Kernel3
#else
// Kernel5
setKernelArg
(
kernel5
,
0
,
sizeof
(
dockpars
.
num_of_atoms
),
&
dockpars
.
num_of_atoms
);
setKernelArg
(
kernel5
,
1
,
sizeof
(
dockpars
.
num_of_atypes
),
&
dockpars
.
num_of_atypes
);
...
...
@@ -620,43 +621,42 @@ filled with clock() */
setKernelArg
(
kernel5
,
10
,
sizeof
(
dockpars
.
coeff_desolv
),
&
dockpars
.
coeff_desolv
);
setKernelArg
(
kernel5
,
11
,
sizeof
(
mem_dockpars_conformations_next
),
&
mem_dockpars_conformations_next
);
setKernelArg
(
kernel5
,
12
,
sizeof
(
mem_dockpars_energies_next
),
&
mem_dockpars_energies_next
);
setKernelArg
(
kernel5
,
13
,
sizeof
(
mem_dockpars_prng_states
),
&
mem_dockpars_prng_states
);
setKernelArg
(
kernel5
,
14
,
sizeof
(
dockpars
.
pop_size
),
&
dockpars
.
pop_size
);
setKernelArg
(
kernel5
,
15
,
sizeof
(
dockpars
.
num_of_genes
),
&
dockpars
.
num_of_genes
);
setKernelArg
(
kernel5
,
16
,
sizeof
(
dockpars
.
lsearch_rate
),
&
dockpars
.
lsearch_rate
);
setKernelArg
(
kernel5
,
17
,
sizeof
(
dockpars
.
num_of_lsentities
),
&
dockpars
.
num_of_lsentities
);
setKernelArg
(
kernel5
,
18
,
sizeof
(
dockpars
.
qasp
),
&
dockpars
.
qasp
);
setKernelArg
(
kernel5
,
19
,
sizeof
(
mem_atom_charges_const
),
&
mem_atom_charges_const
);
setKernelArg
(
kernel5
,
20
,
sizeof
(
mem_atom_types_const
),
&
mem_atom_types_const
);
setKernelArg
(
kernel5
,
21
,
sizeof
(
mem_intraE_contributors_const
),
&
mem_intraE_contributors_const
);
setKernelArg
(
kernel5
,
22
,
sizeof
(
mem_VWpars_AC_const
),
&
mem_VWpars_AC_const
);
setKernelArg
(
kernel5
,
23
,
sizeof
(
mem_VWpars_BD_const
),
&
mem_VWpars_BD_const
);
setKernelArg
(
kernel5
,
24
,
sizeof
(
mem_dspars_S_const
),
&
mem_dspars_S_const
);
setKernelArg
(
kernel5
,
25
,
sizeof
(
mem_dspars_V_const
),
&
mem_dspars_V_const
);
setKernelArg
(
kernel5
,
26
,
sizeof
(
mem_rotlist_const
),
&
mem_rotlist_const
);
setKernelArg
(
kernel5
,
27
,
sizeof
(
mem_ref_coords_x_const
),
&
mem_ref_coords_x_const
);
setKernelArg
(
kernel5
,
28
,
sizeof
(
mem_ref_coords_y_const
),
&
mem_ref_coords_y_const
);
setKernelArg
(
kernel5
,
29
,
sizeof
(
mem_ref_coords_z_const
),
&
mem_ref_coords_z_const
);
setKernelArg
(
kernel5
,
30
,
sizeof
(
mem_rotbonds_moving_vectors_const
),
&
mem_rotbonds_moving_vectors_const
);
setKernelArg
(
kernel5
,
31
,
sizeof
(
mem_rotbonds_unit_vectors_const
),
&
mem_rotbonds_unit_vectors_const
);
setKernelArg
(
kernel5
,
32
,
sizeof
(
mem_ref_orientation_quats_const
),
&
mem_ref_orientation_quats_const
);
setKernelArg
(
kernel5
,
33
,
sizeof
(
mem_rotbonds_const
),
&
mem_rotbonds_const
);
setKernelArg
(
kernel5
,
34
,
sizeof
(
mem_rotbonds_atoms_const
),
&
mem_rotbonds_atoms_const
);
setKernelArg
(
kernel5
,
35
,
sizeof
(
mem_num_rotating_atoms_per_rotbond_const
),
&
mem_num_rotating_atoms_per_rotbond_const
);
setKernelArg
(
kernel5
,
13
,
sizeof
(
mem_dockpars_evals_of_new_entities
),
&
mem_dockpars_evals_of_new_entities
);
setKernelArg
(
kernel5
,
14
,
sizeof
(
mem_dockpars_prng_states
),
&
mem_dockpars_prng_states
);
setKernelArg
(
kernel5
,
15
,
sizeof
(
dockpars
.
pop_size
),
&
dockpars
.
pop_size
);
setKernelArg
(
kernel5
,
16
,
sizeof
(
dockpars
.
num_of_genes
),
&
dockpars
.
num_of_genes
);
setKernelArg
(
kernel5
,
17
,
sizeof
(
dockpars
.
lsearch_rate
),
&
dockpars
.
lsearch_rate
);
setKernelArg
(
kernel5
,
18
,
sizeof
(
dockpars
.
num_of_lsentities
),
&
dockpars
.
num_of_lsentities
);
setKernelArg
(
kernel5
,
19
,
sizeof
(
dockpars
.
qasp
),
&
dockpars
.
qasp
);
setKernelArg
(
kernel5
,
20
,
sizeof
(
mem_atom_charges_const
),
&
mem_atom_charges_const
);
setKernelArg
(
kernel5
,
21
,
sizeof
(
mem_atom_types_const
),
&
mem_atom_types_const
);
setKernelArg
(
kernel5
,
22
,
sizeof
(
mem_intraE_contributors_const
),
&
mem_intraE_contributors_const
);
setKernelArg
(
kernel5
,
23
,
sizeof
(
mem_VWpars_AC_const
),
&
mem_VWpars_AC_const
);
setKernelArg
(
kernel5
,
24
,
sizeof
(
mem_VWpars_BD_const
),
&
mem_VWpars_BD_const
);
setKernelArg
(
kernel5
,
25
,
sizeof
(
mem_dspars_S_const
),
&
mem_dspars_S_const
);
setKernelArg
(
kernel5
,
26
,
sizeof
(
mem_dspars_V_const
),
&
mem_dspars_V_const
);
setKernelArg
(
kernel5
,
27
,
sizeof
(
mem_rotlist_const
),
&
mem_rotlist_const
);
setKernelArg
(
kernel5
,
28
,
sizeof
(
mem_ref_coords_x_const
),
&
mem_ref_coords_x_const
);
setKernelArg
(
kernel5
,
29
,
sizeof
(
mem_ref_coords_y_const
),
&
mem_ref_coords_y_const
);
setKernelArg
(
kernel5
,
30
,
sizeof
(
mem_ref_coords_z_const
),
&
mem_ref_coords_z_const
);
setKernelArg
(
kernel5
,
31
,
sizeof
(
mem_rotbonds_moving_vectors_const
),
&
mem_rotbonds_moving_vectors_const
);
setKernelArg
(
kernel5
,
32
,
sizeof
(
mem_rotbonds_unit_vectors_const
),
&
mem_rotbonds_unit_vectors_const
);
setKernelArg
(
kernel5
,
33
,
sizeof
(
mem_ref_orientation_quats_const
),
&
mem_ref_orientation_quats_const
);
setKernelArg
(
kernel5
,
34
,
sizeof
(
mem_rotbonds_const
),
&
mem_rotbonds_const
);
setKernelArg
(
kernel5
,
35
,
sizeof
(
mem_rotbonds_atoms_const
),
&
mem_rotbonds_atoms_const
);
setKernelArg
(
kernel5
,
36
,
sizeof
(
mem_num_rotating_atoms_per_rotbond_const
),
&
mem_num_rotating_atoms_per_rotbond_const
);
// Specific gradient-minimizer args
setKernelArg
(
kernel5
,
36
,
sizeof
(
gradientpars
.
tolerance
),
&
gradientpars
.
tolerance
);
setKernelArg
(
kernel5
,
37
,
sizeof
(
gradientpars
.
max_num_of_iters
),
&
gradientpars
.
max_num_of_iters
);
setKernelArg
(
kernel5
,
38
,
sizeof
(
gradientpars
.
alpha
),
&
gradientpars
.
alpha
);
setKernelArg
(
kernel5
,
39
,
sizeof
(
gradientpars
.
h
),
&
gradientpars
.
h
);
setKernelArg
(
kernel5
,
40
,
sizeof
(
mem_gradpars_conformation_min_perturbation
),
&
mem_gradpars_conformation_min_perturbation
);
setKernelArg
(
kernel5
,
39
,
sizeof
(
mem_gradpars_conformation_min_perturbation
),
&
mem_gradpars_conformation_min_perturbation
);
kernel5_gxsize
=
blocksPerGridForEachGradMinimizerEntity
*
threadsPerBlock
;
kernel5_lxsize
=
threadsPerBlock
;