diff --git a/scripts/core/kernel.yml b/scripts/core/kernel.yml index 5bd95e1847..6ed5d6c352 100644 --- a/scripts/core/kernel.yml +++ b/scripts/core/kernel.yml @@ -352,7 +352,6 @@ params: desc: "[in][optional] Pointer obtained by USM allocation or virtual memory mapping operation. If null then argument value is considered null." returns: - $X_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX - - $X_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_SIZE --- #-------------------------------------------------------------------------- type: struct desc: "Properties for for $xKernelSetExecInfo." diff --git a/test/conformance/device_code/CMakeLists.txt b/test/conformance/device_code/CMakeLists.txt index a59cbae7b0..b3aae52d26 100644 --- a/test/conformance/device_code/CMakeLists.txt +++ b/test/conformance/device_code/CMakeLists.txt @@ -160,6 +160,7 @@ 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) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/fixed_wg_size.cpp) +add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/max_wg_size.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/sequence.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/standard_types.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/subgroup.cpp) diff --git a/test/conformance/device_code/fixed_wg_size.cpp b/test/conformance/device_code/fixed_wg_size.cpp index db2e8a9250..651a261649 100644 --- a/test/conformance/device_code/fixed_wg_size.cpp +++ b/test/conformance/device_code/fixed_wg_size.cpp @@ -11,7 +11,8 @@ struct KernelFunctor { auto get(sycl::ext::oneapi::experimental::properties_tag) { return sycl::ext::oneapi::experimental::properties{ - sycl::ext::oneapi::experimental::work_group_size<4, 4, 4>}; + sycl::ext::oneapi::experimental::work_group_size<8, 4, 2>, + sycl::ext::oneapi::experimental::sub_group_size<8>}; } }; diff --git a/test/conformance/device_code/max_wg_size.cpp b/test/conformance/device_code/max_wg_size.cpp new file mode 100644 index 0000000000..a5f0c62524 --- /dev/null +++ b/test/conformance/device_code/max_wg_size.cpp @@ -0,0 +1,28 @@ +// 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 + +struct KernelFunctor { + void operator()(sycl::nd_item<3>) const {} + void operator()(sycl::item<3>) const {} + + auto get(sycl::ext::oneapi::experimental::properties_tag) { + return sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::max_work_group_size<8, 4, 2>, + sycl::ext::oneapi::experimental::max_linear_work_group_size<64>}; + } +}; + +int main() { + sycl::queue myQueue; + myQueue.submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::range<3>(8, 8, 8), + KernelFunctor{}); + }); + + myQueue.wait(); + return 0; +} diff --git a/test/conformance/kernel/urKernelGetGroupInfo.cpp b/test/conformance/kernel/urKernelGetGroupInfo.cpp index 2b3c70c22e..f4b3c3934e 100644 --- a/test/conformance/kernel/urKernelGetGroupInfo.cpp +++ b/test/conformance/kernel/urKernelGetGroupInfo.cpp @@ -3,56 +3,156 @@ // See LICENSE.TXT // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +#include "ur_api.h" #include #include -using urKernelGetGroupInfoTest = - uur::urKernelTestWithParam; - -UUR_TEST_SUITE_P( - urKernelGetGroupInfoTest, - ::testing::Values(UR_KERNEL_GROUP_INFO_GLOBAL_WORK_SIZE, - UR_KERNEL_GROUP_INFO_WORK_GROUP_SIZE, - UR_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE, - UR_KERNEL_GROUP_INFO_LOCAL_MEM_SIZE, - UR_KERNEL_GROUP_INFO_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, - UR_KERNEL_GROUP_INFO_PRIVATE_MEM_SIZE, - UR_KERNEL_GROUP_INFO_COMPILE_MAX_WORK_GROUP_SIZE, - UR_KERNEL_GROUP_INFO_COMPILE_MAX_LINEAR_WORK_GROUP_SIZE), - uur::deviceTestWithParamPrinter); - -struct urKernelGetGroupInfoSingleTest : uur::urKernelTest { +struct urKernelGetGroupInfoFixedWorkGroupSizeTest : uur::urKernelTest { void SetUp() override { + program_name = "fixed_wg_size"; UUR_RETURN_ON_FATAL_FAILURE(urKernelTest::SetUp()); } + + // This value correlates to work_group_size<8, 4, 2> in fixed_wg_size.cpp. + // In SYCL, the right-most dimension varies the fastest in linearization. + // In UR, this is on the left, so we reverse the order of these values. + std::array work_group_size{2, 4, 8}; }; -UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urKernelGetGroupInfoSingleTest); +UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urKernelGetGroupInfoFixedWorkGroupSizeTest); -struct urKernelGetGroupInfoWgSizeTest : uur::urKernelTest { +struct urKernelGetGroupInfoMaxWorkGroupSizeTest : uur::urKernelTest { void SetUp() override { - program_name = "fixed_wg_size"; + program_name = "max_wg_size"; UUR_RETURN_ON_FATAL_FAILURE(urKernelTest::SetUp()); } - // This must match the size in fixed_wg_size.cpp - std::array wg_size{4, 4, 4}; + // These values correlate to max_work_group_size<6, 5, 4> and + // max_linear_work_group_size<120> in max_wg_size.cpp. + // In SYCL, the right-most dimension varies the fastest in linearization. + // In UR, this is on the left, so we reverse the order of these values. + std::array max_work_group_size{2, 4, 8}; + size_t max_linear_work_group_size{64}; }; -UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urKernelGetGroupInfoWgSizeTest); +UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urKernelGetGroupInfoMaxWorkGroupSizeTest); -TEST_P(urKernelGetGroupInfoTest, Success) { - auto property_name = getParam(); - size_t property_size = 0; - std::vector property_value; - auto result = urKernelGetGroupInfo(kernel, device, property_name, 0, - nullptr, &property_size); - if (result == UR_RESULT_SUCCESS) { - property_value.resize(property_size); - ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name, - property_size, - property_value.data(), nullptr)); - } else { - ASSERT_EQ_RESULT(result, UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION); +struct urKernelGetGroupInfoTest : uur::urKernelTest { + void SetUp() override { + UUR_RETURN_ON_FATAL_FAILURE(urKernelTest::SetUp()); } +}; +UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urKernelGetGroupInfoTest); + +TEST_P(urKernelGetGroupInfoTest, GlobalWorkSize) { + auto property_name = UR_KERNEL_GROUP_INFO_GLOBAL_WORK_SIZE; + size_t property_size = 0; + UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetGroupInfo( + kernel, device, property_name, 0, nullptr, &property_size)); + ASSERT_EQ(property_size, 3 * sizeof(size_t)); + + std::vector property_value(property_size); + ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name, + property_size, property_value.data(), + nullptr)); +} + +TEST_P(urKernelGetGroupInfoTest, WorkGroupSize) { + auto property_name = UR_KERNEL_GROUP_INFO_WORK_GROUP_SIZE; + size_t property_size = 0; + UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetGroupInfo( + kernel, device, property_name, 0, nullptr, &property_size)); + ASSERT_EQ(property_size, sizeof(size_t)); + + std::vector property_value(property_size); + ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name, + property_size, property_value.data(), + nullptr)); +} + +TEST_P(urKernelGetGroupInfoFixedWorkGroupSizeTest, CompileWorkGroupSize) { + auto property_name = UR_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE; + size_t property_size = 0; + UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetGroupInfo( + kernel, device, property_name, 0, nullptr, &property_size)); + ASSERT_EQ(property_size, 3 * sizeof(size_t)); + + std::array property_value; + ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name, + property_size, property_value.data(), + nullptr)); + + ASSERT_EQ(property_value, work_group_size); +} + +TEST_P(urKernelGetGroupInfoTest, LocalMemSize) { + auto property_name = UR_KERNEL_GROUP_INFO_LOCAL_MEM_SIZE; + size_t property_size = 0; + UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetGroupInfo( + kernel, device, property_name, 0, nullptr, &property_size)); + ASSERT_EQ(property_size, sizeof(size_t)); + + std::vector property_value(property_size); + ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name, + property_size, property_value.data(), + nullptr)); +} + +TEST_P(urKernelGetGroupInfoTest, PreferredWorkGroupSizeMultiple) { + auto property_name = + UR_KERNEL_GROUP_INFO_PREFERRED_WORK_GROUP_SIZE_MULTIPLE; + size_t property_size = 0; + UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetGroupInfo( + kernel, device, property_name, 0, nullptr, &property_size)); + ASSERT_EQ(property_size, sizeof(size_t)); + + std::vector property_value(property_size); + ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name, + property_size, property_value.data(), + nullptr)); +} + +TEST_P(urKernelGetGroupInfoTest, PrivateMemSize) { + auto property_name = UR_KERNEL_GROUP_INFO_PRIVATE_MEM_SIZE; + size_t property_size = 0; + UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetGroupInfo( + kernel, device, property_name, 0, nullptr, &property_size)); + ASSERT_EQ(property_size, sizeof(size_t)); + + std::vector property_value(property_size); + ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name, + property_size, property_value.data(), + nullptr)); +} + +TEST_P(urKernelGetGroupInfoMaxWorkGroupSizeTest, CompileMaxWorkGroupSize) { + auto property_name = UR_KERNEL_GROUP_INFO_COMPILE_MAX_WORK_GROUP_SIZE; + size_t property_size = 0; + UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetGroupInfo( + kernel, device, property_name, 0, nullptr, &property_size)); + ASSERT_EQ(property_size, 3 * sizeof(size_t)); + + std::array property_value; + ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name, + property_size, property_value.data(), + nullptr)); + + ASSERT_EQ(property_value, max_work_group_size); +} + +TEST_P(urKernelGetGroupInfoMaxWorkGroupSizeTest, + CompileMaxLinearWorkGroupSize) { + auto property_name = + UR_KERNEL_GROUP_INFO_COMPILE_MAX_LINEAR_WORK_GROUP_SIZE; + size_t property_size = 0; + UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetGroupInfo( + kernel, device, property_name, 0, nullptr, &property_size)); + ASSERT_EQ(property_size, sizeof(size_t)); + + size_t property_value; + ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name, + property_size, &property_value, + nullptr)); + + ASSERT_EQ(property_value, max_linear_work_group_size); } TEST_P(urKernelGetGroupInfoTest, InvalidNullHandleKernel) { @@ -79,15 +179,7 @@ TEST_P(urKernelGetGroupInfoTest, InvalidEnumeration) { nullptr, &bad_enum_length)); } -TEST_P(urKernelGetGroupInfoWgSizeTest, CompileWorkGroupSize) { - std::array read_dims{1, 1, 1}; - ASSERT_SUCCESS(urKernelGetGroupInfo( - kernel, device, UR_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE, - sizeof(read_dims), read_dims.data(), nullptr)); - ASSERT_EQ(read_dims, wg_size); -} - -TEST_P(urKernelGetGroupInfoSingleTest, CompileWorkGroupSizeEmpty) { +TEST_P(urKernelGetGroupInfoTest, CompileWorkGroupSizeEmpty) { // Returns 0 by default when there is no specific information std::array read_dims{1, 1, 1}; std::array zero{0, 0, 0}; @@ -97,7 +189,7 @@ TEST_P(urKernelGetGroupInfoSingleTest, CompileWorkGroupSizeEmpty) { ASSERT_EQ(read_dims, zero); } -TEST_P(urKernelGetGroupInfoSingleTest, CompileMaxWorkGroupSizeEmpty) { +TEST_P(urKernelGetGroupInfoTest, CompileMaxWorkGroupSizeEmpty) { // Returns 0 by default when there is no specific information std::array read_dims{1, 1, 1}; std::array zero{0, 0, 0}; diff --git a/test/conformance/kernel/urKernelGetInfo.cpp b/test/conformance/kernel/urKernelGetInfo.cpp index e87ab1da13..11ef9b75f1 100644 --- a/test/conformance/kernel/urKernelGetInfo.cpp +++ b/test/conformance/kernel/urKernelGetInfo.cpp @@ -3,70 +3,128 @@ // See LICENSE.TXT // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +#include "ur_api.h" #include -using urKernelGetInfoTest = uur::urKernelTestWithParam; +using urKernelGetInfoTest = uur::urKernelTest; +UUR_INSTANTIATE_KERNEL_TEST_SUITE_P(urKernelGetInfoTest); -UUR_TEST_SUITE_P( - urKernelGetInfoTest, - ::testing::Values(UR_KERNEL_INFO_FUNCTION_NAME, UR_KERNEL_INFO_NUM_ARGS, - UR_KERNEL_INFO_REFERENCE_COUNT, UR_KERNEL_INFO_CONTEXT, - UR_KERNEL_INFO_PROGRAM, UR_KERNEL_INFO_ATTRIBUTES, - UR_KERNEL_INFO_NUM_REGS), - uur::deviceTestWithParamPrinter); +TEST_P(urKernelGetInfoTest, FunctionName) { + auto property_name = UR_KERNEL_INFO_FUNCTION_NAME; + size_t property_size = 0; + + ASSERT_SUCCESS( + urKernelGetInfo(kernel, property_name, 0, nullptr, &property_size)); + ASSERT_GT(property_size, 0U); -using urKernelGetInfoSingleTest = uur::urKernelExecutionTest; -UUR_INSTANTIATE_KERNEL_TEST_SUITE_P(urKernelGetInfoSingleTest); + std::vector property_value(property_size); + ASSERT_SUCCESS(urKernelGetInfo(kernel, property_name, property_size, + property_value.data(), nullptr)); +} -TEST_P(urKernelGetInfoTest, Success) { - auto property_name = getParam(); +TEST_P(urKernelGetInfoTest, NumArgs) { + auto property_name = UR_KERNEL_INFO_NUM_ARGS; size_t property_size = 0; - std::vector property_value; + ASSERT_SUCCESS( urKernelGetInfo(kernel, property_name, 0, nullptr, &property_size)); - property_value.resize(property_size); + ASSERT_EQ(property_size, sizeof(uint32_t)); + + std::vector property_value(property_size); ASSERT_SUCCESS(urKernelGetInfo(kernel, property_name, property_size, property_value.data(), nullptr)); - switch (property_name) { - case UR_KERNEL_INFO_CONTEXT: { - auto returned_context = - reinterpret_cast(property_value.data()); - ASSERT_EQ(context, *returned_context); - break; - } - case UR_KERNEL_INFO_PROGRAM: { - auto returned_program = - reinterpret_cast(property_value.data()); - ASSERT_EQ(program, *returned_program); - break; - } - case UR_KERNEL_INFO_REFERENCE_COUNT: { - auto returned_reference_count = - reinterpret_cast(property_value.data()); - ASSERT_GT(*returned_reference_count, 0U); - break; - } - case UR_KERNEL_INFO_ATTRIBUTES: { - auto returned_attributes = std::string(property_value.data()); - ur_platform_backend_t backend; - ASSERT_SUCCESS(urPlatformGetInfo(platform, UR_PLATFORM_INFO_BACKEND, - sizeof(backend), &backend, nullptr)); - if (backend == UR_PLATFORM_BACKEND_OPENCL || - backend == UR_PLATFORM_BACKEND_LEVEL_ZERO) { - // Older intel drivers don't attach any default attributes and newer ones force walk order to X/Y/Z using special attribute. - ASSERT_TRUE(returned_attributes.empty() || - returned_attributes == - "intel_reqd_workgroup_walk_order(0,1,2)"); - } else { - ASSERT_TRUE(returned_attributes.empty()); - } - break; - } - default: - break; +} + +TEST_P(urKernelGetInfoTest, ReferenceCount) { + auto property_name = UR_KERNEL_INFO_REFERENCE_COUNT; + size_t property_size = 0; + + ASSERT_SUCCESS( + urKernelGetInfo(kernel, property_name, 0, nullptr, &property_size)); + ASSERT_EQ(property_size, sizeof(uint32_t)); + + std::vector property_value(property_size); + ASSERT_SUCCESS(urKernelGetInfo(kernel, property_name, property_size, + property_value.data(), nullptr)); + + auto returned_reference_count = + reinterpret_cast(property_value.data()); + ASSERT_GT(*returned_reference_count, 0U); +} + +TEST_P(urKernelGetInfoTest, Context) { + auto property_name = UR_KERNEL_INFO_CONTEXT; + size_t property_size = 0; + + ASSERT_SUCCESS( + urKernelGetInfo(kernel, property_name, 0, nullptr, &property_size)); + ASSERT_EQ(property_size, sizeof(ur_context_handle_t)); + + std::vector property_value(property_size); + ASSERT_SUCCESS(urKernelGetInfo(kernel, property_name, property_size, + property_value.data(), nullptr)); + + auto returned_context = + reinterpret_cast(property_value.data()); + ASSERT_EQ(context, *returned_context); +} + +TEST_P(urKernelGetInfoTest, Program) { + auto property_name = UR_KERNEL_INFO_PROGRAM; + size_t property_size = 0; + + ASSERT_SUCCESS( + urKernelGetInfo(kernel, property_name, 0, nullptr, &property_size)); + ASSERT_EQ(property_size, sizeof(ur_program_handle_t)); + + std::vector property_value(property_size); + ASSERT_SUCCESS(urKernelGetInfo(kernel, property_name, property_size, + property_value.data(), nullptr)); + + auto returned_program = + reinterpret_cast(property_value.data()); + ASSERT_EQ(program, *returned_program); +} + +TEST_P(urKernelGetInfoTest, Attributes) { + auto property_name = UR_KERNEL_INFO_ATTRIBUTES; + size_t property_size = 0; + + ASSERT_SUCCESS( + urKernelGetInfo(kernel, property_name, 0, nullptr, &property_size)); + + std::vector property_value(property_size); + ASSERT_SUCCESS(urKernelGetInfo(kernel, property_name, property_size, + property_value.data(), nullptr)); + + auto returned_attributes = std::string(property_value.data()); + ur_platform_backend_t backend; + ASSERT_SUCCESS(urPlatformGetInfo(platform, UR_PLATFORM_INFO_BACKEND, + sizeof(backend), &backend, nullptr)); + if (backend == UR_PLATFORM_BACKEND_OPENCL || + backend == UR_PLATFORM_BACKEND_LEVEL_ZERO) { + // Older intel drivers don't attach any default attributes and newer ones force walk order to X/Y/Z using special attribute. + ASSERT_TRUE(returned_attributes.empty() || + returned_attributes == + "intel_reqd_workgroup_walk_order(0,1,2)"); + } else { + ASSERT_TRUE(returned_attributes.empty()); } } +TEST_P(urKernelGetInfoTest, NumRegs) { + auto property_name = UR_KERNEL_INFO_NUM_REGS; + size_t property_size = 0; + + UUR_ASSERT_SUCCESS_OR_UNSUPPORTED( + urKernelGetInfo(kernel, property_name, 0, nullptr, &property_size)); + ASSERT_EQ(property_size, sizeof(uint32_t)); + + std::vector property_value(property_size); + ASSERT_SUCCESS(urKernelGetInfo(kernel, property_name, property_size, + property_value.data(), nullptr)); +} + TEST_P(urKernelGetInfoTest, InvalidNullHandleKernel) { size_t kernel_name_length = 0; ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_NULL_HANDLE, @@ -117,7 +175,7 @@ TEST_P(urKernelGetInfoTest, InvalidNullPointerPropSizeRet) { UR_RESULT_ERROR_INVALID_NULL_POINTER); } -TEST_P(urKernelGetInfoSingleTest, KernelNameCorrect) { +TEST_P(urKernelGetInfoTest, KernelNameCorrect) { size_t name_size = 0; std::vector name_data; ASSERT_SUCCESS(urKernelGetInfo(kernel, UR_KERNEL_INFO_FUNCTION_NAME, 0, @@ -129,7 +187,7 @@ TEST_P(urKernelGetInfoSingleTest, KernelNameCorrect) { ASSERT_STREQ(kernel_name.c_str(), name_data.data()); } -TEST_P(urKernelGetInfoSingleTest, KernelContextCorrect) { +TEST_P(urKernelGetInfoTest, KernelContextCorrect) { ur_context_handle_t info_context; ASSERT_SUCCESS(urKernelGetInfo(kernel, UR_KERNEL_INFO_CONTEXT, sizeof(ur_context_handle_t), &info_context, diff --git a/test/conformance/kernel/urKernelGetSubGroupInfo.cpp b/test/conformance/kernel/urKernelGetSubGroupInfo.cpp index fa4e045483..9c661bd2e4 100644 --- a/test/conformance/kernel/urKernelGetSubGroupInfo.cpp +++ b/test/conformance/kernel/urKernelGetSubGroupInfo.cpp @@ -3,33 +3,76 @@ // See LICENSE.TXT // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +#include "ur_api.h" #include -using urKernelGetSubGroupInfoTest = - uur::urKernelTestWithParam; +struct urKernelGetSubGroupInfoFixedSubGroupSizeTest : uur::urKernelTest { + void SetUp() override { + program_name = "fixed_wg_size"; + UUR_RETURN_ON_FATAL_FAILURE(urKernelTest::SetUp()); + } -UUR_TEST_SUITE_P( - urKernelGetSubGroupInfoTest, - ::testing::Values(UR_KERNEL_SUB_GROUP_INFO_MAX_SUB_GROUP_SIZE, - UR_KERNEL_SUB_GROUP_INFO_MAX_NUM_SUB_GROUPS, - UR_KERNEL_SUB_GROUP_INFO_COMPILE_NUM_SUB_GROUPS, - UR_KERNEL_SUB_GROUP_INFO_SUB_GROUP_SIZE_INTEL), - uur::deviceTestWithParamPrinter); + // This value correlates to sub_group_size<8> in fixed_wg_size.cpp. + uint32_t num_sub_groups{8}; +}; +UUR_INSTANTIATE_DEVICE_TEST_SUITE_P( + urKernelGetSubGroupInfoFixedSubGroupSizeTest); -struct urKernelGetSubGroupInfoSingleTest : uur::urKernelTest { +struct urKernelGetSubGroupInfoTest : uur::urKernelTest { void SetUp() override { UUR_RETURN_ON_FATAL_FAILURE(urKernelTest::SetUp()); } }; -UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urKernelGetSubGroupInfoSingleTest); +UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urKernelGetSubGroupInfoTest); -TEST_P(urKernelGetSubGroupInfoTest, Success) { - auto property_name = getParam(); +TEST_P(urKernelGetSubGroupInfoTest, MaxSubGroupSize) { + auto property_name = UR_KERNEL_SUB_GROUP_INFO_MAX_SUB_GROUP_SIZE; size_t property_size = 0; - std::vector property_value; - ASSERT_SUCCESS(urKernelGetSubGroupInfo(kernel, device, property_name, 0, - nullptr, &property_size)); - property_value.resize(property_size); + UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetSubGroupInfo( + kernel, device, property_name, 0, nullptr, &property_size)); + ASSERT_EQ(property_size, sizeof(uint32_t)); + + std::vector property_value(property_size); + ASSERT_SUCCESS(urKernelGetSubGroupInfo(kernel, device, property_name, + property_size, property_value.data(), + nullptr)); +} + +TEST_P(urKernelGetSubGroupInfoTest, MaxNumSubGroups) { + auto property_name = UR_KERNEL_SUB_GROUP_INFO_MAX_NUM_SUB_GROUPS; + size_t property_size = 0; + UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetSubGroupInfo( + kernel, device, property_name, 0, nullptr, &property_size)); + ASSERT_EQ(property_size, sizeof(uint32_t)); + + std::vector property_value(property_size); + ASSERT_SUCCESS(urKernelGetSubGroupInfo(kernel, device, property_name, + property_size, property_value.data(), + nullptr)); +} + +TEST_P(urKernelGetSubGroupInfoFixedSubGroupSizeTest, CompileNumSubGroups) { + auto property_name = UR_KERNEL_SUB_GROUP_INFO_COMPILE_NUM_SUB_GROUPS; + size_t property_size = 0; + UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetSubGroupInfo( + kernel, device, property_name, 0, nullptr, &property_size)); + ASSERT_EQ(property_size, sizeof(uint32_t)); + + uint32_t property_value; + ASSERT_SUCCESS(urKernelGetSubGroupInfo(kernel, device, property_name, + property_size, &property_value, + nullptr)); + ASSERT_EQ(property_value, num_sub_groups); +} + +TEST_P(urKernelGetSubGroupInfoTest, SubGroupSizeIntel) { + auto property_name = UR_KERNEL_SUB_GROUP_INFO_SUB_GROUP_SIZE_INTEL; + size_t property_size = 0; + UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetSubGroupInfo( + kernel, device, property_name, 0, nullptr, &property_size)); + ASSERT_EQ(property_size, sizeof(uint32_t)); + + std::vector property_value(property_size); ASSERT_SUCCESS(urKernelGetSubGroupInfo(kernel, device, property_name, property_size, property_value.data(), nullptr)); @@ -61,7 +104,7 @@ TEST_P(urKernelGetSubGroupInfoTest, InvalidEnumeration) { 0, nullptr, &bad_enum_length)); } -TEST_P(urKernelGetSubGroupInfoSingleTest, CompileNumSubgroupsIsZero) { +TEST_P(urKernelGetSubGroupInfoTest, CompileNumSubgroupsIsZero) { // Returns 0 by default when there is no specific information size_t subgroups = 1; ASSERT_SUCCESS(urKernelGetSubGroupInfo( diff --git a/test/conformance/kernel/urKernelRelease.cpp b/test/conformance/kernel/urKernelRelease.cpp index 3e2078d98c..cda79617d8 100644 --- a/test/conformance/kernel/urKernelRelease.cpp +++ b/test/conformance/kernel/urKernelRelease.cpp @@ -24,3 +24,24 @@ TEST_P(urKernelReleaseTest, InvalidNullHandleKernel) { ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_NULL_HANDLE, urKernelRelease(nullptr)); } + +TEST_P(urKernelReleaseTest, CheckReferenceCount) { + uint32_t referenceCount = 0; + ASSERT_SUCCESS(urKernelGetInfo(kernel, UR_KERNEL_INFO_REFERENCE_COUNT, + sizeof(referenceCount), &referenceCount, + nullptr)); + ASSERT_EQ(referenceCount, 1); + + ASSERT_SUCCESS(urKernelRetain(kernel)); + ASSERT_SUCCESS(urKernelGetInfo(kernel, UR_KERNEL_INFO_REFERENCE_COUNT, + sizeof(referenceCount), &referenceCount, + nullptr)); + ASSERT_EQ(referenceCount, 2); + + ASSERT_SUCCESS(urKernelRelease(kernel)); + + ASSERT_SUCCESS(urKernelGetInfo(kernel, UR_KERNEL_INFO_REFERENCE_COUNT, + sizeof(referenceCount), &referenceCount, + nullptr)); + ASSERT_EQ(referenceCount, 1); +} diff --git a/test/conformance/kernel/urKernelRetain.cpp b/test/conformance/kernel/urKernelRetain.cpp index cf1958104c..1da0ea0867 100644 --- a/test/conformance/kernel/urKernelRetain.cpp +++ b/test/conformance/kernel/urKernelRetain.cpp @@ -3,6 +3,7 @@ // See LICENSE.TXT // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +#include "uur/checks.h" #include using urKernelRetainTest = uur::urKernelTest; @@ -17,3 +18,27 @@ TEST_P(urKernelRetainTest, InvalidNullHandleKernel) { ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_NULL_HANDLE, urKernelRetain(nullptr)); } + +TEST_P(urKernelRetainTest, CheckReferenceCount) { + uint32_t referenceCount = 0; + ASSERT_SUCCESS(urKernelGetInfo(kernel, UR_KERNEL_INFO_REFERENCE_COUNT, + sizeof(referenceCount), &referenceCount, + nullptr)); + ASSERT_EQ(referenceCount, 1); + + { + ASSERT_SUCCESS(urKernelRetain(kernel)); + } + + ASSERT_SUCCESS(urKernelGetInfo(kernel, UR_KERNEL_INFO_REFERENCE_COUNT, + sizeof(referenceCount), &referenceCount, + nullptr)); + ASSERT_EQ(referenceCount, 2); + + ASSERT_SUCCESS(urKernelRelease(kernel)); + + ASSERT_SUCCESS(urKernelGetInfo(kernel, UR_KERNEL_INFO_REFERENCE_COUNT, + sizeof(referenceCount), &referenceCount, + nullptr)); + ASSERT_EQ(referenceCount, 1); +} diff --git a/test/conformance/kernel/urKernelSetArgMemObj.cpp b/test/conformance/kernel/urKernelSetArgMemObj.cpp index 0a83882465..afa00d3345 100644 --- a/test/conformance/kernel/urKernelSetArgMemObj.cpp +++ b/test/conformance/kernel/urKernelSetArgMemObj.cpp @@ -3,6 +3,7 @@ // See LICENSE.TXT // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +#include "ur_api.h" #include struct urKernelSetArgMemObjTest : uur::urKernelTest { @@ -43,3 +44,13 @@ TEST_P(urKernelSetArgMemObjTest, InvalidKernelArgumentIndex) { UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX, urKernelSetArgMemObj(kernel, num_kernel_args + 1, nullptr, buffer)); } + +TEST_P(urKernelSetArgMemObjTest, InvalidEnumeration) { + ur_kernel_arg_mem_obj_properties_t props{ + UR_STRUCTURE_TYPE_KERNEL_ARG_MEM_OBJ_PROPERTIES, /* stype */ + nullptr, /* pNext */ + UR_MEM_FLAG_FORCE_UINT32 /* memoryAccess */ + }; + ASSERT_EQ_RESULT(urKernelSetArgMemObj(kernel, 0, &props, buffer), + UR_RESULT_ERROR_INVALID_ENUMERATION); +}