From 1b80a22347dbadaa5cf436186f9f68f5b44217d8 Mon Sep 17 00:00:00 2001 From: Ross Brunton Date: Wed, 13 Mar 2024 17:26:10 +0000 Subject: [PATCH] [Testing] Improve "program" testing to better match the DPC++ e2e tests This adds a number of conformance tests to test things that are required for the "program" DPC++ e2e tests. Note that these tests don't all pass. Adapter testing infrastructure for level zero has been added, and a single test has been added to check their specific handling of linker errors (they write it to the build log of the output program). Other than that, test additions are as follows: * A few specialization constant tests to test usage of a kernel with multiple specialization constants defined. * Testing of the default specialization values. * KernelGetInfo outputs the correct kernel name and context pointer. * Compiling an (valid or invalid) program produces a valid (though unspecified) build log. * The "num devices" program info is sensible. --- include/ur_api.h | 4 + scripts/core/program.yml | 2 + source/loader/ur_libapi.cpp | 4 + source/ur_api.cpp | 4 + test/CMakeLists.txt | 5 +- test/adapters/CMakeLists.txt | 13 +++ test/adapters/level_zero/CMakeLists.txt | 27 ++++++ test/adapters/level_zero/urProgramLink.cpp | 31 ++++++ test/conformance/device_code/CMakeLists.txt | 11 +++ .../conformance/device_code/build_failure.cpp | 21 +++++ .../conformance/device_code/spec_constant.cpp | 2 +- .../device_code/spec_constant_multiple.cpp | 43 +++++++++ .../kernel/kernel_adapter_native_cpu.match | 2 + test/conformance/kernel/urKernelGetInfo.cpp | 23 +++++ .../program/program_adapter_cuda.match | 5 + .../program/program_adapter_hip.match | 2 + .../program/program_adapter_native_cpu.match | 10 ++ .../program/program_adapter_opencl.match | 1 + test/conformance/program/urProgramBuild.cpp | 31 ++++++ .../program/urProgramCreateWithBinary.cpp | 10 ++ .../program/urProgramCreateWithIL.cpp | 9 ++ .../program/urProgramGetBuildInfo.cpp | 22 +++++ test/conformance/program/urProgramGetInfo.cpp | 40 ++++++++ .../urProgramSetSpecializationConstants.cpp | 94 +++++++++++++++++++ test/conformance/source/environment.cpp | 2 + 25 files changed, 416 insertions(+), 2 deletions(-) create mode 100644 test/adapters/level_zero/CMakeLists.txt create mode 100644 test/adapters/level_zero/urProgramLink.cpp create mode 100644 test/conformance/device_code/build_failure.cpp create mode 100644 test/conformance/device_code/spec_constant_multiple.cpp diff --git a/include/ur_api.h b/include/ur_api.h index a528c045f4..8680dbeffb 100644 --- a/include/ur_api.h +++ b/include/ur_api.h @@ -4076,6 +4076,8 @@ typedef struct ur_program_properties_t { /// /// @details /// - The application may call this function from simultaneous threads. +/// - The adapter may (but is not required to) perform validation of the +/// provided module during this call. /// /// @remarks /// _Analogues_ @@ -4118,6 +4120,8 @@ urProgramCreateWithIL( /// ::UR_PROGRAM_BINARY_TYPE_LIBRARY for `hDevice`. /// - The device specified by `hDevice` must be device associated with /// context. +/// - The adapter may (but is not required to) perform validation of the +/// provided module during this call. /// /// @remarks /// _Analogues_ diff --git a/scripts/core/program.yml b/scripts/core/program.yml index 65c86e5a5b..45f7710d68 100644 --- a/scripts/core/program.yml +++ b/scripts/core/program.yml @@ -89,6 +89,7 @@ analogue: - "**clCreateProgramWithIL**" details: - "The application may call this function from simultaneous threads." + - "The adapter may (but is not required to) perform validation of the provided module during this call." params: - type: $x_context_handle_t name: hContext @@ -129,6 +130,7 @@ details: - "The application may call this function from simultaneous threads." - "Following a successful call to this entry point, `phProgram` will contain a binary of type $X_PROGRAM_BINARY_TYPE_COMPILED_OBJECT or $X_PROGRAM_BINARY_TYPE_LIBRARY for `hDevice`." - "The device specified by `hDevice` must be device associated with context." + - "The adapter may (but is not required to) perform validation of the provided module during this call." params: - type: $x_context_handle_t name: hContext diff --git a/source/loader/ur_libapi.cpp b/source/loader/ur_libapi.cpp index 2271b8aa09..66f4835c56 100644 --- a/source/loader/ur_libapi.cpp +++ b/source/loader/ur_libapi.cpp @@ -2890,6 +2890,8 @@ ur_result_t UR_APICALL urPhysicalMemRelease( /// /// @details /// - The application may call this function from simultaneous threads. +/// - The adapter may (but is not required to) perform validation of the +/// provided module during this call. /// /// @remarks /// _Analogues_ @@ -2942,6 +2944,8 @@ ur_result_t UR_APICALL urProgramCreateWithIL( /// ::UR_PROGRAM_BINARY_TYPE_LIBRARY for `hDevice`. /// - The device specified by `hDevice` must be device associated with /// context. +/// - The adapter may (but is not required to) perform validation of the +/// provided module during this call. /// /// @remarks /// _Analogues_ diff --git a/source/ur_api.cpp b/source/ur_api.cpp index abaad78548..e6410ee99b 100644 --- a/source/ur_api.cpp +++ b/source/ur_api.cpp @@ -2459,6 +2459,8 @@ ur_result_t UR_APICALL urPhysicalMemRelease( /// /// @details /// - The application may call this function from simultaneous threads. +/// - The adapter may (but is not required to) perform validation of the +/// provided module during this call. /// /// @remarks /// _Analogues_ @@ -2505,6 +2507,8 @@ ur_result_t UR_APICALL urProgramCreateWithIL( /// ::UR_PROGRAM_BINARY_TYPE_LIBRARY for `hDevice`. /// - The device specified by `hDevice` must be device associated with /// context. +/// - The adapter may (but is not required to) perform validation of the +/// provided module during this call. /// /// @remarks /// _Analogues_ diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index d158fcdf24..3df71a081d 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -17,9 +17,12 @@ set(gtest_force_shared_crt ON CACHE BOOL "" FORCE) FetchContent_MakeAvailable(googletest) enable_testing() +# Conformance defines the generate_device_binaries target which should be +# imported first +add_subdirectory(conformance) + add_subdirectory(loader) add_subdirectory(adapters) -add_subdirectory(conformance) add_subdirectory(usm) add_subdirectory(layers) add_subdirectory(unit) diff --git a/test/adapters/CMakeLists.txt b/test/adapters/CMakeLists.txt index 768a70d879..5eff6e357a 100644 --- a/test/adapters/CMakeLists.txt +++ b/test/adapters/CMakeLists.txt @@ -24,6 +24,15 @@ function(add_adapter_test name) target_compile_definitions(${target} PRIVATE ${args_FIXTURE}_ENVIRONMENT) + + if(${args_FIXTURE} STREQUAL "KERNELS") + target_compile_definitions(${target} PRIVATE KERNELS_ENVIRONMENT PRIVATE + KERNELS_DEFAULT_DIR="${UR_CONFORMANCE_DEVICE_BINARIES_DIR}") + target_include_directories(${target} + PRIVATE ${UR_CONFORMANCE_DEVICE_BINARIES_DIR}) + add_dependencies(${target} generate_device_binaries kernel_names_header) + endif() + target_link_libraries(${target} PRIVATE ${PROJECT_NAME}::loader ${PROJECT_NAME}::headers @@ -46,3 +55,7 @@ endif() if(UR_BUILD_ADAPTER_HIP OR UR_BUILD_ADAPTER_ALL) add_subdirectory(hip) endif() + +if(UR_BUILD_ADAPTER_L0 OR UR_BUILD_ADAPTER_ALL) + add_subdirectory(level_zero) +endif() diff --git a/test/adapters/level_zero/CMakeLists.txt b/test/adapters/level_zero/CMakeLists.txt new file mode 100644 index 0000000000..45d68594ed --- /dev/null +++ b/test/adapters/level_zero/CMakeLists.txt @@ -0,0 +1,27 @@ +# Copyright (C) 2024 Intel Corporation +# Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions. +# See LICENSE.TXT +# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +if(NOT UR_DPCXX) + # Tests that require kernels can't be used if we aren't generating + # device binaries + message(WARNING + "UR_DPCXX is not defined, skipping adapter tests for level_zero") +else() + add_adapter_test(level_zero + FIXTURE KERNELS + SOURCES + urProgramLink.cpp + ENVIRONMENT + "UR_ADAPTERS_FORCE_LOAD=\"$\"" + ) + + target_include_directories(test-adapter-level_zero PRIVATE + ${PROJECT_SOURCE_DIR}/source + ${PROJECT_SOURCE_DIR}/source/adapters/level_zero + ) + + add_dependencies(test-adapter-level_zero + generate_device_binaries kernel_names_header) +endif() diff --git a/test/adapters/level_zero/urProgramLink.cpp b/test/adapters/level_zero/urProgramLink.cpp new file mode 100644 index 0000000000..77ce3b8f49 --- /dev/null +++ b/test/adapters/level_zero/urProgramLink.cpp @@ -0,0 +1,31 @@ +// Copyright (C) 2024 Intel Corporation +// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions. +// See LICENSE.TXT +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include + +using urLevelZeroProgramLinkTest = uur::urProgramTest; +UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urLevelZeroProgramLinkTest); + +TEST_P(urLevelZeroProgramLinkTest, InvalidLinkOptionsPrintedInLog) { + ur_program_handle_t linked_program = nullptr; + ASSERT_SUCCESS(urProgramCompile(context, program, "-foo")); + ASSERT_EQ_RESULT( + UR_RESULT_ERROR_PROGRAM_LINK_FAILURE, + urProgramLink(context, 1, &program, "-foo", &linked_program)); + + size_t logSize; + std::vector log; + + ASSERT_SUCCESS(urProgramGetBuildInfo(linked_program, device, + UR_PROGRAM_BUILD_INFO_LOG, 0, nullptr, + &logSize)); + log.resize(logSize); + log[logSize - 1] = 'x'; + ASSERT_SUCCESS(urProgramGetBuildInfo(linked_program, device, + UR_PROGRAM_BUILD_INFO_LOG, logSize, + log.data(), nullptr)); + ASSERT_EQ(log[logSize - 1], '\0'); + ASSERT_NE(std::string{log.data()}.find("-foo"), std::string::npos); +} diff --git a/test/conformance/device_code/CMakeLists.txt b/test/conformance/device_code/CMakeLists.txt index 450733d5ed..ee16b0eb43 100644 --- a/test/conformance/device_code/CMakeLists.txt +++ b/test/conformance/device_code/CMakeLists.txt @@ -67,6 +67,15 @@ macro(add_device_binary SOURCE_FILE) if(${TRIPLE} MATCHES "amd" AND ${KERNEL_NAME} MATCHES "image_copy") continue() endif() + + # This seems to fail to build the SYCL binary due to the invalid asm + if(${TRIPLE} MATCHES "cuda" AND ${KERNEL_NAME} MATCHES "build_failure") + continue() + endif() + if(${TRIPLE} MATCHES "amd" AND ${KERNEL_NAME} MATCHES "build_failure") + continue() + endif() + add_custom_command(OUTPUT ${EXE_PATH} COMMAND ${UR_DPCXX} -fsycl -fsycl-targets=${TRIPLE} -fsycl-device-code-split=off ${AMD_TARGET_BACKEND} ${AMD_OFFLOAD_ARCH} ${AMD_NOGPULIB} @@ -93,10 +102,12 @@ add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/foo.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/image_copy.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/mean.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/spec_constant.cpp) +add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/spec_constant_multiple.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/usm_ll.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/saxpy.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/saxpy_usm.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/indexers_usm.cpp) +add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/build_failure.cpp) set(KERNEL_HEADER ${UR_CONFORMANCE_DEVICE_BINARIES_DIR}/kernel_entry_points.h) add_custom_command(OUTPUT ${KERNEL_HEADER} diff --git a/test/conformance/device_code/build_failure.cpp b/test/conformance/device_code/build_failure.cpp new file mode 100644 index 0000000000..148587c1a6 --- /dev/null +++ b/test/conformance/device_code/build_failure.cpp @@ -0,0 +1,21 @@ +// Copyright (C) 2023 Intel Corporation +// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions. +// See LICENSE.TXT +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include + +int main() { + sycl::queue deviceQueue; + + auto Kernel = []() { +#ifdef __SYCL_DEVICE_ONLY__ + asm volatile("undefined\n"); +#endif // __SYCL_DEVICE_ONLY__ + }; + + deviceQueue.submit( + [&](sycl::handler &cgh) { cgh.single_task(Kernel); }); + + return 0; +} diff --git a/test/conformance/device_code/spec_constant.cpp b/test/conformance/device_code/spec_constant.cpp index fb5c5a13a4..a5908eca10 100644 --- a/test/conformance/device_code/spec_constant.cpp +++ b/test/conformance/device_code/spec_constant.cpp @@ -8,7 +8,7 @@ using namespace sycl; -constexpr specialization_id spec_const; +constexpr specialization_id spec_const{1000}; int main() { queue myQueue; diff --git a/test/conformance/device_code/spec_constant_multiple.cpp b/test/conformance/device_code/spec_constant_multiple.cpp new file mode 100644 index 0000000000..9a8379cdc7 --- /dev/null +++ b/test/conformance/device_code/spec_constant_multiple.cpp @@ -0,0 +1,43 @@ +// Copyright (C) 2024 Intel Corporation +// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions. +// See LICENSE.TXT +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include +#include + +using namespace sycl; + +constexpr specialization_id a_spec; +constexpr specialization_id b_spec; +constexpr specialization_id c_spec; + +int main() { + queue myQueue; + uint64_t out_val = 0; + buffer out(&out_val, sycl::range<1>{1}); + + myQueue.submit([&](handler &cgh) { + accessor out_acc{out, cgh, write_only}; + + cgh.set_specialization_constant(0); + cgh.set_specialization_constant(0); + cgh.set_specialization_constant(false); + + cgh.parallel_for( + out.get_range(), [=](item<1> item_id, kernel_handler h) { + uint32_t a = h.get_specialization_constant(); + uint64_t b = h.get_specialization_constant(); + bool c = h.get_specialization_constant(); + if (c) { + out_acc[0] = b - a; + } else { + out_acc[0] = b + a; + } + }); + }); + + myQueue.wait(); + + return 0; +} diff --git a/test/conformance/kernel/kernel_adapter_native_cpu.match b/test/conformance/kernel/kernel_adapter_native_cpu.match index 4d9e3d3536..54baf7630e 100644 --- a/test/conformance/kernel/kernel_adapter_native_cpu.match +++ b/test/conformance/kernel/kernel_adapter_native_cpu.match @@ -80,6 +80,8 @@ urKernelGetInfoTest.InvalidNullPointerPropSizeRet/SYCL_NATIVE_CPU___SYCL_Native_ urKernelGetInfoTest.InvalidNullPointerPropSizeRet/SYCL_NATIVE_CPU___SYCL_Native_CPU___UR_KERNEL_INFO_PROGRAM urKernelGetInfoTest.InvalidNullPointerPropSizeRet/SYCL_NATIVE_CPU___SYCL_Native_CPU___UR_KERNEL_INFO_ATTRIBUTES urKernelGetInfoTest.InvalidNullPointerPropSizeRet/SYCL_NATIVE_CPU___SYCL_Native_CPU___UR_KERNEL_INFO_NUM_REGS +urKernelGetInfoSingleTest.KernelNameCorrect/SYCL_NATIVE_CPU___SYCL_Native_CPU_ +urKernelGetInfoSingleTest.KernelContextCorrect/SYCL_NATIVE_CPU___SYCL_Native_CPU_ urKernelGetNativeHandleTest.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU_ urKernelGetNativeHandleTest.InvalidNullHandleKernel/SYCL_NATIVE_CPU___SYCL_Native_CPU_ urKernelGetNativeHandleTest.InvalidNullPointerNativeKernel/SYCL_NATIVE_CPU___SYCL_Native_CPU_ diff --git a/test/conformance/kernel/urKernelGetInfo.cpp b/test/conformance/kernel/urKernelGetInfo.cpp index c6b3d04846..69db683a02 100644 --- a/test/conformance/kernel/urKernelGetInfo.cpp +++ b/test/conformance/kernel/urKernelGetInfo.cpp @@ -15,6 +15,9 @@ UUR_TEST_SUITE_P( UR_KERNEL_INFO_NUM_REGS), uur::deviceTestWithParamPrinter); +using urKernelGetInfoSingleTest = uur::urKernelExecutionTest; +UUR_INSTANTIATE_KERNEL_TEST_SUITE_P(urKernelGetInfoSingleTest); + TEST_P(urKernelGetInfoTest, Success) { auto property_name = getParam(); size_t property_size = 0; @@ -88,3 +91,23 @@ TEST_P(urKernelGetInfoTest, InvalidNullPointerPropSizeRet) { urKernelGetInfo(kernel, UR_KERNEL_INFO_NUM_ARGS, 0, nullptr, nullptr), UR_RESULT_ERROR_INVALID_NULL_POINTER); } + +TEST_P(urKernelGetInfoSingleTest, KernelNameCorrect) { + size_t name_size = 0; + std::vector name_data; + ASSERT_SUCCESS(urKernelGetInfo(kernel, UR_KERNEL_INFO_FUNCTION_NAME, 0, + nullptr, &name_size)); + name_data.resize(name_size); + ASSERT_SUCCESS(urKernelGetInfo(kernel, UR_KERNEL_INFO_FUNCTION_NAME, + name_size, name_data.data(), nullptr)); + ASSERT_EQ(name_data[name_size - 1], '\0'); + ASSERT_STREQ(kernel_name.c_str(), name_data.data()); +} + +TEST_P(urKernelGetInfoSingleTest, KernelContextCorrect) { + ur_context_handle_t info_context; + ASSERT_SUCCESS(urKernelGetInfo(kernel, UR_KERNEL_INFO_CONTEXT, + sizeof(ur_context_handle_t), &info_context, + nullptr)); + ASSERT_EQ(context, info_context); +} diff --git a/test/conformance/program/program_adapter_cuda.match b/test/conformance/program/program_adapter_cuda.match index 77f14ed0bb..8b917219ec 100644 --- a/test/conformance/program/program_adapter_cuda.match +++ b/test/conformance/program/program_adapter_cuda.match @@ -1,3 +1,4 @@ +urProgramBuildTest.BuildFailure/NVIDIA_CUDA_BACKEND___{{.*}}_ {{OPT}}urProgramCreateWithNativeHandleTest.InvalidNullHandleContext/NVIDIA_CUDA_BACKEND___{{.*}}_ {{OPT}}urProgramCreateWithNativeHandleTest.InvalidNullPointerProgram/NVIDIA_CUDA_BACKEND___{{.*}}_ {{OPT}}urProgramGetBuildInfoTest.Success/NVIDIA_CUDA_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_BINARY_TYPE @@ -9,6 +10,7 @@ {{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleDevice/NVIDIA_CUDA_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_OPTIONS {{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleDevice/NVIDIA_CUDA_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_LOG {{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleDevice/NVIDIA_CUDA_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_BINARY_TYPE +{{OPT}}urProgramGetBuildInfoSingleTest.LogIsNullTerminated/NVIDIA_CUDA_BACKEND___{{.*}}_ {{OPT}}urProgramGetInfoTest.Success/NVIDIA_CUDA_BACKEND___{{.*}}___UR_PROGRAM_INFO_NUM_KERNELS {{OPT}}urProgramGetInfoTest.Success/NVIDIA_CUDA_BACKEND___{{.*}}___UR_PROGRAM_INFO_KERNEL_NAMES {{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/NVIDIA_CUDA_BACKEND___{{.*}}___UR_PROGRAM_INFO_REFERENCE_COUNT @@ -22,3 +24,6 @@ {{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/NVIDIA_CUDA_BACKEND___{{.*}}___UR_PROGRAM_INFO_KERNEL_NAMES {{OPT}}urProgramLinkTest.Success/NVIDIA_CUDA_BACKEND___{{.*}}_ {{OPT}}urProgramSetSpecializationConstantsTest.Success/NVIDIA_CUDA_BACKEND___{{.*}}_ +{{OPT}}urProgramSetSpecializationConstantsTest.UseDefaultValue/NVIDIA_CUDA_BACKEND___{{.*}}_ +urProgramSetMultipleSpecializationConstantsTest.MultipleCalls/NVIDIA_CUDA_BACKEND___{{.*}}_ +urProgramSetMultipleSpecializationConstantsTest.SingleCall/NVIDIA_CUDA_BACKEND___{{.*}}_ diff --git a/test/conformance/program/program_adapter_hip.match b/test/conformance/program/program_adapter_hip.match index 9aa6e56ef6..684e066ce3 100644 --- a/test/conformance/program/program_adapter_hip.match +++ b/test/conformance/program/program_adapter_hip.match @@ -1,3 +1,4 @@ +urProgramBuildTest.BuildFailure/AMD_HIP_BACKEND___{{.*}}_ {{OPT}}urProgramCreateWithNativeHandleTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}_ {{OPT}}urProgramCreateWithNativeHandleTest.InvalidNullPointerProgram/AMD_HIP_BACKEND___{{.*}}_ {{OPT}}urProgramGetBuildInfoTest.Success/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_BINARY_TYPE @@ -9,6 +10,7 @@ {{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_OPTIONS {{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_LOG {{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_BINARY_TYPE +{{OPT}}urProgramGetBuildInfoSingleTest.LogIsNullTerminated/AMD_HIP_BACKEND___{{.*}}_ {{OPT}}urProgramGetInfoTest.Success/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_NUM_KERNELS {{OPT}}urProgramGetInfoTest.Success/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_KERNEL_NAMES {{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_REFERENCE_COUNT diff --git a/test/conformance/program/program_adapter_native_cpu.match b/test/conformance/program/program_adapter_native_cpu.match index fa17ed17f7..c509c67d3e 100644 --- a/test/conformance/program/program_adapter_native_cpu.match +++ b/test/conformance/program/program_adapter_native_cpu.match @@ -2,6 +2,7 @@ {{OPT}}urProgramBuildTest.SuccessWithOptions/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramBuildTest.InvalidNullHandleContext/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramBuildTest.InvalidNullHandleProgram/SYCL_NATIVE_CPU___SYCL_Native_CPU_ +{{OPT}}urProgramBuildTest.BuildFailure/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramCompileTest.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramCompileTest.InvalidNullHandleContext/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramCompileTest.InvalidNullHandleProgram/SYCL_NATIVE_CPU___SYCL_Native_CPU_ @@ -12,12 +13,14 @@ {{OPT}}urProgramCreateWithBinaryTest.InvalidNullPointerProgram/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramCreateWithBinaryTest.InvalidNullPointerMetadata/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramCreateWithBinaryTest.InvalidSizePropertyCount/SYCL_NATIVE_CPU___SYCL_Native_CPU_ +{{OPT}}urProgramCreateWithBinaryTest.BuildInvalidProgramBinary/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramCreateWithILTest.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramCreateWithILTest.SuccessWithProperties/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramCreateWithILTest.InvalidNullHandle/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramCreateWithILTest.InvalidNullPointerSource/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramCreateWithILTest.InvalidSizeLength/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramCreateWithILTest.InvalidNullPointerProgram/SYCL_NATIVE_CPU___SYCL_Native_CPU_ +{{OPT}}urProgramCreateWithILTest.BuildInvalidProgram/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramCreateWithNativeHandleTest.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramCreateWithNativeHandleTest.InvalidNullHandleContext/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramCreateWithNativeHandleTest.InvalidNullPointerProgram/SYCL_NATIVE_CPU___SYCL_Native_CPU_ @@ -37,6 +40,7 @@ {{OPT}}urProgramGetBuildInfoTest.InvalidEnumeration/SYCL_NATIVE_CPU___SYCL_Native_CPU___UR_PROGRAM_BUILD_INFO_OPTIONS {{OPT}}urProgramGetBuildInfoTest.InvalidEnumeration/SYCL_NATIVE_CPU___SYCL_Native_CPU___UR_PROGRAM_BUILD_INFO_LOG {{OPT}}urProgramGetBuildInfoTest.InvalidEnumeration/SYCL_NATIVE_CPU___SYCL_Native_CPU___UR_PROGRAM_BUILD_INFO_BINARY_TYPE +{{OPT}}urProgramGetBuildInfoSingleTest.LogIsNullTerminated/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramGetFunctionPointerTest.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramGetFunctionPointerTest.InvalidFunctionName/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramGetFunctionPointerTest.InvalidNullHandleDevice/SYCL_NATIVE_CPU___SYCL_Native_CPU_ @@ -112,6 +116,9 @@ {{OPT}}urProgramGetInfoTest.InvalidNullPointerPropValueRet/SYCL_NATIVE_CPU___SYCL_Native_CPU___UR_PROGRAM_INFO_BINARIES {{OPT}}urProgramGetInfoTest.InvalidNullPointerPropValueRet/SYCL_NATIVE_CPU___SYCL_Native_CPU___UR_PROGRAM_INFO_NUM_KERNELS {{OPT}}urProgramGetInfoTest.InvalidNullPointerPropValueRet/SYCL_NATIVE_CPU___SYCL_Native_CPU___UR_PROGRAM_INFO_KERNEL_NAMES +{{OPT}}urProgramGetInfoSingleTest.NumDevicesIsNonzero/SYCL_NATIVE_CPU___SYCL_Native_CPU_ +{{OPT}}urProgramGetInfoSingleTest.NumDevicesMatchesDeviceArray/SYCL_NATIVE_CPU___SYCL_Native_CPU_ +{{OPT}}urProgramGetInfoSingleTest.NumDevicesMatchesContextNumDevices/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramGetNativeHandleTest.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramGetNativeHandleTest.InvalidNullHandleProgram/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramGetNativeHandleTest.InvalidNullPointerNativeProgram/SYCL_NATIVE_CPU___SYCL_Native_CPU_ @@ -125,6 +132,9 @@ {{OPT}}urProgramRetainTest.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramRetainTest.InvalidNullHandleProgram/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramSetSpecializationConstantsTest.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU_ +{{OPT}}urProgramSetSpecializationConstantsTest.UseDefaultValue/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramSetSpecializationConstantsTest.InvalidNullHandleProgram/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramSetSpecializationConstantsTest.InvalidNullPointerSpecConstants/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramSetSpecializationConstantsTest.InvalidSizeCount/SYCL_NATIVE_CPU___SYCL_Native_CPU_ +{{OPT}}urProgramSetMultipleSpecializationConstantsTest.MultipleCalls/SYCL_NATIVE_CPU___SYCL_Native_CPU_ +{{OPT}}urProgramSetMultipleSpecializationConstantsTest.SingleCall/SYCL_NATIVE_CPU___SYCL_Native_CPU_ diff --git a/test/conformance/program/program_adapter_opencl.match b/test/conformance/program/program_adapter_opencl.match index 0d429016ee..d35966d283 100644 --- a/test/conformance/program/program_adapter_opencl.match +++ b/test/conformance/program/program_adapter_opencl.match @@ -1,3 +1,4 @@ +urProgramCreateWithILTest.BuildInvalidProgram/Intel_R__OpenCL___{{.*}}_ urProgramGetFunctionPointerTest.InvalidFunctionName/Intel_R__OpenCL___{{.*}}_ urProgramGetInfoTest.Success/Intel_R__OpenCL___{{.*}}___UR_PROGRAM_INFO_SOURCE urProgramGetInfoTest.Success/Intel_R__OpenCL___{{.*}}___UR_PROGRAM_INFO_BINARIES diff --git a/test/conformance/program/urProgramBuild.cpp b/test/conformance/program/urProgramBuild.cpp index 97e4db77e3..f99b94321e 100644 --- a/test/conformance/program/urProgramBuild.cpp +++ b/test/conformance/program/urProgramBuild.cpp @@ -26,3 +26,34 @@ TEST_P(urProgramBuildTest, InvalidNullHandleProgram) { ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_NULL_HANDLE, urProgramBuild(context, nullptr, nullptr)); } + +TEST_P(urProgramBuildTest, BuildFailure) { + ur_program_handle_t program = nullptr; + std::shared_ptr> il_binary; + uur::KernelsEnvironment::instance->LoadSource("build_failure", 0, + il_binary); + if (!il_binary) { + // The build failure we are testing for happens at SYCL compile time on + // AMD and Nvidia, so no binary exists to check for a build failure + GTEST_SKIP() << "Build failure test not supported on AMD/Nvidia yet"; + return; + } + + // TODO: This seems to fail on opencl/device combination used in the Github + // runners (`2023.16.12.0.12_195853.xmain-hotfix`). It segfaults, so we just + // skip the test so other tests can run + ur_platform_backend_t backend; + ASSERT_SUCCESS(urPlatformGetInfo(platform, UR_PLATFORM_INFO_BACKEND, + sizeof(ur_platform_backend_t), &backend, + nullptr)); + if (backend == UR_PLATFORM_BACKEND_OPENCL) { + GTEST_SKIP() << "Skipping opencl build failure test - segfaults on CI"; + } + + ASSERT_EQ_RESULT(UR_RESULT_SUCCESS, + urProgramCreateWithIL(context, il_binary->data(), + il_binary->size(), nullptr, + &program)); + ASSERT_EQ_RESULT(UR_RESULT_ERROR_PROGRAM_BUILD_FAILURE, + urProgramBuild(context, program, nullptr)); +} diff --git a/test/conformance/program/urProgramCreateWithBinary.cpp b/test/conformance/program/urProgramCreateWithBinary.cpp index 3fb6e3a268..0f525dd293 100644 --- a/test/conformance/program/urProgramCreateWithBinary.cpp +++ b/test/conformance/program/urProgramCreateWithBinary.cpp @@ -94,3 +94,13 @@ TEST_P(urProgramCreateWithBinaryTest, InvalidSizePropertyCount) { binary.data(), &properties, &binary_program)); } + +TEST_P(urProgramCreateWithBinaryTest, BuildInvalidProgramBinary) { + ur_program_handle_t program = nullptr; + uint8_t binary[] = {0, 1, 2, 3, 4}; + auto result = urProgramCreateWithBinary(context, device, 5, binary, nullptr, + &program); + // The driver is not required to reject the binary + ASSERT_TRUE(result == UR_RESULT_ERROR_INVALID_BINARY || + result == UR_RESULT_SUCCESS); +} diff --git a/test/conformance/program/urProgramCreateWithIL.cpp b/test/conformance/program/urProgramCreateWithIL.cpp index 00f41da6ef..800a43cd5d 100644 --- a/test/conformance/program/urProgramCreateWithIL.cpp +++ b/test/conformance/program/urProgramCreateWithIL.cpp @@ -73,3 +73,12 @@ TEST_P(urProgramCreateWithILTest, InvalidNullPointerProgram) { il_binary->size(), nullptr, nullptr)); } + +TEST_P(urProgramCreateWithILTest, BuildInvalidProgram) { + ur_program_handle_t program = nullptr; + char binary[] = {0, 1, 2, 3, 4}; + auto result = urProgramCreateWithIL(context, &binary, 5, nullptr, &program); + // The driver is not required to reject the binary + ASSERT_TRUE(result == UR_RESULT_ERROR_INVALID_BINARY || + result == UR_RESULT_SUCCESS); +} diff --git a/test/conformance/program/urProgramGetBuildInfo.cpp b/test/conformance/program/urProgramGetBuildInfo.cpp index ddc0ff998c..f18e2aadb7 100644 --- a/test/conformance/program/urProgramGetBuildInfo.cpp +++ b/test/conformance/program/urProgramGetBuildInfo.cpp @@ -21,6 +21,14 @@ UUR_TEST_SUITE_P(urProgramGetBuildInfoTest, UR_PROGRAM_BUILD_INFO_BINARY_TYPE), uur::deviceTestWithParamPrinter); +struct urProgramGetBuildInfoSingleTest : uur::urProgramTest { + void SetUp() override { + UUR_RETURN_ON_FATAL_FAILURE(urProgramTest::SetUp()); + ASSERT_SUCCESS(urProgramBuild(this->context, program, nullptr)); + } +}; +UUR_INSTANTIATE_KERNEL_TEST_SUITE_P(urProgramGetBuildInfoSingleTest); + TEST_P(urProgramGetBuildInfoTest, Success) { auto property_name = getParam(); size_t property_size = 0; @@ -60,3 +68,17 @@ TEST_P(urProgramGetBuildInfoTest, InvalidEnumeration) { UR_PROGRAM_BUILD_INFO_FORCE_UINT32, 0, nullptr, &propSizeOut)); } + +TEST_P(urProgramGetBuildInfoSingleTest, LogIsNullTerminated) { + size_t logSize; + std::vector log; + + ASSERT_SUCCESS(urProgramGetBuildInfo( + program, device, UR_PROGRAM_BUILD_INFO_LOG, 0, nullptr, &logSize)); + log.resize(logSize); + log[logSize - 1] = 'x'; + ASSERT_SUCCESS(urProgramGetBuildInfo(program, device, + UR_PROGRAM_BUILD_INFO_LOG, logSize, + log.data(), nullptr)); + ASSERT_EQ(log[logSize - 1], '\0'); +} diff --git a/test/conformance/program/urProgramGetInfo.cpp b/test/conformance/program/urProgramGetInfo.cpp index 2ea3e910b5..09a6dd0302 100644 --- a/test/conformance/program/urProgramGetInfo.cpp +++ b/test/conformance/program/urProgramGetInfo.cpp @@ -23,6 +23,14 @@ UUR_TEST_SUITE_P( UR_PROGRAM_INFO_KERNEL_NAMES), uur::deviceTestWithParamPrinter); +struct urProgramGetInfoSingleTest : uur::urProgramTest { + void SetUp() override { + UUR_RETURN_ON_FATAL_FAILURE(urProgramTest::SetUp()); + ASSERT_SUCCESS(urProgramBuild(this->context, program, nullptr)); + } +}; +UUR_INSTANTIATE_KERNEL_TEST_SUITE_P(urProgramGetInfoSingleTest); + TEST_P(urProgramGetInfoTest, Success) { auto property_name = getParam(); size_t property_size = 0; @@ -122,3 +130,35 @@ TEST_P(urProgramGetInfoTest, InvalidNullPointerPropValueRet) { 0, nullptr, nullptr), UR_RESULT_ERROR_INVALID_NULL_POINTER); } + +TEST_P(urProgramGetInfoSingleTest, NumDevicesIsNonzero) { + uint32_t count; + ASSERT_SUCCESS(urProgramGetInfo(program, UR_PROGRAM_INFO_NUM_DEVICES, + sizeof(uint32_t), &count, nullptr)); + ASSERT_GE(count, 1); +} + +TEST_P(urProgramGetInfoSingleTest, NumDevicesMatchesDeviceArray) { + uint32_t count; + ASSERT_SUCCESS(urProgramGetInfo(program, UR_PROGRAM_INFO_NUM_DEVICES, + sizeof(uint32_t), &count, nullptr)); + + size_t info_devices_size; + ASSERT_SUCCESS(urProgramGetInfo(program, UR_PROGRAM_INFO_DEVICES, 0, + nullptr, &info_devices_size)); + ASSERT_EQ(count, info_devices_size / sizeof(ur_device_handle_t)); +} + +TEST_P(urProgramGetInfoSingleTest, NumDevicesMatchesContextNumDevices) { + uint32_t count; + ASSERT_SUCCESS(urProgramGetInfo(program, UR_PROGRAM_INFO_NUM_DEVICES, + sizeof(uint32_t), &count, nullptr)); + + // The device count either matches the number of devices in the context or + // is 1, depending on how it was built + uint32_t info_context_devices_count; + ASSERT_SUCCESS(urContextGetInfo(context, UR_CONTEXT_INFO_NUM_DEVICES, + sizeof(uint32_t), + &info_context_devices_count, nullptr)); + ASSERT_TRUE(count == 1 || count == info_context_devices_count); +} diff --git a/test/conformance/program/urProgramSetSpecializationConstants.cpp b/test/conformance/program/urProgramSetSpecializationConstants.cpp index 4149711889..74c602cbe3 100644 --- a/test/conformance/program/urProgramSetSpecializationConstants.cpp +++ b/test/conformance/program/urProgramSetSpecializationConstants.cpp @@ -12,11 +12,23 @@ struct urProgramSetSpecializationConstantsTest : uur::urKernelExecutionTest { } uint32_t spec_value = 42; + uint32_t default_spec_value = 1000; // Must match the one in the SYCL source ur_specialization_constant_info_t info = {0, sizeof(spec_value), &spec_value}; }; UUR_INSTANTIATE_KERNEL_TEST_SUITE_P(urProgramSetSpecializationConstantsTest); +struct urProgramSetMultipleSpecializationConstantsTest + : uur::urKernelExecutionTest { + // The types of spec constants in this program are {uint32_t, uint64_t, bool} + void SetUp() override { + program_name = "spec_constant_multiple"; + UUR_RETURN_ON_FATAL_FAILURE(urProgramTest::SetUp()); + } +}; +UUR_INSTANTIATE_KERNEL_TEST_SUITE_P( + urProgramSetMultipleSpecializationConstantsTest); + TEST_P(urProgramSetSpecializationConstantsTest, Success) { ASSERT_SUCCESS(urProgramSetSpecializationConstants(program, 1, &info)); ASSERT_SUCCESS(urProgramBuild(context, program, nullptr)); @@ -31,6 +43,88 @@ TEST_P(urProgramSetSpecializationConstantsTest, Success) { ValidateBuffer(buffer, sizeof(spec_value), spec_value); } +TEST_P(urProgramSetSpecializationConstantsTest, UseDefaultValue) { + ur_platform_backend_t backend; + ASSERT_SUCCESS(urPlatformGetInfo(platform, UR_PLATFORM_INFO_BACKEND, + sizeof(ur_platform_backend_t), &backend, + nullptr)); + if (backend == UR_PLATFORM_BACKEND_CUDA) { + GTEST_FAIL() + << "This test is known to cause crashes on Nvidia; not running."; + } + + ASSERT_SUCCESS(urProgramBuild(context, program, nullptr)); + auto entry_points = + uur::KernelsEnvironment::instance->GetEntryPointNames(program_name); + kernel_name = entry_points[0]; + ASSERT_SUCCESS(urKernelCreate(program, kernel_name.data(), &kernel)); + + ur_mem_handle_t buffer; + AddBuffer1DArg(sizeof(spec_value), &buffer); + Launch1DRange(1); + ValidateBuffer(buffer, sizeof(spec_value), default_spec_value); +} + +TEST_P(urProgramSetMultipleSpecializationConstantsTest, MultipleCalls) { + uint32_t a = 100; + uint64_t b = 200; + bool c = false; + uint64_t output = 0; + + ur_specialization_constant_info_t info_a = {0, sizeof(uint32_t), &a}; + ASSERT_SUCCESS(urProgramSetSpecializationConstants(program, 1, &info_a)); + + ur_specialization_constant_info_t info_c = {2, sizeof(bool), &c}; + ASSERT_SUCCESS(urProgramSetSpecializationConstants(program, 1, &info_c)); + + ur_specialization_constant_info_t info_b = {1, sizeof(uint64_t), &b}; + ASSERT_SUCCESS(urProgramSetSpecializationConstants(program, 1, &info_b)); + + ASSERT_SUCCESS(urProgramBuild(context, program, nullptr)); + auto entry_points = + uur::KernelsEnvironment::instance->GetEntryPointNames(program_name); + kernel_name = entry_points[0]; + ASSERT_SUCCESS(urKernelCreate(program, kernel_name.data(), &kernel)); + + ur_mem_handle_t buffer; + AddBuffer1DArg(sizeof(uint64_t), &buffer); + Launch1DRange(1); + + ASSERT_SUCCESS(urEnqueueMemBufferRead(queue, buffer, true, 0, + sizeof(uint64_t), &output, 0, nullptr, + nullptr)); + ASSERT_EQ(output, 300); +} + +TEST_P(urProgramSetMultipleSpecializationConstantsTest, SingleCall) { + uint32_t a = 200; + uint64_t b = 300; + bool c = true; + uint64_t output = 0; + + ur_specialization_constant_info_t info[3] = { + {0, sizeof(uint32_t), &a}, + {2, sizeof(bool), &c}, + {1, sizeof(uint64_t), &b}, + }; + ASSERT_SUCCESS(urProgramSetSpecializationConstants(program, 3, &info[0])); + + ASSERT_SUCCESS(urProgramBuild(context, program, nullptr)); + auto entry_points = + uur::KernelsEnvironment::instance->GetEntryPointNames(program_name); + kernel_name = entry_points[0]; + ASSERT_SUCCESS(urKernelCreate(program, kernel_name.data(), &kernel)); + + ur_mem_handle_t buffer; + AddBuffer1DArg(sizeof(uint64_t), &buffer); + Launch1DRange(1); + + ASSERT_SUCCESS(urEnqueueMemBufferRead(queue, buffer, true, 0, + sizeof(uint64_t), &output, 0, nullptr, + nullptr)); + ASSERT_EQ(output, 100); +} + TEST_P(urProgramSetSpecializationConstantsTest, InvalidNullHandleProgram) { ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_NULL_HANDLE, urProgramSetSpecializationConstants(nullptr, 1, &info)); diff --git a/test/conformance/source/environment.cpp b/test/conformance/source/environment.cpp index a58b3ecdbd..2e9b7382e3 100644 --- a/test/conformance/source/environment.cpp +++ b/test/conformance/source/environment.cpp @@ -445,6 +445,7 @@ void KernelsEnvironment::LoadSource( if (source_path.empty()) { FAIL() << error; + binary_out = nullptr; } if (cached_kernels.find(source_path) != cached_kernels.end()) { @@ -458,6 +459,7 @@ void KernelsEnvironment::LoadSource( if (!source_file.is_open()) { FAIL() << "failed opening kernel path: " + source_path; + binary_out = nullptr; } size_t source_size = static_cast(source_file.tellg());