diff --git a/bin/hipcc b/bin/hipcc index d001c6febe..7537750ff6 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -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"); diff --git a/bin/hipify b/bin/hipify index d143bdff37..e3b6c64c88 100755 --- a/bin/hipify +++ b/bin/hipify @@ -277,6 +277,8 @@ while (@ARGV) { $ft{'mem'} += s/\bcudaMemcpyKind\b/hipMemcpyKind/g; + $ft{'mem'} += s/\bcudaPointerAttributes\b/hipPointerAttribute_t/g; + #-------- # Memory management: diff --git a/include/hcc_detail/AM.h b/include/hcc_detail/AM.h new file mode 100644 index 0000000000..1cfcf2dfb2 --- /dev/null +++ b/include/hcc_detail/AM.h @@ -0,0 +1,92 @@ +#pragma once + +#include + +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 + diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index 225b065654..a0c676987b 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/include/hcc_detail/hip_runtime_api.h @@ -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 diff --git a/include/hip_runtime_api.h b/include/hip_runtime_api.h index 882103a1f4..41ad338d6d 100644 --- a/include/hip_runtime_api.h +++ b/include/hip_runtime_api.h @@ -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 diff --git a/src/hc_AM.cpp b/src/hc_AM.cpp new file mode 100644 index 0000000000..87e29e4bcc --- /dev/null +++ b/src/hc_AM.cpp @@ -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 +#include +//#include + +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 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 (acc.get_default_view().get_hsa_agent()); + hsa_region_t *alloc_region; + if (flags & amHostPinned) { + alloc_region = static_cast(acc.get_hsa_am_system_region()); + } else { + alloc_region = static_cast(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. diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index f9f3b1b1f5..e9ee4c41dc 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -31,6 +31,8 @@ THE SOFTWARE. #include #include #include +#include + #include #include @@ -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; imemoryType = 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; idevice = 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 @@ -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 diff --git a/tests/src/CMakeLists.txt b/tests/src/CMakeLists.txt index 0ec287b334..bf05fc8407 100644 --- a/tests/src/CMakeLists.txt +++ b/tests/src/CMakeLists.txt @@ -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 " " ) diff --git a/tests/src/hipPointerAttrib.cpp b/tests/src/hipPointerAttrib.cpp new file mode 100644 index 0000000000..9d147d8183 --- /dev/null +++ b/tests/src/hipPointerAttrib.cpp @@ -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 reference(numAllocs); + + HIPASSERT(minSize > 0); + HIPASSERT(maxSize >= minSize); + + int numDevices; + HIPCHECK(hipGetDeviceCount(&numDevices)); + + //--- + //Populate with device and host allocations. + for (int i=0; i 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(); +} diff --git a/tests/src/test_common.cpp b/tests/src/test_common.cpp index d7a108a11b..02deb51c85 100644 --- a/tests/src/test_common.cpp +++ b/tests/src/test_common.cpp @@ -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 { diff --git a/tests/src/test_common.h b/tests/src/test_common.h index fee052c1ad..57d2ebc831 100644 --- a/tests/src/test_common.h +++ b/tests/src/test_common.h @@ -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 {