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 "dbm_multiply.h"
8 : #include "../offload/offload_mempool.h"
9 : #include "../offload/offload_runtime.h"
10 : #include "dbm_hyperparams.h"
11 : #include "dbm_internal.h"
12 : #include "dbm_library.h"
13 : #include "dbm_multiply_comm.h"
14 : #include "dbm_multiply_cpu.h"
15 : #include "dbm_multiply_gpu.h"
16 :
17 : #include <assert.h>
18 : #include <limits.h>
19 : #include <math.h>
20 : #include <omp.h>
21 : #include <stdio.h>
22 : #include <stdlib.h>
23 : #include <string.h>
24 :
25 : /*******************************************************************************
26 : * \brief Private routine for computing the max filter threshold for each row.
27 : * \author Ole Schuett
28 : ******************************************************************************/
29 349938 : static float *compute_rows_max_eps(const bool trans, const dbm_matrix_t *matrix,
30 : const double filter_eps) {
31 349938 : const int nrows = (trans) ? matrix->ncols : matrix->nrows;
32 349938 : int *nblocks_per_row = calloc(nrows, sizeof(int));
33 349938 : float *row_max_eps = malloc(nrows * sizeof(float));
34 349938 : assert((nblocks_per_row != NULL && row_max_eps != NULL) || nrows == 0);
35 :
36 349938 : #pragma omp parallel
37 : {
38 : #pragma omp for
39 : for (int ishard = 0; ishard < dbm_get_num_shards(matrix); ishard++) {
40 : dbm_shard_t *shard = &matrix->shards[ishard];
41 : for (int iblock = 0; iblock < shard->nblocks; iblock++) {
42 : const dbm_block_t *blk = &shard->blocks[iblock];
43 : const int row = (trans) ? blk->col : blk->row;
44 : #pragma omp atomic
45 : ++nblocks_per_row[row];
46 : }
47 : }
48 : #pragma omp master
49 : cp_mpi_sum_int(nblocks_per_row, nrows, matrix->dist->comm);
50 : #pragma omp barrier
51 : #pragma omp for
52 : for (int i = 0; i < nrows; i++) {
53 : const float f =
54 : ((float)filter_eps) / ((float)imax(1, nblocks_per_row[i]));
55 : row_max_eps[i] = f * f;
56 : }
57 : } // end of omp parallel region
58 :
59 349938 : free(nblocks_per_row);
60 349938 : return row_max_eps; // Ownership of row_max_eps transfers to caller.
61 : }
62 :
63 : /*******************************************************************************
64 : * \brief Private struct for storing the context of the multiplication backend.
65 : * \author Ole Schuett
66 : ******************************************************************************/
67 : typedef struct {
68 : #if defined(__OFFLOAD) && !defined(__NO_OFFLOAD_DBM)
69 : dbm_multiply_gpu_context_t gpu;
70 : #endif
71 : int cpu_options; // Binary or'ed dbm_multiply_cpu_options (enum).
72 : } backend_context_t;
73 :
74 : /*******************************************************************************
75 : * \brief Private routine for initializing the multiplication backend.
76 : * \author Ole Schuett
77 : ******************************************************************************/
78 349938 : static backend_context_t *backend_start(const dbm_matrix_t *matrix_c) {
79 349938 : backend_context_t *const ctx = calloc(1, sizeof(backend_context_t));
80 : // BLAS and LIBXS benefit in general from DBM_MULTIPLY_TASK_REORDER.
81 349938 : ctx->cpu_options = DBM_MULTIPLY_TASK_REORDER;
82 :
83 : #if defined(__OFFLOAD) && !defined(__NO_OFFLOAD_DBM)
84 : dbm_multiply_gpu_start(DBM_MAX_BATCH_SIZE, dbm_get_num_shards(matrix_c),
85 : matrix_c->shards, &ctx->gpu);
86 : #else
87 349938 : (void)matrix_c; // mark as used
88 : #endif
89 :
90 349938 : return ctx;
91 : }
92 :
93 : /*******************************************************************************
94 : * \brief Private routine for handing newly arrived packs to the backend.
95 : * \author Ole Schuett
96 : ******************************************************************************/
97 0 : static bool backend_upload_packs(const dbm_pack_t *pack_a,
98 : const dbm_pack_t *pack_b,
99 : backend_context_t *ctx) {
100 : #if defined(__OFFLOAD) && !defined(__NO_OFFLOAD_DBM)
101 : return dbm_multiply_gpu_upload_packs(pack_a, pack_b, &ctx->gpu);
102 : #else
103 0 : (void)pack_a; // mark as used
104 0 : (void)pack_b;
105 0 : (void)ctx;
106 0 : return false;
107 : #endif
108 : }
109 :
110 : /*******************************************************************************
111 : * \brief Private routine for sending a batch to the multiplication backend.
112 : * \author Ole Schuett
113 : ******************************************************************************/
114 385000 : static void backend_process_batch(const int ntasks,
115 : const dbm_task_t batch[ntasks],
116 : const double alpha, const dbm_pack_t *pack_a,
117 : const dbm_pack_t *pack_b, const int kshard,
118 : dbm_shard_t *shard_c, const bool finish,
119 : const bool force_cpu,
120 : backend_context_t *ctx) {
121 385000 : if (NULL != ctx) {
122 : #if defined(__OFFLOAD) && !defined(__NO_OFFLOAD_DBM)
123 : if (!force_cpu) {
124 : dbm_multiply_gpu_process_batch(ntasks, batch, alpha, shard_c, kshard,
125 : finish, &ctx->gpu);
126 : } else
127 : #endif
128 : {
129 385000 : (void)kshard;
130 385000 : (void)finish;
131 385000 : (void)force_cpu;
132 385000 : dbm_multiply_cpu_process_batch(ntasks, batch, alpha, pack_a, pack_b,
133 : shard_c, ctx->cpu_options);
134 : }
135 : } else { // Validate against host (aka CPU).
136 0 : dbm_multiply_cpu_process_batch(ntasks, batch, alpha, pack_a, pack_b,
137 : shard_c, DBM_MULTIPLY_BLAS_LIBRARY);
138 : }
139 385000 : }
140 :
141 : /*******************************************************************************
142 : * \brief Private routine for shutting down the multiplication backend.
143 : * \author Ole Schuett
144 : ******************************************************************************/
145 349938 : static void backend_stop(backend_context_t *ctx) {
146 : #if defined(__OFFLOAD) && !defined(__NO_OFFLOAD_DBM)
147 : dbm_multiply_gpu_stop(&ctx->gpu);
148 : #endif
149 349938 : free(ctx);
150 349938 : }
151 :
152 : /*******************************************************************************
153 : * \brief Private routine returning the seconds since tick, which it resets.
154 : * \author Hans Pabst
155 : ******************************************************************************/
156 3254394 : static double phase_lap(double *tick) {
157 3254394 : const double now = omp_get_wtime();
158 3254394 : const double result = now - *tick;
159 3254394 : *tick = now;
160 3254394 : return result;
161 : }
162 :
163 : /*******************************************************************************
164 : * \brief Private routine for multiplying two packs (C += alpha * A * B).
165 : *
166 : * Blocks in each pack are grouped by shard (free_index % nshards) and sorted
167 : * by sum_index within each group. The algorithm:
168 : * 1. Builds shard-boundary lookup tables for A (rows) and B (cols).
169 : * 2. For each (shard_row, shard_col) pair, determines the contiguous A and B
170 : * block ranges belonging to that shard.
171 : * 3. Performs a merge-join over sum_index: advances A and B cursors in
172 : *lockstep, caching the B sub-range for each sum_index so that multiple A blocks
173 : *with the same sum_index reuse it without rescanning.
174 : * 4. Applies a norm-based filter (alpha^2 * norm_a * norm_b < eps) for early
175 : * rejection before looking up or allocating the C block.
176 : * 5. Accumulates matching pairs into a batched GEMM task list, flushing to the
177 : * backend (CPU or GPU) every DBM_MAX_BATCH_SIZE tasks.
178 : *
179 : * \author Ole Schuett and Hans Pabst
180 : ******************************************************************************/
181 384922 : static void multiply_packs(const bool transa, const bool transb,
182 : const double alpha, const dbm_pack_t *pack_a,
183 : const dbm_pack_t *pack_b,
184 : const dbm_matrix_t *matrix_a,
185 : const dbm_matrix_t *matrix_b, dbm_matrix_t *matrix_c,
186 : const float *rows_max_eps,
187 : const bool retain_sparsity, const bool force_cpu,
188 : int64_t *flop, backend_context_t *ctx) {
189 : // For validation, FLOPS do not count, and relying on ctx is not necessary.
190 384922 : backend_context_t *const context = (NULL != flop ? ctx : NULL);
191 384922 : const float alpha2 = (float)(alpha * alpha);
192 384922 : int64_t flop_sum = 0;
193 :
194 384922 : const int nshard_rows = matrix_c->dist->rows.nshards;
195 384922 : const int nshard_cols = matrix_c->dist->cols.nshards;
196 384922 : int *shard_row_start = calloc(nshard_rows, sizeof(int));
197 384922 : int *shard_col_start = calloc(nshard_cols, sizeof(int));
198 384922 : assert(NULL != shard_row_start && NULL != shard_col_start);
199 :
200 384922 : const int *sum_index_sizes_a =
201 : (transa) ? matrix_a->row_sizes : matrix_a->col_sizes;
202 384922 : const int *sum_index_sizes_b =
203 : (transb) ? matrix_b->col_sizes : matrix_b->row_sizes;
204 384922 : const int *free_index_sizes_a =
205 : (transa) ? matrix_a->col_sizes : matrix_a->row_sizes;
206 384922 : const int *free_index_sizes_b =
207 : (transb) ? matrix_b->row_sizes : matrix_b->col_sizes;
208 :
209 384922 : #pragma omp parallel reduction(+ : flop_sum)
210 : {
211 : // Thread-private array covering given work in piece-wise fashion.
212 : dbm_task_t *batch =
213 : offload_mempool_host_malloc(sizeof(dbm_task_t) * DBM_MAX_BATCH_SIZE);
214 : // Thread-local stats, added to the library's counters once at the end.
215 : int64_t counters[DBM_NUM_COUNTERS] = {0};
216 :
217 : // Blocks are ordered first by shard. Creating lookup tables of boundaries.
218 : #pragma omp for nowait
219 : for (int iblock = 1; iblock < pack_a->nblocks; iblock++) {
220 : const int shard_row = pack_a->blocks[iblock].free_index % nshard_rows;
221 : const int prev_shard_row =
222 : pack_a->blocks[iblock - 1].free_index % nshard_rows;
223 : if (prev_shard_row != shard_row) {
224 : shard_row_start[shard_row] = iblock;
225 : }
226 : }
227 : #pragma omp for
228 : for (int jblock = 1; jblock < pack_b->nblocks; jblock++) {
229 : const int shard_col = pack_b->blocks[jblock].free_index % nshard_cols;
230 : const int prev_shard_col =
231 : pack_b->blocks[jblock - 1].free_index % nshard_cols;
232 : if (prev_shard_col != shard_col) {
233 : shard_col_start[shard_col] = jblock;
234 : }
235 : }
236 :
237 : #pragma omp for collapse(2) DBM_OMP_SCHEDULE
238 : for (int shard_row = 0; shard_row < nshard_rows; shard_row++) {
239 : for (int shard_col = 0; shard_col < nshard_cols; shard_col++) {
240 : const int ishard = shard_row * nshard_cols + shard_col;
241 : dbm_shard_t *const shard_c = &matrix_c->shards[ishard];
242 : int ntasks = 0;
243 :
244 : // Determine contiguous block ranges for this shard in A and B.
245 : // Use a merge-join to find pairs of blocks with matching sum indices.
246 : // This utilizes that blocks within a shard are ordered by sum_index.
247 : const int iblock_start = shard_row_start[shard_row];
248 : int iblock_end = pack_a->nblocks;
249 : for (int t = iblock_start; t < pack_a->nblocks; ++t) {
250 : if (pack_a->blocks[t].free_index % nshard_rows != shard_row) {
251 : iblock_end = t;
252 : break;
253 : }
254 : }
255 : const int jblock_start = shard_col_start[shard_col];
256 : int jblock_end = pack_b->nblocks;
257 : for (int t = jblock_start; t < pack_b->nblocks; ++t) {
258 : if (pack_b->blocks[t].free_index % nshard_cols != shard_col) {
259 : jblock_end = t;
260 : break;
261 : }
262 : }
263 : if (iblock_start >= iblock_end || jblock_start >= jblock_end) {
264 : backend_process_batch(ntasks, batch, alpha, pack_a, pack_b, ishard,
265 : shard_c, true, force_cpu, context);
266 : continue;
267 : }
268 :
269 : // Merge over sum_index (both ranges sorted by sum_index).
270 : // Cache the B sub-range for each sum_index so that multiple A blocks
271 : // sharing the same sum_index reuse it without re-scanning B.
272 : int i = iblock_start, j = jblock_start, last_sum_index = -1;
273 : int b_range_start = -1, b_range_end = -1;
274 :
275 : while (i < iblock_end) {
276 : const dbm_pack_block_t *blk_a = &pack_a->blocks[i];
277 : const int sum_a = blk_a->sum_index;
278 :
279 : // Advance j until sum_b >= sum_a.
280 : while (j < jblock_end && pack_b->blocks[j].sum_index < sum_a) {
281 : ++j;
282 : }
283 : if (j >= jblock_end) {
284 : break; // No more matches possible.
285 : }
286 :
287 : const int sum_b = pack_b->blocks[j].sum_index;
288 : if (sum_b > sum_a) {
289 : ++i;
290 : continue; // Need next A block with higher sum_index.
291 : }
292 :
293 : // sum_a == sum_b: establish (or reuse) B range with this sum_index.
294 : if (sum_a != last_sum_index) {
295 : b_range_start = j;
296 : int t = j + 1;
297 : while (t < jblock_end && pack_b->blocks[t].sum_index == sum_a) {
298 : ++t;
299 : }
300 : b_range_end = t;
301 : last_sum_index = sum_a;
302 : }
303 :
304 : // Iterate over B blocks in current sum_index range.
305 : for (int jb = b_range_start; jb < b_range_end; ++jb) {
306 : const dbm_pack_block_t *const blk_b = &pack_b->blocks[jb];
307 :
308 : // Overlap the two dependent cache misses of upcoming lookups.
309 : if (jb + DBM_PREFETCH_SLOT < b_range_end) {
310 : dbm_shard_prefetch_slot(
311 : shard_c, blk_a->free_index,
312 : pack_b->blocks[jb + DBM_PREFETCH_SLOT].free_index);
313 : }
314 : if (jb + DBM_PREFETCH_BLOCK < b_range_end) {
315 : dbm_shard_prefetch_block(
316 : shard_c, blk_a->free_index,
317 : pack_b->blocks[jb + DBM_PREFETCH_BLOCK].free_index);
318 : }
319 :
320 : // Norm filter first (early reject).
321 : const float result_norm = alpha2 * blk_a->norm * blk_b->norm;
322 : if (result_norm < rows_max_eps[blk_a->free_index]) {
323 : continue;
324 : }
325 :
326 : // Check block sizes.
327 : const int m = free_index_sizes_a[blk_a->free_index];
328 : const int n = free_index_sizes_b[blk_b->free_index];
329 : const int k = sum_index_sizes_a[sum_a];
330 : assert(m == matrix_c->row_sizes[blk_a->free_index]);
331 : assert(n == matrix_c->col_sizes[blk_b->free_index]);
332 : assert(k == sum_index_sizes_b[blk_b->sum_index]);
333 :
334 : if (m == 0 || n == 0 || k == 0) {
335 : continue;
336 : }
337 :
338 : // Get C block.
339 : const int row = blk_a->free_index, col = blk_b->free_index;
340 : dbm_block_t *blk_c = dbm_shard_lookup(shard_c, row, col);
341 : if (blk_c == NULL) {
342 : if (retain_sparsity) {
343 : continue;
344 : }
345 : assert(dbm_get_shard_index(matrix_c, row, col) == ishard);
346 : assert(dbm_get_stored_coordinates(matrix_c, row, col) ==
347 : matrix_c->dist->my_rank);
348 : blk_c = dbm_shard_promise_new_block(shard_c, row, col, m * n);
349 : }
350 :
351 : // Count flops.
352 : const int64_t task_flops = 2LL * m * n * k;
353 : flop_sum += task_flops;
354 : ++counters[dbm_library_counter_index(m, n, k)];
355 :
356 : // Add block multiplication to batch.
357 : dbm_task_t *const tptr = &batch[ntasks];
358 : tptr->offset_a = blk_a->offset;
359 : tptr->offset_b = blk_b->offset;
360 : tptr->offset_c = blk_c->offset;
361 : tptr->m = m;
362 : tptr->n = n;
363 : tptr->k = k;
364 : ++ntasks;
365 :
366 : if (ntasks == DBM_MAX_BATCH_SIZE) {
367 : backend_process_batch(ntasks, batch, alpha, pack_a, pack_b,
368 : ishard, shard_c, false, force_cpu, context);
369 : ntasks = 0;
370 : }
371 : }
372 :
373 : // Advance i; if next A block has same sum_index, B range is reused.
374 : ++i;
375 : }
376 : backend_process_batch(ntasks, batch, alpha, pack_a, pack_b, ishard,
377 : shard_c, true, force_cpu, context);
378 : }
379 : }
380 :
381 : dbm_library_counters_add(counters);
382 : offload_mempool_host_free(batch);
383 : }
384 :
385 384922 : free(shard_row_start);
386 384922 : free(shard_col_start);
387 :
388 384922 : if (NULL != flop) {
389 384922 : *flop += flop_sum;
390 : }
391 384922 : }
392 :
393 : /*******************************************************************************
394 : * \brief Performs a multiplication of two dbm_matrix_t matrices.
395 : * See dbm_matrix.h for details.
396 : * \author Ole Schuett
397 : ******************************************************************************/
398 349938 : void dbm_multiply(const bool transa, const bool transb, const double alpha,
399 : const dbm_matrix_t *matrix_a, const dbm_matrix_t *matrix_b,
400 : const double beta, dbm_matrix_t *matrix_c,
401 : const bool retain_sparsity, const double filter_eps,
402 : int64_t *flop) {
403 349938 : assert(omp_get_num_threads() == 1);
404 349938 : assert(matrix_a != NULL && matrix_b != NULL && matrix_c != NULL);
405 349938 : double phase[DBM_NUM_PHASES] = {0}, tick = omp_get_wtime();
406 349938 : int64_t phase_bytes[DBM_NUM_PHASES] = {0};
407 :
408 : // Throughout the matrix multiplication code the "sum_index" and "free_index"
409 : // denote the summation (aka dummy) and free index from the Einstein notation.
410 349938 : const int num_sum_index_a = (transa) ? matrix_a->nrows : matrix_a->ncols;
411 349938 : const int num_sum_index_b = (transb) ? matrix_b->ncols : matrix_b->nrows;
412 349938 : const int num_free_index_a = (transa) ? matrix_a->ncols : matrix_a->nrows;
413 349938 : const int num_free_index_b = (transb) ? matrix_b->nrows : matrix_b->ncols;
414 :
415 : // Sanity check matrix dimensions.
416 349938 : assert(num_sum_index_a == num_sum_index_b);
417 349938 : assert(num_free_index_a == matrix_c->nrows);
418 349938 : assert(num_free_index_b == matrix_c->ncols);
419 :
420 : // Prepare matrix_c (host).
421 349938 : dbm_scale(matrix_c, beta);
422 :
423 : // Determine if validation shall be performed.
424 349938 : const char *const maxeps_env = getenv("DBM_MULTIPLY_MAXEPS");
425 349938 : const char *const verify_env = getenv("DBM_MULTIPLY_VERIFY");
426 349938 : const double maxeps = (NULL == maxeps_env ? 1E-1 : fabs(atof(maxeps_env)));
427 699876 : const int verify =
428 349938 : (NULL == verify_env ? (NULL == maxeps_env ? 0 : 1) : atoi(verify_env));
429 349938 : dbm_matrix_t *matrix_d = NULL;
430 349938 : phase[DBM_PHASE_SETUP] += phase_lap(&tick);
431 349938 : if (0 != verify) {
432 0 : dbm_distribution_t *const dist_shared = matrix_c->dist;
433 0 : dbm_create(&matrix_d, dist_shared, matrix_c->name, matrix_c->nrows,
434 0 : matrix_c->ncols, matrix_c->row_sizes, matrix_c->col_sizes);
435 0 : dbm_copy(matrix_d, matrix_c);
436 0 : tick = omp_get_wtime(); // validation is not a phase
437 : }
438 :
439 : // Compute filter thresholds for each row.
440 349938 : float *rows_max_eps = compute_rows_max_eps(transa, matrix_a, filter_eps);
441 :
442 : // Start uploading matrix_c to the GPU.
443 349938 : backend_context_t *ctx = backend_start(matrix_c);
444 349938 : phase[DBM_PHASE_SETUP] += phase_lap(&tick);
445 :
446 : // Redistribute matrix_a and matrix_b across MPI ranks.
447 349938 : dbm_comm_iterator_t *iter =
448 349938 : dbm_comm_iterator_start(transa, transb, matrix_a, matrix_b, matrix_c);
449 349938 : const double redistribute = phase_lap(&tick);
450 349938 : phase[DBM_PHASE_ALLTOALL] += iter->seconds_alltoall;
451 349938 : phase[DBM_PHASE_SORT] += iter->seconds_sort;
452 349938 : phase_bytes[DBM_PHASE_ALLTOALL] += iter->bytes_alltoall;
453 349938 : phase[DBM_PHASE_PACK] +=
454 349938 : redistribute - iter->seconds_alltoall - iter->seconds_sort;
455 :
456 : // Count flops if requested.
457 349938 : if (NULL != flop) {
458 349938 : *flop = 0;
459 : }
460 :
461 : // Main loop.
462 : dbm_pack_t *pack_a, *pack_b;
463 734860 : while (dbm_comm_iterator_next(iter, &pack_a, &pack_b)) {
464 384922 : phase[DBM_PHASE_SHIFT] += phase_lap(&tick);
465 384922 : const bool uploaded = backend_upload_packs(pack_a, pack_b, ctx);
466 384922 : (void)uploaded; // mark used
467 384922 : phase[DBM_PHASE_UPLOAD] += phase_lap(&tick);
468 384922 : multiply_packs(transa, transb, alpha, pack_a, pack_b, matrix_a, matrix_b,
469 : matrix_c, rows_max_eps, retain_sparsity, false /*!uploaded*/,
470 : flop, ctx);
471 384922 : phase[DBM_PHASE_MULTIPLY] += phase_lap(&tick);
472 : }
473 :
474 : // Wait for all other MPI ranks to complete, then release ressources.
475 349938 : phase_bytes[DBM_PHASE_SHIFT] += iter->bytes_shift;
476 349938 : dbm_comm_iterator_stop(iter);
477 349938 : phase[DBM_PHASE_SHIFT] += phase_lap(&tick);
478 349938 : backend_stop(ctx);
479 349938 : phase[DBM_PHASE_FINISH] += phase_lap(&tick);
480 :
481 349938 : if (NULL != matrix_d) {
482 0 : ctx = backend_start(matrix_d);
483 0 : iter =
484 0 : dbm_comm_iterator_start(transa, transb, matrix_a, matrix_b, matrix_d);
485 0 : while (dbm_comm_iterator_next(iter, &pack_a, &pack_b)) {
486 0 : multiply_packs(transa, transb, alpha, pack_a, pack_b, matrix_a, matrix_b,
487 : matrix_d, rows_max_eps, retain_sparsity, true, NULL, ctx);
488 : }
489 0 : dbm_comm_iterator_stop(iter);
490 0 : backend_stop(ctx);
491 0 : const double epsilon = dbm_maxeps(matrix_d, matrix_c);
492 0 : if (maxeps < epsilon) {
493 0 : if (1 == verify) {
494 0 : fprintf(stderr, "WARN ACC/LIBDBM: diff=%g\n", epsilon);
495 : } else {
496 0 : fprintf(stderr, "ERROR ACC/LIBDBM: diff=%g\n", epsilon);
497 0 : exit(EXIT_FAILURE);
498 : }
499 : }
500 0 : dbm_release(matrix_d);
501 0 : tick = omp_get_wtime(); // validation is not a phase
502 : }
503 :
504 : // Release filter thresholds.
505 349938 : free(rows_max_eps);
506 :
507 : // Final filter pass.
508 349938 : dbm_filter(matrix_c, filter_eps);
509 349938 : phase[DBM_PHASE_SETUP] += phase_lap(&tick);
510 349938 : dbm_library_phases_add(phase, phase_bytes);
511 349938 : }
512 :
513 : // EOF
|
