#include "setup.cvl" #include #include #include //THIS ASSUMES THE PROGRAMMER HAS ALREADY RAN THE ORIGINAL OPENCL CODE AND IS NOT RESPONSIBLE FOR CHECKING CORRECT USAGE OF METHODS AS OF NOW //Taking a break from this example, will redo later $input int NUM_DEVICES; $input int MAX_NUM_DEVICES; $input int CL_DEVICE_MAX_WORK_GROUP_SIZE; $input int LOCAL; $assume 0 < NUM_DEVICES && NUM_DEVICES < MAX_NUM_DEVICES; $gbarrier gbarrier = $gbarrier_create($here, NUM_DEVICES); //struct goes here typedef struct { int device_id; int workgroup; int global_id; int local_id; //kernel variables bool * input; bool *output; int height; int width; }kernelCL; //kernel goes here void life(int workgroup, int global_id, int local_id, bool * input, bool* output, int height, int width) { int i = global_id; int rowUp = i - width; int rowDown = i + width; bool outOfBounds = (i < width); outOfBounds |= (i > (width * (height-1))); outOfBounds |= (i % width == 0); outOfBounds |= (i % width == width-1); if (outOfBounds) { output[i] = false; return; } int neighbours = input[rowUp-1] + input[rowUp] + input[rowUp+1]; neighbours += input[i-1] + input[i+1]; neighbours += input[rowDown-1] + input[rowDown] + input[rowDown+1]; if (neighbours == 3 || (input[i] && neighbours == 2)) { output[i] = true; } else { output[i] = false; } } void workfunc(size_t local, size_t global, kernelCL param) { for(int i = local * param.workgroup; i < local * param.workgroup + local; i++) { param.local_id = i % local; param.global_id = i; //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); life(param.workgroup, param.global_id, param.local_id, param.input, param.output, param.height, param.width); } } void fail(const char *message) { fprintf(stderr, "%s\n", message); exit(1); } /* Note that the original lines were "__kernel void square( \n" \ " __global int* input, \n" \ " __global int* output, \n" \ " const unsigned int count) \n" \ Any parser must take note of and don't input \n, "", or \ as is __global int * input, __global int * output, int count; */ //variables from __kernel come here bool * input; bool *output; /* int height; int width; */ //from the code before // The board static int width = 128; static int height = 64; static size_t board_size = width * height; static bool board[board_size]; static int queueCL[]; void printBoard(void) { unsigned i = 0; for (unsigned y=0 ; y (INT_MAX / 2); createQueue(queueCL); kernel = createKernelFromSource(device_id, context, kernel_source, "life"); prepareKernel(); printBoard(); int iterations = 0; do { printf("Run for how many iterations (0 to exit)? "); scanf("%d", &iterations); runGame(iterations); printBoard(); } while(iterations > 0); return 0; }