Skip to content

Commit 18c0d92

Browse files
aelovikov-intelaelovikov-intel
committed
[SYCL] Follow-up for #19990
Downstream CI uncovered another scenario where `detail::CompileTimeKernelInfo<KernelName>.Name` might be empty (other than non-SYCL compilation in some of the [unit]tests). That happens when compiling the same TU using multiple offload models (e.g., both SYCL and OMP) and the device compilation for non-SYCL model doesn't use SYCL integration header to provide kernel information. Ideally, we should be just guarding `static_assert`s with some `#if __SYCL_WHATEVER` but we don't set any SYCL-related macro when using 3rd party host compilers. Deal with that by turning compile-time `static_assert` into a run-time `assert` unless we can guarantee that SYCL kernel information is available.
1 parent 9d75ffa commit 18c0d92

File tree

9 files changed

+88
-43
lines changed

9 files changed

+88
-43
lines changed

‎sycl/include/sycl/handler.hpp‎

Lines changed: 10 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -836,26 +836,24 @@ class __SYCL_EXPORT handler {
836836
Dims>());
837837
#endif
838838

839-
// SYCL unittests are built without sycl compiler, so "host" information
840-
// about kernels isn't provided (e.g., via integration headers or compiler
841-
// builtins).
842-
//
843-
// However, some copy/fill USM operation are implemented via SYCL kernels
844-
// and are instantiated resulting in all the `static_assert` checks being
845-
// exercised. Without kernel information that would fail, so we explicitly
846-
// disable such checks when this macro is defined. Note that the unittests
847-
// don't actually execute those operation, that's why disabling
848-
// unconditional `static_asserts`s is enough for now.
849-
#ifndef __SYCL_UNITTESTS_BYPASS_KERNEL_NAME_CHECK
850839
constexpr auto Info = detail::CompileTimeKernelInfo<KernelName>;
851840

841+
#ifdef SYCL_LANGUAGE_VERSION
852842
static_assert(Info.Name != std::string_view{}, "Kernel must have a name!");
843+
#else
844+
// Either non-SYCL compilation (e.g., OpenMP device code generation) or
845+
// using custom host compiler. If the former, the condition will be false,
846+
// but we won't be executing this code. If the latter, the condition must
847+
// hold and we *will* execute this.
848+
assert(Info.Name != std::string_view{} && "Kernel must have a name!");
849+
#endif
853850

854851
// Some host compilers may have different captures from Clang. Currently
855852
// there is no stable way of handling this when extracting the captures,
856853
// so a static assert is made to fail for incompatible kernel lambdas.
857854
static_assert(
858-
sizeof(KernelType) == Info.KernelSize,
855+
Info.Name == std::string_view{} ||
856+
sizeof(KernelType) == Info.KernelSize,
859857
"Unexpected kernel lambda size. This can be caused by an "
860858
"external host compiler producing a lambda with an "
861859
"unexpected layout. This is a limitation of the compiler."
@@ -866,7 +864,6 @@ class __SYCL_EXPORT handler {
866864
"In case of MSVC, passing "
867865
"-fsycl-host-compiler-options='/std:c++latest' "
868866
"might also help.");
869-
#endif
870867

871868
setDeviceKernelInfo<KernelName>((void *)MHostKernel->getPtr());
872869

‎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: 38 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,13 +1,16 @@
1-
// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
1+
// Disable range rounding so that we won't have to emulate integration header
2+
// for that too.
3+
4+
// RUN: %clangxx %fsycl-host-only -D__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__=1 -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
25

36
#include <sycl/sycl.hpp>
47

58
using namespace sycl;
69
namespace syclex = sycl::ext::oneapi::experimental;
710
namespace intelex = sycl::ext::intel::experimental;
811

9-
struct ESIMDKernel {
10-
ESIMDKernel() {}
12+
struct KernelWithProps {
13+
KernelWithProps() {}
1114

1215
void operator()(id<1> i) const {}
1316

@@ -18,18 +21,48 @@ struct ESIMDKernel {
1821
}
1922
};
2023

24+
template <> struct sycl::detail::KernelInfo<KernelWithProps> {
25+
static constexpr kernel_param_desc_t desc{};
26+
static constexpr unsigned getNumParams() { return 0; }
27+
static constexpr const kernel_param_desc_t &getParamDesc(int Idx) {
28+
return desc;
29+
}
30+
static constexpr const char *getName() { return "KernelWithProps"; }
31+
static constexpr bool isESIMD() { return false; }
32+
static constexpr const char *getFileName() { return ""; }
33+
static constexpr const char *getFunctionName() { return ""; }
34+
static constexpr unsigned getLineNumber() { return 0; }
35+
static constexpr unsigned getColumnNumber() { return 0; }
36+
static constexpr int64_t getKernelSize() { return sizeof(KernelWithProps); }
37+
};
38+
39+
template <> struct sycl::detail::KernelInfo<class Kernel> {
40+
static constexpr kernel_param_desc_t desc{};
41+
static constexpr unsigned getNumParams() { return 0; }
42+
static constexpr const kernel_param_desc_t &getParamDesc(int Idx) {
43+
return desc;
44+
}
45+
static constexpr const char *getName() { return "Kernel"; }
46+
static constexpr bool isESIMD() { return false; }
47+
static constexpr const char *getFileName() { return ""; }
48+
static constexpr const char *getFunctionName() { return ""; }
49+
static constexpr unsigned getLineNumber() { return 0; }
50+
static constexpr unsigned getColumnNumber() { return 0; }
51+
static constexpr int64_t getKernelSize() { return 1; }
52+
};
53+
2154
int main(void) {
2255
queue q;
2356
// expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.+}}: Floating point control property is supported for ESIMD kernels only.}}
2457
syclex::properties properties7{
2558
intelex::fp_control<intelex::fp_mode::round_toward_zero |
2659
intelex::fp_mode::denorm_ftz>};
2760
q.submit([&](handler &cgh) {
28-
cgh.single_task<class TestKernel7>(properties7, [=]() {});
61+
cgh.single_task<class Kernel>(properties7, [=]() {});
2962
});
3063

3164
// expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.+}}: Floating point control property is supported for ESIMD kernels only.}}
32-
ESIMDKernel Kern;
65+
KernelWithProps Kern;
3366
q.submit([&](handler &cgh) { cgh.parallel_for(range<1>(1), Kern); });
3467

3568
return 0;

‎sycl/test/lit.cfg.py‎

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -126,9 +126,7 @@
126126
llvm_symbolizer = os.path.join(config.llvm_build_bin_dir, "llvm-symbolizer")
127127
llvm_config.with_environment("LLVM_SYMBOLIZER_PATH", llvm_symbolizer)
128128

129-
sycl_host_only_options = (
130-
"-std=c++17 -Xclang -fsycl-is-host -D__SYCL_UNITTESTS_BYPASS_KERNEL_NAME_CHECK=1"
131-
)
129+
sycl_host_only_options = "-std=c++17 -Xclang -fsycl-is-host"
132130
for include_dir in [
133131
config.sycl_include,
134132
config.level_zero_include_dir,

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

Lines changed: 23 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,25 @@ struct user_defined {
99
float b;
1010
};
1111

12+
struct Kernel {
13+
void operator()() const {}
14+
};
15+
16+
template <> struct sycl::detail::KernelInfo<Kernel> {
17+
static constexpr unsigned getNumParams() { return 0; }
18+
static constexpr const kernel_param_desc_t &getParamDesc(int Idx) {
19+
static constexpr kernel_param_desc_t desc{};
20+
return desc;
21+
}
22+
static constexpr const char *getName() { return "Kernel"; }
23+
static constexpr bool isESIMD() { return false; }
24+
static constexpr const char *getFileName() { return ""; }
25+
static constexpr const char *getFunctionName() { return ""; }
26+
static constexpr unsigned getLineNumber() { return 0; }
27+
static constexpr unsigned getColumnNumber() { return 0; }
28+
static constexpr int64_t getKernelSize() { return sizeof(Kernel); }
29+
};
30+
1231
int main() {
1332
sycl::queue q;
1433

@@ -18,15 +37,15 @@ int main() {
1837
oneapi::properties props_user{oneapi::indirectly_callable_in<user_defined>};
1938

2039
// expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}}
21-
q.single_task(props_empty, [=]() {});
40+
q.single_task(props_empty, Kernel{});
2241
// When both "props_empty" and "props_void" are in use, we won't see the
2342
// static assert firing for the second one, because there will be only one
2443
// instantiation of handler::processProperties.
25-
q.single_task(props_void, [=]() {});
44+
q.single_task(props_void, Kernel{});
2645
// expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}}
27-
q.single_task(props_int, [=]() {});
46+
q.single_task(props_int, Kernel{});
2847
// expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}}
29-
q.single_task(props_user, [=]() {});
48+
q.single_task(props_user, Kernel{});
3049

3150
return 0;
3251
}

‎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: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -9,8 +9,6 @@ endforeach()
99

1010
add_compile_definitions(SYCL2020_DISABLE_DEPRECATION_WARNINGS SYCL_DISABLE_FSYCL_SYCLHPP_WARNING)
1111

12-
add_compile_definitions(__SYCL_UNITTESTS_BYPASS_KERNEL_NAME_CHECK)
13-
1412
# suppress warnings which came from Google Test sources
1513
if (CXX_SUPPORTS_SUGGEST_OVERRIDE_FLAG)
1614
add_compile_options("-Wno-suggest-override")

0 commit comments

Comments
 (0)