Merge pull request #53 from SeisSol/davschneller/memory-fixes

Localize thread Id

Author: davschneller

interfaces/common/Common.h

@@ -50,49 +50,5 @@ constexpr auto mapPercentage(int minval, int maxval, double value) {
   return std::max(std::min(static_cast<int>(std::floor(transformed)), maxval), minval);
 }
 
-class InfoPrinter {
-public:
-  struct InfoPrinterLine {
-    std::shared_ptr<std::ostringstream> stream;
-    InfoPrinter& printer;
-
-    InfoPrinterLine(InfoPrinter& printer) : stream(std::make_shared<std::ostringstream>()), printer(printer) {}
-
-    template<typename T>
-    InfoPrinterLine& operator <<(const T& data) {
-      *stream << data;
-      return *this;
-    }
-
-    ~InfoPrinterLine() {
-      if (printer.rank < 0) {
-        printer.stringCache.emplace_back(stream->str());
-      }
-      else {
-        logInfo() << stream->str().c_str();
-      }
-      stream = nullptr;
-    }
-  };
-
-  InfoPrinterLine printInfo() {
-    return InfoPrinterLine(*this);
-  }
-
-  void setRank(int rank) {
-    this->rank = rank;
-
-    for (const auto& string : stringCache) {
-      logInfo() << string;
-    }
-
-    stringCache.resize(0);
-  }
-
-private:
-  int rank{-1};
-  std::vector<std::string> stringCache;
-};
-
 
 #endif // SEISSOLDEVICE_INTERFACES_COMMON_COMMON_H_

interfaces/cuda/Control.cu

@@ -7,8 +7,10 @@
 #include <cuda.h>
 #include <iostream>
 #include <iomanip>
+#include <mutex>
 #include <sstream>
 #include <string>
+#include <thread>
 
 #ifdef PROFILING_ENABLED
 #include <nvToolsExt.h>
@@ -17,6 +19,16 @@
 #include "CudaWrappedAPI.h"
 #include "Internals.h"
 
+namespace {
+// `static` is a bit out of place here; but we treat the whole class as an effective singleton anyways
+
+#ifdef DEVICE_CONTEXT_GLOBAL
+int currentDeviceId = 0;
+#else
+thread_local int currentDeviceId = 0;
+#endif
+}
+
 using namespace device;
 
 ConcreteAPI::ConcreteAPI() {
@@ -26,18 +38,16 @@ ConcreteAPI::ConcreteAPI() {
 }
 
 void ConcreteAPI::setDevice(int deviceId) {
+  
   currentDeviceId = deviceId;
-  cudaSetDevice(currentDeviceId);
+
+  cudaSetDevice(deviceId);
   CHECK_ERR;
 
   // Note: the following sets the initial CUDA context
   cudaFree(nullptr);
   CHECK_ERR;
 
-  int result;
-  cudaDeviceGetAttribute(&result, cudaDevAttrDirectManagedMemAccessFromHost, currentDeviceId);
-  usmDefault = result != 0;
-
   status[StatusID::DeviceSelected] = true;
 }
 
@@ -52,18 +62,20 @@ void ConcreteAPI::initialize() {
   if (!status[StatusID::InterfaceInitialized]) {
     status[StatusID::InterfaceInitialized] = true;
     cudaStreamCreateWithFlags(&defaultStream, cudaStreamNonBlocking); CHECK_ERR;
-    cudaEventCreate(&defaultStreamEvent); CHECK_ERR;
 
     int result{0};
-    cudaDeviceGetAttribute(&result, cudaDevAttrConcurrentManagedAccess, currentDeviceId);
+    cudaDeviceGetAttribute(&result, cudaDevAttrConcurrentManagedAccess, getDeviceId());
     CHECK_ERR;
     allowedConcurrentManagedAccess = result != 0;
 
+    cudaDeviceGetAttribute(&result, cudaDevAttrDirectManagedMemAccessFromHost, getDeviceId());
+    usmDefault = result != 0;
+
     cudaDeviceGetStreamPriorityRange(&priorityMin, &priorityMax);
     CHECK_ERR;
 
     int canCompressProto = 0;
-    cuDeviceGetAttribute(&canCompressProto, CU_DEVICE_ATTRIBUTE_GENERIC_COMPRESSION_SUPPORTED, currentDeviceId);
+    cuDeviceGetAttribute(&canCompressProto, CU_DEVICE_ATTRIBUTE_GENERIC_COMPRESSION_SUPPORTED, getDeviceId());
     canCompress = canCompressProto != 0;
   }
   else {
@@ -74,9 +86,8 @@ void ConcreteAPI::initialize() {
 void ConcreteAPI::finalize() {
   if (status[StatusID::InterfaceInitialized]) {
     cudaStreamDestroy(defaultStream); CHECK_ERR;
-    cudaEventDestroy(defaultStreamEvent); CHECK_ERR;
     if (!genericStreams.empty()) {
-      printer.printInfo() << "DEVICE::WARNING:" << genericStreams.size()
+      logInfo() << "DEVICE::WARNING:" << genericStreams.size()
                                << "device generic stream(s) were not deleted.";
       for (auto stream : genericStreams) {
         cudaStreamDestroy(stream); CHECK_ERR;
@@ -186,6 +197,6 @@ void ConcreteAPI::popLastProfilingMark() {
 }
 
 void ConcreteAPI::setupPrinting(int rank) {
-  printer.setRank(rank);
+
 }
 

interfaces/cuda/Copy.cu

@@ -80,7 +80,7 @@ void ConcreteAPI::prefetchUnifiedMemTo(Destination type, const void *devPtr, siz
 #endif
   }
   else if (allowedConcurrentManagedAccess) {
-    location.id = currentDeviceId;
+    location.id = getDeviceId();
 #if CUDART_VERSION >= 13000
     location.type = cudaMemLocationTypeDevice;
 #endif

interfaces/cuda/CudaWrappedAPI.h

@@ -9,6 +9,7 @@
 #include <string>
 #include <unordered_map>
 #include <vector>
+#include <thread>
 #include <unordered_set>
 #include <cassert>
 #include <cstdint>
@@ -105,14 +106,15 @@ class ConcreteAPI : public AbstractAPI {
   void createCircularStreamAndEvents();
 
   device::StatusT status{false};
-  int currentDeviceId{-1};
+
+  std::vector<cudaDeviceProp> properties;
+
   bool allowedConcurrentManagedAccess{false};
 
   bool usmDefault{false};
   bool canCompress{false};
 
   cudaStream_t defaultStream{nullptr};
-  cudaEvent_t defaultStreamEvent{};
 
   std::unordered_set<cudaStream_t> genericStreams{};
 
@@ -128,7 +130,6 @@ class ConcreteAPI : public AbstractAPI {
   std::unordered_map<void *, size_t> memToSizeMap{{nullptr, 0}};
 
   int priorityMin, priorityMax;
-  InfoPrinter printer;
 
   std::unordered_map<void *, void *> allocationProperties;
 };

interfaces/cuda/Memory.cu

@@ -76,7 +76,7 @@ void *ConcreteAPI::allocGlobMem(size_t size, bool compress) {
     std::memset(&prop, 0, sizeof(CUmemAllocationProp));
     prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
     prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
-    prop.location.id = currentDeviceId;
+    prop.location.id = getDeviceId();
     prop.allocFlags.compressionType = CU_MEM_ALLOCATION_COMP_GENERIC;
 
     devPtr = driverAllocate(size, prop);
@@ -105,7 +105,7 @@ void *ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint)
 #endif
   }
   else if (allowedConcurrentManagedAccess) {
-    location.id = currentDeviceId;
+    location.id = getDeviceId();
 #if CUDART_VERSION >= 13000
     location.type = cudaMemLocationTypeDevice;
 #endif
@@ -198,7 +198,7 @@ std::string ConcreteAPI::getMemLeaksReport() {
 size_t ConcreteAPI::getMaxAvailableMem() {
   isFlagSet<DeviceSelected>(status);
   cudaDeviceProp property;
-  cudaGetDeviceProperties(&property, currentDeviceId);
+  cudaGetDeviceProperties(&property, getDeviceId());
   CHECK_ERR;
   return property.totalGlobalMem;
 }

interfaces/hip/Control.cpp

@@ -6,8 +6,10 @@
 #include "utils/env.h"
 #include <iostream>
 #include <iomanip>
+#include <mutex>
 #include <sstream>
 #include <string>
+#include <thread>
 #include "hip/hip_runtime.h"
 
 #ifdef PROFILING_ENABLED
@@ -19,40 +21,31 @@
 
 using namespace device;
 
+namespace {
+#ifdef DEVICE_CONTEXT_GLOBAL
+int currentDeviceId = 0;
+#else
+thread_local int currentDeviceId = 0;
+#endif
+}
+
 ConcreteAPI::ConcreteAPI() {
   hipInit(0);
   CHECK_ERR;
   status[StatusID::DriverApiInitialized] = true;
 }
 
 void ConcreteAPI::setDevice(int deviceId) {
+  
   currentDeviceId = deviceId;
-  hipSetDevice(currentDeviceId);
+
+  hipSetDevice(deviceId);
   CHECK_ERR;
 
   // Note: the following sets the initial HIP context
   hipFree(nullptr);
   CHECK_ERR;
 
-  hipDeviceProp_t properties{};
-  hipGetDeviceProperties(&properties, currentDeviceId);
-  CHECK_ERR;
-
-  // NOTE: hipDeviceGetAttribute internally calls hipGetDeviceProperties; hence it doesn't make sense to use it here
-
-  if constexpr (HIP_VERSION >= 60200000) {
-    // cf. https://rocm.docs.amd.com/en/docs-6.2.0/about/release-notes.html
-    // (before 6.2.0, the flag hipDeviceAttributePageableMemoryAccessUsesHostPageTables had effectively the same effect)
-    // (cf. https://github.com/ROCm/clr/commit/7d5b4a8f7a7d34f008d65277f8aae4c98a6da375#diff-596cd550f7fdef76b39f1b7b179b20128313dd9cc9ec662b2eae562efa2b7f33L405 )
-    usmDefault = properties.integrated != 0;
-  }
-  else {
-    usmDefault = properties.directManagedMemAccessFromHost != 0 && properties.pageableMemoryAccessUsesHostPageTables != 0;
-  }
-
-  hipDeviceGetStreamPriorityRange(&priorityMin, &priorityMax);
-  CHECK_ERR;
-
   status[StatusID::DeviceSelected] = true;
 }
 
@@ -68,7 +61,25 @@ void ConcreteAPI::initialize() {
   if (!status[StatusID::InterfaceInitialized]) {
     status[StatusID::InterfaceInitialized] = true;
     hipStreamCreateWithFlags(&defaultStream, hipStreamNonBlocking); CHECK_ERR;
-    hipEventCreate(&defaultStreamEvent); CHECK_ERR;
+
+    hipDeviceProp_t properties{};
+    hipGetDeviceProperties(&properties, getDeviceId());
+    CHECK_ERR;
+
+    // NOTE: hipDeviceGetAttribute internally calls hipGetDeviceProperties; hence it doesn't make sense to use it here
+
+    if constexpr (HIP_VERSION >= 60200000) {
+      // cf. https://rocm.docs.amd.com/en/docs-6.2.0/about/release-notes.html
+      // (before 6.2.0, the flag hipDeviceAttributePageableMemoryAccessUsesHostPageTables had effectively the same effect)
+      // (cf. https://github.com/ROCm/clr/commit/7d5b4a8f7a7d34f008d65277f8aae4c98a6da375#diff-596cd550f7fdef76b39f1b7b179b20128313dd9cc9ec662b2eae562efa2b7f33L405 )
+      usmDefault = properties.integrated != 0;
+    }
+    else {
+      usmDefault = properties.directManagedMemAccessFromHost != 0 && properties.pageableMemoryAccessUsesHostPageTables != 0;
+    }
+
+    hipDeviceGetStreamPriorityRange(&priorityMin, &priorityMax);
+    CHECK_ERR;
   }
   else {
     logWarning() << "Device Interface has already been initialized";
@@ -78,9 +89,8 @@ void ConcreteAPI::initialize() {
 void ConcreteAPI::finalize() {
   if (status[StatusID::InterfaceInitialized]) {
     hipStreamDestroy(defaultStream); CHECK_ERR;
-    hipEventDestroy(defaultStreamEvent); CHECK_ERR;
     if (!genericStreams.empty()) {
-      printer.printInfo() << "DEVICE::WARNING:" << genericStreams.size()
+      logInfo() << "DEVICE::WARNING:" << genericStreams.size()
                                << "device generic stream(s) were not deleted.";
       for (auto stream : genericStreams) {
         hipStreamDestroy(stream); CHECK_ERR;
@@ -189,6 +199,6 @@ void ConcreteAPI::popLastProfilingMark() {
 }
 
 void ConcreteAPI::setupPrinting(int rank) {
-  printer.setRank(rank);
+
 }
 

interfaces/hip/Copy.cpp

@@ -81,7 +81,7 @@ void ConcreteAPI::prefetchUnifiedMemTo(Destination type, const void* devPtr, siz
   hipStream_t stream = (streamPtr == nullptr) ? 0 : (static_cast<hipStream_t>(streamPtr));
   hipMemPrefetchAsync(devPtr,
                        count,
-                       type == Destination::CurrentDevice ? currentDeviceId : hipCpuDeviceId,
+                       type == Destination::CurrentDevice ? getDeviceId() : hipCpuDeviceId,
                        stream);
 }
 

interfaces/hip/HipWrappedAPI.h

@@ -14,6 +14,7 @@
 #include <string>
 #include <unordered_map>
 #include <vector>
+#include <thread>
 #include <unordered_set>
 #include <cassert>
 
@@ -103,12 +104,10 @@ class ConcreteAPI : public AbstractAPI {
 
 private:
   device::StatusT status{false};
-  int currentDeviceId{-1};
 
   bool usmDefault{false};
 
   hipStream_t defaultStream{nullptr};
-  hipEvent_t defaultStreamEvent{};
 
   std::unordered_set<hipStream_t> genericStreams{};
 
@@ -124,7 +123,6 @@ class ConcreteAPI : public AbstractAPI {
   std::unordered_map<void *, size_t> memToSizeMap{{nullptr, 0}};
 
   int priorityMin, priorityMax;
-  InfoPrinter printer;
 };
 } // namespace device
 

interfaces/hip/Memory.cpp

@@ -33,7 +33,7 @@ void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint)
     CHECK_ERR;
   }
   else {
-    hipMemAdvise(devPtr, size, hipMemAdviseSetPreferredLocation, currentDeviceId);
+    hipMemAdvise(devPtr, size, hipMemAdviseSetPreferredLocation, getDeviceId());
     CHECK_ERR;
   }
   statistics.allocatedMemBytes += size;
@@ -105,7 +105,7 @@ std::string ConcreteAPI::getMemLeaksReport() {
 size_t ConcreteAPI::getMaxAvailableMem() {
   isFlagSet<DeviceSelected>(status);
   hipDeviceProp_t property;
-  hipGetDeviceProperties(&property, currentDeviceId); CHECK_ERR;
+  hipGetDeviceProperties(&property, getDeviceId()); CHECK_ERR;
   return property.totalGlobalMem;
 }