diff --git a/catch/unit/memory/hipPointerGetAttributes.cc b/catch/unit/memory/hipPointerGetAttributes.cc index 42ab07e289..4c34a6edc4 100644 --- a/catch/unit/memory/hipPointerGetAttributes.cc +++ b/catch/unit/memory/hipPointerGetAttributes.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2021-2023 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 @@ -28,27 +28,40 @@ Following scenarios are verified for hipPointerGetAttributes API 4. Multi-threaded test with many simul allocs. */ +#include #include +#include +#ifdef __linux__ +#include +#endif #include #include #include - size_t Nbytes = 0; constexpr size_t N{1000000}; - +/** + * @addtogroup hipPointerGetAttributes hipPointerGetAttributes + * @{ + * @ingroup MemoryTest + * `hipError_t hipPointerGetAttributes (hipPointerAttribute_t *attributes, const void *ptr)` - + * Queries the memory pointer attributes. + */ //================================================================================================= // Utility Functions: //================================================================================================= -bool operator==(const hipPointerAttribute_t& lhs, const hipPointerAttribute_t& rhs) { - return ((lhs.hostPointer == rhs.hostPointer) && (lhs.devicePointer == rhs.devicePointer) && +bool operator==(const hipPointerAttribute_t& lhs, + const hipPointerAttribute_t& rhs) { + return ((lhs.hostPointer == rhs.hostPointer) && + (lhs.devicePointer == rhs.devicePointer) && (lhs.type == rhs.type) && (lhs.device == rhs.device) && (lhs.allocationFlags == rhs.allocationFlags)); } -bool operator!=(const hipPointerAttribute_t& lhs, const hipPointerAttribute_t& rhs) { +bool operator!=(const hipPointerAttribute_t& lhs, + const hipPointerAttribute_t& rhs) { return !(lhs == rhs); } @@ -79,7 +92,8 @@ void printAttribs(const hipPointerAttribute_t* attribs) { printf( "hostPointer:%p devicePointer:%p type:%s deviceId:%d isManaged:%d " "allocationFlags:%u\n", - attribs->hostPointer, attribs->devicePointer, memoryTypeToString(attribs->type), + attribs->hostPointer, attribs->devicePointer, + memoryTypeToString(attribs->type), attribs->device, attribs->isManaged, attribs->allocationFlags); } @@ -98,7 +112,8 @@ struct SuperPointerAttribute { // Support function to check result against a reference: -void checkPointer(const SuperPointerAttribute& ref, int major, int minor, void* pointer) { +void checkPointer(const SuperPointerAttribute& ref, int major, + int minor, void* pointer) { hipPointerAttribute_t attribs; resetAttribs(&attribs); @@ -122,8 +137,9 @@ void checkPointer(const SuperPointerAttribute& ref, int major, int minor, void* // we do this in the testMultiThreaded_1 test. void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) { Nbytes = N * sizeof(char); - printf("clusterAllocs numAllocs=%d size=%lu..%lu\n", - numAllocs, static_cast(minSize), static_cast(maxSize)); + printf("clusterAllocs numAllocs=%d size=" "%" PRIu64 ".." "%" PRIu64 "\n", + numAllocs, static_cast(minSize), + static_cast(maxSize)); const int Max_Devices = 256; std::vector reference(numAllocs); @@ -150,14 +166,17 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) { void* ptr; if (isDevice) { - totalDeviceAllocated[reference[i]._attrib.device] += reference[i]._sizeBytes; - HIP_CHECK(hipMalloc(reinterpret_cast(&ptr), reference[i]._sizeBytes)); + totalDeviceAllocated[reference[i]._attrib.device] += + reference[i]._sizeBytes; + HIP_CHECK(hipMalloc(reinterpret_cast(&ptr), + reference[i]._sizeBytes)); reference[i]._attrib.type = hipMemoryTypeDevice; reference[i]._attrib.devicePointer = ptr; reference[i]._attrib.hostPointer = NULL; reference[i]._attrib.allocationFlags = 0; } else { - HIP_CHECK(hipHostMalloc(reinterpret_cast(&ptr), reference[i]._sizeBytes, + HIP_CHECK(hipHostMalloc(reinterpret_cast(&ptr), + reference[i]._sizeBytes, hipHostMallocDefault)); reference[i]._attrib.type = hipMemoryTypeHost; reference[i]._attrib.devicePointer = ptr; @@ -173,10 +192,12 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) { HIP_CHECK(hipMemGetInfo(&free, &total)); printf( " device#%d: hipMemGetInfo: " - "free=%zu (%4.2fMB) totalDevice=%lu (%4.2fMB) total=%zu " + "free=%zu (%4.2fMB) totalDevice=""%" PRIu64 "(%4.2fMB) total=%zu " "(%4.2fMB)\n", - i, free, (free / 1024.0 / 1024.0), static_cast(totalDeviceAllocated[i]), - (totalDeviceAllocated[i]) / 1024.0 / 1024.0, total, (total / 1024.0 / 1024.0)); + i, free, (free / 1024.0 / 1024.0), + static_cast(totalDeviceAllocated[i]), + (totalDeviceAllocated[i]) / 1024.0 / 1024.0, total, + (total / 1024.0 / 1024.0)); REQUIRE(free + totalDeviceAllocated[i] <= total); } @@ -198,12 +219,17 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) { } } } - -//======================================================================== -// Functions to run tests -//======================================================================= -//-- -// Run through a couple simple cases to test lookups host pointer arithmetic: +/** + * Test Description + * ------------------------ + * - Run through a couple simple cases to test lookups host pointer arithmetic +// specified size range + * ------------------------ + * - unit/memory/hipPointerGetAttributes.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.7 + */ TEST_CASE("Unit_hipPointerGetAttributes_Basic") { HIP_CHECK(hipSetDevice(0)); Nbytes = N * sizeof(char); @@ -217,12 +243,15 @@ TEST_CASE("Unit_hipPointerGetAttributes_Basic") { hipError_t e; HIP_CHECK(hipMalloc(&A_d, Nbytes)); - HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_Pinned_h), Nbytes, hipHostMallocDefault)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_Pinned_h), + Nbytes, hipHostMallocDefault)); 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), static_cast(Nbytes), total, (total / 1024.0 / 1024.0)); + printf("hipMemGetInfo: free=%zu (%4.2f) Nbytes=""%" PRIu64 + "total=%zu (%4.2f)\n", free, + (free / 1024.0 / 1024.0), static_cast(Nbytes), + total, (total / 1024.0 / 1024.0)); REQUIRE(free + Nbytes <= total); @@ -242,14 +271,16 @@ TEST_CASE("Unit_hipPointerGetAttributes_Basic") { // Corner case at end of array: resetAttribs(&attribs2); HIP_CHECK(hipPointerGetAttributes(&attribs2, A_d + Nbytes - 1)); - REQUIRE((ptr + Nbytes - 1) == reinterpret_cast(attribs2.devicePointer)); + REQUIRE((ptr + Nbytes - 1) == reinterpret_cast + (attribs2.devicePointer)); // Pointer just beyond array must be invalid or at least a different pointer resetAttribs(&attribs2); e = hipPointerGetAttributes(&attribs2, A_d + Nbytes + 1); if (e != hipErrorInvalidValue) { // We might have strayed into another pointer area. - REQUIRE(reinterpret_cast(ptr) != reinterpret_cast(attribs2.devicePointer)); + REQUIRE(reinterpret_cast(ptr) != + reinterpret_cast(attribs2.devicePointer)); } @@ -274,12 +305,34 @@ TEST_CASE("Unit_hipPointerGetAttributes_Basic") { // OS memory printf("\nOS-allocated memory (malloc)\n"); } +/** + * Test Description + * ------------------------ + * - Test that allocates memory across all devices withing the specified size range + * Test source + * ------------------------ + * - unit/memory/hipPointerGetAttributes.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.7 + */ TEST_CASE("Unit_hipPointerGetAttributes_ClusterAlloc") { srand(0x100); printf("\n=============================================\n"); clusterAllocs(100, 1024 * 1, 1024 * 1024); } +/** + * Test Description + * ------------------------ + * - Test that allocates memory across all devices withing the specified size range + * Test source + * ------------------------ + * - unit/memory/hipPointerGetAttributes.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.7 + */ TEST_CASE("Unit_hipPointerGetAttributes_TinyClusterAlloc") { srand(0x200); @@ -289,7 +342,8 @@ TEST_CASE("Unit_hipPointerGetAttributes_TinyClusterAlloc") { // Multi-threaded test with many simul allocs. // IN : serialize will force the test to run in serial fashion. -#if 0 // FIXME_jatinx These need to be ported to HIP_CHECK_THREAD. Disabling it for now +#if 0 // FIXME_jatinx These need to be ported to HIP_CHECK_THREAD. +Disabling it for now TEST_CASE("Unit_hipPointerGetAttributes_MultiThread") { srand(0x300); auto serialize = 1; @@ -310,55 +364,77 @@ TEST_CASE("Unit_hipPointerGetAttributes_MultiThread") { if (serialize) t4.join(); } #endif +/** + * Test Description + * ------------------------ + * - Perform negative tests for the API hipPointerGetAttributes + * by passing invalid memory pointer and invalid attributes + * Test source + * ------------------------ + * - unit/memory/hipPointerGetAttributes.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.7 + */ TEST_CASE("Unit_hipPointerGetAttributes_Negative") { #if HT_AMD // Nvidia crashed in hipPointerGetAttributes on nullptr SECTION("Invalid Attributes Pointer") { int* dPtr{nullptr}; HIP_CHECK(hipMalloc(&dPtr, sizeof(int))); - HIP_CHECK_ERROR(hipPointerGetAttributes(nullptr, dPtr), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipPointerGetAttributes(nullptr, dPtr), + hipErrorInvalidValue); HIP_CHECK(hipFree(dPtr)); } #endif SECTION("Invalid Device Pointer") { hipPointerAttribute_t attributes{}; - HIP_CHECK_ERROR(hipPointerGetAttributes(&attributes, nullptr), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipPointerGetAttributes(&attributes, nullptr), + hipErrorInvalidValue); } } -// Run this test for all devices for DeviceMemory, HostMemory, ManagedMemory and MappedMemory +/** + * Test Description + * ------------------------ + * - Run this test for all devices for DeviceMemory, + * HostMemory and MappedMemory + * Test source + * ------------------------ + * - unit/memory/hipPointerGetAttributes.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipPointerGetAttributes_GpuIter") { int deviceCount{0}; HIP_CHECK(hipGetDeviceCount(&deviceCount)); REQUIRE(deviceCount != 0); // Memory Types - enum MemoryTypes { DeviceMemory = 0, HostMemory, MappedMemory }; + enum MemoryTypes {DeviceMemory = 0, HostMemory, MappedMemory}; auto MemoryType = - GENERATE(MemoryTypes::DeviceMemory, MemoryTypes::HostMemory, MemoryTypes::MappedMemory); + GENERATE(MemoryTypes::DeviceMemory, MemoryTypes::HostMemory, + MemoryTypes::MappedMemory); INFO("Memory Type: " << MemoryType); int* ptr{nullptr}; for (int i = 0; i < deviceCount; i++) { HIP_CHECK(hipSetDevice(i)); - ptr = nullptr; if (MemoryType == MemoryTypes::DeviceMemory) { HIP_CHECK(hipMalloc(&ptr, sizeof(int))); } else if (MemoryType == MemoryTypes::HostMemory) { - HIP_CHECK(hipHostMalloc(&ptr, sizeof(int))); + HIP_CHECK(hipHostMalloc(&ptr, sizeof(int))); } else if (MemoryType == MemoryTypes::MappedMemory) { - HIP_CHECK(hipHostMalloc(&ptr, sizeof(int), hipHostMallocMapped)); + HIP_CHECK(hipHostMalloc(&ptr, sizeof(int), hipHostMallocMapped)); } - REQUIRE(ptr != nullptr); hipPointerAttribute_t attributes{}; HIP_CHECK(hipPointerGetAttributes(&attributes, ptr)); - REQUIRE(attributes.device == i); // Device Check - // Memory address and type check if (MemoryType == MemoryTypes::DeviceMemory) { REQUIRE(attributes.type == hipMemoryTypeDevice); @@ -369,15 +445,116 @@ TEST_CASE("Unit_hipPointerGetAttributes_GpuIter") { REQUIRE(attributes.hostPointer == ptr); } else if (MemoryType == MemoryTypes::MappedMemory) { int* devicePtr{nullptr}; - HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&devicePtr), ptr, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&devicePtr), + ptr, 0)); REQUIRE(attributes.hostPointer == ptr); REQUIRE(attributes.hostPointer == devicePtr); } if (MemoryType == MemoryTypes::DeviceMemory) { HIP_CHECK(hipFree(ptr)); - } else if (MemoryType == MemoryTypes::MappedMemory || MemoryType == MemoryTypes::HostMemory) { + } else if (MemoryType == MemoryTypes::MappedMemory || + MemoryType == MemoryTypes::HostMemory) { HIP_CHECK(hipHostFree(ptr)); } } } +/** + * Test Description + * ------------------------ + * - Managed memory type attribute verification. + * Test source + * ------------------------ + * - unit/memory/hipPointerGetAttributes.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipPointerGetAttributes_GpuIter_Managed__Memory") { + int deviceCount{0}; + HIP_CHECK(hipGetDeviceCount(&deviceCount)); + REQUIRE(deviceCount != 0); + int* managed_ptr{nullptr}; + for (int i = 0; i < deviceCount; i++) { + HIP_CHECK(hipSetDevice(i)); + int managed_memory = 0; + HIP_CHECK(hipDeviceGetAttribute(&managed_memory, + hipDeviceAttributeManagedMemory, i)); + if (!managed_memory) { + INFO("Managed memory access not supported on the device" << i); + } else { + // Allocating the memory and test changes only if managed + // memory is available + HIP_CHECK(hipMallocManaged(&managed_ptr, sizeof(int))); + REQUIRE(managed_ptr != nullptr); + hipPointerAttribute_t attributes{}; + HIP_CHECK(hipPointerGetAttributes(&attributes, managed_ptr)); + REQUIRE(attributes.device == i); // Device Check + REQUIRE(attributes.isManaged); + REQUIRE(attributes.type == hipMemoryTypeManaged); + REQUIRE(attributes.devicePointer == managed_ptr); + HIP_CHECK(hipFree(managed_ptr)); + } + } +} +/** + * Test Description + * ------------------------ + * - Unregistered memory type attribute verification. + * Test source + * ------------------------ + * - unit/memory/hipPointerGetAttributes.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipPointerGetAttributes_GpuIter_Unregistered_Memory") { + int deviceCount{0}; + HIP_CHECK(hipGetDeviceCount(&deviceCount)); + REQUIRE(deviceCount != 0); + int* ptr1{nullptr}; int* ptr2{nullptr}; int* ptr3{nullptr}; + for (int i = 0; i < deviceCount; i++) { + HIP_CHECK(hipSetDevice(i)); + int N = 5; + ptr1 = new int; + ptr2 = reinterpret_cast(malloc(sizeof(int))); + ptr3 = reinterpret_cast(calloc(N, sizeof(int))); + int ptr4[1000]; + #ifdef __linux__ + void* ptr5{nullptr}; + ptr5 = mmap ( NULL, N*sizeof(int), + PROT_READ | PROT_WRITE, MAP_PRIVATE | MAP_ANONYMOUS, 0, 0); + + if (ptr5 == MAP_FAILED) { + INFO("Mapping Failed\n"); + REQUIRE(false); + } + REQUIRE(ptr5 != nullptr); + hipPointerAttribute_t attributes1{}; + HIP_CHECK(hipPointerGetAttributes(&attributes1, ptr5)); + REQUIRE(attributes1.type == hipMemoryTypeUnregistered); + int err = munmap(ptr5, N*sizeof(int)); + if (err != 0) { + INFO("UnMapping Failed\n"); + REQUIRE(false); + } + #endif + REQUIRE(ptr1 != nullptr); + REQUIRE(ptr2 != nullptr); + REQUIRE(ptr3 != nullptr); + + hipPointerAttribute_t attributes{}; + HIP_CHECK(hipPointerGetAttributes(&attributes, ptr1)); + REQUIRE(attributes.type == hipMemoryTypeUnregistered); + HIP_CHECK(hipPointerGetAttributes(&attributes, ptr2)); + REQUIRE(attributes.type == hipMemoryTypeUnregistered); + HIP_CHECK(hipPointerGetAttributes(&attributes, ptr3)); + REQUIRE(attributes.type == hipMemoryTypeUnregistered); + HIP_CHECK(hipPointerGetAttributes(&attributes, ptr4)); + REQUIRE(attributes.type == hipMemoryTypeUnregistered); + + delete ptr1; + free(ptr2); + free(ptr3); + } +}