diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_launch_bounds.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_launch_bounds.cpp index 96ac3da42a504..75fbf5a18ab73 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_launch_bounds.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_launch_bounds.cpp @@ -4,14 +4,21 @@ #include +constexpr auto Props = sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::max_linear_work_group_size<4>, +}; +struct TestKernelLaunchBounds { + void operator()() const {} + auto get(sycl::ext::oneapi::experimental::properties_tag) { return Props; } +}; + int main() { sycl::queue Q; - constexpr auto Props = sycl::ext::oneapi::experimental::properties{ - sycl::ext::oneapi::experimental::max_linear_work_group_size<4>, - }; // CHECK-IR: spir_kernel void @{{.*}}LaunchBoundsKernel(){{.*}} #[[LaunchBoundsAttrs:[0-9]+]] - Q.single_task(Props, []() {}); + Q.submit([&](sycl::handler &h) { + h.single_task(TestKernelLaunchBounds{}); + }); return 0; } diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_max_work_group_size.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_max_work_group_size.cpp index 924270bb6cafe..e81bd5d20c452 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_max_work_group_size.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_max_work_group_size.cpp @@ -4,26 +4,47 @@ #include +constexpr auto Props1 = sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::max_work_group_size<8>}; +constexpr auto Props2 = sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::max_work_group_size<8, 4>}; +constexpr auto Props3 = sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::max_work_group_size<8, 4, 2>}; + +struct TestKernel_Props1 { + void operator()() const {} + auto get(sycl::ext::oneapi::experimental::properties_tag) { return Props1; } +}; + +struct TestKernel_Props2 { + void operator()() const {} + auto get(sycl::ext::oneapi::experimental::properties_tag) { return Props2; } +}; + +struct TestKernel_Props3 { + void operator()() const {} + auto get(sycl::ext::oneapi::experimental::properties_tag) { return Props3; } +}; + int main() { sycl::queue Q; sycl::event Ev; - constexpr auto Props1 = sycl::ext::oneapi::experimental::properties{ - sycl::ext::oneapi::experimental::max_work_group_size<8>}; - constexpr auto Props2 = sycl::ext::oneapi::experimental::properties{ - sycl::ext::oneapi::experimental::max_work_group_size<8, 4>}; - constexpr auto Props3 = sycl::ext::oneapi::experimental::properties{ - sycl::ext::oneapi::experimental::max_work_group_size<8, 4, 2>}; - // CHECK-IR: spir_kernel void @{{.*}}MaxWGSizeKernel0(){{.*}} #[[MaxWGSizeAttr0:[0-9]+]] // CHECK-IR-SAME: !max_work_group_size ![[MaxWGSizeMD0:[0-9]+]] - Q.single_task(Props1, []() {}); + Q.single_task(TestKernel_Props1{}); // CHECK-IR: spir_kernel void @{{.*}}MaxWGSizeKernel1(){{.*}} #[[MaxWGSizeAttr1:[0-9]+]] // CHECK-IR-SAME: !max_work_group_size ![[MaxWGSizeMD1:[0-9]+]] - Q.single_task(Ev, Props2, []() {}); + Q.submit([&](sycl::handler &h) { + h.depends_on(Ev); + h.single_task(TestKernel_Props2{}); + }); // CHECK-IR: spir_kernel void @{{.*}}MaxWGSizeKernel2(){{.*}} #[[MaxWGSizeAttr2:[0-9]+]] // CHECK-IR-SAME: !max_work_group_size ![[MaxWGSizeMD2:[0-9]+]] - Q.single_task({Ev}, Props3, []() {}); + Q.submit([&](sycl::handler &h) { + h.depends_on({Ev}); + h.single_task(TestKernel_Props3{}); + }); return 0; } diff --git a/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp b/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp index f24e089bed4b1..55ac13958bce9 100644 --- a/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp +++ b/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp @@ -54,128 +54,245 @@ int funcIndirectlyUsingCPU(int a, int b) { return funcUsingCPU(a, b, 1); } SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((device_has)) int funcUsingCPUHasFP64(int a) { return funcIndirectlyUsingCPU(a, 1); } +constexpr auto props_64 = properties{device_has}; +constexpr auto props_16 = properties{device_has}; +constexpr auto props_1664 = properties{device_has}; +constexpr auto props_6416 = properties{device_has}; +constexpr auto props_gpu = properties{device_has}; +constexpr auto props_gpu1664 = + properties{device_has}; +constexpr auto props_1664gpu = + properties{device_has}; +constexpr auto props_emp = properties{}; +constexpr auto props_cpu = properties{device_has}; +constexpr auto props_cpu64 = properties{device_has}; +constexpr auto props_64cpu = properties{device_has}; +constexpr auto props_gpucpu64 = + properties{device_has}; +constexpr auto props_cpu64gpu = + properties{device_has}; + +template struct K_funcIndirectlyUsingFP16 { + T *Props; + K_funcIndirectlyUsingFP16(T Props_param) { Props = &Props_param; }; + void operator()() const { int a = funcIndirectlyUsingFP16(1, 2); } + auto get(properties_tag) { return *Props; } +}; + +template struct K_funcIndirectlyUsingFP16_Warn16 { + T *Props; + K_funcIndirectlyUsingFP16_Warn16(T Props_param) { Props = &Props_param; }; + // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'fp16' not listed in its 'device_has' property}} + void operator()() const { int a = funcIndirectlyUsingFP16(1, 2); } + auto get(properties_tag) { return *Props; } +}; + +template struct K_funcUsingFP16AndFP64 { + T *Props; + K_funcUsingFP16AndFP64(T Props_param) { Props = &Props_param; }; + void operator()() const { int a = funcUsingFP16AndFP64(1, 2); } + auto get(properties_tag) { return *Props; } +}; + +template struct K_funcUsingFP16AndFP64_Warn16 { + T *Props; + K_funcUsingFP16AndFP64_Warn16(T Props_param) { Props = &Props_param; }; + // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'fp16' not listed in its 'device_has' property}} + void operator()() const { int a = funcUsingFP16AndFP64(1, 2); } + auto get(properties_tag) { return *Props; } +}; + +template struct K_funcUsingFP16AndFP64_Warn64 { + T *Props; + K_funcUsingFP16AndFP64_Warn64(T Props_param) { Props = &Props_param; }; + // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'device_has' property}} + void operator()() const { int a = funcUsingFP16AndFP64(1, 2); } + auto get(properties_tag) { return *Props; } +}; + +template struct K_funcUsingFP16AndFP64_Warn1664 { + T *Props; + K_funcUsingFP16AndFP64_Warn1664(T Props_param) { Props = &Props_param; }; + // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'fp16' not listed in its 'device_has' property}} + // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'device_has' property}} + void operator()() const { int a = funcUsingFP16AndFP64(1, 2); } + auto get(properties_tag) { return *Props; } +}; + +template struct K_funcUsingFP16AndFP64_False { + T *Props; + K_funcUsingFP16AndFP64_False(T Props_param) { Props = &Props_param; }; + void operator()() const { + if constexpr (false) { + int a = funcUsingFP16AndFP64(1, 2); + } + } + auto get(properties_tag) { return *Props; } +}; + +template struct K_funcUsingCPUHasFP64 { + T *Props; + K_funcUsingCPUHasFP64(T Props_param) { Props = &Props_param; }; + void operator()() const { int a = funcUsingCPUHasFP64(1); } + auto get(properties_tag) { return *Props; } +}; + +template struct K_funcIndirectlyUsingCPU { + T *Props; + K_funcIndirectlyUsingCPU(T Props_param) { Props = &Props_param; }; + void operator()() const { int a = funcIndirectlyUsingCPU(1, 2); } + auto get(properties_tag) { return *Props; } +}; + +template struct K_funcIndirectlyUsingCPU_WarnCPU { + T *Props; + K_funcIndirectlyUsingCPU_WarnCPU(T Props_param) { Props = &Props_param; }; + // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'cpu' not listed in its 'device_has' property}} + void operator()() const { int a = funcIndirectlyUsingCPU(1, 2); } + auto get(properties_tag) { return *Props; } +}; + +template struct K_funcUsingCPUAndFP64 { + T *Props; + K_funcUsingCPUAndFP64(T Props_param) { Props = &Props_param; }; + void operator()() const { int a = funcUsingCPUAndFP64(1, 2); } + auto get(properties_tag) { return *Props; } +}; + +template struct K_funcUsingCPUAndFP64_WarnCPU { + T *Props; + K_funcUsingCPUAndFP64_WarnCPU(T Props_param) { Props = &Props_param; }; + // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'cpu' not listed in its 'device_has' property}} + void operator()() const { int a = funcUsingCPUAndFP64(1, 2); } + auto get(properties_tag) { return *Props; } +}; + +template struct K_funcUsingCPUAndFP64_Warn64 { + T *Props; + K_funcUsingCPUAndFP64_Warn64(T Props_param) { Props = &Props_param; }; + // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'device_has' property}} + void operator()() const { int a = funcUsingCPUAndFP64(1, 2); } + auto get(properties_tag) { return *Props; } +}; + +template struct K_funcUsingCPUAndFP64_Warn64CPU { + T *Props; + K_funcUsingCPUAndFP64_Warn64CPU(T Props_param) { Props = &Props_param; }; + // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'cpu' not listed in its 'device_has' property}} + // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'device_has' property}} + void operator()() const { int a = funcUsingCPUAndFP64(1, 2); } + auto get(properties_tag) { return *Props; } +}; + +template struct K_funcUsingCPUAndFP64_False { + T *Props; + K_funcUsingCPUAndFP64_False(T Props_param) { Props = &Props_param; }; + void operator()() const { + if constexpr (false) { + int a = funcUsingCPUAndFP64(1, 2); + } + } + auto get(properties_tag) { return *Props; } +}; + int main() { queue Q; Q.submit([&](handler &CGH) { CGH.single_task([=]() { int a = funcUsingFP16HasFP64(1); }); }); Q.submit([&](handler &CGH) { - // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'fp16' not listed in its 'device_has' property}} - CGH.single_task(properties{device_has}, - [=]() { int a = funcIndirectlyUsingFP16(1, 2); }); + CGH.single_task( + K_funcIndirectlyUsingFP16_Warn16(props_64)); }); Q.submit([&](handler &CGH) { - CGH.single_task(properties{device_has}, - [=]() { int a = funcIndirectlyUsingFP16(1, 2); }); + CGH.single_task(K_funcIndirectlyUsingFP16(props_16)); }); Q.submit([&](handler &CGH) { - CGH.single_task(properties{device_has}, - [=]() { int a = funcIndirectlyUsingFP16(1, 2); }); + CGH.single_task( + K_funcIndirectlyUsingFP16(props_1664)); }); Q.submit([&](handler &CGH) { - CGH.single_task(properties{device_has}, - [=]() { int a = funcIndirectlyUsingFP16(1, 2); }); + CGH.single_task( + K_funcIndirectlyUsingFP16(props_6416)); }); Q.submit([&](handler &CGH) { - // expected-warning-re@+3 {{function '{{.*}}' uses aspect 'fp16' not listed in its 'device_has' property}} - // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'device_has' property}} - CGH.single_task(properties{device_has}, - [=]() { int a = funcUsingFP16AndFP64(1, 2); }); + CGH.single_task( + K_funcUsingFP16AndFP64_Warn1664(props_gpu)); }); Q.submit([&](handler &CGH) { - // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'device_has' property}} - CGH.single_task(properties{device_has}, - [=]() { int a = funcUsingFP16AndFP64(1, 2); }); + CGH.single_task( + K_funcUsingFP16AndFP64_Warn64(props_16)); }); Q.submit([&](handler &CGH) { - // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'fp16' not listed in its 'device_has' property}} - CGH.single_task(properties{device_has}, - [=]() { int a = funcUsingFP16AndFP64(1, 2); }); + CGH.single_task( + K_funcUsingFP16AndFP64_Warn16(props_64)); }); Q.submit([&](handler &CGH) { - CGH.single_task(properties{device_has}, [=]() { - if constexpr (false) { - int a = funcUsingFP16AndFP64(1, 2); - } - }); + CGH.single_task( + K_funcUsingFP16AndFP64_False(props_gpu)); }); Q.submit([&](handler &CGH) { - CGH.single_task(properties{device_has}, - [=]() { int a = funcUsingFP16AndFP64(1, 2); }); + CGH.single_task( + K_funcUsingFP16AndFP64_Warn16(props_1664)); }); Q.submit([&](handler &CGH) { - CGH.single_task(properties{device_has}, - [=]() { int a = funcUsingFP16AndFP64(1, 2); }); + CGH.single_task( + K_funcUsingFP16AndFP64_Warn16(props_6416)); }); Q.submit([&](handler &CGH) { CGH.single_task( - properties{device_has}, - [=]() { int a = funcUsingFP16AndFP64(1, 2); }); + K_funcUsingFP16AndFP64_Warn16(props_gpu1664)); }); Q.submit([&](handler &CGH) { CGH.single_task( - properties{device_has}, - [=]() { int a = funcUsingFP16AndFP64(1, 2); }); + K_funcUsingFP16AndFP64_Warn16(props_1664gpu)); }); Q.submit([&](handler &CGH) { - CGH.single_task(properties{}, [=]() { int a = funcUsingCPUHasFP64(1); }); + CGH.single_task(K_funcUsingCPUHasFP64(props_emp)); }); Q.submit([&](handler &CGH) { - // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'cpu' not listed in its 'device_has' property}} - CGH.single_task(properties{device_has}, - [=]() { int a = funcIndirectlyUsingCPU(1, 2); }); + CGH.single_task( + K_funcIndirectlyUsingCPU_WarnCPU(props_64)); }); Q.submit([&](handler &CGH) { - CGH.single_task(properties{device_has}, - [=]() { int a = funcIndirectlyUsingCPU(1, 2); }); + CGH.single_task(K_funcIndirectlyUsingCPU(props_cpu)); }); Q.submit([&](handler &CGH) { - CGH.single_task(properties{device_has}, - [=]() { int a = funcIndirectlyUsingCPU(1, 2); }); + CGH.single_task( + K_funcIndirectlyUsingCPU(props_cpu64)); }); Q.submit([&](handler &CGH) { - CGH.single_task(properties{device_has}, - [=]() { int a = funcIndirectlyUsingCPU(1, 2); }); + CGH.single_task( + K_funcIndirectlyUsingCPU(props_64cpu)); }); Q.submit([&](handler &CGH) { - // expected-warning-re@+3 {{function '{{.*}}' uses aspect 'cpu' not listed in its 'device_has' property}} - // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'device_has' property}} - CGH.single_task(properties{device_has}, - [=]() { int a = funcUsingCPUAndFP64(1, 2); }); + CGH.single_task( + K_funcUsingCPUAndFP64_Warn64CPU(props_gpu)); }); Q.submit([&](handler &CGH) { - // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'device_has' property}} - CGH.single_task(properties{device_has}, - [=]() { int a = funcUsingCPUAndFP64(1, 2); }); + CGH.single_task( + K_funcUsingCPUAndFP64_Warn64(props_cpu)); }); Q.submit([&](handler &CGH) { - // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'cpu' not listed in its 'device_has' property}} - CGH.single_task(properties{device_has}, - [=]() { int a = funcUsingCPUAndFP64(1, 2); }); + CGH.single_task( + K_funcUsingCPUAndFP64_WarnCPU(props_64)); }); Q.submit([&](handler &CGH) { - CGH.single_task(properties{device_has}, [=]() { - if constexpr (false) { - int a = funcUsingCPUAndFP64(1, 2); - } - }); + CGH.single_task( + K_funcUsingCPUAndFP64_False(props_gpu)); }); Q.submit([&](handler &CGH) { - CGH.single_task(properties{device_has}, - [=]() { int a = funcUsingCPUAndFP64(1, 2); }); + CGH.single_task(K_funcUsingCPUAndFP64(props_cpu64)); }); Q.submit([&](handler &CGH) { - CGH.single_task(properties{device_has}, - [=]() { int a = funcUsingCPUAndFP64(1, 2); }); + CGH.single_task(K_funcUsingCPUAndFP64(props_64cpu)); }); Q.submit([&](handler &CGH) { CGH.single_task( - properties{device_has}, - [=]() { int a = funcUsingCPUAndFP64(1, 2); }); + K_funcUsingCPUAndFP64(props_gpucpu64)); }); Q.submit([&](handler &CGH) { CGH.single_task( - properties{device_has}, - [=]() { int a = funcUsingCPUAndFP64(1, 2); }); + K_funcUsingCPUAndFP64(props_cpu64gpu)); }); } diff --git a/sycl/test/extensions/properties/properties_kernel_negative_device.cpp b/sycl/test/extensions/properties/properties_kernel_negative_device.cpp index d451e319a2670..230626e97e248 100644 --- a/sycl/test/extensions/properties/properties_kernel_negative_device.cpp +++ b/sycl/test/extensions/properties/properties_kernel_negative_device.cpp @@ -2,11 +2,6 @@ #include -struct KernelFunctorWithOnlyWGSizeAttr { - // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} - void operator() [[sycl::reqd_work_group_size(32)]] () const {} -}; - template struct KernelFunctorWithWGSizeWithAttr { // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} void operator() [[sycl::reqd_work_group_size(32)]] () const {} @@ -16,11 +11,6 @@ template struct KernelFunctorWithWGSizeWithAttr { } }; -struct KernelFunctorWithOnlyWGSizeHintAttr { - // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} - void operator() [[sycl::work_group_size_hint(32)]] () const {} -}; - template struct KernelFunctorWithWGSizeHintWithAttr { // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} void operator() [[sycl::work_group_size_hint(32)]] () const {} @@ -30,11 +20,6 @@ template struct KernelFunctorWithWGSizeHintWithAttr { } }; -struct KernelFunctorWithOnlySGSizeAttr { - // expected-warning@+1 {{kernel has both attribute 'reqd_sub_group_size' and kernel properties; conflicting properties are ignored}} - void operator() [[sycl::reqd_sub_group_size(32)]] () const {} -}; - template struct KernelFunctorWithSGSizeWithAttr { // expected-warning@+1 {{kernel has both attribute 'reqd_sub_group_size' and kernel properties; conflicting properties are ignored}} void operator() [[sycl::reqd_sub_group_size(32)]] () const {} @@ -44,11 +29,6 @@ template struct KernelFunctorWithSGSizeWithAttr { } }; -struct KernelFunctorWithOnlyDeviceHasAttr { - // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} - void operator() [[sycl::device_has(sycl::aspect::cpu)]] () const {} -}; - template struct KernelFunctorWithDeviceHasWithAttr { // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} void operator() [[sycl::device_has(sycl::aspect::cpu)]] () const {} @@ -61,68 +41,24 @@ template struct KernelFunctorWithDeviceHasWithAttr { void check_work_group_size() { sycl::queue Q; - // expected-warning@+4 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} - Q.single_task( - sycl::ext::oneapi::experimental::properties{ - sycl::ext::oneapi::experimental::work_group_size<1>}, - []() [[sycl::reqd_work_group_size(32)]] {}); - - Q.single_task( - sycl::ext::oneapi::experimental::properties{ - sycl::ext::oneapi::experimental::work_group_size<1>}, - KernelFunctorWithOnlyWGSizeAttr{}); - Q.single_task(KernelFunctorWithWGSizeWithAttr<1>{}); } void check_work_group_size_hint() { sycl::queue Q; - // expected-warning@+4 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} - Q.single_task( - sycl::ext::oneapi::experimental::properties{ - sycl::ext::oneapi::experimental::work_group_size_hint<1>}, - []() [[sycl::work_group_size_hint(32)]] {}); - - Q.single_task( - sycl::ext::oneapi::experimental::properties{ - sycl::ext::oneapi::experimental::work_group_size_hint<1>}, - KernelFunctorWithOnlyWGSizeHintAttr{}); - Q.single_task(KernelFunctorWithWGSizeHintWithAttr<1>{}); } void check_sub_group_size() { sycl::queue Q; - // expected-warning@+4 {{kernel has both attribute 'reqd_sub_group_size' and kernel properties; conflicting properties are ignored}} - Q.single_task( - sycl::ext::oneapi::experimental::properties{ - sycl::ext::oneapi::experimental::sub_group_size<1>}, - []() [[sycl::reqd_sub_group_size(32)]] {}); - - Q.single_task( - sycl::ext::oneapi::experimental::properties{ - sycl::ext::oneapi::experimental::sub_group_size<1>}, - KernelFunctorWithOnlySGSizeAttr{}); - Q.single_task(KernelFunctorWithSGSizeWithAttr<1>{}); } void check_device_has() { sycl::queue Q; - // expected-warning@+4 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} - Q.single_task( - sycl::ext::oneapi::experimental::properties{ - sycl::ext::oneapi::experimental::device_has}, - []() [[sycl::device_has(sycl::aspect::cpu)]] {}); - - Q.single_task( - sycl::ext::oneapi::experimental::properties{ - sycl::ext::oneapi::experimental::device_has}, - KernelFunctorWithOnlyDeviceHasAttr{}); - Q.single_task(KernelFunctorWithDeviceHasWithAttr{}); } diff --git a/sycl/test/virtual-functions/diagnostics-positive.cpp b/sycl/test/virtual-functions/diagnostics-positive.cpp index 1b1e346678e8f..812161457784c 100644 --- a/sycl/test/virtual-functions/diagnostics-positive.cpp +++ b/sycl/test/virtual-functions/diagnostics-positive.cpp @@ -49,46 +49,70 @@ void bar(SubDerived *Ptr) { Ptr->bar(); } -int main() { - sycl::queue q; - - // The exact arguments passed to calls_indirectly property don't matter, - // because we have no way of connecting a virtual call with a particular - // property set, but we test different properties here just in case. - oneapi::properties props_empty{oneapi::assume_indirect_calls}; - oneapi::properties props_void{oneapi::assume_indirect_calls_to}; - oneapi::properties props_int{oneapi::assume_indirect_calls_to}; - oneapi::properties props_base{oneapi::assume_indirect_calls_to}; - - char *Storage = sycl::malloc_device(128, q); - - q.single_task(props_empty, [=]() { +// The exact arguments passed to calls_indirectly property don't matter, +// because we have no way of connecting a virtual call with a particular +// property set, but we test different properties here just in case. +oneapi::properties props_empty{oneapi::assume_indirect_calls}; +oneapi::properties props_void{oneapi::assume_indirect_calls_to}; +oneapi::properties props_int{oneapi::assume_indirect_calls_to}; +oneapi::properties props_base{oneapi::assume_indirect_calls_to}; + +struct TestKernel_props_empty { + char *Storage; + TestKernel_props_empty(char *Storage_param) { Storage = Storage_param; }; + void operator()() const { new (Storage) SubSubDerived; auto *Ptr = reinterpret_cast(Storage); Ptr->foo(); - }); + } + auto get(oneapi::properties_tag) { return props_empty; } +}; - q.single_task(props_void, [=]() { +struct TestKernel_props_void { + char *Storage; + TestKernel_props_void(char *Storage_param) { Storage = Storage_param; }; + void operator()() const { new (Storage) SubDerived; auto *Ptr = reinterpret_cast(Storage); Ptr->bar(); - }); + } + auto get(oneapi::properties_tag) { return props_void; } +}; - q.single_task(props_int, [=]() { +struct TestKernel_props_int { + char *Storage; + TestKernel_props_int(char *Storage_param) { Storage = Storage_param; }; + void operator()() const { new (Storage) Derived; auto *Ptr = reinterpret_cast(Storage); foo(Ptr); - }); + } + auto get(oneapi::properties_tag) { return props_int; } +}; - q.single_task(props_base, [=]() { +struct TestKernel_props_base { + char *Storage; + TestKernel_props_base(char *Storage_param) { Storage = Storage_param; }; + void operator()() const { auto *Ptr = reinterpret_cast(Storage); bar(Ptr); - }); + } + auto get(oneapi::properties_tag) { return props_base; } +}; + +int main() { + sycl::queue q; + + char *Storage = sycl::malloc_device(128, q); + + q.single_task(TestKernel_props_empty(Storage)); + q.single_task(TestKernel_props_void(Storage)); + q.single_task(TestKernel_props_int(Storage)); + q.single_task(TestKernel_props_base(Storage)); sycl::free(Storage, q); return 0; } - diff --git a/sycl/test/virtual-functions/properties-positive.cpp b/sycl/test/virtual-functions/properties-positive.cpp index 3441e91a4d45c..507a284e4de21 100644 --- a/sycl/test/virtual-functions/properties-positive.cpp +++ b/sycl/test/virtual-functions/properties-positive.cpp @@ -41,6 +41,37 @@ class SubSubDerived : public SubDerived { void bar() override {} }; +oneapi::properties props_empty{oneapi::assume_indirect_calls}; +oneapi::properties props_void{oneapi::assume_indirect_calls_to}; +oneapi::properties props_int{oneapi::assume_indirect_calls_to}; +oneapi::properties props_base{oneapi::assume_indirect_calls_to}; +oneapi::properties props_multiple{oneapi::assume_indirect_calls_to}; + +struct TestKernel_props_empty { + void operator()() const {} + auto get(oneapi::properties_tag) { return props_empty; } +}; + +struct TestKernel_props_void { + void operator()() const {} + auto get(oneapi::properties_tag) { return props_void; } +}; + +struct TestKernel_props_int { + void operator()() const {} + auto get(oneapi::properties_tag) { return props_int; } +}; + +struct TestKernel_props_base { + void operator()() const {} + auto get(oneapi::properties_tag) { return props_base; } +}; + +struct TestKernel_props_multiple { + void operator()() const {} + auto get(oneapi::properties_tag) { return props_multiple; } +}; + int main() { sycl::queue q; @@ -51,11 +82,11 @@ int main() { oneapi::properties props_multiple{ oneapi::assume_indirect_calls_to}; - q.single_task(props_empty, [=]() {}); - q.single_task(props_void, [=]() {}); - q.single_task(props_int, [=]() {}); - q.single_task(props_base, [=]() {}); - q.single_task(props_multiple, [=]() {}); + q.single_task(TestKernel_props_empty{}); + q.single_task(TestKernel_props_void{}); + q.single_task(TestKernel_props_int{}); + q.single_task(TestKernel_props_base{}); + q.single_task(TestKernel_props_multiple{}); return 0; }