diff --git a/src/dbm/dbm_internal.h b/src/dbm/dbm_internal.h index abbd40eda7..fb946c0320 100644 --- a/src/dbm/dbm_internal.h +++ b/src/dbm/dbm_internal.h @@ -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); } diff --git a/src/dbm/dbm_matrix.c b/src/dbm/dbm_matrix.c index 44651d34d1..0afdb88e56 100644 --- a/src/dbm/dbm_matrix.c +++ b/src/dbm/dbm_matrix.c @@ -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); diff --git a/src/dbm/dbm_miniapp.c b/src/dbm/dbm_miniapp.c index 1078ae4f20..db45a26ec9 100644 --- a/src/dbm/dbm_miniapp.c +++ b/src/dbm/dbm_miniapp.c @@ -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"); diff --git a/src/dbm/dbm_multiply_opencl.c b/src/dbm/dbm_multiply_opencl.c index fce77c8375..29cd579ed4 100644 --- a/src/dbm/dbm_multiply_opencl.c +++ b/src/dbm/dbm_multiply_opencl.c @@ -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); diff --git a/src/dbm/dbm_multiply_opencl.cl b/src/dbm/dbm_multiply_opencl.cl index 7d35725ae9..366541b6b5 100644 --- a/src/dbm/dbm_multiply_opencl.cl +++ b/src/dbm/dbm_multiply_opencl.cl @@ -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); } } } diff --git a/src/offload/offload_runtime.h b/src/offload/offload_runtime.h index 807583c502..889229578a 100644 --- a/src/offload/offload_runtime.h +++ b/src/offload/offload_runtime.h @@ -22,11 +22,13 @@ #if defined(__OFFLOAD_CUDA) || defined(__OFFLOAD_HIP) || \ defined(__OFFLOAD_OPENCL) -#define __OFFLOAD - #include #include +#if !defined(__OFFLOAD) +#define __OFFLOAD +#endif + #if defined(__OFFLOAD_CUDA) #include #elif defined(__OFFLOAD_HIP)