diff --git a/test_conformance/compiler/CMakeLists.txt b/test_conformance/compiler/CMakeLists.txt index 59fe654b..1090db38 100644 --- a/test_conformance/compiler/CMakeLists.txt +++ b/test_conformance/compiler/CMakeLists.txt @@ -10,6 +10,7 @@ set(${MODULE_NAME}_SOURCES test_image_macro.cpp test_compiler_defines_for_extensions.cpp test_pragma_unroll.cpp + test_unload_platform_compiler.cpp ) include(../CMakeCommon.txt) diff --git a/test_conformance/compiler/main.cpp b/test_conformance/compiler/main.cpp index dd97b084..45b7496d 100644 --- a/test_conformance/compiler/main.cpp +++ b/test_conformance/compiler/main.cpp @@ -1,6 +1,6 @@ // // 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 @@ -15,85 +15,94 @@ // #include "harness/compat.h" +#include "harness/testHarness.h" +#include "procs.h" #include #include -#include "procs.h" -#include "harness/testHarness.h" #if !defined(_WIN32) #include #endif test_definition test_list[] = { - ADD_TEST( load_program_source ), - ADD_TEST( load_multistring_source ), - ADD_TEST( load_two_kernel_source ), - ADD_TEST( load_null_terminated_source ), - ADD_TEST( load_null_terminated_multi_line_source ), - ADD_TEST( load_null_terminated_partial_multi_line_source ), - ADD_TEST( load_discreet_length_source ), - ADD_TEST( get_program_source ), - ADD_TEST( get_program_build_info ), - ADD_TEST( get_program_info ), + ADD_TEST(load_program_source), + ADD_TEST(load_multistring_source), + ADD_TEST(load_two_kernel_source), + ADD_TEST(load_null_terminated_source), + ADD_TEST(load_null_terminated_multi_line_source), + ADD_TEST(load_null_terminated_partial_multi_line_source), + ADD_TEST(load_discreet_length_source), + ADD_TEST(get_program_source), + ADD_TEST(get_program_build_info), + ADD_TEST(get_program_info), - ADD_TEST( large_compile ), - ADD_TEST( async_build ), + ADD_TEST(large_compile), + ADD_TEST(async_build), - ADD_TEST( options_build_optimizations ), - ADD_TEST( options_build_macro ), - ADD_TEST( options_build_macro_existence ), - ADD_TEST( options_include_directory ), - ADD_TEST( options_denorm_cache ), + ADD_TEST(options_build_optimizations), + ADD_TEST(options_build_macro), + ADD_TEST(options_build_macro_existence), + ADD_TEST(options_include_directory), + ADD_TEST(options_denorm_cache), - ADD_TEST( preprocessor_define_udef ), - ADD_TEST( preprocessor_include ), - ADD_TEST( preprocessor_line_error ), - ADD_TEST( preprocessor_pragma ), + ADD_TEST(preprocessor_define_udef), + ADD_TEST(preprocessor_include), + ADD_TEST(preprocessor_line_error), + ADD_TEST(preprocessor_pragma), - ADD_TEST( compiler_defines_for_extensions ), - ADD_TEST( image_macro ), + ADD_TEST(compiler_defines_for_extensions), + ADD_TEST(image_macro), - ADD_TEST( simple_compile_only ), - ADD_TEST( simple_static_compile_only ), - ADD_TEST( simple_extern_compile_only ), - ADD_TEST( simple_compile_with_callback ), - ADD_TEST( simple_embedded_header_compile ), - ADD_TEST( simple_link_only ), - ADD_TEST( two_file_regular_variable_access ), - ADD_TEST( two_file_regular_struct_access ), - ADD_TEST( two_file_regular_function_access ), - ADD_TEST( simple_link_with_callback ), - ADD_TEST( simple_embedded_header_link ), - ADD_TEST( execute_after_simple_compile_and_link ), - ADD_TEST( execute_after_simple_compile_and_link_no_device_info ), - ADD_TEST( execute_after_simple_compile_and_link_with_defines ), - ADD_TEST( execute_after_simple_compile_and_link_with_callbacks ), - ADD_TEST( execute_after_simple_library_with_link ), - ADD_TEST( execute_after_two_file_link ), - ADD_TEST( execute_after_embedded_header_link ), - ADD_TEST( execute_after_included_header_link ), - ADD_TEST( execute_after_serialize_reload_object ), - ADD_TEST( execute_after_serialize_reload_library ), - ADD_TEST( simple_library_only ), - ADD_TEST( simple_library_with_callback ), - ADD_TEST( simple_library_with_link ), - ADD_TEST( two_file_link ), - ADD_TEST( multi_file_libraries ), - ADD_TEST( multiple_files ), - ADD_TEST( multiple_libraries ), - ADD_TEST( multiple_files_multiple_libraries ), - ADD_TEST( multiple_embedded_headers ), + ADD_TEST(simple_compile_only), + ADD_TEST(simple_static_compile_only), + ADD_TEST(simple_extern_compile_only), + ADD_TEST(simple_compile_with_callback), + ADD_TEST(simple_embedded_header_compile), + ADD_TEST(simple_link_only), + ADD_TEST(two_file_regular_variable_access), + ADD_TEST(two_file_regular_struct_access), + ADD_TEST(two_file_regular_function_access), + ADD_TEST(simple_link_with_callback), + ADD_TEST(simple_embedded_header_link), + ADD_TEST(execute_after_simple_compile_and_link), + ADD_TEST(execute_after_simple_compile_and_link_no_device_info), + ADD_TEST(execute_after_simple_compile_and_link_with_defines), + ADD_TEST(execute_after_simple_compile_and_link_with_callbacks), + ADD_TEST(execute_after_simple_library_with_link), + ADD_TEST(execute_after_two_file_link), + ADD_TEST(execute_after_embedded_header_link), + ADD_TEST(execute_after_included_header_link), + ADD_TEST(execute_after_serialize_reload_object), + ADD_TEST(execute_after_serialize_reload_library), + ADD_TEST(simple_library_only), + ADD_TEST(simple_library_with_callback), + ADD_TEST(simple_library_with_link), + ADD_TEST(two_file_link), + ADD_TEST(multi_file_libraries), + ADD_TEST(multiple_files), + ADD_TEST(multiple_libraries), + ADD_TEST(multiple_files_multiple_libraries), + ADD_TEST(multiple_embedded_headers), - ADD_TEST( program_binary_type ), - ADD_TEST( compile_and_link_status_options_log ), + ADD_TEST(program_binary_type), + ADD_TEST(compile_and_link_status_options_log), - ADD_TEST_VERSION( pragma_unroll, Version(2, 0) ), + ADD_TEST_VERSION(pragma_unroll, Version(2, 0)), + + ADD_TEST(unload_valid), + ADD_TEST(unload_invalid), + ADD_TEST(unload_repeated), + ADD_TEST(unload_compile_unload_link), + ADD_TEST(unload_build_unload_create_kernel), + ADD_TEST(unload_link_different), + ADD_TEST(unload_build_threaded), + ADD_TEST(unload_build_info), + ADD_TEST(unload_program_binaries), }; -const int test_num = ARRAY_SIZE( test_list ); +const int test_num = ARRAY_SIZE(test_list); int main(int argc, const char *argv[]) { - return runTestHarness( argc, argv, test_num, test_list, false, false, 0 ); + return runTestHarness(argc, argv, test_num, test_list, false, false, 0); } - diff --git a/test_conformance/compiler/procs.h b/test_conformance/compiler/procs.h index 1ba655a8..05d8bd5b 100644 --- a/test_conformance/compiler/procs.h +++ b/test_conformance/compiler/procs.h @@ -1,6 +1,6 @@ // // 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 @@ -13,76 +13,227 @@ // See the License for the specific language governing permissions and // limitations under the License. // +#include "harness/conversions.h" #include "harness/errorHelpers.h" #include "harness/kernelHelpers.h" -#include "harness/typeWrappers.h" -#include "harness/conversions.h" #include "harness/mt19937.h" +#include "harness/typeWrappers.h" -extern int test_load_program_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_load_multistring_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_load_two_kernel_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_load_null_terminated_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_load_null_terminated_multi_line_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_load_null_terminated_partial_multi_line_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_load_discreet_length_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_get_program_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_get_program_build_info(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_get_program_info(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_load_program_source(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_load_multistring_source(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_load_two_kernel_source(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_load_null_terminated_source(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_load_null_terminated_multi_line_source(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_load_null_terminated_partial_multi_line_source( + cl_device_id deviceID, cl_context context, cl_command_queue queue, + int num_elements); +extern int test_load_discreet_length_source(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_get_program_source(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_get_program_build_info(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_get_program_info(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); -extern int test_large_compile(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_async_build(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_large_compile(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_async_build(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); -extern int test_options_build_optimizations(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_options_build_macro(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_options_build_macro_existence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_options_include_directory(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_options_denorm_cache(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_options_build_optimizations(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_options_build_macro(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_options_build_macro_existence(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_options_include_directory(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_options_denorm_cache(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); -extern int test_preprocessor_define_udef(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_preprocessor_include(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_preprocessor_line_error(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_preprocessor_pragma(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_preprocessor_define_udef(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_preprocessor_include(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_preprocessor_line_error(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_preprocessor_pragma(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); -extern int test_compiler_defines_for_extensions(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems ); -extern int test_image_macro(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_compiler_defines_for_extensions(cl_device_id device, + cl_context context, + cl_command_queue queue, + int n_elems); +extern int test_image_macro(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); -extern int test_simple_compile_only(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_simple_static_compile_only(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_simple_extern_compile_only(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_simple_compile_with_callback(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_simple_embedded_header_compile(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_simple_compile_only(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_simple_static_compile_only(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_simple_extern_compile_only(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_simple_compile_with_callback(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_simple_embedded_header_compile(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); -extern int test_simple_link_only(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_two_file_regular_variable_access(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_two_file_regular_struct_access(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_two_file_regular_function_access(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_simple_link_only(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_two_file_regular_variable_access(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_two_file_regular_struct_access(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_two_file_regular_function_access(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); -extern int test_simple_link_with_callback(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_simple_embedded_header_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_simple_link_with_callback(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_simple_embedded_header_link(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); -extern int test_execute_after_simple_compile_and_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_execute_after_simple_compile_and_link_no_device_info(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_execute_after_simple_compile_and_link_with_defines(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_execute_after_simple_compile_and_link_with_callbacks(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_execute_after_simple_library_with_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_execute_after_two_file_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_execute_after_embedded_header_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_execute_after_included_header_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_execute_after_serialize_reload_object(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_execute_after_serialize_reload_library(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_execute_after_simple_compile_and_link(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_execute_after_simple_compile_and_link_no_device_info( + cl_device_id deviceID, cl_context context, cl_command_queue queue, + int num_elements); +extern int test_execute_after_simple_compile_and_link_with_defines( + cl_device_id deviceID, cl_context context, cl_command_queue queue, + int num_elements); +extern int test_execute_after_simple_compile_and_link_with_callbacks( + cl_device_id deviceID, cl_context context, cl_command_queue queue, + int num_elements); +extern int test_execute_after_simple_library_with_link(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_execute_after_two_file_link(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_execute_after_embedded_header_link(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_execute_after_included_header_link(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_execute_after_serialize_reload_object(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_execute_after_serialize_reload_library(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); -extern int test_simple_library_only(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_simple_library_with_callback(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_simple_library_with_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_two_file_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_multi_file_libraries(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_multiple_libraries(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_multiple_files(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_multiple_files_multiple_libraries(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_multiple_embedded_headers(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_simple_library_only(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_simple_library_with_callback(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_simple_library_with_link(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_two_file_link(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_multi_file_libraries(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_multiple_libraries(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_multiple_files(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_multiple_files_multiple_libraries(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_multiple_embedded_headers(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); -extern int test_program_binary_type(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_compile_and_link_status_options_log(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_program_binary_type(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_compile_and_link_status_options_log(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); -extern int test_pragma_unroll(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_pragma_unroll(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); + +extern int test_unload_valid(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_unload_invalid(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_unload_repeated(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_unload_compile_unload_link(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_unload_build_unload_create_kernel(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_unload_link_different(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_unload_build_threaded(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_unload_build_info(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_unload_program_binaries(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); diff --git a/test_conformance/compiler/test_unload_platform_compiler.cpp b/test_conformance/compiler/test_unload_platform_compiler.cpp new file mode 100644 index 00000000..f88689d9 --- /dev/null +++ b/test_conformance/compiler/test_unload_platform_compiler.cpp @@ -0,0 +1,962 @@ +// +// Copyright (c) 2020 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 "testBase.h" +#include "test_unload_platform_compiler_resources.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace { + +class unload_test_failure : public std::runtime_error { +public: + using std::runtime_error::runtime_error; + + explicit unload_test_failure(const std::string &function, cl_int error) + : std::runtime_error(function + " == " + std::to_string(error)) + {} +}; + +class build_base { +public: + build_base(cl_context context, cl_device_id device) + : m_context{ context }, m_device{ device } + {} + virtual ~build_base() { reset(); } + build_base(const build_base &) = delete; + build_base &operator=(const build_base &) = delete; + + virtual void create() = 0; + + virtual void compile() + { + assert(nullptr != m_program); + + const cl_int err = clCompileProgram(m_program, 1, &m_device, nullptr, 0, + nullptr, nullptr, nullptr, nullptr); + if (CL_SUCCESS != err) + throw unload_test_failure("clCompileProgram()", err); + } + + virtual void link() + { + assert(nullptr != m_program); + + cl_int err = CL_INVALID_PLATFORM; + m_executable = clLinkProgram(m_context, 1, &m_device, nullptr, 1, + &m_program, nullptr, nullptr, &err); + if (CL_SUCCESS != err) + throw unload_test_failure("clLinkProgram()", err); + if (nullptr == m_executable) + throw unload_test_failure("clLinkProgram returned nullptr"); + } + + virtual void verify() + { + assert(nullptr != m_executable); + + cl_int err = CL_INVALID_VALUE; + + const clKernelWrapper kernel = + clCreateKernel(m_executable, "write_kernel", &err); + if (CL_SUCCESS != err) + throw unload_test_failure("clCreateKernel()", err); + + const clCommandQueueWrapper queue = + clCreateCommandQueue(m_context, m_device, 0, &err); + if (CL_SUCCESS != err) + throw unload_test_failure("clCreateCommandQueue()", err); + + const clMemWrapper buffer = clCreateBuffer( + m_context, CL_MEM_READ_WRITE, sizeof(cl_uint), nullptr, &err); + if (CL_SUCCESS != err) + throw unload_test_failure("clCreateBuffer()", err); + + cl_uint value = 0; + + err = clSetKernelArg(kernel, 0, sizeof(buffer), &buffer); + if (CL_SUCCESS != err) + throw unload_test_failure("clSetKernelArg()", err); + + static const size_t work_size = 1; + err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &work_size, + nullptr, 0, nullptr, nullptr); + if (CL_SUCCESS != err) + throw unload_test_failure("clEnqueueNDRangeKernel()", err); + + err = clEnqueueReadBuffer(queue, buffer, CL_BLOCKING, 0, + sizeof(cl_uint), &value, 0, nullptr, nullptr); + if (CL_SUCCESS != err) + throw unload_test_failure("clEnqueueReadBuffer()", err); + + err = clFinish(queue); + if (CL_SUCCESS != err) throw unload_test_failure("clFinish()", err); + + if (42 != value) + { + throw unload_test_failure("Kernel wrote " + std::to_string(value) + + ", expected 42"); + } + } + + void reset() + { + if (m_program) + { + clReleaseProgram(m_program); + m_program = nullptr; + } + if (m_executable) + { + clReleaseProgram(m_executable); + m_executable = nullptr; + } + } + + void build() + { + compile(); + link(); + } + +protected: + const cl_context m_context; + const cl_device_id m_device; + cl_program m_program{}; + cl_program m_executable{}; +}; + +/** + * @brief initializer_list type for constructing loops over build tests. + */ +using build_list = std::initializer_list>; + +class build_with_source : public build_base { +public: + using build_base::build_base; + + void create() final + { + assert(nullptr == m_program); + + static const char *sources[] = { write_kernel_source }; + + cl_int err = CL_INVALID_PLATFORM; + m_program = + clCreateProgramWithSource(m_context, 1, sources, nullptr, &err); + if (CL_SUCCESS != err) + throw unload_test_failure("clCreateProgramWithSource()", err); + if (nullptr == m_program) + throw unload_test_failure( + "clCreateProgramWithSource returned nullptr"); + } +}; + +class build_with_binary : public build_base { +public: + build_with_binary(const cl_context context, const cl_device_id device, + const std::vector &binary) + : build_base{ context, device }, m_binary{ binary } + {} + + build_with_binary(const cl_context context, const cl_device_id device) + : build_base{ context, device } + { + cl_int err = CL_INVALID_VALUE; + + /* Build the program from source */ + static const char *sources[] = { write_kernel_source }; + clProgramWrapper program = + clCreateProgramWithSource(m_context, 1, sources, nullptr, &err); + if (CL_SUCCESS != err) + throw unload_test_failure("clCreateProgramWithSource()", err); + + err = clCompileProgram(program, 1, &m_device, nullptr, 0, nullptr, + nullptr, nullptr, nullptr); + if (CL_SUCCESS != err) + throw unload_test_failure("clCompileProgram()", err); + + const clProgramWrapper executable = + clLinkProgram(m_context, 1, &m_device, nullptr, 1, &program, + nullptr, nullptr, &err); + if (CL_SUCCESS != err) + throw unload_test_failure("clLinkProgram()", err); + + size_t binary_size; + err = clGetProgramInfo(executable, CL_PROGRAM_BINARY_SIZES, + sizeof(binary_size), &binary_size, nullptr); + if (CL_SUCCESS != err) + throw unload_test_failure("clGetProgramInfo()", err); + + m_binary.resize(binary_size); + + /* Grab the program binary */ + unsigned char *binaries[] = { m_binary.data() }; + err = clGetProgramInfo(executable, CL_PROGRAM_BINARIES, + sizeof(unsigned char *), binaries, nullptr); + if (CL_SUCCESS != err) + throw unload_test_failure("clGetProgramInfo()", err); + } + + void create() final + { + assert(nullptr == m_executable); + + const unsigned char *binaries[] = { m_binary.data() }; + const size_t binary_sizes[] = { m_binary.size() }; + + cl_int err = CL_INVALID_PLATFORM; + m_executable = clCreateProgramWithBinary( + m_context, 1, &m_device, binary_sizes, binaries, nullptr, &err); + if (CL_SUCCESS != err) + throw unload_test_failure("clCreateProgramWithBinary()", err); + if (nullptr == m_executable) + throw unload_test_failure( + "clCreateProgramWithBinary returned nullptr"); + } + + void compile() final + { + assert(nullptr != m_executable); + + /* Program created from binary, there is nothing to do */ + } + + void link() final + { + assert(nullptr != m_executable); + + const cl_int err = clBuildProgram(m_executable, 1, &m_device, nullptr, + nullptr, nullptr); + if (CL_SUCCESS != err) + throw unload_test_failure("clBuildProgram()", err); + } + +private: + std::vector m_binary; +}; + +class build_with_il : public build_base { +public: + build_with_il(const cl_context context, const cl_platform_id platform, + const cl_device_id device) + : build_base{ context, device } + { + if (get_device_cl_version(device) >= Version(2, 1)) + { + m_CreateProgramWithIL = clCreateProgramWithIL; + m_enabled = true; + } + else if (is_extension_available(device, "cl_khr_il_program")) + { + m_CreateProgramWithIL = (decltype(m_CreateProgramWithIL)) + clGetExtensionFunctionAddressForPlatform( + platform, "clCreateProgramWithILKHR"); + if (nullptr == m_CreateProgramWithIL) + { + throw unload_test_failure("cl_khr_il_program supported, but " + "function address is nullptr"); + } + m_enabled = true; + } + else + { + /* Disable build_with_il if neither core nor extension functionality + * is available */ + m_enabled = false; + } + + cl_uint address_bits{}; + const cl_int err = + clGetDeviceInfo(device, CL_DEVICE_ADDRESS_BITS, sizeof(cl_uint), + &address_bits, nullptr); + if (CL_SUCCESS != err) + { + throw unload_test_failure("Failure getting device address bits"); + } + + switch (address_bits) + { + case 32: + m_spirv_binary = write_kernel_32_spv.data(); + m_spirv_size = write_kernel_32_spv.size(); + break; + case 64: + m_spirv_binary = write_kernel_64_spv.data(); + m_spirv_size = write_kernel_64_spv.size(); + break; + default: throw unload_test_failure("Invalid address bits"); + } + } + + void create() final + { + if (!m_enabled) return; + + assert(nullptr == m_program); + + cl_int err = CL_INVALID_PLATFORM; + m_program = m_CreateProgramWithIL(m_context, m_spirv_binary, + m_spirv_size, &err); + if (CL_SUCCESS != err) + throw unload_test_failure("clCreateProgramWithIL()", err); + if (nullptr == m_program) + throw unload_test_failure("clCreateProgramWithIL returned nullptr"); + } + + void compile() final + { + if (!m_enabled) return; + build_base::compile(); + } + + void link() final + { + if (!m_enabled) return; + build_base::link(); + } + + void verify() final + { + if (!m_enabled) return; + build_base::verify(); + } + +private: + void *m_spirv_binary; + size_t m_spirv_size; + bool m_enabled; + + using CreateProgramWithIL_fn = decltype(&clCreateProgramWithIL); + CreateProgramWithIL_fn m_CreateProgramWithIL; +}; +} + +static cl_platform_id device_platform(cl_device_id device) +{ + cl_platform_id platform; + const cl_int err = clGetDeviceInfo(device, CL_DEVICE_PLATFORM, + sizeof(platform), &platform, nullptr); + if (CL_SUCCESS != err) + { + log_error("Failure getting platform of tested device\n"); + return nullptr; + } + + return platform; +} + +static void unload_platform_compiler(const cl_platform_id platform) +{ + const cl_int err = clUnloadPlatformCompiler(platform); + if (CL_SUCCESS != err) + throw unload_test_failure("clUnloadPlatformCompiler()", err); +} + +/* Test calling the function with a valid platform */ +int test_unload_valid(cl_device_id device, cl_context, cl_command_queue, int) +{ + const cl_platform_id platform = device_platform(device); + const long int err = clUnloadPlatformCompiler(platform); + + if (CL_SUCCESS != err) + { + log_error("Test failure: clUnloadPlatformCompiler() == %ld\n", err); + return 1; + } + + return 0; +} + +/* Test calling the function with invalid platform */ +int test_unload_invalid(cl_device_id, cl_context, cl_command_queue, int) +{ + const long int err = clUnloadPlatformCompiler(nullptr); + + if (CL_INVALID_PLATFORM != err) + { + log_error("Test failure: clUnloadPlatformCompiler() == %ld\n", err); + return 1; + } + + return 0; +} + +/* Test calling the function multiple times in a row */ +int test_unload_repeated(cl_device_id device, cl_context context, + cl_command_queue, int) +{ + const cl_platform_id platform = device_platform(device); + try + { + build_with_source source(context, device); + build_with_binary binary(context, device); + build_with_il il(context, platform, device); + + for (build_base &test : build_list{ source, binary, il }) + { + unload_platform_compiler(platform); + unload_platform_compiler(platform); + + test.create(); + test.build(); + test.verify(); + } + } catch (const unload_test_failure &e) + { + log_error("Test failure: %s\n", e.what()); + return 1; + } + + return 0; +} + +/* Test calling the function between compilation and linking of programs */ +int test_unload_compile_unload_link(cl_device_id device, cl_context context, + cl_command_queue, int) +{ + const cl_platform_id platform = device_platform(device); + try + { + build_with_source source(context, device); + build_with_binary binary(context, device); + build_with_il il(context, platform, device); + + for (build_base &test : build_list{ source, binary, il }) + { + unload_platform_compiler(platform); + test.create(); + test.compile(); + unload_platform_compiler(platform); + test.link(); + test.verify(); + } + } catch (const unload_test_failure &e) + { + log_error("Test failure: %s\n", e.what()); + return 1; + } + + return 0; +} + +/* Test calling the function between program build and kernel creation */ +int test_unload_build_unload_create_kernel(cl_device_id device, + cl_context context, cl_command_queue, + int) +{ + const cl_platform_id platform = device_platform(device); + try + { + build_with_source source(context, device); + build_with_binary binary(context, device); + build_with_il il(context, platform, device); + + for (build_base &test : build_list{ source, binary, il }) + { + unload_platform_compiler(platform); + test.create(); + test.build(); + unload_platform_compiler(platform); + test.verify(); + } + } catch (const unload_test_failure &e) + { + log_error("Test failure: %s\n", e.what()); + return 1; + } + + return 0; +} + +/* Test linking together two programs that were built with a call to the unload + * function in between */ +int test_unload_link_different(cl_device_id device, cl_context context, + cl_command_queue, int) +{ + const cl_platform_id platform = device_platform(device); + + static const char *sources_1[] = { "unsigned int a() { return 42; }" }; + static const char *sources_2[] = { R"( + unsigned int a(); + kernel void test(global unsigned int *p) + { + *p = a(); + })" }; + + cl_int err = CL_INVALID_PLATFORM; + + /* Create and compile program 1 */ + const clProgramWrapper program_1 = + clCreateProgramWithSource(context, 1, sources_1, nullptr, &err); + if (CL_SUCCESS != err) + { + log_error("Test failure: clCreateProgramWithSource() == %ld\n", + static_cast(err)); + return 1; + } + + err = clCompileProgram(program_1, 1, &device, nullptr, 0, nullptr, nullptr, + nullptr, nullptr); + if (CL_SUCCESS != err) + { + log_error("Test failure: clCompileProgram() == %ld\n", + static_cast(err)); + return 1; + } + + /* Unload the platform compiler */ + err = clUnloadPlatformCompiler(platform); + if (CL_SUCCESS != err) + { + log_error("Test failure: clUnloadPlatformCompiler() == %ld\n", + static_cast(err)); + return 1; + } + + /* Create and compile program 2 with the new compiler context */ + const clProgramWrapper program_2 = + clCreateProgramWithSource(context, 1, sources_2, nullptr, &err); + if (CL_SUCCESS != err) + { + log_error("Test failure: clCreateProgramWithSource() == %ld\n", + static_cast(err)); + return 1; + } + + err = clCompileProgram(program_2, 1, &device, nullptr, 0, nullptr, nullptr, + nullptr, nullptr); + if (CL_SUCCESS != err) + { + log_error("Test failure: clCompileProgram() == %ld\n", + static_cast(err)); + return 1; + } + + /* Link the two programs into an executable program */ + const cl_program compiled_programs[] = { program_1, program_2 }; + + const clProgramWrapper executable = + clLinkProgram(context, 1, &device, nullptr, 2, compiled_programs, + nullptr, nullptr, &err); + if (CL_SUCCESS != err) + { + log_error("Test failure: clLinkProgram() == %ld\n", + static_cast(err)); + return 1; + } + + /* Verify execution of a kernel from the linked executable */ + const clKernelWrapper kernel = clCreateKernel(executable, "test", &err); + if (CL_SUCCESS != err) + { + log_error("Test failure: clCreateKernel() == %ld\n", + static_cast(err)); + return 1; + } + + const clCommandQueueWrapper queue = + clCreateCommandQueue(context, device, 0, &err); + if (CL_SUCCESS != err) + { + log_error("Test failure: clCreateCommandQueue() == %ld\n", + static_cast(err)); + return 1; + } + + const clMemWrapper buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, + sizeof(cl_uint), nullptr, &err); + if (CL_SUCCESS != err) + { + log_error("Test failure: clCreateBuffer() == %ld\n", + static_cast(err)); + return 1; + } + + cl_uint value = 0; + + err = clSetKernelArg(kernel, 0, sizeof(buffer), &buffer); + if (CL_SUCCESS != err) + { + log_error("Test failure: clSetKernelArg() == %ld\n", + static_cast(err)); + return 1; + } + + static const size_t work_size = 1; + err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &work_size, nullptr, + 0, nullptr, nullptr); + if (CL_SUCCESS != err) + { + log_error("Test failure: clEnqueueNDRangeKernel() == %ld\n", + static_cast(err)); + return 1; + } + + err = clEnqueueReadBuffer(queue, buffer, CL_BLOCKING, 0, sizeof(cl_uint), + &value, 0, nullptr, nullptr); + if (CL_SUCCESS != err) + { + log_error("Test failure: clEnqueueReadBuffer() == %ld\n", + static_cast(err)); + return 1; + } + + err = clFinish(queue); + if (CL_SUCCESS != err) throw unload_test_failure("clFinish()", err); + + if (42 != value) + { + log_error("Test failure: Kernel wrote %lu, expected 42)\n", + static_cast(value)); + return 1; + } + + return 0; +} + +/* Test calling the function in a thread while others threads are building + * programs */ +int test_unload_build_threaded(cl_device_id device, cl_context context, + cl_command_queue, int) +{ + using clock = std::chrono::steady_clock; + + const cl_platform_id platform = device_platform(device); + + const auto end = clock::now() + std::chrono::seconds(5); + + const auto unload_thread = [&end, platform] { + bool success = true; + + /* Repeatedly unload the compiler */ + try + { + while (clock::now() < end) + { + unload_platform_compiler(platform); + } + } catch (const unload_test_failure &e) + { + log_error("Test failure: %s\n", e.what()); + success = false; + } + + return success; + }; + + const auto build_thread = [&end](build_base *build) { + bool success = true; + + try + { + while (clock::now() < end) + { + build->create(); + build->build(); + build->verify(); + build->reset(); + } + } catch (unload_test_failure &e) + { + log_error("Test failure: %s\n", e.what()); + success = false; + } + + return success; + }; + + build_with_source build_source(context, device); + build_with_binary build_binary(context, device); + build_with_il build_il(context, platform, device); + + /* Run all threads in parallel and wait for them to finish */ + std::future unload_result = + std::async(std::launch::async, unload_thread); + std::future build_source_result = + std::async(std::launch::async, build_thread, &build_source); + std::future build_binary_result = + std::async(std::launch::async, build_thread, &build_binary); + std::future build_il_result = + std::async(std::launch::async, build_thread, &build_il); + + bool success = true; + if (!unload_result.get()) + { + log_error("unload_thread failed\n"); + success = false; + } + if (!build_source_result.get()) + { + log_error("build_with_source failed\n"); + success = false; + } + if (!build_binary_result.get()) + { + log_error("build_with_binary failed\n"); + success = false; + } + if (!build_il_result.get()) + { + log_error("build_with_il failed\n"); + success = false; + } + + return success ? 0 : 1; +} + +/* Test grabbing program build information after calling the unload function */ +int test_unload_build_info(cl_device_id device, cl_context context, + cl_command_queue, int) +{ + const cl_platform_id platform = device_platform(device); + + static const char *sources[] = { write_kernel_source }; + + cl_int err = CL_INVALID_PLATFORM; + /* Create and build the initial program from source */ + const clProgramWrapper program = + clCreateProgramWithSource(context, 1, sources, nullptr, &err); + if (CL_SUCCESS != err) + { + log_error("Test failure: clCreateProgramWithSource() == %ld\n", + static_cast(err)); + return 1; + } + + static const std::string options("-Dtest"); + + err = + clBuildProgram(program, 1, &device, options.c_str(), nullptr, nullptr); + if (CL_SUCCESS != err) + { + log_error("Test failure: clCompileProgram() == %ld\n", + static_cast(err)); + return 1; + } + + /* Unload the compiler */ + err = clUnloadPlatformCompiler(platform); + if (CL_SUCCESS != err) + { + log_error("Test failure: clUnloadPlatformCompiler() == %ld\n", + static_cast(err)); + return 1; + } + + std::vector infos{ CL_PROGRAM_BUILD_STATUS, + CL_PROGRAM_BUILD_OPTIONS, + CL_PROGRAM_BUILD_LOG, + CL_PROGRAM_BINARY_TYPE }; + + if (get_device_cl_version(device) >= Version(2, 0)) + { + infos.push_back(CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE); + } + + /* Try grabbing the infos after the compiler unload */ + for (cl_program_build_info info : infos) + { + size_t info_size = 0; + err = clGetProgramBuildInfo(program, device, info, 0, nullptr, + &info_size); + if (CL_SUCCESS != err) + { + log_error("Test failure: clGetProgramBuildInfo() == %ld\n", + static_cast(err)); + return 1; + } + + std::vector info_value(info_size); + + size_t written_size = 0; + err = clGetProgramBuildInfo(program, device, info, info_size, + &info_value[0], &written_size); + if (CL_SUCCESS != err) + { + log_error("Test failure: clGetProgramBuildInfo() == %ld\n", + static_cast(err)); + return 1; + } + else if (written_size != info_size) + { + log_error("Test failure: Written info value size (%zu) was " + "different from " + "queried size (%zu).\n", + written_size, info_size); + return 1; + } + + /* Verify the information we know the answer to */ + switch (info) + { + case CL_PROGRAM_BUILD_STATUS: { + constexpr size_t value_size = sizeof(cl_build_status); + if (value_size != info_size) + { + log_error("Test failure: Expected CL_PROGRAM_BUILD_STATUS " + "of size %zu, " + "but got %zu\n", + value_size, info_size); + return 1; + } + cl_build_status value; + memcpy(&value, &info_value[0], value_size); + if (CL_BUILD_SUCCESS != value) + { + log_error( + "Test failure: CL_PROGRAM_BUILD_STATUS did not return " + "CL_BUILD_SUCCESS (%ld), but %ld\n", + static_cast(CL_BUILD_SUCCESS), + static_cast(value)); + return 1; + } + } + break; + + case CL_PROGRAM_BUILD_OPTIONS: { + const size_t value_size = options.length() + 1; + if (value_size != info_size) + { + log_error("Test failure: Expected CL_PROGRAM_BUILD_OPTIONS " + "of size " + "%zu, but got %zu\n", + value_size, info_size); + return 1; + } + else if (options != &info_value[0]) + { + log_error("Test failure: CL_PROGRAM_BUILD_OPTIONS returned " + "\"%s\" " + "instead of \"%s\"\n", + &info_value[0], options.c_str()); + return 1; + } + } + break; + + case CL_PROGRAM_BINARY_TYPE: { + constexpr size_t value_size = sizeof(cl_program_binary_type); + if (value_size != info_size) + { + log_error("Test failure: Expected CL_PROGRAM_BINARY_TYPE " + "of size %zu, " + "but got %zu\n", + value_size, info_size); + return 1; + } + cl_program_binary_type value; + memcpy(&value, &info_value[0], value_size); + if (CL_PROGRAM_BINARY_TYPE_EXECUTABLE != value) + { + log_error( + "Test failure: CL_PROGRAM_BINARY_TYPE did not return " + "CL_PROGRAM_BINARY_TYPE_EXECUTABLE (%ld), but %ld\n", + static_cast( + CL_PROGRAM_BINARY_TYPE_EXECUTABLE), + static_cast(value)); + return 1; + } + } + break; + } + } + + return 0; +} + +/* Test calling the unload function between program building and fetching the + * program binaries */ +int test_unload_program_binaries(cl_device_id device, cl_context context, + cl_command_queue, int) +{ + const cl_platform_id platform = device_platform(device); + + static const char *sources[] = { write_kernel_source }; + + cl_int err = CL_INVALID_PLATFORM; + /* Create and build the initial program from source */ + const clProgramWrapper program = + clCreateProgramWithSource(context, 1, sources, nullptr, &err); + if (CL_SUCCESS != err) + { + log_error("Test failure: clCreateProgramWithSource() == %ld\n", + static_cast(err)); + return 1; + } + + err = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr); + if (CL_SUCCESS != err) + { + log_error("Test failure: clCompileProgram() == %ld\n", + static_cast(err)); + return 1; + } + + /* Unload the compiler */ + err = clUnloadPlatformCompiler(platform); + if (CL_SUCCESS != err) + { + log_error("Test failure: clUnloadPlatformCompiler() == %ld\n", + static_cast(err)); + return 1; + } + + /* Grab the built executable binary after the compiler unload */ + size_t binary_size; + err = clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, + sizeof(binary_size), &binary_size, nullptr); + if (CL_SUCCESS != err) + { + log_error("Test failure: clGetProgramInfo() == %ld\n", + static_cast(err)); + return 1; + } + + std::vector binary(binary_size); + + unsigned char *binaries[] = { binary.data() }; + err = clGetProgramInfo(program, CL_PROGRAM_BINARIES, + sizeof(unsigned char *), binaries, nullptr); + if (CL_SUCCESS != err) + { + log_error("Test failure: clGetProgramInfo() == %ld\n", + static_cast(err)); + return 1; + } + + /* Create a new program from the binary and test its execution */ + try + { + build_with_binary build_binary(context, device, binary); + build_binary.create(); + build_binary.build(); + build_binary.verify(); + } catch (unload_test_failure &e) + { + log_error("Test failure: %s\n", e.what()); + return 1; + } + + return 0; +} diff --git a/test_conformance/compiler/test_unload_platform_compiler_resources.hpp b/test_conformance/compiler/test_unload_platform_compiler_resources.hpp new file mode 100644 index 00000000..82f87ffc --- /dev/null +++ b/test_conformance/compiler/test_unload_platform_compiler_resources.hpp @@ -0,0 +1,50 @@ +#include + +static const char write_kernel_source[] = R"( + kernel void write_kernel(global unsigned int *p) { + *p = 42; + })"; + +/* Assembled SPIR-V 1.0 binary from write_kernel.spvasm64 */ +static std::array write_kernel_64_spv{ + { 0x03, 0x02, 0x23, 0x07, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x07, 0x00, + 0x09, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x02, 0x00, + 0x04, 0x00, 0x00, 0x00, 0x11, 0x00, 0x02, 0x00, 0x06, 0x00, 0x00, 0x00, + 0x0e, 0x00, 0x03, 0x00, 0x02, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, + 0x0f, 0x00, 0x07, 0x00, 0x06, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, + 0x77, 0x72, 0x69, 0x74, 0x65, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, + 0x00, 0x00, 0x00, 0x00, 0x15, 0x00, 0x04, 0x00, 0x02, 0x00, 0x00, 0x00, + 0x20, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, + 0x02, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x2a, 0x00, 0x00, 0x00, + 0x13, 0x00, 0x02, 0x00, 0x04, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, + 0x05, 0x00, 0x00, 0x00, 0x05, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, + 0x21, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, + 0x05, 0x00, 0x00, 0x00, 0x36, 0x00, 0x05, 0x00, 0x04, 0x00, 0x00, 0x00, + 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, + 0x37, 0x00, 0x03, 0x00, 0x05, 0x00, 0x00, 0x00, 0x07, 0x00, 0x00, 0x00, + 0xf8, 0x00, 0x02, 0x00, 0x08, 0x00, 0x00, 0x00, 0x3e, 0x00, 0x05, 0x00, + 0x07, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, + 0x04, 0x00, 0x00, 0x00, 0xfd, 0x00, 0x01, 0x00, 0x38, 0x00, 0x01, 0x00 } +}; + +/* Assembled SPIR-V 1.0 binary from write_kernel.spvasm32 */ +static std::array write_kernel_32_spv{ + { 0x03, 0x02, 0x23, 0x07, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x07, 0x00, + 0x09, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x02, 0x00, + 0x04, 0x00, 0x00, 0x00, 0x11, 0x00, 0x02, 0x00, 0x06, 0x00, 0x00, 0x00, + 0x0e, 0x00, 0x03, 0x00, 0x01, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, + 0x0f, 0x00, 0x07, 0x00, 0x06, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, + 0x77, 0x72, 0x69, 0x74, 0x65, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, + 0x00, 0x00, 0x00, 0x00, 0x15, 0x00, 0x04, 0x00, 0x02, 0x00, 0x00, 0x00, + 0x20, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, + 0x02, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x2a, 0x00, 0x00, 0x00, + 0x13, 0x00, 0x02, 0x00, 0x04, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, + 0x05, 0x00, 0x00, 0x00, 0x05, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, + 0x21, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, + 0x05, 0x00, 0x00, 0x00, 0x36, 0x00, 0x05, 0x00, 0x04, 0x00, 0x00, 0x00, + 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, + 0x37, 0x00, 0x03, 0x00, 0x05, 0x00, 0x00, 0x00, 0x07, 0x00, 0x00, 0x00, + 0xf8, 0x00, 0x02, 0x00, 0x08, 0x00, 0x00, 0x00, 0x3e, 0x00, 0x05, 0x00, + 0x07, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, + 0x04, 0x00, 0x00, 0x00, 0xfd, 0x00, 0x01, 0x00, 0x38, 0x00, 0x01, 0x00 } +}; diff --git a/test_conformance/compiler/write_kernel.spvasm32 b/test_conformance/compiler/write_kernel.spvasm32 new file mode 100644 index 00000000..b6a3fc1f --- /dev/null +++ b/test_conformance/compiler/write_kernel.spvasm32 @@ -0,0 +1,24 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 11 +; Schema: 0 + OpCapability Addresses + OpCapability Kernel +; %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %6 "write_kernel" +; %10 = OpString "kernel_arg_type.write_kernel.uint*," +; OpSource OpenCL_C 200000 +; OpDecorate %7 FuncParamAttr NoCapture + %uint = OpTypeInt 32 0 + %uint_42 = OpConstant %uint 42 + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint + %6 = OpFunction %void None %5 + %7 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %8 = OpLabel + OpStore %7 %uint_42 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/compiler/write_kernel.spvasm64 b/test_conformance/compiler/write_kernel.spvasm64 new file mode 100644 index 00000000..0923bc13 --- /dev/null +++ b/test_conformance/compiler/write_kernel.spvasm64 @@ -0,0 +1,24 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 11 +; Schema: 0 + OpCapability Addresses + OpCapability Kernel +; %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %6 "write_kernel" +; %10 = OpString "kernel_arg_type.write_kernel.uint*," +; OpSource OpenCL_C 200000 +; OpDecorate %7 FuncParamAttr NoCapture + %uint = OpTypeInt 32 0 + %uint_42 = OpConstant %uint 42 + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint + %6 = OpFunction %void None %5 + %7 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %8 = OpLabel + OpStore %7 %uint_42 Aligned 4 + OpReturn + OpFunctionEnd