Alien  1.3.0
Developer documentation
Loading...
Searching...
No Matches
SYCLVectorInternal.h
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#pragma once
9
10#include <alien/kernels/sycl/SYCLPrecomp.h>
11
12#ifdef USE_SYCL2020
13#include <sycl/sycl.hpp>
14#else
15#include <CL/sycl.hpp>
16#endif
17
18#include "SYCLEnv.h"
19#include "SYCLEnvInternal.h"
20/*---------------------------------------------------------------------------*/
21
22namespace Alien::SYCLInternal
23{
24
25/*---------------------------------------------------------------------------*/
26
27#ifndef USE_SYCL2020
28 using namespace cl ;
29#endif
30
31template <typename ValueT = Real>
32class VectorInternal
33{
34 public:
35 // clang-format off
36 typedef ValueT ValueType;
37 typedef VectorInternal<ValueType> ThisType;
38 typedef sycl::buffer<ValueType, 1> ValueBufferType;
39 typedef std::unique_ptr<ValueBufferType> ValueBufferPtrType;
40 // clang-format on
41
42 public:
43 VectorInternal(ValueType const* ptr, std::size_t size)
44 : m_values(ptr, sycl::range<1>(size))
45 {
46 m_values.set_final_data(nullptr);
47 }
48
49 virtual ~VectorInternal() {}
50
51 ValueBufferType& values()
52 {
53 return m_values;
54 }
55
56 ValueBufferType& values() const
57 {
58 return m_values;
59 }
60
61 ValueBufferType& ghostValues(Integer ghost_size) const
62 {
63 if (m_ghost_values.get() == nullptr || ghost_size > m_ghost_size) {
64 m_ghost_size = ghost_size;
65 m_ghost_values.reset(new ValueBufferType(m_ghost_size));
66 }
67 return *m_ghost_values;
68 }
69
70 void copyValuesToHost(std::size_t size, ValueT* ptr)
71 {
72 auto h_values = m_values.get_host_access();
73 for (std::size_t i = 0; i < size; ++i)
74 ptr[i] = h_values[i];
75 }
76
77 void copyValuesToDevice(std::size_t size, ValueT* ptr) const
78 {
79 auto env = SYCLEnv::instance() ;
80 auto& queue = env->internal()->queue() ;
81 auto max_num_treads = env->maxNumThreads() ;
82
83 queue.submit( [&](sycl::handler& cgh)
84 {
85 auto access_x = m_values.template get_access<sycl::access::mode::read>(cgh);
86 std::size_t y_length = size ;
87 cgh.parallel_for<class init_vector_ptr>(sycl::range<1>{max_num_treads}, [=] (sycl::item<1> itemId)
88 {
89 auto id = itemId.get_id(0);
90 for (auto i = id; i < y_length; i += itemId.get_range()[0])
91 ptr[i] = access_x[i];
92 });
93 });
94 queue.wait() ;
95 }
96
97 void setValuesFromDevice(std::size_t size, ValueT const* ptr) const
98 {
99 auto env = SYCLEnv::instance() ;
100 auto& queue = env->internal()->queue() ;
101 auto max_num_treads = env->maxNumThreads() ;
102
103 queue.submit( [&](sycl::handler& cgh)
104 {
105 auto access_x = m_values.template get_access<sycl::access::mode::discard_write>(cgh);
106 std::size_t y_length = size ;
107 cgh.parallel_for<class init_vector_ptr>(sycl::range<1>{max_num_treads}, [=] (sycl::item<1> itemId)
108 {
109 auto id = itemId.get_id(0);
110 for (auto i = id; i < y_length; i += itemId.get_range()[0])
111 access_x[i] = ptr[i];
112 });
113 });
114 queue.wait() ;
115 /*
116 {
117 sycl::host_accessor<ValueT, 1, sycl::access::mode::read> vec_acc(m_values);
118 for(int irow=0;irow<size;++irow)
119 {
120 std::cout<<"VEC["<<irow<<"]"<<vec_acc[irow]<<std::endl;
121 }
122 }*/
123 }
124
125 void setValuesFromHost(std::size_t size, ValueT const* ptr) const
126 {
127 auto env = SYCLEnv::instance() ;
128 auto& queue = env->internal()->queue() ;
129 auto max_num_treads = env->maxNumThreads() ;
130
131 auto rhs = ValueBufferType(ptr,sycl::range<1>(size)) ;
132
133 queue.submit( [&](sycl::handler& cgh)
134 {
135 auto access_x = m_values.template get_access<sycl::access::mode::discard_write>(cgh);
136 auto access_rhs = rhs.template get_access<sycl::access::mode::read>(cgh);
137 std::size_t y_length = size ;
138 cgh.parallel_for<class init_vector_ptr>(sycl::range<1>{max_num_treads}, [=] (sycl::item<1> itemId)
139 {
140 auto id = itemId.get_id(0);
141 for (auto i = id; i < y_length; i += itemId.get_range()[0])
142 access_x[i] = access_rhs[i];
143 });
144 });
145 queue.wait() ;
146 /*
147 {
148 sycl::host_accessor<ValueT, 1, sycl::access::mode::read> vec_acc(m_values);
149 for(int irow=0;irow<size;++irow)
150 {
151 std::cout<<"VEC["<<irow<<"]"<<vec_acc[irow]<<std::endl;
152 }
153 }*/
154 }
155
156 void copy(ValueBufferType& src)
157 {
158 auto env = SYCLEnv::instance() ;
159 env->internal()->queue().submit([&](sycl::handler& cgh)
160 {
161 auto access_x = m_values.template get_access<sycl::access::mode::read_write>(cgh);
162 auto access_src = src.template get_access<sycl::access::mode::read>(cgh);
163 cgh.copy(access_src,access_x) ;
164 }) ;
165 }
166
167 void pointWiseMult(ValueBufferType& y, ValueBufferType& z)
168 {
169 auto env = SYCLEnv::instance() ;
170 auto& queue = env->internal()->queue() ;
171 auto max_num_treads = env->maxNumThreads() ;
172
173 queue.submit( [&](sycl::handler& cgh)
174 {
175 auto access_x = m_values.template get_access<sycl::access::mode::read>(cgh);
176 auto access_y = y.template get_access<sycl::access::mode::read>(cgh);
177 auto access_z = z.template get_access<sycl::access::mode::discard_write>(cgh);
178 std::size_t x_length = m_values.size() ;
179 cgh.parallel_for<class init_vector_ptr>(sycl::range<1>{max_num_treads}, [=] (sycl::item<1> itemId)
180 {
181 auto id = itemId.get_id(0);
182 for (auto i = id; i < x_length; i += itemId.get_range()[0])
183 access_z[i] = access_x[i]*access_y[i];
184 });
185 });
186 }
187
188 void blockMult(std::size_t nrows,
189 int block_size,
190 ValueBufferType& y,
191 ValueBufferType& z)
192 {
193 auto env = SYCLEnv::instance() ;
194 auto& queue = env->internal()->queue() ;
195 auto max_num_treads = env->maxNumThreads() ;
196 int N = block_size ;
197 int NxN = N*N ;
198 assert(m_values.size()>=nrows*NxN) ;
199 assert(y.size()>=nrows*N) ;
200 assert(z.size()>=nrows*N) ;
201 queue.submit( [&](sycl::handler& cgh)
202 {
203 auto access_x = m_values.template get_access<sycl::access::mode::read>(cgh);
204 auto access_y = y.template get_access<sycl::access::mode::read>(cgh);
205 auto access_z = z.template get_access<sycl::access::mode::discard_write>(cgh);
206 cgh.parallel_for<class vector_block_mult>(
207 sycl::range<1>{max_num_treads},
208 [=] (sycl::item<1> itemId)
209 {
210 auto id = itemId.get_id(0);
211 for (auto irow = id; irow < nrows; irow += itemId.get_range()[0])
212 {
213 for(int ieq=0;ieq<N;++ieq)
214 {
215 ValueType value = 0. ;
216 for(int j=0;j<N;++j)
217 {
218 value += access_x[irow*NxN+ieq*N+j]*access_y[irow*N+j] ;
219 }
220 access_z[irow*N+ieq] = value ;
221 }
222 }
223 });
224 });
225#ifdef PRINT_DEBUG_INFO
226 {
227 sycl::host_accessor<ValueT, 1, sycl::access::mode::read> diag_acc(m_values);
228 sycl::host_accessor<ValueT, 1, sycl::access::mode::read> z_acc(z);
229
230 for (std::size_t irow = 0; irow < nrows; ++irow)
231 {
232 std::cout<<"INV DIAG["<<irow<<"]:\n";
233 for(int i=0;i<N;++i)
234 {
235 for(int j=0;j<N;++j)
236 std::cout<<diag_acc[irow*NxN+i*N+j]<<" ";
237 std::cout<<std::endl;
238 }
239 std::cout<<"Y["<<irow<<"]=\n";
240 for(int i=0;i<N;++i)
241 std::cout<<z_acc[irow*N+i]<<std::endl;
242 }
243 }
244#endif
245 }
246
247
248 //VectorInternal<ValueT>* clone() const { return new VectorInternal<ValueT>(*this); }
249
250 // clang-format off
251 mutable ValueBufferType m_values;
252
253 mutable Integer m_ghost_size = 0 ;
254 mutable ValueBufferPtrType m_ghost_values;
255 // clang-format on
256};
257
258/*---------------------------------------------------------------------------*/
259
260} // namespace Alien::SYCLInternal
261
262/*---------------------------------------------------------------------------*/