// // 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 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 T clamp(T x, T a, T b) { return (std::min)(b, (std::max)(a, x)); } template struct convert_cast : public unary_func { static_assert(vector_size::value == vector_size::value, "The operand and result type must have the same number of elements"); typedef typename scalar_type::type in_scalar_type; typedef typename scalar_type::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 \n"; } IN1 min1() { return detail::def_limit(in_min); } IN1 max1() { return detail::def_limit(in_max); } OUT1 operator()(const IN1& x) { OUT1 y; for (size_t i = 0; i < vector_size::value; i++) { in_scalar_type v; if (smode == saturate::on) v = clamp(x.s[i], static_cast((std::numeric_limits::min)()), static_cast((std::numeric_limits::max)()) ); else v = x.s[i]; if (std::is_integral::value) { switch (rmode) { case rounding_mode::rtp: y.s[i] = static_cast(std::ceil(v)); break; case rounding_mode::rtn: y.s[i] = static_cast(std::floor(v)); break; default: y.s[i] = static_cast(v); } } else { y.s[i] = static_cast(v); } } return y; } }; // ----------------------------------------------------------------------------------- // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------ // ----------------------------------------------------------------------------------- #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS) template std::string generate_kernel_convert_cast(func_type func) { std::string in1_value = "input[gid]"; std::string function_call = "convert_" + type_name(); 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() + " *input, global " + type_name() + " *output)\n" "{\n" " size_t gid = get_global_id(0);\n" " output[gid] = " + function_call + ";\n" "}\n"; } #else template 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(); 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 \n" "#include \n" "using namespace cl;\n" "__kernel void test_" + func.str() + "(global_ptr<" + type_name() + "[]> input," "global_ptr<" + type_name() + "[]> output)\n" "{\n" " size_t gid = get_global_id(0);\n" " output[gid] = " + function_call + ";\n" "}\n"; } #endif template 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(device) && type_supported(device))) { return CL_SUCCESS; } std::string code_str = generate_kernel_convert_cast(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 = generate_input(count, op.min1(), op.max1(), op.in_special_cases()); std::vector output = generate_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(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(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().c_str(), type_name().c_str()); } log_info("test_%s %s(%s) passed\n", op.str().c_str(), type_name().c_str(), type_name().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(-100.0f, +100.0f, rounding_mode::rtn, saturate::def))) TEST_CONVERT_CAST_MACRO((convert_cast(0, 255, rounding_mode::def, saturate::def))) // int to int TEST_CONVERT_CAST_MACRO((convert_cast(40000, 40000, rounding_mode::def, saturate::on))) TEST_CONVERT_CAST_MACRO((convert_cast(0, 127, rounding_mode::def, saturate::off))) TEST_CONVERT_CAST_MACRO((convert_cast(-100, 100, rounding_mode::def, saturate::off))) // float to int TEST_CONVERT_CAST_MACRO((convert_cast(-100.0f, +400.0f, rounding_mode::def, saturate::on))) TEST_CONVERT_CAST_MACRO((convert_cast(-127.0, +127.0, rounding_mode::rtp, saturate::off))) TEST_CONVERT_CAST_MACRO((convert_cast(-1000.0f, +10000.0f, rounding_mode::rtp, saturate::on))) TEST_CONVERT_CAST_MACRO((convert_cast(-10000.0f, +70000.0f, rounding_mode::rtn, saturate::on))) // int to float TEST_CONVERT_CAST_MACRO((convert_cast(0, 12345, rounding_mode::def, saturate::def))) TEST_CONVERT_CAST_MACRO((convert_cast(-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