kernel3.cl 15.3 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
perform_LS(	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
38
39

		#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
40
41
42
	        int    dockpars_rotbondlist_length,
		float  dockpars_coeff_elec,
		float  dockpars_coeff_desolv,
Leonardo Solis's avatar
Leonardo Solis committed
43
44

		#if defined (RESTRICT_ARGS)
Leonardo Solis's avatar
Leonardo Solis committed
45
46
47
48
  __global float* restrict dockpars_conformations_next,
  __global float* restrict dockpars_energies_next,
  __global int*   restrict dockpars_evals_of_new_entities,
  __global unsigned int* restrict dockpars_prng_states,
Leonardo Solis's avatar
Leonardo Solis committed
49
		#else
Leonardo Solis's avatar
Leonardo Solis committed
50
51
52
53
  __global float* dockpars_conformations_next,
  __global float* dockpars_energies_next,
  __global int*   dockpars_evals_of_new_entities,
  __global unsigned int* dockpars_prng_states,
Leonardo Solis's avatar
Leonardo Solis committed
54
55
		#endif

Leonardo Solis's avatar
Leonardo Solis committed
56
57
58
59
60
61
62
63
64
65
66
67
		int    dockpars_pop_size,
		int    dockpars_num_of_genes,
		float  dockpars_lsearch_rate,
		unsigned int dockpars_num_of_lsentities,
		float  dockpars_rho_lower_bound,
		float  dockpars_base_dmov_mul_sqrt3,
		float  dockpars_base_dang_mul_sqrt3,
		unsigned int dockpars_cons_limit,
		unsigned int dockpars_max_num_of_iters,
		float  dockpars_qasp,

	  __constant float* atom_charges_const,
Leonardo Solis's avatar
Leonardo Solis committed
68
          __constant char*  atom_types_const,
Leonardo Solis's avatar
Leonardo Solis committed
69
	  __constant char*  intraE_contributors_const,
Leonardo Solis's avatar
Leonardo Solis committed
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
          __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
)
//The GPU global function performs local search on the pre-defined entities of conformations_next.
//The number of blocks which should be started equals to num_of_lsentities*num_of_runs.
//This way the first num_of_lsentities entity of each population will be subjected to local search
//(and each block carries out the algorithm for one entity).
//Since the first entity is always the best one in the current population,
//it is always tested according to the ls probability, and if it not to be
//subjected to local search, the entity with ID num_of_lsentities is selected instead of the first one (with ID 0).
{
	__local float genotype_candidate[ACTUAL_GENOTYPE_LENGTH];
	__local float genotype_deviate  [ACTUAL_GENOTYPE_LENGTH];
	__local float genotype_bias     [ACTUAL_GENOTYPE_LENGTH];
Leonardo Solis's avatar
Leonardo Solis committed
93
        __local float rho;
Leonardo Solis's avatar
Leonardo Solis committed
94
95
96
97
98
99
100
101
102
103
104
105
	__local int   cons_succ;
	__local int   cons_fail;
	__local int   iteration_cnt;
	__local float candidate_energy;
	__local int   evaluation_cnt;
	int gene_counter;

	__local float offspring_genotype[ACTUAL_GENOTYPE_LENGTH];
	__local int run_id;
	__local int entity_id;
	__local float offspring_energy;

Leonardo Solis's avatar
Leonardo Solis committed
106
107
108
109
110
111
112
        // Some OpenCL compilers don't allow local var outside kernels
        // so this local vars are passed from a kernel
	__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
113
114
115
116
117
118
119
120
121
122
	//determining run ID and entity ID, initializing
	if (get_local_id(0) == 0)
	{
		run_id = get_group_id(0) / dockpars_num_of_lsentities;
		entity_id = get_group_id(0) % dockpars_num_of_lsentities;

		//since entity 0 is the best one due to elitism, should be subjected to random selection
		if (entity_id == 0)
			if (100.0f*gpu_randf(dockpars_prng_states) > dockpars_lsearch_rate)
				entity_id = dockpars_num_of_lsentities;	//if entity 0 is not selected according to LS rate,
Leonardo Solis's avatar
Leonardo Solis committed
123
									//choosing an other entity
Leonardo Solis's avatar
Leonardo Solis committed
124
125
126
127
128
129
130
131

		offspring_energy = dockpars_energies_next[run_id*dockpars_pop_size+entity_id];
	}

	barrier(CLK_LOCAL_MEM_FENCE);

#if defined (ASYNC_COPY)
  async_work_group_copy(offspring_genotype,
Leonardo Solis's avatar
Leonardo Solis committed
132
			dockpars_conformations_next+(run_id*dockpars_pop_size+entity_id)*GENOTYPE_LENGTH_IN_GLOBMEM,
Leonardo Solis's avatar
Leonardo Solis committed
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
                        dockpars_num_of_genes,0);
#else
	for (gene_counter=get_local_id(0);
	     gene_counter<dockpars_num_of_genes;
	     gene_counter+=NUM_OF_THREADS_PER_BLOCK)
		   offspring_genotype[gene_counter] = dockpars_conformations_next[(run_id*dockpars_pop_size+entity_id)*GENOTYPE_LENGTH_IN_GLOBMEM+gene_counter];
#endif

	for (gene_counter=get_local_id(0);
	     gene_counter<dockpars_num_of_genes;
	     gene_counter += NUM_OF_THREADS_PER_BLOCK)
		   genotype_bias[gene_counter] = 0.0f;

	if (get_local_id(0) == 0) {
		rho = 1.0f;
		cons_succ = 0;
		cons_fail = 0;
		iteration_cnt = 0;
		evaluation_cnt = 0;
	}

	barrier(CLK_LOCAL_MEM_FENCE);

	while ((iteration_cnt < dockpars_max_num_of_iters) && (rho > dockpars_rho_lower_bound))
	{
		//new random deviate
		for (gene_counter=get_local_id(0);
		     gene_counter<dockpars_num_of_genes;
		     gene_counter += NUM_OF_THREADS_PER_BLOCK)
		{
			genotype_deviate[gene_counter] = rho*(2*gpu_randf(dockpars_prng_states)-1);

165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
// -------------------------------------------------------------------
// L30nardoSV
// Replacing rotation genes: from spherical space to Shoemake space
// gene [0:2]: translation -> kept as original x, y, z
// gene [3:5]: rotation    -> transformed into Shoemake (u1: adimensional, u2&u3: sexagesimal)
// gene [6:N]: torsions	   -> kept as original angles	(all in sexagesimal)

// Shoemake ranges:
// u1: [0, 1]
// u2: [0: 2PI] or [0: 360]

// Random generator in the host is changed:
// LCG (original, myrand()) -> CPP std (rand())
// -------------------------------------------------------------------

// Original code is commented out
/*
			if (gene_counter < 3)
				genotype_deviate[gene_counter] *= dockpars_base_dmov_mul_sqrt3;
			else
				genotype_deviate[gene_counter] *= dockpars_base_dang_mul_sqrt3;
*/

Leonardo Solis's avatar
Leonardo Solis committed
188
189
			if (gene_counter < 3)
				genotype_deviate[gene_counter] *= dockpars_base_dmov_mul_sqrt3;
190
191
192
193
			else if (gene_counter == 3) // u1, FIXME: hardcoded
				genotype_deviate[gene_counter] *=  0.2f;
			else if (gene_counter < 6) // u2&u3, FIXME: harcoded
				genotype_deviate[gene_counter] *= 90.0f;
Leonardo Solis's avatar
Leonardo Solis committed
194
195
			else
				genotype_deviate[gene_counter] *= dockpars_base_dang_mul_sqrt3;
196

Leonardo Solis's avatar
Leonardo Solis committed
197
198
199
200
201
		}

		//generating new genotype candidate
		for (gene_counter=get_local_id(0);
		     gene_counter<dockpars_num_of_genes;
202
		     gene_counter += NUM_OF_THREADS_PER_BLOCK) {
Leonardo Solis's avatar
Leonardo Solis committed
203
204
			   genotype_candidate[gene_counter] = offspring_genotype[gene_counter] + genotype_deviate[gene_counter] + genotype_bias[gene_counter];

205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
///*
			if (gene_counter == 3) {// u1, FIXME: hardcoded
				if (genotype_candidate[gene_counter] > 1.0f) { // clamp it
					offspring_genotype[gene_counter] = 0.9f;
				}
				if (genotype_candidate[gene_counter] < 0.0f) { // clamp it
					offspring_genotype[gene_counter] = 0.1f;
				}
				//printf("LS positive - genotype_candidate[gene_counter]: %f\n", genotype_candidate[gene_counter]);
			}
//*/

		}


Leonardo Solis's avatar
Leonardo Solis committed
220
221
222
223
224
		//evaluating candidate
		barrier(CLK_LOCAL_MEM_FENCE);

		// ==================================================================
		gpu_calc_energy(dockpars_rotbondlist_length,
Leonardo Solis's avatar
Leonardo Solis committed
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
				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_candidate,
				&candidate_energy,
				&run_id,
				// Some OpenCL compilers don't allow local var outside kernels
				// so this local vars are passed from a kernel
				calc_coords_x,
				calc_coords_y,
				calc_coords_z,
				partial_energies,

				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,
				ref_orientation_quats_const);
Leonardo Solis's avatar
Leonardo Solis committed
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
		// =================================================================

		if (get_local_id(0) == 0)
			evaluation_cnt++;

		barrier(CLK_LOCAL_MEM_FENCE);

		if (candidate_energy < offspring_energy)	//if candidate is better, success
		{
			for (gene_counter=get_local_id(0);
			     gene_counter<dockpars_num_of_genes;
			     gene_counter += NUM_OF_THREADS_PER_BLOCK)
			{
				//updating offspring_genotype
				offspring_genotype[gene_counter] = genotype_candidate[gene_counter];

				//updating genotype_bias
				genotype_bias[gene_counter] = 0.6f*genotype_bias[gene_counter] + 0.4f*genotype_deviate[gene_counter];
			}

			//thread 0 will overwrite the shared variables
			//used in the previous if condition,
			//all threads have to be after if
			barrier(CLK_LOCAL_MEM_FENCE);

			if (get_local_id(0) == 0)
			{
				offspring_energy = candidate_energy;
				cons_succ++;
				cons_fail = 0;
			}
		}
		else	//if candidate is worser, check the opposite direction
		{
			//generating the other genotype candidate
			for (gene_counter=get_local_id(0);
			     gene_counter<dockpars_num_of_genes;
297
			     gene_counter += NUM_OF_THREADS_PER_BLOCK) {
Leonardo Solis's avatar
Leonardo Solis committed
298
299
				   genotype_candidate[gene_counter] = offspring_genotype[gene_counter] - genotype_deviate[gene_counter] - genotype_bias[gene_counter];

300
301
302
303
304
305
306
307
308
309
310
311
312
313
///*
				if (gene_counter == 3) {// u1, FIXME: hardcoded
					if (genotype_candidate[gene_counter] > 1.0f) { // clamp it
						offspring_genotype[gene_counter] = 0.9f;
					}
					if (genotype_candidate[gene_counter] < 0.0f) { // clamp it
						offspring_genotype[gene_counter] = 0.1f;
					}
					//printf("LS negative - genotype_candidate[gene_counter]: %f\n", genotype_candidate[gene_counter]);
				}
//*/

			}

Leonardo Solis's avatar
Leonardo Solis committed
314
315
316
317
318
			//evaluating candidate
			barrier(CLK_LOCAL_MEM_FENCE);

			// =================================================================
			gpu_calc_energy(dockpars_rotbondlist_length,
Leonardo Solis's avatar
Leonardo Solis committed
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
					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_candidate,
					&candidate_energy,
					&run_id,
		                        // Some OpenCL compilers don't allow local var outside kernels
					// so this local vars are passed from a kernel
					calc_coords_x,
					calc_coords_y,
					calc_coords_z,
					partial_energies,

					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,
					ref_orientation_quats_const);
Leonardo Solis's avatar
Leonardo Solis committed
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
			// =================================================================

			if (get_local_id(0) == 0)
				evaluation_cnt++;

			barrier(CLK_LOCAL_MEM_FENCE);

			if (candidate_energy < offspring_energy)//if candidate is better, success
			{
				for (gene_counter=get_local_id(0);
				     gene_counter<dockpars_num_of_genes;
			       gene_counter += NUM_OF_THREADS_PER_BLOCK)
				{
					//updating offspring_genotype
					offspring_genotype[gene_counter] = genotype_candidate[gene_counter];

					//updating genotype_bias
					genotype_bias[gene_counter] = 0.6f*genotype_bias[gene_counter] - 0.4f*genotype_deviate[gene_counter];
				}

				//thread 0 will overwrite the shared variables
				//used in the previous if condition,
				//all threads have to be after if
				barrier(CLK_LOCAL_MEM_FENCE);

				if (get_local_id(0) == 0)
				{
					offspring_energy = candidate_energy;
					cons_succ++;
					cons_fail = 0;
				}
			}
			else	//failure in both directions
			{
				for (gene_counter=get_local_id(0);
				     gene_counter<dockpars_num_of_genes;
				     gene_counter += NUM_OF_THREADS_PER_BLOCK)
					   //updating genotype_bias
					   genotype_bias[gene_counter] = 0.5f*genotype_bias[gene_counter];

				if (get_local_id(0) == 0)
				{
					cons_succ = 0;
					cons_fail++;
				}
			}
		}

		//changing rho if needed
		if (get_local_id(0) == 0)
		{
			iteration_cnt++;

			if (cons_succ >= dockpars_cons_limit)
			{
				rho *= LS_EXP_FACTOR;
				cons_succ = 0;
			}
			else
				if (cons_fail >= dockpars_cons_limit)
				{
					rho *= LS_CONT_FACTOR;
					cons_fail = 0;
				}
		}
		barrier(CLK_LOCAL_MEM_FENCE);
	}

	//updating eval counter and energy
	if (get_local_id(0) == 0)
	{
		dockpars_evals_of_new_entities[run_id*dockpars_pop_size+entity_id] += evaluation_cnt;
		dockpars_energies_next[run_id*dockpars_pop_size+entity_id] = offspring_energy;
	}

	//mapping angles
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
// -------------------------------------------------------------------
// L30nardoSV
// Replacing rotation genes: from spherical space to Shoemake space
// gene [0:2]: translation -> kept as original x, y, z
// gene [3:5]: rotation    -> transformed into Shoemake (u1: adimensional, u2&u3: sexagesimal)
// gene [6:N]: torsions	   -> kept as original angles	(all in sexagesimal)

// Shoemake ranges:
// u1: [0, 1]
// u2: [0: 2PI] or [0: 360]

// Random generator in the host is changed:
// LCG (original, myrand()) -> CPP std (rand())
// -------------------------------------------------------------------
/*
Leonardo Solis's avatar
Leonardo Solis committed
445
446
447
448
449
	for (gene_counter=get_local_id(0);
	     gene_counter<dockpars_num_of_genes;
	     gene_counter+=NUM_OF_THREADS_PER_BLOCK)
		   if (gene_counter >=  3)
			    map_angle(&(offspring_genotype[gene_counter]));
450
451
452
453
454
455
456
*/
	for (gene_counter=get_local_id(0);
	     gene_counter<dockpars_num_of_genes;
	     gene_counter+=NUM_OF_THREADS_PER_BLOCK)
		   if (gene_counter >=  4)
			    map_angle(&(offspring_genotype[gene_counter]));

Leonardo Solis's avatar
Leonardo Solis committed
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471

	//updating old offspring in population
	barrier(CLK_LOCAL_MEM_FENCE);

#if defined (ASYNC_COPY)
  async_work_group_copy(dockpars_conformations_next+(run_id*dockpars_pop_size+entity_id)*GENOTYPE_LENGTH_IN_GLOBMEM,
                        offspring_genotype,
                        dockpars_num_of_genes,0);
#else
	for (gene_counter=get_local_id(0);
	     gene_counter<dockpars_num_of_genes;
       gene_counter+=NUM_OF_THREADS_PER_BLOCK)
		   dockpars_conformations_next[(run_id*dockpars_pop_size+entity_id)*GENOTYPE_LENGTH_IN_GLOBMEM+gene_counter] = offspring_genotype[gene_counter];
#endif
}