Merge branch 'privatestaging' of https://github.com/AMDComputeLibraries/HIP-privatestaging into privatestaging
Этот коммит содержится в:
Ben Sander
+24
-22
@@ -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})
|
||||
|
||||
|
||||
|
||||
@@ -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.
|
||||
|
||||
@@ -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++;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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")
|
||||
|
||||
|
||||
@@ -39,7 +39,8 @@ int device;
|
||||
HIPCHECK(hipGetDevice(&device));
|
||||
HIPCHECK(hipGetDeviceProperties(&prop, device));
|
||||
if(prop.canMapHostMemory != 1){
|
||||
//std::cout<<"Exiting..."<<std::endl;
|
||||
std::cout<<"Exiting..."<<std::endl;
|
||||
failed("Does support HostPinned Memory");
|
||||
}
|
||||
|
||||
HIPCHECK(hipHostMalloc((void**)&A, SIZE, hipHostMallocWriteCombined | hipHostMallocMapped));
|
||||
|
||||
@@ -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<iostream>
|
||||
#include<time.h>
|
||||
|
||||
Ссылка в новой задаче
Block a user