CUDA: remove the need of adiosCUDAReduceImpl

ornladios · Feb 3, 2022 · eae3dcc · eae3dcc
1 parent a382cc0
commit eae3dcc
Show file tree

Hide file tree

Showing 4 changed files with 38 additions and 483 deletions.
diff --git a/source/adios2/CMakeLists.txt b/source/adios2/CMakeLists.txt
@@ -3,6 +3,11 @@
 # accompanying file Copyright.txt for details.
 #------------------------------------------------------------------------------#
 
+include_directories(
+    $<BUILD_INTERFACE:${ADIOS2_SOURCE_DIR}/source>
+    $<BUILD_INTERFACE:${ADIOS2_BINARY_DIR}/source>
+    )
+
 add_library(adios2_core
   common/ADIOSTypes.cpp
 
@@ -37,7 +42,6 @@ add_library(adios2_core
   helper/adiosXML.cpp
   helper/adiosXMLUtil.cpp
   helper/adiosYAML.cpp
-  helper/adiosCUDA.cu
   helper/adiosLog.cpp
 
 #engine derived classes
@@ -102,13 +106,28 @@ add_library(adios2_core
   toolkit/burstbuffer/FileDrainer.cpp
   toolkit/burstbuffer/FileDrainerSingleThread.cpp
 )
+
+
 set_property(TARGET adios2_core PROPERTY EXPORT_NAME core)
 set_property(TARGET adios2_core PROPERTY OUTPUT_NAME adios2${ADIOS2_LIBRARY_SUFFIX}_core)
 
+set(maybe_adios2_core_cuda)
 if(ADIOS2_HAVE_CUDA)
   enable_language(CUDA)
-  target_link_libraries(adios2_core PRIVATE CUDA::cudart)
-  set_target_properties(adios2_core PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+
+  add_library(adios2_core_cuda helper/adiosCUDA.cu)
+  set_target_properties(adios2_core_cuda PROPERTIES
+    CUDA_SEPARABLE_COMPILATION ON
+    CUDA_STANDARD 14
+    CUDA_STANDARD_REQUIRED ON
+    CUDA_VISIBILITY_PRESET hidden
+    VISIBILITY_INLINES_HIDDEN ON
+    EXPORT_NAME core_cuda
+    OUTPUT_NAME adios2${ADIOS2_LIBRARY_SUFFIX}_core_cuda
+    )
+
+  target_link_libraries(adios2_core PRIVATE adios2_core_cuda CUDA::cudart CUDA::cuda_driver)
+  set(maybe_adios2_core_cuda adios2_core_cuda)
 endif()
 
 target_include_directories(adios2_core
@@ -358,6 +377,11 @@ install(DIRECTORY core/
   FILES_MATCHING PATTERN "*.h"
 )
 
+install(DIRECTORY core/
+  DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/adios2/core/ COMPONENT adios2_core-development
+  FILES_MATCHING PATTERN "*.cu"
+)
+
 install(DIRECTORY engine/
   DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/adios2/engine COMPONENT adios2_core-development
   FILES_MATCHING PATTERN "*/*.h"
@@ -377,7 +401,7 @@ install(DIRECTORY toolkit/
 )
 
 # Library installation
-install(TARGETS adios2_core ${maybe_adios2_core_mpi} EXPORT adios2Exports
+install(TARGETS adios2_core ${maybe_adios2_core_mpi} ${maybe_adios2_core_cuda} EXPORT adios2Exports
   RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} COMPONENT adios2_core-runtime
   LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} COMPONENT adios2_core-libraries NAMELINK_COMPONENT adios2_core-development
   ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR} COMPONENT adios2_core-development

diff --git a/source/adios2/helper/adiosCUDA.cu b/source/adios2/helper/adiosCUDA.cu
@@ -12,18 +12,22 @@
 #define ADIOS2_HELPER_ADIOSCUDA_CU_
 
 #include "adios2/common/ADIOSMacros.h"
+#include <thrust/device_ptr.h>
+#include <thrust/extrema.h>
 
 #include "adiosCUDA.h"
-#include "adiosCUDAReduceImpl.h"
 
 namespace
 {
-
 template <class T>
 void CUDAMinMaxImpl(const T *values, const size_t size, T &min, T &max)
 {
-    min = reduce<T, MinOp>(size, 1024, 64, 1, values);
-    max = reduce<T, MaxOp>(size, 1024, 64, 1, values);
+    thrust::device_ptr<const T> dev_ptr(values);
+    auto res = thrust::minmax_element(dev_ptr, dev_ptr + size);
+    cudaMemcpy(&min, thrust::raw_pointer_cast(res.first), sizeof(T),
+               cudaMemcpyDeviceToHost);
+    cudaMemcpy(&max, thrust::raw_pointer_cast(res.second), sizeof(T),
+               cudaMemcpyDeviceToHost);
 }
 // types non supported on the device
 void CUDAMinMaxImpl(const long double * /*values*/, const size_t /*size*/,

diff --git a/source/adios2/helper/adiosCUDA.h b/source/adios2/helper/adiosCUDA.h
@@ -21,7 +21,8 @@ namespace helper
  * GPU buffer
  */
 template <class T>
-void CUDAMinMax(const T *values, const size_t size, T &min, T &max);
+__attribute__((visibility("default"))) void
+CUDAMinMax(const T *values, const size_t size, T &min, T &max);
 
 } // helper
 } // adios2