Skip to content

Commit 1918293

Browse files
isaacaultisaacault
authored
Merge pull request #2444 from isaacault/kernel-cts
Reduce gap between Kernel CTS and Specification.
2 parents 3e62cc9 + 9c2b040 commit 1918293

24 files changed

+526
-136
lines changed

‎include/ur_api.h

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -5188,7 +5188,6 @@ typedef struct ur_kernel_arg_pointer_properties_t {
51885188
/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE
51895189
/// + `NULL == hKernel`
51905190
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX
5191-
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_SIZE
51925191
UR_APIEXPORT ur_result_t UR_APICALL
51935192
urKernelSetArgPointer(
51945193
ur_kernel_handle_t hKernel, ///< [in] handle of the kernel object

‎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."

‎source/adapters/level_zero/kernel.cpp

Lines changed: 6 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -813,12 +813,13 @@ ur_result_t urKernelGetGroupInfo(
813813
(ZeKernelDevice, &kernelProperties));
814814
if (ZeResult || workGroupProperties.maxGroupSize == 0) {
815815
return ReturnValue(
816-
uint64_t{Device->ZeDeviceComputeProperties->maxTotalGroupSize});
816+
size_t{Device->ZeDeviceComputeProperties->maxTotalGroupSize});
817817
}
818-
return ReturnValue(workGroupProperties.maxGroupSize);
818+
// Specification states this returns a size_t.
819+
return ReturnValue(size_t{workGroupProperties.maxGroupSize});
819820
} else {
820821
return ReturnValue(
821-
uint64_t{Device->ZeDeviceComputeProperties->maxTotalGroupSize});
822+
size_t{Device->ZeDeviceComputeProperties->maxTotalGroupSize});
822823
}
823824
}
824825
case UR_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE: {
@@ -830,12 +831,12 @@ ur_result_t urKernelGetGroupInfo(
830831
return ReturnValue(WgSize);
831832
}
832833
case UR_KERNEL_GROUP_INFO_LOCAL_MEM_SIZE:
833-
return ReturnValue(uint32_t{Kernel->ZeKernelProperties->localMemSize});
834+
return ReturnValue(size_t{Kernel->ZeKernelProperties->localMemSize});
834835
case UR_KERNEL_GROUP_INFO_PREFERRED_WORK_GROUP_SIZE_MULTIPLE: {
835836
return ReturnValue(size_t{Device->ZeDeviceProperties->physicalEUSimdWidth});
836837
}
837838
case UR_KERNEL_GROUP_INFO_PRIVATE_MEM_SIZE: {
838-
return ReturnValue(uint32_t{Kernel->ZeKernelProperties->privateMemSize});
839+
return ReturnValue(size_t{Kernel->ZeKernelProperties->privateMemSize});
839840
}
840841
case UR_KERNEL_GROUP_INFO_COMPILE_MAX_WORK_GROUP_SIZE:
841842
case UR_KERNEL_GROUP_INFO_COMPILE_MAX_LINEAR_WORK_GROUP_SIZE:

‎source/adapters/level_zero/v2/kernel.cpp

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -512,10 +512,11 @@ ur_result_t urKernelGetGroupInfo(
512512
ZE_CALL_NOCHECK(zeKernelGetProperties, (zeDevice, &kernelProperties));
513513
if (zeResult == ZE_RESULT_SUCCESS &&
514514
workGroupProperties.maxGroupSize != 0) {
515-
return returnValue(workGroupProperties.maxGroupSize);
515+
// Specification states this returns a size_t.
516+
return returnValue(size_t{workGroupProperties.maxGroupSize});
516517
}
517518
return returnValue(
518-
uint64_t{hDevice->ZeDeviceComputeProperties->maxTotalGroupSize});
519+
size_t{hDevice->ZeDeviceComputeProperties->maxTotalGroupSize});
519520
}
520521
case UR_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE: {
521522
auto props = hKernel->getProperties(hDevice);
@@ -527,15 +528,15 @@ ur_result_t urKernelGetGroupInfo(
527528
}
528529
case UR_KERNEL_GROUP_INFO_LOCAL_MEM_SIZE: {
529530
auto props = hKernel->getProperties(hDevice);
530-
return returnValue(uint32_t{props.localMemSize});
531+
return returnValue(size_t{props.localMemSize});
531532
}
532533
case UR_KERNEL_GROUP_INFO_PREFERRED_WORK_GROUP_SIZE_MULTIPLE: {
533534
return returnValue(
534535
size_t{hDevice->ZeDeviceProperties->physicalEUSimdWidth});
535536
}
536537
case UR_KERNEL_GROUP_INFO_PRIVATE_MEM_SIZE: {
537538
auto props = hKernel->getProperties(hDevice);
538-
return returnValue(uint32_t{props.privateMemSize});
539+
return returnValue(size_t{props.privateMemSize});
539540
}
540541
case UR_KERNEL_GROUP_INFO_COMPILE_MAX_WORK_GROUP_SIZE:
541542
case UR_KERNEL_GROUP_INFO_COMPILE_MAX_LINEAR_WORK_GROUP_SIZE:

‎source/loader/ur_libapi.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3983,7 +3983,6 @@ ur_result_t UR_APICALL urKernelRelease(
39833983
/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE
39843984
/// + `NULL == hKernel`
39853985
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX
3986-
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_SIZE
39873986
ur_result_t UR_APICALL urKernelSetArgPointer(
39883987
ur_kernel_handle_t hKernel, ///< [in] handle of the kernel object
39893988
uint32_t argIndex, ///< [in] argument index in range [0, num args - 1]

‎source/ur_api.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3395,7 +3395,6 @@ ur_result_t UR_APICALL urKernelRelease(
33953395
/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE
33963396
/// + `NULL == hKernel`
33973397
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX
3398-
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_SIZE
33993398
ur_result_t UR_APICALL urKernelSetArgPointer(
34003399
ur_kernel_handle_t hKernel, ///< [in] handle of the kernel object
34013400
uint32_t argIndex, ///< [in] argument index in range [0, num args - 1]

‎test/conformance/device_code/CMakeLists.txt

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -159,7 +159,9 @@ add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/saxpy.cpp)
159159
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)
162+
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/fixed_sg_size.cpp)
162163
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/fixed_wg_size.cpp)
164+
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/max_wg_size.cpp)
163165
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/sequence.cpp)
164166
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/standard_types.cpp)
165167
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/subgroup.cpp)

‎test/conformance/device_code/fixed_sg_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::sub_group_size<8>};
15+
}
16+
};
17+
18+
int main() {
19+
sycl::queue myQueue;
20+
myQueue.submit([&](sycl::handler &cgh) {
21+
cgh.parallel_for<class FixedSgSize>(sycl::range<3>(8, 8, 8),
22+
KernelFunctor{});
23+
});
24+
25+
myQueue.wait();
26+
return 0;
27+
}

‎test/conformance/device_code/fixed_wg_size.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,7 @@ 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>};
1515
}
1616
};
1717

‎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+
}

0 commit comments

Comments
 (0)