Skip to content

Commit 4bd7a56

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 b97f2ea commit 4bd7a56

10 files changed

+397
-118
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: 137 additions & 45 deletions
Original file line numberDiff line numberDiff line change
@@ -3,56 +3,156 @@
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};
3935
};
40-
UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urKernelGetGroupInfoWgSizeTest);
36+
UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urKernelGetGroupInfoMaxWorkGroupSizeTest);
4137

42-
TEST_P(urKernelGetGroupInfoTest, Success) {
43-
auto property_name = getParam();
44-
size_t property_size = 0;
45-
std::vector<char> property_value;
46-
auto result = urKernelGetGroupInfo(kernel, device, property_name, 0,
47-
nullptr, &property_size);
48-
if (result == UR_RESULT_SUCCESS) {
49-
property_value.resize(property_size);
50-
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
51-
property_size,
52-
property_value.data(), nullptr));
53-
} else {
54-
ASSERT_EQ_RESULT(result, UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION);
38+
struct urKernelGetGroupInfoTest : uur::urKernelTest {
39+
void SetUp() override {
40+
UUR_RETURN_ON_FATAL_FAILURE(urKernelTest::SetUp());
5541
}
42+
};
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+
UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetGroupInfo(
49+
kernel, device, property_name, 0, nullptr, &property_size));
50+
ASSERT_EQ(property_size, 3 * sizeof(size_t));
51+
52+
std::vector<char> property_value(property_size);
53+
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
54+
property_size, property_value.data(),
55+
nullptr));
56+
}
57+
58+
TEST_P(urKernelGetGroupInfoTest, WorkGroupSize) {
59+
auto property_name = UR_KERNEL_GROUP_INFO_WORK_GROUP_SIZE;
60+
size_t property_size = 0;
61+
UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetGroupInfo(
62+
kernel, device, property_name, 0, nullptr, &property_size));
63+
ASSERT_EQ(property_size, sizeof(size_t));
64+
65+
std::vector<char> property_value(property_size);
66+
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
67+
property_size, property_value.data(),
68+
nullptr));
69+
}
70+
71+
TEST_P(urKernelGetGroupInfoFixedWorkGroupSizeTest, CompileWorkGroupSize) {
72+
auto property_name = UR_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE;
73+
size_t property_size = 0;
74+
UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetGroupInfo(
75+
kernel, device, property_name, 0, nullptr, &property_size));
76+
ASSERT_EQ(property_size, 3 * sizeof(size_t));
77+
78+
std::array<size_t, 3> property_value;
79+
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
80+
property_size, property_value.data(),
81+
nullptr));
82+
83+
ASSERT_EQ(property_value, work_group_size);
84+
}
85+
86+
TEST_P(urKernelGetGroupInfoTest, LocalMemSize) {
87+
auto property_name = UR_KERNEL_GROUP_INFO_LOCAL_MEM_SIZE;
88+
size_t property_size = 0;
89+
UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetGroupInfo(
90+
kernel, device, property_name, 0, nullptr, &property_size));
91+
ASSERT_EQ(property_size, sizeof(size_t));
92+
93+
std::vector<char> property_value(property_size);
94+
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
95+
property_size, property_value.data(),
96+
nullptr));
97+
}
98+
99+
TEST_P(urKernelGetGroupInfoTest, PreferredWorkGroupSizeMultiple) {
100+
auto property_name =
101+
UR_KERNEL_GROUP_INFO_PREFERRED_WORK_GROUP_SIZE_MULTIPLE;
102+
size_t property_size = 0;
103+
UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetGroupInfo(
104+
kernel, device, property_name, 0, nullptr, &property_size));
105+
ASSERT_EQ(property_size, sizeof(size_t));
106+
107+
std::vector<char> property_value(property_size);
108+
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
109+
property_size, property_value.data(),
110+
nullptr));
111+
}
112+
113+
TEST_P(urKernelGetGroupInfoTest, PrivateMemSize) {
114+
auto property_name = UR_KERNEL_GROUP_INFO_PRIVATE_MEM_SIZE;
115+
size_t property_size = 0;
116+
UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetGroupInfo(
117+
kernel, device, property_name, 0, nullptr, &property_size));
118+
ASSERT_EQ(property_size, sizeof(size_t));
119+
120+
std::vector<char> property_value(property_size);
121+
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
122+
property_size, property_value.data(),
123+
nullptr));
124+
}
125+
126+
TEST_P(urKernelGetGroupInfoMaxWorkGroupSizeTest, CompileMaxWorkGroupSize) {
127+
auto property_name = UR_KERNEL_GROUP_INFO_COMPILE_MAX_WORK_GROUP_SIZE;
128+
size_t property_size = 0;
129+
UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetGroupInfo(
130+
kernel, device, property_name, 0, nullptr, &property_size));
131+
ASSERT_EQ(property_size, 3 * sizeof(size_t));
132+
133+
std::array<size_t, 3> property_value;
134+
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
135+
property_size, property_value.data(),
136+
nullptr));
137+
138+
ASSERT_EQ(property_value, max_work_group_size);
139+
}
140+
141+
TEST_P(urKernelGetGroupInfoMaxWorkGroupSizeTest,
142+
CompileMaxLinearWorkGroupSize) {
143+
auto property_name =
144+
UR_KERNEL_GROUP_INFO_COMPILE_MAX_LINEAR_WORK_GROUP_SIZE;
145+
size_t property_size = 0;
146+
UUR_ASSERT_SUCCESS_OR_UNSUPPORTED(urKernelGetGroupInfo(
147+
kernel, device, property_name, 0, nullptr, &property_size));
148+
ASSERT_EQ(property_size, sizeof(size_t));
149+
150+
size_t property_value;
151+
ASSERT_SUCCESS(urKernelGetGroupInfo(kernel, device, property_name,
152+
property_size, &property_value,
153+
nullptr));
154+
155+
ASSERT_EQ(property_value, max_linear_work_group_size);
56156
}
57157

58158
TEST_P(urKernelGetGroupInfoTest, InvalidNullHandleKernel) {
@@ -79,15 +179,7 @@ TEST_P(urKernelGetGroupInfoTest, InvalidEnumeration) {
79179
nullptr, &bad_enum_length));
80180
}
81181

82-
TEST_P(urKernelGetGroupInfoWgSizeTest, CompileWorkGroupSize) {
83-
std::array<size_t, 3> read_dims{1, 1, 1};
84-
ASSERT_SUCCESS(urKernelGetGroupInfo(
85-
kernel, device, UR_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE,
86-
sizeof(read_dims), read_dims.data(), nullptr));
87-
ASSERT_EQ(read_dims, wg_size);
88-
}
89-
90-
TEST_P(urKernelGetGroupInfoSingleTest, CompileWorkGroupSizeEmpty) {
182+
TEST_P(urKernelGetGroupInfoTest, CompileWorkGroupSizeEmpty) {
91183
// Returns 0 by default when there is no specific information
92184
std::array<size_t, 3> read_dims{1, 1, 1};
93185
std::array<size_t, 3> zero{0, 0, 0};
@@ -97,7 +189,7 @@ TEST_P(urKernelGetGroupInfoSingleTest, CompileWorkGroupSizeEmpty) {
97189
ASSERT_EQ(read_dims, zero);
98190
}
99191

100-
TEST_P(urKernelGetGroupInfoSingleTest, CompileMaxWorkGroupSizeEmpty) {
192+
TEST_P(urKernelGetGroupInfoTest, CompileMaxWorkGroupSizeEmpty) {
101193
// Returns 0 by default when there is no specific information
102194
std::array<size_t, 3> read_dims{1, 1, 1};
103195
std::array<size_t, 3> zero{0, 0, 0};

0 commit comments

Comments
 (0)