From 95c96638aa48edf173f201a06d750037ea50ebe9 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Thu, 24 Mar 2016 07:33:24 -0500 Subject: [PATCH] Added canMapMemory feature to device properties --- include/hcc_detail/hip_runtime_api.h | 36 +++++++++++++++++----------- src/hip_hcc.cpp | 7 +++++- tests/src/hipHostAlloc.cpp | 3 ++- 3 files changed, 30 insertions(+), 16 deletions(-) diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index 0ced74a084..92aef3d213 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/include/hcc_detail/hip_runtime_api.h @@ -47,28 +47,36 @@ extern "C" { * @{ */ //! Flags that can be used with hipStreamCreateWithFlags -#define hipStreamDefault 0x00 ///< Default stream creation flags. These are used with hipStreamCreate(). -#define hipStreamNonBlocking 0x01 ///< Stream does not implicitly synchronize with null stream +#define hipStreamDefault 0x00 ///< Default stream creation flags. These are used with hipStreamCreate(). +#define hipStreamNonBlocking 0x01 ///< Stream does not implicitly synchronize with null stream //! Flags that can be used with hipEventCreateWithFlags: -#define hipEventDefault 0x0 ///< Default flags -#define hipEventBlockingSync 0x1 ///< Waiting will yield CPU. Power-friendly and usage-friendly but may increase latency. -#define hipEventDisableTiming 0x2 ///< Disable event's capability to record timing information. May improve performance. -#define hipEventInterprocess 0x4 ///< Event can support IPC. @warning - not supported in HIP. +#define hipEventDefault 0x0 ///< Default flags +#define hipEventBlockingSync 0x1 ///< Waiting will yield CPU. Power-friendly and usage-friendly but may increase latency. +#define hipEventDisableTiming 0x2 ///< Disable event's capability to record timing information. May improve performance. +#define hipEventInterprocess 0x4 ///< Event can support IPC. @warning - not supported in HIP. //! Flags that can be used with hipHostMalloc -#define hipHostMallocDefault 0x0 -#define hipHostMallocPortable 0x1 -#define hipHostMallocMapped 0x2 -#define hipHostMallocWriteCombined 0x4 +#define hipHostMallocDefault 0x0 +#define hipHostMallocPortable 0x1 +#define hipHostMallocMapped 0x2 +#define hipHostMallocWriteCombined 0x4 //! Flags that can be used with hipHostRegister -#define hipHostRegisterDefault 0x0 ///< Memory is Mapped and Portable -#define hipHostRegisterPortable 0x1 ///< Memory is considered registered by all contexts. HIP only supports one context so this is always assumed true. -#define hipHostRegisterMapped 0x2 ///< Map the allocation into the address space for the current device. The device pointer can be obtained with #hipHostGetDevicePointer. -#define hipHostRegisterIoMemory 0x4 ///< Not supported. +#define hipHostRegisterDefault 0x0 ///< Memory is Mapped and Portable +#define hipHostRegisterPortable 0x1 ///< Memory is considered registered by all contexts. HIP only supports one context so this is always assumed true. +#define hipHostRegisterMapped 0x2 ///< Map the allocation into the address space for the current device. The device pointer can be obtained with #hipHostGetDevicePointer. +#define hipHostRegisterIoMemory 0x4 ///< Not supported. + + +#define hipDeviceScheduleAuto 0x0 +#define hipDeviceScheduleSpin 0x1 +#define hipDeviceScheduleYield 0x2 +#define hipDeviceBlockingSync 0x4 +#define hipDeviceMapHost 0x8 +#define hipDeviceLmemResizeToMax 0x16 /** * @warning On AMD devices and recent Nvidia devices, these hints and controls are ignored. diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index ff2876ffb3..0631cc6814 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -658,6 +658,11 @@ hipError_t ihipDevice_t::getProperties(hipDeviceProp_t* prop) prop->arch.hasDynamicParallelism = 0; prop->concurrentKernels = 1; // All ROCR hardware supports executing multiple kernels concurrently + if ( _device_flags | hipDeviceMapHost) { + prop->canMapHostMemory = 1; + } else { + prop->canMapHostMemory = 0; + } return e; } @@ -859,7 +864,7 @@ void ihipInit() //If device is not in visible devices list, ignore continue; } - g_devices[g_deviceCnt].init(g_deviceCnt, accs[i], 0); + g_devices[g_deviceCnt].init(g_deviceCnt, accs[i], hipDeviceMapHost); g_deviceCnt++; } } diff --git a/tests/src/hipHostAlloc.cpp b/tests/src/hipHostAlloc.cpp index 2c72c3e33a..a284d801c6 100644 --- a/tests/src/hipHostAlloc.cpp +++ b/tests/src/hipHostAlloc.cpp @@ -39,7 +39,8 @@ int device; HIPCHECK(hipGetDevice(&device)); HIPCHECK(hipGetDeviceProperties(&prop, device)); if(prop.canMapHostMemory != 1){ -//std::cout<<"Exiting..."<