From dd1d42146ef63532a1f80e183c4efa3d6439c909 Mon Sep 17 00:00:00 2001 From: Sourav Chakraborty Date: Wed, 4 Dec 2019 14:21:05 -0800 Subject: [PATCH] GTEST/UCM/ROCM: test hip runtime malloc/free events --- config/m4/rocm.m4 | 10 +- test/gtest/Makefile.am | 14 +++ test/gtest/common/mem_buffer.cc | 24 ++++ test/gtest/ucm/rocm_hooks.cc | 191 ++++++++++++++++++++++++++++++++ 4 files changed, 237 insertions(+), 2 deletions(-) create mode 100644 test/gtest/ucm/rocm_hooks.cc diff --git a/config/m4/rocm.m4 b/config/m4/rocm.m4 index 567e50ed0950..6e47690a3f13 100644 --- a/config/m4/rocm.m4 +++ b/config/m4/rocm.m4 @@ -49,7 +49,10 @@ AS_IF([test "x$with_rocm" != "xno"], [AC_MSG_NOTICE([ROCm path given as $with_rocm ...]) ROCM_CPPFLAGS="-I$with_rocm/include/hsa -I$with_rocm/include" ROCM_LDFLAGS="-L$with_rocm/hsa/lib -L$with_rocm/lib" - ROCM_LIBS="-lhsa-runtime64"], + ROCM_LIBS="-lhsa-runtime64" + HIP_CPPFLAGS="-D__HIP_PLATFORM_HCC__ -I$with_rocm/include/hip -I$with_rocm/include" + HIP_LDFLAGS="-L$with_rocm/hip/lib -L$with_rocm/lib" + HIP_LIBS="-lhip_hcc"], [AC_MSG_NOTICE([ROCm flags given ...]) ROCM_PARSE_FLAGS([with_rocm], [ROCM_LIBS], [ROCM_LDFLAGS], [ROCM_CPPFLAGS])]) @@ -80,7 +83,10 @@ AS_IF([test "x$with_rocm" != "xno"], AS_IF([test "x$rocm_happy" = "xyes"], [AC_SUBST([ROCM_CPPFLAGS]) AC_SUBST([ROCM_LDFLAGS]) - AC_SUBST([ROCM_LIBS])], + AC_SUBST([ROCM_LIBS]) + AC_SUBST([HIP_CPPFLAGS]) + AC_SUBST([HIP_LDFLAGS]) + AC_SUBST([HIP_LIBS])], [AC_MSG_WARN([ROCm not found])]) ], [AC_MSG_WARN([ROCm was explicitly disabled])] diff --git a/test/gtest/Makefile.am b/test/gtest/Makefile.am index 39d877a00a19..5c2995419d15 100644 --- a/test/gtest/Makefile.am +++ b/test/gtest/Makefile.am @@ -3,6 +3,7 @@ # Copyright (C) UT-Battelle, LLC. 2015. ALL RIGHTS RESERVED. # Copyright (C) The University of Tennessee and the University of Tennessee Research Foundation. 2016. ALL RIGHTS RESERVED. # Copyright (C) Los Alamos National Security, LLC. 2018 ALL RIGHTS RESERVED. +# Copyright (C) Advanced Micro Devices, Inc. 2019. ALL RIGHTS RESERVED. # # See file LICENSE for terms. # @@ -214,6 +215,19 @@ gtest_LDADD += \ $(top_builddir)/src/uct/cuda/libuct_cuda.la endif +if HAVE_ROCM +gtest_SOURCES += \ + ucm/rocm_hooks.cc +gtest_CPPFLAGS += \ + $(HIP_CPPFLAGS) +gtest_CXXFLAGS += \ + -std=gnu++11 +gtest_LDADD += \ + $(HIP_LDFLAGS) \ + $(HIP_LIBS) \ + $(top_builddir)/src/uct/rocm/libuct_rocm.la +endif + noinst_HEADERS = \ common/gtest.h \ common/mem_buffer.h \ diff --git a/test/gtest/common/mem_buffer.cc b/test/gtest/common/mem_buffer.cc index 5b0ce70696d0..73dff1a1892a 100644 --- a/test/gtest/common/mem_buffer.cc +++ b/test/gtest/common/mem_buffer.cc @@ -1,5 +1,6 @@ /** * Copyright (C) Mellanox Technologies Ltd. 2001-2019. ALL RIGHTS RESERVED. + * Copyright (C) Advanced Micro Devices, Inc. 2019. ALL RIGHTS RESERVED. * * See file LICENSE for terms. */ @@ -27,6 +28,19 @@ #endif +#if HAVE_ROCM +# include + +#define ROCM_CALL(_code) \ + do { \ + hipError_t cerr = _code; \ + if (cerr != hipSuccess) { \ + UCS_TEST_ABORT(# _code << " failed"); \ + } \ + } while (0) + +#endif + std::vector mem_buffer::supported_mem_types() { @@ -35,6 +49,10 @@ std::vector mem_buffer::supported_mem_types() #if HAVE_CUDA vec.push_back(UCS_MEMORY_TYPE_CUDA); vec.push_back(UCS_MEMORY_TYPE_CUDA_MANAGED); +#endif +#if HAVE_ROCM + vec.push_back(UCS_MEMORY_TYPE_ROCM); + vec.push_back(UCS_MEMORY_TYPE_ROCM_MANAGED); #endif return vec; } @@ -75,6 +93,12 @@ void mem_buffer::release(void *ptr, ucs_memory_type_t mem_type) case UCS_MEMORY_TYPE_CUDA_MANAGED: CUDA_CALL(cudaFree(ptr)); break; +#endif +#if HAVE_ROCM + case UCS_MEMORY_TYPE_ROCM: + case UCS_MEMORY_TYPE_ROCM_MANAGED: + ROCM_CALL(hipFree(ptr)); + break; #endif default: break; diff --git a/test/gtest/ucm/rocm_hooks.cc b/test/gtest/ucm/rocm_hooks.cc new file mode 100644 index 000000000000..22f7e86d0e95 --- /dev/null +++ b/test/gtest/ucm/rocm_hooks.cc @@ -0,0 +1,191 @@ +/** +* Copyright (C) Mellanox Technologies Ltd. 2001-2017. ALL RIGHTS RESERVED. +* Copyright (C) Advanced Micro Devices, Inc. 2019. ALL RIGHTS RESERVED. +* See file LICENSE for terms. +*/ +#include +#include +#include + +static ucm_event_t alloc_event, free_event; + +static void rocm_mem_alloc_callback(ucm_event_type_t event_type, + ucm_event_t *event, void *arg) +{ + alloc_event.mem_type.address = event->mem_type.address; + alloc_event.mem_type.size = event->mem_type.size; + alloc_event.mem_type.mem_type = event->mem_type.mem_type; +} + +static void rocm_mem_free_callback(ucm_event_type_t event_type, + ucm_event_t *event, void *arg) +{ + free_event.mem_type.address = event->mem_type.address; + free_event.mem_type.size = event->mem_type.size; + free_event.mem_type.mem_type = event->mem_type.mem_type; +} + + +class rocm_hooks : public ucs::test { +protected: + + virtual void init() { + int dev_count; + ucs_status_t result; + hipError_t ret; + ucs::test::init(); + + ret = hipGetDeviceCount(&dev_count); + if (ret != hipSuccess || dev_count < 1) { + UCS_TEST_SKIP_R("no ROCm device detected"); + } + + if (hipSetDevice(0) != hipSuccess) { + UCS_TEST_SKIP_R("can't set ROCm device"); + } + + /* install memory hooks */ + result = ucm_set_event_handler(UCM_EVENT_MEM_TYPE_ALLOC, 0, rocm_mem_alloc_callback, + reinterpret_cast(this)); + ASSERT_UCS_OK(result); + + result = ucm_set_event_handler(UCM_EVENT_MEM_TYPE_FREE, 0, rocm_mem_free_callback, + reinterpret_cast(this)); + ASSERT_UCS_OK(result); + } + + virtual void cleanup() { + ucm_unset_event_handler(UCM_EVENT_MEM_TYPE_ALLOC, rocm_mem_alloc_callback, + reinterpret_cast(this)); + ucm_unset_event_handler(UCM_EVENT_MEM_TYPE_FREE, rocm_mem_free_callback, + reinterpret_cast(this)); + ucs::test::cleanup(); + } + + + void check_mem_alloc_events(void *ptr, size_t size, + int expect_mem_type = UCS_MEMORY_TYPE_ROCM) { + ASSERT_EQ(ptr, alloc_event.mem_type.address); + ASSERT_EQ(size, alloc_event.mem_type.size); + ASSERT_EQ(expect_mem_type, alloc_event.mem_type.mem_type); + } + + void check_mem_free_events(void *ptr, size_t size, + int expect_mem_type = UCS_MEMORY_TYPE_ROCM) { + ASSERT_EQ(ptr, free_event.mem_type.address); + ASSERT_EQ(expect_mem_type, free_event.mem_type.mem_type); + } + +}; + +UCS_TEST_F(rocm_hooks, test_hipMem_Alloc_Free) { + hipError_t ret; + void *dptr, *dptr1; + + /* small allocation */ + ret = hipMalloc(&dptr, 64); + ASSERT_EQ(ret, hipSuccess); + check_mem_alloc_events((void *)dptr, 64); + + ret = hipFree(dptr); + ASSERT_EQ(ret, hipSuccess); + check_mem_free_events((void *)dptr, 64); + + /* large allocation */ + ret = hipMalloc(&dptr, (256 * 1024 *1024)); + ASSERT_EQ(ret, hipSuccess); + check_mem_alloc_events((void *)dptr, (256 * 1024 *1024)); + + ret = hipFree(dptr); + ASSERT_EQ(ret, hipSuccess); + check_mem_free_events((void *)dptr, (256 * 1024 *1024)); + + /* multiple allocations, hipfree in reverse order */ + ret = hipMalloc(&dptr, (1 * 1024 *1024)); + ASSERT_EQ(ret, hipSuccess); + check_mem_alloc_events((void *)dptr, (1 * 1024 *1024)); + + ret = hipMalloc(&dptr1, (1 * 1024 *1024)); + ASSERT_EQ(ret, hipSuccess); + check_mem_alloc_events((void *)dptr1, (1 * 1024 *1024)); + + ret = hipFree(dptr1); + ASSERT_EQ(ret, hipSuccess); + check_mem_free_events((void *)dptr1, (1 * 1024 *1024)); + + ret = hipFree(dptr); + ASSERT_EQ(ret, hipSuccess); + check_mem_free_events((void *)dptr, (1 * 1024 *1024)); +} + +UCS_TEST_F(rocm_hooks, test_hipMallocManaged) { + hipError_t ret; + void * dptr; + + ret = hipMallocManaged(&dptr, 64); + ASSERT_EQ(ret, hipSuccess); + check_mem_alloc_events((void *)dptr, 64, UCS_MEMORY_TYPE_ROCM_MANAGED); + + ret = hipFree(dptr); + ASSERT_EQ(ret, hipSuccess); + check_mem_free_events((void *)dptr, 0, UCS_MEMORY_TYPE_ROCM_MANAGED); +} + +UCS_TEST_F(rocm_hooks, test_hipMallocPitch) { + hipError_t ret; + void * dptr; + size_t pitch; + + ret = hipMallocPitch(&dptr, &pitch, 4, 8); + ASSERT_EQ(ret, hipSuccess); + check_mem_alloc_events((void *)dptr, (128 * 8)); + + ret = hipFree(dptr); + ASSERT_EQ(ret, hipSuccess); + check_mem_free_events((void *)dptr, 0); +} + +UCS_TEST_F(rocm_hooks, test_hip_Malloc_Free) { + hipError_t ret; + void *ptr, *ptr1; + + /* small allocation */ + ret = hipMalloc(&ptr, 64); + ASSERT_EQ(ret, hipSuccess); + check_mem_alloc_events(ptr, 64); + + ret = hipFree(ptr); + ASSERT_EQ(ret, hipSuccess); + check_mem_free_events(ptr, 64); + + /* large allocation */ + ret = hipMalloc(&ptr, (256 * 1024 *1024)); + ASSERT_EQ(ret, hipSuccess); + check_mem_alloc_events(ptr, (256 * 1024 *1024)); + + ret = hipFree(ptr); + ASSERT_EQ(ret, hipSuccess); + check_mem_free_events(ptr, (256 * 1024 *1024)); + + /* multiple allocations, rocmfree in reverse order */ + ret = hipMalloc(&ptr, (1 * 1024 *1024)); + ASSERT_EQ(ret, hipSuccess); + check_mem_alloc_events(ptr, (1 * 1024 *1024)); + + ret = hipMalloc(&ptr1, (1 * 1024 *1024)); + ASSERT_EQ(ret, hipSuccess); + check_mem_alloc_events(ptr1, (1 * 1024 *1024)); + + ret = hipFree(ptr1); + ASSERT_EQ(ret, hipSuccess); + check_mem_free_events(ptr1, (1 * 1024 *1024)); + + ret = hipFree(ptr); + ASSERT_EQ(ret, hipSuccess); + check_mem_free_events(ptr, (1 * 1024 *1024)); + + /* hipFree with NULL */ + ret = hipFree(NULL); + ASSERT_EQ(ret, hipSuccess); +} +