From 85c189c846496f3b6fa331fd49419eebb830a82e Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Mon, 24 Apr 2017 15:24:16 -0500 Subject: [PATCH 01/13] changed arguments for hipPointerGetAttributes Change-Id: Ia7a7c4722c1f7d0a23f0e5cc3dd6dea6c01c1fd8 --- hipamd/include/hip/hcc_detail/hip_runtime_api.h | 4 ++-- hipamd/src/hip_memory.cpp | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index f9bfb5a310..80a0db7e2e 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -853,7 +853,7 @@ hipError_t hipEventQuery(hipEvent_t event) ; * * @see hipGetDeviceCount, hipGetDevice, hipSetDevice, hipChooseDevice */ -hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr); +hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, const void* ptr); /** * @brief Allocate memory on the default accelerator @@ -1922,7 +1922,7 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image); * @param [in] blockDimZ Z grid dimension specified in work-items * @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel. The kernel can access this with HIP_DYNAMIC_SHARED. * @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case th default stream is used with associated synchronization rules. - * @param [in] kernelParams + * @param [in] kernelParams * @param [in] extra Pointer to kernel arguments. These are passed directly to the kernel and must be in the memory layout and alignment expected by the kernel. * * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index da5530349f..f7421f9818 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -133,7 +133,7 @@ void * allocAndSharePtr(const char *msg, size_t sizeBytes, ihipCtx_t *ctx, unsig //_appAllocationFlags : These are flags provided by the user when allocation is performed. They are returned to user in hipHostGetFlags and other APIs. // TODO - add more info here when available. // -hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr) +hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, const void* ptr) { HIP_INIT_API(attributes, ptr); @@ -1268,7 +1268,7 @@ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned hsa_amd_ipc_memory_attach((hsa_amd_ipc_memory_t*)&(iHandle->ipc_handle), iHandle->psize, crit->peerCnt(), crit->peerAgents(), devPtr); if(hsa_status != HSA_STATUS_SUCCESS) hipStatus = hipErrorMapBufferObjectFailed; - } + } #else hipStatus = hipErrorRuntimeOther; #endif From 1f532b06f6ce0d38ae0ccae0839c0279437955a3 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Mon, 24 Apr 2017 15:31:07 -0500 Subject: [PATCH 02/13] fixed build issues with hipPointerGetAttributes Change-Id: I3f5fbc05bdaef720884ba949075928752a070377 --- hipamd/src/hip_memory.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index f7421f9818..b706426efb 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -149,10 +149,10 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, const void attributes->devicePointer = amPointerInfo._devicePointer; attributes->isManaged = 0; if(attributes->memoryType == hipMemoryTypeHost){ - attributes->hostPointer = ptr; + attributes->hostPointer = (void*)ptr; } if(attributes->memoryType == hipMemoryTypeDevice){ - attributes->devicePointer = ptr; + attributes->devicePointer = (void*)ptr; } attributes->allocationFlags = amPointerInfo._appAllocationFlags; attributes->device = amPointerInfo._appId; From fb7eee01ff1a46aaa9610d42b2ced80d4493e8b1 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 24 Apr 2017 11:02:38 -0500 Subject: [PATCH 03/13] Fix hipMalloc to return error code if allocation fails. --- .../include/hip/hcc_detail/hip_runtime_api.h | 2 +- hipamd/src/hip_memory.cpp | 18 +++++++++++------- 2 files changed, 12 insertions(+), 8 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 80a0db7e2e..7a99ff0810 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -863,7 +863,7 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, const void * * If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned. * - * @return #hipSuccess + * @return #hipSuccess, #hipErrorMemoryAllocation, #hipErrorInvalidValue (bad context, null *ptr) * * @see hipMallocPitch, hipFree, hipMallocArray, hipFreeArray, hipMalloc3D, hipMalloc3DArray, hipHostFree, hipHostMalloc */ diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index b706426efb..821f64bc76 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -207,22 +207,26 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) HIP_INIT_API(ptr, sizeBytes); HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; + + auto ctx = ihipGetTlsDefaultCtx(); // return NULL pointer when malloc size is 0 if (sizeBytes == 0) { *ptr = NULL; - return ihipLogStatus(hipSuccess); - } + hip_status = hipSuccess; - auto ctx = ihipGetTlsDefaultCtx(); + } else if ((ctx==nullptr) || (ptr == nullptr)) { + hip_status = hipErrorInvalidValue; - if (ctx) { + } else { auto device = ctx->getWriteableDevice(); *ptr = hip_internal::allocAndSharePtr("device_mem", sizeBytes, ctx, 0/*amFlags*/, 0/*hipFlags*/); - } else { - hip_status = hipErrorMemoryAllocation; - } + if(sizeBytes && (*ptr == NULL)){ + hip_status = hipErrorMemoryAllocation; + } + + } return ihipLogStatus(hip_status); From dfacfbb6412edbf5d5a2ae7c3eff80a900c3dd55 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 24 Apr 2017 11:03:32 -0500 Subject: [PATCH 04/13] Fix hip debug for case where copyAgent is null (host-to-host) --- hipamd/src/hip_hcc.cpp | 16 ++++++++++------ 1 file changed, 10 insertions(+), 6 deletions(-) diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index 35a3e11e71..080d700e63 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -1765,20 +1765,24 @@ void ihipStream_t::resolveHcMemcpyDirection(unsigned hipMemKind, if (HIP_FORCE_P2P_HOST & 0x1) { *forceUnpinnedCopy = true; - tprintf (DB_COPY, "P2P. Copy engine (dev:%d agent=0x%lx) can see src and dst but HIP_FORCE_P2P_HOST=0, forcing copy through staging buffers.\n", - (*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle); + tprintf (DB_COPY, "Copy engine (dev:%d agent=0x%lx) can see src and dst but HIP_FORCE_P2P_HOST=0, forcing copy through staging buffers.\n", + *copyDevice ? (*copyDevice)->getDeviceNum() : -1, + *copyDevice ? (*copyDevice)->getDevice()->_hsaAgent.handle : 0x0); } else { - tprintf (DB_COPY, "P2P. Copy engine (dev:%d agent=0x%lx) can see src and dst.\n", - (*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle); + tprintf (DB_COPY, "Copy engine (dev:%d agent=0x%lx) can see src and dst.\n", + *copyDevice ? (*copyDevice)->getDeviceNum() : -1, + *copyDevice ? (*copyDevice)->getDevice()->_hsaAgent.handle : 0x0); } } else { *forceUnpinnedCopy = true; tprintf (DB_COPY, "P2P: Copy engine(dev:%d agent=0x%lx) cannot see both host and device pointers - forcing copy with unpinned engine.\n", - (*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle); + *copyDevice ? (*copyDevice)->getDeviceNum() : -1, + *copyDevice ? (*copyDevice)->getDevice()->_hsaAgent.handle : 0x0); if (HIP_FAIL_SOC & 0x2) { fprintf (stderr, "HIP_FAIL_SOC: P2P: copy engine(dev:%d agent=0x%lx) cannot see both host and device pointers - forcing copy with unpinned engine.\n", - (*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle); + *copyDevice ? (*copyDevice)->getDeviceNum() : -1, + *copyDevice ? (*copyDevice)->getDevice()->_hsaAgent.handle : 0x0); throw ihipException(hipErrorRuntimeOther); } } From 693e5abc1cdc517c415d2188ccdf2fcbadb647a5 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 24 Apr 2017 11:06:54 -0500 Subject: [PATCH 05/13] Add negative testing for memory full condition. --- .../tests/src/runtimeApi/memory/hipMemoryAllocate.cpp | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemoryAllocate.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemoryAllocate.cpp index 1f7599491a..0a256d6362 100644 --- a/hipamd/tests/src/runtimeApi/memory/hipMemoryAllocate.cpp +++ b/hipamd/tests/src/runtimeApi/memory/hipMemoryAllocate.cpp @@ -56,5 +56,15 @@ int main(){ HIPCHECK_API(hipFree(NULL) , hipSuccess); HIPCHECK_API(hipHostFree(NULL) , hipSuccess); + + { + // Some negative testing - request a too-big allocation and verify it fails: + // Someday when we support virtual memory may need to refactor these: + size_t tooBig = 128LL*1024*1024*1024*1024; // 128 TB; + void *p; + HIPCHECK_API ( hipMalloc(&p, tooBig), hipErrorMemoryAllocation ); + HIPCHECK_API ( hipHostMalloc(&p, tooBig), hipErrorMemoryAllocation ); + } + passed(); } From 6b16f56f38a34381f411e56412f2ce90eb091afe Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 24 Apr 2017 11:53:31 -0500 Subject: [PATCH 06/13] Refactor hipMemcpy test to share mem alloc for multiple copies. --- hipamd/.vimrc | 1 - .../tests/src/runtimeApi/memory/hipMemcpy.cpp | 239 +++++++++++++----- 2 files changed, 182 insertions(+), 58 deletions(-) delete mode 100644 hipamd/.vimrc diff --git a/hipamd/.vimrc b/hipamd/.vimrc deleted file mode 100644 index 019afa57e6..0000000000 --- a/hipamd/.vimrc +++ /dev/null @@ -1 +0,0 @@ -:set makeprg=make\ -C\ build.hcc-LC.db diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp index a320a86022..d50a810a58 100644 --- a/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp +++ b/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp @@ -38,6 +38,130 @@ void printSep() printf ("======================================================================================\n"); } +//------- +template +class DeviceMemory +{ +public: + DeviceMemory(size_t numElements); + ~DeviceMemory(); +public: + T * A_d; + T* B_d; + T* C_d; + T* C_dd; + + size_t _maxNumElements; +}; + +template +DeviceMemory::DeviceMemory(size_t numElements) + : _maxNumElements(numElements) +{ + T ** np = nullptr; + HipTest::initArrays (&A_d, &B_d, &C_d, np, np, np, numElements, 0); + + + size_t sizeElements = numElements * sizeof(T); + + + HIPCHECK ( hipMalloc(&C_dd, sizeElements) ); +} + + +template +DeviceMemory::~DeviceMemory () +{ + T * np = nullptr; + HipTest::freeArrays (A_d, B_d, C_d, np, np, np, 0); + + HIPCHECK (hipFree(C_dd)); + + C_dd = NULL; +}; + + + +//------- +template +class HostMemory +{ +public: + HostMemory(size_t numElements, bool usePinnedHost); + void reset(size_t numElements, bool full=false) ; + ~HostMemory(); +public: + // Host arrays + T * A_h; + T* B_h; + T* C_h; + + // Host arrays, secondary copy + T * A_hh; + T* B_hh; + + size_t _maxNumElements; + bool _usePinnedHost; +}; + +template +HostMemory::HostMemory(size_t numElements, bool usePinnedHost) + : _maxNumElements(numElements), + _usePinnedHost(usePinnedHost) +{ + T ** np = nullptr; + HipTest::initArrays (np, np, np, &A_h, &B_h, &C_h, numElements, usePinnedHost); + + A_hh = NULL; + B_hh = NULL; + + + size_t sizeElements = numElements * sizeof(T); + + if (usePinnedHost) { + HIPCHECK ( hipHostMalloc((void**)&A_hh, sizeElements, hipHostMallocDefault) ); + HIPCHECK ( hipHostMalloc((void**)&B_hh, sizeElements, hipHostMallocDefault) ); + } else { + A_hh = (T*)malloc(sizeElements); + B_hh = (T*)malloc(sizeElements); + } + +} + + +template +void +HostMemory::reset(size_t numElements, bool full) +{ + // Initialize the host data: + for (size_t i=0; i +HostMemory::~HostMemory () +{ + HipTest::freeArraysForHost (A_h, B_h, C_h, _usePinnedHost); + + if (_usePinnedHost) { + HIPCHECK (hipHostFree(A_hh)); + HIPCHECK (hipHostFree(B_hh)); + + } else { + free(A_hh); + free(B_hh); + } + T *A_hh = NULL; + T *B_hh = NULL; + +}; @@ -52,71 +176,55 @@ void printSep() // IN: useMemkindDefault : If true, use memkinddefault (runtime figures out direction). if false, use explicit memcpy direction. // template -void memcpytest2(size_t numElements, bool usePinnedHost, bool useHostToHost, bool useDeviceToDevice, bool useMemkindDefault) +void memcpytest2(DeviceMemory *dmem, HostMemory *hmem, 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__, TYPENAME(T), sizeElements, sizeElements/1024.0/1024.0, - usePinnedHost, useHostToHost, useDeviceToDevice, useMemkindDefault); + hmem->_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); + hmem->reset(numElements); unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements); - T *A_hh = NULL; - T *B_hh = NULL; - T *C_dd = NULL; + assert (numElements <= dmem->_maxNumElements); + assert (numElements <= hmem->_maxNumElements); if (useHostToHost) { - if (usePinnedHost) { - HIPCHECK ( hipHostMalloc((void**)&A_hh, sizeElements, hipHostMallocDefault) ); - HIPCHECK ( hipHostMalloc((void**)&B_hh, sizeElements, hipHostMallocDefault) ); - } 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(hmem->A_hh, hmem->A_h, sizeElements, useMemkindDefault? hipMemcpyDefault : hipMemcpyHostToHost)); + HIPCHECK ( hipMemcpy(hmem->B_hh, hmem->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)); + HIPCHECK ( hipMemcpy(dmem->A_d, hmem->A_hh, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); + HIPCHECK ( hipMemcpy(dmem->B_d, hmem->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)); + HIPCHECK ( hipMemcpy(dmem->A_d, hmem->A_h, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); + HIPCHECK ( hipMemcpy(dmem->B_d, hmem->B_h, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); } - hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d, B_d, C_d, numElements); + hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, dmem->A_d, dmem->B_d, dmem->C_d, numElements); if (useDeviceToDevice) { - HIPCHECK ( hipMalloc(&C_dd, sizeElements) ); + // Do an extra device-to-device copy here to mix things up: + HIPCHECK ( hipMemcpy(dmem->C_dd, dmem->C_d, sizeElements, useMemkindDefault? hipMemcpyDefault : hipMemcpyDeviceToDevice)); - // 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 dmem->C_d: + HIPCHECK ( hipMemset(dmem->C_d, 0x5A, sizeElements)); - //Destroy the original C_d: - HIPCHECK ( hipMemset(C_d, 0x5A, sizeElements)); - - HIPCHECK ( hipMemcpy(C_h, C_dd, sizeElements, useMemkindDefault? hipMemcpyDefault:hipMemcpyDeviceToHost)); + HIPCHECK ( hipMemcpy(hmem->C_h, dmem->C_dd, sizeElements, useMemkindDefault? hipMemcpyDefault:hipMemcpyDeviceToHost)); } else { - HIPCHECK ( hipMemcpy(C_h, C_d, sizeElements, useMemkindDefault? hipMemcpyDefault:hipMemcpyDeviceToHost)); + HIPCHECK ( hipMemcpy(hmem->C_h, dmem->C_d, sizeElements, useMemkindDefault? hipMemcpyDefault:hipMemcpyDeviceToHost)); } HIPCHECK ( hipDeviceSynchronize() ); - HipTest::checkVectorADD(A_h, B_h, C_h, numElements); + HipTest::checkVectorADD(hmem->A_h, hmem->B_h, hmem->C_h, numElements); + - HipTest::freeArrays (A_d, B_d, C_d, A_h, B_h, C_h, usePinnedHost); printf (" %s success\n", __func__); } @@ -129,11 +237,15 @@ void memcpytest2_for_type(size_t numElements) { printSep(); + DeviceMemory memD(numElements); + HostMemory memU(numElements, 0/*usePinnedHost*/); + HostMemory memP(numElements, 1/*usePinnedHost*/); + 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(numElements, usePinnedHost, useHostToHost, useDeviceToDevice, useMemkindDefault); + memcpytest2(&memD, usePinnedHost ? &memP : &memU, numElements, usePinnedHost, useHostToHost, useDeviceToDevice, useMemkindDefault); } } } @@ -156,17 +268,19 @@ void memcpytest2_sizes(size_t maxElem=0, size_t offset=0) HIPCHECK(hipMemGetInfo(&free, &total)); if (maxElem == 0) { - maxElem = free/sizeof(T)/5; + maxElem = free/sizeof(T)/20; } 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); + HIPCHECK ( hipDeviceReset() ); + DeviceMemory memD(maxElem); + HostMemory memU(maxElem, 0/*usePinnedHost*/); + HostMemory memP(maxElem, 1/*usePinnedHost*/); for (size_t elem=64; elem+offset<=maxElem; elem*=2) { - HIPCHECK ( hipDeviceReset() ); - memcpytest2(elem+offset, 0, 1, 1, 0); // unpinned host - HIPCHECK ( hipDeviceReset() ); - memcpytest2(elem+offset, 1, 1, 1, 0); // pinned host + memcpytest2(&memD, &memU, elem+offset, 0, 1, 1, 0); // unpinned host + memcpytest2(&memD, &memP, elem+offset, 1, 1, 1, 0); // pinned host } } @@ -178,13 +292,17 @@ void multiThread_1(bool serialize, bool usePinnedHost) { printSep(); printf ("test: %s<%s> serialize=%d usePinnedHost=%d\n", __func__, TYPENAME(T), serialize, usePinnedHost); - std::thread t1 (memcpytest2,N, usePinnedHost,0,0,0); + DeviceMemory memD(N); + HostMemory mem1(N, usePinnedHost); + HostMemory mem2(N, usePinnedHost); + + std::thread t1 (memcpytest2, &memD, &mem1, N, usePinnedHost,0,0,0); if (serialize) { t1.join(); } - std::thread t2 (memcpytest2,N, usePinnedHost,0,0,0); + std::thread t2 (memcpytest2,&memD, &mem2, N, usePinnedHost,0,0,0); if (serialize) { t2.join(); } @@ -218,24 +336,30 @@ int main(int argc, char *argv[]) if (p_tests & 0x2) { - // Some tests around the 64MB boundary which have historically shown issues: - printf ("\n\n=== tests&0x2 (64MB boundary)\n"); -#if 0 + // Some tests around the 64KB boundary which have historically shown issues: + printf ("\n\n=== tests&0x2 (64KB boundary)\n"); + size_t maxElem = 32*1024*1024; + DeviceMemory memD(maxElem); + HostMemory memU(maxElem, 0/*usePinnedHost*/); + HostMemory memP(maxElem, 0/*usePinnedHost*/); // These all pass: - memcpytest2(15*1024*1024, 1, 0, 0, 0); - memcpytest2(16*1024*1024, 1, 0, 0, 0); - memcpytest2(16*1024*1024+16*1024, 1, 0, 0, 0); -#endif + memcpytest2(&memD, &memP, 15*1024*1024, 1, 0, 0, 0); + memcpytest2(&memD, &memP, 16*1024*1024, 1, 0, 0, 0); + memcpytest2(&memD, &memP, 16*1024*1024+16*1024, 1, 0, 0, 0); + // Just over 64MB: - memcpytest2(16*1024*1024+512*1024, 1, 0, 0, 0); - memcpytest2(17*1024*1024+1024, 1, 0, 0, 0); - memcpytest2(32*1024*1024, 1, 0, 0, 0); - memcpytest2(32*1024*1024, 0, 0, 0, 0); - memcpytest2(32*1024*1024, 1, 1, 1, 0); - memcpytest2(32*1024*1024, 1, 1, 1, 0); + memcpytest2(&memD, &memP, 16*1024*1024+512*1024, 1, 0, 0, 0); + memcpytest2(&memD, &memP, 17*1024*1024+1024, 1, 0, 0, 0); + memcpytest2(&memD, &memP, 32*1024*1024, 1, 0, 0, 0); + memcpytest2(&memD, &memU, 32*1024*1024, 0, 0, 0, 0); + memcpytest2(&memD, &memP, 32*1024*1024, 1, 1, 1, 0); + memcpytest2(&memD, &memP, 32*1024*1024, 1, 1, 1, 0); + + } + if (p_tests & 0x4) { printf ("\n\n=== tests&4 (test sizes and offsets)\n"); HIPCHECK ( hipDeviceReset() ); @@ -270,6 +394,7 @@ int main(int argc, char *argv[]) } + passed(); } From b44a3eefd17aaee1c79a77cfbf8d50060e907f86 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 24 Apr 2017 12:51:17 -0500 Subject: [PATCH 07/13] Add corrected test for offsets --- .../tests/src/runtimeApi/memory/hipMemcpy.cpp | 168 ++++++++++++------ 1 file changed, 115 insertions(+), 53 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp index d50a810a58..ad798d70c1 100644 --- a/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp +++ b/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp @@ -24,6 +24,7 @@ THE SOFTWARE. * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 * RUN_NAMED: %t hipMemcpy-modes --tests 0x1 * RUN_NAMED: %t hipMemcpy-size --tests 0x6 + * RUN_NAMED: %t hipMemcpy-offsets --tests 0x10 * RUN_NAMED: %t hipMemcpy-multithreaded --tests 0x8 * HIT_END */ @@ -45,27 +46,42 @@ class DeviceMemory public: DeviceMemory(size_t numElements); ~DeviceMemory(); -public: - T * A_d; - T* B_d; - T* C_d; - T* C_dd; + + T *A_d() const { return _A_d + _offset; }; + T *B_d() const { return _B_d + _offset; }; + T *C_d() const { return _C_d + _offset; }; + T *C_dd() const { return _C_dd + _offset; }; + + size_t maxNumElements() const { return _maxNumElements; }; + + + void offset(int offset) { _offset = offset; }; + int offset() const { return _offset; }; + +private: + T * _A_d; + T* _B_d; + T* _C_d; + T* _C_dd; + size_t _maxNumElements; + int _offset; }; template DeviceMemory::DeviceMemory(size_t numElements) - : _maxNumElements(numElements) + : _maxNumElements(numElements), + _offset(0) { T ** np = nullptr; - HipTest::initArrays (&A_d, &B_d, &C_d, np, np, np, numElements, 0); + HipTest::initArrays (&_A_d, &_B_d, &_C_d, np, np, np, numElements, 0); size_t sizeElements = numElements * sizeof(T); - HIPCHECK ( hipMalloc(&C_dd, sizeElements) ); + HIPCHECK ( hipMalloc(&_C_dd, sizeElements) ); } @@ -73,11 +89,11 @@ template DeviceMemory::~DeviceMemory () { T * np = nullptr; - HipTest::freeArrays (A_d, B_d, C_d, np, np, np, 0); + HipTest::freeArrays (_A_d, _B_d, _C_d, np, np, np, 0); - HIPCHECK (hipFree(C_dd)); + HIPCHECK (hipFree(_C_dd)); - C_dd = NULL; + _C_dd = NULL; }; @@ -90,6 +106,8 @@ public: HostMemory(size_t numElements, bool usePinnedHost); void reset(size_t numElements, bool full=false) ; ~HostMemory(); + + size_t maxNumElements() const { return _maxNumElements; }; public: // Host arrays T * A_h; @@ -176,21 +194,22 @@ HostMemory::~HostMemory () // IN: useMemkindDefault : If true, use memkinddefault (runtime figures out direction). if false, use explicit memcpy direction. // template -void memcpytest2(DeviceMemory *dmem, HostMemory *hmem, size_t numElements, bool usePinnedHost, bool useHostToHost, bool useDeviceToDevice, bool useMemkindDefault) +void memcpytest2(DeviceMemory *dmem, HostMemory *hmem, size_t numElements, 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", + printf ("test: %s<%s> size=%lu (%6.2fMB) usePinnedHost:%d, useHostToHost:%d, useDeviceToDevice:%d, useMemkindDefault:%d, offsets:%+d\n", __func__, TYPENAME(T), sizeElements, sizeElements/1024.0/1024.0, - hmem->_usePinnedHost, useHostToHost, useDeviceToDevice, useMemkindDefault); + hmem->_usePinnedHost, useHostToHost, useDeviceToDevice, useMemkindDefault, + dmem->offset()); hmem->reset(numElements); unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements); - assert (numElements <= dmem->_maxNumElements); - assert (numElements <= hmem->_maxNumElements); + assert (numElements <= dmem->maxNumElements()); + assert (numElements <= hmem->maxNumElements()); @@ -200,25 +219,25 @@ void memcpytest2(DeviceMemory *dmem, HostMemory *hmem, size_t numElements, HIPCHECK ( hipMemcpy(hmem->B_hh, hmem->B_h, sizeElements, useMemkindDefault? hipMemcpyDefault : hipMemcpyHostToHost)); - HIPCHECK ( hipMemcpy(dmem->A_d, hmem->A_hh, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); - HIPCHECK ( hipMemcpy(dmem->B_d, hmem->B_hh, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); + HIPCHECK ( hipMemcpy(dmem->A_d(), hmem->A_hh, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); + HIPCHECK ( hipMemcpy(dmem->B_d(), hmem->B_hh, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); } else { - HIPCHECK ( hipMemcpy(dmem->A_d, hmem->A_h, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); - HIPCHECK ( hipMemcpy(dmem->B_d, hmem->B_h, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); + HIPCHECK ( hipMemcpy(dmem->A_d(), hmem->A_h, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); + HIPCHECK ( hipMemcpy(dmem->B_d(), hmem->B_h, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); } - hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, dmem->A_d, dmem->B_d, dmem->C_d, numElements); + hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, dmem->A_d(), dmem->B_d(), dmem->C_d(), numElements); if (useDeviceToDevice) { // Do an extra device-to-device copy here to mix things up: - HIPCHECK ( hipMemcpy(dmem->C_dd, dmem->C_d, sizeElements, useMemkindDefault? hipMemcpyDefault : hipMemcpyDeviceToDevice)); + HIPCHECK ( hipMemcpy(dmem->C_dd(), dmem->C_d(), sizeElements, useMemkindDefault? hipMemcpyDefault : hipMemcpyDeviceToDevice)); - //Destroy the original dmem->C_d: - HIPCHECK ( hipMemset(dmem->C_d, 0x5A, sizeElements)); + //Destroy the original dmem->C_d(): + HIPCHECK ( hipMemset(dmem->C_d(), 0x5A, sizeElements)); - HIPCHECK ( hipMemcpy(hmem->C_h, dmem->C_dd, sizeElements, useMemkindDefault? hipMemcpyDefault:hipMemcpyDeviceToHost)); + HIPCHECK ( hipMemcpy(hmem->C_h, dmem->C_dd(), sizeElements, useMemkindDefault? hipMemcpyDefault:hipMemcpyDeviceToHost)); } else { - HIPCHECK ( hipMemcpy(hmem->C_h, dmem->C_d, sizeElements, useMemkindDefault? hipMemcpyDefault:hipMemcpyDeviceToHost)); + HIPCHECK ( hipMemcpy(hmem->C_h, dmem->C_d(), sizeElements, useMemkindDefault? hipMemcpyDefault:hipMemcpyDeviceToHost)); } HIPCHECK ( hipDeviceSynchronize() ); @@ -245,7 +264,7 @@ void memcpytest2_for_type(size_t numElements) for (int useHostToHost =0; useHostToHost<=1; useHostToHost++) { // TODO for (int useDeviceToDevice =0; useDeviceToDevice<=1; useDeviceToDevice++) { for (int useMemkindDefault =0; useMemkindDefault<=1; useMemkindDefault++) { - memcpytest2(&memD, usePinnedHost ? &memP : &memU, numElements, usePinnedHost, useHostToHost, useDeviceToDevice, useMemkindDefault); + memcpytest2(&memD, usePinnedHost ? &memP : &memU, numElements, useHostToHost, useDeviceToDevice, useMemkindDefault); } } } @@ -256,7 +275,7 @@ void memcpytest2_for_type(size_t numElements) //--- //Try many different sizes to memory copy. template -void memcpytest2_sizes(size_t maxElem=0, size_t offset=0) +void memcpytest2_sizes(size_t maxElem=0) { printSep(); printf ("test: %s<%s>\n", __func__, TYPENAME(T)); @@ -268,19 +287,59 @@ void memcpytest2_sizes(size_t maxElem=0, size_t offset=0) HIPCHECK(hipMemGetInfo(&free, &total)); if (maxElem == 0) { - maxElem = free/sizeof(T)/20; + 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); + printf (" device#%d: hipMemGetInfo: free=%zu (%4.2fMB) total=%zu (%4.2fMB) maxSize=%6.1fMB\n", + deviceId, free, (float)(free/1024.0/1024.0), total, (float)(total/1024.0/1024.0), maxElem*sizeof(T)/1024.0/1024.0); HIPCHECK ( hipDeviceReset() ); DeviceMemory memD(maxElem); HostMemory memU(maxElem, 0/*usePinnedHost*/); HostMemory memP(maxElem, 1/*usePinnedHost*/); - for (size_t elem=64; elem+offset<=maxElem; elem*=2) { - memcpytest2(&memD, &memU, elem+offset, 0, 1, 1, 0); // unpinned host - memcpytest2(&memD, &memP, elem+offset, 1, 1, 1, 0); // pinned host + for (size_t elem=1; elem<=maxElem; elem*=2) { + memcpytest2(&memD, &memU, elem, 1, 1, 0); // unpinned host + memcpytest2(&memD, &memP, elem, 1, 1, 0); // pinned host + } +} + + +//--- +//Try many different sizes to memory copy. +template +void memcpytest2_offsets(size_t maxElem) +{ + printSep(); + printf ("test: %s<%s>\n", __func__, TYPENAME(T)); + + int deviceId; + HIPCHECK(hipGetDevice(&deviceId)); + + size_t free, total; + HIPCHECK(hipMemGetInfo(&free, &total)); + + + printf (" device#%d: hipMemGetInfo: free=%zu (%4.2fMB) total=%zu (%4.2fMB) maxSize=%6.1fMB\n", + deviceId, free, (float)(free/1024.0/1024.0), total, (float)(total/1024.0/1024.0), maxElem*sizeof(T)/1024.0/1024.0); + HIPCHECK ( hipDeviceReset() ); + DeviceMemory memD(maxElem); + HostMemory memU(maxElem, 0/*usePinnedHost*/); + HostMemory memP(maxElem, 1/*usePinnedHost*/); + + size_t elem = maxElem / 2; + + for (int offset=0; offset < 512; offset++) { + assert (elem + offset < maxElem); + memD.offset(offset); + memcpytest2(&memD, &memU, elem, 1, 1, 0); // unpinned host + memcpytest2(&memD, &memP, elem, 1, 1, 0); // pinned host + } + + for (int offset=512; offset < maxElem; offset*=2) { + assert (elem + offset < maxElem); + memD.offset(offset); + memcpytest2(&memD, &memU, elem, 1, 1, 0); // unpinned host + memcpytest2(&memD, &memP, elem, 1, 1, 0); // pinned host } } @@ -296,13 +355,13 @@ void multiThread_1(bool serialize, bool usePinnedHost) HostMemory mem1(N, usePinnedHost); HostMemory mem2(N, usePinnedHost); - std::thread t1 (memcpytest2, &memD, &mem1, N, usePinnedHost,0,0,0); + std::thread t1 (memcpytest2, &memD, &mem1, N, 0,0,0); if (serialize) { t1.join(); } - std::thread t2 (memcpytest2,&memD, &mem2, N, usePinnedHost,0,0,0); + std::thread t2 (memcpytest2,&memD, &mem2, N, 0,0,0); if (serialize) { t2.join(); } @@ -343,17 +402,17 @@ int main(int argc, char *argv[]) HostMemory memU(maxElem, 0/*usePinnedHost*/); HostMemory memP(maxElem, 0/*usePinnedHost*/); // These all pass: - memcpytest2(&memD, &memP, 15*1024*1024, 1, 0, 0, 0); - memcpytest2(&memD, &memP, 16*1024*1024, 1, 0, 0, 0); - memcpytest2(&memD, &memP, 16*1024*1024+16*1024, 1, 0, 0, 0); + memcpytest2(&memD, &memP, 15*1024*1024, 0, 0, 0); + memcpytest2(&memD, &memP, 16*1024*1024, 0, 0, 0); + memcpytest2(&memD, &memP, 16*1024*1024+16*1024, 0, 0, 0); // Just over 64MB: - memcpytest2(&memD, &memP, 16*1024*1024+512*1024, 1, 0, 0, 0); - memcpytest2(&memD, &memP, 17*1024*1024+1024, 1, 0, 0, 0); - memcpytest2(&memD, &memP, 32*1024*1024, 1, 0, 0, 0); - memcpytest2(&memD, &memU, 32*1024*1024, 0, 0, 0, 0); - memcpytest2(&memD, &memP, 32*1024*1024, 1, 1, 1, 0); - memcpytest2(&memD, &memP, 32*1024*1024, 1, 1, 1, 0); + memcpytest2(&memD, &memP, 16*1024*1024+512*1024, 0, 0, 0); + memcpytest2(&memD, &memP, 17*1024*1024+1024, 0, 0, 0); + memcpytest2(&memD, &memP, 32*1024*1024, 0, 0, 0); + memcpytest2(&memD, &memU, 32*1024*1024, 0, 0, 0); + memcpytest2(&memD, &memP, 32*1024*1024, 1, 1, 0); + memcpytest2(&memD, &memP, 32*1024*1024, 1, 1, 0); } @@ -361,16 +420,19 @@ int main(int argc, char *argv[]) if (p_tests & 0x4) { - printf ("\n\n=== tests&4 (test sizes and offsets)\n"); + printf ("\n\n=== tests&4 (test sizes)\n"); HIPCHECK ( hipDeviceReset() ); + memcpytest2_sizes(0); printSep(); - memcpytest2_sizes(0,0); - printSep(); - memcpytest2_sizes(0,64); - printSep(); - memcpytest2_sizes(1024*1024, 13); - printSep(); - memcpytest2_sizes(1024*1024, 50); + } + + + if (p_tests & 0x10) { + printf ("\n\n=== tests&4 (test offsets)\n"); + HIPCHECK ( hipDeviceReset() ); + memcpytest2_offsets(256*1024*1024); + memcpytest2_offsets(256*1024*1024); + memcpytest2_offsets(256*1024*1024); } if (p_tests & 0x8) { From d120b2dd123829057736e34a43681998b21f4d4e Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 24 Apr 2017 16:02:22 -0500 Subject: [PATCH 08/13] Add test for non-page-aligned mem copies. --- .../tests/src/runtimeApi/memory/hipMemcpy.cpp | 21 +++++++++++-------- 1 file changed, 12 insertions(+), 9 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp index ad798d70c1..c48f780e44 100644 --- a/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp +++ b/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp @@ -24,7 +24,7 @@ THE SOFTWARE. * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 * RUN_NAMED: %t hipMemcpy-modes --tests 0x1 * RUN_NAMED: %t hipMemcpy-size --tests 0x6 - * RUN_NAMED: %t hipMemcpy-offsets --tests 0x10 + * RUN_NAMED: %t hipMemcpy-dev_offsets --tests 0x10 * RUN_NAMED: %t hipMemcpy-multithreaded --tests 0x8 * HIT_END */ @@ -335,7 +335,7 @@ void memcpytest2_offsets(size_t maxElem) memcpytest2(&memD, &memP, elem, 1, 1, 0); // pinned host } - for (int offset=512; offset < maxElem; offset*=2) { + for (int offset=512; offset < elem; offset*=2) { assert (elem + offset < maxElem); memD.offset(offset); memcpytest2(&memD, &memU, elem, 1, 1, 0); // unpinned host @@ -427,13 +427,6 @@ int main(int argc, char *argv[]) } - if (p_tests & 0x10) { - printf ("\n\n=== tests&4 (test offsets)\n"); - HIPCHECK ( hipDeviceReset() ); - memcpytest2_offsets(256*1024*1024); - memcpytest2_offsets(256*1024*1024); - memcpytest2_offsets(256*1024*1024); - } if (p_tests & 0x8) { printf ("\n\n=== tests&8\n"); @@ -456,6 +449,16 @@ int main(int argc, char *argv[]) } + if (p_tests & 0x10) { + printf ("\n\n=== tests&0x10 (test device offsets)\n"); + HIPCHECK ( hipDeviceReset() ); + size_t maxSize = 256*1024; + memcpytest2_offsets (maxSize); + memcpytest2_offsets (maxSize); + memcpytest2_offsets(maxSize); + } + + passed(); From 3da8e94cbff8fb977a7186c09996fbcdb0b1d0fc Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 24 Apr 2017 16:55:29 -0500 Subject: [PATCH 09/13] Tailor pointer info for src/dst before calling HCC copy routines. HCC sometimes uses the srcPtrInfo or dstPtrInfo to determine the pointer. Make sure these use the actual pointer and not the base of the allocation. --- hipamd/src/hip_hcc.cpp | 66 +++++++++++++++++++++++++++++++++++++++--- 1 file changed, 62 insertions(+), 4 deletions(-) diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index 080d700e63..71d947488d 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -1798,6 +1798,62 @@ void printPointerInfo(unsigned dbFlag, const char *tag, const void *ptr, const h } +// the pointer-info as returned by HC refers to the allocation +// This routine modifies the pointer-info so it appears to refer to the specific ptr and sizeBytes. +// TODO -remove this when HCC uses HSA pointer info functions directly. +void tailorPtrInfo(hc::AmPointerInfo *ptrInfo, const void * ptr, size_t sizeBytes) +{ + const char *ptrc = static_cast (ptr); + if (ptrInfo->_sizeBytes == 0) { + // invalid ptrInfo, don't modify + return; + } else if (ptrInfo->_isInDeviceMem) { + assert (ptrInfo->_devicePointer != nullptr); + std::ptrdiff_t diff = ptrc - static_cast (ptrInfo->_devicePointer); + + //TODO : assert-> runtime assert that only appears in debug mode + assert (diff >= 0); + assert (diff <= ptrInfo->_sizeBytes); + + ptrInfo->_devicePointer = const_cast (ptr); + + if (ptrInfo->_hostPointer != nullptr) { + ptrInfo->_hostPointer = static_cast(ptrInfo->_hostPointer) + diff; + } + + } else { + + assert (ptrInfo->_hostPointer != nullptr); + std::ptrdiff_t diff = ptrc - static_cast (ptrInfo->_hostPointer); + + //TODO : assert-> runtime assert that only appears in debug mode + assert (diff >= 0); + assert (diff <= ptrInfo->_sizeBytes); + + ptrInfo->_hostPointer = const_cast(ptr); + + if (ptrInfo->_devicePointer != nullptr) { + ptrInfo->_devicePointer = static_cast(ptrInfo->_devicePointer) + diff; + } + } + + assert (sizeBytes <= ptrInfo->_sizeBytes); + ptrInfo->_sizeBytes = sizeBytes; +}; + + +bool getTailoredPtrInfo(hc::AmPointerInfo *ptrInfo, const void * ptr, size_t sizeBytes) +{ + bool tracked = (hc::am_memtracker_getinfo(ptrInfo, ptr) == AM_SUCCESS); + + if (tracked) { + tailorPtrInfo(ptrInfo, ptr, sizeBytes); + } + + return tracked; +}; + + // TODO : For registered and host memory, if the portable flag is set, we need to recognize that and perform appropriate copy operation. // What can happen now is that Portable memory is mapped into multiple devices but Peer access is not enabled. i // The peer detection logic doesn't see that the memory is already mapped and so tries to use an unpinned copy algorithm. If this is PinInPlace, then an error can occur. @@ -1816,8 +1872,8 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, hc::accelerator acc; hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0); hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0); - bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS); - bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS); + bool dstTracked = getTailoredPtrInfo(&dstPtrInfo, dst, sizeBytes); + bool srcTracked = getTailoredPtrInfo(&srcPtrInfo, src, sizeBytes); // Some code in HCC and in printPointerInfo uses _sizeBytes==0 as an indication ptr is not valid, so check it here: @@ -1877,6 +1933,7 @@ void ihipStream_t::lockedSymbolCopySync(hc::accelerator &acc, void* dst, void* s void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, size_t offset, unsigned kind) { + // TODO - review - this looks broken , should not be adding pointers to tracker dynamically: if(kind == hipMemcpyHostToDevice) { hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0); bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS); @@ -1903,6 +1960,7 @@ void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, void* } } + void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind) { @@ -1930,8 +1988,8 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes hc::accelerator acc; hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0); hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0); - bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS); - bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS); + bool dstTracked = getTailoredPtrInfo(&dstPtrInfo, dst, sizeBytes); + bool srcTracked = getTailoredPtrInfo(&srcPtrInfo, src, sizeBytes); hc::hcCommandKind hcCopyDir; From 5ba167b82b7aed87fb1a3ae9ba25c3bfdc89925a Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 24 Apr 2017 20:38:37 -0500 Subject: [PATCH 10/13] Refactor hipHostRegister to cover misaligned cases. --- .../src/runtimeApi/memory/hipHostRegister.cpp | 129 ++++++++++++------ 1 file changed, 84 insertions(+), 45 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp b/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp index 1a1319c500..efa23b4068 100644 --- a/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp +++ b/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp @@ -19,87 +19,126 @@ THE SOFTWARE. /* HIT_START * BUILD: %t %s ../../test_common.cpp - * RUN: %t + * RUN: %t --tests 0x1 + * RUN: %t --tests 0x2 * HIT_END */ +// TODO - bug if run both back-to-back + #include"test_common.h" #include __global__ void Inc(hipLaunchParm lp, float *Ad){ -int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; -Ad[tx] = Ad[tx] + float(1); + int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + Ad[tx] = Ad[tx] + float(1); } -int main(){ - float *A, **Ad; - int num_devices; - HIPCHECK(hipGetDeviceCount(&num_devices)); - Ad = new float*[num_devices]; - const size_t size = N * sizeof(float); - A = (float*)malloc(size); - HIPCHECK(hipHostRegister(A, size, 0)); + +template +void doMemCopy(size_t numElements, int offset, T *A, T *Bh, T *Bd) +{ + A = A + offset; + numElements -= offset; + + size_t sizeBytes = numElements * sizeof(T); + + HIPCHECK(hipHostRegister(A, sizeBytes, 0)); - for(int i=0;iOFFSETS_TO_TRY); + for (size_t i=0; i Date: Mon, 24 Apr 2017 21:05:29 -0500 Subject: [PATCH 11/13] Refactor hipHostRegister test. - Add more testing for offsets. - Parse cmdline options and use --tests. --- .../src/runtimeApi/memory/hipHostRegister.cpp | 32 ++++++++++++++----- 1 file changed, 24 insertions(+), 8 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp b/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp index efa23b4068..8cf0979261 100644 --- a/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp +++ b/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp @@ -21,10 +21,11 @@ THE SOFTWARE. * BUILD: %t %s ../../test_common.cpp * RUN: %t --tests 0x1 * RUN: %t --tests 0x2 + * RUN: %t --tests 0x4 * HIT_END */ -// TODO - bug if run both back-to-back +// TODO - bug if run both back-to-back, once fixed should just need one command line #include"test_common.h" #include @@ -36,14 +37,16 @@ __global__ void Inc(hipLaunchParm lp, float *Ad){ template -void doMemCopy(size_t numElements, int offset, T *A, T *Bh, T *Bd) +void doMemCopy(size_t numElements, int offset, T *A, T *Bh, T *Bd, bool internalRegister) { A = A + offset; numElements -= offset; size_t sizeBytes = numElements * sizeof(T); - HIPCHECK(hipHostRegister(A, sizeBytes, 0)); + if (internalRegister) { + HIPCHECK(hipHostRegister(A, sizeBytes, 0)); + } // Reset @@ -67,7 +70,9 @@ void doMemCopy(size_t numElements, int offset, T *A, T *Bh, T *Bd) }; } - HIPCHECK(hipHostUnregister(A)); + if (internalRegister) { + HIPCHECK(hipHostUnregister(A)); + } } @@ -112,7 +117,7 @@ int main(int argc, char *argv[]) } - if (p_tests & 0x2) { + if (p_tests & 0x6) { // Sensitize HIP bug if device does not match where the memory was registered. HIPCHECK(hipSetDevice(0)); @@ -125,11 +130,22 @@ int main(int argc, char *argv[]) Bh = (float*)malloc(size); HIPCHECK(hipMalloc(&Bd, size)); - // TODO - change to 256: + // TODO - set to 128 #define OFFSETS_TO_TRY 1 assert (N>OFFSETS_TO_TRY); - for (size_t i=0; i Date: Mon, 24 Apr 2017 21:22:56 -0500 Subject: [PATCH 12/13] Add host offset checking --- .../tests/src/runtimeApi/memory/hipMemcpy.cpp | 91 +++++++++++++------ 1 file changed, 65 insertions(+), 26 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp index c48f780e44..749ec0de77 100644 --- a/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp +++ b/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp @@ -24,7 +24,8 @@ THE SOFTWARE. * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 * RUN_NAMED: %t hipMemcpy-modes --tests 0x1 * RUN_NAMED: %t hipMemcpy-size --tests 0x6 - * RUN_NAMED: %t hipMemcpy-dev_offsets --tests 0x10 + * RUN_NAMED: %t hipMemcpy-dev-offsets --tests 0x10 + * RUN_NAMED: %t hipMemcpy-host-offsets --tests 0x20 * RUN_NAMED: %t hipMemcpy-multithreaded --tests 0x8 * HIT_END */ @@ -107,28 +108,43 @@ public: void reset(size_t numElements, bool full=false) ; ~HostMemory(); + + T *A_h() const { return _A_h + _offset; }; + T *B_h() const { return _B_h + _offset; }; + T *C_h() const { return _C_h + _offset; }; + + + size_t maxNumElements() const { return _maxNumElements; }; + + void offset(int offset) { _offset = offset; }; + int offset() const { return _offset; }; public: - // Host arrays - T * A_h; - T* B_h; - T* C_h; // Host arrays, secondary copy T * A_hh; T* B_hh; - size_t _maxNumElements; bool _usePinnedHost; +private: + size_t _maxNumElements; + + int _offset; + + // Host arrays + T * _A_h; + T* _B_h; + T* _C_h; }; template HostMemory::HostMemory(size_t numElements, bool usePinnedHost) : _maxNumElements(numElements), - _usePinnedHost(usePinnedHost) + _usePinnedHost(usePinnedHost), + _offset(0) { T ** np = nullptr; - HipTest::initArrays (np, np, np, &A_h, &B_h, &C_h, numElements, usePinnedHost); + HipTest::initArrays (np, np, np, &_A_h, &_B_h, &_C_h, numElements, usePinnedHost); A_hh = NULL; B_hh = NULL; @@ -157,8 +173,8 @@ HostMemory::reset(size_t numElements, bool full) (B_hh)[i] = 1492.0 + i; // Phi if (full) { - (A_h)[i] = 3.146f + i; // Pi - (B_h)[i] = 1.618f + i; // Phi + (_A_h)[i] = 3.146f + i; // Pi + (_B_h)[i] = 1.618f + i; // Phi } } } @@ -166,7 +182,7 @@ HostMemory::reset(size_t numElements, bool full) template HostMemory::~HostMemory () { - HipTest::freeArraysForHost (A_h, B_h, C_h, _usePinnedHost); + HipTest::freeArraysForHost (_A_h, _B_h, _C_h, _usePinnedHost); if (_usePinnedHost) { HIPCHECK (hipHostFree(A_hh)); @@ -197,12 +213,13 @@ template void memcpytest2(DeviceMemory *dmem, HostMemory *hmem, size_t numElements, 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, offsets:%+d\n", + printf ("test: %s<%s> size=%lu (%6.2fMB) usePinnedHost:%d, useHostToHost:%d, useDeviceToDevice:%d, useMemkindDefault:%d, offsets:dev:%+d host:+%d\n", __func__, TYPENAME(T), sizeElements, sizeElements/1024.0/1024.0, hmem->_usePinnedHost, useHostToHost, useDeviceToDevice, useMemkindDefault, - dmem->offset()); + dmem->offset(), hmem->offset() + ); hmem->reset(numElements); @@ -215,15 +232,15 @@ void memcpytest2(DeviceMemory *dmem, HostMemory *hmem, size_t numElements, if (useHostToHost) { // Do some extra host-to-host copies here to mix things up: - HIPCHECK ( hipMemcpy(hmem->A_hh, hmem->A_h, sizeElements, useMemkindDefault? hipMemcpyDefault : hipMemcpyHostToHost)); - HIPCHECK ( hipMemcpy(hmem->B_hh, hmem->B_h, sizeElements, useMemkindDefault? hipMemcpyDefault : hipMemcpyHostToHost)); + HIPCHECK ( hipMemcpy(hmem->A_hh, hmem->A_h(), sizeElements, useMemkindDefault? hipMemcpyDefault : hipMemcpyHostToHost)); + HIPCHECK ( hipMemcpy(hmem->B_hh, hmem->B_h(), sizeElements, useMemkindDefault? hipMemcpyDefault : hipMemcpyHostToHost)); HIPCHECK ( hipMemcpy(dmem->A_d(), hmem->A_hh, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); HIPCHECK ( hipMemcpy(dmem->B_d(), hmem->B_hh, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); } else { - HIPCHECK ( hipMemcpy(dmem->A_d(), hmem->A_h, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); - HIPCHECK ( hipMemcpy(dmem->B_d(), hmem->B_h, sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); + HIPCHECK ( hipMemcpy(dmem->A_d(), hmem->A_h(), sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); + HIPCHECK ( hipMemcpy(dmem->B_d(), hmem->B_h(), sizeElements, useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); } hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, dmem->A_d(), dmem->B_d(), dmem->C_d(), numElements); @@ -235,13 +252,13 @@ void memcpytest2(DeviceMemory *dmem, HostMemory *hmem, size_t numElements, //Destroy the original dmem->C_d(): HIPCHECK ( hipMemset(dmem->C_d(), 0x5A, sizeElements)); - HIPCHECK ( hipMemcpy(hmem->C_h, dmem->C_dd(), sizeElements, useMemkindDefault? hipMemcpyDefault:hipMemcpyDeviceToHost)); + HIPCHECK ( hipMemcpy(hmem->C_h(), dmem->C_dd(), sizeElements, useMemkindDefault? hipMemcpyDefault:hipMemcpyDeviceToHost)); } else { - HIPCHECK ( hipMemcpy(hmem->C_h, dmem->C_d(), sizeElements, useMemkindDefault? hipMemcpyDefault:hipMemcpyDeviceToHost)); + HIPCHECK ( hipMemcpy(hmem->C_h(), dmem->C_d(), sizeElements, useMemkindDefault? hipMemcpyDefault:hipMemcpyDeviceToHost)); } HIPCHECK ( hipDeviceSynchronize() ); - HipTest::checkVectorADD(hmem->A_h, hmem->B_h, hmem->C_h, numElements); + HipTest::checkVectorADD(hmem->A_h(), hmem->B_h(), hmem->C_h(), numElements); @@ -307,7 +324,7 @@ void memcpytest2_sizes(size_t maxElem=0) //--- //Try many different sizes to memory copy. template -void memcpytest2_offsets(size_t maxElem) +void memcpytest2_offsets(size_t maxElem, bool devOffsets, bool hostOffsets) { printSep(); printf ("test: %s<%s>\n", __func__, TYPENAME(T)); @@ -330,14 +347,26 @@ void memcpytest2_offsets(size_t maxElem) for (int offset=0; offset < 512; offset++) { assert (elem + offset < maxElem); - memD.offset(offset); + if (devOffsets) { + memD.offset(offset); + } + if (hostOffsets) { + memU.offset(offset); + memP.offset(offset); + } memcpytest2(&memD, &memU, elem, 1, 1, 0); // unpinned host memcpytest2(&memD, &memP, elem, 1, 1, 0); // pinned host } for (int offset=512; offset < elem; offset*=2) { assert (elem + offset < maxElem); - memD.offset(offset); + if (devOffsets) { + memD.offset(offset); + } + if (hostOffsets) { + memU.offset(offset); + memP.offset(offset); + } memcpytest2(&memD, &memU, elem, 1, 1, 0); // unpinned host memcpytest2(&memD, &memP, elem, 1, 1, 0); // pinned host } @@ -453,9 +482,19 @@ int main(int argc, char *argv[]) printf ("\n\n=== tests&0x10 (test device offsets)\n"); HIPCHECK ( hipDeviceReset() ); size_t maxSize = 256*1024; - memcpytest2_offsets (maxSize); - memcpytest2_offsets (maxSize); - memcpytest2_offsets(maxSize); + memcpytest2_offsets (maxSize, true, false); + memcpytest2_offsets (maxSize, true, false); + memcpytest2_offsets(maxSize, true, false); + } + + + if (p_tests & 0x20) { + printf ("\n\n=== tests&0x10 (test device offsets)\n"); + HIPCHECK ( hipDeviceReset() ); + size_t maxSize = 256*1024; + memcpytest2_offsets (maxSize, false, true); + memcpytest2_offsets (maxSize, false, true); + memcpytest2_offsets(maxSize, false, true); } From 50daa408aa71a546c8b95998542ac81612f9a5d7 Mon Sep 17 00:00:00 2001 From: "Sun, Peng" Date: Tue, 25 Apr 2017 00:13:32 -0500 Subject: [PATCH 13/13] fix hip_complex.h header on NV path Change-Id: Ia95d003ca1b284bab1c76723050e6b3b89178f65 --- hipamd/include/hip/nvcc_detail/hip_complex.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/hipamd/include/hip/nvcc_detail/hip_complex.h b/hipamd/include/hip/nvcc_detail/hip_complex.h index 84afb13e50..20cb24460c 100644 --- a/hipamd/include/hip/nvcc_detail/hip_complex.h +++ b/hipamd/include/hip/nvcc_detail/hip_complex.h @@ -64,7 +64,7 @@ __device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hi } __device__ __host__ static inline float hipCabsf(hipFloatComplex z){ - return cuCabsf(p, q); + return cuCabsf(z); } typedef cuDoubleComplex hipDoubleComplex; @@ -85,7 +85,7 @@ __device__ __host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z){ return cuConj(z); } -__device__ __host__ static inline hipDoubleComplex hipCsqabs(hipDoubleComplex z){ +__device__ __host__ static inline double hipCsqabs(hipDoubleComplex z){ return cuCabs(z) * cuCabs(z); } @@ -123,7 +123,7 @@ __device__ __host__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q return cuCfmaf(p, q, r); } -__device__ __host__ static inline hipDoubleComplex hipCfma(hipComplex p, hipComplex q, hipComplex r){ +__device__ __host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q, hipDoubleComplex r){ return cuCfma(p, q, r); }