test_conformance/clcpp/convert/convert_cast.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_CONVERT_CONVERT_CAST_HPP
 #define TEST_CONFORMANCE_CLCPP_CONVERT_CONVERT_CAST_HPP

 #include "../common.hpp"
 #include "../funcs_test_utils.hpp"

 #include <functional>


 enum class rounding_mode
 {
     def,
     /*rte, not implemented here */
     rtz,
     rtp,
     rtn
 };

 enum class saturate { def, off, on };

 std::string rounding_mode_name(rounding_mode rmode)
 {
     switch (rmode)
     {
         case rounding_mode::rtz: return "rtz";
         case rounding_mode::rtp: return "rtp";
         case rounding_mode::rtn: return "rtn";
         default: return "";
     }
 }

 std::string saturate_name(saturate smode)
 {
     switch (smode)
     {
         case saturate::off: return "off";
         case saturate::on:  return "on";
         default: return "";
     }
 }

 template<class T>
 T clamp(T x, T a, T b)
 {
     return (std::min)(b, (std::max)(a, x));
 }

 template<class IN1, class OUT1>
 struct convert_cast : public unary_func<IN1, OUT1>
 {
     static_assert(vector_size<IN1>::value == vector_size<OUT1>::value, "The operand and result type must have the same number of elements");

     typedef typename scalar_type<IN1>::type in_scalar_type;
     typedef typename scalar_type<OUT1>::type out_scalar_type;

     in_scalar_type in_min;
     in_scalar_type in_max;
     rounding_mode rmode;
     saturate smode;

     convert_cast(in_scalar_type min, in_scalar_type max, rounding_mode rmode, saturate smode)
         : in_min(min), in_max(max), rmode(rmode), smode(smode)
     {
     }

     std::string str()
     {
         return "convert_cast";
     }

     std::string headers()
     {
         return "#include <opencl_convert>\n";
     }

     IN1 min1()
     {
         return detail::def_limit<IN1>(in_min);
     }

     IN1 max1()
     {
         return detail::def_limit<IN1>(in_max);
     }

     OUT1 operator()(const IN1& x)
     {
         OUT1 y;
         for (size_t i = 0; i < vector_size<IN1>::value; i++)
         {
             in_scalar_type v;
             if (smode == saturate::on)
                 v = clamp(x.s[i],
                     static_cast<in_scalar_type>((std::numeric_limits<out_scalar_type>::min)()),
                     static_cast<in_scalar_type>((std::numeric_limits<out_scalar_type>::max)())
                 );
             else
                 v = x.s[i];

             if (std::is_integral<out_scalar_type>::value)
             {
                 switch (rmode)
                 {
                     case rounding_mode::rtp:
                         y.s[i] = static_cast<out_scalar_type>(std::ceil(v));
                         break;
                     case rounding_mode::rtn:
                         y.s[i] = static_cast<out_scalar_type>(std::floor(v));
                         break;
                     default:
                         y.s[i] = static_cast<out_scalar_type>(v);
                 }
             }
             else
             {
                 y.s[i] = static_cast<out_scalar_type>(v);
             }
         }
         return y;
     }
 };

 // -----------------------------------------------------------------------------------
 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
 // -----------------------------------------------------------------------------------
 #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
 template <class func_type, class in_type, class out_type>
 std::string generate_kernel_convert_cast(func_type func)
 {
     std::string in1_value = "input[gid]";
     std::string function_call = "convert_" + type_name<out_type>();
     if (func.smode == saturate::on)
         function_call += "_sat";
     if (func.rmode != rounding_mode::def)
         function_call += "_" + rounding_mode_name(func.rmode);
     function_call += "(" + in1_value + ")";
     return
         "__kernel void test_" + func.str() + "(global " + type_name<in_type>() + " *input, 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 in_type, class out_type>
 std::string generate_kernel_convert_cast(func_type func)
 {
     std::string headers = func.headers();
     std::string in1_value = "input[gid]";
     std::string function_call = "convert_cast<" + type_name<out_type>();
     if (func.rmode != rounding_mode::def)
         function_call += ", rounding_mode::" + rounding_mode_name(func.rmode);
     if (func.smode != saturate::def)
         function_call += ", saturate::" + saturate_name(func.smode);
     function_call += ">(" + in1_value + ")";
     return
         "" + func.defs() +
         "" + headers +
         "#include <opencl_memory>\n"
         "#include <opencl_work_item>\n"
         "using namespace cl;\n"
         "__kernel void test_" + func.str() + "(global_ptr<" + type_name<in_type>() +  "[]> input,"
                                               "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 convert_cast_op>
 int test_convert_cast_func(cl_device_id device, cl_context context, cl_command_queue queue, size_t count, convert_cast_op op)
 {
     cl_mem buffers[2];
     cl_program program;
     cl_kernel kernel;
     size_t work_size[1];
     int error;

     typedef typename convert_cast_op::in_type INPUT;
     typedef typename convert_cast_op::out_type OUTPUT;

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

     std::string code_str = generate_kernel_convert_cast<convert_cast_op, INPUT, OUTPUT>(op);
     std::string kernel_name("test_"); kernel_name += op.str();

 // -----------------------------------------------------------------------------------
 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
 // -----------------------------------------------------------------------------------
 // Only OpenCL C++ to SPIR-V compilation
 #if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
     error = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name);
     RETURN_ON_ERROR(error)
     return error;
 // Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code)
 #elif defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
     error = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name, "-cl-std=CL2.0", false);
     RETURN_ON_ERROR(error)
 #else
     error = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name);
     RETURN_ON_ERROR(error)
 #endif

     std::vector<INPUT> input = generate_input<INPUT>(count, op.min1(), op.max1(), op.in_special_cases());
     std::vector<OUTPUT> output = generate_output<OUTPUT>(count);

     buffers[0] = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(INPUT) * input.size(), NULL, &error);
     RETURN_ON_CL_ERROR(error, "clCreateBuffer")

     buffers[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(OUTPUT) * output.size(), NULL, &error);
     RETURN_ON_CL_ERROR(error, "clCreateBuffer")

     error = clEnqueueWriteBuffer(
         queue, buffers[0], CL_TRUE, 0, sizeof(INPUT) * input.size(),
         static_cast<void *>(input.data()), 0, NULL, NULL
     );
     RETURN_ON_CL_ERROR(error, "clEnqueueWriteBuffer")

     error = clSetKernelArg(kernel, 0, sizeof(buffers[0]), &buffers[0]);
     RETURN_ON_CL_ERROR(error, "clSetKernelArg")
     error = clSetKernelArg(kernel, 1, sizeof(buffers[1]), &buffers[1]);
     RETURN_ON_CL_ERROR(error, "clSetKernelArg")

     work_size[0] = count;
     error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, work_size, NULL, 0, NULL, NULL);
     RETURN_ON_CL_ERROR(error, "clEnqueueNDRangeKernel")

     error = clEnqueueReadBuffer(
         queue, buffers[1], CL_TRUE, 0, sizeof(OUTPUT) * output.size(),
         static_cast<void *>(output.data()), 0, NULL, NULL
     );
     RETURN_ON_CL_ERROR(error, "clEnqueueReadBuffer")

     if (!verify_unary(input, output, op))
     {
         RETURN_ON_ERROR_MSG(-1, "test_%s %s(%s) failed", op.str().c_str(), type_name<OUTPUT>().c_str(), type_name<INPUT>().c_str());
     }
     log_info("test_%s %s(%s) passed\n", op.str().c_str(), type_name<OUTPUT>().c_str(), type_name<INPUT>().c_str());

     clReleaseMemObject(buffers[0]);
     clReleaseMemObject(buffers[1]);
     clReleaseKernel(kernel);
     clReleaseProgram(program);
     return error;
 }


 AUTO_TEST_CASE(test_convert_cast)
 (cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
 {
     int error = CL_SUCCESS;
     int last_error = CL_SUCCESS;

 #define TEST_CONVERT_CAST_MACRO(OP) \
     last_error = test_convert_cast_func( \
         device, context, queue, n_elems, OP \
     ); \
     CHECK_ERROR(last_error) \
     error |= last_error;

     // No-op
     TEST_CONVERT_CAST_MACRO((convert_cast<cl_float2, cl_float2>(-100.0f, +100.0f, rounding_mode::rtn, saturate::def)))
     TEST_CONVERT_CAST_MACRO((convert_cast<cl_uchar2, cl_uchar2>(0, 255, rounding_mode::def, saturate::def)))

     // int to int
     TEST_CONVERT_CAST_MACRO((convert_cast<cl_int4, cl_short4>(40000, 40000, rounding_mode::def, saturate::on)))
     TEST_CONVERT_CAST_MACRO((convert_cast<cl_uchar8, cl_char8>(0, 127, rounding_mode::def, saturate::off)))
     TEST_CONVERT_CAST_MACRO((convert_cast<cl_char8, cl_int8>(-100, 100, rounding_mode::def, saturate::off)))

     // float to int
     TEST_CONVERT_CAST_MACRO((convert_cast<cl_float2, cl_uchar2>(-100.0f, +400.0f, rounding_mode::def, saturate::on)))
     TEST_CONVERT_CAST_MACRO((convert_cast<cl_double4, cl_char4>(-127.0, +127.0, rounding_mode::rtp, saturate::off)))
     TEST_CONVERT_CAST_MACRO((convert_cast<cl_float8, cl_uint8>(-1000.0f, +10000.0f, rounding_mode::rtp, saturate::on)))
     TEST_CONVERT_CAST_MACRO((convert_cast<cl_float16, cl_ushort16>(-10000.0f, +70000.0f, rounding_mode::rtn, saturate::on)))

     // int to float
     TEST_CONVERT_CAST_MACRO((convert_cast<cl_short8, cl_float8>(0, 12345, rounding_mode::def, saturate::def)))
     TEST_CONVERT_CAST_MACRO((convert_cast<cl_long2, cl_float2>(-1000000, +1000000, rounding_mode::rtz, saturate::def)))

 #undef TEST_CONVERT_CAST_MACRO

     if (error != CL_SUCCESS)
     {
         return -1;
     }
     return error;
 }

 #endif // TEST_CONFORMANCE_CLCPP_CONVERT_CONVERT_CAST_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_CONVERT_CONVERT_CAST_HPP
	#define TEST_CONFORMANCE_CLCPP_CONVERT_CONVERT_CAST_HPP

	#include "../common.hpp"
	#include "../funcs_test_utils.hpp"

	#include <functional>


	enum class rounding_mode
	{
	def,
	/rte, not implemented here /
	rtz,
	rtp,
	rtn
	};

	enum class saturate { def, off, on };

	std::string rounding_mode_name(rounding_mode rmode)
	{
	switch (rmode)
	{
	case rounding_mode::rtz: return "rtz";
	case rounding_mode::rtp: return "rtp";
	case rounding_mode::rtn: return "rtn";
	default: return "";
	}
	}

	std::string saturate_name(saturate smode)
	{
	switch (smode)
	{
	case saturate::off: return "off";
	case saturate::on: return "on";
	default: return "";
	}
	}

	template<class T>
	T clamp(T x, T a, T b)
	{
	return (std::min)(b, (std::max)(a, x));
	}

	template<class IN1, class OUT1>
	struct convert_cast : public unary_func<IN1, OUT1>
	{
	static_assert(vector_size<IN1>::value == vector_size<OUT1>::value, "The operand and result type must have the same number of elements");

	typedef typename scalar_type<IN1>::type in_scalar_type;
	typedef typename scalar_type<OUT1>::type out_scalar_type;

	in_scalar_type in_min;
	in_scalar_type in_max;
	rounding_mode rmode;
	saturate smode;

	convert_cast(in_scalar_type min, in_scalar_type max, rounding_mode rmode, saturate smode)
	: in_min(min), in_max(max), rmode(rmode), smode(smode)
	{
	}

	std::string str()
	{
	return "convert_cast";
	}

	std::string headers()
	{
	return "#include <opencl_convert>\n";
	}

	IN1 min1()
	{
	return detail::def_limit<IN1>(in_min);
	}

	IN1 max1()
	{
	return detail::def_limit<IN1>(in_max);
	}

	OUT1 operator()(const IN1& x)
	{
	OUT1 y;
	for (size_t i = 0; i < vector_size<IN1>::value; i++)
	{
	in_scalar_type v;
	if (smode == saturate::on)
	v = clamp(x.s[i],
	static_cast<in_scalar_type>((std::numeric_limits<out_scalar_type>::min)()),
	static_cast<in_scalar_type>((std::numeric_limits<out_scalar_type>::max)())
	);
	else
	v = x.s[i];

	if (std::is_integral<out_scalar_type>::value)
	{
	switch (rmode)
	{
	case rounding_mode::rtp:
	y.s[i] = static_cast<out_scalar_type>(std::ceil(v));
	break;
	case rounding_mode::rtn:
	y.s[i] = static_cast<out_scalar_type>(std::floor(v));
	break;
	default:
	y.s[i] = static_cast<out_scalar_type>(v);
	}
	}
	else
	{
	y.s[i] = static_cast<out_scalar_type>(v);
	}
	}
	return y;
	}
	};

	// -----------------------------------------------------------------------------------
	// ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
	// -----------------------------------------------------------------------------------
	#if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
	template <class func_type, class in_type, class out_type>
	std::string generate_kernel_convert_cast(func_type func)
	{
	std::string in1_value = "input[gid]";
	std::string function_call = "convert_" + type_name<out_type>();
	if (func.smode == saturate::on)
	function_call += "_sat";
	if (func.rmode != rounding_mode::def)
	function_call += "_" + rounding_mode_name(func.rmode);
	function_call += "(" + in1_value + ")";
	return
	"__kernel void test_" + func.str() + "(global " + type_name<in_type>() + " input, 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 in_type, class out_type>
	std::string generate_kernel_convert_cast(func_type func)
	{
	std::string headers = func.headers();
	std::string in1_value = "input[gid]";
	std::string function_call = "convert_cast<" + type_name<out_type>();
	if (func.rmode != rounding_mode::def)
	function_call += ", rounding_mode::" + rounding_mode_name(func.rmode);
	if (func.smode != saturate::def)
	function_call += ", saturate::" + saturate_name(func.smode);
	function_call += ">(" + in1_value + ")";
	return
	"" + func.defs() +
	"" + headers +
	"#include <opencl_memory>\n"
	"#include <opencl_work_item>\n"
	"using namespace cl;\n"
	"__kernel void test_" + func.str() + "(global_ptr<" + type_name<in_type>() + "[]> input,"
	"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 convert_cast_op>
	int test_convert_cast_func(cl_device_id device, cl_context context, cl_command_queue queue, size_t count, convert_cast_op op)
	{
	cl_mem buffers[2];
	cl_program program;
	cl_kernel kernel;
	size_t work_size[1];
	int error;

	typedef typename convert_cast_op::in_type INPUT;
	typedef typename convert_cast_op::out_type OUTPUT;

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

	std::string code_str = generate_kernel_convert_cast<convert_cast_op, INPUT, OUTPUT>(op);
	std::string kernel_name("test_"); kernel_name += op.str();

	// -----------------------------------------------------------------------------------
	// ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
	// -----------------------------------------------------------------------------------
	// Only OpenCL C++ to SPIR-V compilation
	#if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
	error = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name);
	RETURN_ON_ERROR(error)
	return error;
	// Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code)
	#elif defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
	error = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name, "-cl-std=CL2.0", false);
	RETURN_ON_ERROR(error)
	#else
	error = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name);
	RETURN_ON_ERROR(error)
	#endif

	std::vector<INPUT> input = generate_input<INPUT>(count, op.min1(), op.max1(), op.in_special_cases());
	std::vector<OUTPUT> output = generate_output<OUTPUT>(count);

	buffers[0] = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(INPUT) * input.size(), NULL, &error);
	RETURN_ON_CL_ERROR(error, "clCreateBuffer")

	buffers[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(OUTPUT) * output.size(), NULL, &error);
	RETURN_ON_CL_ERROR(error, "clCreateBuffer")

	error = clEnqueueWriteBuffer(
	queue, buffers[0], CL_TRUE, 0, sizeof(INPUT) * input.size(),
	static_cast<void *>(input.data()), 0, NULL, NULL
	);
	RETURN_ON_CL_ERROR(error, "clEnqueueWriteBuffer")

	error = clSetKernelArg(kernel, 0, sizeof(buffers[0]), &buffers[0]);
	RETURN_ON_CL_ERROR(error, "clSetKernelArg")
	error = clSetKernelArg(kernel, 1, sizeof(buffers[1]), &buffers[1]);
	RETURN_ON_CL_ERROR(error, "clSetKernelArg")

	work_size[0] = count;
	error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, work_size, NULL, 0, NULL, NULL);
	RETURN_ON_CL_ERROR(error, "clEnqueueNDRangeKernel")

	error = clEnqueueReadBuffer(
	queue, buffers[1], CL_TRUE, 0, sizeof(OUTPUT) * output.size(),
	static_cast<void *>(output.data()), 0, NULL, NULL
	);
	RETURN_ON_CL_ERROR(error, "clEnqueueReadBuffer")

	if (!verify_unary(input, output, op))
	{
	RETURN_ON_ERROR_MSG(-1, "test_%s %s(%s) failed", op.str().c_str(), type_name<OUTPUT>().c_str(), type_name<INPUT>().c_str());
	}
	log_info("test_%s %s(%s) passed\n", op.str().c_str(), type_name<OUTPUT>().c_str(), type_name<INPUT>().c_str());

	clReleaseMemObject(buffers[0]);
	clReleaseMemObject(buffers[1]);
	clReleaseKernel(kernel);
	clReleaseProgram(program);
	return error;
	}


	AUTO_TEST_CASE(test_convert_cast)
	(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
	{
	int error = CL_SUCCESS;
	int last_error = CL_SUCCESS;

	#define TEST_CONVERT_CAST_MACRO(OP) \
	last_error = test_convert_cast_func( \
	device, context, queue, n_elems, OP \
	); \
	CHECK_ERROR(last_error) \
	error \|= last_error;

	// No-op
	TEST_CONVERT_CAST_MACRO((convert_cast<cl_float2, cl_float2>(-100.0f, +100.0f, rounding_mode::rtn, saturate::def)))
	TEST_CONVERT_CAST_MACRO((convert_cast<cl_uchar2, cl_uchar2>(0, 255, rounding_mode::def, saturate::def)))

	// int to int
	TEST_CONVERT_CAST_MACRO((convert_cast<cl_int4, cl_short4>(40000, 40000, rounding_mode::def, saturate::on)))
	TEST_CONVERT_CAST_MACRO((convert_cast<cl_uchar8, cl_char8>(0, 127, rounding_mode::def, saturate::off)))
	TEST_CONVERT_CAST_MACRO((convert_cast<cl_char8, cl_int8>(-100, 100, rounding_mode::def, saturate::off)))

	// float to int
	TEST_CONVERT_CAST_MACRO((convert_cast<cl_float2, cl_uchar2>(-100.0f, +400.0f, rounding_mode::def, saturate::on)))
	TEST_CONVERT_CAST_MACRO((convert_cast<cl_double4, cl_char4>(-127.0, +127.0, rounding_mode::rtp, saturate::off)))
	TEST_CONVERT_CAST_MACRO((convert_cast<cl_float8, cl_uint8>(-1000.0f, +10000.0f, rounding_mode::rtp, saturate::on)))
	TEST_CONVERT_CAST_MACRO((convert_cast<cl_float16, cl_ushort16>(-10000.0f, +70000.0f, rounding_mode::rtn, saturate::on)))

	// int to float
	TEST_CONVERT_CAST_MACRO((convert_cast<cl_short8, cl_float8>(0, 12345, rounding_mode::def, saturate::def)))
	TEST_CONVERT_CAST_MACRO((convert_cast<cl_long2, cl_float2>(-1000000, +1000000, rounding_mode::rtz, saturate::def)))

	#undef TEST_CONVERT_CAST_MACRO

	if (error != CL_SUCCESS)
	{
	return -1;
	}
	return error;
	}

	#endif // TEST_CONFORMANCE_CLCPP_CONVERT_CONVERT_CAST_HPP