mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-25 08:19:02 +00:00
clUnloadPlatformCompiler tests (#780)
Various tests for the clUnloadPlatformCompiler function. Fixes #442 Signed-off-by: Einar Hov <einar.hov@arm.com>
This commit is contained in:
@@ -10,6 +10,7 @@ set(${MODULE_NAME}_SOURCES
|
|||||||
test_image_macro.cpp
|
test_image_macro.cpp
|
||||||
test_compiler_defines_for_extensions.cpp
|
test_compiler_defines_for_extensions.cpp
|
||||||
test_pragma_unroll.cpp
|
test_pragma_unroll.cpp
|
||||||
|
test_unload_platform_compiler.cpp
|
||||||
)
|
)
|
||||||
|
|
||||||
include(../CMakeCommon.txt)
|
include(../CMakeCommon.txt)
|
||||||
|
|||||||
@@ -1,6 +1,6 @@
|
|||||||
//
|
//
|
||||||
// Copyright (c) 2017 The Khronos Group Inc.
|
// Copyright (c) 2017 The Khronos Group Inc.
|
||||||
//
|
//
|
||||||
// Licensed under the Apache License, Version 2.0 (the "License");
|
// Licensed under the Apache License, Version 2.0 (the "License");
|
||||||
// you may not use this file except in compliance with the License.
|
// you may not use this file except in compliance with the License.
|
||||||
// You may obtain a copy of the License at
|
// You may obtain a copy of the License at
|
||||||
@@ -15,85 +15,94 @@
|
|||||||
//
|
//
|
||||||
#include "harness/compat.h"
|
#include "harness/compat.h"
|
||||||
|
|
||||||
|
#include "harness/testHarness.h"
|
||||||
|
#include "procs.h"
|
||||||
#include <stdio.h>
|
#include <stdio.h>
|
||||||
#include <string.h>
|
#include <string.h>
|
||||||
#include "procs.h"
|
|
||||||
#include "harness/testHarness.h"
|
|
||||||
|
|
||||||
#if !defined(_WIN32)
|
#if !defined(_WIN32)
|
||||||
#include <unistd.h>
|
#include <unistd.h>
|
||||||
#endif
|
#endif
|
||||||
|
|
||||||
test_definition test_list[] = {
|
test_definition test_list[] = {
|
||||||
ADD_TEST( load_program_source ),
|
ADD_TEST(load_program_source),
|
||||||
ADD_TEST( load_multistring_source ),
|
ADD_TEST(load_multistring_source),
|
||||||
ADD_TEST( load_two_kernel_source ),
|
ADD_TEST(load_two_kernel_source),
|
||||||
ADD_TEST( load_null_terminated_source ),
|
ADD_TEST(load_null_terminated_source),
|
||||||
ADD_TEST( load_null_terminated_multi_line_source ),
|
ADD_TEST(load_null_terminated_multi_line_source),
|
||||||
ADD_TEST( load_null_terminated_partial_multi_line_source ),
|
ADD_TEST(load_null_terminated_partial_multi_line_source),
|
||||||
ADD_TEST( load_discreet_length_source ),
|
ADD_TEST(load_discreet_length_source),
|
||||||
ADD_TEST( get_program_source ),
|
ADD_TEST(get_program_source),
|
||||||
ADD_TEST( get_program_build_info ),
|
ADD_TEST(get_program_build_info),
|
||||||
ADD_TEST( get_program_info ),
|
ADD_TEST(get_program_info),
|
||||||
|
|
||||||
ADD_TEST( large_compile ),
|
ADD_TEST(large_compile),
|
||||||
ADD_TEST( async_build ),
|
ADD_TEST(async_build),
|
||||||
|
|
||||||
ADD_TEST( options_build_optimizations ),
|
ADD_TEST(options_build_optimizations),
|
||||||
ADD_TEST( options_build_macro ),
|
ADD_TEST(options_build_macro),
|
||||||
ADD_TEST( options_build_macro_existence ),
|
ADD_TEST(options_build_macro_existence),
|
||||||
ADD_TEST( options_include_directory ),
|
ADD_TEST(options_include_directory),
|
||||||
ADD_TEST( options_denorm_cache ),
|
ADD_TEST(options_denorm_cache),
|
||||||
|
|
||||||
ADD_TEST( preprocessor_define_udef ),
|
ADD_TEST(preprocessor_define_udef),
|
||||||
ADD_TEST( preprocessor_include ),
|
ADD_TEST(preprocessor_include),
|
||||||
ADD_TEST( preprocessor_line_error ),
|
ADD_TEST(preprocessor_line_error),
|
||||||
ADD_TEST( preprocessor_pragma ),
|
ADD_TEST(preprocessor_pragma),
|
||||||
|
|
||||||
ADD_TEST( compiler_defines_for_extensions ),
|
ADD_TEST(compiler_defines_for_extensions),
|
||||||
ADD_TEST( image_macro ),
|
ADD_TEST(image_macro),
|
||||||
|
|
||||||
ADD_TEST( simple_compile_only ),
|
ADD_TEST(simple_compile_only),
|
||||||
ADD_TEST( simple_static_compile_only ),
|
ADD_TEST(simple_static_compile_only),
|
||||||
ADD_TEST( simple_extern_compile_only ),
|
ADD_TEST(simple_extern_compile_only),
|
||||||
ADD_TEST( simple_compile_with_callback ),
|
ADD_TEST(simple_compile_with_callback),
|
||||||
ADD_TEST( simple_embedded_header_compile ),
|
ADD_TEST(simple_embedded_header_compile),
|
||||||
ADD_TEST( simple_link_only ),
|
ADD_TEST(simple_link_only),
|
||||||
ADD_TEST( two_file_regular_variable_access ),
|
ADD_TEST(two_file_regular_variable_access),
|
||||||
ADD_TEST( two_file_regular_struct_access ),
|
ADD_TEST(two_file_regular_struct_access),
|
||||||
ADD_TEST( two_file_regular_function_access ),
|
ADD_TEST(two_file_regular_function_access),
|
||||||
ADD_TEST( simple_link_with_callback ),
|
ADD_TEST(simple_link_with_callback),
|
||||||
ADD_TEST( simple_embedded_header_link ),
|
ADD_TEST(simple_embedded_header_link),
|
||||||
ADD_TEST( execute_after_simple_compile_and_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_no_device_info),
|
||||||
ADD_TEST( execute_after_simple_compile_and_link_with_defines ),
|
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_compile_and_link_with_callbacks),
|
||||||
ADD_TEST( execute_after_simple_library_with_link ),
|
ADD_TEST(execute_after_simple_library_with_link),
|
||||||
ADD_TEST( execute_after_two_file_link ),
|
ADD_TEST(execute_after_two_file_link),
|
||||||
ADD_TEST( execute_after_embedded_header_link ),
|
ADD_TEST(execute_after_embedded_header_link),
|
||||||
ADD_TEST( execute_after_included_header_link ),
|
ADD_TEST(execute_after_included_header_link),
|
||||||
ADD_TEST( execute_after_serialize_reload_object ),
|
ADD_TEST(execute_after_serialize_reload_object),
|
||||||
ADD_TEST( execute_after_serialize_reload_library ),
|
ADD_TEST(execute_after_serialize_reload_library),
|
||||||
ADD_TEST( simple_library_only ),
|
ADD_TEST(simple_library_only),
|
||||||
ADD_TEST( simple_library_with_callback ),
|
ADD_TEST(simple_library_with_callback),
|
||||||
ADD_TEST( simple_library_with_link ),
|
ADD_TEST(simple_library_with_link),
|
||||||
ADD_TEST( two_file_link ),
|
ADD_TEST(two_file_link),
|
||||||
ADD_TEST( multi_file_libraries ),
|
ADD_TEST(multi_file_libraries),
|
||||||
ADD_TEST( multiple_files ),
|
ADD_TEST(multiple_files),
|
||||||
ADD_TEST( multiple_libraries ),
|
ADD_TEST(multiple_libraries),
|
||||||
ADD_TEST( multiple_files_multiple_libraries ),
|
ADD_TEST(multiple_files_multiple_libraries),
|
||||||
ADD_TEST( multiple_embedded_headers ),
|
ADD_TEST(multiple_embedded_headers),
|
||||||
|
|
||||||
ADD_TEST( program_binary_type ),
|
ADD_TEST(program_binary_type),
|
||||||
ADD_TEST( compile_and_link_status_options_log ),
|
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[])
|
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);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|||||||
@@ -1,6 +1,6 @@
|
|||||||
//
|
//
|
||||||
// Copyright (c) 2017 The Khronos Group Inc.
|
// Copyright (c) 2017 The Khronos Group Inc.
|
||||||
//
|
//
|
||||||
// Licensed under the Apache License, Version 2.0 (the "License");
|
// Licensed under the Apache License, Version 2.0 (the "License");
|
||||||
// you may not use this file except in compliance with the License.
|
// you may not use this file except in compliance with the License.
|
||||||
// You may obtain a copy of the License at
|
// You may obtain a copy of the License at
|
||||||
@@ -13,76 +13,227 @@
|
|||||||
// See the License for the specific language governing permissions and
|
// See the License for the specific language governing permissions and
|
||||||
// limitations under the License.
|
// limitations under the License.
|
||||||
//
|
//
|
||||||
|
#include "harness/conversions.h"
|
||||||
#include "harness/errorHelpers.h"
|
#include "harness/errorHelpers.h"
|
||||||
#include "harness/kernelHelpers.h"
|
#include "harness/kernelHelpers.h"
|
||||||
#include "harness/typeWrappers.h"
|
|
||||||
#include "harness/conversions.h"
|
|
||||||
#include "harness/mt19937.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_program_source(cl_device_id deviceID, cl_context context,
|
||||||
extern int test_load_multistring_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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_multistring_source(cl_device_id deviceID,
|
||||||
extern int test_load_null_terminated_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
cl_context context,
|
||||||
extern int test_load_null_terminated_multi_line_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
cl_command_queue queue,
|
||||||
extern int test_load_null_terminated_partial_multi_line_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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_load_two_kernel_source(cl_device_id deviceID,
|
||||||
extern int test_get_program_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
cl_context context,
|
||||||
extern int test_get_program_build_info(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
cl_command_queue queue,
|
||||||
extern int test_get_program_info(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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_large_compile(cl_device_id deviceID, cl_context context,
|
||||||
extern int test_async_build(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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_optimizations(cl_device_id deviceID,
|
||||||
extern int test_options_build_macro(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
cl_context context,
|
||||||
extern int test_options_build_macro_existence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
cl_command_queue queue,
|
||||||
extern int test_options_include_directory(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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_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_define_udef(cl_device_id deviceID,
|
||||||
extern int test_preprocessor_include(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
cl_context context,
|
||||||
extern int test_preprocessor_line_error(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
cl_command_queue queue,
|
||||||
extern int test_preprocessor_pragma(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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_compiler_defines_for_extensions(cl_device_id device,
|
||||||
extern int test_image_macro(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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_compile_only(cl_device_id deviceID, cl_context context,
|
||||||
extern int test_simple_static_compile_only(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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_static_compile_only(cl_device_id deviceID,
|
||||||
extern int test_simple_compile_with_callback(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
cl_context context,
|
||||||
extern int test_simple_embedded_header_compile(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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_simple_link_only(cl_device_id deviceID, cl_context context,
|
||||||
extern int test_two_file_regular_variable_access(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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_variable_access(cl_device_id deviceID,
|
||||||
extern int test_two_file_regular_function_access(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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_link_with_callback(cl_device_id deviceID,
|
||||||
extern int test_simple_embedded_header_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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(cl_device_id deviceID,
|
||||||
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);
|
cl_context context,
|
||||||
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);
|
cl_command_queue queue,
|
||||||
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);
|
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_simple_compile_and_link_no_device_info(
|
||||||
extern int test_execute_after_two_file_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
cl_device_id deviceID, cl_context context, cl_command_queue queue,
|
||||||
extern int test_execute_after_embedded_header_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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_simple_compile_and_link_with_defines(
|
||||||
extern int test_execute_after_serialize_reload_object(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
cl_device_id deviceID, cl_context context, cl_command_queue queue,
|
||||||
extern int test_execute_after_serialize_reload_library(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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_only(cl_device_id deviceID, cl_context context,
|
||||||
extern int test_simple_library_with_callback(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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_simple_library_with_callback(cl_device_id deviceID,
|
||||||
extern int test_two_file_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
cl_context context,
|
||||||
extern int test_multi_file_libraries(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
cl_command_queue queue,
|
||||||
extern int test_multiple_libraries(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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_simple_library_with_link(cl_device_id deviceID,
|
||||||
extern int test_multiple_files_multiple_libraries(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
cl_context context,
|
||||||
extern int test_multiple_embedded_headers(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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_program_binary_type(cl_device_id deviceID, cl_context context,
|
||||||
extern int test_compile_and_link_status_options_log(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
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);
|
||||||
|
|||||||
962
test_conformance/compiler/test_unload_platform_compiler.cpp
Normal file
962
test_conformance/compiler/test_unload_platform_compiler.cpp
Normal file
@@ -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 <cassert>
|
||||||
|
#include <chrono>
|
||||||
|
#include <functional>
|
||||||
|
#include <future>
|
||||||
|
#include <initializer_list>
|
||||||
|
#include <stdexcept>
|
||||||
|
#include <string>
|
||||||
|
#include <thread>
|
||||||
|
#include <vector>
|
||||||
|
|
||||||
|
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<std::reference_wrapper<build_base>>;
|
||||||
|
|
||||||
|
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<unsigned char> &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<unsigned char> 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<long int>(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<long int>(err));
|
||||||
|
return 1;
|
||||||
|
}
|
||||||
|
|
||||||
|
/* Unload the platform compiler */
|
||||||
|
err = clUnloadPlatformCompiler(platform);
|
||||||
|
if (CL_SUCCESS != err)
|
||||||
|
{
|
||||||
|
log_error("Test failure: clUnloadPlatformCompiler() == %ld\n",
|
||||||
|
static_cast<long int>(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<long int>(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<long int>(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<long int>(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<long int>(err));
|
||||||
|
return 1;
|
||||||
|
}
|
||||||
|
|
||||||
|
const clCommandQueueWrapper queue =
|
||||||
|
clCreateCommandQueue(context, device, 0, &err);
|
||||||
|
if (CL_SUCCESS != err)
|
||||||
|
{
|
||||||
|
log_error("Test failure: clCreateCommandQueue() == %ld\n",
|
||||||
|
static_cast<long int>(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<long int>(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<long int>(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<long int>(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<long int>(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<long unsigned>(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<bool> unload_result =
|
||||||
|
std::async(std::launch::async, unload_thread);
|
||||||
|
std::future<bool> build_source_result =
|
||||||
|
std::async(std::launch::async, build_thread, &build_source);
|
||||||
|
std::future<bool> build_binary_result =
|
||||||
|
std::async(std::launch::async, build_thread, &build_binary);
|
||||||
|
std::future<bool> 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<long int>(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<long int>(err));
|
||||||
|
return 1;
|
||||||
|
}
|
||||||
|
|
||||||
|
/* Unload the compiler */
|
||||||
|
err = clUnloadPlatformCompiler(platform);
|
||||||
|
if (CL_SUCCESS != err)
|
||||||
|
{
|
||||||
|
log_error("Test failure: clUnloadPlatformCompiler() == %ld\n",
|
||||||
|
static_cast<long int>(err));
|
||||||
|
return 1;
|
||||||
|
}
|
||||||
|
|
||||||
|
std::vector<cl_program_build_info> 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<long int>(err));
|
||||||
|
return 1;
|
||||||
|
}
|
||||||
|
|
||||||
|
std::vector<char> 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<long int>(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<long int>(CL_BUILD_SUCCESS),
|
||||||
|
static_cast<long int>(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<long int>(
|
||||||
|
CL_PROGRAM_BINARY_TYPE_EXECUTABLE),
|
||||||
|
static_cast<long int>(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<long int>(err));
|
||||||
|
return 1;
|
||||||
|
}
|
||||||
|
|
||||||
|
err = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
|
||||||
|
if (CL_SUCCESS != err)
|
||||||
|
{
|
||||||
|
log_error("Test failure: clCompileProgram() == %ld\n",
|
||||||
|
static_cast<long int>(err));
|
||||||
|
return 1;
|
||||||
|
}
|
||||||
|
|
||||||
|
/* Unload the compiler */
|
||||||
|
err = clUnloadPlatformCompiler(platform);
|
||||||
|
if (CL_SUCCESS != err)
|
||||||
|
{
|
||||||
|
log_error("Test failure: clUnloadPlatformCompiler() == %ld\n",
|
||||||
|
static_cast<long int>(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<long int>(err));
|
||||||
|
return 1;
|
||||||
|
}
|
||||||
|
|
||||||
|
std::vector<unsigned char> 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<long int>(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;
|
||||||
|
}
|
||||||
@@ -0,0 +1,50 @@
|
|||||||
|
#include <array>
|
||||||
|
|
||||||
|
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<unsigned char, 216> 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<unsigned char, 216> 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 }
|
||||||
|
};
|
||||||
24
test_conformance/compiler/write_kernel.spvasm32
Normal file
24
test_conformance/compiler/write_kernel.spvasm32
Normal file
@@ -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
|
||||||
24
test_conformance/compiler/write_kernel.spvasm64
Normal file
24
test_conformance/compiler/write_kernel.spvasm64
Normal file
@@ -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
|
||||||
Reference in New Issue
Block a user