diff --git a/.gitlab/custom-jobs-and-variables.yml b/.gitlab/custom-jobs-and-variables.yml index 3899cbc6a..abdeb5c95 100644 --- a/.gitlab/custom-jobs-and-variables.yml +++ b/.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 diff --git a/.gitlab/jobs/tioga.yml b/.gitlab/jobs/tioga.yml index 38dc44644..df80d0718 100644 --- a/.gitlab/jobs/tioga.yml +++ b/.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 ^hip@6.2.0" 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\" %rocmcc@6.1.1 ^hip@6.1.1" + ASAN_OPTIONS: "detect_leaks=1" + extends: .job_on_tioga diff --git a/src/umpire/util/memory_sanitizers.hpp b/src/umpire/util/memory_sanitizers.hpp index 1da76c835..ab084e5fb 100644 --- a/src/umpire/util/memory_sanitizers.hpp +++ b/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__) diff --git a/tests/tools/sanitizers/CMakeLists.txt b/tests/tools/sanitizers/CMakeLists.txt index 24a6c1c2f..22b8c0354 100644 --- a/tests/tools/sanitizers/CMakeLists.txt +++ b/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) diff --git a/tests/tools/sanitizers/sanitizer_test_runner.py b/tests/tools/sanitizers/sanitizer_test_runner.py index b15fdb412..7f85e6f03 100644 --- a/tests/tools/sanitizers/sanitizer_test_runner.py +++ b/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) diff --git a/tests/tools/sanitizers/sanitizer_tests.cpp b/tests/tools/sanitizers/sanitizer_tests.cpp index 8411acc95..cb458256e 100644 --- a/tests/tools/sanitizers/sanitizer_tests.cpp +++ b/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(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("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("test_allocator", rm.getAllocator(resource_type)); UMPIRE_USE_VAR(pool); - } else if (strategy.find("DynamicPoolList") != std::string::npos) { - auto pool = rm.makeAllocator("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("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; }