Line data Source code
1 : /*----------------------------------------------------------------------------*/
2 : /* CP2K: A general program to perform molecular dynamics simulations */
3 : /* Copyright 2000-2026 CP2K developers group <https://cp2k.org> */
4 : /* */
5 : /* SPDX-License-Identifier: BSD-3-Clause */
6 : /*----------------------------------------------------------------------------*/
7 : #include "grid_library.h"
8 : #include "grid_common.h"
9 : #include "grid_constants.h"
10 :
11 : #include "../../mpiwrap/cp_mpi.h"
12 : #include "../../offload/offload_runtime.h"
13 :
14 : #include <assert.h>
15 : #include <omp.h>
16 : #include <stddef.h>
17 : #include <stdio.h>
18 : #include <stdlib.h>
19 : #include <string.h>
20 :
21 : #define GRID_LIBRARY_PRINT(FN, MSG, OUTPUT_UNIT) \
22 : ((FN)(MSG, (int)strlen(MSG), OUTPUT_UNIT))
23 :
24 : // counter dimensions
25 : #define GRID_NBACKENDS 5
26 : #define GRID_NKERNELS 4
27 : #define GRID_MAX_LP 20
28 :
29 : typedef struct {
30 : grid_sphere_cache sphere_cache;
31 : void *scratch; // thread local scratch, grown on demand
32 : size_t scratch_size;
33 : long counters[GRID_NBACKENDS * GRID_NKERNELS * GRID_MAX_LP];
34 : } grid_library_globals;
35 :
36 : static grid_library_globals **per_thread_globals = NULL;
37 : static bool library_initialized = false;
38 : static int max_threads = 0;
39 : static grid_library_config config = {
40 : .backend = GRID_BACKEND_AUTO, .validate = false, .apply_cutoff = false};
41 :
42 : #if !defined(_OPENMP)
43 : #error "OpenMP is required. Please add -fopenmp to your C compiler flags."
44 : #endif
45 :
46 : #if defined(NDEBUG)
47 : #error \
48 : "Please do not build CP2K with NDEBUG. There is no performance advantage and asserts will save your neck."
49 : #endif
50 :
51 : /*******************************************************************************
52 : * \brief Initializes the grid library.
53 : * \author Ole Schuett
54 : ******************************************************************************/
55 1474 : void grid_library_init(void) {
56 1474 : if (library_initialized) {
57 0 : printf("Error: Grid library was already initialized.\n");
58 0 : abort();
59 : }
60 :
61 : #if defined(__OFFLOAD) && !defined(__NO_OFFLOAD_GRID)
62 : // Reserve global GPU memory for storing the intermediate Cab matrix blocks.
63 : // CUDA does not allow to increase this limit after a kernel was launched.
64 : // Unfortunately, the required memory is hard to predict because we neither
65 : // know which tasks will be run nor how many thread blocks the available GPU
66 : // can execute in parallel... 64 MiB ought to be enough for anybody ;-)
67 : offloadEnsureMallocHeapSize(64 * 1024 * 1024);
68 : #endif
69 :
70 1474 : max_threads = omp_get_max_threads();
71 1474 : per_thread_globals = malloc(max_threads * sizeof(grid_library_globals *));
72 1474 : assert(per_thread_globals != NULL);
73 :
74 : // Using parallel regions to ensure memory is allocated near a thread's core.
75 : #pragma omp parallel default(none) shared(per_thread_globals) \
76 : num_threads(max_threads)
77 : {
78 : const int ithread = omp_get_thread_num();
79 : per_thread_globals[ithread] = malloc(sizeof(grid_library_globals));
80 : assert(per_thread_globals[ithread] != NULL);
81 : memset(per_thread_globals[ithread], 0, sizeof(grid_library_globals));
82 : }
83 :
84 1474 : library_initialized = true;
85 1474 : }
86 :
87 : /*******************************************************************************
88 : * \brief Finalizes the grid library.
89 : * \author Ole Schuett
90 : ******************************************************************************/
91 1474 : void grid_library_finalize(void) {
92 1474 : if (!library_initialized) {
93 0 : printf("Error: Grid library is not initialized.\n");
94 0 : abort();
95 : }
96 :
97 2948 : for (int i = 0; i < max_threads; i++) {
98 1474 : grid_sphere_cache_free(&per_thread_globals[i]->sphere_cache);
99 1474 : free(per_thread_globals[i]->scratch);
100 1474 : free(per_thread_globals[i]);
101 : }
102 1474 : free(per_thread_globals);
103 1474 : per_thread_globals = NULL;
104 1474 : library_initialized = false;
105 1474 : }
106 :
107 : /*******************************************************************************
108 : * \brief Returns a pointer to the thread local sphere cache.
109 : * \author Ole Schuett
110 : ******************************************************************************/
111 181657401 : grid_sphere_cache *grid_library_get_sphere_cache(void) {
112 181657401 : const int ithread = omp_get_thread_num();
113 181657401 : assert(ithread < max_threads);
114 181657401 : return &per_thread_globals[ithread]->sphere_cache;
115 : }
116 :
117 : /*******************************************************************************
118 : * \brief Returns a pointer to a thread local scratch buffer with at least the
119 : * given size. The buffer grows on demand and is reused across calls.
120 : ******************************************************************************/
121 13003937 : void *grid_library_get_scratch(const size_t bytes) {
122 13003937 : const int ithread = omp_get_thread_num();
123 13003937 : assert(ithread < max_threads);
124 13003937 : grid_library_globals *globals = per_thread_globals[ithread];
125 13003937 : if (globals->scratch_size < bytes) {
126 3213 : free(globals->scratch);
127 3213 : globals->scratch = malloc(bytes);
128 3213 : assert(globals->scratch != NULL);
129 3213 : globals->scratch_size = bytes;
130 : }
131 13003937 : return globals->scratch;
132 : }
133 :
134 : /*******************************************************************************
135 : * \brief Configures the grid library.
136 : * \author Ole Schuett
137 : ******************************************************************************/
138 11304 : void grid_library_set_config(const enum grid_backend backend,
139 : const bool validate, const bool apply_cutoff) {
140 11304 : config.backend = backend;
141 11304 : config.validate = validate;
142 11304 : config.apply_cutoff = apply_cutoff;
143 11304 : }
144 :
145 : /*******************************************************************************
146 : * \brief Returns the library config.
147 : * \author Ole Schuett
148 : ******************************************************************************/
149 522629 : grid_library_config grid_library_get_config(void) { return config; }
150 :
151 : /*******************************************************************************
152 : * \brief Adds given increment to counter specified by lp, backend, and kernel.
153 : * \author Ole Schuett
154 : ******************************************************************************/
155 194680174 : void grid_library_counter_add(const int lp, const enum grid_backend backend,
156 : const enum grid_library_kernel kernel,
157 : const int increment) {
158 194680174 : assert(lp >= 0);
159 194680174 : assert(kernel < GRID_NKERNELS);
160 194680174 : const int back = backend - GRID_BACKEND_REF;
161 194680174 : assert(back < GRID_NBACKENDS);
162 194680174 : const int idx = back * GRID_NKERNELS * GRID_MAX_LP + kernel * GRID_MAX_LP +
163 194680174 : imin(lp, GRID_MAX_LP - 1);
164 194680174 : const int ithread = omp_get_thread_num();
165 194680174 : assert(ithread < max_threads);
166 194680174 : per_thread_globals[ithread]->counters[idx] += increment;
167 194680174 : }
168 :
169 : /*******************************************************************************
170 : * \brief Comperator passed to qsort to compare two counters.
171 : * \author Ole Schuett
172 : ******************************************************************************/
173 19033896 : static int compare_counters(const void *a, const void *b) {
174 19033896 : return *(long *)b - *(long *)a;
175 : }
176 :
177 : /*******************************************************************************
178 : * \brief Prints statistics gathered by the grid library.
179 : * \author Ole Schuett
180 : ******************************************************************************/
181 11306 : void grid_library_print_stats(const int fortran_comm,
182 : void (*print_func)(const char *, int, int),
183 11306 : const int output_unit) {
184 11306 : if (!library_initialized) {
185 0 : printf("Error: Grid library is not initialized.\n");
186 0 : abort();
187 : }
188 :
189 11306 : const cp_mpi_comm_t comm = cp_mpi_comm_f2c(fortran_comm);
190 : // Sum all counters across threads and mpi ranks.
191 11306 : const int ncounters = GRID_NBACKENDS * GRID_NKERNELS * GRID_MAX_LP;
192 11306 : long counters[ncounters][2];
193 11306 : memset(counters, 0, ncounters * 2 * sizeof(long));
194 11306 : double total = 0.0;
195 4533706 : for (int i = 0; i < ncounters; i++) {
196 4522400 : counters[i][1] = i; // needed as inverse index after qsort
197 9044800 : for (int j = 0; j < max_threads; j++) {
198 4522400 : counters[i][0] += per_thread_globals[j]->counters[i];
199 : }
200 4522400 : cp_mpi_sum_long(&counters[i][0], 1, comm);
201 4522400 : total += counters[i][0];
202 : }
203 :
204 : // Sort counters.
205 11306 : qsort(counters, ncounters, 2 * sizeof(long), &compare_counters);
206 :
207 : // Determine if anything needs to be printed.
208 11306 : bool print = false;
209 1889932 : for (int i = 0; i < ncounters && !print; i++) {
210 1878626 : if (counters[i][0] != 0) {
211 6626 : print = true;
212 : }
213 : }
214 11306 : if (!print) {
215 4680 : return; // nothing to be printed
216 : }
217 :
218 : // Print counters.
219 6626 : GRID_LIBRARY_PRINT(print_func, "\n", output_unit);
220 6626 : GRID_LIBRARY_PRINT(
221 : print_func,
222 : " ----------------------------------------------------------------"
223 : "---------------\n",
224 : output_unit);
225 6626 : GRID_LIBRARY_PRINT(
226 : print_func,
227 : " - "
228 : " -\n",
229 : output_unit);
230 6626 : GRID_LIBRARY_PRINT(
231 : print_func,
232 : " - GRID STATISTICS "
233 : " -\n",
234 : output_unit);
235 6626 : GRID_LIBRARY_PRINT(
236 : print_func,
237 : " - "
238 : " -\n",
239 : output_unit);
240 6626 : GRID_LIBRARY_PRINT(
241 : print_func,
242 : " ----------------------------------------------------------------"
243 : "---------------\n",
244 : output_unit);
245 6626 : GRID_LIBRARY_PRINT(
246 : print_func,
247 : " LP KERNEL BACKEND "
248 : "COUNT PERCENT\n",
249 : output_unit);
250 :
251 6626 : const char *kernel_names[] = {"collocate ortho", "integrate ortho",
252 : "collocate general", "integrate general"};
253 6626 : const char *backend_names[] = {"REF", "CPU", "DGEMM", "GPU"};
254 :
255 2657026 : for (int i = 0; i < ncounters; i++) {
256 2650400 : if (counters[i][0] == 0)
257 2568058 : continue; // skip empty counters
258 82342 : const double percent = 100.0 * counters[i][0] / total;
259 82342 : const int idx = counters[i][1];
260 82342 : const int backend_stride = GRID_NKERNELS * GRID_MAX_LP;
261 82342 : const int back = idx / backend_stride;
262 82342 : const int kern = (idx % backend_stride) / GRID_MAX_LP;
263 82342 : const int lp = (idx % backend_stride) % GRID_MAX_LP;
264 82342 : char buffer[100];
265 82342 : snprintf(buffer, sizeof(buffer), " %-5i %-17s %-6s %34li %10.2f%%\n", lp,
266 : kernel_names[kern], backend_names[back], counters[i][0], percent);
267 82342 : GRID_LIBRARY_PRINT(print_func, buffer, output_unit);
268 : }
269 :
270 6626 : GRID_LIBRARY_PRINT(
271 : print_func,
272 : " ----------------------------------------------------------------"
273 : "---------------\n",
274 : output_unit);
275 : }
276 :
277 : // EOF
|
