Arcane  v4.1.4.0
Documentation développeur
Chargement...
Recherche...
Aucune correspondance
TestCooperativeLaunch_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#include <gtest/gtest.h>
9
10#include "arccore/base/PlatformUtils.h"
11
12#include "arccore/common/accelerator/Runner.h"
13#include "arccore/common/accelerator/RunQueue.h"
14#include "arccore/common/NumArray.h"
15
17
19
21#include "arccore/accelerator/LocalMemory.h"
22
23/*---------------------------------------------------------------------------*/
24/*---------------------------------------------------------------------------*/
25
26using namespace Arcane;
27using namespace Arcane::Accelerator;
28
29/*---------------------------------------------------------------------------*/
30/*---------------------------------------------------------------------------*/
31
32extern "C++" Int64
33_testCooperativeLaunch(RunQueue queue, SmallSpan<const Int64> c, Int32 nb_thread,
34 Int32 nb_value, Int32 nb_part, Int32 nb_loop)
35{
36 Int64 total_x = 0;
37 if ((nb_value % nb_part) != 0)
38 ARCCORE_FATAL("{0} is not a multiple of {1}", nb_value, nb_part);
39 //Int32 nb_group = 15 * 1200;
40 //Int32 group_size = 128;
41 double x = Platform::getRealTime();
42 // Valeurs partielles par bloc.
43 // Doit être dimensionné au nombre maximum de blocs possibles
44 NumArray<Int64, MDDim1> by_block_partial_sum(queue.memoryResource());
45 by_block_partial_sum.resize(2048);
46 // Pour récupèrer le résultat de la réduction.
48 reduce_result.resize(1);
49 {
50 //nb_loop = 1;
51 //nb_value = 1000000;
52 SmallSpan<const Int64> c_view(c);
53 for (int j = 0; j < nb_loop; ++j) {
54 auto command = makeCommand(queue);
55 CooperativeWorkGroupLoopRange loop_range(nb_value);
56 auto partial_sum_span = viewInOut(command, by_block_partial_sum);
57 auto out_reduce_result = viewOut(command, reduce_result);
58 command << RUNCOMMAND_LAUNCH(iter, loop_range)
59 {
60 auto grid = iter.grid();
61 auto block = iter.block();
62 auto w = iter.workItem();
63
64 Int64 my_v = 0;
65 for (Int32 i : w.linearIndexes())
66 my_v += c_view[i];
67 block.barrier();
68 Int32 nb_block = grid.nbBlock();
69#if defined(ARCCORE_COMPILING_CUDA_OR_HIP) && defined(ARCCORE_DEVICE_CODE)
70 my_v = Arcane::Accelerator::Impl::block_reduce<Arcane::Accelerator::Impl::ReduceFunctorSum<Int64>, 32, Int64>(my_v);
71#endif
72 if (w.rankInBlock() == 0) {
73 //printf("V0=%d %ld\n", block.groupRank(), v);
74 partial_sum_span[block.groupRank()] = my_v;
75 }
76 grid.barrier();
77 if (w.rankInBlock() == 0 && block.groupRank() == 0) {
78 Int64 final_sum = 0;
79 for (Int32 i = 0; i < nb_block; ++i) {
80 Int64 v = partial_sum_span[i];
81 //printf("ADD_V block=%d v=%ld\n", i, v);
82 final_sum += partial_sum_span[i];
83 }
84 partial_sum_span[0] = final_sum;
85 out_reduce_result[0] = final_sum;
86#if !defined(__INTEL_LLVM_COMPILER)
87 // oneDPC++ ne possède pas de printf.
88 //printf("FINAL= nb_block=%d v=%ld\n", nb_block, final_sum);
89#endif
90 }
91 };
92 total_x += reduce_result[0];
93 }
94 }
95 double y = Platform::getRealTime();
96 Int64 nb_byte = c.size() * sizeof(Int64) * nb_loop;
97 Real diff = y - x;
98 Real nb_giga_byte_second = (static_cast<Real>(nb_byte) / 1.0e9) / diff;
99 std::cout << "** TotalCooperativeLaunch=" << total_x
100 << " nb_part=" << nb_part << " nb_value=" << nb_value
101 << " nb_thread=" << nb_thread
102 << " GB/s=" << nb_giga_byte_second << " time=" << diff << "\n";
103 return total_x;
104}
105
106/*---------------------------------------------------------------------------*/
107/*---------------------------------------------------------------------------*/
#define ARCCORE_FATAL(...)
Macro envoyant une exception FatalErrorException.
Types et fonctions pour gérer les synchronisations sur les accélérateurs.
Types et macros pour gérer le parallélisme hiérarchique sur les accélérateurs.
#define RUNCOMMAND_LAUNCH(iter_name, bounds,...)
Macro pour lancer une commande utilisant le parallélisme hiérarchique, éventuellement coopératif.
Intervalle d'itération d'une boucle utilisant le parallélisme hiérarchique collaboratif.
eMemoryResource memoryResource() const
Ressource mémoire utilisée pour les allocations avec cette instance.
Definition RunQueue.cc:372
Tableaux multi-dimensionnels pour les types numériques accessibles sur accélérateurs.
Vue d'un tableau d'éléments de type T.
Definition Span.h:801
constexpr __host__ __device__ SizeType size() const noexcept
Retourne la taille du tableau.
Definition Span.h:325
Espace de nom pour l'utilisation des accélérateurs.
RunCommand makeCommand(const RunQueue &run_queue)
Créé une commande associée à la file run_queue.
auto viewInOut(const ViewBuildInfo &vbi, CellMaterialVariableScalarRef< DataType > &var)
Vue en lecture/écriture pour les variables materiaux scalaire.
auto viewOut(const ViewBuildInfo &vbi, CellMaterialVariableScalarRef< DataType > &var)
Vue en écriture pour les variables materiaux scalaire.
ARCCORE_BASE_EXPORT Real getRealTime()
Temps Real utilisé en secondes.
-*- tab-width: 2; indent-tabs-mode: nil; coding: utf-8-with-signature -*-
@ HostPinned
Alloue sur l'hôte.