kernel1.cl 6.36 KB
Newer Older
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
/*

OCLADock, an OpenCL implementation of AutoDock 4.2 running a Lamarckian Genetic Algorithm
Copyright (C) 2017 TU Darmstadt, Embedded Systems and Applications Group, Germany. All rights reserved.

AutoDock is a Trade Mark of the Scripps Research Institute.

This program is free software; you can redistribute it and/or
modify it under the terms of the GNU General Public License
as published by the Free Software Foundation; either version 2
of the License, or (at your option) any later version.

This program is distributed in the hope that it will be useful,
but WITHOUT ANY WARRANTY; without even the implied warranty of
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
GNU General Public License for more details.

You should have received a copy of the GNU General Public License
along with this program; if not, write to the Free Software
Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301, USA.

*/


Leonardo Solis's avatar
Leonardo Solis committed
25
__kernel void __attribute__ ((reqd_work_group_size(NUM_OF_THREADS_PER_BLOCK,1,1)))
Leonardo Solis's avatar
Leonardo Solis committed
26
27
28
29
30
31
32
gpu_calc_initpop(	char   dockpars_num_of_atoms,
			char   dockpars_num_of_atypes,
			int    dockpars_num_of_intraE_contributors,
			char   dockpars_gridsize_x,
			char   dockpars_gridsize_y,
			char   dockpars_gridsize_z,
			float  dockpars_grid_spacing,
Leonardo Solis's avatar
Leonardo Solis committed
33
34
35
36
37
		#if defined (RESTRICT_ARGS)
		  __global const float* restrict dockpars_fgrids, // cannot be allocated in __constant (too large)
		#else
		  __global const float* dockpars_fgrids,          // cannot be allocated in __constant (too large)
		#endif
Leonardo Solis's avatar
Leonardo Solis committed
38
39
40
			int    dockpars_rotbondlist_length,
			float  dockpars_coeff_elec,
			float  dockpars_coeff_desolv,
Leonardo Solis's avatar
Leonardo Solis committed
41
42

		#if defined (RESTRICT_ARGS)
Leonardo Solis's avatar
Leonardo Solis committed
43
44
			__global const float* restrict dockpars_conformations_current,
			__global float* restrict dockpars_energies_current,
Leonardo Solis's avatar
Leonardo Solis committed
45
46
		        __global int*   restrict dockpars_evals_of_new_entities,
		#else
Leonardo Solis's avatar
Leonardo Solis committed
47
48
			__global const float* dockpars_conformations_current,
			__global float* dockpars_energies_current,
Leonardo Solis's avatar
Leonardo Solis committed
49
50
51
		        __global int*   dockpars_evals_of_new_entities,
		#endif

Leonardo Solis's avatar
Leonardo Solis committed
52
53
			int    dockpars_pop_size,
			float  dockpars_qasp,
Leonardo Solis's avatar
Leonardo Solis committed
54

Leonardo Solis's avatar
Leonardo Solis committed
55
	   __constant float* atom_charges_const,
Leonardo Solis's avatar
Leonardo Solis committed
56
           __constant char*  atom_types_const,
Leonardo Solis's avatar
Leonardo Solis committed
57
	   __constant char*  intraE_contributors_const,
Leonardo Solis's avatar
Leonardo Solis committed
58
59
60
61
62
63
64
65
66
67
68
69
70
71
           __constant float* VWpars_AC_const,
           __constant float* VWpars_BD_const,
           __constant float* dspars_S_const,
           __constant float* dspars_V_const,
           __constant int*   rotlist_const,
           __constant float* ref_coords_x_const,
           __constant float* ref_coords_y_const,
           __constant float* ref_coords_z_const,
           __constant float* rotbonds_moving_vectors_const,
           __constant float* rotbonds_unit_vectors_const,
           __constant float* ref_orientation_quats_const
){
	__local float  genotype[GENOTYPE_LENGTH_IN_GLOBMEM];
	__local float  energy;
Leonardo Solis's avatar
Leonardo Solis committed
72
73
	__local int    run_id;

74
75
76
77
        // 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.
Leonardo Solis's avatar
Leonardo Solis committed
78
79
80
81
	__local float calc_coords_x[MAX_NUM_OF_ATOMS];
	__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];
Leonardo Solis's avatar
Leonardo Solis committed
82
83

	event_t ev = async_work_group_copy(genotype,
Leonardo Solis's avatar
Leonardo Solis committed
84
85
			                   dockpars_conformations_current + GENOTYPE_LENGTH_IN_GLOBMEM*get_group_id(0),
			                   GENOTYPE_LENGTH_IN_GLOBMEM, 0);
Leonardo Solis's avatar
Leonardo Solis committed
86
87
88

	wait_group_events(1,&ev);

Leonardo Solis's avatar
Leonardo Solis committed
89
90
	//determining run ID
	if (get_local_id(0) == 0)
Leonardo Solis's avatar
Leonardo Solis committed
91
92
		run_id = get_group_id(0) / dockpars_pop_size;

93
94
95
96
97
98
99
100
101
	// -------------------------------------------------------------------
	// L30nardoSV
	// Calculate gradients (forces) for intermolecular energy
	// Derived from autodockdev/maps.py
	// -------------------------------------------------------------------
	// 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.
102

103
104
105
106
107
	// Disable gradient calculation for this kernel
	__local bool  is_enabled_gradient_calc;
	if (get_local_id(0) == 0) {
		is_enabled_gradient_calc = false;
	}
108
109
110
111
112
113
114
115
116

	// Variables to store gradient of 
	// the intermolecular energy per each ligand atom
	__local float gradient_inter_x[MAX_NUM_OF_ATOMS];
	__local float gradient_inter_y[MAX_NUM_OF_ATOMS];
	__local float gradient_inter_z[MAX_NUM_OF_ATOMS];
	
	// Final gradient resulting out of gradient calculation
	__local float gradient_genotype[GENOTYPE_LENGTH_IN_GLOBMEM];
117
118
	// -------------------------------------------------------------------

Leonardo Solis's avatar
Leonardo Solis committed
119
	// =============================================================
Leonardo Solis's avatar
Leonardo Solis committed
120
	// WARNING: only energy of work-item=0 will be valid
Leonardo Solis's avatar
Leonardo Solis committed
121
	gpu_calc_energy(dockpars_rotbondlist_length,
Leonardo Solis's avatar
Leonardo Solis committed
122
123
124
125
126
127
128
129
130
131
132
133
134
135
			dockpars_num_of_atoms,
			dockpars_gridsize_x,
			dockpars_gridsize_y,
			dockpars_gridsize_z,
			dockpars_fgrids,
			dockpars_num_of_atypes,
			dockpars_num_of_intraE_contributors,
			dockpars_grid_spacing,
			dockpars_coeff_elec,
			dockpars_qasp,
			dockpars_coeff_desolv,
			genotype,
			&energy,
			&run_id,
136
137
138
139
			// 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.
Leonardo Solis's avatar
Leonardo Solis committed
140
141
142
143
			calc_coords_x,
			calc_coords_y,
			calc_coords_z,
			partial_energies,
Leonardo Solis's avatar
Leonardo Solis committed
144

Leonardo Solis's avatar
Leonardo Solis committed
145
146
147
148
149
150
151
152
153
154
155
156
157
	                atom_charges_const,
		        atom_types_const,
			intraE_contributors_const,
			VWpars_AC_const,
			VWpars_BD_const,
			dspars_S_const,
			dspars_V_const,
			rotlist_const,
			ref_coords_x_const,
			ref_coords_y_const,
			ref_coords_z_const,
			rotbonds_moving_vectors_const,
			rotbonds_unit_vectors_const,
158
159
160
161
162
163
164
165
166
167
168
169
			ref_orientation_quats_const

		 	// -------------------------------------------------------------------
		 	// L30nardoSV
		 	// Gradient-related arguments
		 	// Calculate gradients (forces) for intermolecular energy
		 	// Derived from autodockdev/maps.py
		 	// -------------------------------------------------------------------
			,
			&is_enabled_gradient_calc,
			gradient_inter_x,
			gradient_inter_y,
170
171
172
			gradient_inter_z,

			gradient_genotype
173
174
			);
			// -------------------------------------------------------------------
Leonardo Solis's avatar
Leonardo Solis committed
175
176
	// =============================================================

Leonardo Solis's avatar
Leonardo Solis committed
177
178
179
180
	if (get_local_id(0) == 0) {
		dockpars_energies_current[get_group_id(0)] = energy;
		dockpars_evals_of_new_entities[get_group_id(0)] = 1;
	}
Leonardo Solis's avatar
Leonardo Solis committed
181
}