aboutsummaryrefslogtreecommitdiff
path: root/handcoded-opencl/sql9.cl
blob: 82b464d2d0a534d9ee5ac2820fae10a9d0601a6f (plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39

#define workUnits 64
#define workUnitsM1 63

typedef struct tag_my_struct {
	int4 v;
	} Row;

__kernel void x1_search_kernel(int totalRows,
							   __global Row *data,
							   __global Row *resultArray,
							   __global int *roffsetResult) {

	int i = get_global_id(0);
	size_t offset = i * (totalRows/workUnits);
	size_t endRow = (totalRows/workUnits);
	size_t roffset = offset; 
	int tmp;

	if (i == workUnitsM1) {
		endRow = (totalRows/workUnits) + (totalRows % workUnits);
	}
	do {
		if (data[offset].v.s1 == 0 || data[offset].v.s2 == 0 || data[offset].v.s3 == 0) 
			goto copy;

		goto next;
copy:
		resultArray[roffset].v = data[offset].v;

		roffset++;

next:
		offset++;
		endRow--;
	} while (endRow);

	roffsetResult[i] = roffset- (i * (totalRows/workUnits));
}