#ifndef TEST_CONFORMANCE_CLCPP_UTILS_TEST_BINARY_HPP
#define TEST_CONFORMANCE_CLCPP_UTILS_TEST_BINARY_HPP
#include <type_traits>
#include <algorithm>
#include <string>
#include <cmath>
#include "../common.hpp"
#include "detail/base_func_type.hpp"
#include "generate_inputs.hpp"
#include "compare.hpp"
template<class IN1, class IN2, class OUT1>
struct binary_func : public detail::base_func_type<OUT1>
{
typedef IN1 in1_type;
typedef IN2 in2_type;
typedef OUT1 out_type;
virtual ~binary_func() {};
virtual std::string str() = 0;
std::string decl_str()
{
return type_name<OUT1>() + "(" + type_name<IN1>() + ", " + type_name<IN2>() + ")";
}
bool is_in1_bool()
{
return false;
}
bool is_in2_bool()
{
return false;
}
IN1 min1()
{
return detail::get_min<IN1>();
}
IN1 max1()
{
return detail::get_max<IN1>();
}
IN2 min2()
{
return detail::get_min<IN2>();
}
IN2 max2()
{
return detail::get_max<IN2>();
}
std::vector<IN1> in1_special_cases()
{
return { };
}
std::vector<IN2> in2_special_cases()
{
return { };
}
template<class T>
typename make_vector_type<cl_double, vector_size<T>::value>::type
delta(const IN1& in1, const IN2& in2, const T& expected)
{
typedef
typename make_vector_type<cl_double, vector_size<T>::value>::type
delta_vector_type;
(void) in1;
(void) in2;
auto e = detail::make_value<delta_vector_type>(1e-3);
return detail::multiply<delta_vector_type>(e, expected);
}
};
#if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
template <class func_type, class in1_type, class in2_type, class out_type>
std::string generate_kernel_binary(func_type func)
{
std::string in1_value = "input1[gid]";
if(func.is_in1_bool())
{
std::string i = vector_size<in1_type>::value == 1 ? "" : std::to_string(vector_size<in1_type>::value);
in1_value = "(input1[gid] != (int" + i + ")(0))";
}
std::string in2_value = "input2[gid]";
if(func.is_in2_bool())
{
std::string i = vector_size<in2_type>::value == 1 ? "" : std::to_string(vector_size<in2_type>::value);
in2_value = "(input2[gid] != (int" + i + ")(0))";
}
std::string function_call = func.str() + "(" + in1_value + ", " + in2_value + ")";
if(func.is_out_bool())
{
std::string i = vector_size<out_type>::value == 1 ? "" : std::to_string(vector_size<out_type>::value);
function_call = "convert_int" + i + "(" + func.str() + "(" + in1_value + ", " + in2_value + "))";
}
return
"__kernel void " + func.get_kernel_name() + "(global " + type_name<in1_type>() + " *input1,\n"
" global " + type_name<in2_type>() + " *input2,\n"
" global " + type_name<out_type>() + " *output)\n"
"{\n"
" size_t gid = get_global_id(0);\n"
" output[gid] = " + function_call + ";\n"
"}\n";
}
#else
template <class func_type, class in1_type, class in2_type, class out_type>
std::string generate_kernel_binary(func_type func)
{
std::string headers = func.headers();
std::string in1_value = "input1[gid]";
if(func.is_in1_bool())
{
std::string i = vector_size<in1_type>::value == 1 ? "" : std::to_string(vector_size<in1_type>::value);
in1_value = "(input1[gid] != (int" + i + ")(0))";
}
std::string in2_value = "input2[gid]";
if(func.is_in2_bool())
{
std::string i = vector_size<in2_type>::value == 1 ? "" : std::to_string(vector_size<in2_type>::value);
in2_value = "(input2[gid] != (int" + i + ")(0))";
}
std::string function_call = func.str() + "(" + in1_value + ", " + in2_value + ")";
if(func.is_out_bool())
{
std::string i = vector_size<out_type>::value == 1 ? "" : std::to_string(vector_size<out_type>::value);
function_call = "convert_cast<int" + i + ">(" + func.str() + "(" + in1_value + ", " + in2_value + "))";
}
if(func.is_out_bool() || func.is_in1_bool() || func.is_in2_bool())
{
if(headers.find("#include <opencl_convert>") == std::string::npos)
{
headers += "#include <opencl_convert>\n";
}
}
return
"" + func.defs() +
"" + headers +
"#include <opencl_memory>\n"
"#include <opencl_work_item>\n"
"using namespace cl;\n"
"__kernel void " + func.get_kernel_name() + "(global_ptr<" + type_name<in1_type>() + "[]> input1,\n"
" global_ptr<" + type_name<in2_type>() + "[]> input2,\n"
" global_ptr<" + type_name<out_type>() + "[]> output)\n"
"{\n"
" size_t gid = get_global_id(0);\n"
" output[gid] = " + function_call + ";\n"
"}\n";
}
#endif
template<class INPUT1, class INPUT2, class OUTPUT, class binary_op>
bool verify_binary(const std::vector<INPUT1> &in1,
const std::vector<INPUT2> &in2,
const std::vector<OUTPUT> &out,
binary_op op)
{
for(size_t i = 0; i < in1.size(); i++)
{
auto expected = op(in1[i], in2[i]);
if(!are_equal(expected, out[i], op.delta(in1[i], in2[i], expected), op))
{
print_error_msg(expected, out[i], i, op);
return false;
}
}
return true;
}
template <class binary_op>
int test_binary_func(cl_device_id device, cl_context context, cl_command_queue queue, size_t count, binary_op op)
{
cl_mem buffers[3];
cl_program program;
cl_kernel kernel;
size_t work_size[1];
int err;
typedef typename binary_op::in1_type INPUT1;
typedef typename binary_op::in2_type INPUT2;
typedef typename binary_op::out_type OUTPUT;
if(!(type_supported<INPUT1>(device)
&& type_supported<INPUT2>(device)
&& type_supported<OUTPUT>(device)))
{
return CL_SUCCESS;
}
std::string code_str = generate_kernel_binary<binary_op, INPUT1, INPUT2, OUTPUT>(op);
std::string kernel_name = op.get_kernel_name();
#if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name);
RETURN_ON_ERROR(err)
return err;
#elif defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name, "-cl-std=CL2.0", false);
RETURN_ON_ERROR(err)
#else
err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name);
RETURN_ON_ERROR(err)
#endif
std::vector<INPUT1> in1_spec_cases = op.in1_special_cases();
std::vector<INPUT2> in2_spec_cases = op.in2_special_cases();
prepare_special_cases(in1_spec_cases, in2_spec_cases);
std::vector<INPUT1> input1 = generate_input<INPUT1>(count, op.min1(), op.max1(), in1_spec_cases);
std::vector<INPUT2> input2 = generate_input<INPUT2>(count, op.min2(), op.max2(), in2_spec_cases);
std::vector<OUTPUT> output = generate_output<OUTPUT>(count);
buffers[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
sizeof(INPUT1) * input1.size(), NULL, &err);
RETURN_ON_CL_ERROR(err, "clCreateBuffer")
buffers[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
sizeof(INPUT2) * input2.size(), NULL, &err);
RETURN_ON_CL_ERROR(err, "clCreateBuffer")
buffers[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
sizeof(OUTPUT) * output.size(), NULL, &err);
RETURN_ON_CL_ERROR(err, "clCreateBuffer")
err = clEnqueueWriteBuffer(
queue, buffers[0], CL_TRUE, 0, sizeof(INPUT1) * input1.size(),
static_cast<void *>(input1.data()), 0, NULL, NULL
);
RETURN_ON_CL_ERROR(err, "clEnqueueWriteBuffer")
err = clEnqueueWriteBuffer(
queue, buffers[1], CL_TRUE, 0, sizeof(INPUT2) * input2.size(),
static_cast<void *>(input2.data()), 0, NULL, NULL
);
RETURN_ON_CL_ERROR(err, "clEnqueueWriteBuffer")
err = clSetKernelArg(kernel, 0, sizeof(buffers[0]), &buffers[0]);
err |= clSetKernelArg(kernel, 1, sizeof(buffers[1]), &buffers[1]);
err |= clSetKernelArg(kernel, 2, sizeof(buffers[2]), &buffers[2]);
RETURN_ON_CL_ERROR(err, "clSetKernelArg");
work_size[0] = count;
err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, work_size, NULL, 0, NULL, NULL);
RETURN_ON_CL_ERROR(err, "clEnqueueNDRangeKernel");
err = clEnqueueReadBuffer(
queue, buffers[2], CL_TRUE, 0, sizeof(OUTPUT) * output.size(),
static_cast<void *>(output.data()), 0, NULL, NULL
);
RETURN_ON_CL_ERROR(err, "clEnqueueReadBuffer");
if (!verify_binary(input1, input2, output, op))
{
RETURN_ON_ERROR_MSG(-1,
"test_%s %s(%s, %s) failed", op.str().c_str(),
type_name<OUTPUT>().c_str(), type_name<INPUT1>().c_str(), type_name<INPUT2>().c_str()
);
}
log_info(
"test_%s %s(%s, %s) passed\n", op.str().c_str(),
type_name<OUTPUT>().c_str(), type_name<INPUT1>().c_str(), type_name<INPUT2>().c_str()
);
clReleaseMemObject(buffers[0]);
clReleaseMemObject(buffers[1]);
clReleaseMemObject(buffers[2]);
clReleaseKernel(kernel);
clReleaseProgram(program);
return err;
}
#endif