Permalink
Switch branches/tags
Nothing to show
Find file
Fetching contributors…
Cannot retrieve contributors at this time
309 lines (285 sloc) 8.09 KB
#include <string.h>
#include <ctype.h>
#include <stdlib.h>
#include <stdio.h>
#include <CL/cl.h>
/*
* An element in the grid
*/
typedef struct element_t {
int x, y, r;
char sym;
char dir[4];
} element_t;
/*
* Read all the elements from the file and return them in an array
* will also fill parameters n_elems and grid_dim with the number of
* elements loaded and the grid dimensions
*/
element_t* load_elems(FILE *fp, int *n_elems, int *grid_dim);
/*
* Print out the grid of elements, will print spaces in empty areas
*/
void print_grid(element_t *elems, int n_elems, int grid_dim);
/*
* See if there's an element at some coordinate and get back its symbol
* if there is one, the null terminator will be returned for none
*/
char elem_at(element_t *elems, int n_elems, int x, int y);
/*
* Count the number of active elements in the simulation
*/
int active_elems(element_t *elems, int n_elems);
/*
* Read the entire content of a file and return it
* the char* must be freed later by the caller, NULL is returned
* in the case of an error, will also return the file size
*/
char* file_content(char *f_name, size_t *size);
int main(int argc, char **argv){
if (argc < 2){
printf("Usage: ./prog input.txt\n");
return 1;
}
FILE *fp = fopen(argv[1], "r");
if (!fp){
printf("Failed to open file: %s\n", argv[1]);
return 1;
}
int n_elems, grid_dim;
element_t *elems = load_elems(fp, &n_elems, &grid_dim);
if (!elems){
printf("load_elems failed\n");
fclose(fp);
return 2;
}
fclose(fp);
cl_platform_id platform;
cl_uint num;
int err = clGetPlatformIDs(1, &platform, &num);
if (err < 0){
printf("clGetPlatformIDs error: %d\n", err);
free(elems);
return err;
}
cl_device_id device;
err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
if (err < 0){
printf("clGetDeviceIDs error: %d\n", err);
free(elems);
return err;
}
char name[256];
err = clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(name), name, NULL);
if (err < 0){
printf("clGetDeviceInfo error: %d\n", err);
free(elems);
return err;
}
name[255] = '\0';
printf("Using device: %s\n", name);
cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
if (err < 0){
printf("clCreateContext error: %d\n", err);
free(elems);
return err;
}
size_t prog_size;
char *prog_src = file_content("simulation.cl", &prog_size);
cl_program program = clCreateProgramWithSource(context, 1, (const char**)&prog_src, &prog_size, &err);
if (err < 0){
printf("clCreateProgramWithSource err: %d\n", err);
clReleaseContext(context);
free(prog_src);
free(elems);
return err;
}
free(prog_src);
err = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
if (err < 0){
printf("clBuildProgram err: %d\n", err);
size_t log_size;
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
char *prog_log = (char*)calloc(log_size + 1, sizeof(char));
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, prog_log, NULL);
printf("%s\n", prog_log);
free(prog_log);
clReleaseProgram(program);
clReleaseContext(context);
free(elems);
return err;
}
cl_kernel kernel = clCreateKernel(program, "sim_elem", &err);
if (err < 0){
printf("clCreateKernel err: %d\n", err);
clReleaseProgram(program);
clReleaseContext(context);
free(elems);
return err;
}
cl_command_queue queue = clCreateCommandQueue(context, device, 0, &err);
if (err < 0){
printf("clCreateCommandQueue err: %d\n", err);
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseContext(context);
free(elems);
return err;
}
size_t elem_buf_size = n_elems * sizeof(element_t);
//The in/out element buffers
cl_mem elem_buffs[2];
for (int i = 0; i < 2; ++i){
elem_buffs[i] = clCreateBuffer(context, CL_MEM_READ_WRITE, elem_buf_size, NULL, &err);
if (err < 0){
printf("clCreateBuffer err: %d\n", err);
clReleaseCommandQueue(queue);
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseContext(context);
free(elems);
return err;
}
}
clSetKernelArg(kernel, 0, sizeof(int), &n_elems);
for (int i = 0; i < 2; ++i){
clSetKernelArg(kernel, i + 1, sizeof(cl_mem), &elem_buffs[i]);
}
int step = 0;
printf("Step %d:\n", step);
print_grid(elems, n_elems, grid_dim);
//Activate the first element and write the data
elems[0].sym = 'X';
//Write into buffer 0 and copy that to buffer 1
clEnqueueWriteBuffer(queue, elem_buffs[0], CL_FALSE, 0, elem_buf_size, elems, 0, NULL, NULL);
clEnqueueCopyBuffer(queue, elem_buffs[0], elem_buffs[1], 0, 0, elem_buf_size, 0, NULL, NULL);
/*
* Run the simulation until equilibrium is reached
* Because of how the chain reaction simulation works we really only
* need to simulate the elements activated by the previous step but by re-simulating
* all activated elements we could handle moving particles, although it's unnecessary
* for the challenge
*/
size_t active_prev = 0, active_now = 1;
do {
++step;
printf("Step %d:\n", step);
print_grid(elems, n_elems, grid_dim);
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &active_now, NULL, 0, NULL, NULL);
clEnqueueCopyBuffer(queue, elem_buffs[1], elem_buffs[0], 0, 0, elem_buf_size, 0, NULL, NULL);
//Now read the result back out, completely destroying any performance gains =P
clEnqueueReadBuffer(queue, elem_buffs[1], CL_TRUE, 0, elem_buf_size, elems, 0, NULL, NULL);
active_prev = active_now;
active_now = active_elems(elems, n_elems);
} while (active_now != active_prev);
for (int i = 0; i < 2; ++i){
clReleaseMemObject(elem_buffs[i]);
}
clReleaseCommandQueue(queue);
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseContext(context);
free(elems);
return 0;
}
element_t* load_elems(FILE *fp, int *n_elems, int *grid_dim){
//First line contains the number of elements and the grid dimensions
if (fscanf(fp, "%d %d\n", n_elems, grid_dim) == EOF){
printf("Error reading file header\n");
return NULL;
}
element_t *elems = malloc(*n_elems * sizeof(element_t));
for (int i = 0; i < *n_elems; ++i){
char dir[5];
if (fscanf(fp, "%d %d %d %s\n", &elems[i].x, &elems[i].y, &elems[i].r, dir) == EOF){
printf("Error reading element %d\n", i);
free(elems);
return NULL;
}
dir[4] = '\0';
//'A' is ASCII code 65
elems[i].sym = 65 + i;
//To make it easy to find the direction information make string lower case
//and place the udlr information in that order. so dir[0] will be u if we can go up and null if not
for (char *c = dir; *c; ++c){
*c = tolower(*c);
}
for (int j = 0; j < 4; ++j){
char c = 'u';
switch (j){
case 1:
c = 'd';
break;
case 2:
c = 'l';
break;
case 3:
c = 'r';
break;
default: break;
}
if (strchr(dir, c)){
elems[i].dir[j] = c;
}
else {
elems[i].dir[j] = '\0';
}
}
}
return elems;
}
void print_grid(element_t *elems, int n_elems, int grid_dim){
//E will track which element we should print next, and then we see if we've
//reached its position in the matrix, when we do we print and move to next element
for (int i = 0; i < grid_dim; ++i){
for (int j = 0; j < grid_dim; ++j){
char e = elem_at(elems, n_elems, j, i);
if (e){
printf("%c", e);
}
else {
printf(" ");
}
}
printf("\n");
}
}
char elem_at(element_t *elems, int n_elems, int x, int y){
for (int i = 0; i < n_elems; ++i){
if (elems[i].x == x && elems[i].y == y){
return elems[i].sym;
}
}
return '\0';
}
int active_elems(element_t *elems, int n_elems){
int n_active = 0;
for (int i = 0; i < n_elems; ++i){
if (elems[i].sym == 'X'){
++n_active;
}
}
return n_active;
}
char* file_content(char *f_name, size_t *size){
FILE *fp = fopen(f_name, "rb");
if (!fp){
printf("Failed to open file: %s\n", f_name);
return NULL;
}
fseek(fp, 0, SEEK_END);
*size = ftell(fp);
fseek(fp, 0, SEEK_SET);
char *content = malloc(sizeof(char) * (*size) + sizeof(char));
if (!content){
printf("file_content: Failed to allocate room for file content\n");
return NULL;
}
if (fread(content, sizeof(char), *size, fp) != *size){
printf("file_content: Failed to read all bytes, expected %u\n", *size);
free(content);
return NULL;
}
content[*size] = '\0';
return content;
}