Merge branch 'memtracker' of https://github.com/AMDComputeLibraries/HIP-privatestaging into memtracker
This commit is contained in:
Ben Sander
@@ -19,8 +19,10 @@ MESSAGE ("HIP_PATH=" ${HIP_PATH})
|
||||
|
||||
if (${HIP_PLATFORM} STREQUAL "hcc")
|
||||
MESSAGE ("HIP_PLATFORM=hcc")
|
||||
set (HC_PATH ${HIP_PATH}/hc)
|
||||
set (HSA_PATH /opt/hsa)
|
||||
set (HSA_PATH $ENV{HSA_PATH})
|
||||
if (NOT DEFINED HSA_PATH)
|
||||
set (HSA_PATH /opt/hsa)
|
||||
endif()
|
||||
|
||||
#---
|
||||
# Add HSA library:
|
||||
@@ -30,7 +32,7 @@ if (${HIP_PLATFORM} STREQUAL "hcc")
|
||||
#These includes are used for all files.
|
||||
#Include HIP and HC since the tests need both of these:
|
||||
#Note below HSA path is surgically included only where necessary.
|
||||
include_directories(${HIP_PATH}/include ${HC_PATH}/include)
|
||||
include_directories(${HIP_PATH}/include)
|
||||
|
||||
# hip_hcc.o:
|
||||
add_library(hip_hcc OBJECT ${HIP_PATH}/src/hip_hcc.cpp)
|
||||
@@ -105,6 +107,7 @@ make_hip_executable (hip_brev hip_brev.cpp)
|
||||
make_hip_executable (hip_ffs hip_ffs.cpp)
|
||||
make_hip_executable (hipGetDeviceAttribute hipGetDeviceAttribute.cpp)
|
||||
make_hip_executable (hipMemcpy hipMemcpy.cpp)
|
||||
make_hip_executable (hipMemcpyAsync hipMemcpyAsync.cpp)
|
||||
make_hip_executable (hipMemset hipMemset.cpp)
|
||||
make_hip_executable (hipEventRecord hipEventRecord.cpp)
|
||||
make_hip_executable (hipLanguageExtensions hipLanguageExtensions.cpp)
|
||||
@@ -114,6 +117,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,8 +132,10 @@ 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 " " )
|
||||
make_test(hipMemcpyAsync " " )
|
||||
|
||||
make_test(hipHcc " " )
|
||||
|
||||
|
||||
@@ -23,24 +23,28 @@ THE SOFTWARE.
|
||||
#include "test_common.h"
|
||||
|
||||
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
void printSep()
|
||||
{
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
printf ("======================================================================================\n");
|
||||
}
|
||||
|
||||
//---
|
||||
// Test simple H2D copies and back.
|
||||
// Designed to stress a small number of simple smoke tests
|
||||
void simpleTest1()
|
||||
{
|
||||
printf ("test: %s\n", __func__);
|
||||
size_t Nbytes = N*sizeof(int);
|
||||
|
||||
printf ("N=%zu \n", N);
|
||||
printf ("N=%zu Nbytes=%6.2fMB\n", N, Nbytes/1024.0/1024.0);
|
||||
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
HipTest::initArrays (&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N);
|
||||
|
||||
HipTest::initArrays (&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
printf ("A_d=%p B_d=%p C_d=%p A_h=%p B_h=%p C_h=%p\n", A_d, B_d, C_d, A_h, B_d, C_h);
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
|
||||
HIPCHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
HIPCHECK ( hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
|
||||
|
||||
@@ -50,8 +54,212 @@ int main(int argc, char *argv[])
|
||||
|
||||
HIPCHECK (hipDeviceSynchronize());
|
||||
|
||||
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HipTest::freeArrays (A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK (hipDeviceReset());
|
||||
|
||||
printf (" %s success\n", __func__);
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
// Test many different kinds of memory copies.
|
||||
// THe subroutine allocates memory , copies to device, runs a vector add kernel, copies back, and checks the result.
|
||||
//
|
||||
// IN: numElements controls the number of elements used for allocations.
|
||||
// IN: usePinnedHost : If true, allocate host with hipMallocHost and is pinned ; else allocate host memory with malloc.
|
||||
// IN: useHostToHost : If true, add an extra host-to-host copy.
|
||||
// IN: useDeviceToDevice : If true, add an extra deviceto-device copy after result is produced.
|
||||
// IN: useMemkindDefault : If true, use memkinddefault (runtime figures out direction). if false, use explicit memcpy direction.
|
||||
//
|
||||
template <typename T>
|
||||
void memcpytest2(size_t numElements, bool usePinnedHost, bool useHostToHost, bool useDeviceToDevice, bool useMemkindDefault)
|
||||
{
|
||||
size_t sizeElements = numElements * sizeof(T);
|
||||
printf ("test: %s<%s> size=%lu (%6.2fMB) usePinnedHost:%d, useHostToHost:%d, useDeviceToDevice:%d, useMemkindDefault:%d\n",
|
||||
__func__,
|
||||
typeid(T).name(),
|
||||
sizeElements, sizeElements/1024.0/1024.0,
|
||||
usePinnedHost, useHostToHost, useDeviceToDevice, useMemkindDefault);
|
||||
|
||||
|
||||
T *A_d, *B_d, *C_d;
|
||||
T *A_h, *B_h, *C_h;
|
||||
|
||||
|
||||
HipTest::initArrays (&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, numElements, usePinnedHost);
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
|
||||
|
||||
T *A_hh = NULL;
|
||||
T *B_hh = NULL;
|
||||
T *C_dd = NULL;
|
||||
|
||||
|
||||
|
||||
if (useHostToHost) {
|
||||
if (usePinnedHost) {
|
||||
HIPCHECK ( hipMallocHost(&A_hh, sizeElements) );
|
||||
HIPCHECK ( hipMallocHost(&B_hh, sizeElements) );
|
||||
} else {
|
||||
A_hh = (T*)malloc(sizeElements);
|
||||
B_hh = (T*)malloc(sizeElements);
|
||||
}
|
||||
|
||||
|
||||
// Do some extra host-to-host copies here to mix things up:
|
||||
HIPCHECK ( hipMemcpy(A_hh, A_h, sizeElements, useMemkindDefault? hipMemcpyDefault : hipMemcpyHostToHost));
|
||||
HIPCHECK ( hipMemcpy(B_hh, B_h, sizeElements, useMemkindDefault? hipMemcpyDefault : hipMemcpyHostToHost));
|
||||
|
||||
|
||||
HIPCHECK ( hipMemcpy(A_d, A_hh, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice));
|
||||
HIPCHECK ( hipMemcpy(B_d, B_hh, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice));
|
||||
} else {
|
||||
HIPCHECK ( hipMemcpy(A_d, A_h, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice));
|
||||
HIPCHECK ( hipMemcpy(B_d, B_h, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice));
|
||||
}
|
||||
|
||||
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d, B_d, C_d, numElements);
|
||||
|
||||
if (useDeviceToDevice) {
|
||||
HIPCHECK ( hipMalloc(&C_dd, sizeElements) );
|
||||
|
||||
// Do an extra device-to-device copies here to mix things up:
|
||||
HIPCHECK ( hipMemcpy(C_dd, C_d, sizeElements, useMemkindDefault? hipMemcpyDefault : hipMemcpyDeviceToDevice));
|
||||
|
||||
//Destroy the original C_d:
|
||||
HIPCHECK ( hipMemset(C_d, 0x5A, sizeElements));
|
||||
|
||||
HIPCHECK ( hipMemcpy(C_h, C_dd, sizeElements, useMemkindDefault? hipMemcpyDefault:hipMemcpyDeviceToHost));
|
||||
} else {
|
||||
HIPCHECK ( hipMemcpy(C_h, C_d, sizeElements, useMemkindDefault? hipMemcpyDefault:hipMemcpyDeviceToHost));
|
||||
}
|
||||
|
||||
HIPCHECK ( hipDeviceSynchronize() );
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, numElements);
|
||||
|
||||
HipTest::freeArrays (A_d, B_d, C_d, A_h, B_h, C_h, usePinnedHost);
|
||||
|
||||
printf (" %s success\n", __func__);
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
//Try all the 16 possible combinations to memcpytest2 - usePinnedHost, useHostToHost, useDeviceToDevice, useMemkindDefault
|
||||
template<typename T>
|
||||
void memcpytest2_loop(size_t numElements)
|
||||
{
|
||||
printSep();
|
||||
|
||||
for (int usePinnedHost =0; usePinnedHost<=1; usePinnedHost++) {
|
||||
for (int useHostToHost =0; useHostToHost<=1; useHostToHost++) { // TODO
|
||||
for (int useDeviceToDevice =0; useDeviceToDevice<=1; useDeviceToDevice++) {
|
||||
for (int useMemkindDefault =0; useMemkindDefault<=1; useMemkindDefault++) {
|
||||
memcpytest2<T>(numElements, usePinnedHost, useHostToHost, useDeviceToDevice, useMemkindDefault);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
//Try many different sizes to memory copy.
|
||||
template<typename T>
|
||||
void memcpytest2_sizes(size_t maxElem=0, size_t offset=0)
|
||||
{
|
||||
printSep();
|
||||
printf ("test: %s<%s>\n", __func__, typeid(T).name());
|
||||
|
||||
int deviceId;
|
||||
HIPCHECK(hipGetDevice(&deviceId));
|
||||
|
||||
size_t free, total;
|
||||
HIPCHECK(hipMemGetInfo(&free, &total));
|
||||
|
||||
if (maxElem == 0) {
|
||||
maxElem = free/sizeof(T)/5;
|
||||
}
|
||||
|
||||
printf (" device#%d: hipMemGetInfo: free=%zu (%4.2fMB) total=%zu (%4.2fMB) maxSize=%6.1fMB offset=%lu\n",
|
||||
deviceId, free, (float)(free/1024.0/1024.0), total, (float)(total/1024.0/1024.0), maxElem*sizeof(T)/1024.0/1024.0, offset);
|
||||
|
||||
for (size_t elem=64; elem+offset<=maxElem; elem*=2) {
|
||||
HIPCHECK ( hipDeviceReset() );
|
||||
memcpytest2<T>(elem+offset, 0, 1, 1, 0); // unpinned host
|
||||
HIPCHECK ( hipDeviceReset() );
|
||||
memcpytest2<T>(elem+offset, 1, 1, 1, 0); // pinned host
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
//Create multiple threads to stress multi-thread locking behavior in the allocation/deallocation/tracking logic:
|
||||
template<typename T>
|
||||
void multiThread_1(bool serialize, bool usePinnedHost)
|
||||
{
|
||||
printSep();
|
||||
printf ("test: %s<%s> serialize=%d usePinnedHost=%d\n", __func__, typeid(T).name(), serialize, usePinnedHost);
|
||||
std::thread t1 (memcpytest2<T>,N, usePinnedHost,0,0,0);
|
||||
if (serialize) {
|
||||
t1.join();
|
||||
}
|
||||
|
||||
|
||||
std::thread t2 (memcpytest2<T>,N, usePinnedHost,0,0,0);
|
||||
if (serialize) {
|
||||
t2.join();
|
||||
}
|
||||
|
||||
if (!serialize) {
|
||||
t1.join();
|
||||
t2.join();
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
|
||||
printf ("info: set device to %d\n", p_gpuDevice);
|
||||
HIPCHECK(hipSetDevice(p_gpuDevice));
|
||||
|
||||
|
||||
if (p_tests & 0x1) {
|
||||
HIPCHECK ( hipDeviceReset() );
|
||||
simpleTest1();
|
||||
}
|
||||
|
||||
if (p_tests & 0x2) {
|
||||
HIPCHECK ( hipDeviceReset() );
|
||||
memcpytest2_loop<float>(N);
|
||||
memcpytest2_loop<double>(N);
|
||||
memcpytest2_loop<char>(N);
|
||||
memcpytest2_loop<int>(N);
|
||||
}
|
||||
|
||||
if (p_tests & 0x4) {
|
||||
HIPCHECK ( hipDeviceReset() );
|
||||
printSep();
|
||||
memcpytest2_sizes<float>(0,0);
|
||||
printSep();
|
||||
memcpytest2_sizes<float>(0,64);
|
||||
printSep();
|
||||
memcpytest2_sizes<float>(1024*1024, 13);
|
||||
printSep();
|
||||
memcpytest2_sizes<float>(1024*1024, 50);
|
||||
}
|
||||
|
||||
if (p_tests & 0x8) {
|
||||
HIPCHECK ( hipDeviceReset() );
|
||||
printSep();
|
||||
multiThread_1<float>(true, true);
|
||||
multiThread_1<float>(false, true);
|
||||
multiThread_1<float>(false, false); // TODO
|
||||
}
|
||||
|
||||
passed();
|
||||
|
||||
}
|
||||
|
||||
@@ -0,0 +1,349 @@
|
||||
// Test under-development. Calls async mem-copy API, experiment with functionality.
|
||||
|
||||
#include "hip_runtime.h"
|
||||
#include "test_common.h"
|
||||
|
||||
unsigned p_streams = 2;
|
||||
|
||||
|
||||
void simpleNegTest()
|
||||
{
|
||||
printf ("testing: %s\n",__func__);
|
||||
hipError_t e;
|
||||
float *A_malloc, *A_pinned, *A_d;
|
||||
|
||||
size_t Nbytes = N*sizeof(float);
|
||||
A_malloc = (float*)malloc(Nbytes);
|
||||
HIPCHECK(hipMallocHost(&A_pinned, Nbytes));
|
||||
HIPCHECK(hipMalloc(&A_d, Nbytes));
|
||||
|
||||
|
||||
// Can't use default with async copy
|
||||
e = hipMemcpyAsync(A_pinned, A_d, Nbytes, hipMemcpyDefault, NULL);
|
||||
HIPASSERT (e==hipErrorInvalidMemcpyDirection); // TODO
|
||||
HIPASSERT (e!= hipSuccess);
|
||||
|
||||
|
||||
// Not sure what happens here, the memory must be pinned.
|
||||
e = hipMemcpyAsync(A_malloc, A_d, Nbytes, hipMemcpyHostToDevice, NULL);
|
||||
|
||||
printf (" async memcpy of A_malloc to A_d. Result=%d\n", e);
|
||||
//HIPASSERT (e==hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
class Pinned;
|
||||
class Unpinned;
|
||||
|
||||
template <typename T> struct HostTraits;
|
||||
|
||||
template<>
|
||||
struct HostTraits<Pinned>
|
||||
{
|
||||
static const char *Name() { return "Pinned"; } ;
|
||||
|
||||
static void *Alloc(size_t sizeBytes) {
|
||||
void *p;
|
||||
HIPCHECK(hipMallocHost(&p, sizeBytes));
|
||||
return p;
|
||||
};
|
||||
};
|
||||
|
||||
|
||||
template<typename T>
|
||||
__global__ void
|
||||
addK (hipLaunchParm lp, T *A, T K, size_t numElements)
|
||||
{
|
||||
size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
|
||||
size_t stride = hipBlockDim_x * hipGridDim_x ;
|
||||
|
||||
for (size_t i=offset; i<numElements; i+=stride) {
|
||||
A[i] = A[i] + K;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
//---
|
||||
//Tests propert dependency resolution between H2D and D2H commands in same stream:
|
||||
//IN: numInflight : number of copies inflight at any time:
|
||||
//IN: numPongs = number of iterations to run (iteration)
|
||||
template<typename T, class AllocType>
|
||||
void test_pingpong(hipStream_t stream, size_t numElements, int numInflight, int numPongs, bool doHostSide)
|
||||
{
|
||||
HIPASSERT(numElements % numInflight == 0); // Must be evenly divisible.
|
||||
size_t Nbytes = numElements*sizeof(T);
|
||||
size_t eachCopyElements = numElements / numInflight;
|
||||
size_t eachCopyBytes = eachCopyElements * sizeof(T);
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
|
||||
|
||||
printf ("-----------------------------------------------------------------------------------------------\n");
|
||||
printf ("testing: %s<%s> Nbytes=%zu (%6.1f MB) numPongs=%d numInflight=%d eachCopyElements=%zu eachCopyBytes=%zu\n",
|
||||
__func__, HostTraits<AllocType>::Name(), Nbytes, (double)(Nbytes)/1024.0/1024.0, numPongs, numInflight, eachCopyElements, eachCopyBytes);
|
||||
|
||||
T *A_h;
|
||||
T *A_d;
|
||||
|
||||
A_h = (T*)(HostTraits<AllocType>::Alloc(Nbytes));
|
||||
HIPCHECK(hipMalloc(&A_d, Nbytes));
|
||||
|
||||
// Initialize the host array:
|
||||
const T initValue = 13;
|
||||
const T deviceConst = 2;
|
||||
const T hostConst = 10000;
|
||||
for (size_t i=0; i<numElements; i++) {
|
||||
A_h[i] = initValue + i;
|
||||
}
|
||||
|
||||
|
||||
for (int k=0; k<numPongs; k++ ) {
|
||||
for (int i=0; i<numInflight; i++) {
|
||||
HIPCHECK(hipMemcpyAsync(&A_d[i*eachCopyElements], &A_h[i*eachCopyElements], eachCopyBytes, hipMemcpyHostToDevice, stream));
|
||||
}
|
||||
|
||||
hipLaunchKernel(addK<T>, dim3(blocks), dim3(threadsPerBlock), 0, stream, A_d, 2, numElements);
|
||||
|
||||
for (int i=0; i<numInflight; i++ ) {
|
||||
HIPCHECK(hipMemcpyAsync(&A_h[i*eachCopyElements], &A_d[i*eachCopyElements], eachCopyBytes, hipMemcpyDeviceToHost, stream));
|
||||
}
|
||||
|
||||
if (doHostSide) {
|
||||
assert(0);
|
||||
#if 0
|
||||
hipEvent_t e;
|
||||
HIPCHECK(hipEventCreate(&e));
|
||||
#endif
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
for (size_t i=0; i<numElements; i++) {
|
||||
A_h[i] += hostConst;
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
|
||||
// Verify we copied back all the data correctly:
|
||||
for (size_t i=0; i<numElements; i++) {
|
||||
T gold = initValue + i;
|
||||
// Perform calcs in same order as test above to replicate FP order-of-operations:
|
||||
for (int k=0; k<numPongs; k++) {
|
||||
gold += deviceConst;
|
||||
if (doHostSide) {
|
||||
gold += hostConst;
|
||||
}
|
||||
}
|
||||
|
||||
if (gold != A_h[i]) {
|
||||
std::cout << i << ": gold=" << gold << " out=" << A_h[i] << std::endl;
|
||||
HIPASSERT(gold == A_h[i]);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
HIPCHECK(hipFreeHost(A_h));
|
||||
HIPCHECK(hipFree(A_d));
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
//Send many async copies to the same stream.
|
||||
//This requires runtime to keep track of many outstanding commands, and in the case of HCC requires growing/tracking the signal pool:
|
||||
template<typename T>
|
||||
void test_manyInflightCopies(hipStream_t stream, int numElements, int numCopies, bool syncBetweenCopies)
|
||||
{
|
||||
size_t Nbytes = numElements*sizeof(T);
|
||||
size_t eachCopyElements = numElements / numCopies;
|
||||
size_t eachCopyBytes = eachCopyElements * sizeof(T);
|
||||
|
||||
printf ("-----------------------------------------------------------------------------------------------\n");
|
||||
printf ("testing: %s Nbytes=%zu (%6.1f MB) numCopies=%d eachCopyElements=%zu eachCopyBytes=%zu\n",
|
||||
__func__, Nbytes, (double)(Nbytes)/1024.0/1024.0, numCopies, eachCopyElements, eachCopyBytes);
|
||||
|
||||
T *A_d;
|
||||
T *A_h1, *A_h2;
|
||||
|
||||
HIPCHECK(hipMallocHost(&A_h1, Nbytes));
|
||||
HIPCHECK(hipMallocHost(&A_h2, Nbytes));
|
||||
HIPCHECK(hipMalloc(&A_d, Nbytes));
|
||||
|
||||
for (int i=0; i<numElements; i++) {
|
||||
A_h1[i] = 3.14f + static_cast<T> (i);
|
||||
}
|
||||
|
||||
|
||||
//stream=0; // fixme TODO
|
||||
|
||||
|
||||
for (int i=0; i<numCopies; i++)
|
||||
{
|
||||
HIPCHECK(hipMemcpyAsync(&A_d[i*eachCopyElements], &A_h1[i*eachCopyElements], eachCopyBytes, hipMemcpyHostToDevice, stream));
|
||||
}
|
||||
|
||||
if (syncBetweenCopies) {
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
|
||||
for (int i=0; i<numCopies; i++)
|
||||
{
|
||||
HIPCHECK(hipMemcpyAsync(&A_h2[i*eachCopyElements], &A_d[i*eachCopyElements], eachCopyBytes, hipMemcpyDeviceToHost, stream));
|
||||
}
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
|
||||
// Verify we copied back all the data correctly:
|
||||
for (int i=0; i<numElements; i++) {
|
||||
HIPASSERT(A_h1[i] == A_h2[i]);
|
||||
}
|
||||
|
||||
|
||||
HIPCHECK(hipFreeHost(A_h1));
|
||||
HIPCHECK(hipFreeHost(A_h2));
|
||||
HIPCHECK(hipFree(A_d));
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
//Classic example showing how to overlap data transfer with compute.
|
||||
//We divide the work into "chunks" and create a stream for each chunk.
|
||||
//Each chunk then runs a H2D copy, followed by kernel execution, followed by D2H copyback.
|
||||
//Work in separate streams is independent which enables concurrency.
|
||||
|
||||
// IN: nStreams : number of streams to use for the test
|
||||
// IN :useNullStream - use NULL stream. Synchronizes everything.
|
||||
// IN: useSyncMemcpyH2D - use sync memcpy (no overlap) for H2D
|
||||
// IN: useSyncMemcpyD2H - use sync memcpy (no overlap) for D2H
|
||||
void test_chunkedAsyncExample(int nStreams, bool useNullStream, bool useSyncMemcpyH2D, bool useSyncMemcpyD2H)
|
||||
{
|
||||
|
||||
size_t Nbytes = N*sizeof(int);
|
||||
printf ("testing: %s(useNullStream=%d, useSyncMemcpyH2D=%d, useSyncMemcpyD2H=%d) ",__func__, useNullStream, useSyncMemcpyH2D, useSyncMemcpyD2H);
|
||||
printf ("Nbytes=%zu (%6.1f MB)\n", Nbytes, (double)(Nbytes)/1024.0/1024.0);
|
||||
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
HipTest::initArrays (&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, true);
|
||||
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
|
||||
hipStream_t *stream = (hipStream_t*)malloc(sizeof(hipStream_t) * nStreams);
|
||||
if (useNullStream) {
|
||||
nStreams = 1;
|
||||
stream[0] = NULL;
|
||||
} else {
|
||||
for (int i = 0; i < nStreams; ++i) {
|
||||
HIPCHECK (hipStreamCreate(&stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
size_t workLeft = N;
|
||||
size_t workPerStream = N / nStreams;
|
||||
for (int i = 0; i < nStreams; ++i) {
|
||||
size_t work = (workLeft < workPerStream) ? workLeft : workPerStream;
|
||||
size_t workBytes = work * sizeof(int);
|
||||
|
||||
size_t offset = i*workPerStream;
|
||||
|
||||
if (useSyncMemcpyH2D) {
|
||||
HIPCHECK ( hipMemcpy(&A_d[offset], &A_h[offset], workBytes, hipMemcpyHostToDevice));
|
||||
HIPCHECK ( hipMemcpy(&B_d[offset], &B_h[offset], workBytes, hipMemcpyHostToDevice));
|
||||
} else {
|
||||
HIPCHECK ( hipMemcpyAsync(&A_d[offset], &A_h[offset], workBytes, hipMemcpyHostToDevice, stream[i]));
|
||||
HIPCHECK ( hipMemcpyAsync(&B_d[offset], &B_h[offset], workBytes, hipMemcpyHostToDevice, stream[i]));
|
||||
};
|
||||
|
||||
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i], &A_d[offset], &B_d[offset], &C_d[offset], work);
|
||||
|
||||
if (useSyncMemcpyD2H) {
|
||||
HIPCHECK ( hipMemcpy(&C_h[offset], &C_d[offset], workBytes, hipMemcpyDeviceToHost));
|
||||
} else {
|
||||
HIPCHECK ( hipMemcpyAsync(&C_h[offset], &C_d[offset], workBytes, hipMemcpyDeviceToHost, stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
HIPCHECK (hipDeviceSynchronize());
|
||||
|
||||
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HipTest::freeArrays (A_d, B_d, C_d, A_h, B_h, C_h, true);
|
||||
|
||||
free(stream);
|
||||
};
|
||||
|
||||
|
||||
//---
|
||||
//Parse arguments specific to this test.
|
||||
void parseMyArguments(int argc, char *argv[])
|
||||
{
|
||||
int more_argc = HipTest::parseStandardArguments(argc, argv, false);
|
||||
|
||||
// parse args for this test:
|
||||
for (int i = 1; i < more_argc; i++) {
|
||||
const char *arg = argv[i];
|
||||
|
||||
if (!strcmp(arg, "--streams")) {
|
||||
if (++i >= argc || !HipTest::parseUInt(argv[i], &p_streams)) {
|
||||
failed("Bad streams argument");
|
||||
}
|
||||
} else {
|
||||
failed("Bad argument '%s'", arg);
|
||||
}
|
||||
};
|
||||
};
|
||||
|
||||
|
||||
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
parseMyArguments(argc, argv);
|
||||
|
||||
|
||||
printf ("info: set device to %d\n", p_gpuDevice);
|
||||
HIPCHECK(hipSetDevice(p_gpuDevice));
|
||||
|
||||
if (p_tests & 0x01) {
|
||||
simpleNegTest();
|
||||
}
|
||||
|
||||
if (p_tests & 0x02) {
|
||||
hipStream_t stream;
|
||||
HIPCHECK (hipStreamCreate(&stream));
|
||||
|
||||
test_manyInflightCopies<float>(stream, 1024, 16, true);
|
||||
test_manyInflightCopies<float>(stream, 1024, 4, true); // verify we re-use the same entries instead of growing pool.
|
||||
test_manyInflightCopies<float>(stream, 1024*8, 64, false);
|
||||
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
|
||||
if (p_tests & 0x04) {
|
||||
test_chunkedAsyncExample(p_streams, true, true, true); // Easy sync version
|
||||
test_chunkedAsyncExample(p_streams, false, true, true); // Easy sync version
|
||||
test_chunkedAsyncExample(p_streams, false, false, true); // Some async
|
||||
test_chunkedAsyncExample(p_streams, false, false, false); // All async
|
||||
}
|
||||
|
||||
if (p_tests & 0x08) {
|
||||
hipStream_t stream;
|
||||
HIPCHECK (hipStreamCreate(&stream));
|
||||
|
||||
test_pingpong<int, Pinned>(stream, 1024*1024*32, 1, 1, false);
|
||||
test_pingpong<int, Pinned>(stream, 1024*1024*32, 1, 10, false);
|
||||
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
|
||||
passed();
|
||||
|
||||
}
|
||||
@@ -0,0 +1,524 @@
|
||||
/*
|
||||
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"
|
||||
#include "hc_am.hpp"
|
||||
|
||||
#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(const 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 host pointer arithmetic:
|
||||
void testSimple()
|
||||
{
|
||||
printf ("\n");
|
||||
printf ("===========================================================================\n");
|
||||
printf ("Simple Tests\n");
|
||||
printf ("===========================================================================\n");
|
||||
|
||||
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);
|
||||
|
||||
size_t free, total;
|
||||
HIPCHECK(hipMemGetInfo(&free, &total));
|
||||
printf ("hipMemGetInfo: free=%zu (%4.2f) Nbytes=%lu total=%zu (%4.2f)\n", free, (float)(free/1024.0/1024.0), Nbytes, total, (float)(total/1024.0/1024.0));
|
||||
HIPASSERT(free + Nbytes <= total);
|
||||
|
||||
|
||||
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.
|
||||
}
|
||||
|
||||
//---
|
||||
//Reset the memory tracker (remove allocations from all known devices):
|
||||
//This frees any memory allocated through the runtime.
|
||||
//The routine will not release any
|
||||
void resetTracker ()
|
||||
{
|
||||
if (p_verbose & 0x1) {
|
||||
printf ("info: reset tracker for all devices in platform\n");
|
||||
}
|
||||
|
||||
int numDevices;
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
|
||||
// Clean up:
|
||||
for (int i=0; i<numDevices; i++) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipDeviceReset());
|
||||
};
|
||||
}
|
||||
|
||||
|
||||
// Store the hipPointer attrib and some extra info so can later compare the looked-up info against the reference expectation
|
||||
struct SuperPointerAttribute {
|
||||
void * _pointer;
|
||||
size_t _sizeBytes;
|
||||
hipPointerAttribute_t _attrib;
|
||||
};
|
||||
|
||||
|
||||
//---
|
||||
//Support function to check result against a reference:
|
||||
void checkPointer(SuperPointerAttribute &ref, int major, int minor, void *pointer)
|
||||
{
|
||||
hipPointerAttribute_t attribs;
|
||||
resetAttribs(&attribs);
|
||||
|
||||
hipError_t e = hipPointerGetAttributes(&attribs, pointer);
|
||||
if ((e != hipSuccess) || (attribs != ref._attrib)) {
|
||||
printf("Test %d.%d (err=%d)\n", major, minor, e);
|
||||
HIPCHECK(e);
|
||||
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);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
//Test that allocates memory across all 4 devices withing the specified size range (minSize...maxSize).
|
||||
//Then does lookups to make sure the info reported by the tracker matches expecations
|
||||
//Then deallocates it all.
|
||||
//
|
||||
//Multiple threads can call this funtion and in fact we do this in the testMultiThreaded_1 test.
|
||||
void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize)
|
||||
{
|
||||
printf (" clusterAllocs numAllocs=%d size=%lu..%lu\n", numAllocs, minSize, maxSize);
|
||||
std::vector <SuperPointerAttribute> reference(numAllocs);
|
||||
|
||||
HIPASSERT(minSize > 0);
|
||||
HIPASSERT(maxSize >= minSize);
|
||||
|
||||
int numDevices;
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
|
||||
//---
|
||||
//Populate with device and host allocations.
|
||||
size_t totalDeviceAllocated[numDevices];
|
||||
for (int i =0; i<numDevices; i++) {
|
||||
totalDeviceAllocated[i] = 0;
|
||||
}
|
||||
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) {
|
||||
totalDeviceAllocated[reference[i]._attrib.device] += reference[i]._sizeBytes;
|
||||
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 = 0; // TODO-randomize these.
|
||||
}
|
||||
reference[i]._pointer = ptr;
|
||||
}
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
if (p_verbose & 0x2) {
|
||||
printf ("Tracker after insertions:\n");
|
||||
hc::am_memtracker_print();
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
for (int i =0; i<numDevices; i++) {
|
||||
size_t free, total;
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipMemGetInfo(&free, &total));
|
||||
printf (" device#%d: hipMemGetInfo: free=%zu (%4.2fMB) clusterAllocTotalDevice=%lu (%4.2fMB) total=%zu (%4.2fMB)\n",
|
||||
i, free, (float)(free/1024.0/1024.0), totalDeviceAllocated[i], (float)(totalDeviceAllocated[i])/1024.0/1024.0, total, (float)(total/1024.0/1024.0));
|
||||
HIPASSERT(free + totalDeviceAllocated[i] <= total);
|
||||
}
|
||||
|
||||
|
||||
// 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);
|
||||
}
|
||||
|
||||
if (ref._attrib.memoryType == hipMemoryTypeDevice) {
|
||||
hipFree(ref._pointer);
|
||||
} else {
|
||||
hipFreeHost(ref._pointer);
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
if (p_verbose & 0x2) {
|
||||
printf ("Tracker after cleanup:\n");
|
||||
hc::am_memtracker_print();
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
// Multi-threaded test with many simul allocs.
|
||||
// IN : serialize will force the test to run in serial fashion.
|
||||
// Seems like this does not hit MT corner cases in the tracker very often - testMultiThreaded_2 below seems more effective.
|
||||
void testMultiThreaded_1(bool serialize=false)
|
||||
{
|
||||
printf ("\n===========================================================================\n");
|
||||
printf ("MultiThreaded_1\n");
|
||||
if (serialize) printf ("[SERIALIZE]\n");
|
||||
printf ("===========================================================================\n");
|
||||
std::thread t1(clusterAllocs, 1000, 101, 1000);
|
||||
if (serialize) t1.join();
|
||||
|
||||
std::thread t2(clusterAllocs, 1000, 11, 100);
|
||||
if (serialize) t2.join();
|
||||
|
||||
std::thread t3(clusterAllocs, 1000, 5, 10);
|
||||
if (serialize) t3.join();
|
||||
|
||||
std::thread t4(clusterAllocs, 1000, 1, 4);
|
||||
if (serialize) t4.join();
|
||||
|
||||
if (!serialize) {
|
||||
t1.join();
|
||||
t2.join();
|
||||
t3.join();
|
||||
t4.join();
|
||||
}
|
||||
|
||||
resetTracker();
|
||||
}
|
||||
|
||||
|
||||
///================================================================================================
|
||||
|
||||
//---
|
||||
//Repeatedly query a single entry:
|
||||
void thread_query(void *ptr, const hipPointerAttribute_t *refAttrib)
|
||||
{
|
||||
int count = 0;
|
||||
|
||||
for (int count=0; count< 1000000; count++) {
|
||||
hipPointerAttribute_t a;
|
||||
hipError_t e = hipPointerGetAttributes(&a, ptr);
|
||||
if ((e != hipSuccess) || (a!= *refAttrib)) {
|
||||
printf("Test %d (err=%d)\n", count, e);
|
||||
HIPCHECK(e);
|
||||
|
||||
printf(" ref :: "); printAttribs(refAttrib);
|
||||
printf(" getattr:: "); printAttribs(&a);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
//---
|
||||
// Add pointers to tracker very quickly, then remove them quickly:
|
||||
enum Dir {Up, Down};
|
||||
void thread_noise_generator(int iters, size_t numBuffers, Dir addDir, Dir removeDir)
|
||||
{
|
||||
const size_t bufferSize = 16;
|
||||
size_t maxSize = numBuffers*bufferSize;
|
||||
HIPASSERT((maxSize % bufferSize) == 0); // loop logic assumes this is true
|
||||
|
||||
|
||||
for (int i=0; i<iters; i++) {
|
||||
char * basePtr = (char*)malloc(maxSize);
|
||||
|
||||
auto acc = hc::accelerator();
|
||||
|
||||
if (addDir == Up) {
|
||||
for (char *p = basePtr; p<basePtr + maxSize; p+=bufferSize) {
|
||||
hc::am_memtracker_add(p, bufferSize, acc, false);
|
||||
}
|
||||
} else if (addDir == Down) {
|
||||
for (char *p = basePtr+maxSize-bufferSize; p>=0; p-=bufferSize) {
|
||||
hc::am_memtracker_add(p, bufferSize, acc, false);
|
||||
}
|
||||
}
|
||||
|
||||
if (removeDir == Up) {
|
||||
for (char *p = basePtr; p<basePtr + maxSize; p+=bufferSize) {
|
||||
hc::am_memtracker_remove(p);
|
||||
}
|
||||
} else if (removeDir == Down) {
|
||||
for (char *p = basePtr+maxSize-bufferSize; p>=0; p-=bufferSize) {
|
||||
hc::am_memtracker_remove(p);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
//Multi-thread test that is effective at catching locking errors in the alloc/dealloc/tracker.
|
||||
//The query thread repeately requests information on the same block of memory.
|
||||
//Meanwhile, the thread_noise_generator registers a large number of blocks, and
|
||||
//then unregisters them. This causes a large amount of rebalancing in the tree
|
||||
//structure and will generate errors unless the locks in the tracker are preventing reading
|
||||
//while writing.
|
||||
void testMultiThreaded_2()
|
||||
{
|
||||
std::atomic<int> inflight(2);
|
||||
|
||||
printf ("\n===========================================================================\n");
|
||||
printf ("MultiThreaded_2\n");
|
||||
printf ("===========================================================================\n");
|
||||
|
||||
hipSetDevice(0);
|
||||
hipDeviceReset();
|
||||
|
||||
// Create some entries in the tracker:
|
||||
for (int i=0; i<1000; i++) {
|
||||
void *C_d;
|
||||
HIPCHECK(hipMalloc(&C_d, 32));
|
||||
}
|
||||
|
||||
|
||||
// Allocate a pointer that we will repeatedly lookup:
|
||||
void *A_d;
|
||||
HIPCHECK(hipMalloc(&A_d, 10000));
|
||||
hipPointerAttribute_t attrib1;
|
||||
HIPCHECK(hipPointerGetAttributes(&attrib1, A_d));
|
||||
std::thread t1(thread_query, A_d, &attrib1);
|
||||
|
||||
std::thread t2(thread_noise_generator, 10000, 1000, Up, Up);
|
||||
|
||||
t1.join();
|
||||
t2.join();
|
||||
|
||||
hipSetDevice(0);
|
||||
hipDeviceReset();
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
N= 1000000;
|
||||
HipTest::parseStandardArguments(argc, argv, true);
|
||||
|
||||
|
||||
Nbytes = N*sizeof(char);
|
||||
|
||||
printf ("N=%zu (%6.2f MB) device=%d\n", N, Nbytes/(1024.0*1024.0), p_gpuDevice);
|
||||
|
||||
|
||||
if (p_tests & 0x01) {
|
||||
printf ("info: set device to %d\n", p_gpuDevice);
|
||||
HIPCHECK(hipSetDevice(p_gpuDevice));
|
||||
testSimple();
|
||||
}
|
||||
|
||||
if (p_tests & 0x02) {
|
||||
srand(0x100);
|
||||
printf ("\n===========================================================================\n");
|
||||
clusterAllocs(100, 1024*1, 1024*1024);
|
||||
resetTracker();
|
||||
}
|
||||
|
||||
if (p_tests & 0x04) {
|
||||
srand(0x200);
|
||||
printf ("\n===========================================================================\n");
|
||||
clusterAllocs(1000, 1, 10); // Many tiny allocations;
|
||||
resetTracker();
|
||||
}
|
||||
|
||||
if (p_tests & 0x08) {
|
||||
srand(0x300);
|
||||
testMultiThreaded_1(true);
|
||||
testMultiThreaded_1(false);
|
||||
}
|
||||
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
if (p_tests & 0x10) {
|
||||
srand(0x400);
|
||||
testMultiThreaded_2();
|
||||
resetTracker();
|
||||
}
|
||||
#endif
|
||||
|
||||
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*/
|
||||
|
||||
|
||||
|
||||
@@ -86,7 +88,7 @@ int parseStandardArguments(int argc, char *argv[], bool failOnUndefinedArg)
|
||||
|
||||
if (!strcmp(arg, " ")) {
|
||||
// skip NULL args.
|
||||
} else if (!strcmp(arg, "--N")) {
|
||||
} else if (!strcmp(arg, "--N") || (!strcmp(arg, "-N"))) {
|
||||
if (++i >= argc || !HipTest::parseSize(argv[i], &N)) {
|
||||
failed("Bad N size argument");
|
||||
}
|
||||
@@ -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 {
|
||||
|
||||
@@ -25,7 +25,7 @@
|
||||
printf (__VA_ARGS__);\
|
||||
printf ("\n");\
|
||||
printf ("error: TEST FAILED\n%s", KNRM );\
|
||||
exit(EXIT_FAILURE);
|
||||
abort();
|
||||
|
||||
|
||||
#define HIPCHECK(error) \
|
||||
@@ -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 {
|
||||
|
||||
@@ -86,7 +88,7 @@ vectorADD(hipLaunchParm lp,
|
||||
size_t stride = hipBlockDim_x * hipGridDim_x ;
|
||||
|
||||
for (size_t i=offset; i<N; i+=stride) {
|
||||
C_d[i] = A_d[i] + B_d[i];
|
||||
C_d[i] = A_d[i] + B_d[i];
|
||||
}
|
||||
}
|
||||
|
||||
@@ -94,7 +96,7 @@ vectorADD(hipLaunchParm lp,
|
||||
template <typename T>
|
||||
void initArrays(T **A_d, T **B_d, T **C_d,
|
||||
T **A_h, T **B_h, T **C_h,
|
||||
size_t N)
|
||||
size_t N, bool usePinnedHost=false)
|
||||
{
|
||||
size_t Nbytes = N*sizeof(T);
|
||||
|
||||
@@ -108,14 +110,32 @@ void initArrays(T **A_d, T **B_d, T **C_d,
|
||||
HIPCHECK ( hipMalloc(C_d, Nbytes) );
|
||||
}
|
||||
|
||||
if (A_h)
|
||||
*A_h = (T*)malloc(Nbytes);
|
||||
|
||||
if (B_h)
|
||||
*B_h = (T*)malloc(Nbytes);
|
||||
if (usePinnedHost) {
|
||||
if (A_h) {
|
||||
HIPCHECK ( hipMallocHost(A_h, Nbytes) );
|
||||
}
|
||||
if (B_h) {
|
||||
HIPCHECK ( hipMallocHost(B_h, Nbytes) );
|
||||
}
|
||||
if (C_h) {
|
||||
HIPCHECK ( hipMallocHost(C_h, Nbytes) );
|
||||
}
|
||||
} else {
|
||||
if (A_h) {
|
||||
*A_h = (T*)malloc(Nbytes);
|
||||
HIPASSERT(*A_h != NULL);
|
||||
}
|
||||
|
||||
if (B_h) {
|
||||
*B_h = (T*)malloc(Nbytes);
|
||||
HIPASSERT(*B_h != NULL);
|
||||
}
|
||||
|
||||
if (C_h)
|
||||
*C_h = (T*)malloc(Nbytes);
|
||||
if (C_h) {
|
||||
*C_h = (T*)malloc(Nbytes);
|
||||
HIPASSERT(*C_h != NULL);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
// Initialize the host data:
|
||||
@@ -128,7 +148,43 @@ void initArrays(T **A_d, T **B_d, T **C_d,
|
||||
}
|
||||
|
||||
|
||||
template <typename T>
|
||||
void freeArrays(T *A_d, T *B_d, T *C_d,
|
||||
T *A_h, T *B_h, T *C_h, bool usePinnedHost)
|
||||
{
|
||||
if (A_d) {
|
||||
HIPCHECK ( hipFree(A_d) );
|
||||
}
|
||||
if (B_d) {
|
||||
HIPCHECK ( hipFree(B_d) );
|
||||
}
|
||||
if (C_d) {
|
||||
HIPCHECK ( hipFree(C_d) );
|
||||
}
|
||||
|
||||
if (usePinnedHost) {
|
||||
if (A_h) {
|
||||
HIPCHECK (hipFreeHost(A_h));
|
||||
}
|
||||
if (B_h) {
|
||||
HIPCHECK (hipFreeHost(B_h));
|
||||
}
|
||||
if (C_h) {
|
||||
HIPCHECK (hipFreeHost(C_h));
|
||||
}
|
||||
} else {
|
||||
if (A_h) {
|
||||
free (A_h);
|
||||
}
|
||||
if (B_h) {
|
||||
free (B_h);
|
||||
}
|
||||
if (C_h) {
|
||||
free (C_h);
|
||||
}
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
|
||||
// Assumes C_h contains vector add of A_h + B_h
|
||||
|
||||
Reference in New Issue
Block a user