summaryrefslogtreecommitdiff
path: root/modules/sycl
diff options
context:
space:
mode:
Diffstat (limited to 'modules/sycl')
-rw-r--r--modules/sycl/.nix/derivation.nix43
-rw-r--r--modules/sycl/SConstruct80
-rw-r--r--modules/sycl/c++/SConscript27
-rw-r--r--modules/sycl/c++/common.hpp12
-rw-r--r--modules/sycl/c++/data.hpp961
-rw-r--r--modules/sycl/c++/index.hpp18
-rw-r--r--modules/sycl/c++/lbm.hpp3
-rw-r--r--modules/sycl/c++/sycl.hpp5
-rw-r--r--modules/sycl/tests/SConscript32
-rw-r--r--modules/sycl/tests/data.cpp56
10 files changed, 1237 insertions, 0 deletions
diff --git a/modules/sycl/.nix/derivation.nix b/modules/sycl/.nix/derivation.nix
new file mode 100644
index 0000000..4422739
--- /dev/null
+++ b/modules/sycl/.nix/derivation.nix
@@ -0,0 +1,43 @@
+{ lib
+, stdenv
+, scons
+, clang-tools
+, pname
+, version
+, forstio
+, kel
+, adaptive-cpp
+}:
+
+stdenv.mkDerivation {
+ pname = "${pname}-sycl";
+ inherit version;
+ src = ./..;
+
+ nativeBuildInputs = [
+ scons
+ clang-tools
+ ];
+
+ buildInputs = [
+ forstio.core
+ forstio.async
+ forstio.codec
+ forstio.codec-unit
+ forstio.codec-json
+ forstio.remote
+ #forstio.remote-sycl
+ kel.lbm.core
+ adaptive-cpp
+ ];
+
+ doCheck = true;
+ checkPhase = ''
+ scons test
+ ./bin/tests
+ '';
+
+ preferLocalBuild = true;
+
+ outputs = [ "out" "dev" ];
+}
diff --git a/modules/sycl/SConstruct b/modules/sycl/SConstruct
new file mode 100644
index 0000000..f8bfbcb
--- /dev/null
+++ b/modules/sycl/SConstruct
@@ -0,0 +1,80 @@
+#!/usr/bin/env python3
+
+import sys
+import os
+import os.path
+import glob
+import re
+
+
+if sys.version_info < (3,):
+ def isbasestring(s):
+ return isinstance(s,basestring)
+else:
+ def isbasestring(s):
+ return isinstance(s, (str,bytes))
+
+def add_kel_source_files(self, sources, filetype, lib_env=None, shared=False, target_post=""):
+
+ if isbasestring(filetype):
+ dir_path = self.Dir('.').abspath
+ filetype = sorted(glob.glob(dir_path+"/"+filetype))
+
+ for path in filetype:
+ target_name = re.sub( r'(.*?)(\.cpp|\.c\+\+)', r'\1' + target_post, path )
+ if shared:
+ target_name+='.os'
+ sources.append( self.SharedObject( target=target_name, source=path ) )
+ else:
+ target_name+='.o'
+ sources.append( self.StaticObject( target=target_name, source=path ) )
+ pass
+
+def isAbsolutePath(key, dirname, env):
+ assert os.path.isabs(dirname), "%r must have absolute path syntax" % (key,)
+
+env_vars = Variables(
+ args=ARGUMENTS
+)
+
+env_vars.Add('prefix',
+ help='Installation target location of build results and headers',
+ default='/usr/local/',
+ validator=isAbsolutePath
+)
+
+env=Environment(ENV=os.environ, variables=env_vars, CPPPATH=[],
+ CPPDEFINES=['SAW_UNIX'],
+ CXX = ['syclcc-clang'],
+ CXXFLAGS=[
+ '-std=c++20',
+ '-g',
+ '-O3',
+ '-Wall',
+ '-Wextra',
+ '-isystem', 'AdaptiveCpp'
+ ],
+ LIBS=[
+ 'forstio-core',
+ 'acpp-rt',
+ 'omp'
+ ]
+);
+env.__class__.add_source_files = add_kel_source_files
+env.Tool('compilation_db');
+env.cdb = env.CompilationDatabase('compile_commands.json');
+
+env.objects = [];
+env.sources = [];
+env.headers = [];
+env.targets = [];
+
+Export('env')
+SConscript('c++/SConscript')
+SConscript('tests/SConscript')
+
+env.Alias('cdb', env.cdb);
+env.Alias('all', [env.targets]);
+env.Default('all');
+
+env.Alias('install', '$prefix')
diff --git a/modules/sycl/c++/SConscript b/modules/sycl/c++/SConscript
new file mode 100644
index 0000000..2ed63ba
--- /dev/null
+++ b/modules/sycl/c++/SConscript
@@ -0,0 +1,27 @@
+#!/bin/false
+
+import os
+import os.path
+import glob
+
+
+Import('env')
+
+dir_path = Dir('.').abspath
+
+# Environment for base library
+core_env = env.Clone();
+
+core_env.sources = sorted(glob.glob(dir_path + "/*.cpp"));
+core_env.headers = sorted(glob.glob(dir_path + "/*.hpp"));
+
+env.sources += core_env.sources;
+env.headers += core_env.headers;
+
+## Static lib
+objects = []
+core_env.add_source_files(objects, core_env.sources, shared=False);
+env.library_static = core_env.StaticLibrary('#build/kel-lbm-sycl', [objects]);
+
+env.Install('$prefix/lib/', env.library_static);
+env.Install('$prefix/include/kel/lbm/sycl/', core_env.headers);
diff --git a/modules/sycl/c++/common.hpp b/modules/sycl/c++/common.hpp
new file mode 100644
index 0000000..8ff76fc
--- /dev/null
+++ b/modules/sycl/c++/common.hpp
@@ -0,0 +1,12 @@
+#pragma once
+
+#include <kel/lbm/common.hpp>
+#include <AdaptiveCpp/sycl/sycl.hpp>
+
+namespace kel {
+namespace lbm {
+namespace sycl {
+using namespace acpp::sycl;
+}
+}
+}
diff --git a/modules/sycl/c++/data.hpp b/modules/sycl/c++/data.hpp
new file mode 100644
index 0000000..9f43848
--- /dev/null
+++ b/modules/sycl/c++/data.hpp
@@ -0,0 +1,961 @@
+#include "common.hpp"
+#include <kel/lbm/lbm.hpp>
+
+namespace kel {
+namespace lbm {
+namespace encode {
+template<typename Encode>
+struct Sycl {
+};
+}
+
+/*
+namespace impl {
+template<typename Schema>
+struct struct_has_only_equal_dimension_array{};
+}
+*/
+}
+}
+
+namespace saw {
+template<typename Sch, uint64_t... Dims, typename Encode>
+class data<schema::FixedArray<Sch,Dims...>, kel::lbm::encode::Sycl<Encode>> final {
+public:
+ using Schema = schema::FixedArray<Sch,Dims...>;
+private:
+ acpp::sycl::queue* q_;
+ data<Sch,Encode>* values_;
+
+ SAW_FORBID_COPY(data);
+ SAW_FORBID_MOVE(data);
+public:
+ data(const data<typename meta_schema<Schema>::MetaSchema>& meta__, acpp::sycl::queue& q__):
+ q_{&q__},
+ values_{nullptr}
+ {
+ (void) meta__;
+ SAW_ASSERT(q_);
+ values_ = acpp::sycl::malloc_device<data<Sch,Encode>>(ct_multiply<uint64_t,Dims...>::value,*q_);
+ SAW_ASSERT(values_);
+ }
+
+ data(acpp::sycl::queue& q__):
+ q_{&q__},
+ values_{nullptr}
+ {
+ SAW_ASSERT(q_);
+ values_ = acpp::sycl::malloc_device<data<Sch,Encode>>(ct_multiply<uint64_t,Dims...>::value,*q_);
+ SAW_ASSERT(values_);
+ }
+
+ ~data(){
+ if(not values_){
+ return;
+ }
+ SAW_ASSERT(q_);
+
+ acpp::sycl::free(values_,*q_);
+ values_ = nullptr;
+ }
+
+ static constexpr data<schema::FixedArray<schema::UInt64, sizeof...(Dims)>> get_dims() {
+ return saw::data<schema::FixedArray<schema::UInt64, sizeof...(Dims)>>{{Dims...}};
+ }
+
+ constexpr data<Sch,Encode>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Dims)>>& index){
+ return values_[kel::lbm::flatten_index<schema::UInt64,sizeof...(Dims)>::apply(index,get_dims()).get()];
+ }
+
+ constexpr data<Sch,Encode>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Dims)>>& index) const{
+ return values_[kel::lbm::flatten_index<schema::UInt64,sizeof...(Dims)>::apply(index,get_dims()).get()];
+ }
+
+ constexpr data<Sch,Encode>* flat_data() const {
+ return values_;
+ }
+};
+
+template<typename Sch, uint64_t... Dims, typename Encode>
+class data<schema::Ptr<schema::FixedArray<Sch,Dims...>>, kel::lbm::encode::Sycl<Encode>> final {
+public:
+ using Schema = schema::Ptr<schema::FixedArray<Sch,Dims...>>;
+private:
+ data<Sch,Encode>* values_;
+
+public:
+ SAW_DEFAULT_COPY(data);
+ SAW_DEFAULT_MOVE(data);
+
+ data():
+ values_{nullptr}
+ {}
+
+ data(data<schema::FixedArray<Sch,Dims...>, kel::lbm::encode::Sycl<Encode>>& values__):
+ values_{values__.flat_data()}
+ {}
+
+ data(data<Sch,Encode>* values__):
+ values_{values__}
+ {}
+
+ static constexpr data<schema::FixedArray<schema::UInt64, sizeof...(Dims)>> get_dims() {
+ return saw::data<schema::FixedArray<schema::UInt64, sizeof...(Dims)>>{{Dims...}};
+ }
+
+ constexpr data<Sch,Encode>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Dims)>>& index){
+ return values_[kel::lbm::flatten_index<schema::UInt64,sizeof...(Dims)>::apply(index,get_dims()).get()];
+ }
+
+ constexpr data<Sch,Encode>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Dims)>>& index) const{
+ return values_[kel::lbm::flatten_index<schema::UInt64,sizeof...(Dims)>::apply(index,get_dims()).get()];
+ }
+
+ constexpr data<Sch,Encode>* flat_data() const {
+ return values_;
+ }
+};
+
+template<typename Sch, uint64_t Dims, typename Encode>
+class data<schema::Array<Sch,Dims>, kel::lbm::encode::Sycl<Encode>> final {
+public:
+ using Schema = schema::Array<Sch,Dims>;
+private:
+ static_assert(Dims > 0u, "Zero Dim Arrays make no sense here. If you meant to use this for math style approaches then use Tensor instead of Array");
+
+ data<Sch,Encode>* values_;
+ data<schema::FixedArray<schema::UInt64,Dims>,Encode> meta_;
+ acpp::sycl::queue* q_;
+
+ SAW_FORBID_COPY(data);
+ SAW_FORBID_MOVE(data);
+public:
+ data(const data<typename meta_schema<Schema>::MetaSchema>& meta__, acpp::sycl::queue& q__):
+ values_{nullptr},
+ meta_{meta__},
+ q_{&q__}
+ {
+ SAW_ASSERT(q_);
+ values_ = acpp::sycl::malloc_device<data<Sch,Encode>>(flat_size().get(),*q_);
+ SAW_ASSERT(values_);
+ }
+
+ data(acpp::sycl::queue& q__):
+ values_{nullptr},
+ meta_{},
+ q_{&q__}
+ {
+ SAW_ASSERT(q_);
+ }
+
+ ~data(){
+ if(not values_){
+ return;
+ }
+ SAW_ASSERT(q_);
+
+ acpp::sycl::free(values_,*q_);
+ values_ = nullptr;
+ }
+
+ constexpr data<schema::FixedArray<schema::UInt64, Dims>, Encode> meta() const {
+ return meta_;
+ }
+
+ constexpr data<Sch,Encode>& at(const data<schema::FixedArray<schema::UInt64,Dims>, Encode>& index){
+ return values_[kel::lbm::flatten_index<schema::UInt64,Dims>::apply(index,meta()).get()];
+ }
+
+ constexpr data<Sch,Encode>& at(const data<schema::FixedArray<schema::UInt64,Dims>, Encode>& index) const{
+ return values_[kel::lbm::flatten_index<schema::UInt64,Dims>::apply(index,meta()).get()];
+ }
+
+ constexpr error_or<void> reset_to(const data<typename meta_schema<Schema>::MetaSchema>& meta_arg){
+ SAW_ASSERT(q_);
+ meta_ = meta_arg;
+
+ if(values_){
+ acpp::sycl::free(values_,*q_);
+ }
+ values_ = acpp::sycl::malloc_device<data<Sch,Encode>>(flat_size().get(),*q_);
+ SAW_ASSERT(q_);
+
+ return make_void();
+ }
+
+ constexpr data<Sch,Encode>* flat_data() const {
+ return values_;
+ }
+
+ constexpr data<schema::UInt64,Encode> flat_size() const {
+ data<schema::UInt64> mult{1u};
+
+ for(uint64_t i{0u}; i < Dims; ++i){
+ mult = mult * meta_.at({i});
+ }
+
+ return mult;
+ }
+};
+
+template<typename Sch, uint64_t Dims, typename Encode>
+class data<schema::Ptr<schema::Array<Sch,Dims>>, kel::lbm::encode::Sycl<Encode>> final {
+public:
+ using Schema = schema::Ptr<schema::Array<Sch,Dims>>;
+private:
+ static_assert(Dims > 0u, "Zero Dim Arrays make no sense here. If you meant to use this for math style approaches then use Tensor instead of Array");
+
+ data<Sch,Encode>* values_;
+ data<schema::FixedArray<schema::UInt64,Dims>,Encode> meta_;
+
+public:
+ SAW_DEFAULT_COPY(data);
+ SAW_DEFAULT_MOVE(data);
+
+ data():
+ values_{nullptr}
+ {}
+
+ data(const data<schema::Array<Sch,Dims>, kel::lbm::encode::Sycl<Encode>>& values__):
+ values_{values__.flat_data()},
+ meta_{values__.meta()}
+ {}
+
+ constexpr data<schema::FixedArray<schema::UInt64, Dims>, Encode> meta() const {
+ return meta_;
+ }
+
+ constexpr data<Sch,Encode>& at(const data<schema::FixedArray<schema::UInt64,Dims>, Encode>& index){
+ return values_[kel::lbm::flatten_index<schema::UInt64,Dims>::apply(index,meta()).get()];
+ }
+
+ constexpr data<Sch,Encode>& at(const data<schema::FixedArray<schema::UInt64,Dims>, Encode>& index) const{
+ return values_[kel::lbm::flatten_index<schema::UInt64,Dims>::apply(index,meta()).get()];
+ }
+
+ constexpr data<Sch,Encode>* flat_data() const {
+ return values_;
+ }
+
+ constexpr data<schema::UInt64,Encode> flat_size() const {
+ data<schema::UInt64> mult{1u};
+
+ for(uint64_t i{0u}; i < Dims; ++i){
+ mult = mult * meta_.at({i});
+ }
+
+ return mult;
+ }
+};
+
+template<typename Sch, uint64_t Ghost, uint64_t... Sides, typename Encode>
+class data<kel::lbm::sch::Chunk<Sch,Ghost,Sides...>,kel::lbm::encode::Sycl<Encode>> final {
+public:
+ using Schema = kel::lbm::sch::Chunk<Sch,Ghost,Sides...>;
+private:
+ using InnerSchema = typename Schema::InnerSchema;
+ using ValueSchema = typename InnerSchema::ValueType;
+
+ data<InnerSchema, kel::lbm::encode::Sycl<Encode>> values_;
+public:
+ data(const data<typename meta_schema<Schema>::MetaSchema>& meta__, acpp::sycl::queue& q__):
+ values_{meta__,q__}
+ {}
+
+ data(acpp::sycl::queue& q__):
+ values_{q__}
+ {}
+
+ constexpr data<ValueSchema, Encode>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){
+ return values_.at(index);
+ }
+
+ constexpr data<ValueSchema, Encode>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index) const {
+ return values_.at(index);
+ }
+
+ static constexpr auto get_ghost_dims() {
+ return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::get_dims();
+ }
+
+ static constexpr auto ghost_meta() {
+ return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::meta();
+ }
+
+ data<ValueSchema, Encode>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){
+ std::decay_t<decltype(index)> ind;
+ for(uint64_t i = 0u; i < sizeof...(Sides); ++i){
+ ind.at({i}) = index.at({i}) + Ghost;
+ }
+ return values_.at(ind);
+ }
+
+ data<ValueSchema, Encode>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index) const {
+ std::decay_t<decltype(index)> ind;
+ for(uint64_t i = 0u; i < sizeof...(Sides); ++i){
+ ind.at({i}) = index.at({i}) + Ghost;
+ }
+ return values_.at(ind);
+ }
+
+ static constexpr auto get_dims(){
+ return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}};
+ }
+
+ static constexpr auto meta(){
+ return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}};
+ }
+
+ auto flat_data() const {
+ return values_.flat_data();
+ }
+
+ static constexpr auto flat_size() {
+ return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::flat_size();
+ }
+};
+
+template<typename Sch, uint64_t Ghost, uint64_t... Sides, typename Encode>
+class data<schema::Ptr<kel::lbm::sch::Chunk<Sch,Ghost,Sides...>>,kel::lbm::encode::Sycl<Encode>> final {
+public:
+ using Schema = schema::Ptr<kel::lbm::sch::Chunk<Sch,Ghost,Sides...>>;
+private:
+ using InnerSchema = typename kel::lbm::sch::Chunk<Sch,Ghost,Sides...>::InnerSchema;
+ using ValueSchema = typename InnerSchema::ValueType;
+
+ data<schema::Ptr<InnerSchema>,kel::lbm::encode::Sycl<Encode>> values_;
+
+public:
+ SAW_DEFAULT_MOVE(data);
+ SAW_DEFAULT_COPY(data);
+
+ data():
+ values_{nullptr}
+ {}
+
+ data(const data<kel::lbm::sch::Chunk<Sch,Ghost,Sides...>, kel::lbm::encode::Sycl<Encode>>& values__):
+ values_{values__.flat_data()}
+ {}
+
+ data(data<Sch,Encode>* values__):
+ values_{values__}
+ {}
+
+ constexpr data<ValueSchema, Encode>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){
+ return values_.at(index);
+ }
+
+ constexpr data<ValueSchema, Encode>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index) const {
+ return values_.at(index);
+ }
+
+ static constexpr auto get_ghost_dims() {
+ return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::get_dims();
+ }
+
+ static constexpr auto ghost_meta() {
+ return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::meta();
+ }
+
+ data<ValueSchema, Encode>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){
+ std::decay_t<decltype(index)> ind;
+ for(uint64_t i = 0u; i < sizeof...(Sides); ++i){
+ ind.at({i}) = index.at({i}) + Ghost;
+ }
+ return values_.at(ind);
+ }
+
+ data<ValueSchema, Encode>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index) const {
+ std::decay_t<decltype(index)> ind;
+ for(uint64_t i = 0u; i < sizeof...(Sides); ++i){
+ ind.at({i}) = index.at({i}) + Ghost;
+ }
+ return values_.at(ind);
+ }
+
+ static constexpr auto meta(){
+ return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}};
+ }
+
+ static constexpr auto get_dims(){
+ return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}};
+ }
+
+ auto flat_data() const {
+ return values_.flat_data();
+ }
+
+ static constexpr auto flat_size() {
+ return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::flat_size();
+ }
+};
+
+template<typename... Members, typename Encode>
+struct data<schema::Tuple<Members...>, kel::lbm::encode::Sycl<Encode>> final {
+public:
+ using StorageT = std::tuple<data<Members,kel::lbm::encode::Sycl<Encode>>...>;
+ using Schema = schema::Tuple<Members...>;
+private:
+ StorageT members_;
+
+ /**
+ * A helper constructor to forward the sycl queue to the inner "default" constructors
+ */
+ template<std::size_t... Is>
+ constexpr data(acpp::sycl::queue& q, std::index_sequence<Is...>):
+ members_{(static_cast<void>(Is), q)...}
+ {}
+public:
+ /*
+ data(data<typename meta_schema<Schema>::MetaSchema>& meta__, acpp::sycl::queue& q__):
+ data{q__, std::make_index_sequence<sizeof...(Members)>{}}
+ {}
+ */
+
+ data(acpp::sycl::queue& q__):
+ data{q__, std::make_index_sequence<sizeof...(Members)>{}}
+ {
+ q__.wait();
+ }
+
+ template<size_t i>
+ auto& get(){
+ return std::get<i>(members_);
+ }
+
+ template<size_t i>
+ auto& get() const {
+ return std::get<i>(members_);
+ }
+};
+
+template<typename... Members, typename Encode>
+struct data<schema::Ptr<schema::Tuple<Members...>>, kel::lbm::encode::Sycl<Encode>> final {
+public:
+ using StorageT = std::tuple<data<schema::Ptr<Members>,kel::lbm::encode::Sycl<Encode>>...>;
+ using Schema = schema::Tuple<Members...>;
+private:
+ StorageT members_;
+
+ /**
+ * A helper constructor to forward the sycl queue to the inner "default" constructors
+ */
+public:
+ data() = default;
+
+ template<size_t i>
+ auto& get(){
+ return std::get<i>(members_);
+ }
+
+ template<size_t i>
+ const auto& get() const {
+ return std::get<i>(members_);
+ }
+};
+
+template<typename... Members, typename Encode>
+class data<schema::Struct<Members...>, kel::lbm::encode::Sycl<Encode> > final {
+public:
+ using StorageT = std::tuple<data<typename Members::ValueType,kel::lbm::encode::Sycl<Encode>>...>;
+ using Schema = schema::Struct<Members...>;
+private:
+ /**
+ * @todo Check by static assert that the members all have the same dimensions. Alternatively
+ * Do it here by specializing.
+ */
+ StorageT members_;
+
+ /**
+ * A helper constructor to forward the sycl queue to the inner "default" constructors
+ */
+ template<std::size_t... Is>
+ constexpr data(acpp::sycl::queue& q, std::index_sequence<Is...>):
+ members_{(static_cast<void>(Is), q)...}
+ {}
+public:
+ data(acpp::sycl::queue& q__):
+ data{q__, std::make_index_sequence<sizeof...(Members)>{}}
+ {
+ q__.wait();
+ }
+
+ template<size_t i>
+ auto& get(){
+ return std::get<i>(members_);
+ }
+
+ template<size_t i>
+ auto& get() const {
+ return std::get<i>(members_);
+ }
+
+ template<saw::string_literal K>
+ auto& get(){
+ return std::get<parameter_key_pack_index<K, Members::KeyLiteral...>::value>(members_);
+ }
+
+ template<saw::string_literal K>
+ auto& get() const {
+ return std::get<parameter_key_pack_index<K, Members::KeyLiteral...>::value>(members_);
+ }
+};
+
+template<typename... Sch, saw::string_literal... Keys, typename Encode>
+class data<schema::Ptr<schema::Struct<schema::Member<Sch,Keys>...>>, kel::lbm::encode::Sycl<Encode> > final {
+public:
+ using StorageT = std::tuple<data<schema::Ptr<Sch>,kel::lbm::encode::Sycl<Encode>>...>;
+ using Schema = schema::Struct<schema::Member<schema::Ptr<Sch>,Keys>...>;
+private:
+ /**
+ * @todo Check by static assert that the members all have the same dimensions. Alternatively
+ * Do it here by specializing.
+ */
+ StorageT members_;
+public:
+ data() = default;
+
+ template<size_t i>
+ auto& get(){
+ return std::get<i>(members_);
+ }
+
+ template<size_t i>
+ auto& get() const {
+ return std::get<i>(members_);
+ }
+
+ template<saw::string_literal K>
+ auto& get(){
+ return std::get<parameter_key_pack_index<K, Keys...>::value>(members_);
+ }
+
+ template<saw::string_literal K>
+ auto& get() const {
+ return std::get<parameter_key_pack_index<K, Keys...>::value>(members_);
+ }
+};
+}
+
+namespace kel {
+namespace lbm {
+namespace impl {
+template<typename Sch, typename Encode>
+struct sycl_copy_helper;
+
+template<typename... Members, typename Encode>
+struct sycl_copy_helper<sch::Struct<Members...>, Encode> final {
+ using Schema = sch::Struct<Members...>;
+
+ template<uint64_t i>
+ static saw::error_or<void> copy_to_device_member(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){
+ if constexpr (i < sizeof...(Members)){
+ using M = typename saw::parameter_pack_type<i,Members...>::type;
+ auto& host_member_data = host_data.template get<M::KeyLiteral>();
+ auto& sycl_member_data = sycl_data.template get<M::KeyLiteral>();
+
+ auto eov = sycl_copy_helper<typename M::ValueType,Encode>::copy_to_device(host_member_data,sycl_member_data,q);
+ if(eov.is_error()){
+ return eov;
+ }
+
+ return copy_to_device_member<i+1u>(host_data,sycl_data,q);
+ }
+
+ return saw::make_void();
+ }
+
+ static saw::error_or<void> copy_to_device(saw::data<Schema,Encode>& host_data, saw::data<Schema, encode::Sycl<Encode>>& sycl_data, sycl::queue& q){
+
+ return copy_to_device_member<0u>(host_data, sycl_data, q);
+ }
+
+ template<uint64_t i>
+ static saw::error_or<void> copy_to_host_member(saw::data<Schema,encode::Sycl<Encode>>& sycl_data, saw::data<Schema,Encode>& host_data, sycl::queue& q){
+ if constexpr (i < sizeof...(Members)){
+ using M = typename saw::parameter_pack_type<i,Members...>::type;
+ auto& host_member_data = host_data.template get<M::KeyLiteral>();
+ auto& sycl_member_data = sycl_data.template get<M::KeyLiteral>();
+
+ auto eov = sycl_copy_helper<typename M::ValueType,Encode>::copy_to_host(sycl_member_data,host_member_data,q);
+ if(eov.is_error()){
+ return eov;
+ }
+
+ return copy_to_host_member<i+1u>(sycl_data,host_data,q);
+ }
+
+ return saw::make_void();
+ }
+
+
+ static saw::error_or<void> copy_to_host(saw::data<Schema,encode::Sycl<Encode>>& sycl_data, saw::data<Schema,Encode>& host_data, sycl::queue& q){
+ return copy_to_host_member<0u>(sycl_data, host_data, q);
+ }
+
+ template<uint64_t i>
+ static saw::error_or<void> malloc_on_device_member(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){
+
+ if constexpr (i < sizeof...(Members)){
+ using M = typename saw::parameter_pack_type<i,Members...>::type;
+ auto& host_member_data = host_data.template get<i>();
+ auto& sycl_member_data = sycl_data.template get<i>();
+
+ auto eov = sycl_copy_helper<typename M::ValueType,Encode>::malloc_on_device(host_member_data,sycl_member_data,q);
+ if(eov.is_error()){
+ return eov;
+ }
+
+ return malloc_on_device_member<i+1u>(host_data,sycl_data,q);
+ }
+
+ return saw::make_void();
+ }
+
+ static saw::error_or<void> malloc_on_device(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){
+ return malloc_on_device_member<0u>(host_data,sycl_data,q);
+ }
+};
+
+template<typename... Members, typename Encode>
+struct sycl_copy_helper<sch::Tuple<Members...>, Encode> final {
+ using Schema = sch::Tuple<Members...>;
+
+ template<uint64_t i>
+ static saw::error_or<void> copy_to_device_member(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){
+ if constexpr (i < sizeof...(Members)){
+ using M = typename saw::parameter_pack_type<i,Members...>::type;
+ auto& host_member_data = host_data.template get<i>();
+ auto& sycl_member_data = sycl_data.template get<i>();
+
+ auto eov = sycl_copy_helper<M,Encode>::copy_to_device(host_member_data,sycl_member_data,q);
+ if(eov.is_error()){
+ return eov;
+ }
+
+ return copy_to_device_member<i+1u>(host_data,sycl_data,q);
+ }
+
+ return saw::make_void();
+ }
+
+ static saw::error_or<void> copy_to_device(saw::data<Schema,Encode>& host_data, saw::data<Schema, encode::Sycl<Encode>>& sycl_data, sycl::queue& q){
+ return copy_to_device_member<0u>(host_data, sycl_data, q);
+ }
+
+ template<uint64_t i>
+ static saw::error_or<void> copy_to_host_member(saw::data<Schema,encode::Sycl<Encode>>& sycl_data, saw::data<Schema,Encode>& host_data, sycl::queue& q){
+ if constexpr (i < sizeof...(Members)){
+ using M = typename saw::parameter_pack_type<i,Members...>::type;
+
+ auto& host_member_data = host_data.template get<i>();
+ auto& sycl_member_data = sycl_data.template get<i>();
+
+ auto eov = sycl_copy_helper<M,Encode>::copy_to_host(sycl_member_data,host_member_data,q);
+ if(eov.is_error()){
+ return eov;
+ }
+
+ return copy_to_host_member<i+1u>(sycl_data,host_data,q);
+ }
+
+ return saw::make_void();
+ }
+
+
+ static saw::error_or<void> copy_to_host(
+ saw::data<Schema,Encode>& host_data,
+ saw::data<Schema,encode::Sycl<Encode>>& sycl_data,
+ sycl::queue& q
+ ){
+ return copy_to_host_member<0u>(sycl_data, host_data, q);
+ }
+
+ template<uint64_t i>
+ static saw::error_or<void> malloc_on_device_member(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){
+
+ if constexpr (i < sizeof...(Members)){
+ using M = typename saw::parameter_pack_type<i,Members...>::type;
+ auto& host_member_data = host_data.template get<i>();
+ auto& sycl_member_data = sycl_data.template get<i>();
+
+ auto eov = sycl_copy_helper<M,Encode>::malloc_on_device(host_member_data,sycl_member_data,q);
+ if(eov.is_error()){
+ return eov;
+ }
+
+ return malloc_on_device_member<i+1u>(host_data,sycl_data,q);
+ }
+
+ return saw::make_void();
+ }
+
+ static saw::error_or<void> malloc_on_device(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){
+ return malloc_on_device_member<0u>(host_data,sycl_data,q);
+ }
+};
+
+template<typename Sch, uint64_t... Dims, typename Encode>
+struct sycl_copy_helper<sch::FixedArray<Sch,Dims...>, Encode> final {
+ using Schema = sch::FixedArray<Sch,Dims...>;
+
+ static saw::error_or<void> copy_to_host(saw::data<Schema,encode::Sycl<Encode>>& sycl_data, saw::data<Schema,Encode>& host_data, sycl::queue& q){
+ auto host_ptr = host_data.flat_data();
+ auto sycl_ptr = sycl_data.flat_data();
+
+ static_assert(sizeof(std::decay_t<decltype(sycl_ptr)>) == sizeof(std::decay_t<decltype(host_ptr)>), "Unequal size");
+
+ q.submit([&](acpp::sycl::handler& h){
+ h.copy(sycl_ptr,host_ptr, saw::ct_multiply<uint64_t,Dims...>::value);
+ }).wait();
+ return saw::make_void();
+ }
+
+ static saw::error_or<void> copy_to_device(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){
+ auto host_ptr = host_data.flat_data();
+ auto sycl_ptr = sycl_data.flat_data();
+
+ static_assert(sizeof(std::decay_t<decltype(sycl_ptr)>) == sizeof(std::decay_t<decltype(host_ptr)>), "Unequal size");
+
+ q.submit([&](acpp::sycl::handler& h){
+ h.copy(host_ptr,sycl_ptr, saw::ct_multiply<uint64_t,Dims...>::value);
+ }).wait();
+ return saw::make_void();
+ }
+
+ static saw::error_or<void> malloc_on_device(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){
+ (void) host_data;
+ (void) sycl_data;
+ (void) q;
+ return saw::make_void();
+ }
+};
+
+template<typename Sch, uint64_t Ghost, uint64_t... Dims, typename Encode>
+struct sycl_copy_helper<sch::Chunk<Sch,Ghost,Dims...>, Encode> final {
+ using Schema = sch::Chunk<Sch,Ghost,Dims...>;
+
+ static saw::error_or<void> copy_to_host(saw::data<Schema,encode::Sycl<Encode>>& sycl_data, saw::data<Schema,Encode>& host_data, sycl::queue& q){
+ auto host_ptr = host_data.flat_data();
+ auto sycl_ptr = sycl_data.flat_data();
+
+ static_assert(sizeof(std::decay_t<decltype(sycl_ptr)>) == sizeof(std::decay_t<decltype(host_ptr)>), "Unequal size");
+
+ auto flat_size = host_data.flat_size();
+
+ q.submit([&](acpp::sycl::handler& h){
+ h.copy(sycl_ptr,host_ptr, flat_size.get());
+ }).wait();
+ return saw::make_void();
+ }
+
+ static saw::error_or<void> copy_to_device(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){
+ auto host_ptr = host_data.flat_data();
+ auto sycl_ptr = sycl_data.flat_data();
+
+ static_assert(sizeof(std::decay_t<decltype(sycl_ptr)>) == sizeof(std::decay_t<decltype(host_ptr)>), "Unequal size");
+
+ auto flat_size = host_data.flat_size();
+
+ q.submit([&](acpp::sycl::handler& h){
+ h.copy(host_ptr,sycl_ptr, flat_size.get());
+ }).wait();
+ return saw::make_void();
+ }
+};
+
+template<typename Sch, uint64_t Dims, typename Encode>
+struct sycl_copy_helper<sch::Array<Sch,Dims>, Encode> final {
+ using Schema = sch::Array<Sch,Dims>;
+
+ static saw::error_or<void> copy_to_host(saw::data<Schema,encode::Sycl<Encode>>& sycl_data, saw::data<Schema,Encode>& host_data, sycl::queue& q){
+ auto host_ptr = host_data.flat_data();
+ auto sycl_ptr = sycl_data.flat_data();
+
+ static_assert(sizeof(std::decay_t<decltype(sycl_ptr)>) == sizeof(std::decay_t<decltype(host_ptr)>), "Unequal size");
+
+ SAW_ASSERT(host_data.flat_size() == sycl_data.flat_size());
+ q.submit([&](acpp::sycl::handler& h){
+ h.copy(sycl_ptr,host_ptr, host_data.flat_size().get());
+ }).wait();
+ return saw::make_void();
+ }
+
+ static saw::error_or<void> copy_to_device(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){
+
+ {
+ auto hm = host_data.meta();
+ auto sm = sycl_data.meta();
+ bool equ{true};
+ for(uint64_t i{0u}; i < Dims; ++i){
+ equ &= (hm.at({i}).get() == sm.at({i}).get());
+ }
+ if(not equ){
+ sycl_data.reset_to(hm);
+ }
+ }
+
+ auto host_ptr = host_data.flat_data();
+ auto sycl_ptr = sycl_data.flat_data();
+ static_assert(sizeof(std::decay_t<decltype(sycl_ptr)>) == sizeof(std::decay_t<decltype(host_ptr)>), "Unequal size");
+
+ q.submit([&](acpp::sycl::handler& h){
+ h.copy(host_ptr,sycl_ptr, host_data.flat_size().get());
+ }).wait();
+
+ return saw::make_void();
+ }
+
+ static saw::error_or<void> malloc_on_device(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){
+ sycl_data = {host_data.meta(),q};
+ return saw::make_void();
+ }
+};
+
+
+template<typename Schema, typename Encode>
+struct make_view_helper;
+
+template<typename Sch, uint64_t Dims, typename Encode>
+struct make_view_helper<sch::Array<Sch,Dims>, encode::Sycl<Encode>> {
+public:
+static saw::error_or<void> apply(
+ saw::data<sch::Array<Sch,Dims>, encode::Sycl<Encode>>& dat,
+ saw::data<sch::Ptr<sch::Array<Sch,Dims>>, encode::Sycl<Encode>>& dat_view
+){
+ dat_view = {dat};
+ return saw::make_void();
+}
+};
+
+template<typename Sch, uint64_t... Dims, typename Encode>
+struct make_view_helper<sch::FixedArray<Sch,Dims...>, encode::Sycl<Encode>> {
+public:
+ static saw::error_or<void> apply(
+ saw::data<sch::FixedArray<Sch,Dims...>, encode::Sycl<Encode>>& dat,
+ saw::data<sch::Ptr<sch::FixedArray<Sch,Dims...>>, encode::Sycl<Encode>>& dat_view
+ ){
+ dat_view = {dat};
+ return saw::make_void();
+ }
+};
+
+template<typename Sch, uint64_t Ghost, uint64_t... Dims, typename Encode>
+struct make_view_helper<sch::Chunk<Sch,Ghost,Dims...>, encode::Sycl<Encode>> {
+public:
+ static saw::error_or<void> apply(
+ saw::data<sch::Chunk<Sch,Ghost,Dims...>, encode::Sycl<Encode>>& dat,
+ saw::data<sch::Ptr<sch::Chunk<Sch,Ghost,Dims...>>, encode::Sycl<Encode>>& dat_view
+ ){
+ dat_view = {dat};
+ return saw::make_void();
+ }
+};
+
+template<typename... Sch, saw::string_literal... Keys, typename Encode>
+struct make_view_helper<sch::Struct<sch::Member<Sch,Keys>...>, encode::Sycl<Encode>> {
+private:
+template<uint64_t i>
+static saw::error_or<void> apply_i(
+ saw::data<sch::Struct<sch::Member<Sch,Keys>...>, encode::Sycl<Encode>>& dat,
+ saw::data<sch::Ptr<sch::Struct<sch::Member<Sch,Keys>...>>, encode::Sycl<Encode>>& dat_view
+ ){
+ if constexpr (i < sizeof...(Sch)){
+ using M = typename saw::parameter_pack_type<i,sch::Member<Sch,Keys>...>::type;
+
+ auto eov = make_view_helper<typename M::ValueType,encode::Sycl<Encode>>::apply(dat.template get<M::KeyLiteral>(),dat_view.template get<M::KeyLiteral>());
+ if(eov.is_error()){
+ return eov;
+ }
+
+ return apply_i<i+1u>(dat,dat_view);
+ }
+
+ return saw::make_void();
+ }
+public:
+static saw::error_or<void> apply(
+ saw::data<sch::Struct<sch::Member<Sch,Keys>...>, encode::Sycl<Encode>>& dat,
+ saw::data<sch::Ptr<sch::Struct<sch::Member<Sch,Keys>...>>, encode::Sycl<Encode>>& dat_view
+){
+ return apply_i<0u>(dat,dat_view);
+}
+};
+
+template<typename... Sch, typename Encode>
+struct make_view_helper<sch::Tuple<Sch...>, encode::Sycl<Encode>> {
+private:
+ template<uint64_t i>
+ static saw::error_or<void> apply_i(
+ saw::data<sch::Tuple<Sch...>, encode::Sycl<Encode>>& dat,
+ saw::data<sch::Ptr<sch::Tuple<Sch...>>, encode::Sycl<Encode>>& dat_view
+ ){
+ if constexpr (i < sizeof...(Sch)){
+ using M = typename saw::parameter_pack_type<i,Sch...>::type;
+
+ auto eov = make_view_helper<M,encode::Sycl<Encode>>::apply(dat.template get<i>(), dat_view.template get<i>());
+ if(eov.is_error()){
+ return eov;
+ }
+
+ return apply_i<i+1u>(dat,dat_view);
+ }
+
+ return saw::make_void();
+ }
+public:
+ static saw::error_or<void> apply(
+ saw::data<sch::Tuple<Sch...>, encode::Sycl<Encode>>& dat,
+ saw::data<sch::Ptr<sch::Tuple<Sch...>>, encode::Sycl<Encode>> dat_view
+ ){
+ return apply_i<0u>(dat,dat_view);
+ }
+};
+}
+
+// Ptr => Ptr<Chunk<T,Ghost,Dims...>>> => Ptr<FixedArray<T,Dims+Ghost...>
+template<typename Schema,typename Encode>
+auto make_view(saw::data<Schema,Encode>& dat){
+ saw::data<sch::Ptr<Schema>,Encode> dat_view;
+ auto eov = impl::make_view_helper<Schema,Encode>::apply(dat,dat_view);
+ (void) eov;
+
+ return dat_view;
+}
+
+class device final {
+private:
+ sycl::queue q_;
+
+ SAW_FORBID_COPY(device);
+ SAW_FORBID_MOVE(device);
+public:
+ device() = default;
+ ~device() = default;
+
+ template<typename Sch, typename Encode>
+ saw::error_or<void> copy_to_device(saw::data<Sch,Encode>& host_data, saw::data<Sch,encode::Sycl<Encode>>& sycl_data){
+ return impl::sycl_copy_helper<Sch,Encode>::copy_to_device(host_data, sycl_data, q_);
+ }
+
+ template<typename Sch, typename Encode>
+ saw::error_or<void> copy_to_host(saw::data<Sch,encode::Sycl<Encode>>& sycl_data, saw::data<Sch,Encode>& host_data){
+ return impl::sycl_copy_helper<Sch,Encode>::copy_to_host(sycl_data, host_data, q_);
+ }
+
+ template<typename Sch, typename Encode>
+ saw::error_or<void> malloc_on_device(
+ saw::data<Sch,Encode>& host_data,
+ saw::data<Sch,encode::Sycl<Encode>>& sycl_data
+ ){
+ auto eov = impl::sycl_copy_helper<Sch,Encode>::malloc_on_device(host_data, sycl_data, q_);
+ q_.wait();
+ return eov;
+ }
+
+ auto& get_handle(){
+ return q_;
+ }
+};
+}
+}
diff --git a/modules/sycl/c++/index.hpp b/modules/sycl/c++/index.hpp
new file mode 100644
index 0000000..0d2c035
--- /dev/null
+++ b/modules/sycl/c++/index.hpp
@@ -0,0 +1,18 @@
+#pragma once
+
+#include "common.hpp"
+
+namespace kel {
+namespace lbm {
+template<typename D>
+saw::data<sch::FixedArray<sch::UInt64,D>> sycl_to_saw_index(const acpp::sycl::id<D>& idx){
+ saw::data<sch::FixedArray<sch::UInt64,D>> trans_index;
+
+ for(uint64_t i{0u}; i < D; ++i){
+ trans_index.at({{i}}).set(idx[i]);
+ }
+
+ return trans_index;
+}
+}
+}
diff --git a/modules/sycl/c++/lbm.hpp b/modules/sycl/c++/lbm.hpp
new file mode 100644
index 0000000..3e67ae6
--- /dev/null
+++ b/modules/sycl/c++/lbm.hpp
@@ -0,0 +1,3 @@
+#pragma once
+
+#include "data.hpp"
diff --git a/modules/sycl/c++/sycl.hpp b/modules/sycl/c++/sycl.hpp
new file mode 100644
index 0000000..973b511
--- /dev/null
+++ b/modules/sycl/c++/sycl.hpp
@@ -0,0 +1,5 @@
+#pragma once
+
+#include "common.hpp"
+#include "data.hpp"
+#include "index.hpp"
diff --git a/modules/sycl/tests/SConscript b/modules/sycl/tests/SConscript
new file mode 100644
index 0000000..d1b381e
--- /dev/null
+++ b/modules/sycl/tests/SConscript
@@ -0,0 +1,32 @@
+#!/bin/false
+
+import os
+import os.path
+import glob
+
+
+Import('env')
+
+dir_path = Dir('.').abspath
+
+# Environment for base library
+test_cases_env = env.Clone();
+
+test_cases_env.Append(LIBS=['forstio-test']);
+
+test_cases_env.sources = sorted(glob.glob(dir_path + "/*.cpp"))
+test_cases_env.headers = sorted(glob.glob(dir_path + "/*.hpp"))
+
+env.sources += test_cases_env.sources;
+env.headers += test_cases_env.headers;
+
+objects_static = []
+test_cases_env.add_source_files(objects_static, test_cases_env.sources, shared=False);
+test_cases_env.program = test_cases_env.Program('#bin/tests', [objects_static]);
+# , env.library_static]);
+
+# Set Alias
+env.Alias('test', test_cases_env.program);
+env.Alias('check', test_cases_env.program);
+
+env.targets += ['test','check'];
diff --git a/modules/sycl/tests/data.cpp b/modules/sycl/tests/data.cpp
new file mode 100644
index 0000000..4321a0d
--- /dev/null
+++ b/modules/sycl/tests/data.cpp
@@ -0,0 +1,56 @@
+#include <forstio/test/suite.hpp>
+
+#include "../c++/lbm.hpp"
+
+namespace {
+
+namespace sch {
+using namespace kel::lbm::sch;
+using TestObjSchema = Tuple<
+ Member<FixedArray<UInt64,2u,2u>, "foo">,
+ Member<Array<Float32>, "bar">,
+ Member<
+ Array<
+ Struct<
+ Member<FixedArray<Float32,2u>,"pos">
+ >
+ >,
+ "baz"
+ >
+>;
+}
+
+SAW_TEST("Sycl Data Compilation"){
+ acpp::sycl::queue q;
+ saw::data<
+ saw::schema::Struct<
+ saw::schema::Member<
+ kel::lbm::sch::Chunk<saw::schema::UInt8,1u,1u,1u>,
+ "test"
+ >
+ >,
+ kel::lbm::encode::Sycl<saw::encode::Native>
+ > dat{q};
+
+ auto& test_f = dat.template get<"test">();
+
+ // test_f.at({}).set(1);
+ // SAW_EXPECT(test_f.at({}).get() == 1, "Value check failed");
+}
+
+SAW_TEST("Sycl Data Array of Struct"){
+ acpp::sycl::queue q;
+
+ saw::data<sch::Array<sch::Float64>, kel::lbm::encode::Sycl<saw::encode::Native>> a{{{2u}},q};
+}
+
+/*
+SAW_TEST("Sycl Data Compilation for Particle Similacrum"){
+ acpp::sycl::queue q;
+
+ saw::data<
+ sch::TestObjSchema
+ > a;
+}
+*/
+}
generated by cgit v1.2.3 (git 2.25.1) at 2026-07-21 21:13:19 +0000