Skip to content

Commit f4a279b

Browse files
isaacaultisaacault
committed
Align Kernel CTS and Specification.
* 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
1 parent 545a01b commit f4a279b

File tree

10 files changed

+412
-109
lines changed

10 files changed

+412
-109
lines changed

‎scripts/core/kernel.yml

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -352,7 +352,6 @@ params:
352352
desc: "[in][optional] Pointer obtained by USM allocation or virtual memory mapping operation. If null then argument value is considered null."
353353
returns:
354354
- $X_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX
355-
- $X_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_SIZE
356355
--- #--------------------------------------------------------------------------
357356
type: struct
358357
desc: "Properties for for $xKernelSetExecInfo."

‎test/conformance/device_code/CMakeLists.txt

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -160,6 +160,7 @@ add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/saxpy_usm.cpp)
160160
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/indexers_usm.cpp)
161161
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/build_failure.cpp)
162162
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/fixed_wg_size.cpp)
163+
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/max_wg_size.cpp)
163164
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/sequence.cpp)
164165
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/standard_types.cpp)
165166
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/subgroup.cpp)

‎test/conformance/device_code/fixed_wg_size.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,8 @@ struct KernelFunctor {
1111

1212
auto get(sycl::ext::oneapi::experimental::properties_tag) {
1313
return sycl::ext::oneapi::experimental::properties{
14-
sycl::ext::oneapi::experimental::work_group_size<4, 4, 4>};
14+
sycl::ext::oneapi::experimental::work_group_size<8, 4, 2>,
15+
sycl::ext::oneapi::experimental::sub_group_size<8>};
1516
}
1617
};
1718

‎test/conformance/device_code/max_wg_size.cpp

Lines changed: 28 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,28 @@
1+
// Copyright (C) 2024 Intel Corporation
2+
// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions.
3+
// See LICENSE.TXT
4+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
5+
6+
#include <sycl/sycl.hpp>
7+
8+
struct KernelFunctor {
9+
void operator()(sycl::nd_item<3>) const {}
10+
void operator()(sycl::item<3>) const {}
11+
12+
auto get(sycl::ext::oneapi::experimental::properties_tag) {
13+
return sycl::ext::oneapi::experimental::properties{
14+
sycl::ext::oneapi::experimental::max_work_group_size<8, 4, 2>,
15+
sycl::ext::oneapi::experimental::max_linear_work_group_size<64>};
16+
}
17+
};
18+
19+
int main() {
20+
sycl::queue myQueue;
21+
myQueue.submit([&](sycl::handler &cgh) {
22+
cgh.parallel_for<class MaxWgSize>(sycl::range<3>(8, 8, 8),
23+
KernelFunctor{});
24+
});
25+
26+
myQueue.wait();
27+
return 0;
28+
}

‎test/conformance/kernel/urKernelGetGroupInfo.cpp

Lines changed: 146 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -3,56 +3,174 @@
33
// See LICENSE.TXT
44
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
55

6+
#include "ur_api.h"
67
#include <array>
78
#include <uur/fixtures.h>
89

9-
using urKernelGetGroupInfoTest =
10-
uur::urKernelTestWithParam<ur_kernel_group_info_t>;
11-
12-
UUR_TEST_SUITE_P(
13-
urKernelGetGroupInfoTest,
14-
::testing::Values(UR_KERNEL_GROUP_INFO_GLOBAL_WORK_SIZE,
15-
UR_KERNEL_GROUP_INFO_WORK_GROUP_SIZE,
16-
UR_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE,
17-
UR_KERNEL_GROUP_INFO_LOCAL_MEM_SIZE,
18-
UR_KERNEL_GROUP_INFO_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
19-
UR_KERNEL_GROUP_INFO_PRIVATE_MEM_SIZE,
20-
UR_KERNEL_GROUP_INFO_COMPILE_MAX_WORK_GROUP_SIZE,
21-
UR_KERNEL_GROUP_INFO_COMPILE_MAX_LINEAR_WORK_GROUP_SIZE),
22-
uur::deviceTestWithParamPrinter<ur_kernel_group_info_t>);
23-
24-
struct urKernelGetGroupInfoSingleTest : uur::urKernelTest {
10+
struct urKernelGetGroupInfoFixedWorkGroupSizeTest : uur::urKernelTest {
2511
void SetUp() override {
12+
program_name = "fixed_wg_size";
2613
UUR_RETURN_ON_FATAL_FAILURE(urKernelTest::SetUp());
2714
}
15+
16+
// This value correlates to work_group_size<8, 4, 2> in fixed_wg_size.cpp.
17+
// In SYCL, the right-most dimension varies the fastest in linearization.
18+
// In UR, this is on the left, so we reverse the order of these values.
19+
std::array<size_t, 3> work_group_size{2, 4, 8};
2820
};
29-
UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urKernelGetGroupInfoSingleTest);
21+
UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urKernelGetGroupInfoFixedWorkGroupSizeTest);
3022

31-
struct urKernelGetGroupInfoWgSizeTest : uur::urKernelTest {
23+
struct urKernelGetGroupInfoMaxWorkGroupSizeTest : uur::urKernelTest {
3224
void SetUp() override {
33-
program_name = "fixed_wg_size";
25+
program_name = "max_wg_size";
3426
UUR_RETURN_ON_FATAL_FAILURE(urKernelTest::SetUp());
3527
}
3628

37-
// This must match the size in fixed_wg_size.cpp
38-
std::array<size_t, 3> wg_size{4, 4, 4};
29+
// These values correlate to max_work_group_size<6, 5, 4> and
30+
// max_linear_work_group_size<120> in max_wg_size.cpp.
31+
// In SYCL, the right-most dimension varies the fastest in linearization.
32+
// In UR, this is on the left, so we reverse the order of these values.
33+
std::array<size_t, 3> max_work_group_size{2, 4, 8};
34+
size_t max_linear_work_group_size{64};
35+
};
36+
UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urKernelGetGroupInfoMaxWorkGroupSizeTest);
37+
38+
struct urKernelGetGroupInfoTest : uur::urKernelTest {
39+
void SetUp() override {
40+
UUR_RETURN_ON_FATAL_FAILURE(urKernelTest::SetUp());
41+
}
3942
};
40-
UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urKernelGetGroupInfoWgSizeTest);
43+
UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urKernelGetGroupInfoTest);
44+
45+
TEST_P(urKernelGetGroupInfoTest, GlobalWorkSize) {
46+
auto property_name = UR_KERNEL_GROUP_INFO_GLOBAL_WORK_SIZE;
47+
size_t property_size = 0;
48+
ASSERT_SUCCESS_OR_OPTIONAL_QUERY(
49+
urKernelGetGroupInfo(kernel, device, property_name, 0, nullptr,
50+
&property_size),
51+
property_name);
52+
ASSERT_EQ(property_size, 3 * sizeof(size_t));
53+
54+
std::vector<char> property_value(property_size);
55+
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
56+
property_size, property_value.data(),
57+
nullptr));
58+
}
59+
60+
TEST_P(urKernelGetGroupInfoTest, WorkGroupSize) {
61+
auto property_name = UR_KERNEL_GROUP_INFO_WORK_GROUP_SIZE;
62+
size_t property_size = 0;
63+
ASSERT_SUCCESS_OR_OPTIONAL_QUERY(
64+
urKernelGetGroupInfo(kernel, device, property_name, 0, nullptr,
65+
&property_size),
66+
property_name);
67+
ASSERT_EQ(property_size, sizeof(size_t));
68+
69+
std::vector<char> property_value(property_size);
70+
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
71+
property_size, property_value.data(),
72+
nullptr));
73+
}
74+
75+
TEST_P(urKernelGetGroupInfoFixedWorkGroupSizeTest, CompileWorkGroupSize) {
76+
auto property_name = UR_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE;
77+
size_t property_size = 0;
78+
ASSERT_SUCCESS_OR_OPTIONAL_QUERY(
79+
urKernelGetGroupInfo(kernel, device, property_name, 0, nullptr,
80+
&property_size),
81+
property_name);
82+
ASSERT_EQ(property_size, 3 * sizeof(size_t));
83+
84+
std::array<size_t, 3> property_value;
85+
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
86+
property_size, property_value.data(),
87+
nullptr));
88+
89+
ASSERT_EQ(property_value, work_group_size);
90+
}
91+
92+
TEST_P(urKernelGetGroupInfoTest, LocalMemSize) {
93+
auto property_name = UR_KERNEL_GROUP_INFO_LOCAL_MEM_SIZE;
94+
size_t property_size = 0;
95+
ASSERT_SUCCESS_OR_OPTIONAL_QUERY(
96+
urKernelGetGroupInfo(kernel, device, property_name, 0, nullptr,
97+
&property_size),
98+
property_name);
99+
ASSERT_EQ(property_size, sizeof(size_t));
100+
101+
std::vector<char> property_value(property_size);
102+
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
103+
property_size, property_value.data(),
104+
nullptr));
105+
}
106+
107+
TEST_P(urKernelGetGroupInfoTest, PreferredWorkGroupSizeMultiple) {
108+
auto property_name =
109+
UR_KERNEL_GROUP_INFO_PREFERRED_WORK_GROUP_SIZE_MULTIPLE;
110+
size_t property_size = 0;
111+
ASSERT_SUCCESS_OR_OPTIONAL_QUERY(
112+
urKernelGetGroupInfo(kernel, device, property_name, 0, nullptr,
113+
&property_size),
114+
property_name);
115+
ASSERT_EQ(property_size, sizeof(size_t));
41116

42-
TEST_P(urKernelGetGroupInfoTest, Success) {
43-
auto property_name = getParam();
117+
std::vector<char> property_value(property_size);
118+
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
119+
property_size, property_value.data(),
120+
nullptr));
121+
}
122+
123+
TEST_P(urKernelGetGroupInfoTest, PrivateMemSize) {
124+
auto property_name = UR_KERNEL_GROUP_INFO_PRIVATE_MEM_SIZE;
44125
size_t property_size = 0;
45-
std::vector<char> property_value;
46126
ASSERT_SUCCESS_OR_OPTIONAL_QUERY(
47127
urKernelGetGroupInfo(kernel, device, property_name, 0, nullptr,
48128
&property_size),
49129
property_name);
50-
property_value.resize(property_size);
130+
ASSERT_EQ(property_size, sizeof(size_t));
131+
132+
std::vector<char> property_value(property_size);
51133
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
52134
property_size, property_value.data(),
53135
nullptr));
54136
}
55137

138+
TEST_P(urKernelGetGroupInfoMaxWorkGroupSizeTest, CompileMaxWorkGroupSize) {
139+
auto property_name = UR_KERNEL_GROUP_INFO_COMPILE_MAX_WORK_GROUP_SIZE;
140+
size_t property_size = 0;
141+
ASSERT_SUCCESS_OR_OPTIONAL_QUERY(
142+
urKernelGetGroupInfo(kernel, device, property_name, 0, nullptr,
143+
&property_size),
144+
property_name);
145+
ASSERT_EQ(property_size, 3 * sizeof(size_t));
146+
147+
std::array<size_t, 3> property_value;
148+
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
149+
property_size, property_value.data(),
150+
nullptr));
151+
152+
ASSERT_EQ(property_value, max_work_group_size);
153+
}
154+
155+
TEST_P(urKernelGetGroupInfoMaxWorkGroupSizeTest,
156+
CompileMaxLinearWorkGroupSize) {
157+
auto property_name =
158+
UR_KERNEL_GROUP_INFO_COMPILE_MAX_LINEAR_WORK_GROUP_SIZE;
159+
size_t property_size = 0;
160+
ASSERT_SUCCESS_OR_OPTIONAL_QUERY(
161+
urKernelGetGroupInfo(kernel, device, property_name, 0, nullptr,
162+
&property_size),
163+
property_name);
164+
ASSERT_EQ(property_size, sizeof(size_t));
165+
166+
size_t property_value;
167+
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
168+
property_size, &property_value,
169+
nullptr));
170+
171+
ASSERT_EQ(property_value, max_linear_work_group_size);
172+
}
173+
56174
TEST_P(urKernelGetGroupInfoTest, InvalidNullHandleKernel) {
57175
size_t work_group_size = 0;
58176
ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_NULL_HANDLE,
@@ -77,15 +195,7 @@ TEST_P(urKernelGetGroupInfoTest, InvalidEnumeration) {
77195
nullptr, &bad_enum_length));
78196
}
79197

80-
TEST_P(urKernelGetGroupInfoWgSizeTest, CompileWorkGroupSize) {
81-
std::array<size_t, 3> read_dims{1, 1, 1};
82-
ASSERT_SUCCESS(urKernelGetGroupInfo(
83-
kernel, device, UR_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE,
84-
sizeof(read_dims), read_dims.data(), nullptr));
85-
ASSERT_EQ(read_dims, wg_size);
86-
}
87-
88-
TEST_P(urKernelGetGroupInfoSingleTest, CompileWorkGroupSizeEmpty) {
198+
TEST_P(urKernelGetGroupInfoTest, CompileWorkGroupSizeEmpty) {
89199
// Returns 0 by default when there is no specific information
90200
std::array<size_t, 3> read_dims{1, 1, 1};
91201
std::array<size_t, 3> zero{0, 0, 0};
@@ -95,7 +205,7 @@ TEST_P(urKernelGetGroupInfoSingleTest, CompileWorkGroupSizeEmpty) {
95205
ASSERT_EQ(read_dims, zero);
96206
}
97207

98-
TEST_P(urKernelGetGroupInfoSingleTest, CompileMaxWorkGroupSizeEmpty) {
208+
TEST_P(urKernelGetGroupInfoTest, CompileMaxWorkGroupSizeEmpty) {
99209
// Returns 0 by default when there is no specific information
100210
std::array<size_t, 3> read_dims{1, 1, 1};
101211
std::array<size_t, 3> zero{0, 0, 0};

0 commit comments

Comments
 (0)