From 750a8ce92eb4d66cc7f67243be27f0889d7074b7 Mon Sep 17 00:00:00 2001 From: Georg Hopp Date: Wed, 15 Jul 2015 18:27:26 +0200 Subject: [PATCH] checkin of the relevant examples of the OpenCL cookbook --- Makefile | 16 +++++++++++ README.md | 4 +++ part1.c | 60 ++++++++++++++++++++++++++++++++++++++++ part2.c | 78 +++++++++++++++++++++++++++++++++++++++++++++++++++ part3.c | 49 ++++++++++++++++++++++++++++++++ part4.c | 57 ++++++++++++++++++++++++++++++++++++++ part4.cl | 20 ++++++++++++++ part5.c | 71 +++++++++++++++++++++++++++++++++++++++++++++++ part6.c | 72 +++++++++++++++++++++++++++++++++++++++++++++++ part8.c | 83 +++++++++++++++++++++++++++++++++++++++++++++++++++++++ part8.cl | 11 ++++++++ 11 files changed, 521 insertions(+) create mode 100644 Makefile create mode 100644 README.md create mode 100644 part1.c create mode 100644 part2.c create mode 100644 part3.c create mode 100644 part4.c create mode 100644 part4.cl create mode 100644 part5.c create mode 100644 part6.c create mode 100644 part8.c create mode 100644 part8.cl diff --git a/Makefile b/Makefile new file mode 100644 index 0000000..ec1d749 --- /dev/null +++ b/Makefile @@ -0,0 +1,16 @@ +CFLAGS += -O2 -march=native -std=c99 -I/usr/local/include +LIBS += -lcl -L/usr/local/lib64/beignet +CC = cc + +BINARIES = part1 part2 part3 part4 part5 part6 part8 + +all: $(BINARIES) + +%: %.c + $(CC) $(CFLAGS) $(LIBS) -o $@ $< + strip $@ + +.PHONY: clean + +clean: + rm $(BINARIES) diff --git a/README.md b/README.md new file mode 100644 index 0000000..b92eb9a --- /dev/null +++ b/README.md @@ -0,0 +1,4 @@ +OpenCL tutorial notes +===================== + +URL: http://dhruba.name/2012/08/21/opencl-cookbook-series-reference/ diff --git a/part1.c b/part1.c new file mode 100644 index 0000000..5a2e5cd --- /dev/null +++ b/part1.c @@ -0,0 +1,60 @@ +#include +#include +#ifdef __APPLE__ +#include +#else +#include +#endif + +int main() { + + int i, j; + char* info; + size_t infoSize; + + cl_uint platformCount; + cl_platform_id *platforms; + + const char* attributeNames[5] = { + "Name", "Vendor", "Version", "Profile", "Extensions" }; + + const cl_platform_info attributeTypes[5] = { + CL_PLATFORM_NAME, CL_PLATFORM_VENDOR, CL_PLATFORM_VERSION, + CL_PLATFORM_PROFILE, CL_PLATFORM_EXTENSIONS }; + + const int attributeCount = sizeof(attributeNames) / sizeof(char*); + + // get platform count + clGetPlatformIDs(5, NULL, &platformCount); + // get all platforms + platforms = (cl_platform_id*) malloc(sizeof(cl_platform_id) * platformCount); + clGetPlatformIDs(platformCount, platforms, NULL); + + // for each platform print all attributes + for (i=0; i +#include +#ifdef __APPLE__ +#include +#else +#include +#endif + +int main() { + + int i, j; + char* value; + size_t valueSize; + cl_uint platformCount; + cl_platform_id* platforms; + cl_uint deviceCount; + cl_device_id* devices; + cl_uint maxComputeUnits; + + // get all platforms + clGetPlatformIDs(0, NULL, &platformCount); + platforms = (cl_platform_id*) malloc(sizeof(cl_platform_id) * platformCount); + clGetPlatformIDs(platformCount, platforms, NULL); + + for (i = 0; i < platformCount; i++) { + + // get all devices + clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, 0, NULL, &deviceCount); + devices = (cl_device_id*) malloc(sizeof(cl_device_id) * deviceCount); + clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, deviceCount, devices, NULL); + + // for each device print critical attributes + for (j = 0; j < deviceCount; j++) { + + // print device name + clGetDeviceInfo(devices[j], CL_DEVICE_NAME, 0, NULL, &valueSize); + value = (char*) malloc(valueSize); + clGetDeviceInfo(devices[j], CL_DEVICE_NAME, valueSize, value, NULL); + printf("%d. Device: %s\n", j+1, value); + free(value); + + // print hardware device version + clGetDeviceInfo(devices[j], CL_DEVICE_VERSION, 0, NULL, &valueSize); + value = (char*) malloc(valueSize); + clGetDeviceInfo(devices[j], CL_DEVICE_VERSION, valueSize, value, NULL); + printf(" %d.%d Hardware version: %s\n", j+1, 1, value); + free(value); + + // print software driver version + clGetDeviceInfo(devices[j], CL_DRIVER_VERSION, 0, NULL, &valueSize); + value = (char*) malloc(valueSize); + clGetDeviceInfo(devices[j], CL_DRIVER_VERSION, valueSize, value, NULL); + printf(" %d.%d Software version: %s\n", j+1, 2, value); + free(value); + + // print c version supported by compiler for device + clGetDeviceInfo(devices[j], CL_DEVICE_OPENCL_C_VERSION, 0, NULL, &valueSize); + value = (char*) malloc(valueSize); + clGetDeviceInfo(devices[j], CL_DEVICE_OPENCL_C_VERSION, valueSize, value, NULL); + printf(" %d.%d OpenCL C version: %s\n", j+1, 3, value); + free(value); + + // print parallel compute units + clGetDeviceInfo(devices[j], CL_DEVICE_MAX_COMPUTE_UNITS, + sizeof(maxComputeUnits), &maxComputeUnits, NULL); + printf(" %d.%d Parallel compute units: %d\n", j+1, 4, maxComputeUnits); + + } + + free(devices); + + } + + free(platforms); + return 0; +} + +// vim: set ft=c ts=4 sw=4: diff --git a/part3.c b/part3.c new file mode 100644 index 0000000..81dbf74 --- /dev/null +++ b/part3.c @@ -0,0 +1,49 @@ +#include +#include +#ifdef __APPLE__ +#include +#else +#include +#endif + +int main() { + + cl_platform_id platform; + cl_device_id device; + cl_context context; + cl_uint refCount; + + // get first available platform + clGetPlatformIDs(1, &platform, NULL); + + // get first available gpu device + clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL); + + // create context + context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL); + + // get context reference count + clGetContextInfo(context, CL_CONTEXT_REFERENCE_COUNT, + sizeof(refCount), &refCount, NULL); + printf("Ref count: %u ", refCount); + + // increment reference count + clRetainContext(context); + clGetContextInfo(context, CL_CONTEXT_REFERENCE_COUNT, + sizeof(refCount), &refCount, NULL); + printf(">> %u ", refCount); + + // decrement reference count + clReleaseContext(context); + clGetContextInfo(context, CL_CONTEXT_REFERENCE_COUNT, + sizeof(refCount), &refCount, NULL); + printf(">> %u ", refCount); + + // finally release context + clReleaseContext(context); + printf(">> 0\n"); + return 0; + +} + +// vim: set ft=c ts=4 sw=4: diff --git a/part4.c b/part4.c new file mode 100644 index 0000000..76defb4 --- /dev/null +++ b/part4.c @@ -0,0 +1,57 @@ +#include +#include +#include +#ifdef __APPLE__ +#include +#else +#include +#endif + +#define KERNEL "part4.cl" + +int main() { + + cl_platform_id platform; + cl_device_id device; + cl_context context; + cl_program program; + + FILE* programHandle; + size_t programSize, kernelSourceSize; + char *programBuffer, *kernelSource; + + // get first available platform and gpu and create context + clGetPlatformIDs(1, &platform, NULL); + clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL); + context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL); + + // get size of kernel source + programHandle = fopen(KERNEL, "r"); + fseek(programHandle, 0, SEEK_END); + programSize = ftell(programHandle); + rewind(programHandle); + + // read kernel source into buffer + programBuffer = (char*) malloc(programSize + 1); + programBuffer[programSize] = '\0'; + assert (programSize == fread(programBuffer, sizeof(char), programSize, programHandle)); + fclose(programHandle); + + // create program from buffer + program = clCreateProgramWithSource(context, 1, + (const char**) &programBuffer, &programSize, NULL); + free(programBuffer); + + // read kernel source back in from program to check + clGetProgramInfo(program, CL_PROGRAM_SOURCE, 0, NULL, &kernelSourceSize); + kernelSource = (char*) malloc(kernelSourceSize); + clGetProgramInfo(program, CL_PROGRAM_SOURCE, kernelSourceSize, kernelSource, NULL); + printf("\nKernel source:\n\n%s\n", kernelSource); + free(kernelSource); + + clReleaseContext(context); + return 0; + +} + +// vim: set ft=c ts=4 sw=4: diff --git a/part4.cl b/part4.cl new file mode 100644 index 0000000..431e0be --- /dev/null +++ b/part4.cl @@ -0,0 +1,20 @@ +__kernel void hello(__global char* string) { + + string[0] = 'H'; + string[1] = 'e'; + string[2] = 'l'; + string[3] = 'l'; + string[4] = 'o'; + string[5] = ','; + string[6] = ' '; + string[7] = 'W'; + string[8] = 'o'; + string[9] = 'r'; + string[10] = 'l'; + string[11] = 'd'; + string[12] = '!'; + string[13] = 0; + +} + +// vim: ft=c ts=4 sw=4: diff --git a/part5.c b/part5.c new file mode 100644 index 0000000..e3239a7 --- /dev/null +++ b/part5.c @@ -0,0 +1,71 @@ +#include +#include +#include +#ifdef __APPLE__ +#include +#else +#include +#endif + +#define KERNEL "part4.cl" + +int main() { + + cl_platform_id platform; cl_device_id device; cl_context context; + cl_program program; cl_int error; cl_build_status status; + + FILE* programHandle; + char *programBuffer; char *programLog; + size_t programSize; size_t logSize; + + // get first available platform and gpu and create context + clGetPlatformIDs(1, &platform, NULL); + clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL); + context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL); + + // get size of kernel source + programHandle = fopen(KERNEL, "r"); + fseek(programHandle, 0, SEEK_END); + programSize = ftell(programHandle); + rewind(programHandle); + + // read kernel source into buffer + programBuffer = (char*) malloc(programSize + 1); + programBuffer[programSize] = '\0'; + assert (programSize == fread(programBuffer, sizeof(char), programSize, programHandle)); + + fclose(programHandle); + + // create program from buffer + program = clCreateProgramWithSource(context, 1, + (const char**) &programBuffer, &programSize, NULL); + free(programBuffer); + + // build program + const char options[] = "-Werror -cl-std=CL1.1"; + error = clBuildProgram(program, 1, &device, options, NULL, NULL); + + // build failed + if (error != CL_SUCCESS) { + + // check build error and build status first + clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_STATUS, + sizeof(cl_build_status), &status, NULL); + + // check build log + clGetProgramBuildInfo(program, device, + CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize); + programLog = (char*) calloc (logSize+1, sizeof(char)); + clGetProgramBuildInfo(program, device, + CL_PROGRAM_BUILD_LOG, logSize+1, programLog, NULL); + printf("Build failed; error=%d, status=%d, programLog:\n\n%s", + error, status, programLog); + free(programLog); + + } + + clReleaseContext(context); + return 0; +} + +// vim: set ts=4 sw=4: diff --git a/part6.c b/part6.c new file mode 100644 index 0000000..eb27563 --- /dev/null +++ b/part6.c @@ -0,0 +1,72 @@ +#include +#include +#include +#ifdef __APPLE__ +#include +#else +#include +#endif + +#define KERNEL "part4.cl" + +int main() { + + cl_platform_id platform; cl_device_id device; cl_context context; + cl_program program; cl_kernel kernel; cl_command_queue queue; + cl_mem kernelBuffer; + + FILE* programHandle; char *programBuffer; char *programLog; + size_t programSize; char hostBuffer[32]; + + // get first available sdk and gpu and create context + clGetPlatformIDs(1, &platform, NULL); + clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL); + context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL); + + // get size of kernel source + programHandle = fopen(KERNEL, "r"); + fseek(programHandle, 0, SEEK_END); + programSize = ftell(programHandle); + rewind(programHandle); + + // read kernel source into buffer + programBuffer = (char*) malloc(programSize + 1); + programBuffer[programSize] = '\0'; + assert (programSize == fread(programBuffer, sizeof(char), programSize, programHandle)); + + fclose(programHandle); + + // create and build program + program = clCreateProgramWithSource(context, 1, + (const char**) &programBuffer, &programSize, NULL); + free(programBuffer); + clBuildProgram(program, 1, &device, "-Werror -cl-std=CL1.1", NULL, NULL); + + // create kernel and command queue + kernel = clCreateKernel(program, "hello", NULL); + queue = clCreateCommandQueue(context, device, 0, NULL); + + // create kernel argument buffer and set it into kernel + kernelBuffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, + 32 * sizeof(char), NULL, NULL); + clSetKernelArg(kernel, 0, sizeof(cl_mem), &kernelBuffer); + + // execute kernel, read back the output and print to screen + clEnqueueTask(queue, kernel, 0, NULL, NULL); + clEnqueueReadBuffer(queue, kernelBuffer, CL_TRUE, 0, + 32 * sizeof(char), hostBuffer, 0, NULL, NULL); + puts(hostBuffer); + + clFlush(queue); + clFinish(queue); + clReleaseKernel(kernel); + clReleaseProgram(program); + clReleaseMemObject(kernelBuffer); + clReleaseCommandQueue(queue); + clReleaseContext(context); + + return 0; + +} + +// vim: set ft=c ts=4 sw=4: diff --git a/part8.c b/part8.c new file mode 100644 index 0000000..fe01657 --- /dev/null +++ b/part8.c @@ -0,0 +1,83 @@ +#include +#include +#include +#ifdef __APPLE__ +#include +#else +#include +#endif + +#define KERNEL "part8.cl" + +void cpu_3d_loop (int x, int y, int z) { + for (int i = 0; i < x; i++) { + for (int j = 0; j < y; j++) { + for (int k = 0; k < z; k++) { + printf("CPU %d,%d,%d\n", i, j, k); + } + } + } +} + +int main() { + + cl_platform_id platform; cl_device_id device; cl_context context; + cl_program program; cl_kernel kernel; cl_command_queue queue; + cl_mem kernelBuffer; + + FILE* programHandle; char *programBuffer; char *programLog; + size_t programSize; char hostBuffer[32]; + + int x = 4; + int y = 3; + int z = 2; + + cpu_3d_loop(x, y, z); + + // get first available sdk and gpu and create context + clGetPlatformIDs(1, &platform, NULL); + clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL); + context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL); + + // get size of kernel source + programHandle = fopen(KERNEL, "r"); + fseek(programHandle, 0, SEEK_END); + programSize = ftell(programHandle); + rewind(programHandle); + + // read kernel source into buffer + programBuffer = (char*) malloc(programSize + 1); + programBuffer[programSize] = '\0'; + assert (programSize == fread(programBuffer, sizeof(char), programSize, programHandle)); + + fclose(programHandle); + + // create and build program + program = clCreateProgramWithSource(context, 1, + (const char**) &programBuffer, &programSize, NULL); + free(programBuffer); + clBuildProgram(program, 1, &device, "-Werror -cl-std=CL1.1", NULL, NULL); + + // create kernel and command queue + kernel = clCreateKernel(program, "ndrange_parallelism", NULL); + queue = clCreateCommandQueue(context, device, 0, NULL); + + size_t globalWorkSize[3] = {x, y, z}; + + // execute kernel, read back the output and print to screen + clEnqueueNDRangeKernel( + queue, kernel, 3, NULL, globalWorkSize, NULL, 0, NULL, NULL); + + clFlush(queue); + clFinish(queue); + clReleaseKernel(kernel); + clReleaseProgram(program); + clReleaseMemObject(kernelBuffer); + clReleaseCommandQueue(queue); + clReleaseContext(context); + + return 0; + +} + +// vim: set ft=c ts=4 sw=4: diff --git a/part8.cl b/part8.cl new file mode 100644 index 0000000..09cce3e --- /dev/null +++ b/part8.cl @@ -0,0 +1,11 @@ +__kernel void ndrange_parallelism () { + + int i = get_global_id(0); + int j = get_global_id(1); + int k = get_global_id(2); + + printf("GPU %d,%d,%d\n", i, j, k); + +} + +// vim: set ft=c ts=4 sw=4: