test_conformance/clcpp/utils_test/ternary.hpp - external/github.com/KhronosGroup/OpenCL-CTS - Git at Google

 //
 // Copyright (c) 2017 The Khronos Group Inc.
 //
 // Licensed under the Apache License, Version 2.0 (the "License");
 // you may not use this file except in compliance with the License.
 // You may obtain a copy of the License at
 //
 //    http://www.apache.org/licenses/LICENSE-2.0
 //
 // Unless required by applicable law or agreed to in writing, software
 // distributed under the License is distributed on an "AS IS" BASIS,
 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 // See the License for the specific language governing permissions and
 // limitations under the License.
 //
 #ifndef TEST_CONFORMANCE_CLCPP_UTILS_TEST_TERNARY_HPP
 #define TEST_CONFORMANCE_CLCPP_UTILS_TEST_TERNARY_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 IN3, class OUT1>
 struct ternary_func : public detail::base_func_type<OUT1>
 {
     typedef IN1 in1_type;
     typedef IN2 in2_type;
     typedef IN3 in3_type;
     typedef OUT1 out_type;

     virtual ~ternary_func() {};
     virtual std::string str() = 0;

     std::string decl_str()
     {
         return type_name<OUT1>() + "(" + type_name<IN1>() + ", " + type_name<IN2>()+  ", " + type_name<IN3>() + ")";
     }

     bool is_in1_bool()
     {
         return false;
     }

     bool is_in2_bool()
     {
         return false;
     }

     bool is_in3_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>();
     }

     IN3 min3()
     {
         return detail::get_min<IN3>();
     }

     IN3 max3()
     {
         return detail::get_max<IN3>();
     }

     std::vector<IN1> in1_special_cases()
     {
         return { };
     }

     std::vector<IN2> in2_special_cases()
     {
         return { };
     }

     std::vector<IN3> in3_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 IN3& in3, const T& expected)
     {
         typedef
             typename make_vector_type<cl_double, vector_size<T>::value>::type
             delta_vector_type;
         // Take care of unused variable warning
         (void) in1;
         (void) in2;
         (void) in3;
         auto e = detail::make_value<delta_vector_type>(1e-3);
         return detail::multiply<delta_vector_type>(e, expected);
     }
 };

 // -----------------------------------------------------------------------------------
 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
 // -----------------------------------------------------------------------------------
 #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
 template <class func_type, class in1_type, class in2_type, class in3_type, class out_type>
 std::string generate_kernel_ternary(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 in3_value = "input3[gid]";
     if(func.is_in3_bool())
     {
         std::string i = vector_size<in3_type>::value == 1 ? "" : std::to_string(vector_size<in3_type>::value);
         in3_value = "(input3[gid] != (int" + i + ")(0))";
     }
     std::string function_call = func.str() + "(" + in1_value + ", " + in2_value + ", " + in3_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 + ", " + in3_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<in3_type>() + " *input3,\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 in3_type, class out_type>
 std::string generate_kernel_ternary(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 in3_value = "input3[gid]";
     if(func.is_in3_bool())
     {
         std::string i = vector_size<in3_type>::value == 1 ? "" : std::to_string(vector_size<in3_type>::value);
         in3_value = "(input3[gid] != (int" + i + ")(0))";
     }
     std::string function_call = func.str() + "(" + in1_value + ", " + in2_value + ", " + in3_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 + ", " + in3_value + "))";
     }
     if(func.is_out_bool() || func.is_in1_bool() || func.is_in2_bool() || func.is_in3_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<in3_type>() + "[]> input3,\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 INPUT3, class OUTPUT, class ternary_op>
 bool verify_ternary(const std::vector<INPUT1> &in1,
                     const std::vector<INPUT2> &in2,
                     const std::vector<INPUT3> &in3,
                     const std::vector<OUTPUT> &out,
                     ternary_op op)
 {
     for(size_t i = 0; i < in1.size(); i++)
     {
         auto expected = op(in1[i], in2[i], in3[i]);
         if(!are_equal(expected, out[i], op.delta(in1[i], in2[i], in3[i], expected), op))
         {
             print_error_msg(expected, out[i], i, op);
             return false;
         }
     }
     return true;
 }

 template <class ternary_op>
 int test_ternary_func(cl_device_id device, cl_context context, cl_command_queue queue, size_t count, ternary_op op)
 {
     cl_mem buffers[4];
     cl_program program;
     cl_kernel kernel;
     size_t work_size[1];
     int err;

     typedef typename ternary_op::in1_type INPUT1;
     typedef typename ternary_op::in2_type INPUT2;
     typedef typename ternary_op::in3_type INPUT3;
     typedef typename ternary_op::out_type OUTPUT;

     // Don't run test for unsupported types
     if(!(type_supported<INPUT1>(device)
          && type_supported<INPUT2>(device)
          && type_supported<INPUT3>(device)
          && type_supported<OUTPUT>(device)))
     {
         return CL_SUCCESS;
     }

     std::string code_str = generate_kernel_ternary<ternary_op, INPUT1, INPUT2, INPUT3, OUTPUT>(op);
     std::string kernel_name = op.get_kernel_name();

 // -----------------------------------------------------------------------------------
 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
 // -----------------------------------------------------------------------------------
 // Only OpenCL C++ to SPIR-V compilation
 #if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
     err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name);
     RETURN_ON_ERROR(err)
     return err;
 // Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code)
 #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();
     std::vector<INPUT3> in3_spec_cases = op.in3_special_cases();
     prepare_special_cases(in1_spec_cases, in2_spec_cases, in3_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<INPUT3> input3 = generate_input<INPUT3>(count, op.min3(), op.max3(), in3_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(INPUT3) * input3.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer")

     buffers[3] = 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 = clEnqueueWriteBuffer(
         queue, buffers[2], CL_TRUE, 0, sizeof(INPUT3) * input3.size(),
         static_cast<void *>(input3.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]);
     err |= clSetKernelArg(kernel, 3, sizeof(buffers[3]), &buffers[3]);
     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[3], CL_TRUE, 0, sizeof(OUTPUT) * output.size(),
         static_cast<void *>(output.data()), 0, NULL, NULL
     );
     RETURN_ON_CL_ERROR(err, "clEnqueueReadBuffer");

     if (!verify_ternary(input1, input2, input3, output, op))
     {
         RETURN_ON_ERROR_MSG(-1,
             "test_%s %s(%s, %s, %s) failed", op.str().c_str(),
             type_name<OUTPUT>().c_str(),
             type_name<INPUT1>().c_str(),
             type_name<INPUT2>().c_str(),
             type_name<INPUT3>().c_str()
         );
     }
     log_info(
         "test_%s %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(),
         type_name<INPUT3>().c_str()
     );

     clReleaseMemObject(buffers[0]);
     clReleaseMemObject(buffers[1]);
     clReleaseMemObject(buffers[2]);
     clReleaseMemObject(buffers[3]);
     clReleaseKernel(kernel);
     clReleaseProgram(program);
     return err;
 }

 #endif // TEST_CONFORMANCE_CLCPP_UTILS_TEST_TERNARY_HPP
	//
	// Copyright (c) 2017 The Khronos Group Inc.
	//
	// Licensed under the Apache License, Version 2.0 (the "License");
	// you may not use this file except in compliance with the License.
	// You may obtain a copy of the License at
	//
	// http://www.apache.org/licenses/LICENSE-2.0
	//
	// Unless required by applicable law or agreed to in writing, software
	// distributed under the License is distributed on an "AS IS" BASIS,
	// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
	// See the License for the specific language governing permissions and
	// limitations under the License.
	//
	#ifndef TEST_CONFORMANCE_CLCPP_UTILS_TEST_TERNARY_HPP
	#define TEST_CONFORMANCE_CLCPP_UTILS_TEST_TERNARY_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 IN3, class OUT1>
	struct ternary_func : public detail::base_func_type<OUT1>
	{
	typedef IN1 in1_type;
	typedef IN2 in2_type;
	typedef IN3 in3_type;
	typedef OUT1 out_type;

	virtual ~ternary_func() {};
	virtual std::string str() = 0;

	std::string decl_str()
	{
	return type_name<OUT1>() + "(" + type_name<IN1>() + ", " + type_name<IN2>()+ ", " + type_name<IN3>() + ")";
	}

	bool is_in1_bool()
	{
	return false;
	}

	bool is_in2_bool()
	{
	return false;
	}

	bool is_in3_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>();
	}

	IN3 min3()
	{
	return detail::get_min<IN3>();
	}

	IN3 max3()
	{
	return detail::get_max<IN3>();
	}

	std::vector<IN1> in1_special_cases()
	{
	return { };
	}

	std::vector<IN2> in2_special_cases()
	{
	return { };
	}

	std::vector<IN3> in3_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 IN3& in3, const T& expected)
	{
	typedef
	typename make_vector_type<cl_double, vector_size<T>::value>::type
	delta_vector_type;
	// Take care of unused variable warning
	(void) in1;
	(void) in2;
	(void) in3;
	auto e = detail::make_value<delta_vector_type>(1e-3);
	return detail::multiply<delta_vector_type>(e, expected);
	}
	};

	// -----------------------------------------------------------------------------------
	// ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
	// -----------------------------------------------------------------------------------
	#if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
	template <class func_type, class in1_type, class in2_type, class in3_type, class out_type>
	std::string generate_kernel_ternary(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 in3_value = "input3[gid]";
	if(func.is_in3_bool())
	{
	std::string i = vector_size<in3_type>::value == 1 ? "" : std::to_string(vector_size<in3_type>::value);
	in3_value = "(input3[gid] != (int" + i + ")(0))";
	}
	std::string function_call = func.str() + "(" + in1_value + ", " + in2_value + ", " + in3_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 + ", " + in3_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<in3_type>() + " *input3,\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 in3_type, class out_type>
	std::string generate_kernel_ternary(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 in3_value = "input3[gid]";
	if(func.is_in3_bool())
	{
	std::string i = vector_size<in3_type>::value == 1 ? "" : std::to_string(vector_size<in3_type>::value);
	in3_value = "(input3[gid] != (int" + i + ")(0))";
	}
	std::string function_call = func.str() + "(" + in1_value + ", " + in2_value + ", " + in3_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 + ", " + in3_value + "))";
	}
	if(func.is_out_bool() \|\| func.is_in1_bool() \|\| func.is_in2_bool() \|\| func.is_in3_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<in3_type>() + "[]> input3,\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 INPUT3, class OUTPUT, class ternary_op>
	bool verify_ternary(const std::vector<INPUT1> &in1,
	const std::vector<INPUT2> &in2,
	const std::vector<INPUT3> &in3,
	const std::vector<OUTPUT> &out,
	ternary_op op)
	{
	for(size_t i = 0; i < in1.size(); i++)
	{
	auto expected = op(in1[i], in2[i], in3[i]);
	if(!are_equal(expected, out[i], op.delta(in1[i], in2[i], in3[i], expected), op))
	{
	print_error_msg(expected, out[i], i, op);
	return false;
	}
	}
	return true;
	}

	template <class ternary_op>
	int test_ternary_func(cl_device_id device, cl_context context, cl_command_queue queue, size_t count, ternary_op op)
	{
	cl_mem buffers[4];
	cl_program program;
	cl_kernel kernel;
	size_t work_size[1];
	int err;

	typedef typename ternary_op::in1_type INPUT1;
	typedef typename ternary_op::in2_type INPUT2;
	typedef typename ternary_op::in3_type INPUT3;
	typedef typename ternary_op::out_type OUTPUT;

	// Don't run test for unsupported types
	if(!(type_supported<INPUT1>(device)
	&& type_supported<INPUT2>(device)
	&& type_supported<INPUT3>(device)
	&& type_supported<OUTPUT>(device)))
	{
	return CL_SUCCESS;
	}

	std::string code_str = generate_kernel_ternary<ternary_op, INPUT1, INPUT2, INPUT3, OUTPUT>(op);
	std::string kernel_name = op.get_kernel_name();

	// -----------------------------------------------------------------------------------
	// ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
	// -----------------------------------------------------------------------------------
	// Only OpenCL C++ to SPIR-V compilation
	#if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
	err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name);
	RETURN_ON_ERROR(err)
	return err;
	// Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code)
	#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();
	std::vector<INPUT3> in3_spec_cases = op.in3_special_cases();
	prepare_special_cases(in1_spec_cases, in2_spec_cases, in3_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<INPUT3> input3 = generate_input<INPUT3>(count, op.min3(), op.max3(), in3_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(INPUT3) * input3.size(), NULL, &err);
	RETURN_ON_CL_ERROR(err, "clCreateBuffer")

	buffers[3] = 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 = clEnqueueWriteBuffer(
	queue, buffers[2], CL_TRUE, 0, sizeof(INPUT3) * input3.size(),
	static_cast<void *>(input3.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]);
	err \|= clSetKernelArg(kernel, 3, sizeof(buffers[3]), &buffers[3]);
	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[3], CL_TRUE, 0, sizeof(OUTPUT) * output.size(),
	static_cast<void *>(output.data()), 0, NULL, NULL
	);
	RETURN_ON_CL_ERROR(err, "clEnqueueReadBuffer");

	if (!verify_ternary(input1, input2, input3, output, op))
	{
	RETURN_ON_ERROR_MSG(-1,
	"test_%s %s(%s, %s, %s) failed", op.str().c_str(),
	type_name<OUTPUT>().c_str(),
	type_name<INPUT1>().c_str(),
	type_name<INPUT2>().c_str(),
	type_name<INPUT3>().c_str()
	);
	}
	log_info(
	"test_%s %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(),
	type_name<INPUT3>().c_str()
	);

	clReleaseMemObject(buffers[0]);
	clReleaseMemObject(buffers[1]);
	clReleaseMemObject(buffers[2]);
	clReleaseMemObject(buffers[3]);
	clReleaseKernel(kernel);
	clReleaseProgram(program);
	return err;
	}

	#endif // TEST_CONFORMANCE_CLCPP_UTILS_TEST_TERNARY_HPP