Adding asan sanitizer support for hip

.gitlab/custom-jobs-and-variables.yml

@@ -45,7 +45,7 @@ variables:
 
 # Tioga
 # Arguments for top level allocation
-  TIOGA_SHARED_ALLOC: "--queue=pci --exclusive --time-limit=15m --nodes=1"
+  TIOGA_SHARED_ALLOC: "--queue=pci --exclusive --time-limit=60m --nodes=1"
 # Arguments for job level allocation
   TIOGA_JOB_ALLOC: "--nodes=1 --begin-time=+5s"
 # Project specific variants for tioga

.gitlab/jobs/tioga.yml

@@ -38,3 +38,8 @@ rocmcc_6_2_0_hip_openmp_device_alloc:
     SPEC: "~shared +fortran +openmp +rocm +device_alloc tests=basic amdgpu_target=gfx90a %rocmcc@=6.2.0 ^[email protected]"
   extends: .job_on_tioga
 
+rocmcc_6_1_1_hip_memleak:
+  variables:
+    SPEC: "~shared +asan +sanitizer_tests +tools +rocm tests=basic amdgpu_target=gfx90a:xnack+ cxxflags==\"-fsanitize=address\" %[email protected] ^[email protected]"
+    ASAN_OPTIONS: "detect_leaks=1"
+  extends: .job_on_tioga

src/umpire/util/memory_sanitizers.hpp

@@ -34,14 +34,14 @@
 
 #if defined(__UMPIRE_USE_MEMORY_SANITIZER__)
 
-#define UMPIRE_POISON_MEMORY_REGION(allocator, ptr, size)   \
-  if (allocator->getPlatform() == umpire::Platform::host) { \
-    ASAN_POISON_MEMORY_REGION((ptr), (size));               \
+#define UMPIRE_POISON_MEMORY_REGION(allocator, ptr, size)                                                        \
+  if (allocator->getPlatform() == umpire::Platform::host || allocator->getPlatform() == umpire::Platform::hip) { \
+    ASAN_POISON_MEMORY_REGION((ptr), (size));                                                                    \
   }
 
-#define UMPIRE_UNPOISON_MEMORY_REGION(allocator, ptr, size) \
-  if (allocator->getPlatform() == umpire::Platform::host) { \
-    ASAN_UNPOISON_MEMORY_REGION((ptr), (size));             \
+#define UMPIRE_UNPOISON_MEMORY_REGION(allocator, ptr, size)                                                      \
+  if (allocator->getPlatform() == umpire::Platform::host || allocator->getPlatform() == umpire::Platform::hip) { \
+    ASAN_UNPOISON_MEMORY_REGION((ptr), (size));                                                                  \
   }
 
 #else // !defined(__UMPIRE_USE_MEMORY_SANITIZER__)

tests/tools/sanitizers/CMakeLists.txt

@@ -4,11 +4,18 @@
 #
 # SPDX-License-Identifier: (MIT)
 ##############################################################################
+set (sanitizer_depends umpire umpire_util)
+
+if (UMPIRE_ENABLE_HIP)
+  set (sanitizer_depends
+      ${sanitizer_depends}
+      blt::hip)
+endif ()
 
 blt_add_executable(
   NAME sanitizer_tests
   SOURCES sanitizer_tests.cpp
-  DEPENDS_ON umpire)
+  DEPENDS_ON ${sanitizer_depends})
 
 include(FindPythonInterp)
 

tests/tools/sanitizers/sanitizer_test_runner.py

@@ -31,13 +31,14 @@ def check_output(file_object, expected):
         print("{BLUE}[      OK]{END} Found \"{expected}\"".format(expected=expected, **formatters))
 
 
-def run_sanitizer_test(strategy, kind):
+def run_sanitizer_test(strategy, kind, resource):
     import subprocess
     import os
 
     cmd_args = ['./sanitizer_tests']
     cmd_args.append(strategy)
     cmd_args.append(kind)
+    cmd_args.append(resource)
 
     test_program = subprocess.Popen(cmd_args,
             stdout=subprocess.PIPE,
@@ -57,9 +58,9 @@ def run_sanitizer_test(strategy, kind):
     import sys
 
     print("{BLUE}[--------]{END}".format(**formatters))
-    run_sanitizer_test('DynamicPoolList', 'read')
-    run_sanitizer_test('DynamicPoolList', 'write')
-    run_sanitizer_test('QuickPool', 'read')
-    run_sanitizer_test('QuickPool', 'write')
+    run_sanitizer_test('DynamicPoolList', 'read', 'HOST')
+    run_sanitizer_test('DynamicPoolList', 'write', 'HOST')
+    run_sanitizer_test('QuickPool', 'read', 'HOST')
+    run_sanitizer_test('QuickPool', 'write', 'HOST')
     print("{BLUE}[--------]{END}".format(**formatters))
     sys.exit(errors)

tests/tools/sanitizers/sanitizer_tests.cpp

@@ -9,24 +9,27 @@
 #include "umpire/ResourceManager.hpp"
 #include "umpire/strategy/DynamicPoolList.hpp"
 #include "umpire/strategy/QuickPool.hpp"
+#include "umpire/util/Macros.hpp"
+#include "umpire/util/error.hpp"
 
-void test_read_after_free()
+#if defined(UMPIRE_ENABLE_HIP)
+__global__ void test_read_for_hip(double* data_ptr, std::size_t INDEX)
 {
-  auto& rm = umpire::ResourceManager::getInstance();
-  auto allocator = rm.getAllocator("test_allocator");
-
-  const std::size_t SIZE = 1356;
-  const std::size_t INDEX = SIZE / 2;
-  double* data = static_cast<double*>(allocator.allocate(SIZE * sizeof(double)));
-
-  data[INDEX] = 100;
-  std::cout << "data[INDEX] = " << data[INDEX] << std::endl;
-
-  allocator.deallocate(data);
-  std::cout << "data[256] = " << data[256] << std::endl;
+  if (threadIdx.x == 0) {
+    double dummy = data_ptr[INDEX];
+    dummy = dummy + 42;
+  }
 }
+__global__ void test_write_for_hip(double* data_ptr, std::size_t INDEX)
+{
+  if (threadIdx.x == 0) {
+    data_ptr[INDEX] = 256;
+    data_ptr[INDEX] = INDEX + 42;
+  }
+}
+#endif
 
-void test_write_after_free()
+void sanitizer_test(const std::string test_type)
 {
   auto& rm = umpire::ResourceManager::getInstance();
   auto allocator = rm.getAllocator("test_allocator");
@@ -37,36 +40,72 @@ void test_write_after_free()
 
   data[INDEX] = 100;
   std::cout << "data[INDEX] = " << data[INDEX] << std::endl;
-
   allocator.deallocate(data);
-  data[INDEX] = -1;
-  std::cout << "data[INDEX] = " << data[INDEX] << std::endl;
+
+  if (test_type.find("read") != std::string::npos) {
+#if defined(UMPIRE_ENABLE_HIP)
+    hipLaunchKernelGGL(test_read_for_hip, dim3(1), dim3(16), 0, 0, data, INDEX);
+    hipDeviceSynchronize();
+#endif
+    std::cout << "data[256] = " << data[256] << std::endl;
+  } else {
+    if (test_type.find("write") == std::string::npos) {
+      std::cout << "Test type did not match either option - using write" << std::endl;
+    }
+#if defined(UMPIRE_ENABLE_HIP)
+    hipLaunchKernelGGL(test_write_for_hip, dim3(1), dim3(16), 0, 0, data, INDEX);
+    hipDeviceSynchronize();
+#endif
+    data[INDEX] = -1;
+    std::cout << "data[INDEX] = " << data[INDEX] << std::endl;
+  }
 }
 
 int main(int argc, char* argv[])
 {
-  if (argc < 3) {
-    std::cout << argv[0] << " requires 2 arguments, test type and allocation strategy" << std::endl;
+  if (argc < 4) {
+    std::cout << "Usage: requires 3 arguments." << std::endl;
+    std::cout << "First, an allocation strategy (QuickPool or DynamicPoolList)." << std::endl;
+    std::cout << "Second, a test type (read or write)." << std::endl;
+    std::cout << "Third, a resource type (DEVICE, HOST, or UM)." << std::endl;
+    return 0;
   }
 
-  const std::string strategy{argv[1]};
-  const std::string test_type{argv[2]};
+  std::string strategy{argv[1]};
+  std::string test_type{argv[2]};
+  std::string resource_type{argv[3]};
 
   auto& rm = umpire::ResourceManager::getInstance();
 
-  if (strategy.find("QuickPool") != std::string::npos) {
-    auto pool = rm.makeAllocator<umpire::strategy::QuickPool>("test_allocator", rm.getAllocator("HOST"));
+  if ((resource_type.find("DEVICE") != std::string::npos) || (resource_type.find("UM") != std::string::npos)) {
+#if !defined(UMPIRE_ENABLE_HIP)
+    UMPIRE_ERROR(umpire::runtime_error,
+                 fmt::format("The resource, \"{}\", can't be used if HIP is not enabled.", resource_type));
+#endif
+  } else {
+    if (resource_type.find("HOST") == std::string::npos) {
+      std::cout << "Resource type did not match any available options - using HOST." << std::endl;
+      resource_type = "HOST";
+    }
+  }
+
+  if (strategy.find("DynamicPoolList") != std::string::npos) {
+    auto pool = rm.makeAllocator<umpire::strategy::DynamicPoolList>("test_allocator", rm.getAllocator(resource_type));
     UMPIRE_USE_VAR(pool);
-  } else if (strategy.find("DynamicPoolList") != std::string::npos) {
-    auto pool = rm.makeAllocator<umpire::strategy::DynamicPoolList>("test_allocator", rm.getAllocator("HOST"));
+  } else {
+    if (strategy.find("QuickPool") == std::string::npos) {
+      std::cout << "Allocation strategy did not match either option - using QuickPool." << std::endl;
+      strategy = "QuickPool";
+    }
+    auto pool = rm.makeAllocator<umpire::strategy::QuickPool>("test_allocator", rm.getAllocator(resource_type));
     UMPIRE_USE_VAR(pool);
   }
 
-  if (test_type.find("read") != std::string::npos) {
-    test_read_after_free();
-  } else if (test_type.find("write") != std::string::npos) {
-    test_write_after_free();
-  }
+  std::cout << " Conducting sanitizer test with " << strategy << " strategy, " << test_type << " test type, and the "
+            << resource_type << " resource." << std::endl;
+
+  // Conduct the test
+  sanitizer_test(test_type);
 
   return 0;
 }