Arcane  4.1.12.0
Developer documentation
Loading...
Searching...
No Matches
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,
34 Int32 nb_value, 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 // Partial values per block.
43 // Must be sized to the maximum number of possible blocks
44 NumArray<Int64, MDDim1> by_block_partial_sum(queue.memoryResource());
45 by_block_partial_sum.resize(2048);
46 // To retrieve the reduction result.
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++ does not have 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_value=" << nb_value
101 << " GB/s=" << nb_giga_byte_second << " time=" << diff << "\n";
102 return total_x;
103}
104
105/*---------------------------------------------------------------------------*/
106/*---------------------------------------------------------------------------*/
Types and functions for managing synchronizations on accelerators.
Types and macros for managing hierarchical parallelism on accelerators.
#define RUNCOMMAND_LAUNCH(iter_name, bounds,...)
Macro to launch a command using hierarchical, possibly cooperative, parallelism.
Iteration range of a loop using cooperative hierarchical parallelism.
eMemoryResource memoryResource() const
Memory resource used for allocations with this instance.
Definition RunQueue.cc:372
Multi-dimensional arrays for numerical types accessible on accelerators.
View of an array of elements of type T.
Definition Span.h:805
constexpr __host__ __device__ SizeType size() const noexcept
Returns the size of the array.
Definition Span.h:327
RunCommand makeCommand(const RunQueue &run_queue)
Creates a command associated with the queue run_queue.
auto viewInOut(const ViewBuildInfo &vbi, CellMaterialVariableScalarRef< DataType > &var)
Read/write view for scalar material variables.
auto viewOut(const ViewBuildInfo &vbi, CellMaterialVariableScalarRef< DataType > &var)
Write view for scalar material variables.
Real getRealTime()
Real time used in seconds.
-- tab-width: 2; indent-tabs-mode: nil; coding: utf-8-with-signature --
@ HostPinned
Allocates on the host.