From 119f5a1e67d3b20cd9d46b1dc979970ab09d7ff2 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 [ROCm/clr commit: 85c189c846496f3b6fa331fd49419eebb830a82e] --- projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h | 4 ++-- projects/clr/hipamd/src/hip_memory.cpp | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index f9bfb5a310..80a0db7e2e 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/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/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index da5530349f..f7421f9818 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/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 8d9e22c9752067d0a0bf706ba4310b231d7fd4f3 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 [ROCm/clr commit: 1f532b06f6ce0d38ae0ccae0839c0279437955a3] --- projects/clr/hipamd/src/hip_memory.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index f7421f9818..b706426efb 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/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 05c136e7b8ef105631d07f7a884c6cb74b7510ae 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. [ROCm/clr commit: fb7eee01ff1a46aaa9610d42b2ced80d4493e8b1] --- .../include/hip/hcc_detail/hip_runtime_api.h | 2 +- projects/clr/hipamd/src/hip_memory.cpp | 18 +++++++++++------- 2 files changed, 12 insertions(+), 8 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 80a0db7e2e..7a99ff0810 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/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/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index b706426efb..821f64bc76 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/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 eef8282c967da82f36eaea058dae3a7c4b4d457c 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) [ROCm/clr commit: dfacfbb6412edbf5d5a2ae7c3eff80a900c3dd55] --- projects/clr/hipamd/src/hip_hcc.cpp | 16 ++++++++++------ 1 file changed, 10 insertions(+), 6 deletions(-) diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 35a3e11e71..080d700e63 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/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 88fa7488938a190280ca189cb52c1a9bcbcb662a 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. [ROCm/clr commit: 693e5abc1cdc517c415d2188ccdf2fcbadb647a5] --- .../tests/src/runtimeApi/memory/hipMemoryAllocate.cpp | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemoryAllocate.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemoryAllocate.cpp index 1f7599491a..0a256d6362 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemoryAllocate.cpp +++ b/projects/clr/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 04082c14a7b4e94842ba0c964473ba1cce3ca1fb 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. [ROCm/clr commit: 6b16f56f38a34381f411e56412f2ce90eb091afe] --- projects/clr/hipamd/.vimrc | 1 - .../tests/src/runtimeApi/memory/hipMemcpy.cpp | 239 +++++++++++++----- 2 files changed, 182 insertions(+), 58 deletions(-) delete mode 100644 projects/clr/hipamd/.vimrc diff --git a/projects/clr/hipamd/.vimrc b/projects/clr/hipamd/.vimrc deleted file mode 100644 index 019afa57e6..0000000000 --- a/projects/clr/hipamd/.vimrc +++ /dev/null @@ -1 +0,0 @@ -:set makeprg=make\ -C\ build.hcc-LC.db diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp index a320a86022..d50a810a58 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp +++ b/projects/clr/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 37fcb3bd133f4689195b105f2ee0df61ddc2159e 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 [ROCm/clr commit: b44a3eefd17aaee1c79a77cfbf8d50060e907f86] --- .../tests/src/runtimeApi/memory/hipMemcpy.cpp | 168 ++++++++++++------ 1 file changed, 115 insertions(+), 53 deletions(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp index d50a810a58..ad798d70c1 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp +++ b/projects/clr/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 67383ddb5c537bbcd5ae0aa333db22a4b011b410 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. [ROCm/clr commit: d120b2dd123829057736e34a43681998b21f4d4e] --- .../tests/src/runtimeApi/memory/hipMemcpy.cpp | 21 +++++++++++-------- 1 file changed, 12 insertions(+), 9 deletions(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp index ad798d70c1..c48f780e44 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp +++ b/projects/clr/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 0e042a1529ab3f1900951c9d895f570e3c6ac2cb 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. [ROCm/clr commit: 3da8e94cbff8fb977a7186c09996fbcdb0b1d0fc] --- projects/clr/hipamd/src/hip_hcc.cpp | 66 +++++++++++++++++++++++++++-- 1 file changed, 62 insertions(+), 4 deletions(-) diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 080d700e63..71d947488d 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/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 aa4c89380eef37486ec482345b3ea47737916825 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. [ROCm/clr commit: 5ba167b82b7aed87fb1a3ae9ba25c3bfdc89925a] --- .../src/runtimeApi/memory/hipHostRegister.cpp | 129 ++++++++++++------ 1 file changed, 84 insertions(+), 45 deletions(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp index 1a1319c500..efa23b4068 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp +++ b/projects/clr/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. [ROCm/clr commit: e7af4ef641709ad8c91d4c9ebd701a4bb38d571e] --- .../src/runtimeApi/memory/hipHostRegister.cpp | 32 ++++++++++++++----- 1 file changed, 24 insertions(+), 8 deletions(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp index efa23b4068..8cf0979261 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp +++ b/projects/clr/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 [ROCm/clr commit: fbf28a87281041dfbe2d6e27877dae0c51b2ae7a] --- .../tests/src/runtimeApi/memory/hipMemcpy.cpp | 91 +++++++++++++------ 1 file changed, 65 insertions(+), 26 deletions(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp index c48f780e44..749ec0de77 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpy.cpp +++ b/projects/clr/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 a82ecbcc131383f518f8153040caf60af73a7c67 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 [ROCm/clr commit: 50daa408aa71a546c8b95998542ac81612f9a5d7] --- projects/clr/hipamd/include/hip/nvcc_detail/hip_complex.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_complex.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_complex.h index 84afb13e50..20cb24460c 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_complex.h +++ b/projects/clr/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); }