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 <victoria.holodovsky@arm.com>
Co-authored-by: Victoria Holodovsky <victoria.holodovsky@arm.com>
diff --git a/test_conformance/CMakeLists.txt b/test_conformance/CMakeLists.txt
index 6714e23..b9b87c1 100644
--- a/test_conformance/CMakeLists.txt
+++ b/test_conformance/CMakeLists.txt
@@ -21,6 +21,7 @@
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 0000000..c917eb3
--- /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 0000000..fd397c3
--- /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 0000000..4b03b54
--- /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<typename T>
+ 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<long unsigned>(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 0000000..0376081
--- /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<long unsigned>(value),
+ static_cast<long unsigned>(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 0000000..5e8c14a
--- /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 0000000..5665e01
--- /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*/
