Arcane  v3.14.10.0
Documentation utilisateur
Chargement...
Recherche...
Aucune correspondance
MemoryCopier.cc
1// -*- tab-width: 2; indent-tabs-mode: nil; coding: utf-8-with-signature -*-
2//-----------------------------------------------------------------------------
3// Copyright 2000-2024 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/* MemoryCopier.cc (C) 2000-2024 */
9/* */
10/* Fonctions diverses de copie mémoire. */
11/*---------------------------------------------------------------------------*/
12/*---------------------------------------------------------------------------*/
13
14#include "arcane/accelerator/AcceleratorGlobal.h"
15
16#include "arcane/utils/Ref.h"
17#include "arcane/utils/FixedArray.h"
18#include "arcane/utils/NotSupportedException.h"
19#include "arcane/utils/internal/SpecificMemoryCopyList.h"
20
21#include "arcane/accelerator/core/RunQueue.h"
23
24/*---------------------------------------------------------------------------*/
25/*---------------------------------------------------------------------------*/
26
27namespace Arcane::Accelerator::impl
28{
29using IndexedMemoryCopyArgs = Arcane::impl::IndexedMemoryCopyArgs;
30using IndexedMultiMemoryCopyArgs = Arcane::impl::IndexedMultiMemoryCopyArgs;
31
32/*---------------------------------------------------------------------------*/
33/*---------------------------------------------------------------------------*/
34
35template <typename DataType, typename Extent>
37: public Arcane::impl::SpecificMemoryCopyBase<DataType, Extent>
38{
39 using BaseClass = Arcane::impl::SpecificMemoryCopyBase<DataType, Extent>;
40 using BaseClass::_toTrueType;
41
42 public:
43
44 using BaseClass::m_extent;
45
46 public:
47
48 void copyFrom(const IndexedMemoryCopyArgs& args) override
49 {
50 _copyFrom(args.m_queue, args.m_indexes, _toTrueType(args.m_source), _toTrueType(args.m_destination));
51 }
52
53 void copyTo(const IndexedMemoryCopyArgs& args) override
54 {
55 _copyTo(args.m_queue, args.m_indexes, _toTrueType(args.m_source), _toTrueType(args.m_destination));
56 }
57
58 void fill(const IndexedMemoryCopyArgs& args) override
59 {
60 _fill(args.m_queue, args.m_indexes, _toTrueType(args.m_source), _toTrueType(args.m_destination));
61 }
62
63 void copyFrom(const IndexedMultiMemoryCopyArgs& args) override
64 {
65 _copyFrom(args.m_queue, args.m_indexes, args.m_multi_memory, _toTrueType(args.m_source_buffer));
66 }
67
68 void copyTo(const IndexedMultiMemoryCopyArgs& args) override
69 {
70 _copyTo(args.m_queue, args.m_indexes, args.m_const_multi_memory, _toTrueType(args.m_destination_buffer));
71 }
72
73 void fill(const IndexedMultiMemoryCopyArgs& args) override
74 {
75 _fill(args.m_queue, args.m_indexes, args.m_multi_memory, _toTrueType(args.m_source_buffer));
76 }
77
78 public:
79
80 void _copyFrom(RunQueue* queue, SmallSpan<const Int32> indexes,
81 Span<const DataType> source, Span<DataType> destination)
82 {
84
85 ARCANE_CHECK_ACCESSIBLE_POINTER(queue, indexes.data());
87 ARCANE_CHECK_ACCESSIBLE_POINTER(queue, destination.data());
88
89 Int32 nb_index = indexes.size();
90 const Int64 sub_size = m_extent.v;
91
92 auto command = makeCommand(queue);
93 command << RUNCOMMAND_LOOP1(iter, nb_index)
94 {
95 auto [i] = iter();
96 Int64 zindex = i * sub_size;
97 Int64 zci = indexes[i] * sub_size;
98 for (Int32 z = 0; z < sub_size; ++z)
99 destination[zindex + z] = source[zci + z];
100 };
101 }
102
103 void _copyFrom(RunQueue* queue, SmallSpan<const Int32> indexes, SmallSpan<Span<std::byte>> multi_views,
105 {
107
108 if (arcaneIsCheck()) {
111 ARCANE_CHECK_ACCESSIBLE_POINTER_ALWAYS(queue, multi_views.data());
112 // Idéalement il faudrait tester les valeurs des éléments de multi_views
113 // mais si on fait cela on peut potentiellement faire des transferts
114 // entre l'accélérateur et le CPU.
115 }
116 const Int32 nb_index = indexes.size() / 2;
117 // On devrait pouvoir utiliser 'm_extent.v' mais avec CUDA 12.1 cela génère
118 // une erreur lors de l'exécution: error 98 : invalid device function
119 const Int32 sub_size = m_extent.v;
120
121 auto command = makeCommand(queue);
122 command << RUNCOMMAND_LOOP1(iter, nb_index)
123 {
124 auto [i] = iter();
125 Int32 index0 = indexes[i * 2];
126 Int32 index1 = indexes[(i * 2) + 1];
127 Span<std::byte> orig_view_bytes = multi_views[index0];
128 auto* orig_view_data = reinterpret_cast<DataType*>(orig_view_bytes.data());
129 // Utilise un span pour tester les débordements de tableau mais on
130 // pourrait directement utiliser 'orig_view_data' pour plus de performances
131 Span<DataType> orig_view = { orig_view_data, orig_view_bytes.size() / (Int64)sizeof(DataType) };
132 Int64 zci = ((Int64)(index1)) * sub_size;
133 Int64 z_index = (Int64)i * sub_size;
134 for (Int32 z = 0, n = sub_size; z < n; ++z)
135 orig_view[zci + z] = source[z_index + z];
136 };
137 }
138
139 /*!
140 * \brief Remplit les valeurs d'indices spécifiés par \a indexes.
141 *
142 * Si \a indexes est vide, remplit toutes les valeurs.
143 */
145 Span<DataType> destination)
146 {
148
149 ARCANE_CHECK_ACCESSIBLE_POINTER(queue, indexes.data());
150 ARCANE_CHECK_ACCESSIBLE_POINTER(queue, destination.data());
152
153 Int32 nb_index = indexes.size();
154 const Int32 sub_size = m_extent.v;
155 constexpr Int32 max_size = 24;
156
157 // Pour l'instant on limite la taille de DataType en dur.
158 // A terme, il faudrait allouer sur le device et désallouer en fin
159 // d'exécution (via cudaMallocAsync/cudaFreeAsync pour gérer l'asynchronisme)
160 if (sub_size > max_size)
161 ARCANE_THROW(NotSupportedException, "sizeof(type) is too big (v={0} max={1})",
162 sizeof(DataType) * sub_size, sizeof(DataType) * max_size);
164 for (Int32 z = 0; z < sub_size; ++z)
165 local_source[z] = source[z];
166 for (Int32 z = sub_size; z < max_size; ++z)
167 local_source[z] = {};
168
169 auto command = makeCommand(queue);
170 // Si \a nb_index vaut 0, on remplit tous les éléments
171 if (nb_index == 0) {
172 Int32 nb_value = CheckedConvert::toInt32(destination.size() / sub_size);
173 command << RUNCOMMAND_LOOP1(iter, nb_value)
174 {
175 auto [i] = iter();
176 Int64 zci = i * sub_size;
177 for (Int32 z = 0; z < sub_size; ++z)
178 destination[zci + z] = local_source[z];
179 };
180 }
181 else {
182 command << RUNCOMMAND_LOOP1(iter, nb_index)
183 {
184 auto [i] = iter();
185 Int64 zci = indexes[i] * sub_size;
186 for (Int32 z = 0; z < sub_size; ++z)
187 destination[zci + z] = local_source[z];
188 };
189 }
190 }
191
192 void _fill(RunQueue* queue, SmallSpan<const Int32> indexes, SmallSpan<Span<std::byte>> multi_views,
194 {
196
197 if (arcaneIsCheck()) {
200 ARCANE_CHECK_ACCESSIBLE_POINTER_ALWAYS(queue, multi_views.data());
201 // Idéalement il faudrait tester les valeurs des éléments de multi_views
202 // mais si on fait cela on peut potentiellement faire des transferts
203 // entre l'accélérateur et le CPU.
204 }
205 const Int32 nb_index = indexes.size() / 2;
206 // On devrait pouvoir utiliser 'm_extent.v' mais avec CUDA 12.1 cela génère
207 // une erreur lors de l'exécution: error 98 : invalid device function
208 const Int32 sub_size = m_extent.v;
209 constexpr Int32 max_size = 24;
210
211 // Pour l'instant on limite la taille de DataType en dur.
212 // A terme, il faudrait allouer sur le device et désallouer en fin
213 // d'exécution (via cudaMallocAsync/cudaFreeAsync pour gérer l'asynchronisme)
214 if (sub_size > max_size)
215 ARCANE_THROW(NotSupportedException, "sizeof(type) is too big (v={0} max={1})",
216 sizeof(DataType) * sub_size, sizeof(DataType) * max_size);
218 for (Int32 z = 0; z < sub_size; ++z)
219 local_source[z] = source[z];
220 for (Int32 z = sub_size; z < max_size; ++z)
221 local_source[z] = {};
222
223 if (nb_index == 0) {
224 // Remplit toutes les valeurs du tableau avec la source.
225 // Comme le nombre d'éléments de la deuxième dimension dépend de la première,
226 // on utilise un noyau par dimension.
227 RunQueue::ScopedAsync sc(queue);
228 const Int32 nb_dim1 = multi_views.size();
229 for (Int32 zz = 0; zz < nb_dim1; ++zz) {
230 Span<DataType> orig_view = Arccore::asSpan<DataType>(multi_views[zz]);
231 Int32 nb_value = CheckedConvert::toInt32(orig_view.size() / sub_size);
232 auto command = makeCommand(queue);
233 command << RUNCOMMAND_LOOP1(iter, nb_value)
234 {
235 auto [i] = iter();
236 orig_view[i] = local_source[i % sub_size];
237 };
238 }
239 }
240 else {
241 auto command = makeCommand(queue);
242 command << RUNCOMMAND_LOOP1(iter, nb_index)
243 {
244 auto [i] = iter();
245 Int32 index0 = indexes[i * 2];
246 Int32 index1 = indexes[(i * 2) + 1];
247 Span<std::byte> orig_view_bytes = multi_views[index0];
248 auto* orig_view_data = reinterpret_cast<DataType*>(orig_view_bytes.data());
249 // Utilise un span pour tester les débordements de tableau mais on
250 // pourrait directement utiliser 'orig_view_data' pour plus de performances
251 Span<DataType> orig_view = { orig_view_data, orig_view_bytes.size() / (Int64)sizeof(DataType) };
252 Int64 zci = ((Int64)(index1)) * sub_size;
253 for (Int32 z = 0, n = sub_size; z < n; ++z)
254 orig_view[zci + z] = local_source[z];
255 };
256 }
257 }
258
259 void _copyTo(RunQueue* queue, SmallSpan<const Int32> indexes, Span<const DataType> source,
260 Span<DataType> destination)
261 {
263
264 ARCANE_CHECK_ACCESSIBLE_POINTER(queue, indexes.data());
265 ARCANE_CHECK_ACCESSIBLE_POINTER(queue, source.data());
266 ARCANE_CHECK_ACCESSIBLE_POINTER(queue, destination.data());
267
268 Int32 nb_index = indexes.size();
269 const Int64 sub_size = m_extent.v;
270
271 auto command = makeCommand(queue);
272 command << RUNCOMMAND_LOOP1(iter, nb_index)
273 {
274 auto [i] = iter();
275 Int64 zindex = i * sub_size;
276 Int64 zci = indexes[i] * sub_size;
277 for (Int32 z = 0; z < sub_size; ++z)
278 destination[zci + z] = source[zindex + z];
279 };
280 }
281
282 void _copyTo(RunQueue* queue, SmallSpan<const Int32> indexes, SmallSpan<const Span<const std::byte>> multi_views,
283 Span<DataType> destination)
284 {
286
287 if (arcaneIsCheck()) {
288 ARCANE_CHECK_ACCESSIBLE_POINTER_ALWAYS(queue, indexes.data());
289 ARCANE_CHECK_ACCESSIBLE_POINTER_ALWAYS(queue, destination.data());
290 ARCANE_CHECK_ACCESSIBLE_POINTER_ALWAYS(queue, multi_views.data());
291 // Idéalement il faudrait tester les valeurs des éléments de multi_views
292 // mais si on fait cela on peut potentiellement faire des transferts
293 // entre l'accélérateur et le CPU.
294 }
295
296 const Int32 nb_index = indexes.size() / 2;
297 // On devrait pouvoir utiliser 'm_extent.v' mais avec CUDA 12.1 cela génère
298 // une erreur lors de l'exécution: error 98 : invalid device function
299 const Int32 sub_size = m_extent.v;
300
301 auto command = makeCommand(queue);
302 command << RUNCOMMAND_LOOP1(iter, nb_index)
303 {
304 auto [i] = iter();
305 Int32 index0 = indexes[i * 2];
306 Int32 index1 = indexes[(i * 2) + 1];
307 Span<const std::byte> orig_view_bytes = multi_views[index0];
308 auto* orig_view_data = reinterpret_cast<const DataType*>(orig_view_bytes.data());
309 // Utilise un span pour tester les débordements de tableau mais on
310 // pourrait directement utiliser 'orig_view_data' pour plus de performances
311 Span<const DataType> orig_view = { orig_view_data, orig_view_bytes.size() / (Int64)sizeof(DataType) };
312 Int64 zci = ((Int64)(index1)) * sub_size;
313 Int64 z_index = (Int64)i * sub_size;
314 for (Int32 z = 0, n = sub_size; z < n; ++z)
315 destination[z_index + z] = orig_view[zci + z];
316 };
317 }
318};
319
320/*---------------------------------------------------------------------------*/
321/*---------------------------------------------------------------------------*/
322
324{
325 public:
326
327 using InterfaceType = Arcane::impl::ISpecificMemoryCopy;
328 template <typename DataType, typename Extent> using SpecificType = AcceleratorSpecificMemoryCopy<DataType, Extent>;
329 using RefType = Arcane::impl::SpecificMemoryCopyRef<AcceleratorIndexedCopyTraits>;
330};
331
332/*---------------------------------------------------------------------------*/
333/*---------------------------------------------------------------------------*/
334
336{
337 public:
338
340 {
341 Arcane::impl::ISpecificMemoryCopyList::setDefaultCopyListIfNotSet(&m_copy_list);
342 }
343 Arcane::impl::SpecificMemoryCopyList<AcceleratorIndexedCopyTraits> m_copy_list;
344};
345
346namespace
347{
348 AcceleratorSpecificMemoryCopyList global_copy_list;
349}
350
351/*---------------------------------------------------------------------------*/
352/*---------------------------------------------------------------------------*/
353
354} // namespace Arcane::Accelerator::impl
355
356/*---------------------------------------------------------------------------*/
357/*---------------------------------------------------------------------------*/
#define ARCANE_CHECK_ACCESSIBLE_POINTER(queue_or_runner_or_policy, ptr)
Macro qui vérifie en mode check si ptr est accessible pour une RunQueue ou un Runner.
#define ARCANE_CHECK_ACCESSIBLE_POINTER_ALWAYS(queue_or_runner_or_policy, ptr)
Macro qui vérifie si ptr est accessible pour une RunQueue ou un Runner.
#define ARCANE_CHECK_POINTER(ptr)
Macro retournant le pointeur ptr s'il est non nul ou lancant une exception s'il est nul.
#define ARCANE_THROW(exception_class,...)
Macro pour envoyer une exception avec formattage.
Types et macros pour gérer les boucles sur les accélérateurs.
#define RUNCOMMAND_LOOP1(iter_name, x1,...)
Boucle sur accélérateur avec arguments supplémentaires pour les réductions.
Permet de modifier l'asynchronisme de la file pendant la durée de vie de l'instance.
File d'exécution pour un accélérateur.
void _fill(RunQueue *queue, SmallSpan< const Int32 > indexes, Span< const DataType > source, Span< DataType > destination)
Remplit les valeurs d'indices spécifiés par indexes.
Tableau 1D de taille fixe.
Definition FixedArray.h:45
Vue d'un tableau d'éléments de type T.
Definition Span.h:670
constexpr __host__ __device__ SizeType size() const noexcept
Retourne la taille du tableau.
Definition Span.h:209
constexpr __host__ __device__ pointer data() const noexcept
Pointeur sur le début de la vue.
Definition Span.h:419
Vue d'un tableau d'éléments de type T.
Definition Span.h:510
RunCommand makeCommand(const RunQueue &run_queue)
Créé une commande associée à la file run_queue.
@ Sequential
Politique d'exécution séquentielle.
Int32 toInt32(Int64 v)
Converti un Int64 en un Int32.
bool arcaneIsCheck()
Vrai si on est en mode vérification.
Definition Misc.cc:151
std::int64_t Int64
Type entier signé sur 64 bits.
std::int32_t Int32
Type entier signé sur 32 bits.