From 62500b0a468a23f0af8602e4535c28007c427a2d Mon Sep 17 00:00:00 2001 From: Marc Henry de Frahan Date: Thu, 11 Jun 2026 21:18:30 +0000 Subject: [PATCH] Update OpenACC timing and kernel tuning --- CMakeLists.txt | 12 +++- swm_c/c_OpenACC/CMakeLists.txt | 4 +- swm_c/c_OpenACC/shallow_swap.acc.Tile.c | 93 +++++++++++++++++-------- swm_c/c_OpenACC/shallow_swap.acc.c | 74 ++++++++++++++------ 4 files changed, 130 insertions(+), 53 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 2c493d2..a40b66f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -25,6 +25,16 @@ if(SWM_CUDA AND SWM_HIP) message(FATAL_ERROR "SWM_CUDA and SWM_HIP cannot both be ON. Please enable only one.") endif() +if(SWM_OPENACC AND SWM_DEVICE STREQUAL "gpu") + if(NOT (CMAKE_C_COMPILER_ID STREQUAL "NVHPC" OR CMAKE_C_COMPILER_ID STREQUAL "NVIDIA")) + message(FATAL_ERROR + "OpenACC GPU build requires NVIDIA HPC SDK compiler (nvc).\n" + "Current C compiler: ${CMAKE_C_COMPILER_ID}\n" + "Use: cmake -DCMAKE_C_COMPILER=nvc -DCMAKE_CXX_COMPILER=nvc++ ..." + ) + endif() +endif() + if(SWM_DEVICE STREQUAL "cpu" AND SWM_CUDA) message(WARNING "SWM_CUDA is ON but SWM_DEVICE is set to 'cpu'. Disabling SWM_CUDA.") set(SWM_CUDA OFF CACHE BOOL "Enable the CUDA versions of the mini-app" FORCE) @@ -49,4 +59,4 @@ endif() if(SWM_KOKKOS) add_subdirectory(swm_kokkos) -endif() \ No newline at end of file +endif() diff --git a/swm_c/c_OpenACC/CMakeLists.txt b/swm_c/c_OpenACC/CMakeLists.txt index a9b01bf..a0bdff3 100644 --- a/swm_c/c_OpenACC/CMakeLists.txt +++ b/swm_c/c_OpenACC/CMakeLists.txt @@ -8,10 +8,10 @@ if(CMAKE_C_COMPILER_ID STREQUAL "GNU") set(OPENACC_LINK_FLAGS -fopenacc) elseif(CMAKE_C_COMPILER_ID STREQUAL "NVHPC" OR CMAKE_C_COMPILER_ID STREQUAL "NVIDIA") if(SWM_DEVICE STREQUAL "cpu") - set(OPENACC_FLAGS -acc=multicore -Minfo=accel -Mnofma) + set(OPENACC_COMPILE_FLAGS -acc=multicore -Minfo=accel) set(OPENACC_LINK_FLAGS -acc=multicore) elseif(SWM_DEVICE STREQUAL "gpu") - set(OPENACC_FLAGS -acc=gpu -Minfo=accel -Mnofma ) + set(OPENACC_COMPILE_FLAGS -acc=gpu -Minfo=accel) set(OPENACC_LINK_FLAGS -acc=gpu) endif() else() diff --git a/swm_c/c_OpenACC/shallow_swap.acc.Tile.c b/swm_c/c_OpenACC/shallow_swap.acc.Tile.c index e21d38f..f72cf12 100644 --- a/swm_c/c_OpenACC/shallow_swap.acc.Tile.c +++ b/swm_c/c_OpenACC/shallow_swap.acc.Tile.c @@ -27,6 +27,16 @@ #include #ifdef _OPENACC #include +#if defined(__has_include) +#if __has_include() +#include +#define SWM_HAVE_NVTX 1 +#endif +#endif +#endif +#ifndef SWM_HAVE_NVTX +#define nvtxRangePush(name) ((void)0) +#define nvtxRangePop() ((void)0) #endif #define MIN(x,y) ((x)>(y)?(y):(x)) #define MAX(x,y) ((x)>(y)?(x):(y)) @@ -40,7 +50,8 @@ #define SIZE ((M_LEN)*(N_LEN)) #define ITMAX 4000 #define L_OUT TRUE -#define TILE 32 +#define TILE_I 64 +#define TILE_J 4 extern double wtime(); extern void dswap(double **a, double **b); @@ -86,7 +97,10 @@ extern void dswap(double **a, double **b); int main(int argc, char **argv) { - double *u,*v,*p,*unew,*vnew,*pnew,*uold,*vold,*pold,*cu,*cv,*z,*h,*psi; + double * restrict u, * restrict v, * restrict p; + double * restrict unew, * restrict vnew, * restrict pnew; + double * restrict uold, * restrict vold, * restrict pold; + double * restrict cu, * restrict cv, * restrict z, * restrict h, * restrict psi; u = (double *)malloc(sizeof(double)*M_LEN*N_LEN); v = (double *)malloc(sizeof(double)*M_LEN*N_LEN); @@ -126,10 +140,9 @@ int main(int argc, char **argv) { int i,j; // timer variables - double mfs100,mfs200,mfs300,mfs310; + double mfs100,mfs200,mfs300; double t100,t200,t300; - double tstart,ctime,tcyc,time,ptime; - double t100i,t200i,t300i; + double tstart,tcyc,time,ptime; double c1,c2; // ** Initialisations ** @@ -148,17 +161,16 @@ int main(int argc, char **argv) { alpha = .001; el = N * dx; - pi = 4. * atanf(1.); + pi = 4. * atan(1.); tpi = pi + pi; di = tpi / M; dj = tpi / N; pcf = pi * pi * a * a / (el * el); -#pragma acc enter data copyin(dt,tdt,dx,dy,a,alpha,el,pi,tpi,di,dj,pcf,tdts8,tdtsdx,tdtsdy,fsdx,fsdy,p[:SIZE],u[:SIZE], \ - v[:SIZE],pnew[:SIZE],unew[:SIZE],vnew[:SIZE]) +#pragma acc enter data copyin(p[:SIZE],u[:SIZE],v[:SIZE]) create(pnew[:SIZE],unew[:SIZE],vnew[:SIZE]) // Initial values of the stream function and p -#pragma acc parallel loop tile(TILE,TILE) independent present(p[:SIZE]) deviceptr(psi)//private(a,di,dj,pcf) +#pragma acc parallel loop tile(TILE_I,TILE_J) independent present(p[:SIZE]) deviceptr(psi)//private(a,di,dj,pcf) for (i=0;i %d, vnew %d -> %d , pnew %d -> %d\n",M*N_LEN+j,j,j+1,M*N_LEN +j + 1,j,M*N_LEN +j); unew[j] = unew[M*N_LEN+j]; vnew[M*N_LEN +j + 1] = vnew[j + 1]; @@ -337,12 +354,21 @@ int main(int argc, char **argv) { vnew[M*N_LEN] = vnew[N]; pnew[M*N_LEN+N] = pnew[0]; } +#pragma acc wait + c2 = wtime(); + t200 = t200 + (c2 - c1); +#ifdef _OPENACC + nvtxRangePop(); +#endif time = time + dt; // Time smoothing and update for next cycle if ( ncycle > 1 ) { +#ifdef _OPENACC + nvtxRangePush("UpdateOldVariables"); +#endif c1 = wtime(); #pragma acc parallel loop collapse(2) independent present(unew[:SIZE],vnew[:SIZE],pnew[:SIZE],u[:SIZE],v[:SIZE],p[:SIZE]) deviceptr(uold,vold,pold) for (i=0;i 0 ) { mfs200 = ITMAX * 26. * M * N / t200 / 1000000; } if ( t300 > 0 ) { mfs300 = ITMAX * 15. * M * N / t300 / 1000000; } - c2 = wtime(); - ctime = c2 - tstart; - tcyc = ctime / ITMAX; + tcyc = loop_time / ITMAX; - printf(" cycle number %d total computer time %f time per cycle %f\n", ITMAX, ctime, tcyc); + printf(" cycle number %d total computer time %f time per cycle %f\n", ITMAX, loop_time, tcyc); printf(" time and megaflops for loop 100 %.6f %.6f\n", t100, mfs100); printf(" time and megaflops for loop 200 %.6f %.6f\n", t200, mfs200); printf(" time and megaflops for loop 300 %.6f %.6f\n", t300, mfs300); @@ -487,6 +516,10 @@ int main(int argc, char **argv) { free((void *) psi); #endif +#ifdef _OPENACC + acc_shutdown(acc_get_device_type()); +#endif + return(0); } void dswap(double **pA, double **pB) diff --git a/swm_c/c_OpenACC/shallow_swap.acc.c b/swm_c/c_OpenACC/shallow_swap.acc.c index 4f6e97a..cbfbb2c 100644 --- a/swm_c/c_OpenACC/shallow_swap.acc.c +++ b/swm_c/c_OpenACC/shallow_swap.acc.c @@ -27,6 +27,16 @@ #include #ifdef _OPENACC #include +#if defined(__has_include) +#if __has_include() +#include +#define SWM_HAVE_NVTX 1 +#endif +#endif +#endif +#ifndef SWM_HAVE_NVTX +#define nvtxRangePush(name) ((void)0) +#define nvtxRangePop() ((void)0) #endif #define MIN(x,y) ((x)>(y)?(y):(x)) #define MAX(x,y) ((x)>(y)?(x):(y)) @@ -85,7 +95,10 @@ extern void dswap(double **a, double **b); int main(int argc, char **argv) { - double *u,*v,*p,*unew,*vnew,*pnew,*uold,*vold,*pold,*cu,*cv,*z,*h,*psi; + double * restrict u, * restrict v, * restrict p; + double * restrict unew, * restrict vnew, * restrict pnew; + double * restrict uold, * restrict vold, * restrict pold; + double * restrict cu, * restrict cv, * restrict z, * restrict h, * restrict psi; u = (double *)malloc(sizeof(double)*M_LEN*N_LEN); v = (double *)malloc(sizeof(double)*M_LEN*N_LEN); @@ -127,7 +140,7 @@ int main(int argc, char **argv) { // timer variables double mfs100,mfs200,mfs300; double t100,t200,t300; - double tstart,ctime,tcyc,time,ptime; + double tstart,tcyc,time,ptime; double c1,c2; // ** Initialisations ** @@ -146,15 +159,14 @@ int main(int argc, char **argv) { alpha = .001; el = N * dx; - pi = 4. * atanf(1.); + pi = 4. * atan(1.); tpi = pi + pi; di = tpi / M; dj = tpi / N; pcf = pi * pi * a * a / (el * el); -#pragma acc enter data copyin(dt,tdt,dx,dy,a,alpha,el,pi,tpi,di,dj,pcf,tdts8,tdtsdx,tdtsdy,fsdx,fsdy,p[:SIZE],u[:SIZE], \ - v[:SIZE],pnew[:SIZE],unew[:SIZE],vnew[:SIZE]) +#pragma acc enter data copyin(p[:SIZE],u[:SIZE],v[:SIZE]) create(pnew[:SIZE],unew[:SIZE],vnew[:SIZE]) // Initial values of the stream function and p #pragma acc parallel loop collapse(2) present(p[:SIZE]) deviceptr(psi)//private(a,di,dj,pcf) for (i=0;i 1 ) { +#ifdef _OPENACC + nvtxRangePush("UpdateOldVariables"); +#endif c1 = wtime(); #pragma acc parallel loop collapse(2) independent present(unew[:SIZE],vnew[:SIZE],pnew[:SIZE],u[:SIZE],v[:SIZE],p[:SIZE]) deviceptr(uold,vold,pold) for (i=0;i 0 ) { mfs200 = ITMAX * 26. * M * N / t200 / 1000000; } if ( t300 > 0 ) { mfs300 = ITMAX * 15. * M * N / t300 / 1000000; } - c2 = wtime(); - ctime = c2 - tstart; - tcyc = ctime / ITMAX; + tcyc = loop_time / ITMAX; - printf(" cycle number %d total computer time %f time per cycle %f\n", ITMAX, ctime, tcyc); + printf(" cycle number %d total computer time %f time per cycle %f\n", ITMAX, loop_time, tcyc); printf(" time and megaflops for loop 100 %.6f %.6f\n", t100, mfs100); printf(" time and megaflops for loop 200 %.6f %.6f\n", t200, mfs200); printf(" time and megaflops for loop 300 %.6f %.6f\n", t300, mfs300); @@ -486,6 +516,10 @@ int main(int argc, char **argv) { free((void *) psi); #endif +#ifdef _OPENACC + acc_shutdown(acc_get_device_type()); +#endif + return(0); } void dswap(double **pA, double **pB)