From ffa75c37ce2c65217e2cea2fe473e91c0ef3dc57 Mon Sep 17 00:00:00 2001 From: Anastasia Stulova <38433336+AnastasiaStulova@users.noreply.github.com> Date: Tue, 12 Jan 2021 10:17:13 +0000 Subject: [PATCH] Test cl_ext_cxx_for_opencl (#1095) Add tests for API extension for compilation of kernels in C++ for OpenCL language. * Test that -cl-std=CLC++ is accepted and a basic kernel with C++ features is compiled. * Test that API extension reports the same language version as __OPENCL_CPP_VERSION__. This commit also adds a separate folder for extension tests. Signed-off-by: Victoria Holodovsky Co-authored-by: Victoria Holodovsky --- test_conformance/CMakeLists.txt | 1 + test_conformance/extensions/CMakeLists.txt | 1 + .../cl_ext_cxx_for_opencl/CMakeLists.txt | 9 ++ .../cxx_for_opencl_ext.cpp | 105 ++++++++++++++++++ .../cxx_for_opencl_ver.cpp | 102 +++++++++++++++++ .../extensions/cl_ext_cxx_for_opencl/main.cpp | 28 +++++ .../extensions/cl_ext_cxx_for_opencl/procs.h | 26 +++++ 7 files changed, 272 insertions(+) create mode 100644 test_conformance/extensions/CMakeLists.txt create mode 100644 test_conformance/extensions/cl_ext_cxx_for_opencl/CMakeLists.txt create mode 100644 test_conformance/extensions/cl_ext_cxx_for_opencl/cxx_for_opencl_ext.cpp create mode 100644 test_conformance/extensions/cl_ext_cxx_for_opencl/cxx_for_opencl_ver.cpp create mode 100644 test_conformance/extensions/cl_ext_cxx_for_opencl/main.cpp create mode 100644 test_conformance/extensions/cl_ext_cxx_for_opencl/procs.h diff --git a/test_conformance/CMakeLists.txt b/test_conformance/CMakeLists.txt index 6714e234..b9b87c1d 100644 --- a/test_conformance/CMakeLists.txt +++ b/test_conformance/CMakeLists.txt @@ -21,6 +21,7 @@ if(D3D11_IS_SUPPORTED) endif(D3D11_IS_SUPPORTED) add_subdirectory( device_partition ) add_subdirectory( events ) +add_subdirectory( extensions ) add_subdirectory( geometrics ) if(GL_IS_SUPPORTED) add_subdirectory( gl ) diff --git a/test_conformance/extensions/CMakeLists.txt b/test_conformance/extensions/CMakeLists.txt new file mode 100644 index 00000000..c917eb30 --- /dev/null +++ b/test_conformance/extensions/CMakeLists.txt @@ -0,0 +1 @@ +add_subdirectory( cl_ext_cxx_for_opencl ) diff --git a/test_conformance/extensions/cl_ext_cxx_for_opencl/CMakeLists.txt b/test_conformance/extensions/cl_ext_cxx_for_opencl/CMakeLists.txt new file mode 100644 index 00000000..fd397c31 --- /dev/null +++ b/test_conformance/extensions/cl_ext_cxx_for_opencl/CMakeLists.txt @@ -0,0 +1,9 @@ +set(MODULE_NAME CL_EXT_CXX_FOR_OPENCL) + +set(${MODULE_NAME}_SOURCES + main.cpp + cxx_for_opencl_ext.cpp + cxx_for_opencl_ver.cpp +) + +include(../../CMakeCommon.txt) diff --git a/test_conformance/extensions/cl_ext_cxx_for_opencl/cxx_for_opencl_ext.cpp b/test_conformance/extensions/cl_ext_cxx_for_opencl/cxx_for_opencl_ext.cpp new file mode 100644 index 00000000..4b03b54a --- /dev/null +++ b/test_conformance/extensions/cl_ext_cxx_for_opencl/cxx_for_opencl_ext.cpp @@ -0,0 +1,105 @@ +// +// Copyright (c) 2021 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 "procs.h" + + +int test_cxx_for_opencl(cl_device_id device, cl_context context, + cl_command_queue queue) +{ + cl_int error; + clProgramWrapper program; + clKernelWrapper kernel1; + clKernelWrapper kernel2; + clMemWrapper in_buffer; + clMemWrapper out_buffer; + cl_int value = 7; + + const char *kernel_sstr = + R"( + __global int x; + template + void execute(T &a, const T &b) { + a = b * 2; + } + __kernel void k1(__global int *p) { + execute(x, *p); + } + __kernel void k2(__global int *p) { + execute(*p, x); + })"; + + error = create_single_kernel_helper_with_build_options( + context, &program, &kernel1, 1, &kernel_sstr, "k1", "-cl-std=CLC++", + false); + test_error(error, "Failed to create k1 kernel"); + + kernel2 = clCreateKernel(program, "k2", &error); + test_error(error, "Failed to create k2 kernel"); + + in_buffer = + clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, + sizeof(value), &value, &error); + test_error(error, "clCreateBuffer failed"); + + out_buffer = + clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, + sizeof(value), &value, &error); + test_error(error, "clCreateBuffer failed"); + + error = clSetKernelArg(kernel1, 0, sizeof(in_buffer), &in_buffer); + test_error(error, "clSetKernelArg failed"); + + error = clSetKernelArg(kernel2, 0, sizeof(out_buffer), &out_buffer); + test_error(error, "clSetKernelArg failed"); + + size_t global_size = 1; + error = clEnqueueNDRangeKernel(queue, kernel1, 1, nullptr, &global_size, + nullptr, 0, nullptr, nullptr); + test_error(error, "clEnqueueNDRangeKernel failed"); + + error = clEnqueueNDRangeKernel(queue, kernel2, 1, nullptr, &global_size, + nullptr, 0, nullptr, nullptr); + test_error(error, "clEnqueueNDRangeKernel failed"); + + error = clEnqueueReadBuffer(queue, out_buffer, CL_BLOCKING, 0, + sizeof(value), &value, 0, nullptr, nullptr); + test_error(error, "clEnqueueReadBuffer failed"); + + error = clFinish(queue); + test_error(error, "clFinish failed"); + + if (value != 28) + { + log_error("ERROR: Kernel wrote %lu, expected 28\n", + static_cast(value)); + return TEST_FAIL; + } + + return TEST_PASS; +} + +int test_cxx_for_opencl_ext(cl_device_id device, cl_context context, + cl_command_queue queue, int) +{ + if (!is_extension_available(device, "cl_ext_cxx_for_opencl")) + { + log_info("Device does not support 'cl_ext_cxx_for_opencl'. Skipping " + "the test.\n"); + return TEST_SKIPPED_ITSELF; + } + + return test_cxx_for_opencl(device, context, queue); +} diff --git a/test_conformance/extensions/cl_ext_cxx_for_opencl/cxx_for_opencl_ver.cpp b/test_conformance/extensions/cl_ext_cxx_for_opencl/cxx_for_opencl_ver.cpp new file mode 100644 index 00000000..03760814 --- /dev/null +++ b/test_conformance/extensions/cl_ext_cxx_for_opencl/cxx_for_opencl_ver.cpp @@ -0,0 +1,102 @@ +// +// Copyright (c) 2021 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 "procs.h" + + +int test_cxx_for_opencl_version(cl_device_id device, cl_context context, + cl_command_queue queue) +{ + cl_int cxx4opencl_version; + cl_int cxx4opencl_expected_version; + clProgramWrapper program; + clKernelWrapper kernel; + cl_int error; + cl_int value = 0; + const char *kernel_sstr = + R"( + __kernel void k(__global int* buf) { + buf[0] = __OPENCL_CPP_VERSION__; + })"; + const size_t lengths[1] = { std::string{ kernel_sstr }.size() }; + + clProgramWrapper writer_program = + clCreateProgramWithSource(context, 1, &kernel_sstr, lengths, &error); + test_error(error, "Failed to create program with source"); + + error = clCompileProgram(writer_program, 1, &device, "-cl-std=CLC++", 0, + nullptr, nullptr, nullptr, nullptr); + test_error(error, "Failed to compile program"); + + cl_program progs[1] = { writer_program }; + program = clLinkProgram(context, 1, &device, "", 1, progs, 0, 0, &error); + test_error(error, "Failed to link program"); + + clMemWrapper out = + clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, + sizeof(cxx4opencl_version), &cxx4opencl_version, &error); + test_error(error, "clCreateBuffer failed"); + + kernel = clCreateKernel(program, "k", &error); + test_error(error, "Failed to create k kernel"); + + error = clSetKernelArg(kernel, 0, sizeof(out), &out); + test_error(error, "clSetKernelArg failed"); + + size_t global_size = 1; + error = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &global_size, + nullptr, 0, nullptr, nullptr); + test_error(error, "clEnqueueNDRangeKernel failed"); + + error = clEnqueueReadBuffer(queue, out, CL_BLOCKING, 0, + sizeof(cxx4opencl_version), &cxx4opencl_version, + 0, nullptr, nullptr); + test_error(error, "clEnqueueReadBuffer failed"); + + error = clFinish(queue); + test_error(error, "clFinish failed"); + + error = + clGetDeviceInfo(device, CL_DEVICE_CXX_FOR_OPENCL_NUMERIC_VERSION_EXT, + sizeof(value), &value, nullptr); + test_error(error, "Failed to get device info"); + + cxx4opencl_expected_version = CL_VERSION_MAJOR_KHR(value) * 100 + + CL_VERSION_MINOR_KHR(value) * 10 + CL_VERSION_PATCH_KHR(value); + + if (cxx4opencl_version != cxx4opencl_expected_version) + { + log_error("ERROR: C++ for OpenCL version mismatch - returned %lu, " + "expected %lu\n", + static_cast(value), + static_cast(cxx4opencl_expected_version)); + return TEST_FAIL; + } + + return TEST_PASS; +} + +int test_cxx_for_opencl_ver(cl_device_id device, cl_context context, + cl_command_queue queue, int) +{ + if (!is_extension_available(device, "cl_ext_cxx_for_opencl")) + { + log_info("Device does not support 'cl_ext_cxx_for_opencl'. Skipping " + "the test.\n"); + return TEST_SKIPPED_ITSELF; + } + + return test_cxx_for_opencl_version(device, context, queue); +} diff --git a/test_conformance/extensions/cl_ext_cxx_for_opencl/main.cpp b/test_conformance/extensions/cl_ext_cxx_for_opencl/main.cpp new file mode 100644 index 00000000..5e8c14af --- /dev/null +++ b/test_conformance/extensions/cl_ext_cxx_for_opencl/main.cpp @@ -0,0 +1,28 @@ +// +// Copyright (c) 2021 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 "procs.h" + +test_definition test_list[] = { + ADD_TEST_VERSION(cxx_for_opencl_ext, Version(2, 0)), + ADD_TEST_VERSION(cxx_for_opencl_ver, Version(2, 0)) +}; + +int main(int argc, const char *argv[]) +{ + return runTestHarnessWithCheck(argc, argv, ARRAY_SIZE(test_list), test_list, + false, 0, nullptr); +} diff --git a/test_conformance/extensions/cl_ext_cxx_for_opencl/procs.h b/test_conformance/extensions/cl_ext_cxx_for_opencl/procs.h new file mode 100644 index 00000000..5665e012 --- /dev/null +++ b/test_conformance/extensions/cl_ext_cxx_for_opencl/procs.h @@ -0,0 +1,26 @@ +// +// Copyright (c) 2021 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 _procs_h +#define _procs_h + +#include "harness/typeWrappers.h" + +extern int test_cxx_for_opencl_ext(cl_device_id device, cl_context context, + cl_command_queue queue, int); +extern int test_cxx_for_opencl_ver(cl_device_id device, cl_context context, + cl_command_queue queue, int); + +#endif /*_procs_h*/