Skip to content

Commit

Permalink
Align Kernel CTS and Specification.
Browse files Browse the repository at this point in the history 
* Separate parameterized tests. Progresses oneapi-src#2290.
* urKernelGetGroupInfo CompileWorkGroupSize test modified to use distinct
  dimensions.
* urKernelSetArgPointer no longer returns invalid argument size.
* Testcases added:
    * urKernelGetGroupInfo CompileMaxWorkGroupSize
    * urKernelGetGroupInfo CompileMaxLinearWorkGroupSize
    * urKernelGetSubGroupInfo CompileNumSubGroups
    * urKernelRetain CheckReferenceCount
    * urKernelRelease CheckReferenceCount
    * urKernelSetArgMemObj InvalidEnumeration
  • Loading branch information
@isaacault
isaacault committed Jan 6, 2025
1 parent 545a01b commit f4a279b
Show file tree
Hide file tree
Showing 10 changed files with 412 additions and 109 deletions.
1 change: 0 additions & 1 deletion scripts/core/kernel.yml
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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."
Expand Down
1 change: 1 addition & 0 deletions test/conformance/device_code/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
3 changes: 2 additions & 1 deletion test/conformance/device_code/fixed_wg_size.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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>};
}
};

Expand Down
28 changes: 28 additions & 0 deletions test/conformance/device_code/max_wg_size.cpp
View file
Original file line number Diff line number Diff line change
@@ -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 <sycl/sycl.hpp>

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<class MaxWgSize>(sycl::range<3>(8, 8, 8),
KernelFunctor{});
});

myQueue.wait();
return 0;
}
182 changes: 146 additions & 36 deletions test/conformance/kernel/urKernelGetGroupInfo.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -3,56 +3,174 @@
// See LICENSE.TXT
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception

#include "ur_api.h"
#include <array>
#include <uur/fixtures.h>

using urKernelGetGroupInfoTest =
uur::urKernelTestWithParam<ur_kernel_group_info_t>;

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<ur_kernel_group_info_t>);

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<size_t, 3> 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<size_t, 3> 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<size_t, 3> max_work_group_size{2, 4, 8};
size_t max_linear_work_group_size{64};
};
UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urKernelGetGroupInfoMaxWorkGroupSizeTest);

struct urKernelGetGroupInfoTest : uur::urKernelTest {
void SetUp() override {
UUR_RETURN_ON_FATAL_FAILURE(urKernelTest::SetUp());
}
};
UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urKernelGetGroupInfoWgSizeTest);
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;
ASSERT_SUCCESS_OR_OPTIONAL_QUERY(
urKernelGetGroupInfo(kernel, device, property_name, 0, nullptr,
&property_size),
property_name);
ASSERT_EQ(property_size, 3 * sizeof(size_t));

std::vector<char> 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;
ASSERT_SUCCESS_OR_OPTIONAL_QUERY(
urKernelGetGroupInfo(kernel, device, property_name, 0, nullptr,
&property_size),
property_name);
ASSERT_EQ(property_size, sizeof(size_t));

std::vector<char> 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;
ASSERT_SUCCESS_OR_OPTIONAL_QUERY(
urKernelGetGroupInfo(kernel, device, property_name, 0, nullptr,
&property_size),
property_name);
ASSERT_EQ(property_size, 3 * sizeof(size_t));

std::array<size_t, 3> 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;
ASSERT_SUCCESS_OR_OPTIONAL_QUERY(
urKernelGetGroupInfo(kernel, device, property_name, 0, nullptr,
&property_size),
property_name);
ASSERT_EQ(property_size, sizeof(size_t));

std::vector<char> 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;
ASSERT_SUCCESS_OR_OPTIONAL_QUERY(
urKernelGetGroupInfo(kernel, device, property_name, 0, nullptr,
&property_size),
property_name);
ASSERT_EQ(property_size, sizeof(size_t));

TEST_P(urKernelGetGroupInfoTest, Success) {
auto property_name = getParam();
std::vector<char> 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;
std::vector<char> property_value;
ASSERT_SUCCESS_OR_OPTIONAL_QUERY(
urKernelGetGroupInfo(kernel, device, property_name, 0, nullptr,
&property_size),
property_name);
property_value.resize(property_size);
ASSERT_EQ(property_size, sizeof(size_t));

std::vector<char> 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;
ASSERT_SUCCESS_OR_OPTIONAL_QUERY(
urKernelGetGroupInfo(kernel, device, property_name, 0, nullptr,
&property_size),
property_name);
ASSERT_EQ(property_size, 3 * sizeof(size_t));

std::array<size_t, 3> 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;
ASSERT_SUCCESS_OR_OPTIONAL_QUERY(
urKernelGetGroupInfo(kernel, device, property_name, 0, nullptr,
&property_size),
property_name);
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) {
size_t work_group_size = 0;
ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_NULL_HANDLE,
Expand All @@ -77,15 +195,7 @@ TEST_P(urKernelGetGroupInfoTest, InvalidEnumeration) {
nullptr, &bad_enum_length));
}

TEST_P(urKernelGetGroupInfoWgSizeTest, CompileWorkGroupSize) {
std::array<size_t, 3> 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<size_t, 3> read_dims{1, 1, 1};
std::array<size_t, 3> zero{0, 0, 0};
Expand All @@ -95,7 +205,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<size_t, 3> read_dims{1, 1, 1};
std::array<size_t, 3> zero{0, 0, 0};
Expand Down
Loading

0 comments on commit f4a279b

Please sign in to comment.