/*
* Copyright © 2012 Intel Corporation
*
* This library is free software; you can redistribute it and/or
* modify it under the terms of the GNU Lesser General Public
* License as published by the Free Software Foundation; either
* version 2.1 of the License, or (at your option) any later version.
*
* This library is distributed in the hope that it will be useful,
* but WITHOUT ANY WARRANTY; without even the implied warranty of
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
* Lesser General Public License for more details.
*
* You should have received a copy of the GNU Lesser General Public
* License along with this library. If not, see .
*
* Author: Benjamin Segovia
*/
/**
* \file utest_helper.hpp
*
* \author Benjamin Segovia
*/
#ifndef __UTEST_HELPER_HPP__
#define __UTEST_HELPER_HPP__
#include "CL/cl.h"
#include "CL/cl_ext.h"
#include "CL/cl_intel.h"
#include "utest.hpp"
#include "utest_assert.hpp"
#include "utest_error.h"
#include
#include
#include
#if defined(__ANDROID__)
#define __thread
#endif
#ifdef HAS_GL_EGL_X11
#define EGL_WINDOW_WIDTH 256
#define EGL_WINDOW_HEIGHT 256
#define GL_GLEXT_PROTOTYPES
#include
#include
#include
#include
#include
extern EGLDisplay eglDisplay;
extern EGLContext eglContext;
extern EGLSurface eglSurface;
extern Display *xDisplay;
extern Window xWindow;
#endif
union uint32_cast {
uint32_t _uint;
float _float;
};
#define OCL_THROW_ERROR(FN, STATUS) \
do { \
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, NULL); \
} while (0)
#define OCL_DESTROY_KERNEL_KEEP_PROGRAM(KEEP_PROGRAM) \
do { \
cl_kernel_destroy(!(KEEP_PROGRAM)); \
} while(0)
#define OCL_CREATE_KERNEL_FROM_FILE(FILE_NAME, KERNEL_NAME) \
do { \
OCL_CALL(cl_kernel_init, FILE_NAME".cl", KERNEL_NAME, SOURCE, NULL); \
} while (0)
#define OCL_FLUSH() \
do { \
OCL_CALL(clFlush, queue); \
} while(0)
#define OCL_FINISH() \
do { \
OCL_CALL(clFinish, queue); \
} while(0)
#define OCL_CALL2(FN, RET, ...) \
do { \
cl_int status; \
RET = FN(__VA_ARGS__, &status);\
if (status != CL_SUCCESS) OCL_THROW_ERROR(FN, status); \
} while (0)
#define OCL_CREATE_BUFFER(BUFFER, FLAGS, SIZE, DATA) \
OCL_CALL2(clCreateBuffer, BUFFER, ctx, FLAGS, SIZE, DATA)
#define OCL_CREATE_USER_EVENT(EVENT) \
OCL_CALL2(clCreateUserEvent, EVENT, ctx)
#define OCL_SET_USER_EVENT_STATUS(EVENT, STATUS) \
OCL_CALL(clSetUserEventStatus, EVENT, STATUS)
#define OCL_CREATE_IMAGE(IMAGE, FLAGS, FORMAT, DESC, DATA) \
OCL_CALL2(clCreateImage, IMAGE, ctx, FLAGS, FORMAT, DESC, DATA)
#define OCL_READ_IMAGE(IMAGE, ORIGIN, REGION, DATA) \
OCL_CALL(clEnqueueReadImage, queue, IMAGE, CL_TRUE, ORIGIN, REGION, 0, 0, DATA, 0, NULL, NULL)
#define OCL_WRITE_IMAGE(IMAGE, ORIGIN, REGION, DATA) \
OCL_CALL(clEnqueueWriteImage, queue, IMAGE, CL_TRUE, ORIGIN, REGION, 0, 0, DATA, 0, NULL, NULL)
#define OCL_CREATE_GL_IMAGE(IMAGE, FLAGS, TARGET, LEVEL, TEXTURE) \
OCL_CALL2(clCreateFromGLTexture, IMAGE, ctx, FLAGS, TARGET, LEVEL, TEXTURE)
#define OCL_ENQUEUE_ACQUIRE_GL_OBJECTS(ID) \
OCL_CALL(clEnqueueAcquireGLObjects, queue, 1, &buf[ID], 0, 0, 0)
#define OCL_ENQUEUE_RELEASE_GL_OBJECTS(ID) \
OCL_CALL(clEnqueueReleaseGLObjects, queue, 1, &buf[ID], 0, 0, 0)
#define OCL_SWAP_EGL_BUFFERS() \
eglSwapBuffers(eglDisplay, eglSurface);
#define OCL_CREATE_SAMPLER(SAMPLER, ADDRESS_MODE, FILTER_MODE) \
OCL_CALL2(clCreateSampler, SAMPLER, ctx, 0, ADDRESS_MODE, FILTER_MODE)
#define OCL_CALL_MAP(FN, ID, RET, ...) \
do { \
cl_int status; \
size_t size = 0; \
status = clGetMemObjectInfo(buf[ID], CL_MEM_SIZE, sizeof(size), &size, NULL);\
if (status != CL_SUCCESS) OCL_THROW_ERROR(FN, status); \
RET = FN(__VA_ARGS__, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, size, 0, NULL, NULL, &status);\
if (status != CL_SUCCESS) OCL_THROW_ERROR(FN, status); \
} while (0)
#define OCL_MAP_BUFFER(ID) \
OCL_CALL_MAP(clEnqueueMapBuffer, ID, buf_data[ID], queue, buf[ID])
#define OCL_UNMAP_BUFFER(ID) \
do { \
if (buf[ID] != NULL) { \
OCL_CALL (clEnqueueUnmapMemObject, queue, buf[ID], buf_data[ID], 0, NULL, NULL); \
buf_data[ID] = NULL; \
} \
} while (0)
#define OCL_CALL_MAP_GTT(FN, ID, RET, ...) \
do { \
cl_int status; \
size_t image_row_pitch = 0; \
status = clGetImageInfo(buf[ID], CL_IMAGE_ROW_PITCH, sizeof(image_row_pitch), &image_row_pitch, NULL);\
if (status != CL_SUCCESS) OCL_THROW_ERROR(FN, status); \
size_t image_slice_pitch = 0; \
status = clGetImageInfo(buf[ID], CL_IMAGE_ROW_PITCH, sizeof(image_slice_pitch), &image_slice_pitch, NULL);\
if (status != CL_SUCCESS) OCL_THROW_ERROR(FN, status); \
size_t image_width = 0; \
status = clGetImageInfo(buf[ID], CL_IMAGE_WIDTH, sizeof(image_width), &image_width, NULL);\
if (status != CL_SUCCESS) OCL_THROW_ERROR(FN, status); \
size_t image_height = 0; \
status = clGetImageInfo(buf[ID], CL_IMAGE_HEIGHT, sizeof(image_height), &image_height, NULL);\
if (status != CL_SUCCESS) OCL_THROW_ERROR(FN, status); \
size_t image_depth= 0; \
status = clGetImageInfo(buf[ID], CL_IMAGE_DEPTH, sizeof(image_depth), &image_depth, NULL);\
if (status != CL_SUCCESS) OCL_THROW_ERROR(FN, status); \
if(image_depth == 0) image_depth = 1; \
if(image_height == 0) image_height = 1; \
size_t origin[3] = {0, 0, 0}; \
size_t region[3] = {image_width, image_height, image_depth}; \
RET = FN(__VA_ARGS__, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, origin, region, &image_row_pitch, &image_slice_pitch, 0, NULL, NULL, &status);\
if (status != CL_SUCCESS) OCL_THROW_ERROR(FN, status); \
} while (0)
#define OCL_MAP_BUFFER_GTT(ID) \
OCL_CALL_MAP_GTT(clEnqueueMapImage, ID, buf_data[ID], queue, buf[ID])
#define OCL_UNMAP_BUFFER_GTT(ID) \
do { \
if (buf[ID] != NULL) { \
OCL_CALL (clEnqueueUnmapMemObject, queue, buf[ID], buf_data[ID], 0, NULL, NULL); \
buf_data[ID] = NULL; \
} \
} while (0)
#define OCL_NDRANGE(DIM_N) \
OCL_CALL (clEnqueueNDRangeKernel, queue, kernel, DIM_N, NULL, globals, locals, 0, NULL, NULL)
#define OCL_SET_ARG(ID, SIZE, ARG) \
OCL_CALL (clSetKernelArg, kernel, ID, SIZE, ARG)
#define OCL_CHECK_IMAGE(DATA, W, H, FILENAME) \
if (cl_check_image(DATA, W, H, FILENAME) == 0) \
OCL_ASSERTM(false, "image mismatch")
enum { MAX_BUFFER_N = 16 };
extern cl_platform_id platform;
extern cl_device_id device;
extern cl_context ctx;
extern __thread cl_program program;
extern __thread cl_kernel kernel;
extern cl_command_queue queue;
extern __thread cl_mem buf[MAX_BUFFER_N];
extern __thread void* buf_data[MAX_BUFFER_N];
extern __thread size_t globals[3];
extern __thread size_t locals[3];
extern float ULPSIZE_FAST_MATH;
enum {
SOURCE = 0,
LLVM = 1,
BIN = 2
};
/* The SF is float type spliter*/
typedef struct
{
unsigned int mantissa:23;
unsigned int exponent:8;
unsigned int sign:1;
} FLOAT;
typedef union
{
float f;
unsigned int i;
FLOAT spliter;
} SF;
/* Init OpenCL */
extern int cl_ocl_init(void);
/* Init program and kernel for the test */
extern int cl_kernel_init(const char *file_name,
const char *kernel_name, int format, const char * build_opt);
extern int cl_kernel_compile(const char *file_name, const char *kernel_name,
const char * compile_opt);
extern int cl_kernel_link(const char *file_name, const char *kernel_name,
const char * link_opt);
/* Get the file path */
extern char* cl_do_kiss_path(const char *file, cl_device_id device);
/* init the bunch of global varaibles here */
extern int cl_test_init(const char *file_name, const char *kernel_name, int format);
/* 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);
/* Release kernel and program */
extern void cl_kernel_destroy(bool needDestroyProgram = true);
/* Release everything allocated in cl_test_init */
extern void cl_test_destroy(void);
/* Nicely output the performance counters */
extern void cl_report_perf_counters(cl_mem perf);
/* Read a bmp from file */
extern int *cl_read_bmp(const char *filename, int *width, int *height);
/* Write a bmp to a file */
extern void cl_write_bmp(const int *data, int width, int height, const char *filename);
/* Check data from img against bmp file located at "bmp" */
extern int cl_check_image(const int *img, int w, int h, const char *bmp);
/* Calculator ULP of each FLOAT value */
extern float cl_FLT_ULP(float float_number);
/* Calculator ULP of each INT value */
extern int cl_INT_ULP(int int_number);
/* subtract the time */
double time_subtract(struct timeval *y, struct timeval *x, struct timeval *result);
/* check ulpsize */
float select_ulpsize(float ULPSIZE_FAST_MATH, float ULPSIZE_NO_FAST_MATH);
/* Check is FP64 enabled. */
extern int cl_check_double(void);
/* Check is beignet device. */
extern int cl_check_beignet(void);
/* Check is intel subgroups enabled. */
extern int cl_check_subgroups(void);
/* Check is intel_media_block_io enabled. */
extern int cl_check_media_block_io(void);
typedef cl_int(clGetKernelSubGroupInfoKHR_cb)(cl_kernel, cl_device_id,
cl_kernel_sub_group_info, size_t,
const void *, size_t, void *,
size_t *);
extern clGetKernelSubGroupInfoKHR_cb* utestclGetKernelSubGroupInfoKHR;
/* Check if cl_intel_motion_estimation enabled. */
extern int cl_check_motion_estimation(void);
/* Check is cl version 2.0 or Beignet extension. */
extern int cl_check_ocl20(bool or_beignet = true);
/* Check is FP16 enabled. */
extern int cl_check_half(void);
/* Helper function for half type numbers */
extern uint32_t __half_to_float(uint16_t h, bool* isInf = NULL, bool* infSign = NULL);
extern uint16_t __float_to_half(uint32_t x);
extern float as_float(uint32_t i);
extern uint32_t as_uint(float f);
/* Check is intel subgroups short enabled. */
extern int cl_check_subgroups_short(void);
#endif /* __UTEST_HELPER_HPP__ */