mirror of
https://github.com/cp2k/cp2k.git
synced 2026-07-27 21:55:16 -04:00
DBM: more accurate matrix checksum
- Miniapp: fixed conditional compilation (#4052). - Calculate checksum more accurately. - Miniapp: expected value uses uint64_t. - Kahan's summation is needed. - OpenCL - Revised SLM-variant of DBM-kernel. - Adjusted internal parameters - Make use of NDEBUG. - Cleanup offload/offload_library.h.
This commit is contained in:
Hans Pabst
parent
7c193b409d
commit
d2e611afe3
6 changed files with 58 additions and 45 deletions
|
|
@ -9,7 +9,7 @@
|
|||
#define DBM_INTERNAL_H
|
||||
|
||||
/*******************************************************************************
|
||||
* \brief Returns the larger of two given integer (missing from the C standard)
|
||||
* \brief Returns the larger of two given integers (missing from the C standard)
|
||||
* \author Ole Schuett
|
||||
******************************************************************************/
|
||||
static inline int imax(int x, int y) { return (x > y ? x : y); }
|
||||
|
|
|
|||
|
|
@ -505,16 +505,32 @@ void dbm_iterator_next_block(dbm_iterator_t *iter, int *row, int *col,
|
|||
******************************************************************************/
|
||||
void dbm_iterator_stop(dbm_iterator_t *iter) { free(iter); }
|
||||
|
||||
/*******************************************************************************
|
||||
* \brief Private routine to accumulate using Kahan's summation.
|
||||
* \author Hans Pabst
|
||||
******************************************************************************/
|
||||
static double kahan_sum(double value, double *accumulator,
|
||||
double *compensation) {
|
||||
double r, c;
|
||||
assert(NULL != accumulator && NULL != compensation);
|
||||
c = value - *compensation;
|
||||
r = *accumulator + c;
|
||||
*compensation = (r - *accumulator) - c;
|
||||
*accumulator = r;
|
||||
return r;
|
||||
}
|
||||
|
||||
/*******************************************************************************
|
||||
* \brief Computes a checksum of the given matrix.
|
||||
* \author Ole Schuett
|
||||
******************************************************************************/
|
||||
double dbm_checksum(const dbm_matrix_t *matrix) {
|
||||
double checksum = 0.0;
|
||||
double checksum = 0.0, compensation = 0.0;
|
||||
for (int ishard = 0; ishard < dbm_get_num_shards(matrix); ishard++) {
|
||||
const dbm_shard_t *shard = &matrix->shards[ishard];
|
||||
for (int i = 0; i < shard->data_size; i++) {
|
||||
checksum += shard->data[i] * shard->data[i];
|
||||
const double value = shard->data[i];
|
||||
kahan_sum(value * value, &checksum, &compensation);
|
||||
}
|
||||
}
|
||||
dbm_mpi_sum_double(&checksum, 1, matrix->dist->comm);
|
||||
|
|
|
|||
|
|
@ -32,7 +32,7 @@ static void print_func(char *message, int output_unit) {
|
|||
}
|
||||
|
||||
/*******************************************************************************
|
||||
* \brief Returns the smaller of two given integer (missing from the C standard)
|
||||
* \brief Returns the smaller of the two integers (missing from the C standard).
|
||||
* \author Ole Schuett
|
||||
******************************************************************************/
|
||||
static inline int imin(int x, int y) { return (x < y ? x : y); }
|
||||
|
|
@ -168,8 +168,7 @@ static void set_all_blocks(dbm_matrix_t *matrix) {
|
|||
dbm_iterator_next_block(iter, &row, &col, &block, &row_size, &col_size);
|
||||
const int block_size = row_size * col_size;
|
||||
for (int i = 0; i < block_size; i++) {
|
||||
#if defined(DBM_VALIDATE_AGAINST_LIBXSMM) && defined(__LIBXSMM) && \
|
||||
defined(__OFFLOAD) && !defined(__NO_OFFLOAD_DBM)
|
||||
#if defined(DBM_VALIDATE_AGAINST_LIBXSMM) && defined(__LIBXSMM)
|
||||
block[i] = 1.0 / (i + 1);
|
||||
#else
|
||||
block[i] = 1.0;
|
||||
|
|
@ -185,8 +184,7 @@ static void set_all_blocks(dbm_matrix_t *matrix) {
|
|||
******************************************************************************/
|
||||
void benchmark_multiply(const int M, const int N, const int K, const int m,
|
||||
const int n, const int k, const dbm_mpi_comm_t comm) {
|
||||
#if defined(DBM_VALIDATE_AGAINST_LIBXSMM) && defined(__LIBXSMM) && \
|
||||
defined(__OFFLOAD) && !defined(__NO_OFFLOAD_DBM)
|
||||
#if defined(DBM_VALIDATE_AGAINST_LIBXSMM) && defined(__LIBXSMM)
|
||||
dbm_matrix_t *matrix_a = create_some_matrix(M, K, 1, m, k, k, comm);
|
||||
dbm_matrix_t *matrix_b = create_some_matrix(K, N, k, k, 1, n, comm);
|
||||
#else
|
||||
|
|
@ -212,12 +210,10 @@ void benchmark_multiply(const int M, const int N, const int K, const int m,
|
|||
|
||||
// Validate checksum.
|
||||
// Since all matrix elements were set to 1.0 the checksum is an integer.
|
||||
#if defined(DBM_VALIDATE_AGAINST_LIBXSMM) && defined(__LIBXSMM) && \
|
||||
defined(__OFFLOAD) && !defined(__NO_OFFLOAD_DBM)
|
||||
#if defined(DBM_VALIDATE_AGAINST_LIBXSMM) && defined(__LIBXSMM)
|
||||
const double expected = 0, checksum = 0;
|
||||
#else
|
||||
const double expected = (int64_t)M * (int64_t)m * (int64_t)N * (int64_t)n *
|
||||
(int64_t)K * (int64_t)K * (int64_t)k * (int64_t)k;
|
||||
const double expected = (uint64_t)M * m * N * n * K * K * k * k;
|
||||
const double checksum = dbm_checksum(matrix_c);
|
||||
#endif
|
||||
|
||||
|
|
@ -326,17 +322,19 @@ int main(int argc, char *argv[]) {
|
|||
continue;
|
||||
}
|
||||
if (0 < mnk[0]) { /* valid MxNxK? */
|
||||
int nm = (NULL == arg ? 0 : atoi(arg)), nn, nk;
|
||||
if (0 < nm) {
|
||||
const int m = mnk[0];
|
||||
const int n = (0 < mnk[1] ? mnk[1] : m);
|
||||
const int k = (0 < mnk[2] ? mnk[2] : m);
|
||||
int M = (NULL == arg ? 0 : atoi(arg)), N, K;
|
||||
if (0 < M) {
|
||||
arg = strtok(NULL, delims);
|
||||
nn = (NULL == arg ? 1 : atoi(arg));
|
||||
N = (NULL == arg ? 1 : atoi(arg));
|
||||
arg = strtok(NULL, delims);
|
||||
nk = (NULL == arg ? 1 : atoi(arg));
|
||||
K = (NULL == arg ? 1 : atoi(arg));
|
||||
} else { /* default */
|
||||
nm = nn = nk = 128;
|
||||
M = N = K = 128;
|
||||
}
|
||||
benchmark_multiply(nm, nn, nk, mnk[0], 0 < mnk[1] ? mnk[1] : mnk[0],
|
||||
0 < mnk[2] ? mnk[2] : mnk[0], comm);
|
||||
benchmark_multiply(M, N, K, m, n, k, comm);
|
||||
mnk[0] = mnk[1] = mnk[2] = 0;
|
||||
} else {
|
||||
fprintf(stderr, "ERROR: invalid argument(s)\n");
|
||||
|
|
|
|||
|
|
@ -12,7 +12,7 @@
|
|||
#include "dbm_multiply_opencl.cl.h"
|
||||
|
||||
void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream,
|
||||
const int mnk_range[3][2], double alpha,
|
||||
const int mnk[3][2], double alpha,
|
||||
int ntasks, const dbm_task_t *tasks,
|
||||
const double *pack_a_data,
|
||||
const double *pack_b_data,
|
||||
|
|
@ -27,17 +27,14 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream,
|
|||
cl_event event, *const perf_event =
|
||||
((0 <= verbosity && 2 >= verbosity) ? NULL : &event);
|
||||
const c_dbcsr_acc_opencl_stream_t *const str = ACC_OPENCL_STREAM(stream);
|
||||
const size_t max_m = mnk_range[0][1], work_tasks = ntasks;
|
||||
const size_t max_m = mnk[0][1], work_tasks = ntasks;
|
||||
size_t work_size[] = {1, 1, 1}, ibatch = 0, iadata = 0, ibdata = 0,
|
||||
icdata = 0;
|
||||
c_dbcsr_acc_opencl_info_memptr_t adata, bdata, cdata, batch;
|
||||
assert(NULL != pack_a_data && NULL != pack_b_data && NULL != shard_c_data);
|
||||
assert(0 < mnk_range[0][0] && 0 < mnk_range[0][1] &&
|
||||
mnk_range[0][0] <= mnk_range[0][1]);
|
||||
assert(0 < mnk_range[1][0] && 0 < mnk_range[1][1] &&
|
||||
mnk_range[1][0] <= mnk_range[1][1]);
|
||||
assert(0 < mnk_range[2][0] && 0 < mnk_range[2][1] &&
|
||||
mnk_range[2][0] <= mnk_range[2][1]);
|
||||
assert(0 < mnk[0][0] && 0 < mnk[0][1] && mnk[0][0] <= mnk[0][1]);
|
||||
assert(0 < mnk[1][0] && 0 < mnk[1][1] && mnk[1][0] <= mnk[1][1]);
|
||||
assert(0 < mnk[2][0] && 0 < mnk[2][1] && mnk[2][0] <= mnk[2][1]);
|
||||
assert(NULL != str && NULL != str->queue);
|
||||
assert(0 < ntasks && NULL != tasks);
|
||||
#if defined(OPENCL_DBM_SOURCE_MULTIPLY)
|
||||
|
|
@ -55,9 +52,10 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream,
|
|||
const int bn0 = (0 == c_dbcsr_acc_opencl_config.device.nv
|
||||
? (0 == c_dbcsr_acc_opencl_config.device.amd ? 4 : 8)
|
||||
: 2);
|
||||
int bn = LIBXSMM_CLMP(NULL == bn_env ? bn0 : atoi(bn_env), 1, 32);
|
||||
int lu = LIBXSMM_CLMP(NULL == lu_env ? 0 : atoi(lu_env), -2, 1);
|
||||
int sm = (NULL == sm_env ? 0 /*default*/ : atoi(sm_env));
|
||||
int bn = LIBXSMM_CLMP(
|
||||
NULL == bn_env ? (0 == sm ? bn0 : (bn0 * 2)) : atoi(bn_env), 1, 32);
|
||||
int gen = ((NULL == lu_env && NULL == bn_env && NULL == sm_env &&
|
||||
NULL == wg_env)
|
||||
? (NULL == gen_env ? 1 /*default*/ : atoi(gen_env))
|
||||
|
|
@ -70,8 +68,10 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream,
|
|||
const size_t wgsize0 = c_dbcsr_acc_opencl_config.device.wgsize[0];
|
||||
const size_t wgsize1 = c_dbcsr_acc_opencl_config.device.wgsize[1];
|
||||
size_t wgsize2 = c_dbcsr_acc_opencl_config.device.wgsize[2];
|
||||
size_t offset =
|
||||
(0 == c_dbcsr_acc_opencl_config.debug ? strlen(params) : 0);
|
||||
size_t offset = ((0 == c_dbcsr_acc_opencl_config.debug &&
|
||||
0 == c_dbcsr_acc_opencl_config.dump)
|
||||
? strlen(params)
|
||||
: 0);
|
||||
offset += (size_t)c_dbcsr_acc_opencl_flags_atomics(
|
||||
&c_dbcsr_acc_opencl_config.device, c_dbcsr_acc_opencl_atomic_fp_64,
|
||||
extensions, &nextensions, params + offset, sizeof(params) - offset);
|
||||
|
|
@ -209,14 +209,14 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream,
|
|||
EXIT_SUCCESS == clGetEventProfilingInfo(*perf_event,
|
||||
CL_PROFILING_COMMAND_END,
|
||||
sizeof(cl_ulong), &end, NULL)) {
|
||||
const size_t flops = max_m * mnk_range[1][1] * mnk_range[2][1] * ntasks;
|
||||
const double flops = 1E-6 * mnk[0][1] * mnk[1][1] * mnk[2][1] * ntasks;
|
||||
const double dkrnl = 1E-9 * LIBXSMM_DELTA(begin, end);
|
||||
const double dtotl = 1E+3 * LIBXSMM_MAX(dkrnl, dhost);
|
||||
fprintf(stderr,
|
||||
"INFO ACC/LIBDBM: DBM-kernel mnk=%ix%ix%i ntasks=%i "
|
||||
"kernel_ms=%.2g total_ms=%.2g gflops=%.1f\n",
|
||||
mnk_range[0][1], mnk_range[1][1], mnk_range[2][1], ntasks, dkrnl,
|
||||
dtotl, 1E-6 * flops / dtotl);
|
||||
mnk[0][1], mnk[1][1], mnk[2][1], ntasks, dkrnl, dtotl,
|
||||
flops / dtotl);
|
||||
}
|
||||
}
|
||||
OFFLOAD_CHECK(result);
|
||||
|
|
|
|||
|
|
@ -19,12 +19,6 @@
|
|||
#endif
|
||||
#define BCST_NO(V) (V)
|
||||
|
||||
#if defined(SM) && (0 < SM)
|
||||
#define TLS(MEM) (MEM)[get_local_id(0)]
|
||||
#else
|
||||
#define TLS(MEM) (MEM)
|
||||
#endif
|
||||
|
||||
#define SINT short
|
||||
|
||||
#define X(T, I) (T)->I /* task can be taken by value or by pointer */
|
||||
|
|
@ -81,7 +75,7 @@ dbm_multiply(double alpha, int itask, int ntasks, int size,
|
|||
global const double *restrict bmat, global double *restrict cmat) {
|
||||
const int i = (int)get_global_id(0);
|
||||
#if defined(SM) && (0 < SM)
|
||||
local double cvec[WG][BN + SM - 1];
|
||||
local double tls[WG][BN + SM - 1], *const cvec = &tls[get_local_id(0)];
|
||||
#else
|
||||
double cvec[BN];
|
||||
#endif
|
||||
|
|
@ -92,15 +86,18 @@ dbm_multiply(double alpha, int itask, int ntasks, int size,
|
|||
const int max_m = size / ntasks, tid = i / max_m;
|
||||
const SINT m = i - tid * max_m;
|
||||
global const dbm_task_t *const task = &tasks[itask + tid];
|
||||
if (m < XM(task)) { /* valid task */
|
||||
#if !defined(NDEBUG)
|
||||
if (m < XM(task))
|
||||
#endif
|
||||
{ /* valid task */
|
||||
bmat += XB(task);
|
||||
#if defined(BCST_WG)
|
||||
if (XM(task) <= XN(task)) { /* BCST_WG to broadcast B-values */
|
||||
DBM_MULTIPLY(alpha, task, amat, bmat, cmat, TLS(cvec), m, BN, BCST_WG);
|
||||
DBM_MULTIPLY(alpha, task, amat, bmat, cmat, cvec, m, BN, BCST_WG);
|
||||
} else
|
||||
#endif
|
||||
{
|
||||
DBM_MULTIPLY(alpha, task, amat, bmat, cmat, TLS(cvec), m, BN, BCST_NO);
|
||||
DBM_MULTIPLY(alpha, task, amat, bmat, cmat, cvec, m, BN, BCST_NO);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -22,11 +22,13 @@
|
|||
|
||||
#if defined(__OFFLOAD_CUDA) || defined(__OFFLOAD_HIP) || \
|
||||
defined(__OFFLOAD_OPENCL)
|
||||
#define __OFFLOAD
|
||||
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
#if !defined(__OFFLOAD)
|
||||
#define __OFFLOAD
|
||||
#endif
|
||||
|
||||
#if defined(__OFFLOAD_CUDA)
|
||||
#include <cuda_runtime.h>
|
||||
#elif defined(__OFFLOAD_HIP)
|
||||
|
|
|
|||
Loading…
Add table
Add a link
Reference in a new issue