Alien  1.3.0
User documentation
Loading...
Searching...
No Matches
SYCLSendRecvOp.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#pragma once
8
9#include <vector>
10
11#include <alien/kernels/sycl/SYCLPrecomp.h>
12#include <alien/utils/Precomp.h>
13#include <alien/utils/Trace.h>
14
15#include <arccore/message_passing/Messages.h>
16#include <arccore/message_passing/Request.h>
17
18#include <alien/kernels/simple_csr/SendRecvOp.h>
19
20#include <alien/kernels/sycl/data/SYCLEnv.h>
21#include <alien/kernels/sycl/data/SYCLEnvInternal.h>
22
23namespace Alien::SYCLInternal
24{
25
26#ifndef USE_SYCL2020
27 using namespace cl ;
28#endif
29
30template <typename ValueT>
31class SYCLSendRecvOp : public Alien::SimpleCSRInternal::IASynchOp
32{
33 public:
34 // clang-format off
35 typedef ValueT ValueType ;
36 typedef sycl::buffer<ValueType, 1> ValueBufferType ;
37
38 typedef sycl::buffer<int> IndexBufferType ;
39 typedef std::unique_ptr<IndexBufferType> IndexBufferPtrType ;
40 // clang-format on
41
42 SYCLSendRecvOp(ValueBufferType& send_buffer,
43 const Alien::SimpleCSRInternal::CommInfo& send_info,
44 IndexBufferType& send_ids,
45 Alien::SimpleCSRInternal::CommProperty::ePolicyType send_policy,
46 ValueBufferType& recv_buffer,
47 const Alien::SimpleCSRInternal::CommInfo& recv_info,
48 IndexBufferType& recv_ids,
49 Alien::SimpleCSRInternal::CommProperty::ePolicyType recv_policy,
50 IMessagePassingMng* mng,
51 Arccore::ITraceMng* trace_mng)
52 : m_send_buffer(send_buffer)
53 , m_send_info(send_info)
54 , m_send_ids(send_ids)
55 , m_send_policy(send_policy)
56 , m_recv_buffer(recv_buffer)
57 , m_recv_info(recv_info)
58 , m_recv_ids(recv_ids)
59 , m_recv_policy(recv_policy)
60 , m_parallel_mng(mng)
61 , m_trace(trace_mng)
62 {}
63
64 virtual ~SYCLSendRecvOp()
65 {
66#ifdef USE_SYCL_USM
67 auto& queue = SYCLEnv::instance()->internal()->queue();
68 sycl::free(m_rbuffer, queue);
69 sycl::free(m_sbuffer, queue);
70#endif
71 }
72
73 void start([[maybe_unused]] bool insitu = false)
74 {
75 //alien_debug([&] {cout() << "SYCLSendRecvOP START "<<m_send_policy;});
76 //Universe().traceMng()->flush() ;
77
78 // clang-format off
79 auto env = SYCLEnv::instance();
80 auto& queue = env->internal()->queue();
81 auto total_threads = env->maxNumThreads() ;
82 // clang-format on
83 if (m_recv_policy == Alien::SimpleCSRInternal::CommProperty::ASynch) {
84 m_recv_request.resize(m_recv_info.m_num_neighbours);
85 Integer total_nb_recv_ids = m_recv_info.m_ids_offset[m_recv_info.m_num_neighbours] - m_recv_info.m_ids_offset[0];
86#ifdef USE_SYCL_USM
87 m_rbuffer = sycl::malloc_shared<ValueT>(total_nb_recv_ids, queue);
88#else
89 m_rbuffer.resize(total_nb_recv_ids);
90#endif
91 for (Integer i = 0; i < m_recv_info.m_num_neighbours; ++i) {
92 Integer off = m_recv_info.m_ids_offset[i] - m_recv_info.m_ids_offset[0];
93 Integer size = m_recv_info.m_ids_offset[i + 1] - off;
94#ifdef USE_SYCL_USM
95 ValueT* ptr = m_rbuffer + off;
96#else
97 ValueT* ptr = m_rbuffer.data() + off;
98#endif
99 Integer rank = m_recv_info.m_ranks[i];
100
101 m_recv_request[i] =
102 Arccore::MessagePassing::mpReceive(m_parallel_mng,
103 ArrayView<ValueT>(size, ptr),
104 rank,
105 false);
106 }
107 }
108 if (m_send_policy == Alien::SimpleCSRInternal::CommProperty::ASynch)
109 m_send_request.resize(m_send_info.m_num_neighbours);
110 if (m_send_info.m_ids.size()) {
111 std::size_t total_nb_send_ids = m_send_info.m_ids_offset[m_send_info.m_num_neighbours] - m_send_info.m_ids_offset[0];
112#ifdef USE_SYCL_USM
113 m_sbuffer = sycl::malloc_shared<ValueT>(total_nb_send_ids, queue);
114#else
115 m_sbuffer.resize(total_nb_send_ids);
116#endif
117 {
118#ifdef USE_SYCL_USM
119 sycl::buffer<ValueType> sbuffer{ { sycl::buffer_allocation::empty_view(m_sbuffer, queue.get_device()) }, total_nb_send_ids };
120#else
121 sycl::buffer<ValueType> sbuffer(m_sbuffer.data(), total_nb_send_ids);
122#endif
123 // clang-format off
124 queue.submit([&](sycl::handler& cgh)
125 {
126 auto access_send_buffer = m_send_buffer.template get_access<sycl::access::mode::read>(cgh);
127 auto access_ids = m_send_ids.template get_access<sycl::access::mode::read>(cgh);
128 auto access_sbuffer = sycl::accessor(sbuffer, cgh, sycl::write_only, sycl::property::no_init{});
129
130 cgh.parallel_for<class vector_mult_send>(sycl::range<1>{total_threads},
131 [=] (sycl::item<1> itemId)
132 {
133 auto id = itemId.get_id(0);
134 for( auto i=id; i<total_nb_send_ids; i+=total_threads)
135 access_sbuffer[i] = access_send_buffer[access_ids[i]];
136 });
137 });
138 // clang-format on
139 }
140 //for (Integer i = 0; i < total_nb_send_ids; ++i) {
141 // alien_debug([&] {cout() << "MPI SEND["<<i<<"]="<<m_sbuffer[i];}) ;;
142 //}
143 }
144 for (Integer i = 0; i < m_send_info.m_num_neighbours; ++i) {
145 Integer off = m_send_info.m_ids_offset[i];
146 Integer nb_send_ids = m_send_info.m_ids_offset[i + 1] - off;
147#ifdef USE_SYCL_USM
148 ValueT const* ptr = m_sbuffer + off;
149#else
150 ValueT const* ptr = m_sbuffer.data() + off;
151#endif
152 Integer rank = m_send_info.m_ranks[i];
153 if (m_send_policy == Alien::SimpleCSRInternal::CommProperty::ASynch)
154 m_send_request[i] = Arccore::MessagePassing::mpSend(m_parallel_mng,
155 ConstArrayView<ValueT>(nb_send_ids, ptr), rank, false);
156 else
157 Arccore::MessagePassing::mpSend(m_parallel_mng,
158 ConstArrayView<ValueT>(nb_send_ids, ptr), rank);
159 }
160
161 //alien_debug([&] {cout()<<"END SYCLSendRecvOP START" ; });
162 //Universe().traceMng()->flush() ;
163 }
164
165 void end([[maybe_unused]] bool insitu = false)
166 {
167 //alien_debug([&] {cout() << "SYCLSendRecvOP END : "<<m_recv_policy;});
168 //Universe().traceMng()->flush() ;
169
170 // clang-format off
171 auto env = SYCLEnv::instance();
172 auto& queue = env->internal()->queue();
173 auto total_threads = env->maxNumThreads() ;
174 // clang-format on
175 if (m_recv_policy == Alien::SimpleCSRInternal::CommProperty::ASynch) {
176 Arccore::MessagePassing::mpWaitAll(m_parallel_mng, m_recv_request);
177
178 //Arccore::Integer total_recv_ids = m_recv_info.m_ids_offset[m_recv_info.m_num_neighbours] - m_recv_info.m_ids_offset[0];
179 //for (Integer i = 0; i < total_recv_ids; ++i) {
180 // alien_debug([&] {cout() << "MPI RECV["<<i<<"]="<<m_rbuffer[i];});
181 //}
182 }
183 else {
184 if (m_recv_info.m_ids.size()) {
185 Arccore::Integer total_recv_ids = m_recv_info.m_ids_offset[m_recv_info.m_num_neighbours] - m_recv_info.m_ids_offset[0];
186#ifdef USE_SYCL_USM
187 m_rbuffer = sycl::malloc_shared<ValueT>(total_recv_ids, queue);
188#else
189 m_rbuffer.resize(total_recv_ids);
190#endif
191 }
192 for (Integer i = 0; i < m_recv_info.m_num_neighbours; ++i) {
193 Integer off = m_recv_info.m_ids_offset[i];
194 Integer size = m_recv_info.m_ids_offset[i + 1] - off;
195#ifdef USE_SYCL_USM
196 ValueT* ptr = m_rbuffer + off;
197#else
198 ValueT* ptr = m_rbuffer.data() + off;
199#endif
200 Integer rank = m_recv_info.m_ranks[i];
201 Arccore::MessagePassing::mpReceive(m_parallel_mng, ArrayView<ValueT>(size, ptr), rank);
202 }
203 }
204 if (m_recv_info.m_ids.size()) {
205 std::size_t total_nb_recv_ids = m_recv_info.m_ids_offset[m_recv_info.m_num_neighbours] - m_recv_info.m_ids_offset[0];
206
207 {
208#ifdef USE_SYCL_USM
209 sycl::buffer<ValueType> rbuffer{ { sycl::buffer_allocation::view(m_rbuffer, queue.get_device()) }, total_nb_recv_ids };
210#else
211 sycl::buffer<ValueType> rbuffer(m_rbuffer.data(), total_nb_recv_ids);
212#endif
213 // clang-format off
214 queue.submit([&](sycl::handler& cgh)
215 {
216 //auto access_recv_buffer = m_recv_buffer.template get_access<sycl::access::mode::read_write>(cgh);
217 //auto access_ids = m_recv_ids.template get_access<sycl::access::mode::read>(cgh);
218#ifdef USE_HIPSYCL
219 sycl::accessor<ValueType> access_recv_buffer{m_recv_buffer, cgh, sycl::write_only, sycl::property::no_init{}};
220#endif
221#ifdef USE_ONEAPI
222 auto access_recv_buffer = sycl::accessor(m_recv_buffer, cgh, sycl::write_only, sycl::property::no_init{});
223#endif
224#ifdef USE_ACPPSYCL
225 auto access_recv_buffer = sycl::accessor(m_recv_buffer, cgh, sycl::write_only, sycl::property::no_init{});
226#endif
227
228
229 sycl::accessor<ValueType> access_rbuffer{rbuffer, cgh};
230
231 cgh.parallel_for<class vector_mult_recv>(sycl::range<1>{total_threads},
232 [=] (sycl::item<1> itemId)
233 {
234 auto id = itemId.get_id(0);
235 for(auto i=id;i<total_nb_recv_ids;i += total_threads)
236 access_recv_buffer[i] = access_rbuffer[i];
237 });
238 });
239 // clang-format on
240 }
241 }
242 if (m_send_policy == Alien::SimpleCSRInternal::CommProperty::ASynch) {
243 Arccore::MessagePassing::mpWaitAll(m_parallel_mng, m_send_request);
244 }
245
246 //alien_debug([&] {cout() << "AFTER SYCLSendRecvOP END : "<<m_recv_policy;});
247 //Universe().traceMng()->flush() ;
248 }
249
250 void upperRecv([[maybe_unused]] bool insitu = true)
251 {
252 }
253
254 void upperSend()
255 {
256 }
257
258 void lowerRecv([[maybe_unused]] bool insitu = true)
259 {
260 }
261
262 void lowerSend()
263 {
264
265 }
266 private:
267 // clang-format off
268 ValueBufferType& m_send_buffer;
269 const Alien::SimpleCSRInternal::CommInfo& m_send_info;
270 IndexBufferType& m_send_ids;
271 Alien::SimpleCSRInternal::CommProperty::ePolicyType m_send_policy;
272
273 ValueBufferType& m_recv_buffer;
274 const Alien::SimpleCSRInternal::CommInfo& m_recv_info;
275 IndexBufferType& m_recv_ids;
276 Alien::SimpleCSRInternal::CommProperty::ePolicyType m_recv_policy;
277#ifdef USE_SYCL_USM
278 ValueT* m_rbuffer = nullptr;
279 ValueT* m_sbuffer = nullptr;
280#else
281 std::vector<ValueT> m_rbuffer ;
282 std::vector<ValueT> m_sbuffer ;
283#endif
284 Arccore::Integer m_unknowns_num = 0;
285 Arccore::MessagePassing::IMessagePassingMng* m_parallel_mng = nullptr;
286 Arccore::ITraceMng* m_trace = nullptr;
287 Arccore::UniqueArray<Arccore::MessagePassing::Request> m_recv_request;
288 Arccore::UniqueArray<Arccore::MessagePassing::Request> m_send_request;
289 // clang-format on
290};
291
292} // namespace Alien::SYCLInternal