kernel2.cl 2.14 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)))
26
27
gpu_sum_evals(	       
		       uint pop_size,
Leonardo Solis's avatar
Leonardo Solis committed
28
	      __global int* restrict dockpars_evals_of_new_entities,
Leonardo Solis's avatar
Leonardo Solis committed
29
              __global int* restrict evals_of_runs
Leonardo Solis's avatar
Leonardo Solis committed
30
31
32
33
34
35
36
)
//The GPU global function sums the evaluation counter states
//which are stored in evals_of_new_entities array foreach entity,
//calculates the sums for each run and stores it in evals_of_runs array.
//The number of blocks which should be started equals to num_of_runs,
//since each block performs the summation for one run.
{
37
38
39
40
41
42
43
44
45
	int entity_counter;
	int sum_evals;
	__local int partsum_evals[NUM_OF_THREADS_PER_BLOCK];

	partsum_evals[get_local_id(0)] = 0;

  	for (entity_counter = get_local_id(0);
	     entity_counter < pop_size;
	     entity_counter+= NUM_OF_THREADS_PER_BLOCK) {
lvs's avatar
lvs committed
46
		partsum_evals[get_local_id(0)] += dockpars_evals_of_new_entities[get_group_id(0)*pop_size + entity_counter];
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
	}

	barrier(CLK_LOCAL_MEM_FENCE);

	if (get_local_id(0) == 0) {
		sum_evals = partsum_evals[0];

		for (entity_counter = 1;
		     entity_counter < NUM_OF_THREADS_PER_BLOCK;
		     entity_counter++) {
			sum_evals += partsum_evals[entity_counter];
		}

		evals_of_runs[get_group_id(0)] += sum_evals;
	}
Leonardo Solis's avatar
Leonardo Solis committed
62
}