Arcane  4.1.12.0
Developer documentation
Loading...
Searching...
No Matches
IncrementalComponentModifier_Accelerator.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/* IncrementalComponentModifier_Accelerator.cc (C) 2000-2024 */
9/* */
10/* Incremental modification of constituents. */
11/*---------------------------------------------------------------------------*/
12/*---------------------------------------------------------------------------*/
13
14#include "arcane/materials/internal/IncrementalComponentModifier.h"
15
16#include "arcane/utils/ITraceMng.h"
17#include "arcane/utils/FunctorUtils.h"
18
19#include "arcane/core/internal/ItemGroupImplInternal.h"
20#include "arcane/core/materials/IMeshMaterialVariable.h"
21
22#include "arcane/materials/internal/MeshMaterialMng.h"
23#include "arcane/materials/internal/MaterialModifierOperation.h"
24#include "arcane/materials/internal/ConstituentConnectivityList.h"
25#include "arcane/materials/internal/AllEnvData.h"
26
27#include "arcane/accelerator/RunCommandLoop.h"
28#include "arcane/accelerator/Filter.h"
29#include "arcane/accelerator/Reduce.h"
30#include "arcane/accelerator/core/ProfileRegion.h"
31
32/*---------------------------------------------------------------------------*/
33/*---------------------------------------------------------------------------*/
34
35namespace Arcane::Materials
36{
37
38/*---------------------------------------------------------------------------*/
39/*---------------------------------------------------------------------------*/
40
41void IncrementalComponentModifier::
42apply(MaterialModifierOperation* operation)
43{
44 const char* str_add = "ApplyConstituentOperationAdd";
45 const char* str_remove = "ApplyConstituentOperationRemove";
46 bool is_add = operation->isAdd();
47 Int32 color = (is_add) ? 0x00FFFF : 0x007FFF;
48 Accelerator::ProfileRegion ps(m_queue, is_add ? str_add : str_remove, color);
49
50 IMeshMaterial* mat = operation->material();
51 SmallSpan<const Int32> orig_ids = operation->ids();
52 SmallSpan<const Int32> ids = orig_ids;
53
54 auto* true_mat = ARCANE_CHECK_POINTER(dynamic_cast<MeshMaterial*>(mat));
55
56 const IMeshEnvironment* env = mat->environment();
57 MeshEnvironment* true_env = true_mat->trueEnvironment();
58 const Integer nb_mat = env->nbMaterial();
59
60 info(4) << "-- ** -- Using optimisation updateMaterialDirect is_add=" << is_add
61 << " mat=" << mat->name() << " nb_item=" << orig_ids.size()
62 << " mat_id=" << mat->id() << " env_id=" << env->id();
63
64 ConstituentConnectivityList* connectivity = m_all_env_data->componentConnectivityList();
65 const bool check_if_present = !m_queue.isAcceleratorPolicy();
66
67 const bool is_device = m_queue.isAcceleratorPolicy();
68
69 // Fills the arrays indicating if a constituent is affected by
70 // the current modification. If not, we can avoid testing it
71 // in the constituent loop.
72 {
73 m_work_info.m_is_materials_modified.fillHost(false);
74 m_work_info.m_is_environments_modified.fillHost(false);
75 m_work_info.m_is_materials_modified.sync(is_device);
76 m_work_info.m_is_environments_modified.sync(is_device);
77
78 {
79 auto mat_modifier = m_work_info.m_is_materials_modified.modifier(is_device);
80 auto env_modifier = m_work_info.m_is_environments_modified.modifier(is_device);
81 connectivity->fillModifiedConstituents(orig_ids, mat_modifier.view(), env_modifier.view(), mat->id(), is_add, m_queue);
82 if (m_is_debug)
83 connectivity->printConstituents(orig_ids);
84 }
85 {
86 auto is_mat_modified = m_work_info.m_is_materials_modified.view(false);
87 auto is_env_modified = m_work_info.m_is_environments_modified.view(false);
88 info(4) << "ModifiedInfosAfter: mats=" << is_mat_modified << " envs=" << is_env_modified;
89 }
90 }
91
92 if (nb_mat != 1) {
93
94 // If it is possible to have multiple materials per environment, it must be handled
95 // for each cell if the environment changes following the addition/removal of material.
96 // The two cases are:
97 // - in case of addition, the environment changes for a cell if there was no
98 // material before. In this case, the environment is added to the cell.
99 // - in case of removal, the environment changes in the cell if there was
100 // only 1 material before. In this case, the environment is removed from the cell.
101
102 UniqueArray<Int32>& cells_changed_in_env = m_work_info.cells_changed_in_env;
103 UniqueArray<Int32>& cells_unchanged_in_env = m_work_info.cells_unchanged_in_env;
104 UniqueArray<Int16>& cells_current_nb_material = m_work_info.m_cells_current_nb_material;
105 const Int32 nb_id = ids.size();
106 cells_unchanged_in_env.resize(nb_id);
107 cells_changed_in_env.resize(nb_id);
108 cells_current_nb_material.resize(nb_id);
109 const Int32 ref_nb_mat = is_add ? 0 : 1;
110 const Int16 env_id = true_env->componentId();
111 info(4) << "Using optimisation updateMaterialDirect is_add?=" << is_add;
112
113 connectivity->fillCellsNbMaterial(ids, env_id, cells_current_nb_material.view(), m_queue);
114
115 {
116 Accelerator::GenericFilterer filterer(m_queue);
117 SmallSpan<Int32> cells_unchanged_in_env_view = cells_unchanged_in_env.view();
118 SmallSpan<Int32> cells_changed_in_env_view = cells_changed_in_env.view();
119 SmallSpan<const Int16> cells_current_nb_material_view = m_work_info.m_cells_current_nb_material.view();
120 {
121 auto select_lambda = [=] ARCCORE_HOST_DEVICE(Int32 index) -> bool {
122 Int16 current_cell_nb_mat = cells_current_nb_material_view[index];
123 return current_cell_nb_mat != ref_nb_mat;
124 };
125 auto setter_lambda = [=] ARCCORE_HOST_DEVICE(Int32 input_index, Int32 output_index) {
126 cells_unchanged_in_env_view[output_index] = ids[input_index];
127 };
128 filterer.applyWithIndex(nb_id, select_lambda, setter_lambda, A_FUNCINFO);
129 cells_unchanged_in_env.resize(filterer.nbOutputElement());
130 }
131 {
132 auto select_lambda = [=] ARCCORE_HOST_DEVICE(Int32 index) -> bool {
133 Int16 current_cell_nb_mat = cells_current_nb_material_view[index];
134 return current_cell_nb_mat == ref_nb_mat;
135 };
136 auto setter_lambda = [=] ARCCORE_HOST_DEVICE(Int32 input_index, Int32 output_index) {
137 cells_changed_in_env_view[output_index] = ids[input_index];
138 };
139 filterer.applyWithIndex(nb_id, select_lambda, setter_lambda, A_FUNCINFO);
140 cells_changed_in_env.resize(filterer.nbOutputElement());
141 }
142 }
143
144 Integer nb_unchanged_in_env = cells_unchanged_in_env.size();
145 info(4) << "Cells unchanged in environment n=" << nb_unchanged_in_env;
146
147 Int16 mat_id = true_mat->componentId();
148 if (is_add) {
149 mat->cells().addItems(cells_unchanged_in_env, check_if_present);
150 connectivity->addCellsToMaterial(mat_id, cells_unchanged_in_env.view(), m_queue);
151 _addItemsToEnvironment(true_env, true_mat, cells_unchanged_in_env.view(), false);
152 }
153 else {
154 flagRemovedCells(cells_unchanged_in_env, true);
155 _removeItemsInGroup(mat->cells(), cells_unchanged_in_env);
156 connectivity->removeCellsToMaterial(mat_id, cells_unchanged_in_env.view(), m_queue);
157 _removeItemsFromEnvironment(true_env, true_mat, cells_unchanged_in_env.view(), false);
158 flagRemovedCells(cells_unchanged_in_env, false);
159 }
160
161 // Takes for \a ids only the list of cells
162 // that did not yet belong to the environment in which we
163 // are adding the material.
164 ids = cells_changed_in_env.view();
165 }
166
167 // Updates the number of environments and materials for each cell.
168 // NOTE: the operation must first be performed on the environments before
169 // the materials.
170 {
171 Int16 env_id = true_env->componentId();
172 Int16 mat_id = true_mat->componentId();
173 if (is_add) {
174 connectivity->addCellsToEnvironment(env_id, ids, m_queue);
175 connectivity->addCellsToMaterial(mat_id, ids, m_queue);
176 }
177 else {
178 connectivity->removeCellsToEnvironment(env_id, ids, m_queue);
179 connectivity->removeCellsToMaterial(mat_id, ids, m_queue);
180 }
181 }
182
183 // Since we have added/removed material cells in the environment,
184 // we must transform pure cells into partial cells (in case
185 // of addition) or partial cells into pure cells (in case of
186 // removal).
187 info(4) << "Transform PartialPure for material name=" << true_mat->name();
188 _switchCellsForMaterials(true_mat, orig_ids);
189 info(4) << "Transform PartialPure for environment name=" << env->name();
190 _switchCellsForEnvironments(env, orig_ids);
191
192 // If I am mono-mat, then mat->cells()<=>env->cells() and only one
193 // of the two groups needs to be updated.
194 bool need_update_env = (nb_mat != 1);
195
196 if (is_add) {
197 mat->cells().addItems(ids.smallView(), check_if_present);
198 if (need_update_env)
199 env->cells().addItems(ids.smallView(), check_if_present);
200 _addItemsToEnvironment(true_env, true_mat, ids, need_update_env);
201 }
202 else {
203 flagRemovedCells(ids, true);
204 _removeItemsInGroup(mat->cells(), ids);
205 if (need_update_env)
206 _removeItemsInGroup(env->cells(), ids);
207 _removeItemsFromEnvironment(true_env, true_mat, ids, need_update_env);
208 // Reset \a removed_local_ids_filter to the initial value for subsequent operations
209 flagRemovedCells(ids, false);
210 }
211}
212
213/*---------------------------------------------------------------------------*/
214/*---------------------------------------------------------------------------*/
215
222{
223 ConstituentConnectivityList* connectivity = m_all_env_data->componentConnectivityList();
224 ConstArrayView<Int16> cells_nb_env = connectivity->cellsNbEnvironment();
225 const bool is_add = m_work_info.isAdd();
226 SmallSpan<bool> transformed_cells = m_work_info.transformedCells();
227
228 const Int32 n = ids.size();
229 auto command = makeCommand(m_queue);
230 Accelerator::ReducerSum2<Int32> sum_transformed(command);
231 command << RUNCOMMAND_LOOP1(iter, n, sum_transformed)
232 {
233 auto [i] = iter();
234 Int32 lid = ids[i];
235 bool do_transform = false;
236 // In case of addition, we switch from pure to partial if there are multiple environments.
237 // In case of removal, we switch from partial to pure if we are the only environment.
238 if (is_add)
239 do_transform = cells_nb_env[lid] > 1;
240 else
241 do_transform = cells_nb_env[lid] == 1;
242 if (do_transform) {
243 transformed_cells[lid] = do_transform;
244 sum_transformed.combine(1);
245 }
246 };
247 Int32 total_transformed = sum_transformed.reducedValue();
248 return total_transformed;
249}
250
251/*---------------------------------------------------------------------------*/
252/*---------------------------------------------------------------------------*/
253
254void IncrementalComponentModifier::
255_computeItemsToAdd(ComponentItemListBuilder& list_builder, SmallSpan<const Int32> local_ids)
256{
257 SmallSpan<const bool> cells_is_partial = m_work_info.m_cells_is_partial;
258
259 Accelerator::GenericFilterer filterer(m_queue);
260
261 MeshMaterialVariableIndexer* var_indexer = list_builder.indexer();
262
263 const Int32 nb_id = local_ids.size();
264
265 SmallSpan<Int32> pure_indexes = list_builder.pureIndexes();
266 SmallSpan<Int32> partial_indexes = list_builder.partialIndexes();
267 SmallSpan<Int32> partial_local_ids = list_builder.partialLocalIds();
268 Int32 nb_pure_added = 0;
269 Int32 nb_partial_added = 0;
270 Int32 index_in_partial = var_indexer->maxIndexInMultipleArray();
271
272 // TODO: for now we fill in two passes, but it would be possible to do it in one go using the Partition algorithm.
273 // We would then need to reverse the elements of the second list to have
274 // the same traversal order as before going to the accelerator.
275
276 // Fills the list of pure cells
277 {
278 auto select_lambda = [=] ARCCORE_HOST_DEVICE(Int32 index) -> bool {
279 return !cells_is_partial[index];
280 };
281 auto setter_lambda = [=] ARCCORE_HOST_DEVICE(Int32 input_index, Int32 output_index) {
282 Int32 local_id = local_ids[input_index];
283 pure_indexes[output_index] = local_id;
284 };
285 filterer.applyWithIndex(nb_id, select_lambda, setter_lambda, A_FUNCINFO);
286 nb_pure_added = filterer.nbOutputElement();
287 }
288 // Fills the list of partial cells
289 {
290 auto select_lambda = [=] ARCCORE_HOST_DEVICE(Int32 index) -> bool {
291 return cells_is_partial[index];
292 };
293 auto setter_lambda = [=] ARCCORE_HOST_DEVICE(Int32 input_index, Int32 output_index) {
294 Int32 local_id = local_ids[input_index];
295 partial_indexes[output_index] = index_in_partial + output_index;
296 partial_local_ids[output_index] = local_id;
297 };
298 filterer.applyWithIndex(nb_id, select_lambda, setter_lambda, A_FUNCINFO);
299 nb_partial_added = filterer.nbOutputElement();
300 }
301
302 list_builder.resize(nb_pure_added, nb_partial_added);
303}
304
305/*---------------------------------------------------------------------------*/
306/*---------------------------------------------------------------------------*/
307
308void IncrementalComponentModifier::
309flagRemovedCells(SmallSpan<const Int32> local_ids, bool value_to_set)
310{
311 const Int32 nb_item = local_ids.size();
312 SmallSpan<bool> removed_cells = m_work_info.removedCells();
313 auto command = makeCommand(m_queue);
314
315 ARCANE_CHECK_ACCESSIBLE_POINTER(m_queue, local_ids.data());
316 ARCANE_CHECK_ACCESSIBLE_POINTER(m_queue, removed_cells.data());
317
318 command << RUNCOMMAND_LOOP1(iter, nb_item)
319 {
320 auto [i] = iter();
321 removed_cells[local_ids[i]] = value_to_set;
322 };
323}
324
325/*---------------------------------------------------------------------------*/
326/*---------------------------------------------------------------------------*/
327
328void IncrementalComponentModifier::
329_resetTransformedCells(SmallSpan<const Int32> local_ids)
330{
331 const Int32 nb_item = local_ids.size();
332 auto command = makeCommand(m_queue);
333 SmallSpan<bool> transformed_cells = m_work_info.transformedCells();
334 command << RUNCOMMAND_LOOP1(iter, nb_item)
335 {
336 auto [i] = iter();
337 Int32 lid = local_ids[i];
338 transformed_cells[lid] = false;
339 };
340}
341
342/*---------------------------------------------------------------------------*/
343/*---------------------------------------------------------------------------*/
344
345void IncrementalComponentModifier::
346_removeItemsInGroup(ItemGroup cells, SmallSpan<const Int32> removed_ids)
347{
348 const Int32 nb_removed = removed_ids.size();
349 if (nb_removed == 0)
350 return;
351
352 const bool do_old = false;
353 if (do_old) {
354 cells.removeItems(removed_ids.smallView(), false);
355 }
356 else {
357 // Filters the entities of the group \a cells considering that
358 // m_work_info.removedCells() is true for the cells that
359 // must be deleted.
360 ItemGroupImplInternal* impl_internal = cells._internalApi();
361 SmallSpan<Int32> items_local_id(impl_internal->itemsLocalId());
362
363 // During the application of the filter, the input and output array
364 // is the same (this is normally supported by the GenericFilterer).
365 SmallSpan<Int32> input_ids(items_local_id);
366 SmallSpan<const bool> filtered_cells(m_work_info.removedCells());
367 Accelerator::GenericFilterer filterer(m_queue);
368 auto select_filter = [=] ARCCORE_HOST_DEVICE(Int32 local_id) -> bool {
369 return !filtered_cells[local_id];
370 };
371 filterer.applyIf(input_ids, select_filter, A_FUNCINFO);
372
373 Int32 current_nb_item = items_local_id.size();
374 Int32 nb_remaining = filterer.nbOutputElement();
375 if ((nb_remaining + nb_removed) != current_nb_item)
376 ARCANE_FATAL("Internal error in removing nb_remaining={0} nb_removed={1} original_size={2}",
377 nb_remaining, nb_removed, current_nb_item);
378 impl_internal->notifyDirectRemoveItems(removed_ids, nb_remaining);
379 }
380}
381
382/*---------------------------------------------------------------------------*/
383/*---------------------------------------------------------------------------*/
384
392{
394
395 auto output_indexes = args.m_local_ids;
396 auto input_indexes = args.m_indexes_in_multiple;
397 const bool is_global_to_partial = args.m_is_global_to_partial;
398
399 if (is_global_to_partial)
400 std::swap(output_indexes, input_indexes);
401 SmallSpan<const CopyBetweenDataInfo> host_copy_data(m_work_info.m_host_variables_copy_data);
402 SmallSpan<const CopyBetweenDataInfo> copy_data(m_work_info.m_variables_copy_data.to1DSmallSpan());
403 const Int32 nb_value = input_indexes.size();
404 if (nb_value != output_indexes.size())
405 ARCANE_FATAL("input_indexes ({0}) and output_indexes ({1}) are different", nb_value, output_indexes);
406
407 const Int32 nb_copy = copy_data.size();
408
409 for (Int32 i = 0; i < nb_copy; ++i) {
410 ARCANE_CHECK_ACCESSIBLE_POINTER(queue, host_copy_data[i].m_output.data());
411 ARCANE_CHECK_ACCESSIBLE_POINTER(queue, host_copy_data[i].m_input.data());
412 }
413 ARCANE_CHECK_ACCESSIBLE_POINTER(queue, input_indexes.data());
414 ARCANE_CHECK_ACCESSIBLE_POINTER(queue, output_indexes.data());
415
416 // TODO: Handle the copy in a way that allows using coalescence
417 // TODO: Make specializations if dim2_size is 4 or 8
418 // (or a multiple) to avoid the internal loop.
419 auto command = makeCommand(queue);
420 command << RUNCOMMAND_LOOP2(iter, nb_copy, nb_value)
421 {
422 auto [icopy, i] = iter();
423 auto input = copy_data[icopy].m_input;
424 auto output = copy_data[icopy].m_output;
425 Int32 dim2_size = copy_data[icopy].m_data_size;
426 Int32 output_base = output_indexes[i] * dim2_size;
427 Int32 input_base = input_indexes[i] * dim2_size;
428 for (Int32 j = 0; j < dim2_size; ++j)
429 output[output_base + j] = input[input_base + j];
430 };
431}
432
433/*---------------------------------------------------------------------------*/
434/*---------------------------------------------------------------------------*/
435
436void IncrementalComponentModifier::
437_applyInitializeWithZero(const InitializeWithZeroArgs& args)
438{
440
441 const RunQueue& queue = args.m_queue;
442 auto output_indexes = args.m_indexes_in_multiple;
443
445 SmallSpan<const CopyBetweenDataInfo> copy_data(m_work_info.m_variables_copy_data.to1DSmallSpan());
446 const Int32 nb_value = output_indexes.size();
447 const Int32 nb_copy = copy_data.size();
448
449 for (Int32 i = 0; i < nb_copy; ++i) {
450 ARCANE_CHECK_ACCESSIBLE_POINTER(queue, host_copy_data[i].m_output.data());
451 }
452 ARCANE_CHECK_ACCESSIBLE_POINTER(queue, output_indexes.data());
453
454 // TODO: Handle the copy in a way that allows using coalescence
455 // TODO: Make specializations if dim2_size is 4 or 8
456 // (or a multiple) to avoid the internal loop.
457 auto command = makeCommand(queue);
458 command << RUNCOMMAND_LOOP2(iter, nb_copy, nb_value)
459 {
460 auto [icopy, i] = iter();
461 auto output = copy_data[icopy].m_output;
462 Int32 dim2_size = copy_data[icopy].m_data_size;
463 Int32 output_base = output_indexes[i] * dim2_size;
464 for (Int32 j = 0; j < dim2_size; ++j)
465 output[output_base + j] = {};
466 };
467}
468
469/*---------------------------------------------------------------------------*/
470/*---------------------------------------------------------------------------*/
471
480{
481 SmallSpan<const CopyBetweenDataInfo> host_copy_data(m_work_info.m_host_variables_copy_data);
482 SmallSpan<const CopyBetweenDataInfo> copy_data(m_work_info.m_variables_copy_data.to1DSmallSpan());
483
484 const Int32 nb_copy = host_copy_data.size();
485 if (nb_copy == 0)
486 return;
487
488 // Assumes that all views have the same sizes.
489 // This is the case because the views are composed of 'ArrayView<>' and 'Array2View' and these
490 // two classes have the same size.
491 // TODO: it would be preferable to take the max of the sizes and in the command
492 // only perform the copy if we do not exceed the size.
493 Int32 nb_value = host_copy_data[0].m_input.size();
494
495 for (Int32 i = 0; i < nb_copy; ++i) {
496 const CopyBetweenDataInfo& h = host_copy_data[i];
497 ARCANE_CHECK_ACCESSIBLE_POINTER(queue, h.m_output.data());
498 ARCANE_CHECK_ACCESSIBLE_POINTER(queue, h.m_input.data());
499 if (h.m_input.size() != nb_value)
500 ARCANE_FATAL("Invalid nb_value '{0} i={1} expected={2}", h.m_input.size(), nb_value);
501 }
502
503 auto command = makeCommand(queue);
504 command << RUNCOMMAND_LOOP2(iter, nb_copy, nb_value)
505 {
506 auto [icopy, i] = iter();
507 auto input = copy_data[icopy].m_input;
508 auto output = copy_data[icopy].m_output;
509 output[i] = input[i];
510 };
511}
512
513/*---------------------------------------------------------------------------*/
514/*---------------------------------------------------------------------------*/
515
516} // End namespace Arcane::Materials
517
518/*---------------------------------------------------------------------------*/
519/*---------------------------------------------------------------------------*/
#define ARCANE_CHECK_POINTER(ptr)
Macro returning the pointer ptr if it is not null or throwing an exception if it is null.
#define ARCANE_FATAL(...)
Macro throwing a FatalErrorException.
#define RUNCOMMAND_LOOP2(iter_name, x1, x2)
2D loop on accelerator
#define RUNCOMMAND_LOOP1(iter_name, x1,...)
1D loop on accelerator with additional arguments.
Constant view of an array of type T.
Helper class for building a list of ComponentItems for a MeshMaterialVariableIndexer.
Management of constituent connectivity lists.
UniqueArray< CopyBetweenDataInfo > m_host_variables_copy_data
Information for copies between partial and global values.
NumArray< CopyBetweenDataInfo, MDDim1 > m_variables_copy_data
Information for copies between partial and global values.
Arguments for methods copying between partial and global values.
void _switchCellsForMaterials(const MeshMaterial *modified_mat, SmallSpan< const Int32 > ids)
Transforms entities for an environment.
Int32 _computeCellsToTransformForEnvironments(SmallSpan< const Int32 > ids)
Calculates the cells to transform when modifying the cells of an environment.
void _switchCellsForEnvironments(const IMeshEnvironment *modified_env, SmallSpan< const Int32 > ids)
Transforms entities for environments.
void _applyCopyVariableViews(RunQueue &queue)
Performs the copy of views for variables.
void _addItemsToEnvironment(MeshEnvironment *env, MeshMaterial *mat, SmallSpan< const Int32 > local_ids, bool update_env_indexer)
Adds the cells of an environment material.
void _applyCopyBetweenPartialsAndGlobals(const CopyBetweenPartialAndGlobalArgs &args, RunQueue &queue)
Performs the copy between partial and global values.
void _removeItemsFromEnvironment(MeshEnvironment *env, MeshMaterial *mat, SmallSpan< const Int32 > local_ids, bool update_env_indexer)
Removes cells of a material from the environment.
Arguments for methods copying between partial and global values.
Operation to add or remove cells from a material.
UniqueArray< CopyBetweenDataInfo > * m_copy_data
Copy information if only one command is used.
View of an array of elements of type T.
Definition Span.h:805
constexpr __host__ __device__ pointer data() const noexcept
Pointer to the start of the view.
Definition Span.h:539
constexpr __host__ __device__ SizeType size() const noexcept
Returns the size of the array.
Definition Span.h:327
TraceMessage info() const
Flow for an information message.
RunCommand makeCommand(const RunQueue &run_queue)
Creates a command associated with the queue run_queue.
Always enables tracing in Arcane parts concerning materials.
Int32 Integer
Type representing an integer.
std::int16_t Int16
Signed integer type of 16 bits.
std::int32_t Int32
Signed integer type of 32 bits.
Information for copying between two memory regions.