From 764404f24367103a633d6be04866426c8ef9c589 Mon Sep 17 00:00:00 2001 From: Giuseppe Bilotta Date: Wed, 31 Aug 2011 10:37:35 +0200 Subject: [PATCH] First draft --- .gitignore | 1 + Makefile | 2 + babel.cc | 232 +++++++++++++++++++++++++++++++++++++++++++++++++++++ babel.cl | 18 +++++ babel.h | 46 +++++++++++ 5 files changed, 299 insertions(+) create mode 100644 .gitignore create mode 100644 Makefile create mode 100644 babel.cc create mode 100644 babel.cl create mode 100644 babel.h diff --git a/.gitignore b/.gitignore new file mode 100644 index 0000000..98f6593 --- /dev/null +++ b/.gitignore @@ -0,0 +1 @@ +babel diff --git a/Makefile b/Makefile new file mode 100644 index 0000000..a9c143f --- /dev/null +++ b/Makefile @@ -0,0 +1,2 @@ +babel: babel.cc babel.h + $(CXX) -o babel babel.cc -lOpenCL diff --git a/babel.cc b/babel.cc new file mode 100644 index 0000000..259dcb5 --- /dev/null +++ b/babel.cc @@ -0,0 +1,232 @@ +/* usual C/C++ includes */ +#include +#include // for ceil() +#include + +/* OpenCL includes */ +#ifdef __APPLE__ +#include +#else +#include +#endif + +void check_ocl_error(const cl_int &error, const char *message) { + if (error != CL_SUCCESS) { + fprintf(stderr, "error %d %s\n", error, message); + exit(1); + } +} + +char *read_file(const char *fname) { + size_t fsize, readsize; + char *buff; + + FILE *fd = fopen(fname, "rb"); + if (!fd) { + fprintf(stderr, "%s not found\n", fname); + return NULL; + } + + fseek(fd, 0, SEEK_END); + fsize = ftell(fd); + + buff = (char *)malloc(fsize+1); + rewind(fd); + readsize = fread(buff, 1, fsize, fd); + if (fsize != readsize) { + fprintf(stderr, "could only read %lu/%lu bytes from %s\n", + readsize, fsize, fname); + free(buff); + return NULL; + } + buff[fsize] = '\0'; + + printf("read %lu bytes from %s\n", fsize, fname); + + return buff; +} + +#include "babel.h" + +int main(int argc, char **argv) { + + /* sanity check */ + if (SYMBOLS + 1 != sizeof(alphabet)) { + fprintf(stderr, "Wrong alphabet: %s has %lu symbols, expected %d\n", + alphabet, sizeof(alphabet) - 1, SYMBOLS); + return 2; + } + + char page[CHARS_PER_PAGE+1]; + page[CHARS_PER_PAGE] = '\0'; + + + /* auxiliary buffer to read platform and device info */ + char buffer[1024]; + + /* platform selection */ + cl_uint num_platforms = 0; + cl_platform_id *platform_list = NULL; + cl_platform_id platform = NULL; + + clGetPlatformIDs(0, NULL, &num_platforms); // retrieve number of platform IDs + platform_list = (cl_platform_id *)calloc(num_platforms, sizeof(cl_platform_id)); + cl_int error = clGetPlatformIDs(num_platforms, platform_list, NULL); // retrieve the actual platform IDs + + /* a quicker way if we are only interested in the first/default platform, ID, would be to have: + clGetPlatformIDs(1, &platform, NULL); + */ + + check_ocl_error(error, "getting platform IDs"); + + printf("%d OpenCL platforms found:\n", num_platforms); + + for (cl_uint i = 0; i < num_platforms; ++i) { + /* last param: actual size of the query result */ + error = clGetPlatformInfo(platform_list[i], CL_PLATFORM_NAME, sizeof(buffer), buffer, NULL); + check_ocl_error(error, "getting platform name"); + printf("\tplatform %u: %s ", i, buffer); + error = clGetPlatformInfo(platform_list[i], CL_PLATFORM_VENDOR, sizeof(buffer), buffer, NULL); + check_ocl_error(error, "getting platform vendor"); + printf(" (%s)\n", buffer); + } + + cl_uint platnum = 0; + if (argc > 1) + platnum = atoi(argv[1]); + platform = platform_list[platnum]; + printf("using platform %u\n", platnum); + + /* device selection */ + + cl_uint num_devs = 0; + cl_device_id *device_list = NULL; + cl_device_id device = NULL; + + /* possible types: CPU, GPU, ACCELERATOR, DEFAULT, ALL */ + clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, NULL, &num_devs); + device_list = (cl_device_id *)calloc(num_devs, sizeof(cl_device_id)); + error = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, num_devs, device_list, NULL); + + check_ocl_error(error, "getting device IDs"); + + printf("%d devices found:\n", num_devs); + + for (cl_uint i = 0; i < num_devs; ++i) { + /* last param: actual size of the query result */ + error = clGetDeviceInfo(device_list[i], CL_DEVICE_NAME, sizeof(buffer), buffer, NULL); + check_ocl_error(error, "getting device name"); + printf("\tdevice %u: %s\n", i, buffer); + } + + cl_uint devnum = 0; + if (argc > 2) + devnum = atoi(argv[2]); + device = device_list[devnum]; + printf("using device %u\n", devnum); + + /* creating a context for one devices */ + + cl_context_properties ctx_prop[] = { + CL_CONTEXT_PLATFORM, (cl_context_properties)platform, + 0 + }; + + cl_context ctx = clCreateContext(ctx_prop, 1, &device, NULL, NULL, &error); + check_ocl_error(error, "creating context"); + + /* and a command queue to go with it */ + cl_command_queue queue = clCreateCommandQueue(ctx, device, CL_QUEUE_PROFILING_ENABLE, &error); + check_ocl_error(error, "creating command queue"); + + + /* allocate device memory */ + cl_mem page_d; + + page_d = clCreateBuffer(ctx, CL_MEM_WRITE_ONLY, CHARS_PER_PAGE, NULL, &error); + check_ocl_error(error, "allocating device page memory buffer"); + + /* load and build program */ + char *prog_source = read_file("babel.cl"); + if (prog_source == NULL) + exit(1); + + cl_program program = clCreateProgramWithSource(ctx, 1, (const char **)&prog_source, NULL, &error); + check_ocl_error(error, "creating program"); + + /* AMD APP doesn't include the current directory by default, apparently */ + const char *clopts = "-I."; + error = clBuildProgram(program, + 1, &device, // device(s) + clopts, // compiler options + NULL, // callback + NULL); + if (error == CL_BUILD_PROGRAM_FAILURE) { + size_t logSize = 0; + char *log; + error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize); + check_ocl_error(error, "getting program build info size"); + log = (char *)malloc(logSize); + error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, logSize, log, NULL); + check_ocl_error(error, "getting program build info"); + fputs(log, stderr); + fputs("\n", stderr); + exit(1); + } else + check_ocl_error(error, "building program"); + + + /* loading the kernel */ + cl_kernel fillpageKernel = clCreateKernel(program, "fillpage", &error); + check_ocl_error(error, "creating kernel"); + + error = clSetKernelArg(fillpageKernel, 0, sizeof(page_d), &page_d); + check_ocl_error(error, "setting kernel param 0"); + + /* group size */ + size_t group_size = CHARS_PER_LINE; + if (argc > 3) + group_size = atoi(argv[3]); + + /* work_size must be a multiple of group_size */ + size_t work_size = ceil(float(CHARS_PER_PAGE)/group_size)*group_size; + + /* launch kernel, with an event to collect profiling info */ + cl_ulong startTime, endTime; + cl_event evt; + + clEnqueueNDRangeKernel(queue, fillpageKernel, + 1, + NULL, &work_size, &group_size, + 0, NULL, + &evt); + + error = clFinish(queue); // sync on queue + check_ocl_error(error, "finishing queue"); + + clGetEventProfilingInfo(evt, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &startTime, NULL); + error = clGetEventProfilingInfo(evt, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &endTime, NULL); + check_ocl_error(error, "getting profiling info"); + + printf("Kernel runtime: %gms\n", double(endTime-startTime)/1000000); + + /* copy memory down */ + error = clEnqueueReadBuffer(queue, page_d, true, 0, CHARS_PER_PAGE, page, 0, NULL, NULL); + check_ocl_error(error, "getting results"); + + clFinish(queue); + + clReleaseMemObject(page_d); + + clReleaseProgram(program); + clReleaseCommandQueue(queue); + clReleaseContext(ctx); + + for(size_t line = 0; line < LINES_PER_PAGE; ++line) { + for(size_t col = 0; col < CHARS_PER_LINE; ++col) { + putchar(page[line*CHARS_PER_LINE + col]); + } + puts(""); + } + fflush(stdout); +} diff --git a/babel.cl b/babel.cl new file mode 100644 index 0000000..02bfcfa --- /dev/null +++ b/babel.cl @@ -0,0 +1,18 @@ +#include "babel.h" + +__kernel +void +fillpage(__global char* restrict page) +{ + /* linearized index in the page */ + size_t gix = get_global_id(0); + if (gix >= CHARS_PER_PAGE) + return; + + /* TODO get a character with a specific algorithm */ + size_t cix = gix; + cix %= SYMBOLS; + + page[gix] = alphabet[cix]; + +} diff --git a/babel.h b/babel.h new file mode 100644 index 0000000..73d076f --- /dev/null +++ b/babel.h @@ -0,0 +1,46 @@ +#ifndef BABEL_H +#define BABEL_H + +/* Veinte anaqueles, a cinco largos anaqueles por lado, cubren todos los + * lados menos dos + */ +#define SIDES_PER_CELL 4 +#define SHELVES_PER_SIDE 5 +#define SHELVES_PER_CELL (SHELVES_PER_SIDE*SIDES_PER_CELL) + +/* cada anaquel encierra treinta y dos libros de formato uniforme */ +#define BOOKS_PER_SHELF 32 +/* cada libro es de cuatrocientas diez páginas */ +#define PAGES_PER_BOOK 410 +/* cada página, de cuarenta renglones */ +#define LINES_PER_PAGE 40 +/* cada renglón, de unas ochenta letras */ +#define CHARS_PER_LINE 80 + +/* También hay letras en el dorso de cada libro */ +/* FIXME: the amount of characters on the spine is not specified in the + * text. If we assume a square character matrix, then the maximum number + * of characters on the spine is possibly about the same as the number + * of lines per page, but that's a rather wild assumption (it's likely + * to be about twice as much) + */ +/* TODO +#define CHARS_PER_SPINE LINES_PER_PAGE +*/ + +#define CHARS_PER_PAGE (CHARS_PER_LINE*LINES_PER_PAGE) +#define CHARS_PER_BOOK (CHARS_PER_PAGE*PAGES_PER_BOOK) + +/* el espacio, el punto, la coma, las veintidós letras del alfabeto */ + +/* to get an alphabet with 22 symbols, we use the letters of the latin + * alphabet which are also preset in the Esperanto alphabet (i.e. the + * Esperanto letters without diacritics). If this makes Zamenhof or + * Borges and his Kabalistic reference spin in his grave, we can use + * them to produce energy. + */ + +#define SYMBOLS 25 +const char alphabet[] = " .,ABCDEFGHIJKLMNOPRSTUVZ"; + +#endif // BABEL_H -- 2.32.0.93.g670b81a890