| // |
| // 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_PS_CTORS_DTORS_COMMON_HPP |
| #define TEST_CONFORMANCE_CLCPP_PS_CTORS_DTORS_COMMON_HPP |
| |
| #include <algorithm> |
| |
| #include "../common.hpp" |
| #include "../funcs_test_utils.hpp" |
| |
| #define RUN_PS_CTORS_DTORS_TEST_MACRO(TEST_CLASS) \ |
| last_error = run_ps_ctor_dtor_test( \ |
| device, context, queue, count, TEST_CLASS \ |
| ); \ |
| CHECK_ERROR(last_error) \ |
| error |= last_error; |
| |
| // Base class for all tests for kernels with program scope object with |
| // non-trivial ctors and/or dtors |
| struct ps_ctors_dtors_test_base : public detail::base_func_type<cl_uint> |
| { |
| // ctor is true, if and only if OpenCL program of this test contains program |
| // scope variable with non-trivial ctor. |
| // dtor is true, if and only if OpenCL program of this test contains program |
| // scope variable with non-trivial dtor. |
| ps_ctors_dtors_test_base(const bool ctor, |
| const bool dtor) |
| : m_ctor(ctor), m_dtor(dtor) |
| { |
| |
| } |
| virtual ~ps_ctors_dtors_test_base() { }; |
| // Returns test name |
| virtual std::string str() = 0; |
| // Returns OpenCL program source |
| virtual std::string generate_program() = 0; |
| // Returns kernel names IN ORDER |
| virtual std::vector<std::string> get_kernel_names() |
| { |
| // Typical case, that is, only one kernel |
| return { this->get_kernel_name() }; |
| } |
| // Returns value that is expected to be in output_buffer[i] |
| virtual cl_uint operator()(size_t i) = 0; |
| // Executes kernels |
| // Typical case: execute every kernel once, every kernel has only |
| // one argument, that is, output buffer |
| virtual cl_int execute(const std::vector<cl_kernel>& kernels, |
| cl_mem& output_buffer, |
| cl_command_queue& queue, |
| size_t work_size) |
| { |
| cl_int err; |
| for(auto& k : kernels) |
| { |
| err = clSetKernelArg(k, 0, sizeof(output_buffer), &output_buffer); |
| RETURN_ON_CL_ERROR(err, "clSetKernelArg"); |
| |
| err = clEnqueueNDRangeKernel( |
| queue, k, 1, |
| NULL, &work_size, NULL, |
| 0, NULL, NULL |
| ); |
| RETURN_ON_CL_ERROR(err, "clEnqueueNDRangeKernel"); |
| } |
| return err; |
| } |
| // This method check if queries for CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT |
| // and CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT using clGetProgramInfo() |
| // return correct values |
| virtual cl_int ctors_dtors_present_queries(cl_program program) |
| { |
| cl_int error = CL_SUCCESS; |
| #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS) |
| return error; |
| #else |
| // CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT cl_bool |
| // This indicates that the program object contains non-trivial constructor(s) that will be |
| // executed by runtime before any kernel from the program is executed. |
| |
| // CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT cl_bool |
| // This indicates that the program object contains non-trivial destructor(s) that will be |
| // executed by runtime when program is destroyed. |
| |
| // CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT |
| cl_bool ctors_present; |
| size_t cl_bool_size; |
| error = clGetProgramInfo( |
| program, |
| CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT, |
| sizeof(cl_bool), |
| static_cast<void*>(&ctors_present), |
| &cl_bool_size |
| ); |
| RETURN_ON_CL_ERROR(error, "clGetProgramInfo") |
| if(cl_bool_size != sizeof(cl_bool)) |
| { |
| error = -1; |
| CHECK_ERROR_MSG( |
| error, |
| "Test failed, param_value_size_ret != sizeof(cl_bool) (%lu != %lu).\n", |
| cl_bool_size, |
| sizeof(cl_bool) |
| ); |
| } |
| |
| // CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT |
| cl_bool dtors_present = 0; |
| error = clGetProgramInfo( |
| program, |
| CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT, |
| sizeof(cl_bool), |
| static_cast<void*>(&ctors_present), |
| &cl_bool_size |
| ); |
| RETURN_ON_CL_ERROR(error, "clGetProgramInfo") |
| if(cl_bool_size != sizeof(cl_bool)) |
| { |
| error = -1; |
| CHECK_ERROR_MSG( |
| error, |
| "Test failed, param_value_size_ret != sizeof(cl_bool) (%lu != %lu).\n", |
| cl_bool_size, |
| sizeof(cl_bool) |
| ); |
| } |
| |
| // check constructors |
| if(m_ctor && ctors_present != CL_TRUE) |
| { |
| error = -1; |
| CHECK_ERROR_MSG( |
| error, |
| "Test failed, CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT: 0, should be: 1.\n" |
| ); |
| } |
| else if(!m_ctor && ctors_present == CL_TRUE) |
| { |
| error = -1; |
| CHECK_ERROR_MSG( |
| error, |
| "Test failed, CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT: 1, should be: 0.\n" |
| ); |
| } |
| |
| // check destructors |
| if(m_dtor && dtors_present != CL_TRUE) |
| { |
| error = -1; |
| CHECK_ERROR_MSG( |
| error, |
| "Test failed, CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT: 0, should be: 1.\n" |
| ); |
| } |
| else if(!m_dtor && dtors_present == CL_TRUE) |
| { |
| error = -1; |
| CHECK_ERROR_MSG( |
| error, |
| "Test failed, CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT: 1, should be: 0.\n" |
| ); |
| } |
| return error; |
| #endif |
| } |
| |
| private: |
| bool m_ctor; |
| bool m_dtor; |
| }; |
| |
| template <class ps_ctor_dtor_test> |
| int run_ps_ctor_dtor_test(cl_device_id device, cl_context context, cl_command_queue queue, size_t count, ps_ctor_dtor_test op) |
| { |
| cl_mem buffers[1]; |
| cl_program program; |
| std::vector<cl_kernel> kernels; |
| size_t work_size[1]; |
| cl_int err; |
| |
| std::string code_str = op.generate_program(); |
| std::vector<std::string> kernel_names = op.get_kernel_names(); |
| if(kernel_names.empty()) |
| { |
| RETURN_ON_ERROR_MSG(-1, "No kernel to run"); |
| } |
| kernels.resize(kernel_names.size()); |
| |
| // ----------------------------------------------------------------------------------- |
| // ------------- 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, &(kernels[0]), code_str, kernel_names[0]); |
| 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, &(kernels[0]), code_str, kernel_names[0], "-cl-std=CL2.0", false); |
| RETURN_ON_ERROR(err) |
| for(size_t i = 1; i < kernels.size(); i++) |
| { |
| kernels[i] = clCreateKernel(program, kernel_names[i].c_str(), &err); |
| RETURN_ON_CL_ERROR(err, "clCreateKernel"); |
| } |
| #else |
| err = create_opencl_kernel(context, &program, &(kernels[0]), code_str, kernel_names[0]); |
| RETURN_ON_ERROR(err) |
| for(size_t i = 1; i < kernels.size(); i++) |
| { |
| kernels[i] = clCreateKernel(program, kernel_names[i].c_str(), &err); |
| RETURN_ON_CL_ERROR(err, "clCreateKernel"); |
| } |
| #endif |
| |
| work_size[0] = count; |
| // host output vector |
| std::vector<cl_uint> output = generate_output<cl_uint>(work_size[0], 9999); |
| |
| // device output buffer |
| buffers[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, |
| sizeof(cl_uint) * output.size(), NULL, &err); |
| RETURN_ON_CL_ERROR(err, "clCreateBuffer") |
| |
| // Execute test |
| err = op.execute(kernels, buffers[0], queue, work_size[0]); |
| RETURN_ON_ERROR(err) |
| |
| // Check if queries returns correct values |
| err = op.ctors_dtors_present_queries(program); |
| RETURN_ON_ERROR(err); |
| |
| // Release kernels and program |
| // Destructors should be called now |
| for(auto& k : kernels) |
| { |
| err = clReleaseKernel(k); |
| RETURN_ON_CL_ERROR(err, "clReleaseKernel"); |
| } |
| err = clReleaseProgram(program); |
| RETURN_ON_CL_ERROR(err, "clReleaseProgram"); |
| |
| // Finish |
| err = clFinish(queue); |
| RETURN_ON_CL_ERROR(err, "clFinish"); |
| |
| err = clEnqueueReadBuffer( |
| queue, buffers[0], CL_TRUE, 0, sizeof(cl_uint) * output.size(), |
| static_cast<void *>(output.data()), 0, NULL, NULL |
| ); |
| RETURN_ON_CL_ERROR(err, "clEnqueueReadBuffer"); |
| |
| // Check output values |
| for(size_t i = 0; i < output.size(); i++) |
| { |
| cl_uint v = op(i); |
| if(!(are_equal(v, output[i], detail::make_value<cl_uint>(0), op))) |
| { |
| RETURN_ON_ERROR_MSG(-1, |
| "test_%s(%s) failed. Expected: %s, got: %s", op.str().c_str(), type_name<cl_uint>().c_str(), |
| format_value(v).c_str(), format_value(output[i]).c_str() |
| ); |
| } |
| } |
| log_info("test_%s(%s) passed\n", op.str().c_str(), type_name<cl_uint>().c_str()); |
| |
| clReleaseMemObject(buffers[0]); |
| return err; |
| } |
| |
| #endif // TEST_CONFORMANCE_CLCPP_PS_CTORS_DTORS_COMMON_HPP |