From 2d0fade1f73f938675924aa5ddee969a282c4d03 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Tue, 22 Mar 2016 02:30:10 -0500 Subject: [PATCH] hipHostRegister and hipHostMalloc refactor. Note hipHostMalloc (not hipHostAlloc or hipMallocHost). - the hipHost* is used for all HIP APIs dealing with Host memory. (including hipHostMalloc, hipHostFree, hipHostRegister, hipHostUnregister, hipHostGetFlags, hipHostGetDevicePointer). - hipMallocHost is consistent with "hipMalloc" for allocating device memory. Enumerations hipHostMalloc* also used as optional flags parm to hipHostMalloc. --- hipamd/RELEASE.md | 2 +- hipamd/bin/hipify | 12 +++++------ hipamd/include/hcc_detail/hip_runtime_api.h | 17 ++++++++-------- hipamd/include/hip_runtime_api.h | 4 ++-- hipamd/include/nvcc_detail/hip_runtime_api.h | 4 ++-- .../hipBusBandwidth/hipBusBandwidth.cpp | 20 +++++++++---------- hipamd/src/hip_hcc.cpp | 14 ++++++++++--- hipamd/tests/src/hipHostAlloc.cpp | 6 +++--- hipamd/tests/src/hipHostGetFlags.cpp | 14 ++++++------- hipamd/tests/src/hipMemcpy.cpp | 6 +++--- hipamd/tests/src/hipMemcpyAsync.cpp | 8 ++++---- hipamd/tests/src/hipMemcpy_simple.cpp | 4 ++-- hipamd/tests/src/hipMultiThreadStreams2.cpp | 18 ++++++++--------- hipamd/tests/src/hipPointerAttrib.cpp | 6 +++--- hipamd/tests/src/hipStream.h | 4 ++-- hipamd/tests/src/test_common.h | 8 ++++---- hipamd/util/vim/hip.vim | 18 ++++++++++++++++- 17 files changed, 95 insertions(+), 70 deletions(-) diff --git a/hipamd/RELEASE.md b/hipamd/RELEASE.md index 04102ed0c1..d22452b954 100644 --- a/hipamd/RELEASE.md +++ b/hipamd/RELEASE.md @@ -15,7 +15,7 @@ Stay tuned - the work for many of these features is already in-flight. Next: - Deprecate hipDeviceGetProp, replace with hipGetDeviceProp -- Deprecate hipMallocHost (Replace with hipHostAlloc with hipHostAllocDefault as last parameter). +- Deprecate hipMallocHost (Replace with hipHostMalloc) - Deprecate hipFreeHost (Replace with hipHostFree). diff --git a/hipamd/bin/hipify b/hipamd/bin/hipify index 8ff04030fd..5c135f653a 100755 --- a/hipamd/bin/hipify +++ b/hipamd/bin/hipify @@ -283,15 +283,15 @@ while (@ARGV) { #-------- # Memory management: $ft{'mem'} += s/\bcudaMalloc\b/hipMalloc/g; - $ft{'mem'} += s/\bcudaMallocHost\b/hipMallocHost/g; + $ft{'mem'} += s/\bcudaMallocHost\b/hipHostMalloc/g; # note conversion to standard hipHost* naming convention $ft{'mem'} += s/\bcudaFree\b/hipFree/g; - $ft{'mem'} += s/\bcudaFreeHost\b/hipFreeHost/g; + $ft{'mem'} += s/\bcudaFreeHost\b/hipHostFree/g; # note conversion to standard hipHost* naming convention $ft{'mem'} += s/\bcudaHostAlloc\b/hipHostAlloc/g; $ft{'mem'} += s/\bcudaHostGetDevicePointer\b/hipHostGetDevicePointer/g; - $ft{'mem'} += s/\bcudaHostAllocDefault\b/hipHostAllocDefault/g; - $ft{'mem'} += s/\bcudaHostAllocPortable\b/hipHostAllocPortable/g; - $ft{'mem'} += s/\bcudaHostAllocMapped\b/hipHostAllocMapped/g; - $ft{'mem'} += s/\bcudaHostAllocWriteCombined\b/hipHostAllocWriteCombined/g; + $ft{'mem'} += s/\bcudaHostAllocDefault\b/hipHostMallocDefault/g; + $ft{'mem'} += s/\bcudaHostAllocPortable\b/hipHostMallocPortable/g; + $ft{'mem'} += s/\bcudaHostAllocMapped\b/hipHostMallocMapped/g; + $ft{'mem'} += s/\bcudaHostAllocWriteCombined\b/hipHostMallocWriteCombined/g; #-------- diff --git a/hipamd/include/hcc_detail/hip_runtime_api.h b/hipamd/include/hcc_detail/hip_runtime_api.h index 00e28788b7..ef5ad4927c 100644 --- a/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hcc_detail/hip_runtime_api.h @@ -56,10 +56,10 @@ extern "C" { #define hipEventInterprocess 0x4 ///< Event can support IPC. @warning - not supported in HIP. -#define hipHostAllocDefault 0x0 -#define hipHostAllocPortable 0x1 -#define hipHostAllocMapped 0x2 -#define hipHostAllocWriteCombined 0x4 +#define hipHostMallocDefault 0x0 +#define hipHostMallocPortable 0x1 +#define hipHostMallocMapped 0x2 +#define hipHostMallocWriteCombined 0x4 #define hipHostRegisterDefault 0x0 #define hipHostRegisterPortable 0x1 @@ -670,7 +670,7 @@ hipError_t hipMalloc(void** ptr, size_t size) ; * @param[in] size Requested memory size * @return Error code */ -hipError_t hipMallocHost(void** ptr, size_t size) __attribute__((deprecated("use hipHostAlloc instead"))) ; +hipError_t hipMallocHost(void** ptr, size_t size) __attribute__((deprecated("use hipHostMalloc instead"))) ; /** * @brief Allocate device accessible page locked host memory @@ -680,7 +680,8 @@ hipError_t hipMallocHost(void** ptr, size_t size) __attribute__((deprecated("use * @param[in] flags Type of host memory allocation * @return Error code */ -hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags) ; +hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags) ; +hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags) __attribute__((deprecated("use hipHostMalloc instead"))) ;; /** * @brief Get Device pointer from Host Pointer allocated through hipHostAlloc @@ -696,7 +697,7 @@ hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, size_t size) ; * @brief Get flags associated with host pointer * * @param[out] flagsPtr Memory location to store flags - * @param[in] hostPtr Host Pointer allocated through hipHostAlloc + * @param[in] hostPtr Host Pointer allocated through hipHostMalloc * @return Error code */ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) ; @@ -786,7 +787,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t siz /** * @brief Copy data from src to dst asynchronously. * - * @warning If host or dest are not pinned, the memory copy will be performed synchronously. For best performance, use hipHostAlloc to + * @warning If host or dest are not pinned, the memory copy will be performed synchronously. For best performance, use hipHostMalloc to * allocate host memory that is transferred asynchronously. * * @param[out] dst Data being copy to diff --git a/hipamd/include/hip_runtime_api.h b/hipamd/include/hip_runtime_api.h index 9d6aa6d3ca..c7d9f3fa8b 100644 --- a/hipamd/include/hip_runtime_api.h +++ b/hipamd/include/hip_runtime_api.h @@ -220,8 +220,8 @@ static inline hipError_t hipMalloc ( T** devPtr, size_t size) // Provide an override to automatically typecast the pointer type from void**, and also provide a default for the flags. template -static inline hipError_t hipHostAlloc( T** ptr, size_t size, unsigned int flags = hipHostAllocDefault) +static inline hipError_t hipHostMalloc( T** ptr, size_t size, unsigned int flags = hipHostMallocDefault) { - return hipHostAlloc((void**)ptr, size, flags); + return hipHostMalloc((void**)ptr, size, flags); } #endif diff --git a/hipamd/include/nvcc_detail/hip_runtime_api.h b/hipamd/include/nvcc_detail/hip_runtime_api.h index 7c37e7f71d..edf2568687 100644 --- a/hipamd/include/nvcc_detail/hip_runtime_api.h +++ b/hipamd/include/nvcc_detail/hip_runtime_api.h @@ -119,13 +119,13 @@ inline static hipError_t hipFree(void* ptr) { return hipCUDAErrorTohipError(cudaFree(ptr)); } -inline static hipError_t hipMallocHost(void** ptr, size_t size) __attribute__((deprecated("use hipHostAlloc instead"))); +inline static hipError_t hipMallocHost(void** ptr, size_t size) __attribute__((deprecated("use hipHostMalloc instead"))); inline static hipError_t hipMallocHost(void** ptr, size_t size) { return hipCUDAErrorTohipError(cudaMallocHost(ptr, size)); } -inline static hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags){ +inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags){ return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags)); } diff --git a/hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index 25b14305dc..2230a8eef7 100644 --- a/hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -108,7 +108,7 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) float *hostMem = NULL; if (p_pinned) { - hipHostAlloc((void**)&hostMem, sizeof(float) * numMaxFloats, hipHostAllocDefault); + hipHostMalloc((void**)&hostMem, sizeof(float) * numMaxFloats); while (hipGetLastError() != hipSuccess) { // drop the size and try again @@ -120,7 +120,7 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) return; } numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; - hipHostAlloc((void**)&hostMem, sizeof(float) * numMaxFloats, hipHostAllocDefault); + hipHostMalloc((void**)&hostMem, sizeof(float) * numMaxFloats); } } else @@ -232,7 +232,7 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) CHECK_HIP_ERROR(); if (p_pinned) { - hipFreeHost((void*)hostMem); + hipHostFree((void*)hostMem); CHECK_HIP_ERROR(); } else @@ -259,15 +259,15 @@ void RunBenchmark_D2H(ResultDatabase &resultDB) float *hostMem2; if (p_pinned) { - hipHostAlloc((void**)&hostMem1, sizeof(float)*numMaxFloats, hipHostAllocDefault); + hipHostMalloc((void**)&hostMem1, sizeof(float)*numMaxFloats); hipError_t err1 = hipGetLastError(); - hipHostAlloc((void**)&hostMem2, sizeof(float)*numMaxFloats, hipHostAllocDefault); + hipHostMalloc((void**)&hostMem2, sizeof(float)*numMaxFloats); hipError_t err2 = hipGetLastError(); while (err1 != hipSuccess || err2 != hipSuccess) { // free the first buffer if only the second failed if (err1 == hipSuccess) - hipFreeHost((void*)hostMem1); + hipHostFree((void*)hostMem1); // drop the size and try again if (p_verbose) std::cout << " - dropping size allocating pinned mem\n"; @@ -383,9 +383,9 @@ void RunBenchmark_D2H(ResultDatabase &resultDB) CHECK_HIP_ERROR(); if (p_pinned) { - hipFreeHost((void*)hostMem1); + hipHostFree((void*)hostMem1); CHECK_HIP_ERROR(); - hipFreeHost((void*)hostMem2); + hipHostFree((void*)hostMem2); CHECK_HIP_ERROR(); } else @@ -523,8 +523,8 @@ void RunBenchmark_Bidir(ResultDatabase &resultDB) CHECK_HIP_ERROR(); if (p_pinned) { - hipFreeHost((void*)hostMem[0]); - hipFreeHost((void*)hostMem[1]); + hipHostFree((void*)hostMem[0]); + hipHostFree((void*)hostMem[1]); CHECK_HIP_ERROR(); } else diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index 91204147e2..7aa7ae4482 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -2085,7 +2085,7 @@ hipError_t hipMallocHost(void** ptr, size_t sizeBytes) } -hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags){ +hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags){ std::call_once(hip_initialized, ihipInit); hipError_t hip_status = hipSuccess; @@ -2093,7 +2093,7 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags){ auto device = ihipGetTlsDefaultDevice(); if(device){ - if(flags == hipHostAllocDefault){ + if(flags == hipHostMallocDefault){ *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned); if(sizeBytes && (*ptr == NULL)){ hip_status = hipErrorMemoryAllocation; @@ -2101,7 +2101,7 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags){ hc::am_memtracker_update(*ptr, device->_device_index, 0); } tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr); - } else if(flags & hipHostAllocMapped){ + } else if(flags & hipHostMallocMapped){ *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned); if(sizeBytes && (*ptr == NULL)){ hip_status = hipErrorMemoryAllocation; @@ -2115,6 +2115,14 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags){ } +// TODO - remove me, this is deprecated. +hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) +{ + return hipHostMalloc(ptr, sizeBytes, flags); +}; + + + hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, size_t size){ std::call_once(hip_initialized, ihipInit); diff --git a/hipamd/tests/src/hipHostAlloc.cpp b/hipamd/tests/src/hipHostAlloc.cpp index fe82a0c78f..2c72c3e33a 100644 --- a/hipamd/tests/src/hipHostAlloc.cpp +++ b/hipamd/tests/src/hipHostAlloc.cpp @@ -42,9 +42,9 @@ if(prop.canMapHostMemory != 1){ //std::cout<<"Exiting..."< static void *Alloc(size_t sizeBytes) { void *p; - HIPCHECK(hipHostAlloc((void**)&p, sizeBytes, hipHostAllocDefault)); + HIPCHECK(hipHostMalloc((void**)&p, sizeBytes, hipHostMallocDefault)); return p; }; }; @@ -181,8 +181,8 @@ void test_manyInflightCopies(hipStream_t stream, int numElements, int numCopies, T *A_d; T *A_h1, *A_h2; - HIPCHECK(hipHostAlloc((void**)&A_h1, Nbytes, hipHostAllocDefault)); - HIPCHECK(hipHostAlloc((void**)&A_h2, Nbytes, hipHostAllocDefault)); + HIPCHECK(hipHostMalloc((void**)&A_h1, Nbytes, hipHostMallocDefault)); + HIPCHECK(hipHostMalloc((void**)&A_h2, Nbytes, hipHostMallocDefault)); HIPCHECK(hipMalloc(&A_d, Nbytes)); for (int i=0; i