source: CIVL/examples/translation/openclversion2/life.cvl@ 054e457

1.23 2.0 acw/focus-triggers main test-branch
Last change on this file since 054e457 was 4246e94, checked in by Jacob Trieu <fuufusuu@…>, 12 years ago

git-svn-id: svn://vsl.cis.udel.edu/civl/trunk@1204 fb995dde-84ed-4084-dfe6-e5aef3e2452c

  • Property mode set to 100644
File size: 4.9 KB
Line 
1#include "setup.cvl"
2#include <stdio.h>
3#include <stdlib.h>
4#include <string.h>
5
6//THIS ASSUMES THE PROGRAMMER HAS ALREADY RAN THE ORIGINAL OPENCL CODE AND IS NOT RESPONSIBLE FOR CHECKING CORRECT USAGE OF METHODS AS OF NOW
7//Taking a break from this example, will redo later
8
9$input int NUM_DEVICES;
10$input int MAX_NUM_DEVICES;
11$input int CL_DEVICE_MAX_WORK_GROUP_SIZE;
12$input int LOCAL;
13$assume 0 < NUM_DEVICES && NUM_DEVICES < MAX_NUM_DEVICES;
14$gbarrier gbarrier = $gbarrier_create($here, NUM_DEVICES);
15//struct goes here
16
17
18typedef struct
19{
20 int device_id;
21 int workgroup;
22 int global_id;
23 int local_id;
24
25 //kernel variables
26 bool * input;
27 bool *output;
28 int height;
29 int width;
30}kernelCL;
31
32//kernel goes here
33void life(int workgroup,
34 int global_id,
35 int local_id,
36 bool * input,
37 bool* output,
38 int height,
39 int width)
40{
41 int i = global_id;
42 int rowUp = i - width;
43 int rowDown = i + width;
44 bool outOfBounds = (i < width);
45 outOfBounds |= (i > (width * (height-1)));
46 outOfBounds |= (i % width == 0);
47 outOfBounds |= (i % width == width-1);
48 if (outOfBounds)
49 {
50 output[i] = false; return;
51 }
52 int neighbours = input[rowUp-1] + input[rowUp] + input[rowUp+1];
53 neighbours += input[i-1] + input[i+1];
54 neighbours += input[rowDown-1] + input[rowDown] + input[rowDown+1];
55 if (neighbours == 3 || (input[i] && neighbours == 2))
56 {
57 output[i] = true;
58 }
59 else
60 {
61 output[i] = false;
62 }
63}
64
65void workfunc(size_t local, size_t global, kernelCL param)
66{
67
68 for(int i = local * param.workgroup; i < local * param.workgroup + local; i++)
69 {
70 param.local_id = i % local;
71 param.global_id = i;
72 //printf("My workgroup id is %d, my global id is %d, my local id is %d\n", param.workgroup, param.global_id, param.local_id);
73 life(param.workgroup, param.global_id, param.local_id, param.input, param.output, param.height, param.width);
74 }
75
76}
77
78void fail(const char *message)
79{
80 fprintf(stderr, "%s\n", message);
81 exit(1);
82}
83
84/*
85Note that the original lines were "__kernel void square( \n" \
86" __global int* input, \n" \
87" __global int* output, \n" \
88" const unsigned int count) \n" \
89
90Any parser must take note of and don't input \n, "", or \ as is
91__global int * input, __global int * output, int count;
92*/
93
94//variables from __kernel come here
95bool * input;
96bool *output;
97/*
98int height;
99int width;
100*/
101//from the code before
102// The board
103static int width = 128;
104static int height = 64;
105static size_t board_size = width * height;
106static bool board[board_size];
107static int queueCL[];
108
109void printBoard(void)
110{
111 unsigned i = 0;
112
113 for (unsigned y=0 ; y<height ; y++)
114 {
115 for (unsigned x=0 ; x<width ; x++)
116 {
117 putc(board[i++] ? '*' : ' ', stdout);
118 }
119 puts("");
120 }
121 puts("\n");
122}
123
124void createQueue(int queueCL[])
125{
126 int device_id[NUM_DEVICES];
127 //put device_ids
128 for(int i = 0; i < NUM_DEVICES; i++)
129 {
130 device_id[i] = i;
131 }
132 //clCreateCommandQueue
133 queueCL[0] = device_id[0];
134 //presents all sorts of strange problems
135 //How does the parser know to make this pass in queueCL and not (void)?
136
137}
138
139void prepareKernel(void)
140{
141 input = (int *)malloc(sizeof(board));
142 output = (int *)malloc(sizeof(board));
143
144 kernelCL param[board_size/local];
145 for(int i = 0; i < board_size/local; i++)
146 {
147 //Also picks the device to be used
148 param[i].device_id = queueCL[0];
149 //other parts of the struct
150 param[i].input = input;
151 param[i].output = output;
152 param[i].height = height;
153 param[i].width = width;
154 }
155}
156
157void runGame(int iterations)
158{
159 if (iterations == 0) {return;}
160 for(int i = 0; i < board_size/local; i++)
161 {
162 param[i].input = board;
163 }
164 for(int i = 0; i < iterations; i++)
165 {
166 for(int i = 0; i < board_size/local; i++)
167 {
168 param[i].workgroup = i;
169 //procs[i] = $spawn square(param[i].global_id, param[i].input, param[i].output, param[i].count);
170 procs[i] = $spawn workfunc(local, global, param[i]);
171 }
172
173 for(int i = 0; i < board_size/local; i++)
174 {
175 $wait(procs[i]);
176 }
177 }
178}
179
180int main(int argc, char** argv)
181{
182 bool * input;
183 bool *output;
184 int height;
185 int width;
186 queueCL[1];
187 int local = LOCAL;
188 $assert(board_size%local == 0);
189 //How do we even know what size to use?
190
191 //get the number from clGetDeviceIDs 3rd parameter
192 $assert (LOCAL < CL_DEVICE_MAX_WORK_GROUP_SIZE);
193
194
195 srandomdev();
196 for(unsigned int i=0 ; i<board_size ; i++)
197 board[i] = random() > (INT_MAX / 2);
198
199 createQueue(queueCL);
200
201 kernel = createKernelFromSource(device_id, context,
202 kernel_source, "life");
203
204 prepareKernel();
205
206 printBoard();
207 int iterations = 0;
208 do
209 {
210 printf("Run for how many iterations (0 to exit)? ");
211 scanf("%d", &iterations);
212 runGame(iterations);
213 printBoard();
214 } while(iterations > 0);
215
216 return 0;
217}
Note: See TracBrowser for help on using the repository browser.