Skip to content

Commit ca4a3cb

Browse files
aelovikov-intelaelovikov-intel
committed
[NFCI] Add handler::setDeviceKernelInfo helpers
They seem to provide a better abstraction than separate `MKernelName` assignment and `setDeviceKernelInfoPtr` that describe implementation rather than the intent.
1 parent e86363f commit ca4a3cb

File tree

8 files changed

+68
-69
lines changed

8 files changed

+68
-69
lines changed

‎sycl/include/sycl/handler.hpp

Lines changed: 48 additions & 51 deletions
Original file line numberDiff line numberDiff line change
@@ -491,6 +491,25 @@ class __SYCL_EXPORT handler {
491491
"a single kernel or explicit memory operation.");
492492
}
493493

494+
template <class Kernel> void setDeviceKernelInfo(void *KernelFuncPtr) {
495+
constexpr auto Info = detail::CompileTimeKernelInfo<Kernel>;
496+
MKernelName = Info.Name;
497+
// TODO support ESIMD in no-integration-header case too.
498+
setKernelInfo(KernelFuncPtr, Info.NumParams, Info.ParamDescGetter,
499+
Info.IsESIMD, Info.HasSpecialCaptures);
500+
setDeviceKernelInfoPtr(&detail::getDeviceKernelInfo<Kernel>());
501+
setType(detail::CGType::Kernel);
502+
}
503+
504+
void setDeviceKernelInfo(kernel &&Kernel) {
505+
MKernel = detail::getSyclObjImpl(std::move(Kernel));
506+
MKernelName = getKernelName();
507+
setType(detail::CGType::Kernel);
508+
509+
// If any extra actions are added here make sure that logic around
510+
// `lambdaAndKernelHaveEqualName` calls can handle that.
511+
}
512+
494513
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
495514
// TODO: Those functions are not used anymore, remove it in the next
496515
// ABI-breaking window.
@@ -823,18 +842,27 @@ class __SYCL_EXPORT handler {
823842
detail::GetInstantiateKernelOnHostPtr<KernelType, LambdaArgType,
824843
Dims>());
825844
#endif
845+
846+
// SYCL unittests are built without sycl compiler, so "host" information
847+
// about kernels isn't provided (e.g., via integration headers or compiler
848+
// builtins).
849+
//
850+
// However, some copy/fill USM operation are implemented via SYCL kernels
851+
// and are instantiated resulting in all the `static_assert` checks being
852+
// exercised. Without kernel information that would fail, so we explicitly
853+
// disable such checks when this macro is defined. Note that the unittests
854+
// don't actually execute those operation, that's why disabling
855+
// unconditional `static_asserts`s is enough for now.
856+
#ifndef __SYCL_UNITTESTS
826857
constexpr auto Info = detail::CompileTimeKernelInfo<KernelName>;
827858

828-
constexpr bool KernelHasName = (Info.Name != std::string_view{});
859+
static_assert(Info.Name != std::string_view{}, "Kernel must have a name!");
829860

830861
// Some host compilers may have different captures from Clang. Currently
831862
// there is no stable way of handling this when extracting the captures,
832863
// so a static assert is made to fail for incompatible kernel lambdas.
833-
834-
// TODO remove the ifdef once the kernel size builtin is supported.
835-
#ifdef __INTEL_SYCL_USE_INTEGRATION_HEADERS
836864
static_assert(
837-
!KernelHasName || sizeof(KernelType) == Info.KernelSize,
865+
sizeof(KernelType) == Info.KernelSize,
838866
"Unexpected kernel lambda size. This can be caused by an "
839867
"external host compiler producing a lambda with an "
840868
"unexpected layout. This is a limitation of the compiler."
@@ -846,25 +874,8 @@ class __SYCL_EXPORT handler {
846874
"-fsycl-host-compiler-options='/std:c++latest' "
847875
"might also help.");
848876
#endif
849-
// Empty name indicates that the compilation happens without integration
850-
// header, so don't perform things that require it.
851-
if constexpr (KernelHasName) {
852-
// TODO support ESIMD in no-integration-header case too.
853-
854-
// Force hasSpecialCaptures to be evaluated at compile-time.
855-
setKernelInfo((void *)MHostKernel->getPtr(), Info.NumParams,
856-
Info.ParamDescGetter, Info.IsESIMD,
857-
Info.HasSpecialCaptures);
858-
859-
MKernelName = Info.Name;
860-
setDeviceKernelInfoPtr(&detail::getDeviceKernelInfo<KernelName>());
861-
} else {
862-
// In case w/o the integration header it is necessary to process
863-
// accessors from the list(which are associated with this handler) as
864-
// arguments. We must copy the associated accessors as they are checked
865-
// later during finalize.
866-
setArgsToAssociatedAccessors();
867-
}
877+
878+
setDeviceKernelInfo<KernelName>((void *)MHostKernel->getPtr());
868879

869880
// If the kernel lambda is callable with a kernel_handler argument, manifest
870881
// the associated kernel handler.
@@ -1306,7 +1317,6 @@ class __SYCL_EXPORT handler {
13061317
setNDRangeDescriptor(RoundedRange);
13071318
StoreLambda<KName, decltype(Wrapper), Dims, TransformedArgType>(
13081319
std::move(Wrapper));
1309-
setType(detail::CGType::Kernel);
13101320
#endif
13111321
} else
13121322
#endif // !__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ &&
@@ -1328,7 +1338,6 @@ class __SYCL_EXPORT handler {
13281338
setNDRangeDescriptor(std::move(UserRange));
13291339
StoreLambda<NameT, KernelType, Dims, TransformedArgType>(
13301340
std::move(KernelFunc));
1331-
setType(detail::CGType::Kernel);
13321341
#endif
13331342
#else
13341343
(void)KernelFunc;
@@ -1350,13 +1359,11 @@ class __SYCL_EXPORT handler {
13501359
[[maybe_unused]] kernel Kernel) {
13511360
#ifndef __SYCL_DEVICE_ONLY__
13521361
throwIfActionIsCreated();
1353-
MKernel = detail::getSyclObjImpl(std::move(Kernel));
1362+
setDeviceKernelInfo(std::move(Kernel));
13541363
detail::checkValueRange<Dims>(NumWorkItems);
13551364
setNDRangeDescriptor(std::move(NumWorkItems));
13561365
processLaunchProperties<PropertiesT>(Props);
1357-
setType(detail::CGType::Kernel);
13581366
extractArgsAndReqs();
1359-
MKernelName = getKernelName();
13601367
#endif
13611368
}
13621369

@@ -1375,13 +1382,11 @@ class __SYCL_EXPORT handler {
13751382
[[maybe_unused]] kernel Kernel) {
13761383
#ifndef __SYCL_DEVICE_ONLY__
13771384
throwIfActionIsCreated();
1378-
MKernel = detail::getSyclObjImpl(std::move(Kernel));
1385+
setDeviceKernelInfo(std::move(Kernel));
13791386
detail::checkValueRange<Dims>(NDRange);
13801387
setNDRangeDescriptor(std::move(NDRange));
13811388
processLaunchProperties(Props);
1382-
setType(detail::CGType::Kernel);
13831389
extractArgsAndReqs();
1384-
MKernelName = getKernelName();
13851390
#endif
13861391
}
13871392

@@ -1408,7 +1413,6 @@ class __SYCL_EXPORT handler {
14081413
}
14091414
throwIfActionIsCreated();
14101415
verifyUsedKernelBundleInternal(Info.Name);
1411-
setType(detail::CGType::Kernel);
14121416

14131417
detail::checkValueRange<Dims>(params...);
14141418
if constexpr (SetNumWorkGroups) {
@@ -1454,7 +1458,6 @@ class __SYCL_EXPORT handler {
14541458
// kernel.
14551459
setHandlerKernelBundle(Kernel);
14561460
verifyUsedKernelBundleInternal(Info.Name);
1457-
setType(detail::CGType::Kernel);
14581461

14591462
detail::checkValueRange<Dims>(params...);
14601463
if constexpr (SetNumWorkGroups) {
@@ -1464,12 +1467,11 @@ class __SYCL_EXPORT handler {
14641467
setNDRangeDescriptor(std::move(params)...);
14651468
}
14661469

1467-
MKernel = detail::getSyclObjImpl(std::move(Kernel));
1468-
if (!lambdaAndKernelHaveEqualName<NameT>()) {
1469-
extractArgsAndReqs();
1470-
MKernelName = getKernelName();
1471-
} else {
1470+
setDeviceKernelInfo(std::move(Kernel));
1471+
if (lambdaAndKernelHaveEqualName<NameT>()) {
14721472
StoreLambda<NameT, KernelType, Dims, ElementType>(std::move(KernelFunc));
1473+
} else {
1474+
extractArgsAndReqs();
14731475
}
14741476
processProperties<Info.IsESIMD, PropertiesT>(Props);
14751477
#endif
@@ -1849,10 +1851,8 @@ class __SYCL_EXPORT handler {
18491851
// No need to check if range is out of INT_MAX limits as it's compile-time
18501852
// known constant
18511853
setNDRangeDescriptor(range<1>{1});
1852-
MKernel = detail::getSyclObjImpl(std::move(Kernel));
1853-
setType(detail::CGType::Kernel);
1854+
setDeviceKernelInfo(std::move(Kernel));
18541855
extractArgsAndReqs();
1855-
MKernelName = getKernelName();
18561856
}
18571857

18581858
void parallel_for(range<1> NumWorkItems, kernel Kernel) {
@@ -1885,12 +1885,10 @@ class __SYCL_EXPORT handler {
18851885
[[maybe_unused]] kernel Kernel) {
18861886
#ifndef __SYCL_DEVICE_ONLY__
18871887
throwIfActionIsCreated();
1888-
MKernel = detail::getSyclObjImpl(std::move(Kernel));
1888+
setDeviceKernelInfo(std::move(Kernel));
18891889
detail::checkValueRange<Dims>(NumWorkItems, WorkItemOffset);
18901890
setNDRangeDescriptor(std::move(NumWorkItems), std::move(WorkItemOffset));
1891-
setType(detail::CGType::Kernel);
18921891
extractArgsAndReqs();
1893-
MKernelName = getKernelName();
18941892
#endif
18951893
}
18961894

@@ -1928,13 +1926,12 @@ class __SYCL_EXPORT handler {
19281926
// No need to check if range is out of INT_MAX limits as it's compile-time
19291927
// known constant
19301928
setNDRangeDescriptor(range<1>{1});
1931-
MKernel = detail::getSyclObjImpl(std::move(Kernel));
1932-
setType(detail::CGType::Kernel);
1933-
if (!lambdaAndKernelHaveEqualName<NameT>()) {
1934-
extractArgsAndReqs();
1935-
MKernelName = getKernelName();
1936-
} else
1929+
setDeviceKernelInfo(std::move(Kernel));
1930+
if (lambdaAndKernelHaveEqualName<NameT>()) {
19371931
StoreLambda<NameT, KernelType, /*Dims*/ 1, void>(std::move(KernelFunc));
1932+
} else {
1933+
extractArgsAndReqs();
1934+
}
19381935
#else
19391936
detail::CheckDeviceCopyable<KernelType>();
19401937
#endif

‎sycl/test/basic_tests/fp-accuracy.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx -%fsycl-host-only -c -ffp-accuracy=high -faltmathlib=SVMLAltMathLibrary -fno-math-errno %s
1+
// RUN: %clangxx -fsycl -c -ffp-accuracy=high -faltmathlib=SVMLAltMathLibrary -fno-math-errno %s
22

33
#include <sycl/sycl.hpp>
44

‎sycl/test/basic_tests/single_task_error_message.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
1+
// RUN: %clangxx -fsycl-device-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
22
#include <iostream>
33
#include <sycl/sycl.hpp>
44
int main() {
@@ -11,7 +11,7 @@ int main() {
1111
myQueue
1212
.single_task([&](sycl::handler &cgh) {
1313
// expected-error-re@sycl/queue.hpp:* {{static assertion failed due to requirement '{{.*}}': sycl::queue.single_task() requires a kernel instead of command group.{{.*}} Use queue.submit() instead}}
14-
// expected-error-re@sycl/detail/cg_types.hpp:* {{no matching function for call to object of type '(lambda at {{.*}}single_task_error_message.cpp:{{.*}})'}}
14+
// expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{no matching function for call to object of type 'const (lambda at {{.*}}single_task_error_message.cpp:{{.*}})'}}
1515
})
1616
.wait();
1717
}
@@ -27,7 +27,7 @@ int main() {
2727
.single_task(e,
2828
[&](sycl::handler &cgh) {
2929
// expected-error-re@sycl/queue.hpp:* {{static assertion failed due to requirement '{{.*}}': sycl::queue.single_task() requires a kernel instead of command group.{{.*}} Use queue.submit() instead}}
30-
// expected-error-re@sycl/detail/cg_types.hpp:* {{no matching function for call to object of type '(lambda at {{.*}}single_task_error_message.cpp:{{.*}})'}}
30+
// expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{no matching function for call to object of type 'const (lambda at {{.*}}single_task_error_message.cpp:{{.*}})'}}
3131
})
3232
.wait();
3333
}
@@ -43,7 +43,7 @@ int main() {
4343
.single_task(vector_event,
4444
[&](sycl::handler &cgh) {
4545
// expected-error-re@sycl/queue.hpp:* {{static assertion failed due to requirement '{{.*}}': sycl::queue.single_task() requires a kernel instead of command group.{{.*}} Use queue.submit() instead}}
46-
// expected-error-re@sycl/detail/cg_types.hpp:* {{no matching function for call to object of type '(lambda at {{.*}}single_task_error_message.cpp:{{.*}})'}}
46+
// expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{no matching function for call to object of type 'const (lambda at {{.*}}single_task_error_message.cpp:{{.*}})'}}
4747
})
4848
.wait();
4949
}

‎sycl/test/extensions/properties/non_esimd_kernel_fp_control.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
1+
// RUN: %clangxx -D__SYCL_UNITTESTS %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
22

33
#include <sycl/sycl.hpp>
44

‎sycl/test/virtual-functions/properties-negative.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
1+
// RUN: %clangxx -D__SYCL_UNITTESTS %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
22

33
#include <sycl/sycl.hpp>
44

‎sycl/test/warnings/deprecated_get_backend_info.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
1+
// RUN: %clangxx -fsycl-device-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
22
#include <iostream>
33
#include <sycl/detail/core.hpp>
44
#include <sycl/kernel_bundle.hpp>

‎sycl/test/warnings/sycl_2020_deprecations.cpp

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx %fsycl-host-only -fsyntax-only -ferror-limit=0 -sycl-std=2020 -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
1+
// RUN: %clangxx -fsycl-device-only -fsyntax-only -ferror-limit=0 -sycl-std=2020 -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
22

33
// expected-warning@CL/sycl.hpp:* {{CL/sycl.hpp is deprecated, use sycl/sycl.hpp}}
44
#include <CL/sycl.hpp>
@@ -283,23 +283,23 @@ int main() {
283283

284284
// expected-warning@+8{{'get_pointer' is deprecated: accessor::get_pointer() is deprecated, please use get_multi_ptr()}}
285285
// expected-warning@+7{{'get_pointer<sycl::access::target::device, void>' is deprecated: accessor::get_pointer() is deprecated, please use get_multi_ptr()}}
286-
// expected-warning@+4{{'make_ptr<int, sycl::access::address_space::global_space, sycl::access::decorated::legacy, void>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
286+
// expected-warning@+4{{'make_ptr<int, sycl::access::address_space::global_space, sycl::access::decorated::legacy>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
287287
sycl::multi_ptr<int, sycl::access::address_space::global_space,
288288
sycl::access::decorated::legacy>
289289
LegacyGlobalMptr =
290290
sycl::make_ptr<int, sycl::access::address_space::global_space,
291291
sycl::access::decorated::legacy>(
292292
GlobalAcc.get_pointer());
293293
// expected-warning@+7{{'get_pointer' is deprecated: local_accessor::get_pointer() is deprecated, please use get_multi_ptr()}}
294-
// expected-warning@+4{{'make_ptr<int, sycl::access::address_space::local_space, sycl::access::decorated::legacy, void>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
294+
// expected-warning@+4{{'make_ptr<int, sycl::access::address_space::local_space, sycl::access::decorated::legacy>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
295295
sycl::multi_ptr<int, sycl::access::address_space::local_space,
296296
sycl::access::decorated::legacy>
297297
LegacyLocalMptr =
298298
sycl::make_ptr<int, sycl::access::address_space::local_space,
299299
sycl::access::decorated::legacy>(
300300
LocalAcc.get_pointer());
301301

302-
// expected-warning@+4{{'make_ptr<int, sycl::access::address_space::private_space, sycl::access::decorated::legacy, void>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
302+
// expected-warning@+4{{'make_ptr<int, sycl::access::address_space::private_space, sycl::access::decorated::legacy>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
303303
sycl::multi_ptr<int, sycl::access::address_space::private_space,
304304
sycl::access::decorated::legacy>
305305
LegacyPrivateMptr =
@@ -329,27 +329,27 @@ int main() {
329329
sycl::access::decorated::yes>
330330
UndecoratedPrivateMptr = DecoratedPrivateMptr;
331331

332-
// expected-warning@+2{{'operator int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
332+
// expected-warning@+2{{'operator __global int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
333333
auto DecoratedGlobalPtr =
334334
static_cast<typename decltype(DecoratedGlobalMptr)::pointer>(
335335
DecoratedGlobalMptr);
336-
// expected-warning@+2{{'operator int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
336+
// expected-warning@+2{{'operator __local int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
337337
auto DecoratedLocalPtr =
338338
static_cast<typename decltype(DecoratedLocalMptr)::pointer>(
339339
DecoratedLocalMptr);
340-
// expected-warning@+2{{'operator int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
340+
// expected-warning@+2{{'operator __private int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
341341
auto DecoratedPrivatePtr =
342342
static_cast<typename decltype(DecoratedPrivateMptr)::pointer>(
343343
DecoratedPrivateMptr);
344-
// expected-warning@+2{{'operator int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
344+
// expected-warning@+2{{'operator __global int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
345345
auto UndecoratedGlobalPtr =
346346
static_cast<typename decltype(UndecoratedGlobalMptr)::pointer>(
347347
UndecoratedGlobalMptr);
348-
// expected-warning@+2{{'operator int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
348+
// expected-warning@+2{{'operator __local int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
349349
auto UndecoratedLocalPtr =
350350
static_cast<typename decltype(UndecoratedLocalMptr)::pointer>(
351351
UndecoratedLocalMptr);
352-
// expected-warning@+2{{'operator int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
352+
// expected-warning@+2{{'operator __private int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
353353
auto UndecoratedPrivatePtr =
354354
static_cast<typename decltype(UndecoratedPrivateMptr)::pointer>(
355355
UndecoratedPrivateMptr);

‎sycl/unittests/CMakeLists.txt

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,8 @@ endforeach()
99

1010
add_compile_definitions(SYCL2020_DISABLE_DEPRECATION_WARNINGS SYCL_DISABLE_FSYCL_SYCLHPP_WARNING)
1111

12+
add_compile_definitions(__SYCL_UNITTESTS)
13+
1214
# suppress warnings which came from Google Test sources
1315
if (CXX_SUPPORTS_SUGGEST_OVERRIDE_FLAG)
1416
add_compile_options("-Wno-suggest-override")

0 commit comments

Comments
 (0)