diff --git a/sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp b/sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp new file mode 100644 index 0000000000000..c55847b9e2735 --- /dev/null +++ b/sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp @@ -0,0 +1,113 @@ +// UNSUPPORTED: cuda, hip, acc +// FIXME: replace unsupported with an aspect check once we have it +// +// RUN: %{build} -o %t.out -Xclang -fsycl-allow-virtual-functions %helper-includes +// RUN: %{run} %t.out + +#include + +#include "helpers.hpp" + +#include + +namespace oneapi = sycl::ext::oneapi::experimental; + +class Base { +public: + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + virtual void increment(int *) { /* do nothhing */ + } + + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + virtual void multiply(int *) { /* do nothhing */ + } + + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + virtual void substract(int *) { /* do nothhing */ + } +}; + +class IncrementBy1 : public Base { +public: + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + void increment(int *Data) override { *Data += 1; } +}; + +class IncrementBy1AndSubstractBy2 : public IncrementBy1 { +public: + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + void substract(int *Data) override { *Data -= 2; } +}; + +class MultiplyBy2 : public Base { +public: + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + void multiply(int *Data) override { *Data *= 2; } +}; + +class MultiplyBy2AndIncrementBy8 : public MultiplyBy2 { +public: + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + void increment(int *Data) override { *Data += 8; } +}; + +class SubstractBy4 : public Base { +public: + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + void substract(int *Data) override { *Data -= 4; } +}; + +class SubstractBy4AndMultiplyBy4 : public SubstractBy4 { +public: + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + void multiply(int *Data) override { *Data *= 4; } +}; + +void applyOp(int *DataPtr, Base *ObjPtr) { + ObjPtr->increment(DataPtr); + ObjPtr->substract(DataPtr); + ObjPtr->multiply(DataPtr); +} + +int main() try { + using storage_t = obj_storage_t; + storage_t HostStorage; + sycl::buffer DeviceStorage(sycl::range{1}); + + auto asyncHandler = [](sycl::exception_list list) { + for (auto &e : list) + std::rethrow_exception(e); + }; + + sycl::queue q(asyncHandler); + + constexpr oneapi::properties props{oneapi::calls_indirectly<>}; + for (unsigned TestCase = 0; TestCase < 6; ++TestCase) { + int HostData = 42; + int Data = HostData; + sycl::buffer DataStorage(&Data, sycl::range{1}); + + q.submit([&](sycl::handler &CGH) { + sycl::accessor StorageAcc(DeviceStorage, CGH, sycl::write_only); + sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only); + CGH.single_task(props, [=]() { + auto *Ptr = StorageAcc[0].construct(TestCase); + applyOp(DataAcc.get_multi_ptr().get(), + Ptr); + }); + }); + + Base *Ptr = HostStorage.construct(TestCase); + applyOp(&HostData, Ptr); + + sycl::host_accessor HostAcc(DataStorage); + assert(HostAcc[0] == HostData); + } + + return 0; +} catch (sycl::exception &e) { + std::cout << "Unexpected exception was thrown: " << e.what() << std::endl; + return 1; +} diff --git a/sycl/test-e2e/VirtualFunctions/2/1/1/more-complex-hierarchy.cpp b/sycl/test-e2e/VirtualFunctions/2/1/1/more-complex-hierarchy.cpp new file mode 100644 index 0000000000000..fce036b890294 --- /dev/null +++ b/sycl/test-e2e/VirtualFunctions/2/1/1/more-complex-hierarchy.cpp @@ -0,0 +1,93 @@ +// UNSUPPORTED: cuda, hip, acc +// FIXME: replace unsupported with an aspect check once we have it +// +// RUN: %{build} -o %t.out -Xclang -fsycl-allow-virtual-functions %helper-includes +// RUN: %{run} %t.out + +#include + +#include "helpers.hpp" + +#include + +namespace oneapi = sycl::ext::oneapi::experimental; + +class AbstractOp { +public: + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + virtual void applyOp(int *) = 0; +}; + +class IncrementOp : public AbstractOp { + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + void applyOp(int *Data) final override { increment(Data); } + + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + virtual void increment(int *) = 0; +}; + +class IncrementBy1 : public IncrementOp { + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + void increment(int *Data) override { *Data += 1; } +}; + +class IncrementBy2 : public IncrementOp { + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + void increment(int *Data) override { *Data += 2; } +}; + +class IncrementBy4 : public IncrementOp { + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + void increment(int *Data) override { *Data += 4; } +}; + +class IncrementBy8 : public IncrementOp { + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + void increment(int *Data) override { *Data += 8; } +}; + +void applyOp(int *Data, AbstractOp *Obj) { Obj->applyOp(Data); } + +int main() try { + using storage_t = + obj_storage_t; + + storage_t HostStorage; + sycl::buffer DeviceStorage(sycl::range{1}); + + auto asyncHandler = [](sycl::exception_list list) { + for (auto &e : list) + std::rethrow_exception(e); + }; + + sycl::queue q(asyncHandler); + + constexpr oneapi::properties props{oneapi::calls_indirectly<>}; + for (unsigned TestCase = 0; TestCase < 4; ++TestCase) { + int HostData = 42; + int Data = HostData; + sycl::buffer DataStorage(&Data, sycl::range{1}); + + q.submit([&](sycl::handler &CGH) { + sycl::accessor StorageAcc(DeviceStorage, CGH, sycl::write_only); + sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only); + CGH.single_task(props, [=]() { + auto *Ptr = + StorageAcc[0].construct(TestCase); + applyOp(DataAcc.get_multi_ptr().get(), + Ptr); + }); + }); + + auto *Ptr = HostStorage.construct(TestCase); + Ptr->applyOp(&HostData); + + sycl::host_accessor HostAcc(DataStorage); + assert(HostAcc[0] == HostData); + } + + return 0; +} catch (sycl::exception &e) { + std::cout << "Unexpected exception was thrown: " << e.what() << std::endl; + return 1; +} diff --git a/sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp b/sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp new file mode 100644 index 0000000000000..83ec49ee3482d --- /dev/null +++ b/sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp @@ -0,0 +1,78 @@ +// UNSUPPORTED: cuda, hip, acc +// FIXME: replace unsupported with an aspect check once we have it +// +// RUN: %{build} -o %t.out -Xclang -fsycl-allow-virtual-functions %helper-includes +// RUN: %{run} %t.out + +#include + +#include "helpers.hpp" + +#include + +namespace oneapi = sycl::ext::oneapi::experimental; + +class BaseIncrement { +public: + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + virtual void increment(int *Data) { *Data += 1; } +}; + +class IncrementBy2 : public BaseIncrement { + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + void increment(int *Data) override { *Data += 2; } +}; + +class IncrementBy4 : public BaseIncrement { + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + void increment(int *Data) override { *Data += 4; } +}; + +class IncrementBy8 : public BaseIncrement { + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>) + void increment(int *Data) override { *Data += 8; } +}; + +int main() try { + using storage_t = + obj_storage_t; + + storage_t HostStorage; + sycl::buffer DeviceStorage(sycl::range{1}); + + auto asyncHandler = [](sycl::exception_list list) { + for (auto &e : list) + std::rethrow_exception(e); + }; + + sycl::queue q(asyncHandler); + + constexpr oneapi::properties props{oneapi::calls_indirectly<>}; + for (unsigned TestCase = 0; TestCase < 4; ++TestCase) { + int HostData = 42; + int Data = HostData; + sycl::buffer DataStorage(&Data, sycl::range{1}); + + q.submit([&](sycl::handler &CGH) { + sycl::accessor StorageAcc(DeviceStorage, CGH, sycl::write_only); + sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only); + CGH.single_task(props, [=]() { + auto *Ptr = + StorageAcc[0].construct(TestCase); + Ptr->increment( + DataAcc.get_multi_ptr().get()); + }); + }); + + auto *Ptr = HostStorage.construct(TestCase); + Ptr->increment(&HostData); + + sycl::host_accessor HostAcc(DataStorage); + assert(HostAcc[0] == HostData); + } + + return 0; +} catch (sycl::exception &e) { + std::cout << "Unexpected exception was thrown: " << e.what() << std::endl; + return 1; +} diff --git a/sycl/test-e2e/VirtualFunctions/README.md b/sycl/test-e2e/VirtualFunctions/README.md new file mode 100644 index 0000000000000..8eef65107eec7 --- /dev/null +++ b/sycl/test-e2e/VirtualFunctions/README.md @@ -0,0 +1,7 @@ +# E2E tests for `sycl_ext_oneapi_virtual_functions` extension + +Note about naming convention and files organization for this folder: the tests, +files and directories are named and organized in a way that resembles their +description in the corresponding test plan document: link to be inserted here +later, but for now look into +[intel/llvm#10540](https://github.com/intel/llvm/pull/10540) PR. diff --git a/sycl/test-e2e/VirtualFunctions/helpers.hpp b/sycl/test-e2e/VirtualFunctions/helpers.hpp new file mode 100644 index 0000000000000..7c5748182e6ea --- /dev/null +++ b/sycl/test-e2e/VirtualFunctions/helpers.hpp @@ -0,0 +1,53 @@ +#include +#include + +// TODO: strictly speaking, selecting a max alignment here may not be always +// valid, but for test cases that we have now we expect alignment of all types +// to be the same. +// std::aligned_storage uses double under the hood which prevents us from +// using it on some HW. Therefore we use a custom implementation. +template struct aligned_storage { + static constexpr size_t Len = std::max({sizeof(T)...}); + static constexpr size_t Align = std::max({alignof(T)...}); + + struct type { + alignas(Align) unsigned char data[Len]; + }; +}; + +// Helper data structure that automatically creates a right (in terms of size +// and alignment) storage to accomodate a value of any of types T... +template struct obj_storage_t { + static_assert(std::max({alignof(T)...}) == std::min({alignof(T)...}), + "Unsupported alignment of input types"); + using type = typename aligned_storage::type; + static constexpr size_t size = std::max({sizeof(T)...}); + + type storage; + + template RetT *construct(const unsigned int TypeIndex) { + if (TypeIndex >= sizeof...(T)) { +#ifndef __SYCL_DEVICE_ONLY__ + assert(false && "Type index is invalid"); +#endif + return nullptr; + } + + return constructHelper(TypeIndex, 0); + } + +private: + template RetT *constructHelper(const int, const int) { + // Won't be ever called, but required to compile + return nullptr; + } + + template + RetT *constructHelper(const int TargetIndex, const int CurIndex) { + if (TargetIndex != CurIndex) + return constructHelper(TargetIndex, CurIndex + 1); + + RetT *Ptr = new (reinterpret_cast(&storage)) Type; + return Ptr; + } +}; diff --git a/sycl/test-e2e/VirtualFunctions/lit.local.cfg b/sycl/test-e2e/VirtualFunctions/lit.local.cfg new file mode 100644 index 0000000000000..f74079fb0725a --- /dev/null +++ b/sycl/test-e2e/VirtualFunctions/lit.local.cfg @@ -0,0 +1,6 @@ +import os + +# Tests are sharing some common header, but we don't won't to use relative +# paths like "../../../helper.hpp" in them, so let's just register a +# substitution to add directory with helper headers into include search path +config.substitutions.append(("%helper-includes", "-I {}".format(os.path.dirname(os.path.abspath(__file__)))))