From d2b79ceec379be88ba87c33ee3a0902f76ac0fce Mon Sep 17 00:00:00 2001 From: mitchellhansen Date: Thu, 11 Aug 2016 21:39:06 -0700 Subject: [PATCH] Added a compilation routine, probably will abstract all this out into it's own class / function. Added a small kernel that I got from a tutorial to test the compilers error codes. Added a small notes file with error codes. Added some error checking for the error codes --- kernels/kernel.txt | 32 +++++++++++++ notes/opencl_error_codes.txt | 47 +++++++++++++++++++ src/main.cpp | 88 +++++++++++++++++++++++++++++++++--- 3 files changed, 161 insertions(+), 6 deletions(-) create mode 100644 kernels/kernel.txt create mode 100644 notes/opencl_error_codes.txt diff --git a/kernels/kernel.txt b/kernels/kernel.txt new file mode 100644 index 0000000..bf5518c --- /dev/null +++ b/kernels/kernel.txt @@ -0,0 +1,32 @@ +__constant sampler_t sampler = + CLK_NORMALIZED_COORDS_FALSE + | CLK_ADDRESS_CLAMP_TO_EDGE + | CLK_FILTER_NEAREST; + +__constant int FILTER_SIZE = 10; + +float FilterValue (__constant const float* filterWeights, + const int x, const int y) +{ + return filterWeights[(x+FILTER_SIZE) + (y+FILTER_SIZE)*(FILTER_SIZE*2 + 1)]; +} + +__kernel void Filter ( + __read_only image2d_t input, + __constant float* filterWeights, + __write_only image2d_t output) +{ + const int2 pos = {get_global_id(0), get_global_id(1)}; + + float4 sum = (float4)(0.0f); + for(int y = -FILTER_SIZE; y <= FILTER_SIZE; y++) { + for(int x = -FILTER_SIZE; x <= FILTER_SIZE; x++) { + sum += FilterValue(filterWeights, x, y) + * read_imagef(input, sampler, pos + (int2)(x,y)); + } + } + + write_imagef (output, (int2)(pos.x, pos.y), sum); +} + + diff --git a/notes/opencl_error_codes.txt b/notes/opencl_error_codes.txt new file mode 100644 index 0000000..61ead52 --- /dev/null +++ b/notes/opencl_error_codes.txt @@ -0,0 +1,47 @@ +CL_SUCCESS 0 +CL_DEVICE_NOT_FOUND -1 +CL_DEVICE_NOT_AVAILABLE -2 +CL_COMPILER_NOT_AVAILABLE -3 +CL_MEM_OBJECT_ALLOCATION_FAILURE -4 +CL_OUT_OF_RESOURCES -5 +CL_OUT_OF_HOST_MEMORY -6 +CL_PROFILING_INFO_NOT_AVAILABLE -7 +CL_MEM_COPY_OVERLAP -8 +CL_IMAGE_FORMAT_MISMATCH -9 +CL_IMAGE_FORMAT_NOT_SUPPORTED -10 +CL_BUILD_PROGRAM_FAILURE -11 +CL_MAP_FAILURE -12 + +CL_INVALID_VALUE -30 +CL_INVALID_DEVICE_TYPE -31 +CL_INVALID_PLATFORM -32 +CL_INVALID_DEVICE -33 +CL_INVALID_CONTEXT -34 +CL_INVALID_QUEUE_PROPERTIES -35 +CL_INVALID_COMMAND_QUEUE -36 +CL_INVALID_HOST_PTR -37 +CL_INVALID_MEM_OBJECT -38 +CL_INVALID_IMAGE_FORMAT_DESCRIPTOR -39 +CL_INVALID_IMAGE_SIZE -40 +CL_INVALID_SAMPLER -41 +CL_INVALID_BINARY -42 +CL_INVALID_BUILD_OPTIONS -43 +CL_INVALID_PROGRAM -44 +CL_INVALID_PROGRAM_EXECUTABLE -45 +CL_INVALID_KERNEL_NAME -46 +CL_INVALID_KERNEL_DEFINITION -47 +CL_INVALID_KERNEL -48 +CL_INVALID_ARG_INDEX -49 +CL_INVALID_ARG_VALUE -50 +CL_INVALID_ARG_SIZE -51 +CL_INVALID_KERNEL_ARGS -52 +CL_INVALID_WORK_DIMENSION -53 +CL_INVALID_WORK_GROUP_SIZE -54 +CL_INVALID_WORK_ITEM_SIZE -55 +CL_INVALID_GLOBAL_OFFSET -56 +CL_INVALID_EVENT_WAIT_LIST -57 +CL_INVALID_EVENT -58 +CL_INVALID_OPERATION -59 +CL_INVALID_GL_OBJECT -60 +CL_INVALID_BUFFER_SIZE -61 +CL_INVALID_MIP_LEVEL -62 diff --git a/src/main.cpp b/src/main.cpp index 86f4eee..43619f3 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -6,6 +6,8 @@ #include #include "Curses.h" # include +#include +#include #ifdef linux @@ -18,19 +20,37 @@ # include # include + + #endif const int WINDOW_X = 150; const int WINDOW_Y = 150; +std::string read_file(std::string file_name){ + std::ifstream input_file(file_name); -int main(){ + if (!input_file.is_open()){ + std::cout << file_name << " could not be opened" << std::endl; + return nullptr; + } - // ===================================================================== // - // ==== Opencl + std::stringstream buf; + buf << input_file.rdbuf(); + return buf.str(); +} +int main(){ + char buffer[256]; + char *val = getcwd(buffer, sizeof(buffer)); + if (val) { + std::cout << buffer << std::endl; + } + // ===================================================================== // + // ==== Opencl setup + int error = 0; // Get the number of platforms @@ -44,8 +64,7 @@ int main(){ // get the number of devices, fetch them, choose the first one cl_uint deviceIdCount = 0; - std::vector deviceIds (deviceIdCount); - + std::vector deviceIds; // Try to get a GPU first error = clGetDeviceIDs (platformIds [0], CL_DEVICE_TYPE_GPU, 0, nullptr, &deviceIdCount); @@ -53,12 +72,18 @@ int main(){ if (deviceIdCount == 0) { std::cout << "couldn't aquire a GPU, falling back to CPU" << std::endl; error = clGetDeviceIDs(platformIds[0], CL_DEVICE_TYPE_CPU, 0, nullptr, &deviceIdCount); + deviceIds.resize(deviceIdCount); error = clGetDeviceIDs(platformIds[0], CL_DEVICE_TYPE_CPU, deviceIdCount, deviceIds.data(), NULL); } else { std::cout << "aquired GPU cl target" << std::endl; + deviceIds.resize(deviceIdCount); clGetDeviceIDs (platformIds[0], CL_DEVICE_TYPE_GPU, deviceIdCount, deviceIds.data (), nullptr); } + if (error != 0){ + std::cout << "Err: clGetDeviceIDs returned: " << error << std::endl; + return error; + } // Hurray for standards! // Setup the context properties to grab the current GL context @@ -98,11 +123,62 @@ int main(){ &error ); + if (error != 0){ + std::cout << "Err: clCreateContext returned: " << error << std::endl; + return error; + } + // And the cl command queue - auto commandQueue = clCreateCommandQueue(context, deviceIds[0], 0, NULL); + auto commandQueue = clCreateCommandQueue(context, deviceIds[0], 0, &error); + if (error != 0){ + std::cout << "Err: clCreateCommandQueue returned: " << error << std::endl; + return error; + } // At this point the shared GL/CL context is up and running + // ====================================================================== // + // ========== Kernel setup & compilation + + // Load in the kernel, and c stringify it + std::string kernel_source; + kernel_source = read_file("../kernels/kernel.txt"); + const char* kernel_source_c_str = kernel_source.c_str(); + size_t kernel_source_size = strlen(kernel_source_c_str); + + + // Load the source into CL's data structure + cl_program kernel_program = clCreateProgramWithSource( + context, 1, + &kernel_source_c_str, + &kernel_source_size, &error + ); + + if (error != 0){ + std::cout << "Err: clCreateProgramWithSource returned: " << error << std::endl; + return error; + } + + + // Try and build the program + error = clBuildProgram(kernel_program, 1, &deviceIds[0], NULL, NULL, NULL); + + // Check to see if it errored out + if (error == CL_BUILD_PROGRAM_FAILURE){ + + // Get the size of the queued log + size_t log_size; + clGetProgramBuildInfo(kernel_program, deviceIds[0], CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size); + char *log = new char[log_size]; + + // Grab the log + clGetProgramBuildInfo(kernel_program, deviceIds[0], CL_PROGRAM_BUILD_LOG, log_size, log, NULL); + + + std::cout << "Err: clBuildProgram returned: " << error << std::endl; + std::cout << log << std::endl; + return error; + } };