Skip to content

Commit 396fb20

Browse files
kbenziekbenzie
authored
Merge pull request #1501 from RossBrunton/ross/kerneltests
[Testing] Spec clarifications and testing updates for kernel
2 parents 719bb9c + e2ffea6 commit 396fb20

21 files changed

+347
-10
lines changed

‎include/ur_api.h

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4772,7 +4772,8 @@ typedef enum ur_kernel_group_info_t {
47724772
UR_KERNEL_GROUP_INFO_GLOBAL_WORK_SIZE = 0, ///< [size_t[3]] Return Work Group maximum global size
47734773
UR_KERNEL_GROUP_INFO_WORK_GROUP_SIZE = 1, ///< [size_t] Return maximum Work Group size
47744774
UR_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE = 2, ///< [size_t[3]] Return Work Group size required by the source code, such
4775-
///< as __attribute__((required_work_group_size(X,Y,Z))
4775+
///< as __attribute__((required_work_group_size(X,Y,Z)), or (0, 0, 0) if
4776+
///< unspecified
47764777
UR_KERNEL_GROUP_INFO_LOCAL_MEM_SIZE = 3, ///< [size_t] Return local memory required by the Kernel
47774778
UR_KERNEL_GROUP_INFO_PREFERRED_WORK_GROUP_SIZE_MULTIPLE = 4, ///< [size_t] Return preferred multiple of Work Group size for launch
47784779
UR_KERNEL_GROUP_INFO_PRIVATE_MEM_SIZE = 5, ///< [size_t] Return minimum amount of private memory in bytes used by each
@@ -4788,7 +4789,8 @@ typedef enum ur_kernel_group_info_t {
47884789
typedef enum ur_kernel_sub_group_info_t {
47894790
UR_KERNEL_SUB_GROUP_INFO_MAX_SUB_GROUP_SIZE = 0, ///< [uint32_t] Return maximum SubGroup size
47904791
UR_KERNEL_SUB_GROUP_INFO_MAX_NUM_SUB_GROUPS = 1, ///< [uint32_t] Return maximum number of SubGroup
4791-
UR_KERNEL_SUB_GROUP_INFO_COMPILE_NUM_SUB_GROUPS = 2, ///< [uint32_t] Return number of SubGroup required by the source code
4792+
UR_KERNEL_SUB_GROUP_INFO_COMPILE_NUM_SUB_GROUPS = 2, ///< [uint32_t] Return number of SubGroup required by the source code or 0
4793+
///< if unspecified
47924794
UR_KERNEL_SUB_GROUP_INFO_SUB_GROUP_SIZE_INTEL = 3, ///< [uint32_t] Return SubGroup size required by Intel
47934795
/// @cond
47944796
UR_KERNEL_SUB_GROUP_INFO_FORCE_UINT32 = 0x7fffffff
@@ -5989,6 +5991,7 @@ urEventSetCallback(
59895991
/// - ::UR_RESULT_ERROR_INVALID_WORK_DIMENSION
59905992
/// - ::UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE
59915993
/// - ::UR_RESULT_ERROR_INVALID_VALUE
5994+
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS - "The kernel argument values have not been specified."
59925995
/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY
59935996
/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES
59945997
UR_APIEXPORT ur_result_t UR_APICALL

‎scripts/core/enqueue.yml

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -65,6 +65,8 @@ returns:
6565
- $X_RESULT_ERROR_INVALID_WORK_DIMENSION
6666
- $X_RESULT_ERROR_INVALID_WORK_GROUP_SIZE
6767
- $X_RESULT_ERROR_INVALID_VALUE
68+
- $X_RESULT_ERROR_INVALID_KERNEL_ARGS
69+
- "The kernel argument values have not been specified."
6870
- $X_RESULT_ERROR_OUT_OF_HOST_MEMORY
6971
- $X_RESULT_ERROR_OUT_OF_RESOURCES
7072
--- #--------------------------------------------------------------------------

‎scripts/core/kernel.yml

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -135,7 +135,7 @@ etors:
135135
- name: WORK_GROUP_SIZE
136136
desc: "[size_t] Return maximum Work Group size"
137137
- name: COMPILE_WORK_GROUP_SIZE
138-
desc: "[size_t[3]] Return Work Group size required by the source code, such as __attribute__((required_work_group_size(X,Y,Z))"
138+
desc: "[size_t[3]] Return Work Group size required by the source code, such as __attribute__((required_work_group_size(X,Y,Z)), or (0, 0, 0) if unspecified"
139139
- name: LOCAL_MEM_SIZE
140140
desc: "[size_t] Return local memory required by the Kernel"
141141
- name: PREFERRED_WORK_GROUP_SIZE_MULTIPLE
@@ -154,7 +154,7 @@ etors:
154154
- name: MAX_NUM_SUB_GROUPS
155155
desc: "[uint32_t] Return maximum number of SubGroup"
156156
- name: COMPILE_NUM_SUB_GROUPS
157-
desc: "[uint32_t] Return number of SubGroup required by the source code"
157+
desc: "[uint32_t] Return number of SubGroup required by the source code or 0 if unspecified"
158158
- name: SUB_GROUP_SIZE_INTEL
159159
desc: "[uint32_t] Return SubGroup size required by Intel"
160160
--- #--------------------------------------------------------------------------

‎source/adapters/opencl/common.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -85,6 +85,8 @@ ur_result_t mapCLErrorToUR(cl_int Result) {
8585
return UR_RESULT_ERROR_IN_EVENT_LIST_EXEC_STATUS;
8686
case CL_DEVICE_NOT_AVAILABLE:
8787
return UR_RESULT_ERROR_DEVICE_NOT_AVAILABLE;
88+
case CL_INVALID_KERNEL_ARGS:
89+
return UR_RESULT_ERROR_INVALID_KERNEL_ARGS;
8890
default:
8991
return UR_RESULT_ERROR_UNKNOWN;
9092
}

‎source/loader/ur_libapi.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4768,6 +4768,7 @@ ur_result_t UR_APICALL urEventSetCallback(
47684768
/// - ::UR_RESULT_ERROR_INVALID_WORK_DIMENSION
47694769
/// - ::UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE
47704770
/// - ::UR_RESULT_ERROR_INVALID_VALUE
4771+
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS - "The kernel argument values have not been specified."
47714772
/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY
47724773
/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES
47734774
ur_result_t UR_APICALL urEnqueueKernelLaunch(

‎source/ur_api.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4040,6 +4040,7 @@ ur_result_t UR_APICALL urEventSetCallback(
40404040
/// - ::UR_RESULT_ERROR_INVALID_WORK_DIMENSION
40414041
/// - ::UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE
40424042
/// - ::UR_RESULT_ERROR_INVALID_VALUE
4043+
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS - "The kernel argument values have not been specified."
40434044
/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY
40444045
/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES
40454046
ur_result_t UR_APICALL urEnqueueKernelLaunch(

‎test/conformance/device/device_adapter_native_cpu.match

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,7 @@
11
urDeviceCreateWithNativeHandleTest.InvalidNullHandlePlatform
22
urDeviceCreateWithNativeHandleTest.InvalidNullPointerDevice
33
{{OPT}}urDeviceGetGlobalTimestampTest.SuccessSynchronizedTime
4+
urDeviceGetInfoSingleTest.MaxWorkGroupSizeIsNonzero
45
{{OPT}}urDeviceSelectBinaryTest.Success
56
urDeviceGetInfoTest.Success/UR_DEVICE_INFO_DEVICE_ID
67
urDeviceGetInfoTest.Success/UR_DEVICE_INFO_MEMORY_CLOCK_RATE

‎test/conformance/device/urDeviceGetInfo.cpp

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

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

@@ -242,6 +243,12 @@ INSTANTIATE_TEST_SUITE_P(
242243
return ss.str();
243244
});
244245

246+
struct urDeviceGetInfoSingleTest : uur::urAllDevicesTest {
247+
void SetUp() override {
248+
UUR_RETURN_ON_FATAL_FAILURE(uur::urAllDevicesTest::SetUp());
249+
}
250+
};
251+
245252
bool doesReturnArray(ur_device_info_t info_type) {
246253
if (info_type == UR_DEVICE_INFO_SUPPORTED_PARTITIONS ||
247254
info_type == UR_DEVICE_INFO_PARTITION_TYPE) {
@@ -284,15 +291,15 @@ TEST_P(urDeviceGetInfoTest, Success) {
284291
}
285292
}
286293

287-
TEST_P(urDeviceGetInfoTest, InvalidNullHandleDevice) {
294+
TEST_F(urDeviceGetInfoSingleTest, InvalidNullHandleDevice) {
288295
ur_device_type_t device_type;
289296
ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_NULL_HANDLE,
290297
urDeviceGetInfo(nullptr, UR_DEVICE_INFO_TYPE,
291298
sizeof(ur_device_type_t), &device_type,
292299
nullptr));
293300
}
294301

295-
TEST_P(urDeviceGetInfoTest, InvalidEnumerationInfoType) {
302+
TEST_F(urDeviceGetInfoSingleTest, InvalidEnumerationInfoType) {
296303
for (auto device : devices) {
297304
ur_device_type_t device_type;
298305
ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_ENUMERATION,
@@ -302,7 +309,7 @@ TEST_P(urDeviceGetInfoTest, InvalidEnumerationInfoType) {
302309
}
303310
}
304311

305-
TEST_P(urDeviceGetInfoTest, InvalidSizePropSize) {
312+
TEST_F(urDeviceGetInfoSingleTest, InvalidSizePropSize) {
306313
for (auto device : devices) {
307314
ur_device_type_t device_type;
308315
ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_SIZE,
@@ -311,7 +318,7 @@ TEST_P(urDeviceGetInfoTest, InvalidSizePropSize) {
311318
}
312319
}
313320

314-
TEST_P(urDeviceGetInfoTest, InvalidSizePropSizeSmall) {
321+
TEST_F(urDeviceGetInfoSingleTest, InvalidSizePropSizeSmall) {
315322
for (auto device : devices) {
316323
ur_device_type_t device_type;
317324
ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_SIZE,
@@ -321,7 +328,7 @@ TEST_P(urDeviceGetInfoTest, InvalidSizePropSizeSmall) {
321328
}
322329
}
323330

324-
TEST_P(urDeviceGetInfoTest, InvalidNullPointerPropValue) {
331+
TEST_F(urDeviceGetInfoSingleTest, InvalidNullPointerPropValue) {
325332
for (auto device : devices) {
326333
ur_device_type_t device_type;
327334
ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_NULL_POINTER,
@@ -331,10 +338,30 @@ TEST_P(urDeviceGetInfoTest, InvalidNullPointerPropValue) {
331338
}
332339
}
333340

334-
TEST_P(urDeviceGetInfoTest, InvalidNullPointerPropSizeRet) {
341+
TEST_F(urDeviceGetInfoSingleTest, InvalidNullPointerPropSizeRet) {
335342
for (auto device : devices) {
336343
ASSERT_EQ_RESULT(
337344
UR_RESULT_ERROR_INVALID_NULL_POINTER,
338345
urDeviceGetInfo(device, UR_DEVICE_INFO_TYPE, 0, nullptr, nullptr));
339346
}
340347
}
348+
349+
TEST_F(urDeviceGetInfoSingleTest, MaxWorkGroupSizeIsNonzero) {
350+
for (auto device : devices) {
351+
size_t max_global_size;
352+
353+
ASSERT_SUCCESS(
354+
urDeviceGetInfo(device, UR_DEVICE_INFO_MAX_WORK_GROUP_SIZE,
355+
sizeof(size_t), &max_global_size, nullptr));
356+
ASSERT_NE(max_global_size, 0);
357+
358+
std::array<size_t, 3> max_work_group_sizes;
359+
ASSERT_SUCCESS(urDeviceGetInfo(device,
360+
UR_DEVICE_INFO_MAX_WORK_GROUPS_3D,
361+
sizeof(max_work_group_sizes),
362+
max_work_group_sizes.data(), nullptr));
363+
for (size_t i = 0; i < 3; i++) {
364+
ASSERT_NE(max_work_group_sizes[i], 0);
365+
}
366+
}
367+
}

‎test/conformance/device_code/CMakeLists.txt

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -93,6 +93,13 @@ macro(add_device_binary SOURCE_FILE)
9393
continue()
9494
endif()
9595

96+
# HIP doesn't seem to provide the symbol
97+
# `_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E11FixedSgSize` which
98+
# causes a build failure here
99+
if(${TRIPLE} MATCHES "amd" AND ${KERNEL_NAME} MATCHES "subgroup")
100+
continue()
101+
endif()
102+
96103
add_custom_command(OUTPUT "${BIN_PATH}"
97104
COMMAND ${UR_DPCXX} -fsycl -fsycl-targets=${TRIPLE} -fsycl-device-code-split=off
98105
${AMD_TARGET_BACKEND} ${AMD_OFFLOAD_ARCH} ${AMD_NOGPULIB}
@@ -139,6 +146,9 @@ add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/saxpy.cpp)
139146
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/saxpy_usm.cpp)
140147
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/indexers_usm.cpp)
141148
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/build_failure.cpp)
149+
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/fixed_wg_size.cpp)
150+
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/standard_types.cpp)
151+
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/subgroup.cpp)
142152

143153
set(KERNEL_HEADER ${UR_CONFORMANCE_DEVICE_BINARIES_DIR}/kernel_entry_points.h)
144154
add_custom_command(OUTPUT ${KERNEL_HEADER}

‎test/conformance/device_code/fixed_wg_size.cpp

Lines changed: 27 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,27 @@
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::work_group_size<4, 4, 4>};
15+
}
16+
};
17+
18+
int main() {
19+
sycl::queue myQueue;
20+
myQueue.submit([&](sycl::handler &cgh) {
21+
cgh.parallel_for<class FixedWgSize>(sycl::range<3>(8, 8, 8),
22+
KernelFunctor{});
23+
});
24+
25+
myQueue.wait();
26+
return 0;
27+
}

0 commit comments

Comments
 (0)