From bb2dec8173bfab26c3f6b29db4e58df3adf91f79 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_cuda.match | 2 + .../kernel/kernel_adapter_level_zero.match | 2 + .../kernel/kernel_adapter_native_cpu.match | 2 + test/conformance/kernel/urKernelGetInfo.cpp | 23 +++++ .../program/program_adapter_cuda.match | 7 ++ .../program/program_adapter_hip.match | 4 + .../program/program_adapter_level_zero.match | 2 + .../program/program_adapter_native_cpu.match | 10 ++ .../program/program_adapter_opencl.match | 1 + test/conformance/program/urProgramBuild.cpp | 16 ++++ .../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 + 28 files changed, 411 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 d4ea8d3404..16647c5075 100644 --- a/include/ur_api.h +++ b/include/ur_api.h @@ -4077,6 +4077,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_ @@ -4117,6 +4119,8 @@ urProgramCreateWithIL( /// - Following a successful call to this entry point, `phProgram` will /// contain a binary of type ::UR_PROGRAM_BINARY_TYPE_COMPILED_OBJECT or /// ::UR_PROGRAM_BINARY_TYPE_LIBRARY for `hDevice`. +/// - 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 98dcc1d267..cffb7783dc 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 @@ -128,6 +129,7 @@ analogue: 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 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 01cc285752..5a03b15f9f 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_ @@ -2940,6 +2942,8 @@ ur_result_t UR_APICALL urProgramCreateWithIL( /// - Following a successful call to this entry point, `phProgram` will /// contain a binary of type ::UR_PROGRAM_BINARY_TYPE_COMPILED_OBJECT or /// ::UR_PROGRAM_BINARY_TYPE_LIBRARY for `hDevice`. +/// - 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 e3b1ba0481..17c0e5d4db 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_ @@ -2503,6 +2505,8 @@ ur_result_t UR_APICALL urProgramCreateWithIL( /// - Following a successful call to this entry point, `phProgram` will /// contain a binary of type ::UR_PROGRAM_BINARY_TYPE_COMPILED_OBJECT or /// ::UR_PROGRAM_BINARY_TYPE_LIBRARY for `hDevice`. +/// - 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_cuda.match b/test/conformance/kernel/kernel_adapter_cuda.match index 992b184a29..dc14c9fbe8 100644 --- a/test/conformance/kernel/kernel_adapter_cuda.match +++ b/test/conformance/kernel/kernel_adapter_cuda.match @@ -1,3 +1,5 @@ +urKernelGetInfoSingleTest.KernelNameCorrect/NVIDIA_CUDA_BACKEND___{{.*}}_ +urKernelGetInfoSingleTest.KernelContextCorrect/NVIDIA_CUDA_BACKEND___{{.*}}_ {{OPT}}urKernelSetArgLocalTest.Success/NVIDIA_CUDA_BACKEND___{{.*}}_ {{OPT}}urKernelSetArgLocalTest.InvalidNullHandleKernel/NVIDIA_CUDA_BACKEND___{{.*}}_ {{OPT}}urKernelSetArgLocalTest.InvalidKernelArgumentIndex/NVIDIA_CUDA_BACKEND___{{.*}}_ diff --git a/test/conformance/kernel/kernel_adapter_level_zero.match b/test/conformance/kernel/kernel_adapter_level_zero.match index 2668b6821a..5f09314cae 100644 --- a/test/conformance/kernel/kernel_adapter_level_zero.match +++ b/test/conformance/kernel/kernel_adapter_level_zero.match @@ -6,6 +6,8 @@ urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_ urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_PROGRAM urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_ATTRIBUTES urKernelGetInfoTest.InvalidSizeSmall/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}___UR_KERNEL_INFO_NUM_REGS +urKernelGetInfoSingleTest.KernelNameCorrect/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urKernelGetInfoSingleTest.KernelContextCorrect/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urKernelSetArgLocalTest.InvalidKernelArgumentIndex/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urKernelSetArgMemObjTest.InvalidKernelArgumentIndex/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urKernelSetArgPointerTest.SuccessHost/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ 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 add18cd27d..b36cedd915 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::urBaseKernelTest; +UUR_INSTANTIATE_KERNEL_TEST_SUITE_P(urKernelGetInfoSingleTest); + TEST_P(urKernelGetInfoTest, Success) { auto property_name = getParam(); size_t property_size = 0; @@ -66,3 +69,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..ab4b669cd0 100644 --- a/test/conformance/program/program_adapter_cuda.match +++ b/test/conformance/program/program_adapter_cuda.match @@ -1,5 +1,8 @@ +urProgramBuildTest.BuildFailure/NVIDIA_CUDA_BACKEND___{{.*}}_ {{OPT}}urProgramCreateWithNativeHandleTest.InvalidNullHandleContext/NVIDIA_CUDA_BACKEND___{{.*}}_ {{OPT}}urProgramCreateWithNativeHandleTest.InvalidNullPointerProgram/NVIDIA_CUDA_BACKEND___{{.*}}_ +urProgramCreateWithBinaryTest.BuildInvalidProgramBinary/NVIDIA_CUDA_BACKEND___{{.*}}_ +urProgramCreateWithILTest.BuildInvalidProgram/NVIDIA_CUDA_BACKEND___{{.*}}_ {{OPT}}urProgramGetBuildInfoTest.Success/NVIDIA_CUDA_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_BINARY_TYPE {{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleProgram/NVIDIA_CUDA_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_STATUS {{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleProgram/NVIDIA_CUDA_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_OPTIONS @@ -9,6 +12,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 +26,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..8dc6cbbebe 100644 --- a/test/conformance/program/program_adapter_hip.match +++ b/test/conformance/program/program_adapter_hip.match @@ -1,3 +1,5 @@ +urProgramBuildTest.BuildFailure/AMD_HIP_BACKEND___{{.*}}_ +urProgramCreateWithBinaryTest.BuildInvalidProgramBinary/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 +11,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 @@ -20,6 +23,7 @@ {{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_BINARIES {{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_NUM_KERNELS {{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_KERNEL_NAMES +urProgramGetInfoSingleTest.NumDevicesMatchesDeviceArray/AMD_HIP_BACKEND___{{.*}}_ {{OPT}}urProgramLinkTest.Success/AMD_HIP_BACKEND___{{.*}}_ {{OPT}}urProgramSetSpecializationConstantsTest.Success/AMD_HIP_BACKEND___{{.*}}_ {{OPT}}{{Segmentation fault|Aborted}} diff --git a/test/conformance/program/program_adapter_level_zero.match b/test/conformance/program/program_adapter_level_zero.match index 5bbdfd554c..4ca3dec4a3 100644 --- a/test/conformance/program/program_adapter_level_zero.match +++ b/test/conformance/program/program_adapter_level_zero.match @@ -1,3 +1,5 @@ +urProgramCreateWithBinaryTest.BuildInvalidProgramBinary/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero__{{.*}}_ +urProgramCreateWithILTest.BuildInvalidProgram/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero__{{.*}}_ urProgramCreateWithNativeHandleTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urProgramCreateWithNativeHandleTest.InvalidNullHandleContext/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urProgramCreateWithNativeHandleTest.InvalidNullPointerProgram/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ diff --git a/test/conformance/program/program_adapter_native_cpu.match b/test/conformance/program/program_adapter_native_cpu.match index fa17ed17f7..a3ef89c1d0 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_ +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..1c6f9b37d1 100644 --- a/test/conformance/program/urProgramBuild.cpp +++ b/test/conformance/program/urProgramBuild.cpp @@ -26,3 +26,19 @@ 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) { + return; + } + 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 80d00072e7..f4d9e0b23d 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; @@ -74,3 +82,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());