8#include <gtest/gtest.h>
10#include "arccore/base/PlatformUtils.h"
12#include "arccore/common/accelerator/Runner.h"
13#include "arccore/common/accelerator/RunQueue.h"
14#include "arccore/common/NumArray.h"
21#include "arccore/accelerator/LocalMemory.h"
34 Int32 nb_value, Int32 nb_loop)
40 by_block_partial_sum.resize(2048);
43 reduce_result.resize(1);
49 for (
int j = 0; j < nb_loop; ++j) {
52 loop_range.setBlockSize(command);
55 auto partial_sum_span =
viewInOut(command, by_block_partial_sum);
56 auto out_reduce_result =
viewOut(command, reduce_result);
59 auto grid = iter.grid();
60 auto block = iter.block();
61 auto w = iter.workItem();
62 auto block_partial_sum_span = block_partial_sum.span();
67 for (Int32 i : w.linearIndexes())
69 block_partial_sum_span[w.rankInBlock()] = my_v;
74 if (w.rankInBlock() == 0) {
75 Int32 nb_local = block_partial_sum_span.size();
77 for (Int32 i = 0; i < nb_local; ++i)
78 block_v += block_partial_sum_span[i];
79 partial_sum_span[block.groupRank()] = block_v;
84 if (w.rankInBlock() == 0 && block.groupRank() == 0) {
86 Int32 nb_block = grid.nbBlock();
87 for (Int32 i = 0; i < nb_block; ++i) {
88 final_sum += partial_sum_span[i];
90 out_reduce_result[0] = final_sum;
93 total_x += reduce_result[0];
97 Int64 nb_byte = c.
size() *
sizeof(Int64) * nb_loop;
99 Real nb_giga_byte_second = (
static_cast<Real
>(nb_byte) / 1.0e9) / diff;
100 std::cout <<
"** TotalCooperativeLaunch2=" << total_x
101 <<
" nb_value=" << nb_value
102 <<
" GB/s=" << nb_giga_byte_second <<
" time=" << diff <<
"\n";
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.
Mémoire locale (shared) à une RunCommand.
File d'exécution pour un accélérateur.
eMemoryResource memoryResource() const
Ressource mémoire utilisée pour les allocations avec cette instance.
bool isAcceleratorPolicy() const
Indique si l'instance est associée à un accélérateur.
Tableaux multi-dimensionnels pour les types numériques accessibles sur accélérateurs.
Vue d'un tableau d'éléments de type T.
constexpr __host__ __device__ SizeType size() const noexcept
Retourne la taille du tableau.
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.
-*- tab-width: 2; indent-tabs-mode: nil; coding: utf-8-with-signature -*-
@ HostPinned
Alloue sur l'hôte.