Alien  1.3.0
Developer documentation
Loading...
Searching...
No Matches
SYCLVector.cc
1// -*- tab-width: 2; indent-tabs-mode: nil; coding: utf-8-with-signature -*-
2//-----------------------------------------------------------------------------
3// Copyright 2000-2026 CEA (www.cea.fr) IFPEN (www.ifpenergiesnouvelles.com)
4// See the top-level COPYRIGHT file for details.
5// SPDX-License-Identifier: Apache-2.0
6//-----------------------------------------------------------------------------
7
8#include <cassert>
9
10#include "SYCLVector.h"
11#include "SYCLVectorInternal.h"
12
13/*---------------------------------------------------------------------------*/
14/*---------------------------------------------------------------------------*/
15
16using namespace Arccore;
17
18namespace Alien
19{
20
22 template <typename ValueT>
23 SYCLVector<ValueT>::SYCLVector()
24 : IVectorImpl(nullptr, AlgebraTraits<BackEnd::tag::sycl>::name())
25 {}
26
28 template <typename ValueT>
29 SYCLVector<ValueT>::SYCLVector(const MultiVectorImpl* multi_impl)
30 : IVectorImpl(multi_impl, AlgebraTraits<BackEnd::tag::sycl>::name())
31 {}
32
33
34 template <typename ValueT>
35 SYCLVector<ValueT>::~SYCLVector()
36 {}
37
38 template <typename ValueT>
39 void SYCLVector<ValueT>::allocate()
40 {
41 //delete m_internal;
42 auto block_size = blockSize() ;
43 m_h_values.resize(m_local_size*block_size);
44 m_internal.reset(new VectorInternal(m_h_values.data(), m_local_size*block_size));
45 }
46
47 template <typename ValueT>
48 void SYCLVector<ValueT>::resize(Integer alloc_size) const
49 {
50 //delete m_internal;
51 m_h_values.resize(alloc_size);
52 m_internal.reset(new VectorInternal(m_h_values.data(), alloc_size));
53 }
54
55 template <typename ValueT>
56 void SYCLVector<ValueT>::clear()
57 {
58 //delete m_internal;
59 //m_internal = nullptr;
60 m_internal.reset() ;
61 std::vector<ValueType>().swap(m_h_values);
62 }
63
64 template <typename ValueT>
65 void SYCLVector<ValueT>::setValuesFromHost()
66 {
67 //delete m_internal;
68 m_internal.reset(new VectorInternal(m_h_values.data(), m_local_size*blockSize()));
69 }
70
71 template <typename ValueT>
72 void SYCLVector<ValueT>::setValues(std::size_t size, ValueType const* ptr)
73 {
74 //delete m_internal;
75 m_h_values.resize(m_local_size*blockSize());
76 std::copy(ptr, ptr + size, m_h_values.begin());
77 m_internal.reset(new VectorInternal(m_h_values.data(), m_local_size*blockSize()));
78 }
79
80 template <typename ValueT>
81 void SYCLVector<ValueT>::setValuesFromDevice(std::size_t size, ValueType const* ptr)
82 {
83 assert(m_internal.get());
84 m_internal->setValuesFromDevice(size,ptr);
85 }
86
87 template <typename ValueT>
88 void SYCLVector<ValueT>::setValuesFromHost(std::size_t size, ValueType const* ptr)
89 {
90 assert(m_internal.get());
91 m_internal->setValuesFromHost(size,ptr);
92 }
93
94
95 template <typename ValueT>
96 void SYCLVector<ValueT>::copyValuesTo(std::size_t size, ValueType* ptr) const
97 {
98 if (m_internal.get())
99 m_internal->copyValuesToHost(size, ptr);
100 }
101
102 template <typename ValueT>
103 void
104 SYCLVector<ValueT>::copyValuesToDevice(std::size_t size, ValueT* ptr) const
105 {
106 if (m_internal.get())
107 m_internal->copyValuesToDevice(size, ptr);
108 }
109
110
111
112 template <typename ValueT>
113 void SYCLVector<ValueT>::initDevicePointers(int** rows, ValueType** values) const
114 {
115 if(m_internal.get()==nullptr)
116 return ;
117
118 auto env = SYCLEnv::instance() ;
119 auto& queue = env->internal()->queue() ;
120 auto max_num_treads = env->maxNumThreads() ;
121
122 auto alloc_size = m_local_size*blockSize() ;
123 auto values_ptr = malloc_device<ValueT>(alloc_size, queue);
124 auto rows_ptr = malloc_device<IndexType>(alloc_size, queue);
125
126 queue.submit( [&](sycl::handler& cgh)
127 {
128 auto access_x = m_internal->m_values.template get_access<sycl::access::mode::read>(cgh);
129 std::size_t y_length = alloc_size ;
130 cgh.parallel_for<class init_vector_ptr>(sycl::range<1>{max_num_treads}, [=] (sycl::item<1> itemId)
131 {
132 auto id = itemId.get_id(0);
133 for (auto i = id; i < y_length; i += itemId.get_range()[0])
134 {
135 values_ptr[i] = access_x[i];
136 rows_ptr[i] = IndexType(i) ;
137 }
138 });
139 });
140 queue.wait() ;
141 *values = values_ptr;
142 *rows = rows_ptr;
143 }
144
145 template <typename ValueT>
146 void SYCLVector<ValueT>::freeDevicePointers(int* rows, ValueType* values)
147 {
148 auto env = SYCLEnv::instance() ;
149 auto& queue = env->internal()->queue() ;
150 sycl::free(values,queue) ;
151 sycl::free(rows,queue) ;
152 }
153
154 template <typename ValueT>
155 void SYCLVector<ValueT>::freeDevicePointers(ValueType* values)
156 {
157 auto env = SYCLEnv::instance() ;
158 auto& queue = env->internal()->queue() ;
159 sycl::free(values,queue) ;
160 }
161
162 template <typename ValueT>
163 void SYCLVector<ValueT>::allocateDevicePointers(std::size_t alloc_size,
164 int** rows,
165 ValueType** values)
166 {
167
168 auto env = SYCLEnv::instance() ;
169 auto& queue = env->internal()->queue() ;
170 auto max_num_treads = env->maxNumThreads() ;
171
172 auto values_ptr = malloc_device<ValueT>(alloc_size, queue);
173 auto rows_ptr = malloc_device<IndexType>(alloc_size, queue);
174 queue.submit( [&](sycl::handler& cgh)
175 {
176 std::size_t y_length = alloc_size ;
177 cgh.parallel_for<class init2_vector_ptr>(sycl::range<1>{max_num_treads}, [=] (sycl::item<1> itemId)
178 {
179 auto id = itemId.get_id(0);
180 for (auto i = id; i < y_length; i += itemId.get_range()[0])
181 {
182 values_ptr[i] = 0.;
183 rows_ptr[i] = IndexType(i) ;
184 }
185 });
186 });
187
188 queue.wait() ;
189 *values = values_ptr;
190 *rows = rows_ptr;
191 }
192
193
194 template <typename ValueT>
195 void SYCLVector<ValueT>::allocateDevicePointers(std::size_t alloc_size,
196 ValueType** values)
197 {
198
199 auto env = SYCLEnv::instance() ;
200 auto& queue = env->internal()->queue() ;
201 auto max_num_treads = env->maxNumThreads() ;
202
203 auto values_ptr = malloc_device<ValueT>(alloc_size, queue);
204 queue.submit( [&](sycl::handler& cgh)
205 {
206 std::size_t y_length = alloc_size ;
207 cgh.parallel_for<class init2_vector_ptr>(sycl::range<1>{max_num_treads}, [=] (sycl::item<1> itemId)
208 {
209 auto id = itemId.get_id(0);
210 for (auto i = id; i < y_length; i += itemId.get_range()[0])
211 {
212 values_ptr[i] = 0.;
213 }
214 });
215 });
216
217 queue.wait() ;
218 *values = values_ptr;
219 }
220
221
222 template <typename ValueT>
223 void SYCLVector<ValueT>::initDevicePointers(std::size_t alloc_size,
224 ValueType const* host_values,
225 int** rows,
226 ValueType** values)
227 {
228
229 auto env = SYCLEnv::instance() ;
230 auto& queue = env->internal()->queue() ;
231 auto max_num_treads = env->maxNumThreads() ;
232
233 auto values_ptr = malloc_device<ValueT>(alloc_size, queue);
234 auto rows_ptr = malloc_device<IndexType>(alloc_size, queue);
235 sycl::buffer<ValueT, 1> values_buf(host_values, sycl::range<1>(alloc_size)) ;
236 queue.submit( [&](sycl::handler& cgh)
237 {
238 auto access_x = values_buf.template get_access<sycl::access::mode::read>(cgh);
239 std::size_t y_length = alloc_size ;
240 cgh.parallel_for<class init3_vector_ptr>(sycl::range<1>{max_num_treads}, [=] (sycl::item<1> itemId)
241 {
242 auto id = itemId.get_id(0);
243 for (auto i = id; i < y_length; i += itemId.get_range()[0])
244 {
245 values_ptr[i] = access_x[i];
246 rows_ptr[i] = IndexType(i) ;
247 }
248 });
249 });
250
251 queue.wait() ;
252 *values = values_ptr;
253 *rows = rows_ptr;
254 }
255
256
257 template <typename ValueT>
258 void SYCLVector<ValueT>::copyDeviceToHost(std::size_t alloc_size,
259 ValueType const* device_values,
260 ValueType* host_values)
261 {
262
263 auto env = SYCLEnv::instance() ;
264 auto& queue = env->internal()->queue() ;
265 auto max_num_treads = env->maxNumThreads() ;
266
267 sycl::buffer<ValueT, 1> values_buf(host_values, sycl::range<1>(alloc_size)) ;
268 queue.submit( [&](sycl::handler& cgh)
269 {
270 auto access_x = sycl::accessor { values_buf, cgh, sycl::write_only, sycl::property::no_init{}};
271 //auto access_x = values_buf.template get_access<sycl::access::mode::write>(cgh);
272 std::size_t y_length = alloc_size ;
273 cgh.parallel_for<class copy_vector_ptr>(sycl::range<1>{max_num_treads}, [=] (sycl::item<1> itemId)
274 {
275 auto id = itemId.get_id(0);
276 for (auto i = id; i < y_length; i += itemId.get_range()[0])
277 {
278 access_x[i] = device_values[i];
279 }
280 });
281 });
282
283 queue.wait() ;
284 }
285
286 template <typename ValueT>
287 void SYCLVector<ValueT>::pointWiseMult(SYCLVector const& y, SYCLVector& z) const
288 {
289 auto block_size = blockSize() ;
290 auto block_size_y = y.blockSize() ;
291 auto block_size_z = z.blockSize() ;
292 assert(block_size_y==block_size_z) ;
293 if(block_size==1)
294 m_internal->pointWiseMult(y.m_internal->m_values,
295 z.m_internal->m_values) ;
296 else
297 {
298 assert(block_size==block_size_z*block_size_z) ;
299 m_internal->blockMult(m_local_size,
300 block_size_z,
301 y.m_internal->m_values,
302 z.m_internal->m_values) ;
303 }
304 }
305 /*---------------------------------------------------------------------------*/
306
307 template class ALIEN_EXPORT SYCLVector<Real>;
308
309 /*---------------------------------------------------------------------------*/
310
311} // namespace Alien
312
313/*---------------------------------------------------------------------------*/
314/*---------------------------------------------------------------------------*/
Alien::IVectorImpl::IVectorImpl
IVectorImpl(const MultiVectorImpl *multi_impl, BackEndId backend="")
Constructor.
Definition IVectorImpl.cc:40
Alien::SYCLVector::SYCLVector
SYCLVector()
Constructeur sans association un MultiImpl.
Definition SYCLVector.cc:23
Alien::SYCLVector::clear
void clear()
Wipe out internal data.
Definition SYCLVector.cc:56
Alien
-- tab-width: 2; indent-tabs-mode: nil; coding: utf-8-with-signature --
Definition BackEnd.h:17