From 14430f2092f27a4fda9d8d80be90373388a06801 Mon Sep 17 00:00:00 2001 From: Benjamin Segovia Date: Thu, 3 May 2012 17:12:18 +0000 Subject: [PATCH] Finished the small test framework --- utests/compiler_copy_buffer.cpp | 18 +++++-------- utests/compiler_copy_buffer_row.cpp | 21 ++++----------- utests/compiler_write_only.cpp | 17 ++++++------ utests/utest.cpp | 5 ++++ utests/utest.hpp | 12 ++++++--- utests/utest_assert.cpp | 3 +-- utests/utest_assert.hpp | 12 +++++++++ utests/utest_helper.cpp | 49 +++++++++++++++------------------ utests/utest_helper.hpp | 54 +++++++++++++++++++++---------------- utests/utest_run.cpp | 20 +++++++++----- 10 files changed, 112 insertions(+), 99 deletions(-) diff --git a/utests/compiler_copy_buffer.cpp b/utests/compiler_copy_buffer.cpp index 9380ace..ee1dbf1 100644 --- a/utests/compiler_copy_buffer.cpp +++ b/utests/compiler_copy_buffer.cpp @@ -22,9 +22,9 @@ static void compiler_copy_buffer(void) { const size_t n = 8192 * 4; - int status = 0; - CALL (cl_test_init, "test_copy_buffer.cl", "test_copy_buffer", SOURCE); + // Setup kernel and buffers + OCL_CREATE_KERNEL("test_copy_buffer"); buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * n); for (uint32_t i = 0; i < n; ++i) ((uint32_t*)buf_data[0])[i] = i; OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), buf_data[0]); @@ -32,25 +32,19 @@ static void compiler_copy_buffer(void) free(buf_data[0]); buf_data[0] = NULL; + // Run the kernel OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); - globals[0] = n; locals[0] = 16; OCL_NDRANGE(1); + // Check result OCL_MAP_BUFFER(0); OCL_MAP_BUFFER(1); for (uint32_t i = 0; i < n; ++i) - assert(((uint32_t*)buf_data[0])[i] == ((uint32_t*)buf_data[1])[i]); - OCL_UNMAP_BUFFER(0); - OCL_UNMAP_BUFFER(1); - -error: - cl_release_buffers(); - cl_report_error(status); - cl_test_destroy(); + OCL_ASSERT(((uint32_t*)buf_data[0])[i] == ((uint32_t*)buf_data[1])[i]); } -UTEST_REGISTER(compiler_copy_buffer) +MAKE_UTEST_FROM_FUNCTION(compiler_copy_buffer); diff --git a/utests/compiler_copy_buffer_row.cpp b/utests/compiler_copy_buffer_row.cpp index 308e6e3..cfaea0e 100644 --- a/utests/compiler_copy_buffer_row.cpp +++ b/utests/compiler_copy_buffer_row.cpp @@ -26,25 +26,21 @@ static void compiler_copy_buffer_row(void) const int row = 8192; const int row_n = 2; const int n = row * row_n; - int status = 0; - CALL (cl_test_init, "test_copy_buffer_row.cl", "test_copy_buffer_row", SOURCE); - - // Create the input data + // Setup kernel and buffers + OCL_CREATE_KERNEL("test_copy_buffer_row"); src_buffer = (uint32_t *) malloc(sizeof(uint32_t) * n); for (int32_t i = 0; i < n; ++i) src_buffer[i] = i; data_buffer = (int *) malloc(sizeof(int) * 2); data_buffer[0] = row; data_buffer[1] = n; - - // Allocate all buffers OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(uint32_t), src_buffer); OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); OCL_CREATE_BUFFER(buf[2], CL_MEM_COPY_HOST_PTR, 2 * sizeof(uint32_t), data_buffer); free(src_buffer); free(data_buffer); - // Run the code + // Run the kernel OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]); @@ -56,15 +52,8 @@ static void compiler_copy_buffer_row(void) OCL_MAP_BUFFER(0); OCL_MAP_BUFFER(1); for (int32_t i = 0; i < n; ++i) - assert(((uint32_t*)buf_data[0])[i] == ((uint32_t*)buf_data[1])[i]); - OCL_UNMAP_BUFFER(0); - OCL_UNMAP_BUFFER(1); -error: - cl_release_buffers(); - cl_report_error(status); - cl_test_destroy(); - + OCL_ASSERT(((uint32_t*)buf_data[0])[i] == ((uint32_t*)buf_data[1])[i]); } -UTEST_REGISTER(compiler_copy_buffer_row) +MAKE_UTEST_FROM_FUNCTION(compiler_copy_buffer_row); diff --git a/utests/compiler_write_only.cpp b/utests/compiler_write_only.cpp index f738253..b42bfd2 100644 --- a/utests/compiler_write_only.cpp +++ b/utests/compiler_write_only.cpp @@ -22,23 +22,22 @@ void compiler_write_only(void) { const size_t n = 2048; - int status = 0; - CALL (cl_test_init, "test_write_only.cl", "test_write_only", SOURCE); + // Setup kernel and buffers + OCL_CREATE_KERNEL("test_write_only"); OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + + // Run the kernel globals[0] = n; locals[0] = 16; OCL_NDRANGE(1); OCL_MAP_BUFFER(0); - for (uint32_t i = 0; i < n; ++i) assert(((uint32_t*)buf_data[0])[i] == i); - OCL_UNMAP_BUFFER(0); -error: - cl_release_buffers(); - cl_report_error(status); - cl_test_destroy(); + // Check results + for (uint32_t i = 0; i < n; ++i) + OCL_ASSERT(((uint32_t*)buf_data[0])[i] == i); } -UTEST_REGISTER(compiler_write_only); +MAKE_UTEST_FROM_FUNCTION(compiler_write_only); diff --git a/utests/utest.cpp b/utests/utest.cpp index 421c3b6..fc3467e 100644 --- a/utests/utest.cpp +++ b/utests/utest.cpp @@ -22,6 +22,7 @@ * \author Benjamin Segovia */ #include "utest.hpp" +#include "utest_helper.hpp" #include #include #include @@ -56,6 +57,8 @@ void UTest::run(const char *name) { std::cout << utest.name << ":" << std::endl; (utest.fn)(); std::cout << std::endl; + cl_kernel_destroy(); + cl_buffer_destroy(); } } } @@ -68,6 +71,8 @@ void UTest::runAll(void) { std::cout << utest.name << ":" << std::endl; (utest.fn)(); std::cout << std::endl; + cl_kernel_destroy(); + cl_buffer_destroy(); } } diff --git a/utests/utest.hpp b/utests/utest.hpp index e444c72..338a4dc 100644 --- a/utests/utest.hpp +++ b/utests/utest.hpp @@ -27,8 +27,9 @@ #ifndef __UTEST_UTEST_HPP__ #define __UTEST_UTEST_HPP__ +#include "utest_exception.hpp" #include -#include +#include /*! Quick and dirty unit test system with registration */ struct UTest @@ -51,9 +52,14 @@ struct UTest static void runAll(void); }; -/*! RegisterData a new unit test */ +/*! Register a new unit test */ #define UTEST_REGISTER(FN) static const UTest __##FN##__(FN, #FN); +/*! Turn a function into a unit test */ +#define MAKE_UTEST_FROM_FUNCTION(FN) \ + static void __ANON__##FN##__(void) { UTEST_EXPECT_SUCCESS(FN()); } \ + static const UTest __##FN##__(__ANON__##FN##__, #FN); + /*! No assert is expected */ #define UTEST_EXPECT_SUCCESS(EXPR) \ do { \ @@ -61,7 +67,7 @@ struct UTest EXPR; \ std::cout << " " << #EXPR << " [SUCCESS]" << std::endl; \ } \ - catch (gbe::Exception e) { \ + catch (Exception e) { \ std::cout << " " << #EXPR << " [FAILED]" << std::endl; \ std::cout << " " << e.what() << std::endl; \ } \ diff --git a/utests/utest_assert.cpp b/utests/utest_assert.cpp index da656e8..f3b9a00 100644 --- a/utests/utest_assert.cpp +++ b/utests/utest_assert.cpp @@ -31,12 +31,11 @@ void onFailedAssertion(const char *msg, const char *file, const char *fn, int li char lineString[256]; sprintf(lineString, "%i", line); assert(msg != NULL && file != NULL && fn != NULL); - const std::string str = "Compiler error: " + const std::string str = "Error: " + std::string(msg) + "\n at file " + std::string(file) + ", function " + std::string(fn) + ", line " + std::string(lineString); - // assert(0); throw Exception(str); } diff --git a/utests/utest_assert.hpp b/utests/utest_assert.hpp index d124e9f..f93f9ac 100644 --- a/utests/utest_assert.hpp +++ b/utests/utest_assert.hpp @@ -28,5 +28,17 @@ /*! To ensure that condition truth. Optional message is supported */ void onFailedAssertion(const char *msg, const char *file, const char *fn, int line); +#define OCL_ASSERT(EXPR) \ + do { \ + if (!(EXPR)) \ + onFailedAssertion(#EXPR, __FILE__, __FUNCTION__, __LINE__); \ + } while (0) + +#define OCL_ASSERTM(EXPR, MSG) \ + do { \ + if (!(EXPR)) \ + onFailedAssertion(MSG, __FILE__, __FUNCTION__, __LINE__); \ + } while (0) + #endif /* __OCL_ASSERT_HPP__ */ diff --git a/utests/utest_helper.cpp b/utests/utest_helper.cpp index 3423b68..0083618 100644 --- a/utests/utest_helper.cpp +++ b/utests/utest_helper.cpp @@ -160,7 +160,7 @@ cl_kernel_init(const char *file_name, const char *kernel_name, int format) } /* OCL requires to build the program even if it is created from a binary */ - CALL (clBuildProgram, program, 1, &device, NULL, NULL, NULL); + OCL_CALL (clBuildProgram, program, 1, &device, NULL, NULL, NULL); /* Create a kernel from the program */ kernel = clCreateKernel(program, kernel_name, &status); @@ -186,30 +186,30 @@ cl_ocl_init(void) size_t i; /* Get the platform number */ - CALL (clGetPlatformIDs, 0, NULL, &platform_n); + OCL_CALL (clGetPlatformIDs, 0, NULL, &platform_n); printf("platform number %u\n", platform_n); assert(platform_n >= 1); /* Get a valid platform */ - CALL (clGetPlatformIDs, 1, &platform, &platform_n); - CALL (clGetPlatformInfo, platform, CL_PLATFORM_PROFILE, sizeof(name), name, NULL); + OCL_CALL (clGetPlatformIDs, 1, &platform, &platform_n); + OCL_CALL (clGetPlatformInfo, platform, CL_PLATFORM_PROFILE, sizeof(name), name, NULL); printf("platform_profile \"%s\"\n", name); - CALL (clGetPlatformInfo, platform, CL_PLATFORM_NAME, sizeof(name), name, NULL); + OCL_CALL (clGetPlatformInfo, platform, CL_PLATFORM_NAME, sizeof(name), name, NULL); printf("platform_name \"%s\"\n", name); - CALL (clGetPlatformInfo, platform, CL_PLATFORM_VENDOR, sizeof(name), name, NULL); + OCL_CALL (clGetPlatformInfo, platform, CL_PLATFORM_VENDOR, sizeof(name), name, NULL); printf("platform_vendor \"%s\"\n", name); - CALL (clGetPlatformInfo, platform, CL_PLATFORM_VERSION, sizeof(name), name, NULL); + OCL_CALL (clGetPlatformInfo, platform, CL_PLATFORM_VERSION, sizeof(name), name, NULL); printf("platform_version \"%s\"\n", name); /* Get the device (only GPU device is supported right now) */ - CALL (clGetDeviceIDs, platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL); - CALL (clGetDeviceInfo, device, CL_DEVICE_PROFILE, sizeof(name), name, NULL); + OCL_CALL (clGetDeviceIDs, platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL); + OCL_CALL (clGetDeviceInfo, device, CL_DEVICE_PROFILE, sizeof(name), name, NULL); printf("device_profile \"%s\"\n", name); - CALL (clGetDeviceInfo, device, CL_DEVICE_NAME, sizeof(name), name, NULL); + OCL_CALL (clGetDeviceInfo, device, CL_DEVICE_NAME, sizeof(name), name, NULL); printf("device_name \"%s\"\n", name); - CALL (clGetDeviceInfo, device, CL_DEVICE_VENDOR, sizeof(name), name, NULL); + OCL_CALL (clGetDeviceInfo, device, CL_DEVICE_VENDOR, sizeof(name), name, NULL); printf("device_vendor \"%s\"\n", name); - CALL (clGetDeviceInfo, device, CL_DEVICE_VERSION, sizeof(name), name, NULL); + OCL_CALL (clGetDeviceInfo, device, CL_DEVICE_VERSION, sizeof(name), name, NULL); printf("device_version \"%s\"\n", name); /* Now create a context */ @@ -260,8 +260,8 @@ error: void cl_kernel_destroy(void) { - clReleaseKernel(kernel); - clReleaseProgram(program); + if (kernel) clReleaseKernel(kernel); + if (program) clReleaseProgram(program); kernel = NULL; program = NULL; } @@ -283,26 +283,19 @@ cl_test_destroy(void) } void -cl_release_buffers(void) +cl_buffer_destroy(void) { int i; - for (i = 0; i < MAX_BUFFER_N; ++i) + for (i = 0; i < MAX_BUFFER_N; ++i) { + if (buf_data[i] != NULL) { + clIntelUnmapBuffer(buf[i]); + buf_data[i] = NULL; + } if (buf[i] != NULL) { clReleaseMemObject(buf[i]); buf[i] = NULL; } -} - -void -cl_report_error(cl_int err) -{ - if (err > 0) - return; - if ((uint32_t) -err > err_msg_n) - return; - if (err == CL_SUCCESS) - return; - fprintf(stderr, "error %s\n", err_msg[-err]); + } } void diff --git a/utests/utest_helper.hpp b/utests/utest_helper.hpp index f014b37..bf0c3ab 100644 --- a/utests/utest_helper.hpp +++ b/utests/utest_helper.hpp @@ -28,47 +28,58 @@ #include "CL/cl.h" #include "CL/cl_intel.h" #include "utest.hpp" +#include "utest_assert.hpp" +#include "utest_error.h" #include #include #include -#define CALL(FN, ...) \ +#define OCL_THROW_ERROR(FN, STATUS) \ do { \ - status = FN(__VA_ARGS__); \ - if (status != CL_SUCCESS) { \ - fprintf(stderr, "error calling %s \n", #FN); \ - goto error; \ - } \ + char msg[2048]; \ + sprintf(msg, "error calling %s with error%s \n", #FN, err_msg[-STATUS]); \ + OCL_ASSERTM(false, msg); \ + } while (0) + +#define OCL_CALL(FN, ...) \ + do { \ + int status = FN(__VA_ARGS__); \ + if (status != CL_SUCCESS) OCL_THROW_ERROR(FN, status); \ + } while (0) + +#define OCL_CREATE_KERNEL(NAME) \ + do { \ + OCL_CALL (cl_kernel_init, NAME".cl", NAME, SOURCE); \ } while (0) #define OCL_CREATE_BUFFER(BUFFER, FLAGS, SIZE, DATA) \ do { \ cl_int status; \ BUFFER = clCreateBuffer(ctx, FLAGS, SIZE, DATA, &status); \ - if (status != CL_SUCCESS) goto error; \ + if (status != CL_SUCCESS) OCL_THROW_ERROR(FN, status); \ } while (0) #define OCL_MAP_BUFFER(ID) \ do { \ cl_int status; \ buf_data[ID] = (int *) clIntelMapBuffer(buf[ID], &status); \ - if (status != CL_SUCCESS) goto error; \ + if (status != CL_SUCCESS) OCL_THROW_ERROR(FN, status); \ } while (0) #define OCL_UNMAP_BUFFER(ID) \ - do { \ - CALL (clIntelUnmapBuffer, buf[ID]); \ - buf_data[ID] = NULL; \ - } while (0) + do { \ + OCL_CALL (clIntelUnmapBuffer, buf[ID]); \ + buf_data[ID] = NULL; \ + } while (0) #define OCL_NDRANGE(DIM_N) \ do { \ - CALL (clEnqueueNDRangeKernel, queue, kernel, DIM_N, NULL, globals, locals, 0, NULL, NULL); \ + OCL_CALL (clEnqueueNDRangeKernel, queue, kernel, DIM_N, NULL, globals, locals, 0, NULL, NULL); \ } while (0) #define OCL_SET_ARG(ID, SIZE, ARG) \ do { \ - CALL (clSetKernelArg, kernel, ID, SIZE, ARG); \ + OCL_CALL (clSetKernelArg, kernel, ID, SIZE, ARG); \ } while (0) enum { MAX_BUFFER_N = 16 }; @@ -78,10 +89,10 @@ extern cl_context ctx; extern cl_program program; extern cl_kernel kernel; extern cl_command_queue queue; -extern cl_mem buf[MAX_BUFFER_N]; // initialized at NULL -extern void* buf_data[MAX_BUFFER_N]; // initialized at NULL -extern size_t globals[3]; // initialized at zero -extern size_t locals[3]; // initialized at zero +extern cl_mem buf[MAX_BUFFER_N]; +extern void* buf_data[MAX_BUFFER_N]; +extern size_t globals[3]; +extern size_t locals[3]; enum { SOURCE = 0, @@ -98,8 +109,8 @@ extern int cl_kernel_init(const char *file_name, const char *kernel_name, int fo /* init the bunch of global varaibles here */ extern int cl_test_init(const char *file_name, const char *kernel_name, int format); -/* Release all the created buffers */ -extern void cl_release_buffers(void); +/* Unmap and release all the created buffers */ +extern void cl_buffer_destroy(void); /* Release OCL queue, context and device */ extern void cl_ocl_destroy(void); @@ -110,9 +121,6 @@ extern void cl_kernel_destroy(void); /* Release everything allocated in cl_test_init */ extern void cl_test_destroy(void); -/* Properly report the error in stderr */ -extern void cl_report_error(cl_int err); - /* Nicely output the performance counters */ extern void cl_report_perf_counters(cl_mem perf); diff --git a/utests/utest_run.cpp b/utests/utest_run.cpp index 9721c75..e577b7b 100644 --- a/utests/utest_run.cpp +++ b/utests/utest_run.cpp @@ -23,14 +23,22 @@ * * Just run the unit tests. The user can possibly provides the subset of it */ -#include "utest.hpp" +#include "utest_helper.hpp" +#include "utest_exception.hpp" +#include int main(int argc, char *argv[]) { - if (argc >= 2) - for (int i = 1; i < argc; ++i) - UTest::run(argv[i]); - else - UTest::runAll(); + try { + cl_ocl_init(); + if (argc >= 2) + for (int i = 1; i < argc; ++i) + UTest::run(argv[i]); + else + UTest::runAll(); + cl_ocl_destroy(); + } catch (Exception e) { + std::cout << " " << e.what() << " [SUCCESS]" << std::endl; + } } -- 2.7.4