Create address tracker for am_alloc.

Tracks device where memory is allocated, pinned-host or device, and more. Uses memory-range-based lookups - so pointers that exist anywhere in the range of hostPtr + size will find the associated AmPointerInfo. The insertions and lookups use a self-balancing binary tree and should support O(logN) lookup speed. [ROCm/hip commit: 4ee2a5229b]
2016-02-10 11:52:42 -06:00
parent 0a6e6e3b7e
commit d4a90f8afd
11 changed files with 743 additions and 11 deletions
@@ -71,6 +71,7 @@ if ($HIP_PLATFORM eq "hcc") {
 	$HIPLDFLAGS .= " -L$HSA_PATH/lib -lhsa-runtime64 -lhc_am";
 	# Add C++ libs for GCC.
 	$HIPLDFLAGS .= " -lstdc++";
+	$HIPLDFLAGS .= " -lm";

    if ($verbose & 0x2) {
        print ("HSA_PATH=$HSA_PATH\n");
@@ -277,6 +277,8 @@ while (@ARGV) {

        $ft{'mem'} += s/\bcudaMemcpyKind\b/hipMemcpyKind/g;

+        $ft{'mem'} += s/\bcudaPointerAttributes\b/hipPointerAttribute_t/g;
+

        #--------
        # Memory management:
@@ -0,0 +1,92 @@
+#pragma once
+
+#include <hc_am.hpp>
+
+typedef int am_status_t;
+#define AM_SUCCESS                           0
+// TODO - provide better mapping of HSA error conditions to HC error codes.
+#define AM_ERROR_MISC                       -1 /** Misellaneous error */
+
+// Flags for am_alloc API:
+#define amHostPinned 0x1
+
+
+namespace hc {
+
+// This is the data that is maintained for each pointer:
+struct AmPointerInfo {
+    bool        _isDeviceMem;
+    void *      _hostPointer;
+    void *      _devicePointer;
+    size_t      _sizeBytes;
+    hc::accelerator _acc;
+    unsigned    _allocationFlags;
+
+    AmPointerInfo() {};
+
+    AmPointerInfo(bool isDeviceMem, void *hostPointer, void *devicePointer, size_t sizeBytes, hc::accelerator acc, unsigned allocationFlags) :
+        _isDeviceMem(isDeviceMem),
+        _hostPointer(hostPointer),
+        _devicePointer(devicePointer),
+        _sizeBytes(sizeBytes),
+        _acc(acc),
+        _allocationFlags(allocationFlags)  {};
+};
+}
+
+
+
+namespace hc {
+
+
+/**
+ * Allocates a block of @p size bytes of memory on the specified @p acc.
+ *
+ * The contents of the newly allocated block of memory are not initialized.
+ *
+ * If @p size == 0, 0 is returned.
+ *
+ * Flags must be 0.
+ *
+ * @returns : On success, pointer to the newly allocated memory is returned.
+ * The pointer is typecast to the desired return type.
+ *
+ * If an error occurred trying to allocate the requested memory, 0 is returned.
+ *
+ * @see am_free, am_copy
+ */
+auto_voidp AM_alloc(size_t size, hc::accelerator acc, unsigned flags);
+
+/**
+ * Frees a block of memory previously allocated with am_alloc.
+ *
+ * @see am_alloc, am_copy
+ */
+am_status_t AM_free(void*  ptr);
+
+
+/**
+ * Copies @p size bytes of memory from @p src to @ dst.  The memory areas (src+size and dst+size) must not overlap.
+ *
+ * @returns AM_SUCCESS on error or AM_ERROR_MISC if an error occurs.
+ * @see am_alloc, am_free
+ */
+am_status_t AM_copy(void*  dst, const void*  src, size_t size);
+
+am_status_t AM_get_pointer_info(hc::AmPointerInfo *info, void *ptr);
+
+
+// TODO-implement these:
+//am_status_t AM_track_pointer(void* ptr, size_t size, bool isDeviceMem=false, unsigned allocationFlags=0x0);
+//am_status_t AM_untrack_pointer(void* ptr);
+
+/**
+ * Prints the contents of the memory tracker table to stderr
+ *
+ * Intended primarily for debug purposes.
+ **/
+void AM_print_tracker();
+
+
+}; // namespace hc
+
@@ -105,6 +105,8 @@ enum hipMemcpyKind {
 } ;


+
+
 // Doxygen end group GlobalDefs
 /**  @} */

@@ -128,6 +130,7 @@ typedef struct hipEvent_t {



+
 #ifdef __cplusplus
 } /* extern "C" */
 #endif
@@ -634,6 +637,11 @@ hipError_t hipEventQuery(hipEvent_t event) ;
 */


+/**
+ *  @brief Return attributes for the specified pointer
+ */
+hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr) ;
+

 /**
 *  @brief Allocate memory on the default accelerator
@@ -97,6 +97,30 @@ typedef struct hipDeviceProp_t {
 } hipDeviceProp_t;


+/**
+ * Memory type (for pointer attributes)
+ */
+enum hipMemoryType {
+    hipMemoryTypeHost,   ///< Memory is physically located on host
+    hipMemoryTypeDevice  ///< Memory is physically located on device. (see deviceId for specific device)
+};
+
+
+
+/**
+ * Pointer attributes
+ */
+typedef struct hipPointerAttribute_t {
+    enum hipMemoryType memoryType;
+    int device;
+    void *devicePointer;
+    void *hostPointer;
+    int isManaged;
+    unsigned allocationFlags; /* flags specified when memory was allocated*/
+    /* peers? */
+} hipPointerAttribute_t;
+
+
 // hack to get these to show up in Doxygen:
 /**
 *     @defgroup GlobalDefs Global enum and defines
@@ -0,0 +1,219 @@
+
+#include "hc_am.hpp"
+#include "hsa.h"
+
+
+#include "hcc_detail/AM.h" // TODO - Remove me.
+
+#define DB_TRACKER 1
+
+#if DB_TRACKER 
+#define mprintf( ...) {\
+        fprintf (stderr, __VA_ARGS__);\
+        };
+#else
+#define mprintf( ...) 
+#endif
+
+//=========================================================================================================
+// Pointer Tracker Structures:
+//=========================================================================================================
+#include <map>
+#include <iostream>
+//#include <shared_mutex>
+
+struct AmMemoryRange {
+    void * _basePointer;
+    void * _endPointer;
+    AmMemoryRange(void *basePointer, size_t sizeBytes) :
+        _basePointer(basePointer), _endPointer((unsigned char*)basePointer + sizeBytes - 1) {};
+};
+
+// Functor to compare ranges:
+struct AmMemoryRangeCompare {
+    // Return true is LHS range is less than RHS - used to order the 
+    bool operator()(const AmMemoryRange &lhs, const AmMemoryRange &rhs) const
+    {
+        return lhs._endPointer < rhs._basePointer;
+    }
+
+};
+
+
+std::ostream &operator<<(std::ostream &os, const hc::AmPointerInfo &ap)
+{
+    os << "hostPointer:" << ap._hostPointer << " devicePointer:"<< ap._devicePointer << " sizeBytes:" << ap._sizeBytes
+       << " isDeviceMem:" << ap._isDeviceMem  << " allocFlags:" << ap._allocationFlags;
+    return os;
+}
+
+
+
+// This structure tracks information for each pointer.
+// Uses memory-range-based lookups - so pointers that exist anywhere in the range of hostPtr + size will find the associated AmPointerInfo.
+// The insertions and lookups use a self-balancing binary tree and should support O(logN) lookup speed.
+// The structure is thread-safe - writers obtain a mutex before modifying the tree.  Multiple simulatenous readers are supported.
+class AmPointerTracker {
+typedef std::map<AmMemoryRange, hc::AmPointerInfo, AmMemoryRangeCompare> MapTrackerType;
+public:
+
+    void insert(void *pointer, const hc::AmPointerInfo &p);
+    int remove(void *pointer);
+
+    MapTrackerType::iterator find(void *hostPtr);
+
+    MapTrackerType::iterator end() { return _tracker.end(); };
+
+    std::ostream & print (std::ostream &os);
+private:
+    MapTrackerType  _tracker;
+    //std::shared_timed_mutex _mut;
+};
+
+
+//---
+void AmPointerTracker::insert (void *pointer, const hc::AmPointerInfo &p)
+{
+    // TODO-mutex - write lock.
+    mprintf ("insert: %p + %zu\n", pointer, p._sizeBytes);
+    _tracker.insert(std::make_pair(AmMemoryRange(pointer, p._sizeBytes), p));
+
+
+}
+
+
+//---
+// Return 1 if removed or 0 if not found.
+int AmPointerTracker::remove (void *pointer)
+{
+    // TODO-mutex - write lock.
+    mprintf ("remove: %p\n", pointer);
+    return _tracker.erase(AmMemoryRange(pointer,1));
+}
+
+
+//---
+AmPointerTracker::MapTrackerType::iterator  AmPointerTracker::find (void *pointer)
+{
+    // TODO-mutex- read lock
+    auto iter = _tracker.find(AmMemoryRange(pointer,1));
+    mprintf ("find: %p\n", pointer);
+    return iter;
+}
+
+
+std::ostream & AmPointerTracker::print (std::ostream &os)
+{
+    for (auto iter = _tracker.begin() ; iter != _tracker.end(); iter++) {
+        os << "  " << iter->first._basePointer << "..." << iter->first._endPointer << "::  ";
+        os << iter->second << std::endl;
+    }
+
+    return os;
+}
+
+
+//=========================================================================================================
+// Global var defs:
+//=========================================================================================================
+AmPointerTracker g_amPointerTracker;  // Track all am pointer allocations.
+
+
+//=========================================================================================================
+// API Definitions.
+//=========================================================================================================
+//
+//
+
+namespace hc {
+
+// Allocate accelerator memory, return NULL if memory could not be allocated:
+auto_voidp AM_alloc(size_t sizeBytes, hc::accelerator acc, unsigned flags) 
+{
+
+    void *ptr = NULL;
+
+    if (sizeBytes != 0 ) {
+        if (acc.is_hsa_accelerator()) {
+            hsa_agent_t *hsa_agent = static_cast<hsa_agent_t*> (acc.get_default_view().get_hsa_agent());
+            hsa_region_t *alloc_region;
+            if (flags & amHostPinned) {
+               alloc_region = static_cast<hsa_region_t*>(acc.get_hsa_am_system_region());
+            } else {
+               alloc_region = static_cast<hsa_region_t*>(acc.get_hsa_am_region());
+            }
+
+            if (alloc_region->handle != -1) {
+
+                hsa_status_t s1 = hsa_memory_allocate(*alloc_region, sizeBytes, &ptr);
+                hsa_status_t s2 = hsa_memory_assign_agent(ptr, *hsa_agent, HSA_ACCESS_PERMISSION_RW);
+
+                if ((s1 != HSA_STATUS_SUCCESS) || (s2 != HSA_STATUS_SUCCESS)) {
+                    ptr = NULL;
+                } else {
+                    if (flags & amHostPinned) {
+                        g_amPointerTracker.insert(ptr, 
+                                hc::AmPointerInfo(false/*isDevice*/, ptr/*hostPointer*/,  ptr /*devicePointer*/, sizeBytes, acc, flags));
+                    } else {
+                        g_amPointerTracker.insert(ptr, 
+                                hc::AmPointerInfo(true/*isDevice*/, NULL/*hostPointer*/,  ptr /*devicePointer*/, sizeBytes, acc, flags));
+                    }
+                }
+            }
+        }
+    }
+
+    return ptr;
+};
+
+
+am_status_t AM_free(void* ptr) 
+{
+    am_status_t status = AM_SUCCESS;
+
+    if (ptr != NULL) {
+        hsa_memory_free(ptr);
+
+        size_t numRemoved = g_amPointerTracker.remove(ptr) ;
+        if (numRemoved == 0) {
+            status = AM_ERROR_MISC;
+        }
+    }
+    return status;
+}
+
+
+
+am_status_t AM_copy(void*  dst, const void*  src, size_t sizeBytes)
+{
+    am_status_t am_status = AM_ERROR_MISC;
+    hsa_status_t err = hsa_memory_copy(dst, src, sizeBytes);
+
+    if (err == HSA_STATUS_SUCCESS) {
+        am_status = AM_SUCCESS;
+    } else {
+        am_status = AM_ERROR_MISC;
+    }
+
+    return am_status;
+}
+
+
+am_status_t AM_get_pointer_info(hc::AmPointerInfo *info, void *ptr)
+{
+    auto infoI = g_amPointerTracker.find(ptr);
+    if (infoI != g_amPointerTracker.end()) {
+        *info = infoI->second;
+        return AM_SUCCESS;
+    } else {
+        return AM_ERROR_MISC;
+    }
+}
+
+void AM_print_tracker()
+{
+    g_amPointerTracker.print(std::cerr);
+}
+
+
+} // end namespace hc.
@@ -31,6 +31,8 @@ THE SOFTWARE.
 #include <list>
 #include <sys/types.h>
 #include <unistd.h>
+#include <unordered_map>
+
 #include <hc.hpp>
 #include <hc_am.hpp>

@@ -38,6 +40,9 @@ THE SOFTWARE.

 #include "hsa_ext_amd.h"

+
+#include "hc_AM.cpp"
+
 #define USE_PINNED_HOST (__hcc_workweek__ >= 1601)

 #define USE_ASYNC_COPY  0
@@ -466,7 +471,8 @@ void ihipInit()
    g_devices.reserve(accs.size());
    for (int i=0; i<accs.size(); i++) {
        if (! accs[i].get_is_emulated()) {
-           g_devices.emplace_back(ihipDevice_t(g_devices.size(), accs[i]));
+            int deviceId = g_devices.size();
+           g_devices.emplace_back(ihipDevice_t(deviceId, accs[i]));
        }
    }

@@ -1262,6 +1268,53 @@ hipError_t hipEventQuery(hipEvent_t event)
 // Memory
 //
 //
+//
+
+//---
+/**
+ * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDevice
+ */
+hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr) 
+{
+    std::call_once(hip_initialized, ihipInit);
+
+    hipError_t e = hipSuccess;
+
+    hc::AmPointerInfo amPointerInfo;
+    am_status_t status = hc::AM_get_pointer_info(&amPointerInfo, ptr);
+    if (status == AM_SUCCESS) {
+
+        attributes->memoryType    = amPointerInfo._isDeviceMem ? hipMemoryTypeDevice: hipMemoryTypeHost;
+        attributes->hostPointer   = amPointerInfo._hostPointer;
+        attributes->devicePointer = amPointerInfo._devicePointer;
+        attributes->isManaged     = 0;
+        attributes->allocationFlags = amPointerInfo._allocationFlags;
+
+
+        attributes->device        = -1;
+        e = hipErrorInvalidDevice;
+        for (int i=0; i<g_devices.size(); i++) {
+            if (g_devices[i]._acc == amPointerInfo._acc) {
+                attributes->device = i;
+                e = hipSuccess;
+                break;
+            }
+        }
+    } else {
+        attributes->memoryType    = hipMemoryTypeDevice;
+        attributes->hostPointer   = 0;
+        attributes->devicePointer = 0;
+        attributes->device        = -1;
+        attributes->isManaged     = 0;
+        attributes->allocationFlags = 0;
+
+        e = hipErrorInvalidValue;
+    }
+
+    return ihipLogStatus(e);
+}
+
+

 // kernel for launching memcpy operations:
 template <typename T>
@@ -1345,9 +1398,9 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
    hipError_t  hip_status = hipSuccess;

    const unsigned am_flags = 0;
-    *ptr = hc::am_alloc(sizeBytes, ihipGetTlsDefaultDevice()->_acc, am_flags);
+    *ptr = hc::AM_alloc(sizeBytes, ihipGetTlsDefaultDevice()->_acc, am_flags);

-    if (*ptr == NULL) {
+    if (sizeBytes && (*ptr == NULL)) {
        hip_status = hipErrorMemoryAllocation;
    } else {
        hip_status = hipSuccess;
@@ -1367,9 +1420,9 @@ hipError_t hipMallocHost(void** ptr, size_t sizeBytes)

    const unsigned am_flags = amHostPinned;

-    *ptr = hc::am_alloc(sizeBytes, ihipGetTlsDefaultDevice()->_acc, am_flags);
+    *ptr = hc::AM_alloc(sizeBytes, ihipGetTlsDefaultDevice()->_acc, am_flags);
    hipError_t  hip_status = hipSuccess;
-    if (*ptr == NULL) {
+    if (sizeBytes && (*ptr == NULL)) {
        hip_status = hipErrorMemoryAllocation;
    } else {
        hip_status = hipSuccess;
@@ -1444,7 +1497,7 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind

 #else
    // TODO-hsart - what synchronization does hsa_copy provide?
-    hc::am_copy(dst, src, sizeBytes);
+    hc::AM_copy(dst, src, sizeBytes);
    e = hipSuccess;
 #endif

@@ -1475,7 +1528,7 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp

    // TODO-hsart This routine needs to ensure that dst and src are mapped on the GPU.
    // This is a synchronous copy - remove and replace with code below when we have appropriate LOCK APIs.
-    hc::am_copy(dst, src, sizeBytes);
+    hc::AM_copy(dst, src, sizeBytes);

 #if 0

@@ -1592,7 +1645,7 @@ hipError_t hipFree(void* ptr)
    ihipWaitAllStreams(ihipGetTlsDefaultDevice());

    if (ptr) {
-        hc::am_free(ptr);
+        hc::AM_free(ptr);
    }

    return ihipLogStatus(hipSuccess);
@@ -1606,7 +1659,7 @@ hipError_t hipFreeHost(void* ptr)
    if (ptr) {
 #if USE_PINNED_HOST
        tprintf (TRACE_MEM, "  %s: %p\n", __func__, ptr);
-        hc::am_free(ptr);
+        hc::AM_free(ptr);
 #else
        free(ptr);
 #endif
@@ -114,6 +114,7 @@ make_hip_executable (hipSimpleAtomicsTest hipSimpleAtomicsTest.cpp)
 make_hip_executable (hipMathFunctionsHost hipMathFunctions.cpp hipSinglePrecisionMathHost.cpp hipDoublePrecisionMathHost.cpp)
 make_hip_executable (hipMathFunctionsDevice hipMathFunctions.cpp hipSinglePrecisionMathDevice.cpp hipDoublePrecisionMathDevice.cpp)
 make_hip_executable (hipIntrinsics hipMathFunctions.cpp hipSinglePrecisionIntrinsics.cpp hipDoublePrecisionIntrinsics.cpp hipIntegerIntrinsics.cpp)
+make_hip_executable (hipPointerAttrib hipPointerAttrib.cpp) 
 target_link_libraries(hipMathFunctionsHost m)

 make_test(hip_ballot " " )
@@ -128,6 +129,7 @@ make_test(hipMemset --N 10    --memsetval 0x42 )  # small copy, just 10 bytes.
 make_test(hipMemset --N 10013 --memsetval 0x5a )  # oddball size.
 make_test(hipMemset --N 256M  --memsetval 0xa6 )  # big copy
 make_test(hipGridLaunch " " )
+make_test(hipPointerAttrib " " )

 make_test(hipMemcpy  " " )

@@ -0,0 +1,319 @@
+/*
+Copyright (c) 2015-2016 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.
+*/
+
+
+// Test pointer tracking logic: allocate memory and retrieve stats with hipPointerGetAttributes
+
+#include "hip_runtime.h"
+#include "test_common.h"
+
+#ifdef __HIP_PLATFORM_HCC__
+#include "hcc_detail/AM.h"
+#endif
+
+size_t Nbytes = 0;
+
+//=================================================================================================
+// Utility Functions:
+//=================================================================================================
+
+bool operator==(const hipPointerAttribute_t &lhs, const hipPointerAttribute_t &rhs) 
+{
+    return ((lhs.hostPointer == rhs.hostPointer) &&
+            (lhs.devicePointer == rhs.devicePointer) &&
+            (lhs.memoryType == rhs.memoryType) &&
+            (lhs.device == rhs.device) &&
+            (lhs.allocationFlags == rhs.allocationFlags)
+            ) ;
+
+};
+
+
+bool operator!=(const hipPointerAttribute_t &lhs, const hipPointerAttribute_t &rhs) 
+{
+    return ! (lhs == rhs);
+}
+
+
+const char *memoryTypeToString(hipMemoryType memoryType)
+{
+    switch (memoryType) {
+        case hipMemoryTypeHost   : return "[Host]";
+        case hipMemoryTypeDevice : return "[Device]";
+        default:                   return "[Unknown]";
+    };
+}
+
+
+void resetAttribs(hipPointerAttribute_t *attribs) 
+{
+    attribs->hostPointer = (void*) (-1);
+    attribs->devicePointer = (void*) (-1);
+    attribs->memoryType = hipMemoryTypeHost;
+    attribs->device = -2;
+    attribs->isManaged = -1;
+    attribs->allocationFlags = 0xffff;
+};
+
+
+void printAttribs(hipPointerAttribute_t *attribs) 
+{
+    printf ("hostPointer:%p devicePointer:%p  memoryType:%s deviceId:%d isManaged:%d allocationFlags:%u\n", 
+            attribs->hostPointer,
+            attribs->devicePointer,
+            memoryTypeToString(attribs->memoryType),
+            attribs->device,
+            attribs->isManaged,
+            attribs->allocationFlags
+            );
+};
+
+
+inline int zrand(int max) 
+{
+    return rand() % max;
+}
+
+
+//=================================================================================================
+// Functins to run tests
+//=================================================================================================
+//
+//Run through a couple simple cases to test lookups and hostd pointer arithmetic:
+void simpleTests() 
+{
+    char *A_d;
+    char *A_Pinned_h;
+    char *A_OSAlloc_h;
+    hipError_t e;
+
+    HIPCHECK ( hipMalloc(&A_d, Nbytes) );
+    HIPCHECK ( hipMallocHost(&A_Pinned_h, Nbytes) );
+    A_OSAlloc_h = (char*)malloc(Nbytes);
+
+
+    hipPointerAttribute_t attribs;
+    hipPointerAttribute_t attribs2;
+
+    // Device memory
+    printf ("\nDevice memory (hipMalloc)\n");
+    HIPCHECK( hipPointerGetAttributes(&attribs, A_d));
+    printf("getAttr:%-20s", "A_d"); printAttribs(&attribs);
+
+    // Check pointer arithmetic cases:
+    resetAttribs(&attribs2);
+    HIPCHECK( hipPointerGetAttributes(&attribs2, A_d+100));
+    printf("getAttr:%-20s", "A_d+100"); printAttribs(&attribs2);
+    HIPASSERT(attribs == attribs2);
+
+    // Corner case at end of array:
+    resetAttribs(&attribs2);
+    HIPCHECK( hipPointerGetAttributes(&attribs2, A_d+Nbytes-1));
+    printf("getAttr:%-20s", "A_d+NBytes-1"); printAttribs(&attribs2);
+    HIPASSERT(attribs == attribs2);
+
+    // Pointer just beyond array - must be invalid or at least a different pointer
+    resetAttribs(&attribs2);
+    e = hipPointerGetAttributes(&attribs2, A_d+Nbytes+1);
+    printf("getAttr:%-20s err=%d (%s), neg-test expected\n", "A_d+NBytes", e, hipGetErrorString(e)); 
+    if (e != hipErrorInvalidValue) {
+        // We might have strayed into another pointer area.
+        printf("getAttr:%-20s", "A_d+NBytes"); printAttribs(&attribs2);
+        HIPASSERT(attribs.devicePointer != attribs2.devicePointer);
+    }
+
+
+    resetAttribs(&attribs2);
+    e = hipPointerGetAttributes(&attribs2, A_d+Nbytes);
+    if (e != hipErrorInvalidValue) {
+        printf("%-20s", "A_d+Nbytes"); printAttribs(&attribs2);
+        HIPASSERT(attribs.devicePointer != attribs2.devicePointer);
+    }
+
+    hipFree(A_d);
+    e = hipPointerGetAttributes(&attribs, A_d);
+    HIPASSERT(e == hipErrorInvalidValue); // Just freed the pointer, this should return an error.
+
+
+    // Device-visible host memory
+    printf ("\nDevice-visible host memory (hipMallocHost)\n");
+    HIPCHECK( hipPointerGetAttributes(&attribs, A_Pinned_h));
+    printf("getAttr:%-20s", "A_pinned_h"); printAttribs(&attribs);
+
+    resetAttribs(&attribs2);
+    HIPCHECK( hipPointerGetAttributes(&attribs2, A_Pinned_h+Nbytes/2));
+    printf("getAttr:%-20s", "A_pinned_h+NBytes/2"); printAttribs(&attribs2);
+    HIPASSERT(attribs == attribs2);
+
+
+    hipFreeHost(A_Pinned_h);
+    e = hipPointerGetAttributes(&attribs, A_Pinned_h);
+    HIPASSERT(e == hipErrorInvalidValue); // Just freed the pointer, this should return an error.
+    printf("getAttr:%-20s err=%d (%s), neg-test expected\n", "A_d+NBytes", e, hipGetErrorString(e)); 
+
+
+    // OS memory
+    printf ("\nOS-allocated memory (malloc)\n");
+    e = hipPointerGetAttributes(&attribs, A_OSAlloc_h);
+    printf("getAttr:%-20s err=%d (%s), neg-test expected\n", "A_OSAlloc_h", e, hipGetErrorString(e)); 
+    HIPASSERT(e == hipErrorInvalidValue); // OS-allocated pointers should return hipErrorInvalidValue.
+}
+
+
+
+
+struct SuperPointerAttribute {
+    void *                  _pointer;
+    size_t                  _sizeBytes;
+    hipPointerAttribute_t   _attrib;
+};
+
+
+
+void checkPointer(SuperPointerAttribute &ref, int major, int minor, void *pointer)
+{
+    hipPointerAttribute_t attribs;
+    resetAttribs(&attribs);
+
+    HIPCHECK(hipPointerGetAttributes(&attribs, pointer));
+    if (attribs != ref._attrib) {
+        printf("Test %d.%d", major, minor);
+        printf("  ref    ::  "); printAttribs(&ref._attrib);
+        printf("  getattr::  "); printAttribs(&attribs);
+        
+        HIPASSERT(attribs == ref._attrib);
+    } else {
+        if (p_verbose & 0x1) {
+            printf("#%4d.%d GOOD:%p getattr ::  ",major, minor, pointer); printAttribs(&attribs);
+        }
+    }
+}
+
+
+void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize)
+{
+    printf ("===========================================================================\n");
+    printf ("clusterAllocs numAllocs=%d size=%lu..%lu\n", numAllocs, minSize, maxSize);
+    printf ("===========================================================================\n");
+    std::vector <SuperPointerAttribute> reference(numAllocs);
+
+    HIPASSERT(minSize > 0);
+    HIPASSERT(maxSize >= minSize);
+
+    int numDevices;
+    HIPCHECK(hipGetDeviceCount(&numDevices));
+
+    //---
+    //Populate with device and host allocations.
+    for (int i=0; i<numAllocs; i++) {
+        bool isDevice = rand() & 0x1;
+        reference[i]._sizeBytes = zrand(maxSize-minSize) + minSize;
+
+        reference[i]._attrib.device = zrand(numDevices);
+        HIPCHECK(hipSetDevice(reference[i]._attrib.device));
+        reference[i]._attrib.isManaged = 0;
+
+        void * ptr;
+        if (isDevice) {
+            HIPCHECK(hipMalloc(&ptr, reference[i]._sizeBytes));
+            reference[i]._attrib.memoryType    = hipMemoryTypeDevice;
+            reference[i]._attrib.devicePointer = ptr;
+            reference[i]._attrib.hostPointer   = NULL;
+            reference[i]._attrib.allocationFlags = 0; // TODO-randomize these.
+        } else {
+            HIPCHECK(hipMallocHost(&ptr, reference[i]._sizeBytes));
+            reference[i]._attrib.memoryType    = hipMemoryTypeHost;
+            reference[i]._attrib.devicePointer = ptr;
+            reference[i]._attrib.hostPointer   = ptr;
+            reference[i]._attrib.allocationFlags = 1; // TODO-randomize these.
+        }
+        reference[i]._pointer = ptr;
+    }
+
+#ifdef __HIP_PLATFORM_HCC__
+    if (p_verbose & 0x2) {
+        hc::AM_print_tracker();
+    }
+#endif
+
+
+    // Now look up each pointer we inserted and verify we can find it:
+    for (int i=0; i<numAllocs; i++) {
+        SuperPointerAttribute &ref = reference[i];
+        checkPointer(ref, i, 0, ref._pointer);
+        checkPointer(ref, i, 1, (char *)ref._pointer + ref._sizeBytes/2);
+        if (ref._sizeBytes > 1) {
+            checkPointer(ref, i, 2, (char *)ref._pointer + ref._sizeBytes-1);
+        }
+
+    }
+}
+
+
+void testMultiThreaded()
+{
+    std::thread t1(clusterAllocs, 1000, 101, 1000);
+    std::thread t2(clusterAllocs, 1000,  11,  100);
+    std::thread t3(clusterAllocs, 1000,   5,  10);
+    std::thread t4(clusterAllocs, 1000,   1,   4);
+
+    t1.join();
+    t2.join();
+    t3.join();
+    t4.join();
+}
+
+
+int main(int argc, char *argv[])
+{
+
+    N= 1000000;
+    HipTest::parseStandardArguments(argc, argv, true);
+
+    HIPCHECK(hipSetDevice(p_gpuDevice));
+
+    Nbytes = N*sizeof(char);
+
+    printf ("N=%zu (%6.2f MB) device=%d\n", N, Nbytes/(1024.0*1024.0), p_gpuDevice);
+
+
+    if (p_tests & 0x1) {
+        simpleTests();
+    }
+
+    if (p_tests & 0x2) {
+        srand(0x100);
+        clusterAllocs(100, 1024*1, 1024*1024);
+    }
+
+    if (p_tests & 0x4) {
+        srand(0x200);
+        clusterAllocs(1000, 1, 10); //  Many tiny allocations;
+    }
+
+    if (p_tests & 0x8) {
+        testMultiThreaded();
+    }
+
+    printf ("\n");
+    passed();
+}
@@ -28,6 +28,8 @@ int iterations = 1;
 unsigned blocksPerCU = 6; // to hide latency
 unsigned threadsPerBlock = 256; 
 int p_gpuDevice = 0;
+unsigned p_verbose  = 0;
+int p_tests = -1; /*which tests to run. Interpretation is left to each test.  default:all*/



@@ -114,8 +116,16 @@ int parseStandardArguments(int argc, char *argv[], bool failOnUndefinedArg)
               failed("Bad gpuDevice argument"); 
            }

-        } 
-        else {
+        } else if (!strcmp(arg, "--verbose") || (!strcmp(arg, "-v"))) {
+            if (++i >= argc || !HipTest::parseUInt(argv[i], &p_verbose)) {
+               failed("Bad verbose argument"); 
+            }
+        } else if (!strcmp(arg, "--tests") || (!strcmp(arg, "-t"))) {
+            if (++i >= argc || !HipTest::parseInt(argv[i], &p_tests)) {
+               failed("Bad tests argument"); 
+            }
+
+        } else {
            if (failOnUndefinedArg) {
                failed("Bad argument '%s'", arg);
            } else {
@@ -53,6 +53,8 @@ extern int iterations;
 extern unsigned blocksPerCU;
 extern unsigned threadsPerBlock;
 extern int p_gpuDevice;
+extern unsigned p_verbose;
+extern int p_tests;

 namespace HipTest {