diff --git a/configure.ac b/configure.ac index 64fd9c27..ca018e1e 100644 --- a/configure.ac +++ b/configure.ac @@ -78,69 +78,225 @@ AS_IF([test "$ac_cv_header_gpfs_h" = "yes" -o "$ac_cv_header_gpfs_fcntl_h" = "ye ]) ]) -# Check for CUDA +# --------------------------------------------------------------------------- +# GPU runtime: CUDA or HIP +# --------------------------------------------------------------------------- + +# CUDA runtime AC_ARG_WITH([cuda], - [AS_HELP_STRING([--with-cuda], - [support configurable CUDA @<:@default=check@:>@])], + [AS_HELP_STRING([--with-cuda@<:@=PREFIX@:>@], + [enable CUDA runtime support @<:@default=check@:>@])], [], [with_cuda=check]) +have_cuda=no AS_IF([test "x$with_cuda" != xno], [ - LDFLAGS="$LDFLAGS -L$with_cuda/lib64 -L$with_cuda/lib -Wl,--enable-new-dtags -Wl,-rpath=$with_cuda/lib64:$with_cuda/lib" - CPPFLAGS="$CPPFLAGS -I$with_cuda/include" + AS_IF([test "x$with_cuda" != xyes && test "x$with_cuda" != xcheck], [ + CUDA_CPPFLAGS="-I$with_cuda/include" + CUDA_LDFLAGS="-L$with_cuda/lib64 -L$with_cuda/lib -Wl,--enable-new-dtags -Wl,-rpath=$with_cuda/lib64:$with_cuda/lib" + ]) + + save_CPPFLAGS="$CPPFLAGS"; save_LDFLAGS="$LDFLAGS"; save_LIBS="$LIBS" + CPPFLAGS="$CPPFLAGS $CUDA_CPPFLAGS" + LDFLAGS="$LDFLAGS $CUDA_LDFLAGS" - AC_CHECK_HEADERS([cuda_runtime.h], [AC_DEFINE([HAVE_CUDA], [], [CUDA GPU API found])], [ + AC_CHECK_HEADERS([cuda_runtime.h], [], [ if test "x$with_cuda" != xcheck; then AC_MSG_FAILURE([--with-cuda was given, not found]) fi ]) -AS_IF([test "$ac_cv_header_cuda_runtime_h" = "yes"], [ - AC_SEARCH_LIBS([cudaMalloc], [cudart cudart_static], [], - [AC_MSG_ERROR([Library containing cudaMalloc symbol not found])]) - ]) -]) -AM_CONDITIONAL([HAVE_CUDA], [test x$ac_cv_search_cudaMalloc != x ]) -AM_COND_IF([HAVE_CUDA],[AC_DEFINE([HAVE_CUDA], [], [CUDA GPU API found])]) + AS_IF([test "$ac_cv_header_cuda_runtime_h" = "yes"], [ + AC_SEARCH_LIBS([cudaMalloc], [cudart cudart_static], [ + have_cuda=yes + AS_IF([test "x$ac_cv_search_cudaMalloc" = "xnone required"], + [CUDA_LIBS=""], [CUDA_LIBS="$ac_cv_search_cudaMalloc"]) + ], [AC_MSG_ERROR([Library containing cudaMalloc symbol not found])]) + ]) -# Check for GPUDirect -AC_ARG_WITH([gpuDirect], - [AS_HELP_STRING([--with-gpuDirect], - [support configurable GPUDirect @<:@default=check@:>@])], - [], [with_gpuDirect=check]) + CPPFLAGS="$save_CPPFLAGS"; LDFLAGS="$save_LDFLAGS"; LIBS="$save_LIBS" +]) +AC_SUBST([CUDA_CPPFLAGS]) +AC_SUBST([CUDA_LDFLAGS]) +AC_SUBST([CUDA_LIBS]) +AM_CONDITIONAL([HAVE_CUDA], [test "x$have_cuda" = xyes]) +AM_COND_IF([HAVE_CUDA], [AC_DEFINE([HAVE_CUDA], [1], [CUDA GPU runtime available])]) + +# HIP runtime +AC_ARG_WITH([hip], + [AS_HELP_STRING([--with-hip@<:@=PREFIX@:>@], + [enable HIP runtime support @<:@default=check@:>@])], + [], [with_hip=check]) + +have_hip=no +AS_IF([test "x$with_hip" != xno], [ + AS_IF([test "x$with_hip" != xyes && test "x$with_hip" != xcheck], [ + HIP_CPPFLAGS="-I$with_hip/include -D__HIP_PLATFORM_AMD__" + HIP_LDFLAGS="-L$with_hip/lib -L$with_hip/lib64 -Wl,--enable-new-dtags -Wl,-rpath=$with_hip/lib:$with_hip/lib64" + ],[ + HIP_CPPFLAGS="-D__HIP_PLATFORM_AMD__" + ]) -AS_IF([test "x$with_gpuDirect" != xno], [ - LDFLAGS="$LDFLAGS -L$with_gpuDirect/lib64 -L$with_gpuDirect/lib -Wl,--enable-new-dtags -Wl,-rpath=$with_gpuDirect/lib64:$with_gpuDirect/lib" - CPPFLAGS="$CPPFLAGS -I$with_gpuDirect/include" + save_CPPFLAGS="$CPPFLAGS"; save_LDFLAGS="$LDFLAGS"; save_LIBS="$LIBS" + CPPFLAGS="$CPPFLAGS $HIP_CPPFLAGS" + LDFLAGS="$LDFLAGS $HIP_LDFLAGS" - AC_CHECK_HEADERS([cufile.h], [], [ - if test "x$with_gpuDirect" != xcheck; then - AC_MSG_FAILURE([--with-gpuDirect was given, not found]) + AC_CHECK_HEADERS([hip/hip_runtime.h], [], [ + if test "x$with_hip" != xcheck; then + AC_MSG_FAILURE([--with-hip was given, not found]) fi - with_gpuDirect=no ]) - AS_IF([test "$ac_cv_header_cufile_h" = "yes"], [ - AC_SEARCH_LIBS([cuFileDriverOpen], [cufile], [], [ - AC_MSG_WARN([Library containing cuFileDriverOpen symbol not found]) - with_gpuDirect=no - ]) + AS_IF([test "$ac_cv_header_hip_hip_runtime_h" = "yes"], [ + AC_SEARCH_LIBS([hipMalloc], [amdhip64], [ + have_hip=yes + AS_IF([test "x$ac_cv_search_hipMalloc" = "xnone required"], + [HIP_LIBS=""], [HIP_LIBS="$ac_cv_search_hipMalloc"]) + ], [AC_MSG_WARN([Library containing hipMalloc symbol not found])]) ]) + + CPPFLAGS="$save_CPPFLAGS"; LDFLAGS="$save_LDFLAGS"; LIBS="$save_LIBS" +]) +AC_SUBST([HIP_CPPFLAGS]) +AC_SUBST([HIP_LDFLAGS]) +AC_SUBST([HIP_LIBS]) + +# HIPCC compiler +AC_ARG_VAR([HIPCC], [HIP compiler]) +AC_PATH_PROG([HIPCC], [hipcc]) +AM_CONDITIONAL([HAVE_HIPCC], [test -n "$HIPCC"]) + +AS_IF([test "x$have_hip" = xyes && test "x$HIPCC" = x], [ + AC_MSG_FAILURE([HIP runtime support requires hipcc for utilities-gpu.hip]) ]) -# Check for NVCC +AS_IF([test "x$have_cuda" = xyes && test "x$have_hip" = xyes], [ + AC_MSG_WARN([Both CUDA and HIP detected; CUDA takes precedence.]) + have_hip=no +]) +AM_CONDITIONAL([HAVE_HIP], [test "x$have_hip" = xyes]) +AM_COND_IF([HAVE_HIP], [AC_DEFINE([HAVE_HIP], [1], [HIP GPU runtime available])]) + +AM_CONDITIONAL([HAVE_GPU], [test "x$have_cuda" = xyes || test "x$have_hip" = xyes]) +AM_COND_IF([HAVE_GPU], [AC_DEFINE([HAVE_GPU], [1], [GPU runtime available])]) + +# NVCC compiler NVCC=nvcc AC_ARG_VAR(varNVCC, "Name/path of the NVIDIA compiler") AC_ARG_WITH([nvcc], AS_HELP_STRING([--with-nvcc], [Use the NVCC as specified]), [AS_IF([test "$with_nvcc" != "yes"], NVCC=$with_nvcc)], [], [with_nvcc=check]) - AC_PATH_PROG(NVCC, nvcc) -AM_CONDITIONAL([HAVE_NVCC], test -n "$NVCC") +AM_CONDITIONAL([HAVE_NVCC], [test -n "$NVCC"]) -# Use GPU Direct only if NVCC and HAVE_GPU_DIRECT is supported -AM_CONDITIONAL([HAVE_GPU_DIRECT], test x$with_gpuDirect != xno -a x$NVCC != x ) -AM_COND_IF([HAVE_GPU_DIRECT],[ - AC_DEFINE([HAVE_GPU_DIRECT], [], [GPUDirect API found]) - ]) +# CUDA runtime requires NVCC for utilities-gpu.cu +AS_IF([test "x$have_cuda" = xyes && test "x$NVCC" = x], [ + AC_MSG_FAILURE([CUDA runtime support requires nvcc for utilities-gpu.cu]) +]) + +# --------------------------------------------------------------------------- +# GPU Direct Storage backend: cuFile or hipFile +# --------------------------------------------------------------------------- + +# Backend selector +AC_ARG_WITH([gpu-direct-backend], + [AS_HELP_STRING([--with-gpu-direct-backend=cufile|hipfile|none], + [select one GPU Direct Storage backend @<:@default=none@:>@])], + [gpu_direct_backend="$withval"], [gpu_direct_backend=none]) + +# Legacy alias: --with-gpuDirect -> cufile +AC_ARG_WITH([gpuDirect], + [AS_HELP_STRING([--with-gpuDirect@<:@=PREFIX@:>@], + [deprecated alias for --with-gpu-direct-backend=cufile])], + [legacy_gpu_direct=yes + AS_CASE([$withval], + [no], [gpu_direct_backend=none], + [yes], [gpu_direct_backend=cufile], + [check], [gpu_direct_backend=cufile], + [*], [gpu_direct_backend=cufile + cufile_prefix="$withval"])], + [legacy_gpu_direct=no]) + +AS_CASE([$gpu_direct_backend], + [cufile], [], + [hipfile], [], + [none], [], + [AC_MSG_FAILURE([--with-gpu-direct-backend must be cufile, hipfile, or none])]) + +# cuFile backend +have_cufile=no +AS_IF([test "x$gpu_direct_backend" = xcufile], [ + AS_IF([test "x$have_cuda" != xyes], [ + AC_MSG_FAILURE([cuFile backend requires CUDA runtime support]) + ]) + AC_ARG_WITH([cufile], + [AS_HELP_STRING([--with-cufile@<:@=PREFIX@:>@], + [cuFile install prefix])], + [cufile_prefix="$withval"], []) + AS_IF([test "x$cufile_prefix" != x && test "x$cufile_prefix" != xyes], [ + CUFILE_CPPFLAGS="-I$cufile_prefix/include" + CUFILE_LDFLAGS="-L$cufile_prefix/lib64 -L$cufile_prefix/lib -Wl,--enable-new-dtags -Wl,-rpath=$cufile_prefix/lib64:$cufile_prefix/lib" + ]) + save_CPPFLAGS="$CPPFLAGS"; save_LDFLAGS="$LDFLAGS"; save_LIBS="$LIBS" + CPPFLAGS="$CPPFLAGS $CUFILE_CPPFLAGS" + LDFLAGS="$LDFLAGS $CUFILE_LDFLAGS" + AC_CHECK_HEADERS([cufile.h], [], [ + AC_MSG_FAILURE([cuFile backend selected, but cufile.h not found]) + ]) + AS_IF([test "$ac_cv_header_cufile_h" = "yes"], [ + AC_SEARCH_LIBS([cuFileDriverOpen], [cufile], [ + have_cufile=yes + AS_IF([test "x$ac_cv_search_cuFileDriverOpen" = "xnone required"], + [CUFILE_LIBS=""], [CUFILE_LIBS="$ac_cv_search_cuFileDriverOpen"]) + ], [ + AC_MSG_FAILURE([cuFile backend selected, but libcuFile not found]) + ]) + ]) + CPPFLAGS="$save_CPPFLAGS"; LDFLAGS="$save_LDFLAGS"; LIBS="$save_LIBS" +]) +AC_SUBST([CUFILE_CPPFLAGS]) +AC_SUBST([CUFILE_LDFLAGS]) +AC_SUBST([CUFILE_LIBS]) +AM_CONDITIONAL([HAVE_CUFILE], [test "x$have_cufile" = xyes]) +AM_COND_IF([HAVE_CUFILE], [AC_DEFINE([HAVE_CUFILE], [1], [cuFile GPU Direct Storage backend])]) + +# hipFile backend +have_hipfile=no +AS_IF([test "x$gpu_direct_backend" = xhipfile], [ + AS_IF([test "x$have_hip" != xyes], [ + AC_MSG_FAILURE([hipFile backend requires HIP runtime support]) + ]) + AC_ARG_WITH([hipfile], + [AS_HELP_STRING([--with-hipfile@<:@=PREFIX@:>@], + [hipFile install prefix])], + [hipfile_prefix="$withval"], []) + AS_IF([test "x$hipfile_prefix" != x && test "x$hipfile_prefix" != xyes], [ + HIPFILE_CPPFLAGS="-I$hipfile_prefix/include" + HIPFILE_LDFLAGS="-L$hipfile_prefix/lib -L$hipfile_prefix/lib64 -Wl,--enable-new-dtags -Wl,-rpath=$hipfile_prefix/lib:$hipfile_prefix/lib64" + ]) + save_CPPFLAGS="$CPPFLAGS"; save_LDFLAGS="$LDFLAGS"; save_LIBS="$LIBS" + CPPFLAGS="$CPPFLAGS $HIPFILE_CPPFLAGS" + LDFLAGS="$LDFLAGS $HIPFILE_LDFLAGS" + AC_CHECK_HEADERS([hipfile.h], [], [ + AC_MSG_FAILURE([hipFile backend selected, but hipfile.h not found]) + ]) + AS_IF([test "$ac_cv_header_hipfile_h" = "yes"], [ + AC_SEARCH_LIBS([hipFileDriverOpen], [hipfile], [ + have_hipfile=yes + AS_IF([test "x$ac_cv_search_hipFileDriverOpen" = "xnone required"], + [HIPFILE_LIBS=""], [HIPFILE_LIBS="$ac_cv_search_hipFileDriverOpen"]) + ], [ + AC_MSG_FAILURE([hipFile backend selected, but libhipfile not found]) + ]) + ]) + CPPFLAGS="$save_CPPFLAGS"; LDFLAGS="$save_LDFLAGS"; LIBS="$save_LIBS" +]) +AC_SUBST([HIPFILE_CPPFLAGS]) +AC_SUBST([HIPFILE_LDFLAGS]) +AC_SUBST([HIPFILE_LIBS]) +AM_CONDITIONAL([HAVE_HIPFILE], [test "x$have_hipfile" = xyes]) +AM_COND_IF([HAVE_HIPFILE], [AC_DEFINE([HAVE_HIPFILE], [1], [hipFile GPU Direct Storage backend])]) + +AM_CONDITIONAL([HAVE_GPU_DIRECT], [test "x$have_cufile" = xyes || test "x$have_hipfile" = xyes]) +AM_COND_IF([HAVE_GPU_DIRECT], [AC_DEFINE([HAVE_GPU_DIRECT], [1], [GPU Direct Storage backend available])]) AC_CHECK_FUNCS diff --git a/src/Makefile.am b/src/Makefile.am index 2cd2d569..59487f41 100755 --- a/src/Makefile.am +++ b/src/Makefile.am @@ -5,7 +5,7 @@ if USE_CAPS bin_PROGRAMS += IOR MDTEST MD-WORKBENCH endif -noinst_HEADERS = ior.h utilities.h parse_options.h aiori.h iordef.h ior-internal.h option.h mdtest.h aiori-debug.h aiori-POSIX.h md-workbench.h +noinst_HEADERS = ior.h utilities.h parse_options.h aiori.h iordef.h ior-internal.h option.h mdtest.h aiori-debug.h aiori-POSIX.h md-workbench.h gpu_runtime.h gpu_io.h lib_LIBRARIES = libaiori.a libaiori_a_SOURCES = ior.c mdtest.c utilities.c parse_options.c ior-output.c option.c md-workbench.c @@ -33,7 +33,7 @@ mdtest_CPPFLAGS = if USE_HDFS_AIORI # TBD: figure out how to find the appropriate -I and -L dirs. Maybe we can # get them from the corresponding bin/ dir in $PATH, or pick an -# environment var to use (and set it in modulefiles), or provide a +# environment var to use (and set it up in modulefiles), or provide a # config-flag, to set a variable we use here. extraSOURCES += aiori-HDFS.c extraCPPFLAGS += -I/opt/hadoop-2.2.0/include @@ -42,12 +42,31 @@ extraLDADD += -lhdfs endif if HAVE_CUDA -extraLDADD += -lcudart +extraSOURCES += cuda_runtime.c utilities-gpu.cu +extraCPPFLAGS += @CUDA_CPPFLAGS@ +extraLDFLAGS += @CUDA_LDFLAGS@ +extraLDADD += @CUDA_LIBS@ endif -if HAVE_GPU_DIRECT -extraLDADD += -lcufile -lstdc++ -extraSOURCES += utilities-gpu.cu +if HAVE_HIP +extraSOURCES += hip_runtime.c utilities-gpu.hip +extraCPPFLAGS += @HIP_CPPFLAGS@ +extraLDFLAGS += @HIP_LDFLAGS@ +extraLDADD += @HIP_LIBS@ +endif + +if HAVE_CUFILE +extraSOURCES += gpu_io_cufile.c +extraCPPFLAGS += @CUFILE_CPPFLAGS@ +extraLDFLAGS += @CUFILE_LDFLAGS@ +extraLDADD += @CUFILE_LIBS@ -lstdc++ +endif + +if HAVE_HIPFILE +extraSOURCES += gpu_io_hipfile.c +extraCPPFLAGS += @HIPFILE_CPPFLAGS@ +extraLDFLAGS += @HIPFILE_LDFLAGS@ +extraLDADD += @HIPFILE_LIBS@ -lstdc++ endif if USE_HDF5_AIORI @@ -197,4 +216,7 @@ build.conf: @echo CFLAGS=$(CFLAGS) $(extraCPPFLAGS) >> build.conf .cu.o: - $(NVCC) $(CFLAGS) -c -o $@ $< + $(NVCC) $(CPPFLAGS) $(extraCPPFLAGS) -c -o $@ $< + +.hip.o: + $(HIPCC) $(CPPFLAGS) $(extraCPPFLAGS) -c -o $@ $< diff --git a/src/aiori-POSIX.c b/src/aiori-POSIX.c index 66c7e4e3..80ea3d06 100755 --- a/src/aiori-POSIX.c +++ b/src/aiori-POSIX.c @@ -53,15 +53,14 @@ #include "aiori-POSIX.h" #ifdef HAVE_GPU_DIRECT -typedef long long loff_t; -# include -# include +# include "gpu_io.h" #endif typedef struct { int fd; #ifdef HAVE_GPU_DIRECT - CUfileHandle_t cf_handle; + gpu_io_file_t *gpu_file; + int gpu_file_registered; #endif } posix_fd; @@ -79,22 +78,20 @@ typedef struct { #endif #ifdef HAVE_GPU_DIRECT -static const char* cuFileGetErrorString(CUfileError_t status){ - if(IS_CUDA_ERR(status)){ - return cudaGetErrorString(status.err); - } - return strerror(status.err); -} +static void init_gpu_direct(posix_fd * pfd){ + char errbuf[256]; + gpu_io_status_t status; -static void init_cufile(posix_fd * pfd){ - CUfileDescr_t cf_descr = (CUfileDescr_t){ - .handle.fd = pfd->fd, - .type = CU_FILE_HANDLE_TYPE_OPAQUE_FD - }; - CUfileError_t status = cuFileHandleRegister(& pfd->cf_handle, & cf_descr); - if(status.err != CU_FILE_SUCCESS){ - WARNF("Could not register handle %s", cuFileGetErrorString(status)); + pfd->gpu_file = NULL; + pfd->gpu_file_registered = 0; + + status = gpu_io_register_fd(& pfd->gpu_file, pfd->fd); + if(! gpu_io_status_ok(status)){ + ERRF("Could not register GPU Direct handle: %s", + gpu_io_strerror(status, errbuf, sizeof(errbuf))); } + + pfd->gpu_file_registered = 1; } #endif @@ -643,7 +640,7 @@ aiori_fd_t *POSIX_Create(char *testFileName, int flags, aiori_mod_opt_t * param) #endif #ifdef HAVE_GPU_DIRECT if(o->gpuDirect){ - init_cufile(pfd); + init_gpu_direct(pfd); } #endif return (aiori_fd_t*) pfd; @@ -708,7 +705,7 @@ aiori_fd_t *POSIX_Open(char *testFileName, int flags, aiori_mod_opt_t * param) #endif #ifdef HAVE_GPU_DIRECT if(o->gpuDirect){ - init_cufile(pfd); + init_gpu_direct(pfd); } #endif return (aiori_fd_t*) pfd; @@ -766,16 +763,25 @@ static IOR_offset_t POSIX_Xfer(int access, aiori_fd_t *file, IOR_size_t * buffer } #ifdef HAVE_GPU_DIRECT if(o->gpuDirect){ - rc = cuFileWrite(pfd->cf_handle, ptr, remaining, offset + mem_offset, mem_offset); + gpu_io_result_t gio_res; + char errbuf[256]; + gio_res = gpu_io_write(pfd->gpu_file, ptr, remaining, + offset + mem_offset, mem_offset); + rc = gio_res.nbytes; + if(rc < 0){ + WARNF("gpu_io_write(%d, %p, %lld) failed: %s", + fd, (void*)ptr, remaining, + gpu_io_strerror(gio_res.status, errbuf, sizeof(errbuf))); + } }else{ #endif rc = write(fd, ptr, remaining); + if (rc < 0){ + WARNF("write(%d, %p, %lld) failed %s", fd, (void*)ptr, remaining, strerror(errno)); + } #ifdef HAVE_GPU_DIRECT } #endif - if (rc < 0){ - WARNF("write(%d, %p, %lld) failed %s", fd, (void*)ptr, remaining, strerror(errno)); - } if (hints->fsyncPerWrite == TRUE){ POSIX_Fsync((aiori_fd_t*) &fd, param); } @@ -787,7 +793,16 @@ static IOR_offset_t POSIX_Xfer(int access, aiori_fd_t *file, IOR_size_t * buffer } #ifdef HAVE_GPU_DIRECT if(o->gpuDirect){ - rc = cuFileRead(pfd->cf_handle, ptr, remaining, offset + mem_offset, mem_offset); + gpu_io_result_t gio_res; + char errbuf[256]; + gio_res = gpu_io_read(pfd->gpu_file, ptr, remaining, + offset + mem_offset, mem_offset); + rc = gio_res.nbytes; + if(rc < 0){ + WARNF("gpu_io_read(%d, %p, %lld) failed: %s", + fd, (void*)ptr, remaining, + gpu_io_strerror(gio_res.status, errbuf, sizeof(errbuf))); + } }else{ #endif rc = read(fd, ptr, remaining); @@ -800,7 +815,10 @@ static IOR_offset_t POSIX_Xfer(int access, aiori_fd_t *file, IOR_size_t * buffer } if (rc < 0){ - WARNF("read(%d, %p, %lld) failed %s", fd, (void*)ptr, remaining, strerror(errno)); +#ifdef HAVE_GPU_DIRECT + if(!o->gpuDirect) +#endif + WARNF("read(%d, %p, %lld) failed %s", fd, (void*)ptr, remaining, strerror(errno)); return length - remaining; } } @@ -815,7 +833,9 @@ static IOR_offset_t POSIX_Xfer(int access, aiori_fd_t *file, IOR_size_t * buffer return length - remaining; } } - assert(rc >= 0); + if (rc < 0) { + return length - remaining; + } assert(rc <= remaining); remaining -= rc; ptr += rc; @@ -868,8 +888,8 @@ void POSIX_Close(aiori_fd_t *afd, aiori_mod_opt_t * param) posix_options_t * o = (posix_options_t*) param; int fd = ((posix_fd*) afd)->fd; #ifdef HAVE_GPU_DIRECT - if(o->gpuDirect){ - cuFileHandleDeregister(((posix_fd*) afd)->cf_handle); + if(o->gpuDirect && ((posix_fd*) afd)->gpu_file_registered){ + gpu_io_deregister_fd(& ((posix_fd*) afd)->gpu_file); } #endif if (close(fd) != 0){ @@ -920,13 +940,7 @@ IOR_offset_t POSIX_GetFileSize(aiori_mod_opt_t * test, char *testFileName) } void POSIX_Initialize(aiori_mod_opt_t * options){ -#ifdef HAVE_GPU_DIRECT - CUfileError_t err = cuFileDriverOpen(); -#endif } void POSIX_Finalize(aiori_mod_opt_t * options){ -#ifdef HAVE_GPU_DIRECT - CUfileError_t err = cuFileDriverClose(); -#endif } diff --git a/src/cuda_runtime.c b/src/cuda_runtime.c new file mode 100644 index 00000000..74fd03f3 --- /dev/null +++ b/src/cuda_runtime.c @@ -0,0 +1,89 @@ +#ifdef HAVE_CONFIG_H +# include "config.h" +#endif + +#ifdef HAVE_CUDA + +#include +#include "gpu_runtime.h" + +int +gpu_runtime_ok(gpu_runtime_status_t status) +{ + return status.code == 0; +} + +const char * +gpu_runtime_strerror(gpu_runtime_status_t status) +{ + return cudaGetErrorString((cudaError_t) status.code); +} + +gpu_runtime_status_t +gpu_runtime_get_device_count(int *count) +{ + gpu_runtime_status_t st; + st.code = (int) cudaGetDeviceCount(count); + return st; +} + +gpu_runtime_status_t +gpu_runtime_set_device(int device) +{ + gpu_runtime_status_t st; + st.code = (int) cudaSetDevice(device); + return st; +} + +gpu_runtime_status_t +gpu_runtime_malloc(void **ptr, size_t size) +{ + gpu_runtime_status_t st; + st.code = (int) cudaMalloc(ptr, size); + return st; +} + +gpu_runtime_status_t +gpu_runtime_malloc_managed(void **ptr, size_t size) +{ + gpu_runtime_status_t st; + st.code = (int) cudaMallocManaged(ptr, size, cudaMemAttachGlobal); + return st; +} + +gpu_runtime_status_t +gpu_runtime_free(void *ptr) +{ + gpu_runtime_status_t st; + st.code = (int) cudaFree(ptr); + return st; +} + +gpu_runtime_status_t +gpu_runtime_memset(void *ptr, int value, size_t count) +{ + gpu_runtime_status_t st; + st.code = (int) cudaMemset(ptr, value, count); + return st; +} + +gpu_runtime_status_t +gpu_runtime_memcpy(void *dst, const void *src, size_t count, + gpu_memcpy_kind_t kind) +{ + gpu_runtime_status_t st; + enum cudaMemcpyKind cuda_kind; + + switch (kind) { + case GPU_MEMCPY_HOST_TO_DEVICE: cuda_kind = cudaMemcpyHostToDevice; break; + case GPU_MEMCPY_DEVICE_TO_HOST: cuda_kind = cudaMemcpyDeviceToHost; break; + case GPU_MEMCPY_DEVICE_TO_DEVICE: cuda_kind = cudaMemcpyDeviceToDevice; break; + case GPU_MEMCPY_HOST_TO_HOST: cuda_kind = cudaMemcpyHostToHost; break; + default: cuda_kind = cudaMemcpyDefault; break; + } + + st.code = (int) cudaMemcpy(dst, src, count, cuda_kind); + return st; +} + +#endif diff --git a/src/gpu_io.h b/src/gpu_io.h new file mode 100644 index 00000000..0b3217d6 --- /dev/null +++ b/src/gpu_io.h @@ -0,0 +1,48 @@ +#ifndef GPU_IO_H +#define GPU_IO_H + +#include +#include +#include + +#ifdef __cplusplus +extern "C" { +#endif + +typedef struct gpu_io_file gpu_io_file_t; + +typedef struct { + int ok; + int errnum; + long backend_error; + long runtime_error; + const char *message; +} gpu_io_status_t; + +typedef struct { + ssize_t nbytes; + gpu_io_status_t status; +} gpu_io_result_t; + +int gpu_io_status_ok(gpu_io_status_t status); +const char *gpu_io_strerror(gpu_io_status_t status, char *buf, size_t buflen); + +gpu_io_status_t gpu_io_driver_open(void); +gpu_io_status_t gpu_io_driver_close(void); + +gpu_io_status_t gpu_io_register_fd(gpu_io_file_t **file, int fd); +void gpu_io_deregister_fd(gpu_io_file_t **file); + +gpu_io_result_t gpu_io_read(gpu_io_file_t *file, void *buffer, + size_t size, int64_t file_offset, + int64_t buffer_offset); + +gpu_io_result_t gpu_io_write(gpu_io_file_t *file, const void *buffer, + size_t size, int64_t file_offset, + int64_t buffer_offset); + +#ifdef __cplusplus +} +#endif + +#endif diff --git a/src/gpu_io_cufile.c b/src/gpu_io_cufile.c new file mode 100644 index 00000000..2a04085c --- /dev/null +++ b/src/gpu_io_cufile.c @@ -0,0 +1,209 @@ +#ifdef HAVE_CONFIG_H +# include "config.h" +#endif + +#ifdef HAVE_CUFILE + +#include +#include +#include +#include + +#include +#include + +#include "gpu_io.h" +#include "aiori-debug.h" + +struct gpu_io_file { + CUfileHandle_t handle; +}; + +int +gpu_io_status_ok(gpu_io_status_t status) +{ + return status.ok; +} + +const char * +gpu_io_strerror(gpu_io_status_t status, char *buf, size_t buflen) +{ + if (status.ok) + return "success"; + + if (status.message) + snprintf(buf, buflen, "%s", status.message); + else if (status.runtime_error) + snprintf(buf, buflen, "CUDA error: %s", + cudaGetErrorString((cudaError_t) status.runtime_error)); + else if (status.backend_error) + snprintf(buf, buflen, "cuFile error: %s (err=%ld)", + strerror((int) status.backend_error), status.backend_error); + else if (status.errnum) + snprintf(buf, buflen, "POSIX error: %s", strerror(status.errnum)); + else + snprintf(buf, buflen, "unknown GPU I/O error"); + + return buf; +} + +static gpu_io_status_t +cufile_to_status(CUfileError_t err) +{ + gpu_io_status_t st; + memset(&st, 0, sizeof(st)); + + if (err.err == CU_FILE_SUCCESS) { + st.ok = 1; + return st; + } + + st.ok = 0; + + if (IS_CUDA_ERR(err)) { + st.runtime_error = (long) err.err; + st.message = cudaGetErrorString(err.err); + } else { + st.backend_error = (long) err.err; + st.message = strerror(err.err); + } + + return st; +} + +gpu_io_status_t +gpu_io_driver_open(void) +{ + CUfileError_t err = cuFileDriverOpen(); + return cufile_to_status(err); +} + +gpu_io_status_t +gpu_io_driver_close(void) +{ + CUfileError_t err = cuFileDriverClose(); + return cufile_to_status(err); +} + +static int gpu_io_driver_opened = 0; + +gpu_io_status_t +gpu_io_register_fd(gpu_io_file_t **file, int fd) +{ + CUfileDescr_t descr; + CUfileError_t err; + + if (!gpu_io_driver_opened) { + gpu_io_status_t open_st = gpu_io_driver_open(); + if (!gpu_io_status_ok(open_st)) + return open_st; + gpu_io_driver_opened = 1; + } + + *file = (gpu_io_file_t *) malloc(sizeof(gpu_io_file_t)); + if (*file == NULL) { + gpu_io_status_t st; + memset(&st, 0, sizeof(st)); + st.ok = 0; + st.errnum = ENOMEM; + return st; + } + + memset(&descr, 0, sizeof(descr)); + descr.handle.fd = fd; + descr.type = CU_FILE_HANDLE_TYPE_OPAQUE_FD; + + err = cuFileHandleRegister(&(*file)->handle, &descr); + if (err.err != CU_FILE_SUCCESS) { + gpu_io_status_t st = cufile_to_status(err); + free(*file); + *file = NULL; + return st; + } + + { + gpu_io_status_t ok; + memset(&ok, 0, sizeof(ok)); + ok.ok = 1; + return ok; + } +} + +void +gpu_io_deregister_fd(gpu_io_file_t **file) +{ + if (file == NULL || *file == NULL) + return; + + cuFileHandleDeregister((*file)->handle); + free(*file); + *file = NULL; +} + +/* + * cuFileRead/Write return: + * >= 0 : bytes transferred + * < 0 : negative CUfileOpError_t. Use CUFILE_ERRSTR(-rc). + */ +static void +cufile_xfer_decode(gpu_io_result_t *res) +{ + if (res->nbytes < 0) { + int raw_err = (int) -res->nbytes; + + res->status.ok = 0; + res->status.backend_error = (long) raw_err; +#ifdef CUFILE_ERRSTR + res->status.message = CUFILE_ERRSTR(raw_err); +#else + res->status.message = strerror(raw_err); +#endif + res->nbytes = -1; + } else { + res->status.ok = 1; + } +} + +gpu_io_result_t +gpu_io_read(gpu_io_file_t *file, void *buffer, + size_t size, int64_t file_offset, int64_t buffer_offset) +{ + gpu_io_result_t res; + memset(&res, 0, sizeof(res)); + + if (file == NULL) { + res.nbytes = -1; + res.status.ok = 0; + res.status.errnum = EBADF; + return res; + } + + res.nbytes = (ssize_t) cuFileRead(file->handle, buffer, size, + file_offset, buffer_offset); + cufile_xfer_decode(&res); + + return res; +} + +gpu_io_result_t +gpu_io_write(gpu_io_file_t *file, const void *buffer, + size_t size, int64_t file_offset, int64_t buffer_offset) +{ + gpu_io_result_t res; + memset(&res, 0, sizeof(res)); + + if (file == NULL) { + res.nbytes = -1; + res.status.ok = 0; + res.status.errnum = EBADF; + return res; + } + + res.nbytes = (ssize_t) cuFileWrite(file->handle, buffer, size, + file_offset, buffer_offset); + cufile_xfer_decode(&res); + + return res; +} + +#endif diff --git a/src/gpu_io_hipfile.c b/src/gpu_io_hipfile.c new file mode 100644 index 00000000..8b1a39a6 --- /dev/null +++ b/src/gpu_io_hipfile.c @@ -0,0 +1,215 @@ +/* + * hipFile return values: + * >= 0 : bytes transferred + * == -1 : POSIX error (errno) + * < -1 : -(hipFileOpError_t), use HIPFILE_ERRSTR(-rc) + */ +#ifdef HAVE_CONFIG_H +# include "config.h" +#endif + +#ifdef HAVE_HIPFILE + +#include +#include +#include +#include + +#include +#include + +#include "gpu_io.h" +#include "aiori-debug.h" + +struct gpu_io_file { + hipFileHandle_t handle; +}; + +int +gpu_io_status_ok(gpu_io_status_t status) +{ + return status.ok; +} + +const char * +gpu_io_strerror(gpu_io_status_t status, char *buf, size_t buflen) +{ + if (status.ok) + return "success"; + + if (status.message) + snprintf(buf, buflen, "%s", status.message); + else if (status.runtime_error) + snprintf(buf, buflen, "HIP error: %s", + hipGetErrorString((hipError_t) status.runtime_error)); + else if (status.backend_error) + snprintf(buf, buflen, "hipFile error: %s (err=%ld)", + HIPFILE_ERRSTR((int) status.backend_error), + status.backend_error); + else if (status.errnum) + snprintf(buf, buflen, "POSIX error: %s", strerror(status.errnum)); + else + snprintf(buf, buflen, "unknown GPU I/O error"); + + return buf; +} + +static gpu_io_status_t +hipfile_to_status(hipFileError_t err) +{ + gpu_io_status_t st; + memset(&st, 0, sizeof(st)); + + if (err.err == hipFileSuccess) { + st.ok = 1; + return st; + } + + st.ok = 0; + + if (IS_HIP_DRV_ERR(err)) { + st.runtime_error = (long) HIP_DRV_ERR(err); + st.message = hipGetErrorString(HIP_DRV_ERR(err)); + } else { + st.backend_error = (long) err.err; + st.message = HIPFILE_ERRSTR(err.err); + } + + return st; +} + +gpu_io_status_t +gpu_io_driver_open(void) +{ + hipFileError_t err = hipFileDriverOpen(); + return hipfile_to_status(err); +} + +gpu_io_status_t +gpu_io_driver_close(void) +{ + hipFileError_t err = hipFileDriverClose(); + return hipfile_to_status(err); +} + +/* hipFileHandleRegister() auto-initializes the library on first call. */ +gpu_io_status_t +gpu_io_register_fd(gpu_io_file_t **file, int fd) +{ + hipFileDescr_t descr; + hipFileError_t err; + + *file = (gpu_io_file_t *) malloc(sizeof(gpu_io_file_t)); + if (*file == NULL) { + gpu_io_status_t st; + memset(&st, 0, sizeof(st)); + st.ok = 0; + st.errnum = ENOMEM; + return st; + } + + memset(&descr, 0, sizeof(descr)); + descr.type = hipFileHandleTypeOpaqueFD; + descr.handle.fd = fd; + descr.fs_ops = NULL; + + err = hipFileHandleRegister(&(*file)->handle, &descr); + if (err.err != hipFileSuccess) { + gpu_io_status_t st = hipfile_to_status(err); + free(*file); + *file = NULL; + return st; + } + + { + gpu_io_status_t ok; + memset(&ok, 0, sizeof(ok)); + ok.ok = 1; + return ok; + } +} + +void +gpu_io_deregister_fd(gpu_io_file_t **file) +{ + if (file == NULL || *file == NULL) + return; + + hipFileHandleDeregister((*file)->handle); + free(*file); + *file = NULL; +} + +gpu_io_result_t +gpu_io_read(gpu_io_file_t *file, void *buffer, + size_t size, int64_t file_offset, int64_t buffer_offset) +{ + gpu_io_result_t res; + memset(&res, 0, sizeof(res)); + + if (file == NULL) { + res.nbytes = -1; + res.status.ok = 0; + res.status.errnum = EBADF; + return res; + } + + res.nbytes = (ssize_t) hipFileRead(file->handle, buffer, size, + (hoff_t) file_offset, + (hoff_t) buffer_offset); + + if (res.nbytes == -1) { + res.status.ok = 0; + res.status.errnum = errno; + res.status.message = strerror(errno); + res.nbytes = -1; + } else if (res.nbytes < -1) { + int raw_err = (int) -res.nbytes; + res.status.ok = 0; + res.status.backend_error = (long) raw_err; + res.status.message = HIPFILE_ERRSTR(raw_err); + res.nbytes = -1; + } else { + res.status.ok = 1; + } + + return res; +} + +gpu_io_result_t +gpu_io_write(gpu_io_file_t *file, const void *buffer, + size_t size, int64_t file_offset, int64_t buffer_offset) +{ + gpu_io_result_t res; + memset(&res, 0, sizeof(res)); + + if (file == NULL) { + res.nbytes = -1; + res.status.ok = 0; + res.status.errnum = EBADF; + return res; + } + + res.nbytes = (ssize_t) hipFileWrite(file->handle, buffer, size, + (hoff_t) file_offset, + (hoff_t) buffer_offset); + + if (res.nbytes == -1) { + res.status.ok = 0; + res.status.errnum = errno; + res.status.message = strerror(errno); + res.nbytes = -1; + } else if (res.nbytes < -1) { + int raw_err = (int) -res.nbytes; + res.status.ok = 0; + res.status.backend_error = (long) raw_err; + res.status.message = HIPFILE_ERRSTR(raw_err); + res.nbytes = -1; + } else { + res.status.ok = 1; + } + + return res; +} + +#endif diff --git a/src/gpu_runtime.h b/src/gpu_runtime.h new file mode 100644 index 00000000..0ed83ab2 --- /dev/null +++ b/src/gpu_runtime.h @@ -0,0 +1,38 @@ +#ifndef GPU_RUNTIME_H +#define GPU_RUNTIME_H + +#include + +#ifdef __cplusplus +extern "C" { +#endif + +typedef enum { + GPU_MEMCPY_HOST_TO_DEVICE, + GPU_MEMCPY_DEVICE_TO_HOST, + GPU_MEMCPY_DEVICE_TO_DEVICE, + GPU_MEMCPY_HOST_TO_HOST +} gpu_memcpy_kind_t; + +typedef struct { + int code; +} gpu_runtime_status_t; + +int gpu_runtime_ok(gpu_runtime_status_t status); +const char *gpu_runtime_strerror(gpu_runtime_status_t status); + +gpu_runtime_status_t gpu_runtime_get_device_count(int *count); +gpu_runtime_status_t gpu_runtime_set_device(int device); + +gpu_runtime_status_t gpu_runtime_malloc(void **ptr, size_t size); +gpu_runtime_status_t gpu_runtime_malloc_managed(void **ptr, size_t size); +gpu_runtime_status_t gpu_runtime_free(void *ptr); +gpu_runtime_status_t gpu_runtime_memset(void *ptr, int value, size_t count); +gpu_runtime_status_t gpu_runtime_memcpy(void *dst, const void *src, + size_t count, gpu_memcpy_kind_t kind); + +#ifdef __cplusplus +} +#endif + +#endif diff --git a/src/hip_runtime.c b/src/hip_runtime.c new file mode 100644 index 00000000..13fbb662 --- /dev/null +++ b/src/hip_runtime.c @@ -0,0 +1,89 @@ +#ifdef HAVE_CONFIG_H +# include "config.h" +#endif + +#ifdef HAVE_HIP + +#include +#include "gpu_runtime.h" + +int +gpu_runtime_ok(gpu_runtime_status_t status) +{ + return status.code == 0; +} + +const char * +gpu_runtime_strerror(gpu_runtime_status_t status) +{ + return hipGetErrorString((hipError_t) status.code); +} + +gpu_runtime_status_t +gpu_runtime_get_device_count(int *count) +{ + gpu_runtime_status_t st; + st.code = (int) hipGetDeviceCount(count); + return st; +} + +gpu_runtime_status_t +gpu_runtime_set_device(int device) +{ + gpu_runtime_status_t st; + st.code = (int) hipSetDevice(device); + return st; +} + +gpu_runtime_status_t +gpu_runtime_malloc(void **ptr, size_t size) +{ + gpu_runtime_status_t st; + st.code = (int) hipMalloc(ptr, size); + return st; +} + +gpu_runtime_status_t +gpu_runtime_malloc_managed(void **ptr, size_t size) +{ + gpu_runtime_status_t st; + st.code = (int) hipMallocManaged(ptr, size, hipMemAttachGlobal); + return st; +} + +gpu_runtime_status_t +gpu_runtime_free(void *ptr) +{ + gpu_runtime_status_t st; + st.code = (int) hipFree(ptr); + return st; +} + +gpu_runtime_status_t +gpu_runtime_memset(void *ptr, int value, size_t count) +{ + gpu_runtime_status_t st; + st.code = (int) hipMemset(ptr, value, count); + return st; +} + +gpu_runtime_status_t +gpu_runtime_memcpy(void *dst, const void *src, size_t count, + gpu_memcpy_kind_t kind) +{ + gpu_runtime_status_t st; + enum hipMemcpyKind hip_kind; + + switch (kind) { + case GPU_MEMCPY_HOST_TO_DEVICE: hip_kind = hipMemcpyHostToDevice; break; + case GPU_MEMCPY_DEVICE_TO_HOST: hip_kind = hipMemcpyDeviceToHost; break; + case GPU_MEMCPY_DEVICE_TO_DEVICE: hip_kind = hipMemcpyDeviceToDevice; break; + case GPU_MEMCPY_HOST_TO_HOST: hip_kind = hipMemcpyHostToHost; break; + default: hip_kind = hipMemcpyDefault; break; + } + + st.code = (int) hipMemcpy(dst, src, count, hip_kind); + return st; +} + +#endif diff --git a/src/ior.c b/src/ior.c index 43bb1395..5f701cd2 100755 --- a/src/ior.c +++ b/src/ior.c @@ -33,8 +33,8 @@ # include /* uname() */ #endif -#ifdef HAVE_CUDA -#include +#ifdef HAVE_GPU +#include "gpu_runtime.h" #endif #include @@ -126,7 +126,7 @@ static int test_initialize(IOR_test_t * test){ backend = test->params.backend; if(test->params.gpuMemoryFlags != IOR_MEMORY_TYPE_CPU){ - initCUDA(test->params.tasksBlockMapping, rank, test->params.numNodes, test->params.numTasksOnNode0, test->params.gpuID); + initGPU(test->params.tasksBlockMapping, rank, test->params.numNodes, test->params.numTasksOnNode0, test->params.gpuID); } diff --git a/src/md-workbench.c b/src/md-workbench.c index 95420ca6..6b3a59a6 100644 --- a/src/md-workbench.c +++ b/src/md-workbench.c @@ -832,7 +832,7 @@ static option_help options [] = { {'W', "stonewall-wear-out", "Stop with stonewall after specified time and use a soft wear-out phase -- all processes perform the same number of iterations", OPTION_FLAG, 'd', & o.stonewall_timer_wear_out}, {'X', "verify-read", "Verify the data on read", OPTION_FLAG, 'd', & o.verify_read}, {0, "dataPacketType", "type of packet that will be created [offset|incompressible|timestamp|random|o|i|t|r]", OPTION_OPTIONAL_ARGUMENT, 's', & o.packetTypeStr}, -#ifdef HAVE_CUDA +#ifdef HAVE_GPU {0, "allocateBufferOnGPU", "Allocate I/O buffers on the GPU: X=1 uses managed memory - verifications are run on CPU; X=2 managed memory - verifications on GPU; X=3 device memory with verifications on GPU.", OPTION_OPTIONAL_ARGUMENT, 'd', & o.gpuMemoryFlags}, {0, "GPUid", "Select the GPU to use, use -1 for round-robin among local procs.", OPTION_OPTIONAL_ARGUMENT, 'd', & o.gpuID}, #ifdef HAVE_GPU_DIRECT @@ -954,7 +954,7 @@ mdworkbench_results_t* md_workbench_run(int argc, char ** argv, MPI_Comm world_c int numNodes = GetNumNodes(o.com); int numTasksOnNode0 = GetNumTasksOnNode0(o.com); if(o.gpuMemoryFlags != IOR_MEMORY_TYPE_CPU){ - initCUDA(tasksBlockMapping, o.rank, numNodes, numTasksOnNode0, o.gpuID); + initGPU(tasksBlockMapping, o.rank, numNodes, numTasksOnNode0, o.gpuID); } int current_index = 0; diff --git a/src/mdtest.c b/src/mdtest.c index 0631ba67..458cc2a7 100644 --- a/src/mdtest.c +++ b/src/mdtest.c @@ -2359,7 +2359,7 @@ mdtest_results_t * mdtest_run(int argc, char **argv, MPI_Comm world_com, FILE * {0, "dataPacketType", "type of packet that will be created [offset|incompressible|timestamp|random|o|i|t|r]", OPTION_OPTIONAL_ARGUMENT, 's', & packetType}, {0, "run-cmd-before-phase", "call this external command before each phase (excluded from the timing)", OPTION_OPTIONAL_ARGUMENT, 's', & o.prologue}, {0, "run-cmd-after-phase", "call this external command after each phase (included in the timing)", OPTION_OPTIONAL_ARGUMENT, 's', & o.epilogue}, -#ifdef HAVE_CUDA +#ifdef HAVE_GPU {0, "allocateBufferOnGPU", "Allocate I/O buffers on the GPU: X=1 uses managed memory - verifications are run on CPU; X=2 managed memory - verifications on GPU; X=3 device memory with verifications on GPU.", OPTION_OPTIONAL_ARGUMENT, 'd', & o.gpuMemoryFlags}, {0, "GPUid", "Select the GPU to use, use -1 for round-robin among local procs.", OPTION_OPTIONAL_ARGUMENT, 'd', & o.gpuID}, #ifdef HAVE_GPU_DIRECT @@ -2478,7 +2478,7 @@ mdtest_results_t * mdtest_run(int argc, char **argv, MPI_Comm world_com, FILE * int tasksBlockMapping = QueryNodeMapping(testComm, true); if(o.gpuMemoryFlags != IOR_MEMORY_TYPE_CPU){ - initCUDA(tasksBlockMapping, rank, numNodes, numTasksOnNode0, o.gpuID); + initGPU(tasksBlockMapping, rank, numNodes, numTasksOnNode0, o.gpuID); } /* setup total number of items and number of items per dir */ diff --git a/src/parse_options.c b/src/parse_options.c index 084f3c52..32ef090e 100755 --- a/src/parse_options.c +++ b/src/parse_options.c @@ -426,7 +426,7 @@ option_help * createGlobalOptions(IOR_param_t * params){ {.help=" -O stoneWallingWearOutIterations=N -- stop after processing this number of iterations, needed for reading data back written with stoneWallingWearOut", .arg = OPTION_OPTIONAL_ARGUMENT}, {.help=" -O stoneWallingStatusFile=FILE -- this file keeps the number of iterations from stonewalling during write and allows to use them for read", .arg = OPTION_OPTIONAL_ARGUMENT}, {.help=" -O minTimeDuration=0 -- minimum Runtime for the run (will repeat from beginning of the file if time is not yet over)", .arg = OPTION_OPTIONAL_ARGUMENT}, -#ifdef HAVE_CUDA +#ifdef HAVE_GPU {.help=" -O allocateBufferOnGPU=X -- allocate I/O buffers on the GPU: X=1 uses managed memory - verifications are run on CPU; X=2 managed memory - verifications on GPU; X=3 device memory with verifications on GPU.", .arg = OPTION_OPTIONAL_ARGUMENT}, {.help=" -O GPUid=X -- select the GPU to use, use -1 for round-robin among local procs.", .arg = OPTION_OPTIONAL_ARGUMENT}, #ifdef HAVE_GPU_DIRECT diff --git a/src/utilities-gpu.cu b/src/utilities-gpu.cu index 75a71c3d..085b6665 100644 --- a/src/utilities-gpu.cu +++ b/src/utilities-gpu.cu @@ -1,14 +1,15 @@ /* - This file contains CUDA code for creating and checking memory patterns on the device. + This file contains GPU code for creating and checking memory patterns + on the device. Built with nvcc under HAVE_CUDA. */ -#include - #ifdef HAVE_CONFIG_H #include "config.h" #endif #include +#include +#include "gpu_runtime.h" #include "iordef.h" #define RANDALGO_GOLDEN_RATIO_PRIME 0x9e37fffffffc0001UL @@ -69,14 +70,14 @@ extern "C" int verify_memory_pattern_gpu(uint64_t item, char * buffer, size_t by size_t blocks = (bytes+2047)/2048; size_t threads = 256; int * derror_found; - cudaMalloc(&derror_found, sizeof(int)); - cudaMemcpy(derror_found, & errors, sizeof(int), cudaMemcpyHostToDevice); + gpu_runtime_malloc((void**) & derror_found, sizeof(int)); + gpu_runtime_memcpy(derror_found, & errors, sizeof(int), GPU_MEMCPY_HOST_TO_DEVICE); if(dataPacketType == DATA_TIMESTAMP){ cu_verify_memory_timestamp<<>>(item, (uint64_t*) buffer, bytes/sizeof(uint64_t), rand_seed, ((uint64_t) pretendRank) << 32, derror_found); }else if(dataPacketType == DATA_INCOMPRESSIBLE){ } - cudaMemcpy(& errors, derror_found, sizeof(int), cudaMemcpyDeviceToHost); - cudaFree(derror_found); + gpu_runtime_memcpy(& errors, derror_found, sizeof(int), GPU_MEMCPY_DEVICE_TO_HOST); + gpu_runtime_free(derror_found); return errors; } diff --git a/src/utilities-gpu.hip b/src/utilities-gpu.hip new file mode 100644 index 00000000..66f70358 --- /dev/null +++ b/src/utilities-gpu.hip @@ -0,0 +1,83 @@ +/* + This file contains GPU code for creating and checking memory patterns + on the device. Built with hipcc under HAVE_HIP. +*/ +#ifdef HAVE_CONFIG_H +#include "config.h" +#endif + +#include + +#include +#include "gpu_runtime.h" +#include "iordef.h" + +#define RANDALGO_GOLDEN_RATIO_PRIME 0x9e37fffffffc0001UL + +__global__ +void cu_generate_memory_timestamp(uint64_t * buf, size_t length, int rand_seed, uint64_t pretendRank){ + size_t pos = blockIdx.x * blockDim.x + threadIdx.x; + if(pos < length){ + buf[pos] = pretendRank | rand_seed + pos; + } +} + +__global__ +void cu_generate_memory_incompressible(uint64_t * buf, size_t length, uint64_t seed){ + size_t pos = blockIdx.x * blockDim.x + threadIdx.x; + if(pos < length){ + buf[pos] = seed | pos; + } +} + +__global__ +void cu_verify_memory_timestamp(uint64_t item, uint64_t * buf, size_t length, int rand_seed, uint64_t pretendRank, int * errors){ + size_t pos = blockIdx.x * blockDim.x + threadIdx.x; + if(pos < length){ + int correct = buf[pos] == (pretendRank | rand_seed + pos); + if(! correct){ + *errors = 1; // it isn't thread safe but one error reported is enough + } + } +} + +extern "C" void generate_memory_pattern_gpu(char * buf, size_t bytes, int rand_seed, int pretendRank, ior_dataPacketType_e dataPacketType){ + size_t blocks = (bytes+2047)/2048; + size_t threads = 256; + switch(dataPacketType){ + case(DATA_RANDOM): + // Nothing to do, will work on updates + break; + case(DATA_INCOMPRESSIBLE):{ + cu_generate_memory_incompressible<<>>((uint64_t*) buf, bytes/sizeof(uint64_t), rand_seed + pretendRank); + break; + }case(DATA_OFFSET):{ + }case(DATA_TIMESTAMP):{ + cu_generate_memory_timestamp<<>>((uint64_t*) buf, bytes/sizeof(uint64_t), rand_seed, ((uint64_t) pretendRank) << 32); + break; + } + } +} + +extern "C" void update_write_memory_pattern_gpu(uint64_t item, char * buf, size_t bytes, int rand_seed, int rank, ior_dataPacketType_e dataPacketType){ + // nothing to do for dataPacketType == DATA_TIMESTAMP, i.e., won't be called for this parameter + size_t blocks = (bytes+2047)/2048; + size_t threads = 256; +} + +extern "C" int verify_memory_pattern_gpu(uint64_t item, char * buffer, size_t bytes, int rand_seed, int pretendRank, ior_dataPacketType_e dataPacketType){ + int errors = 0; + size_t blocks = (bytes+2047)/2048; + size_t threads = 256; + int * derror_found; + gpu_runtime_malloc((void**) & derror_found, sizeof(int)); + gpu_runtime_memcpy(derror_found, & errors, sizeof(int), GPU_MEMCPY_HOST_TO_DEVICE); + if(dataPacketType == DATA_TIMESTAMP){ + cu_verify_memory_timestamp<<>>(item, (uint64_t*) buffer, bytes/sizeof(uint64_t), rand_seed, ((uint64_t) pretendRank) << 32, derror_found); + }else if(dataPacketType == DATA_INCOMPRESSIBLE){ + + } + gpu_runtime_memcpy(& errors, derror_found, sizeof(int), GPU_MEMCPY_DEVICE_TO_HOST); + gpu_runtime_free(derror_found); + return errors; +} diff --git a/src/utilities.c b/src/utilities.c index 2228e5ef..345d6d54 100755 --- a/src/utilities.c +++ b/src/utilities.c @@ -37,8 +37,8 @@ #include #include -#ifdef HAVE_CUDA -#include +#ifdef HAVE_GPU +#include "gpu_runtime.h" #endif #ifndef _WIN32 @@ -95,7 +95,7 @@ void update_write_memory_pattern(uint64_t item, char * buf, size_t bytes, int ra if (dataPacketType == DATA_TIMESTAMP || bytes < 8) return; -#ifdef HAVE_GPU_DIRECT +#ifdef HAVE_GPU if(type == IOR_MEMORY_TYPE_GPU_DEVICE_ONLY || type == IOR_MEMORY_TYPE_GPU_MANAGED_CHECK_GPU){ update_write_memory_pattern_gpu(item, buf, bytes, rand_seed, pretendRank, dataPacketType); return; @@ -159,7 +159,7 @@ uint64_t rand64(void) { * @param dataPacketType identifier to designate pattern to fill buffer */ void generate_memory_pattern(char * buf, size_t bytes, int rand_seed, int pretendRank, ior_dataPacketType_e dataPacketType, ior_memory_flags type){ -#ifdef HAVE_GPU_DIRECT +#ifdef HAVE_GPU if(type == IOR_MEMORY_TYPE_GPU_DEVICE_ONLY || type == IOR_MEMORY_TYPE_GPU_MANAGED_CHECK_GPU){ generate_memory_pattern_gpu(buf, bytes, rand_seed, pretendRank, dataPacketType); return; @@ -187,7 +187,7 @@ void generate_memory_pattern(char * buf, size_t bytes, int rand_seed, int preten } } } - + for(size_t i=size*8; i < bytes; i++){ buf[i] = (char) i; } @@ -195,17 +195,17 @@ void generate_memory_pattern(char * buf, size_t bytes, int rand_seed, int preten void invalidate_buffer_pattern(char * buffer, size_t bytes, ior_memory_flags type){ if(type == IOR_MEMORY_TYPE_GPU_DEVICE_ONLY || type == IOR_MEMORY_TYPE_GPU_MANAGED_CHECK_GPU){ -#ifdef HAVE_GPU_DIRECT - cudaMemset(buffer, 0x42, bytes > 512 ? 512 : bytes); +#ifdef HAVE_GPU + gpu_runtime_memset(buffer, 0x42, bytes > 512 ? 512 : bytes); #endif }else{ buffer[0] = ~buffer[0]; // changes the buffer, no memset to reduce the memory pressure } } -int verify_memory_pattern(uint64_t item, char * buffer, size_t bytes, int rand_seed, int pretendRank, ior_dataPacketType_e dataPacketType, ior_memory_flags type){ +int verify_memory_pattern(uint64_t item, char * buffer, size_t bytes, int rand_seed, int pretendRank, ior_dataPacketType_e dataPacketType, ior_memory_flags type){ int error = 0; -#ifdef HAVE_GPU_DIRECT +#ifdef HAVE_GPU if(type == IOR_MEMORY_TYPE_GPU_DEVICE_ONLY || type == IOR_MEMORY_TYPE_GPU_MANAGED_CHECK_GPU){ error = verify_memory_pattern_gpu(item, buffer, bytes, rand_seed, pretendRank, dataPacketType); return error; @@ -213,17 +213,17 @@ int verify_memory_pattern(uint64_t item, char * buffer, size_t bytes, int rand_s #endif // always read all data to ensure that performance numbers stay the same uint64_t * buffi = (uint64_t*) buffer; - + // the first 8 bytes are set to item number - int k=1; - + int k=1; + uint64_t rand_state_local; unsigned seed = rand_seed + pretendRank + item; rand_state_local = rand_r(&seed); const size_t size = bytes / 8; for(size_t i=0; i < size; i++){ uint64_t exp; - + switch(dataPacketType){ case(DATA_RANDOM): rand_state_local *= RANDALGO_GOLDEN_RATIO_PRIME; @@ -255,7 +255,7 @@ int verify_memory_pattern(uint64_t item, char * buffer, size_t bytes, int rand_s error = 1; } } - + return error; } @@ -295,7 +295,7 @@ OpTimer* OpTimerInit(char * filename, int size){ void OpTimerFlush(OpTimer* ot){ if(ot == NULL) { return; - } + } for(int i=0; i < ot->pos; i++){ fprintf(ot->fd, "%.8e,%.8e,%e\n", ot->time[i], ot->value[i], ot->size/ot->value[i]); } @@ -305,7 +305,7 @@ void OpTimerFlush(OpTimer* ot){ void OpTimerValue(OpTimer* ot, double now, double runTime){ if(ot == NULL) { return; - } + } ot->time[ot->pos] = now; ot->value[ot->pos++] = runTime; if(ot->pos == OP_BUFFER_SIZE){ @@ -528,18 +528,14 @@ int QueryNodeMapping(MPI_Comm comm, int print_nodemap) { return ret; } -void initCUDA(int blockMapping, int rank, int numNodes, int tasksPerNode, int useGPUID){ -#ifdef HAVE_CUDA +void initGPU(int blockMapping, int rank, int numNodes, int tasksPerNode, int useGPUID){ +#ifdef HAVE_GPU int device_count; - cudaError_t cret = cudaGetDeviceCount(& device_count); - if(cret != cudaSuccess){ - ERRF("cudaGetDeviceCount() error: %d %s", (int) cret, cudaGetErrorString(cret)); - } - //if (rank == 0){ - // char val[20]; - // sprintf(val, "%d", device_count); - // PrintKeyVal("cudaDevices", val); - //} + gpu_runtime_status_t st = gpu_runtime_get_device_count(& device_count); + if(! gpu_runtime_ok(st)){ + ERRF("gpu_runtime_get_device_count() error: %s", + gpu_runtime_strerror(st)); + } // if set to -1 use round robin per task if(useGPUID == -1){ int device = 0; @@ -548,12 +544,13 @@ void initCUDA(int blockMapping, int rank, int numNodes, int tasksPerNode, int us }else{ device = (rank / numNodes) % device_count; } - cret = cudaSetDevice(device); + st = gpu_runtime_set_device(device); }else{ - cret = cudaSetDevice(useGPUID); - } - if(cret != cudaSuccess){ - WARNF("cudaSetDevice(%d) error: %s", useGPUID, cudaGetErrorString(cret)); + st = gpu_runtime_set_device(useGPUID); + } + if(! gpu_runtime_ok(st)){ + WARNF("gpu_runtime_set_device(%d) error: %s", useGPUID, + gpu_runtime_strerror(st)); } #endif } @@ -1167,23 +1164,23 @@ void *aligned_buffer_alloc(size_t size, ior_memory_flags type) char *aligned; if(type == IOR_MEMORY_TYPE_GPU_MANAGED_CHECK_CPU || type == IOR_MEMORY_TYPE_GPU_MANAGED_CHECK_GPU){ -#ifdef HAVE_CUDA - // use unified memory here to allow drop-in-replacement - if (cudaMallocManaged((void**) & buf, size, cudaMemAttachGlobal) != cudaSuccess){ - ERR("Cannot allocate buffer on GPU"); +#ifdef HAVE_GPU + /* use unified memory here to allow drop-in-replacement */ + if (! gpu_runtime_ok(gpu_runtime_malloc_managed((void**) & buf, size))){ + ERR("Cannot allocate managed buffer on GPU"); } return buf; #else - ERR("No CUDA supported, cannot allocate on the GPU"); + ERR("No GPU runtime supported, cannot allocate managed memory"); #endif }else if(type == IOR_MEMORY_TYPE_GPU_DEVICE_ONLY){ -#ifdef HAVE_GPU_DIRECT - if (cudaMalloc((void**) & buf, size) != cudaSuccess){ +#ifdef HAVE_GPU + if (! gpu_runtime_ok(gpu_runtime_malloc((void**) & buf, size))){ ERR("Cannot allocate buffer on GPU"); } return buf; #else - ERR("No GPUDirect supported, cannot allocate on the GPU"); + ERR("No GPU runtime supported, cannot allocate on the GPU"); #endif } @@ -1212,13 +1209,13 @@ void *aligned_buffer_alloc(size_t size, ior_memory_flags type) void aligned_buffer_free(void *buf, ior_memory_flags gpu) { if(gpu){ -#ifdef HAVE_CUDA - if (cudaFree(buf) != cudaSuccess){ +#ifdef HAVE_GPU + if (! gpu_runtime_ok(gpu_runtime_free(buf))){ WARN("Cannot free buffer on GPU"); } return; #else - ERR("No CUDA supported, cannot free on the GPU"); + ERR("No GPU runtime supported, cannot free on the GPU"); #endif } free(*(void **)((char *)buf - sizeof(char *))); diff --git a/src/utilities.h b/src/utilities.h index 025ed726..1c334394 100755 --- a/src/utilities.h +++ b/src/utilities.h @@ -48,7 +48,7 @@ int verify_memory_pattern(uint64_t item, char * buffer, size_t bytes, int rand_s int verify_memory_pattern_gpu(uint64_t item, char * buffer, size_t bytes, int rand_seed, int pretendRank, ior_dataPacketType_e dataPacketType); void PrintKeyVal(char * key, char * value); -void initCUDA(int blockMapping, int rank, int numNodes, int tasksPerNode, int useGPUID); +void initGPU(int blockMapping, int rank, int numNodes, int tasksPerNode, int useGPUID); char *CurrentTimeString(void); int Regex(char *, char *); void ShowFileSystemSize(char * filename, const struct ior_aiori * backend, void * backend_options);