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_part, Int32 nb_loop)
37 if ((nb_value % nb_part) != 0)
38 ARCCORE_FATAL(
"{0} is not a multiple of {1}", nb_value, nb_part);
45 by_block_partial_sum.resize(2048);
48 reduce_result.resize(1);
53 for (
int j = 0; j < nb_loop; ++j) {
56 auto partial_sum_span =
viewInOut(command, by_block_partial_sum);
57 auto out_reduce_result =
viewOut(command, reduce_result);
60 auto grid = iter.grid();
61 auto block = iter.block();
62 auto w = iter.workItem();
65 for (Int32 i : w.linearIndexes())
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);
72 if (w.rankInBlock() == 0) {
74 partial_sum_span[block.groupRank()] = my_v;
77 if (w.rankInBlock() == 0 && block.groupRank() == 0) {
79 for (Int32 i = 0; i < nb_block; ++i) {
80 Int64 v = partial_sum_span[i];
82 final_sum += partial_sum_span[i];
84 partial_sum_span[0] = final_sum;
85 out_reduce_result[0] = final_sum;
86#if !defined(__INTEL_LLVM_COMPILER)
92 total_x += reduce_result[0];
96 Int64 nb_byte = c.
size() *
sizeof(Int64) * nb_loop;
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";
#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.
File d'exécution pour un accélérateur.
eMemoryResource memoryResource() const
Ressource mémoire utilisée pour les allocations avec cette instance.
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.