From 91dbc3114d8b43bbaca344e4818009da62a2e8bf Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Sun, 6 Mar 2016 12:17:30 -0600 Subject: [PATCH] added feature for hipHostGetFlags for CUDA and HIP [ROCm/hip commit: 75952029d6197181fecdcfe6ae03697b4165fb9f] --- .../hip/include/hcc_detail/hip_runtime_api.h | 9 +++ .../hip/include/nvcc_detail/hip_runtime_api.h | 4 ++ projects/hip/src/hip_hcc.cpp | 35 +++++++++-- projects/hip/tests/src/CMakeLists.txt | 3 +- projects/hip/tests/src/hipHostGetFlags.cpp | 59 +++++++++++++++++++ 5 files changed, 103 insertions(+), 7 deletions(-) create mode 100644 projects/hip/tests/src/hipHostGetFlags.cpp diff --git a/projects/hip/include/hcc_detail/hip_runtime_api.h b/projects/hip/include/hcc_detail/hip_runtime_api.h index 1588f73a5a..0d83dedd64 100644 --- a/projects/hip/include/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hcc_detail/hip_runtime_api.h @@ -687,6 +687,15 @@ hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags) ; */ 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 + * @return Error code + */ +hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) ; + /** * @brief Free memory allocated by the hcc hip memory allocation API. * This API performs an implicit hipDeviceSynchronize() call. diff --git a/projects/hip/include/nvcc_detail/hip_runtime_api.h b/projects/hip/include/nvcc_detail/hip_runtime_api.h index 98e7ca004c..783750d44d 100644 --- a/projects/hip/include/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/nvcc_detail/hip_runtime_api.h @@ -128,6 +128,10 @@ inline static hipError_t hipHostGetDevicePointer(void** devPtr, void* hostPtr, u return hipCUDAErrorTohipError(cudaHostGetDevicePointer(devPtr, hostPtr, flags)); } +inline static hipError_ hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr){ + return hipCUDAErrorTohipError(cudaHostGetFlags(flagsPtr, hostPtr)); +} + inline static hipError_t hipFreeHost(void* ptr) { return hipCUDAErrorTohipError(cudaFreeHost(ptr)); } diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 8f12dd14c6..a86616a14d 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -2008,7 +2008,7 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags){ auto device = ihipGetTlsDefaultDevice(); if(device){ - if(flags | hipHostAllocDefault){ + if(flags & hipHostAllocDefault){ const unsigned am_flags = amHostPinned; *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags); @@ -2021,7 +2021,7 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags){ } tprintf(TRACE_MEM, " %s: pinned ptr=%p\n", __func__, *ptr); } - if(flags | hipHostAllocMapped && device->_props.canMapHostMemory == 1){ + if(flags & hipHostAllocMapped){ const unsigned am_flags = amHostPinned; *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags); @@ -2030,10 +2030,10 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags){ }else{ #if USE_AM_TRACKER hc::am_memtracker_update(*ptr, device->_device_index, flags); - void *srcPtr; - hsa_status_t hsa_status = hsa_amd_memory_lock((*ptr), sizeBytes, &device->_hsa_agent, 1, &srcPtr); - assert(hsa_status == HSA_STATUS_SUCCESS); - hc::am_memtracker_add(srcPtr, sizeBytes, device->_acc, false); +// void *srcPtr; +// hsa_status_t hsa_status = hsa_amd_memory_lock((*ptr), sizeBytes, &device->_hsa_agent, 1, &srcPtr); +// assert(hsa_status == HSA_STATUS_SUCCESS); +// hc::am_memtracker_add(srcPtr, sizeBytes, device->_acc, false); #endif } tprintf(TRACE_MEM, " %s: pinned ptr=%p\n", __func__, *ptr); @@ -2067,6 +2067,29 @@ hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, size_t size){ return ihipLogStatus(hip_status); } +hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) +{ + std::call_once(hip_initialized, ihipInit); + hipError_t hip_status = hipSuccess; + + hc::accelerator acc; + hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); + am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr); + if(status == AM_SUCCESS){ + *flagsPtr = amPointerInfo._appAllocationFlags; + if(*flagsPtr == 0){ + hip_status = hipErrorInvalidValue; + } + else{ + hip_status = hipSuccess; + } + tprintf(TRACE_MEM, " %s: host ptr=%p\n", __func__, hostPtr); + }else{ + hip_status = hipErrorInvalidValue; + } + return ihipLogStatus(hip_status); +} + //--- hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind) { diff --git a/projects/hip/tests/src/CMakeLists.txt b/projects/hip/tests/src/CMakeLists.txt index 4454ac73fb..0b69839e72 100644 --- a/projects/hip/tests/src/CMakeLists.txt +++ b/projects/hip/tests/src/CMakeLists.txt @@ -128,6 +128,7 @@ make_hip_executable (hipMultiThreadStreams1 hipMultiThreadStreams1.cpp) make_hip_executable (hipMultiThreadStreams2 hipMultiThreadStreams2.cpp) make_hip_executable (hipHostAlloc hipHostAlloc.cpp) make_hip_executable (hipStreamL5 hipStreamL5.cpp) +make_hip_executable (hipHostGetFlags hipHostGetFlags.cpp) target_link_libraries(hipMathFunctionsHost m) make_test(hip_ballot " " ) @@ -149,7 +150,7 @@ make_test(hipMultiThreadStreams2 " " ) make_test(hipHostAlloc " ") make_test(hipMemcpy " " ) make_test(hipMemcpyAsync " " ) - +make_test(hipHostGetFlags " ") make_test(hipHcc " " ) make_test(hipStreamL5 " ") diff --git a/projects/hip/tests/src/hipHostGetFlags.cpp b/projects/hip/tests/src/hipHostGetFlags.cpp new file mode 100644 index 0000000000..648c1f2343 --- /dev/null +++ b/projects/hip/tests/src/hipHostGetFlags.cpp @@ -0,0 +1,59 @@ +#include"test_common.h" +#include + +#define LEN 1024*1024 +#define SIZE LEN*sizeof(float) + +__global__ void Add(hipLaunchParm lp, float *Ad, float *Bd, float *Cd){ +int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; +Cd[tx] = Ad[tx] + Bd[tx]; +} + +int main(){ +float *A, *B, *C, *D; +float *Ad, *Bd, *Cd, *Dd; +unsigned int FlagA, FlagB, FlagC; +FlagA = hipHostAllocWriteCombined | hipHostAllocMapped; +FlagB = hipHostAllocWriteCombined | hipHostAllocMapped; +FlagC = hipHostAllocMapped; +hipDeviceProp_t prop; +int device; +HIPCHECK(hipGetDevice(&device)); +HIPCHECK(hipGetDeviceProperties(&prop, device)); +if(prop.canMapHostMemory != 1){ +std::cout<<"Exiting..."<