From 95c96638aa48edf173f201a06d750037ea50ebe9 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Thu, 24 Mar 2016 07:33:24 -0500 Subject: [PATCH 1/3] 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..."< Date: Thu, 24 Mar 2016 11:17:57 -0500 Subject: [PATCH 2/3] corrected cmake --- CMakeLists.txt | 46 ++++++++++++++++++++------------------ tests/src/CMakeLists.txt | 7 +++--- tests/src/hipHostAlloc.cpp | 2 +- 3 files changed, 28 insertions(+), 27 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 5aedd12d77..1011afd7a0 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -2,32 +2,34 @@ cmake_minimum_required(VERSION 2.6) project(hip_hcc) # preserve HCC_HOME env in the generated -set (HCC_HOME $ENV{HCC_HOME}) -if (NOT DEFINED HCC_HOME) - set (HCC_HOME /opt/hcc) + +if (NOT DEFINED HCC_DIR) + set (HCC_DIR "/opt/hcc") endif() +message(STATUS ${HCC_DIR}) -MESSAGE ("HCC_HOME=" ${HCC_HOME}) - -set(HSA_PATH "/opt/hsa") - - -set (HIP_INSTALL_PATH $ENV{HIP_PATH}) -if (NOT DEFINED HIP_INSTALL_PATH) - set (HIP_INSTALLPATH /opt/hip) +if(NOT DEFINED HIP_INSTALL_DIR) + set(HIP_INSTALL_DIR "/opt/hip") endif() +message(STATUS ${HIP_INSTALL_DIR}) -if (CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT) - set(CMAKE_INSTALL_PREFIX "${HIP_INSTALL_PATH}" CACHE PATH "Default installation path of hip" FORCE) -endif () +if(NOT DEFINED HSA_DIR) + set(HSA_DIR "/opt/hsa") +endif() +message(STATUS ${HSA_DIR}) + + +#if (CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT) +# set(CMAKE_INSTALL_PREFIX "${HIP_INSTALL_PATH}" CACHE PATH "Default installation path of hip" FORCE) +#endif () include_directories(${PROJECT_SOURCE_DIR}/include) -set(CMAKE_CXX_COMPILER "${HCC_HOME}/bin/hcc") -set(CMAKE_C_COMPILER "${HCC_HOME}/bin/hcc") +set(CMAKE_CXX_COMPILER "${HCC_DIR}/bin/hcc") +set(CMAKE_C_COMPILER "${HCC_DIR}/bin/hcc") -set(CMAKE_CXX_FLAGS " -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I${HIP_PATH}/include -stdlib=libc++ ") -set(CMAKE_C_FLAGS " -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I${HIP_PATH}/include -stdlib=libc++ ") +set(CMAKE_CXX_FLAGS " -hc -I${HCC_DIR}/include -I${HSA_DIR}/include -stdlib=libc++ ") +set(CMAKE_C_FLAGS " -hc -I${HCC_DIR}/include -I${HSA_DIR}/include -stdlib=libc++ ") set(SOURCE_FILES src/hip_hcc.cpp src/hip_device.cpp @@ -40,9 +42,9 @@ src/staging_buffer.cpp) add_library(hip_hcc STATIC ${SOURCE_FILES}) -install(TARGETS hip_hcc DESTINATION lib) -install(DIRECTORY src DESTINATION .) -install(DIRECTORY bin DESTINATION . USE_SOURCE_PERMISSIONS) -install(DIRECTORY include DESTINATION .) +install(TARGETS hip_hcc DESTINATION ${HIP_INSTALL_DIR}/lib) +install(DIRECTORY src DESTINATION ${HIP_INSTALL_DIR}) +install(DIRECTORY bin DESTINATION ${HIP_INSTALL_DIR} USE_SOURCE_PERMISSIONS) +install(DIRECTORY include DESTINATION ${HIP_INSTALL_DIR}) diff --git a/tests/src/CMakeLists.txt b/tests/src/CMakeLists.txt index b470956254..3846c1c287 100644 --- a/tests/src/CMakeLists.txt +++ b/tests/src/CMakeLists.txt @@ -8,7 +8,8 @@ include_directories( ${PROJECT_SOURCE_DIR}/include ) set (HIP_Unit_Test_VERSION_MAJOR 1) set (HIP_Unit_Test_VERSION_MINOR 0) -set (HIP_PATH $ENV{HIP_PATH}) +set(HIP_PATH $ENV{HIP_PATH}) +MESSAGE("HIP_PATH=" ${HIP_PATH}) if (NOT DEFINED HIP_PATH) set (HIP_PATH ../..) endif() @@ -44,9 +45,7 @@ if (${HIP_PLATFORM} STREQUAL "hcc") # This will create a subdir "hip_hcc" in the test build directory # Any changes to hip_hcc source will be detected and force the library and then the tests to be rebuilt. - add_subdirectory(${HIP_PATH} hip_hcc) - link_directories(${CMAKE_CURRENT_BINARY_DIR}/hip_hcc) # search the local hip_hcc for libhip_hcc.a - + elseif (${HIP_PLATFORM} STREQUAL "nvcc") MESSAGE ("HIP_PLATFORM=nvcc") diff --git a/tests/src/hipHostAlloc.cpp b/tests/src/hipHostAlloc.cpp index a284d801c6..01ca04b311 100644 --- a/tests/src/hipHostAlloc.cpp +++ b/tests/src/hipHostAlloc.cpp @@ -40,7 +40,7 @@ HIPCHECK(hipGetDevice(&device)); HIPCHECK(hipGetDeviceProperties(&prop, device)); if(prop.canMapHostMemory != 1){ std::cout<<"Exiting..."< Date: Thu, 24 Mar 2016 21:48:27 -0500 Subject: [PATCH 3/3] updated CR --- tests/src/hipPerfMemcpy.cpp | 19 +++++++++++++++++++ 1 file changed, 19 insertions(+) diff --git a/tests/src/hipPerfMemcpy.cpp b/tests/src/hipPerfMemcpy.cpp index d351dbc868..b528ed0e10 100644 --- a/tests/src/hipPerfMemcpy.cpp +++ b/tests/src/hipPerfMemcpy.cpp @@ -1,3 +1,22 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + #include"test_common.h" #include #include