1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
|
#pragma once
#include "common.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<uint64_t... Meta, typename... Sch, string_literal... Keys, typename Encode>
class data<schema::Struct<schema::Member<schema::Chunk<Sch,Ghost,Meta...>, Keys>...>, kel::lbm::encode::Sycl<Encode>> final {
public:
static constexpr data<schema::FixedArray meta = {{Meta...}};
using StorageT = std::tuple<data<Members::Type::InnerSchema,Encode>*...>;
private:
/**
* @todo Check by static assert that the members all have the same dimensions. Alternatively
* Do it here by specializing.
*/
StorageT members_;
kel::lbm::sycl::queue* q_;
public:
data():
members_{},
q_{nullptr}
{}
data(StorageT members__, kel::lbm::sycl::queue& q__):
members_{members__},
q_{&q__}
{}
~data(){
SAW_ASSERT(q_){
return;
}
std::visit([this](auto arg){
if(not arg){
return;
}
sycl::free(arg,*q_);
arg = nullptr;
},members_);
}
template<saw::string_literal K>
auto* get_ptr(){
return std::get<parameter_key_pack_index<K, Keys...>::value(members_);
}
template<saw::string_literal K>
auto& get(){
auto ptr = get_ptr<K>();
SAW_ASSERT(ptr);
return *ptr;
}
};
namespace kel {
namespace lbm {
namespace impl {
template<typename Sch, typename Encode>
struct sycl_malloc_struct_helper;
template<typename... Members, typename Encode>
struct sycl_malloc_struct_helper<sch::Struct<Members...>, Encode> final {
template<uint64_t i>
static saw::error_or<void> allocate_on_device_member(typename data<Sch,encode::Sycl<Encode>>::StorageT& storage, sycl::queue& q){
if constexpr (i < sizeof...(Members)){
using M = typename saw::parameter_pack_type<i,Members...>::type;
auto& ptr = std::get<i>(storage);
ptr = sycl::malloc_device<M::ValueType::InnerSchema>(,q);
return allocate_on_device_member<i+1u>(storage,q);
}
return saw::make_void();
}
static saw::error_or<void> allocate_on_device(data<Sch,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){
typename data<Sch,encode::Sycl<Encode>>::StorageT storage;
auto eov = allocate_on_device_member<0u>(storage,q);
sycl_data = {storage, q};
return eov;
}
};
}
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> allocate_on_device(data<Sch,encode::Sycl<Encode>>& sycl_data){
auto eov = sycl_malloc_struct_helper<Sch,Encode>::allocate_on_device(sycl_data, q_);
if(eov.is_error()){
return eov;
}
sycl_data.set_queue(q_);
}
};
}
}
|
