From aabcb6488ea8f6d5b368b87e22d575a2e7d3c8e4 Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi <53337087+satyanveshd@users.noreply.github.com> Date: Wed, 23 Mar 2022 17:15:21 +0530 Subject: [PATCH] SWDEV-292714 - [catch2][dtest] Add unit test for hipPointerGetAttribute (#2485) Change-Id: Ic7490596f4f1ee641d9af2861a18380c4bfccd69 --- catch/unit/memory/CMakeLists.txt | 2 + catch/unit/memory/hipPtrGetAttribute.cc | 156 ++++++++++++++++++++++++ 2 files changed, 158 insertions(+) create mode 100644 catch/unit/memory/hipPtrGetAttribute.cc diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index 2a531232a8..5a5f48df72 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -62,6 +62,7 @@ set(TEST_SRC hipMemcpyFromSymbolAsync.cc hipMemcpyToSymbol.cc hipMemcpyToSymbolAsync.cc + hipPtrGetAttribute.cc ) else() set(TEST_SRC @@ -103,6 +104,7 @@ set(TEST_SRC hipMemcpyFromSymbolAsync.cc hipMemcpyToSymbol.cc hipMemcpyToSymbolAsync.cc + hipPtrGetAttribute.cc ) endif() diff --git a/catch/unit/memory/hipPtrGetAttribute.cc b/catch/unit/memory/hipPtrGetAttribute.cc new file mode 100644 index 0000000000..99dbc5ab8d --- /dev/null +++ b/catch/unit/memory/hipPtrGetAttribute.cc @@ -0,0 +1,156 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* + Run through few sanity tests to verify different attributes of hipPointerGetAttribute +*/ +#include +#include +#include +#include + +// Run few simple cases including host pointer arithmetic: +TEST_CASE("Unit_hipPtrGetAttribute_Simple") { + HIP_CHECK(hipSetDevice(0)); + size_t Nbytes = 0; + constexpr size_t N {1000000}; + Nbytes = N * sizeof(char); + printf("\n"); + + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + char* A_d; + char* A_Pinned_h; + char* A_Hmm; + + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_Pinned_h), Nbytes, + hipHostMallocDefault)); + HIP_CHECK(hipMallocManaged(&A_Hmm, Nbytes)); + + size_t free, total; + HIP_CHECK(hipMemGetInfo(&free, &total)); + printf("hipMemGetInfo: free=%zu (%4.2f) Nbytes=%lu total=%zu (%4.2f)\n", free, + (free / 1024.0 / 1024.0), Nbytes, total, + (total / 1024.0 / 1024.0)); + REQUIRE(free + Nbytes <= total); + + hipDeviceptr_t data = 0; + + // Device memory + printf("\nDevice memory (hipMalloc)\n"); + HIP_CHECK(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_DEVICE_POINTER, + reinterpret_cast(A_d))); + char *ptr1 = reinterpret_cast(data); + + // Check pointer arithmetic cases: + HIP_CHECK(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_DEVICE_POINTER, + reinterpret_cast(A_d + 100))); + char *ptr2 = reinterpret_cast(data); + REQUIRE(ptr2 == ptr1+100); + + // Corner case at end of array: + HIP_CHECK(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_DEVICE_POINTER, + reinterpret_cast(A_d + Nbytes - 1))); + ptr2 = reinterpret_cast(data); + REQUIRE(ptr2 == (ptr1 + Nbytes -1)); + + // Device-visible host memory + printf("\nDevice-visible host memory (hipHostMalloc)\n"); + HIP_CHECK(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_HOST_POINTER, + reinterpret_cast(A_Pinned_h))); + ptr1 = reinterpret_cast(data); + + // Check pointer arithmetic cases: + HIP_CHECK(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_HOST_POINTER, + reinterpret_cast(A_Pinned_h + 100))); + ptr2 = reinterpret_cast(data); + REQUIRE(ptr2 == ptr1+100); + + // Corner case at end of array: + HIP_CHECK(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_HOST_POINTER, + reinterpret_cast(A_Pinned_h + Nbytes - 1))); + ptr2 = reinterpret_cast(data); + REQUIRE(ptr2 == (ptr1 + Nbytes -1)); + + // HIP_POINTER_ATTRIBUTE_MEMORY_TYPE + unsigned int datatype; + HIP_CHECK(hipPointerGetAttribute(&datatype, HIP_POINTER_ATTRIBUTE_MEMORY_TYPE, + reinterpret_cast(A_d))); +#ifdef __HIP_PLATFORM_NVCC__ + REQUIRE(datatype == CU_MEMORYTYPE_DEVICE); +#else + REQUIRE(datatype == hipMemoryTypeDevice); +#endif + + HIP_CHECK(hipPointerGetAttribute(&datatype, HIP_POINTER_ATTRIBUTE_MEMORY_TYPE, + reinterpret_cast(A_Pinned_h))); +#ifdef __HIP_PLATFORM_NVCC__ + REQUIRE(datatype == CU_MEMORYTYPE_HOST); +#else + REQUIRE(datatype == hipMemoryTypeHost); +#endif + + // HIP_POINTER_ATTRIBUTE_IS_MANAGED + bool isHmm; + HIP_CHECK(hipPointerGetAttribute(&isHmm, HIP_POINTER_ATTRIBUTE_IS_MANAGED, + reinterpret_cast(A_Hmm))); + REQUIRE(isHmm == 1); + + HIP_CHECK(hipPointerGetAttribute(&isHmm, HIP_POINTER_ATTRIBUTE_IS_MANAGED, + reinterpret_cast(A_Pinned_h))); + REQUIRE(isHmm == 0); + + // HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL + if (numDevices > 1) { + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipPointerGetAttribute(&datatype, HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL, + reinterpret_cast(A_d))); + REQUIRE(datatype == 0); + } + + // HIP_POINTER_ATTRIBUTE_MAPPED + bool isMapped; + HIP_CHECK(hipPointerGetAttribute(&isMapped, HIP_POINTER_ATTRIBUTE_MAPPED, + reinterpret_cast(A_d))); + REQUIRE(isMapped == 1); + + // HIP_POINTER_ATTRIBUTE_RANGE_START_ADDR + HIP_CHECK(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_RANGE_START_ADDR, + reinterpret_cast(A_d+2))); + char *ptr3 = reinterpret_cast(data); + REQUIRE(ptr3 == A_d); + + // HIP_POINTER_ATTRIBUTE_RANGE_SIZE + HIP_CHECK(hipPointerGetAttribute(&datatype, HIP_POINTER_ATTRIBUTE_RANGE_SIZE, + reinterpret_cast(A_d))); + REQUIRE(datatype == Nbytes); + + unsigned int bufId1, bufId2; + // HIP_POINTER_ATTRIBUTE_BUFFER_ID + HIP_CHECK(hipPointerGetAttribute(&bufId1, HIP_POINTER_ATTRIBUTE_BUFFER_ID, + reinterpret_cast(A_d))); + HIP_CHECK(hipPointerGetAttribute(&bufId2, HIP_POINTER_ATTRIBUTE_BUFFER_ID, + reinterpret_cast(A_Pinned_h))); + REQUIRE(bufId1 != bufId2); + +}