// // 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_ADDRESS_SPACES_COMMON_HPP #define TEST_CONFORMANCE_CLCPP_ADDRESS_SPACES_COMMON_HPP #include "../common.hpp" #include "../funcs_test_utils.hpp" #define RUN_ADDRESS_SPACES_TEST_MACRO(TEST_CLASS) \ last_error = run_address_spaces_test( \ device, context, queue, n_elems, TEST_CLASS \ ); \ CHECK_ERROR(last_error) \ error |= last_error; // This is a base class for address spaces tests. template struct address_spaces_test : public detail::base_func_type { // output buffer type typedef T type; virtual ~address_spaces_test() {}; // 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 get_kernel_names() { // Typical case, that is, only one kernel return { this->get_kernel_name() }; } // Return value that is expected to be in output_buffer[i] virtual T operator()(size_t i, size_t work_group_size) = 0; // If local size has to be set in clEnqueueNDRangeKernel() // this should return true; otherwise - false; virtual bool set_local_size() { return false; } // Calculates maximal work-group size (one dim) virtual size_t get_max_local_size(const std::vector& kernels, cl_device_id device, size_t work_group_size, // default work-group size cl_int& error) { size_t wg_size = work_group_size; for(auto&k : kernels) { size_t max_wg_size; error = clGetKernelWorkGroupInfo(k, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &max_wg_size, NULL); RETURN_ON_CL_ERROR(error, "clGetKernelWorkGroupInfo") wg_size = (std::min)(max_wg_size, wg_size); } return wg_size; } // This covers typical case: each kernel is executed once, every kernel // has only one argument which is output buffer virtual cl_int execute(const std::vector& kernels, cl_mem& output_buffer, cl_command_queue& queue, size_t work_size, size_t work_group_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, this->set_local_size() ? &work_group_size : NULL, 0, NULL, NULL ); RETURN_ON_CL_ERROR(err, "clEnqueueNDRangeKernel"); } return err; } }; template int run_address_spaces_test(cl_device_id device, cl_context context, cl_command_queue queue, size_t count, address_spaces_test op) { cl_mem buffers[1]; cl_program program; std::vector kernels; size_t wg_size; size_t work_size[1]; cl_int err; typedef typename address_spaces_test::type TYPE; // Don't run test for unsupported types if(!(type_supported(device))) { return CL_SUCCESS; } std::string code_str = op.generate_program(); std::vector 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 // Find the max possible wg size for among all the kernels wg_size = op.get_max_local_size(kernels, device, 1024, err); RETURN_ON_ERROR(err); work_size[0] = count; if(op.set_local_size()) { size_t wg_number = static_cast( std::ceil(static_cast(count) / wg_size) ); work_size[0] = wg_number * wg_size; } // output on host std::vector output = generate_output(work_size[0], 9999); // output buffer buffers[0] = clCreateBuffer (context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(TYPE) * output.size(), NULL, &err ); RETURN_ON_CL_ERROR(err, "clCreateBuffer") // Execute test err = op.execute(kernels, buffers[0], queue, work_size[0], wg_size); RETURN_ON_ERROR(err) err = clEnqueueReadBuffer( queue, buffers[0], CL_TRUE, 0, sizeof(TYPE) * output.size(), static_cast(output.data()), 0, NULL, NULL ); RETURN_ON_CL_ERROR(err, "clEnqueueReadBuffer"); for(size_t i = 0; i < output.size(); i++) { TYPE v = op(i, wg_size); if(!(are_equal(v, output[i], detail::make_value(0), op))) { RETURN_ON_ERROR_MSG(-1, "test_%s(%s) failed. Expected: %s, got: %s", op.str().c_str(), type_name().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().c_str()); clReleaseMemObject(buffers[0]); for(auto& k : kernels) clReleaseKernel(k); clReleaseProgram(program); return err; } #endif // TEST_CONFORMANCE_CLCPP_ADDRESS_SPACES_COMMON_HPP