From f0f887d07f4bae1f6f6f7f009ceebf6d8a93bab2 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Tue, 4 Apr 2023 09:12:05 +0200 Subject: [PATCH] printf test cases for cl_khr_command_buffer extenstion (#1592) * Added printf test case for cl_khr_command_buffer extenstion according to point 1.6 from: https://github.com/KhronosGroup/OpenCL-CTS/issues/1369 * Corrections related to macos and windows build * Added missing header * Added corrections related to code review of draft PR: https://github.com/KhronosGroup/OpenCL-CTS/pull/1592 refers to issue #1369, p.1.6: https://github.com/KhronosGroup/OpenCL-CTS/issues/1369 * minor fix for previous commit * printf capability condition moved to ::Skip method (#1369, p.1.6): https://github.com/KhronosGroup/OpenCL-CTS/issues/1369 * Applied minor corrections related to code review (#1369, p.1.6): https://github.com/KhronosGroup/OpenCL-CTS/issues/1369 * Added correction to replace string literals with sequence of characters printf arguments (issue #1369, printf) * Added clang formatting. * Added header related to android build (issue #1369, printf) * Fixed memory leak * Fixed memory leak in test_printf * Added cosmetic fix for last commit --- test_common/harness/os_helpers.cpp | 32 ++ test_common/harness/os_helpers.h | 2 + .../cl_khr_command_buffer/CMakeLists.txt | 1 + .../command_buffer_printf.cpp | 537 ++++++++++++++++++ .../extensions/cl_khr_command_buffer/main.cpp | 2 + .../extensions/cl_khr_command_buffer/procs.h | 4 + test_conformance/printf/test_printf.cpp | 36 +- 7 files changed, 589 insertions(+), 25 deletions(-) create mode 100644 test_conformance/extensions/cl_khr_command_buffer/command_buffer_printf.cpp diff --git a/test_common/harness/os_helpers.cpp b/test_common/harness/os_helpers.cpp index 628a206e..3989edf6 100644 --- a/test_common/harness/os_helpers.cpp +++ b/test_common/harness/os_helpers.cpp @@ -30,8 +30,17 @@ #if defined(__ANDROID__) #include +#include "harness/mt19937.h" #endif +#if !defined(_WIN32) +#if defined(__APPLE__) +#include +#endif +#include +#endif + + #define CHECK_PTR(ptr) \ if ((ptr) == NULL) \ { \ @@ -556,4 +565,27 @@ char* get_exe_dir() } // get_exe_dir +char* get_temp_filename() +{ + char gFileName[256] = ""; + // Create a unique temporary file to allow parallel executed tests. +#if (defined(__linux__) || defined(__APPLE__)) && (!defined(__ANDROID__)) + sprintf(gFileName, "/tmp/tmpfile.XXXXXX"); + int fd = mkstemp(gFileName); + if (fd == -1) return strdup(gFileName); + close(fd); +#elif defined(_WIN32) + UINT ret = GetTempFileName(".", "tmp", 0, gFileName); + if (ret == 0) return gFileName; +#else + MTdata d = init_genrand((cl_uint)time(NULL)); + sprintf(gFileName, "tmpfile.%u", genrand_int32(d)); +#endif + + char* fn = strdup(gFileName); + CHECK_PTR(fn); + return fn; +} + + // end of file // diff --git a/test_common/harness/os_helpers.h b/test_common/harness/os_helpers.h index aa3080d9..0ab8507f 100644 --- a/test_common/harness/os_helpers.h +++ b/test_common/harness/os_helpers.h @@ -41,5 +41,7 @@ char* get_err_msg(int err); // Returns system error message. Subject to free. char* get_dir_sep(); // Returns dir separator. Subject to free. char* get_exe_path(); // Returns path of current executable. Subject to free. char* get_exe_dir(); // Returns dir of current executable. Subject to free. +char* get_temp_filename(); // returns temporary file name + #endif // __os_helpers_h__ diff --git a/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt b/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt index 9f162f69..4b9968c3 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt +++ b/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt @@ -3,6 +3,7 @@ set(MODULE_NAME CL_KHR_COMMAND_BUFFER) set(${MODULE_NAME}_SOURCES main.cpp basic_command_buffer.cpp + command_buffer_printf.cpp command_buffer_get_command_buffer_info.cpp command_buffer_set_kernel_arg.cpp command_buffer_event_sync.cpp diff --git a/test_conformance/extensions/cl_khr_command_buffer/command_buffer_printf.cpp b/test_conformance/extensions/cl_khr_command_buffer/command_buffer_printf.cpp new file mode 100644 index 00000000..eef3e355 --- /dev/null +++ b/test_conformance/extensions/cl_khr_command_buffer/command_buffer_printf.cpp @@ -0,0 +1,537 @@ +// +// Copyright (c) 2022 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. +// + +#include + +#include "basic_command_buffer.h" +#include "procs.h" + +#if !defined(_WIN32) +#if defined(__APPLE__) +#include +#endif +#include +#define streamDup(fd1) dup(fd1) +#define streamDup2(fd1, fd2) dup2(fd1, fd2) +#endif +#include +#include + +#if defined(_WIN32) +#include +#define streamDup(fd1) _dup(fd1) +#define streamDup2(fd1, fd2) _dup2(fd1, fd2) +#endif + +#include +#include +#include +#include +#include + +namespace { + +//////////////////////////////////////////////////////////////////////////////// +// printf tests for cl_khr_command_buffer which handles below cases: +// -test cases for device side printf +// -test cases for device side printf with a simultaneous use command-buffer + +template +struct CommandBufferPrintfTest : public BasicCommandBufferTest +{ + CommandBufferPrintfTest(cl_device_id device, cl_context context, + cl_command_queue queue) + : BasicCommandBufferTest(device, context, queue), + trigger_event(nullptr), wait_event(nullptr), file_descriptor(0), + printf_use_support(false) + { + simultaneous_use_requested = simul_use; + if (simul_use) + { + buffer_size_multiplier = num_test_iters; + } + } + + //-------------------------------------------------------------------------- + void ReleaseOutputStream(int fd) + { + fflush(stdout); + streamDup2(fd, fileno(stdout)); + close(fd); + } + + //-------------------------------------------------------------------------- + int AcquireOutputStream(int* error) + { + int fd = streamDup(fileno(stdout)); + *error = 0; + if (!freopen(temp_filename.c_str(), "wt", stdout)) + { + ReleaseOutputStream(fd); + *error = -1; + } + return fd; + } + + //-------------------------------------------------------------------------- + void GetAnalysisBuffer(std::stringstream& buffer) + { + std::ifstream fp(temp_filename, std::ios::in); + if (fp.is_open()) + { + buffer << fp.rdbuf(); + } + } + + //-------------------------------------------------------------------------- + void PurgeTempFile() + { + std::ofstream ofs(temp_filename, + std::ofstream::out | std::ofstream::trunc); + ofs.close(); + } + + //-------------------------------------------------------------------------- + bool Skip() override + { + // Query if device supports kernel printf use + cl_device_command_buffer_capabilities_khr capabilities; + cl_int error = + clGetDeviceInfo(device, CL_DEVICE_COMMAND_BUFFER_CAPABILITIES_KHR, + sizeof(capabilities), &capabilities, NULL); + test_error(error, + "Unable to query CL_DEVICE_COMMAND_BUFFER_CAPABILITIES_KHR"); + + printf_use_support = + (capabilities & CL_COMMAND_BUFFER_CAPABILITY_KERNEL_PRINTF_KHR) + != 0; + + if (!printf_use_support) return true; + return BasicCommandBufferTest::Skip() + || (simultaneous_use_requested && !simultaneous_use_support); + } + + //-------------------------------------------------------------------------- + cl_int SetUpKernel() override + { + cl_int error = CL_SUCCESS; + + const char* kernel_str = + R"( + __kernel void print(__global char* in, __global char* out, __global int* offset) + { + size_t id = get_global_id(0); + int ind = offset[0] + offset[1] * id; + for(int i=0; i& pattern, + std::vector& output_data) + { + cl_int error = CL_SUCCESS; + auto in_mem_size = sizeof(cl_char) * pattern.size(); + error = clEnqueueWriteBuffer(queue, in_mem, CL_TRUE, 0, in_mem_size, + &pattern[0], 0, nullptr, nullptr); + test_error(error, "clEnqueueWriteBuffer failed"); + + cl_int offset[] = { 0, pattern.size() - 1 }; + error = clEnqueueWriteBuffer(queue, off_mem, CL_TRUE, 0, sizeof(offset), + offset, 0, nullptr, nullptr); + test_error(error, "clEnqueueWriteBuffer failed"); + + // redirect output stream to temporary file + file_descriptor = AcquireOutputStream(&error); + if (error != 0) + { + log_error("Error while redirection stdout to file"); + return TEST_FAIL; + } + + // enqueue command buffer with kernel containing printf command + error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0, + nullptr, &wait_event); + test_error_release_stdout(error, "clEnqueueCommandBufferKHR failed"); + + fflush(stdout); + + // Wait until kernel finishes its execution and (thus) the output + // printed from the kernel is immediately printed + error = clWaitForEvents(1, &wait_event); + test_error(error, "clWaitForEvents failed"); + + // output buffer contains pattern to be compared with printout + error = clEnqueueReadBuffer(queue, out_mem, CL_FALSE, 0, data_size(), + output_data.data(), 0, nullptr, nullptr); + test_error_release_stdout(error, "clEnqueueReadBuffer failed"); + + error = clFinish(queue); + test_error_release_stdout(error, "clFinish failed"); + + ReleaseOutputStream(file_descriptor); + + // copy content of temporary file into string stream + std::stringstream sstr; + GetAnalysisBuffer(sstr); + if (sstr.str().size() != num_elements * offset[1]) + { + log_error("GetAnalysisBuffer failed\n"); + return TEST_FAIL; + } + + // verify the result - compare printout and output buffer + for (size_t i = 0; i < num_elements * offset[1]; i++) + { + CHECK_VERIFICATION_ERROR(sstr.str().at(i), output_data[i], i); + } + + return CL_SUCCESS; + } + + //-------------------------------------------------------------------------- + cl_int RunSingle() + { + cl_int error = CL_SUCCESS; + std::vector output_data(num_elements * max_pattern_length); + + for (unsigned i = 0; i < num_test_iters; i++) + { + unsigned pattern_length = + std::max(min_pattern_length, rand() % max_pattern_length); + char pattern_character = 'a' + rand() % 26; + std::vector pattern(pattern_length + 1, pattern_character); + pattern[pattern_length] = '\0'; + error = EnqueueSinglePass(pattern, output_data); + test_error(error, "EnqueueSinglePass failed"); + + output_data.assign(output_data.size(), 0); + PurgeTempFile(); + } + + return CL_SUCCESS; + } + + //-------------------------------------------------------------------------- + struct SimulPassData + { + // null terminated character buffer + std::vector pattern; + // 0-command buffer offset, 1-pattern offset + cl_int offset[2]; + std::vector output_buffer; + }; + + //-------------------------------------------------------------------------- + cl_int EnqueueSimultaneousPass(SimulPassData& pd) + { + // write current pattern to device memory + auto in_mem_size = sizeof(cl_char) * pd.pattern.size(); + cl_int error = + clEnqueueWriteBuffer(queue, in_mem, CL_FALSE, 0, in_mem_size, + &pd.pattern[0], 0, nullptr, nullptr); + test_error_release_stdout(error, "clEnqueueWriteBuffer failed"); + + // refresh offsets for current enqueuing + error = + clEnqueueWriteBuffer(queue, off_mem, CL_FALSE, 0, sizeof(pd.offset), + pd.offset, 0, nullptr, nullptr); + test_error_release_stdout(error, "clEnqueueWriteBuffer failed"); + + // create user event to block simultaneous command buffers + if (!trigger_event) + { + trigger_event = clCreateUserEvent(context, &error); + test_error_release_stdout(error, "clCreateUserEvent failed"); + } + + error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 1, + &trigger_event, nullptr); + test_error_release_stdout(error, "clEnqueueCommandBufferKHR failed"); + + // output buffer contains pattern to be compared with printout + error = clEnqueueReadBuffer( + queue, out_mem, CL_FALSE, pd.offset[0] * sizeof(cl_char), + pd.output_buffer.size() * sizeof(cl_char), pd.output_buffer.data(), + 0, nullptr, nullptr); + test_error_release_stdout(error, "clEnqueueReadBuffer failed"); + + return CL_SUCCESS; + } + + + //-------------------------------------------------------------------------- + cl_int RunSimultaneous() + { + cl_int error = CL_SUCCESS; + cl_int offset = static_cast(num_elements * max_pattern_length); + + std::vector simul_passes(num_test_iters); + + const int pattern_chars_range = 26; + std::list pattern_chars; + for (size_t i = 0; i < pattern_chars_range; i++) + pattern_chars.push_back(cl_char('a' + i)); + + test_assert_error(pattern_chars.size() >= num_test_iters, + "Number of simultaneous launches must be lower than " + "size of characters container"); + + cl_int total_pattern_coverage = 0; + for (unsigned i = 0; i < num_test_iters; i++) + { + // random character pattern unique for each iteration + auto it = pattern_chars.begin(); + std::advance(it, rand() % pattern_chars.size()); + char pattern_character = *it; + unsigned pattern_length = + std::max(min_pattern_length, rand() % max_pattern_length); + + std::vector pattern(pattern_length + 1, pattern_character); + pattern[pattern_length] = '\0'; + simul_passes[i] = { pattern, + { cl_int(i * offset), cl_int(pattern_length) }, + std::vector(num_elements + * pattern_length) }; + total_pattern_coverage += simul_passes[i].output_buffer.size(); + pattern_chars.erase(it); + }; + + // takeover stdout stream + file_descriptor = AcquireOutputStream(&error); + if (error != 0) + { + log_error("Error while redirection stdout to file"); + return TEST_FAIL; + } + + // enqueue read/write and command buffer operations + for (auto&& pass : simul_passes) + { + error = EnqueueSimultaneousPass(pass); + test_error_release_stdout(error, "EnqueueSimultaneousPass failed"); + } + + // execute command buffers + error = clSetUserEventStatus(trigger_event, CL_COMPLETE); + test_error_release_stdout(error, "clSetUserEventStatus failed"); + + // flush streams + fflush(stdout); + + // finish command queue + error = clFinish(queue); + test_error_release_stdout(error, "clFinish failed\n"); + + ReleaseOutputStream(file_descriptor); + + std::stringstream sstr; + GetAnalysisBuffer(sstr); + if (sstr.str().size() != total_pattern_coverage) + { + log_error("GetAnalysisBuffer failed\n"); + return TEST_FAIL; + } + + // verify the result - compare printout and output buffer + std::map counters_map; + for (int i = 0; i < total_pattern_coverage; i++) + counters_map[sstr.str().at(i)]++; + + if (counters_map.size() != simul_passes.size()) + { + log_error("printout inconsistent with input data\n"); + return TEST_FAIL; + } + + for (auto&& pass : simul_passes) + { + auto& res_data = pass.output_buffer; + + if (res_data.empty() + || res_data.size() != counters_map[res_data.front()]) + { + log_error("output buffer inconsistent with printout\n"); + return TEST_FAIL; + } + + // verify consistency of output buffer + for (size_t i = 0; i < res_data.size(); i++) + { + CHECK_VERIFICATION_ERROR(res_data.front(), res_data[i], i); + } + } + + return CL_SUCCESS; + } + + //-------------------------------------------------------------------------- + clEventWrapper trigger_event = nullptr; + clEventWrapper wait_event = nullptr; + + std::string temp_filename; + int file_descriptor; + + bool printf_use_support; + + // specifies max test length for printf pattern + const unsigned max_pattern_length = 6; + // specifies min test length for printf pattern + const unsigned min_pattern_length = 1; + // specifies number of command-buffer enqueue iterations + const unsigned num_test_iters = 3; +}; + +} // anonymous namespace + +int test_basic_printf(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest>(device, context, + queue, num_elements); +} + +int test_simultaneous_printf(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest>(device, context, queue, + num_elements); +} diff --git a/test_conformance/extensions/cl_khr_command_buffer/main.cpp b/test_conformance/extensions/cl_khr_command_buffer/main.cpp index 43d2725d..4eefc8ab 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/main.cpp +++ b/test_conformance/extensions/cl_khr_command_buffer/main.cpp @@ -50,6 +50,8 @@ test_definition test_list[] = { ADD_TEST(copy_image_to_buffer), ADD_TEST(copy_buffer_rect), ADD_TEST(barrier_wait_list), + ADD_TEST(basic_printf), + ADD_TEST(simultaneous_printf), ADD_TEST(basic_set_kernel_arg), ADD_TEST(pending_set_kernel_arg), ADD_TEST(event_info_command_type), diff --git a/test_conformance/extensions/cl_khr_command_buffer/procs.h b/test_conformance/extensions/cl_khr_command_buffer/procs.h index b54b01e7..63e004a7 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/procs.h +++ b/test_conformance/extensions/cl_khr_command_buffer/procs.h @@ -29,6 +29,10 @@ extern int test_explicit_flush(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements); extern int test_out_of_order(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements); +extern int test_basic_printf(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_simultaneous_printf(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements); extern int test_info_queues(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements); extern int test_info_ref_count(cl_device_id device, cl_context context, diff --git a/test_conformance/printf/test_printf.cpp b/test_conformance/printf/test_printf.cpp index c4b6a0ba..e789e0ca 100644 --- a/test_conformance/printf/test_printf.cpp +++ b/test_conformance/printf/test_printf.cpp @@ -13,6 +13,7 @@ // See the License for the specific language governing permissions and // limitations under the License. // +#include "harness/os_helpers.h" #include #include @@ -113,29 +114,6 @@ static char gFileName[256]; // Static helper functions definition //----------------------------------------- -//----------------------------------------- -// getTempFileName -//----------------------------------------- -static int getTempFileName() -{ - // Create a unique temporary file to allow parallel executed tests. -#if (defined(__linux__) || defined(__APPLE__)) && (!defined( __ANDROID__ )) - sprintf(gFileName, "/tmp/tmpfile.XXXXXX"); - int fd = mkstemp(gFileName); - if (fd == -1) - return -1; - close(fd); -#elif defined(_WIN32) - UINT ret = GetTempFileName(".", "tmp", 0, gFileName); - if (ret == 0) - return -1; -#else - MTdata d = init_genrand((cl_uint)time(NULL)); - sprintf(gFileName, "tmpfile.%u", genrand_int32(d)); -#endif - return 0; -} - //----------------------------------------- // acquireOutputStream //----------------------------------------- @@ -1062,9 +1040,17 @@ int main(int argc, const char* argv[]) } } - if (getTempFileName() == -1) + char* pcTempFname = get_temp_filename(); + if (pcTempFname != nullptr) { - log_error("getTempFileName failed\n"); + strncpy(gFileName, pcTempFname, sizeof(gFileName)); + } + + free(pcTempFname); + + if (strlen(gFileName) == 0) + { + log_error("get_temp_filename failed\n"); return -1; }