diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt index 1011afd7a0..2f11a2642f 100644 --- a/projects/clr/hipamd/CMakeLists.txt +++ b/projects/clr/hipamd/CMakeLists.txt @@ -1,27 +1,20 @@ cmake_minimum_required(VERSION 2.6) project(hip_hcc) -# preserve HCC_HOME env in the generated +set (HCC_DIR "/opt/hcc" CACHE PATH "Path to which HCC has been installed") +message(STATUS "Looking for HCC in: " ${HCC_DIR}) -if (NOT DEFINED HCC_DIR) - set (HCC_DIR "/opt/hcc") +set(HSA_DIR "/opt/hsa" CACHE PATH "Path to which HSA runtime has been installed") +message(STATUS "Looking for HSA runtime in: " ${HSA_DIR}) + +if(CMAKE_BUILD_TYPE MATCHES Debug) + set(HIP_INSTALL_DIR ${CMAKE_SOURCE_DIR} CACHE PATH "Installation path for HIP") +else() + set(HIP_INSTALL_DIR "/opt/hip" CACHE PATH "Installation path for HIP") endif() -message(STATUS ${HCC_DIR}) +message(STATUS "HIP will be installed in: " ${HIP_INSTALL_DIR}) -if(NOT DEFINED HIP_INSTALL_DIR) - set(HIP_INSTALL_DIR "/opt/hip") -endif() -message(STATUS ${HIP_INSTALL_DIR}) - -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 () +set(CMAKE_INSTALL_PREFIX "${HIP_INSTALL_PATH}" CACHE INTERNAL "Installation path for HIP" FORCE) include_directories(${PROJECT_SOURCE_DIR}/include) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index ee510efe47..df287cf437 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -120,7 +120,7 @@ class ihipDevice_t; // #include CPP files to produce one object file -#define ONE_OBJECT_FILE 1 +#define ONE_OBJECT_FILE 0 // Compile support for trace markers that are displayed on CodeXL GUI at start/stop of each function boundary. diff --git a/projects/clr/hipamd/include/hcc_detail/trace_helper.h b/projects/clr/hipamd/include/hcc_detail/trace_helper.h index db7cc07073..8ba877f7ec 100644 --- a/projects/clr/hipamd/include/hcc_detail/trace_helper.h +++ b/projects/clr/hipamd/include/hcc_detail/trace_helper.h @@ -16,7 +16,7 @@ LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ - +#pragma once #include #include @@ -35,7 +35,7 @@ THE SOFTWARE. // Building block functions: template -std::string ToHexString(T v) +inline std::string ToHexString(T v) { std::ostringstream ss; ss << "0x" << std::hex << v; @@ -48,7 +48,7 @@ std::string ToHexString(T v) // This is the default which works for most types: template -std::string ToString(T v) +inline std::string ToString(T v) { std::ostringstream ss; ss << v; @@ -77,7 +77,8 @@ std::string ToString(hipStream_t v) // hipMemcpyKind specialization template <> -std::string ToString(hipMemcpyKind v) { +inline std::string ToString(hipMemcpyKind v) +{ switch(v) { CASE_STR(hipMemcpyHostToHost); CASE_STR(hipMemcpyHostToDevice); @@ -90,13 +91,15 @@ std::string ToString(hipMemcpyKind v) { template <> -std::string ToString(hipError_t v) { +inline std::string ToString(hipError_t v) +{ return ihipErrorString(v); }; // Catch empty arguments case -std::string ToString() { +inline std::string ToString() +{ return (""); } @@ -105,6 +108,7 @@ std::string ToString() { // C++11 variadic template - peels off first argument, converts to string, and calls itself again to peel the next arg. // Strings are automatically separated by comma+space. template -std::string ToString(T first, Args... args) { +inline std::string ToString(T first, Args... args) +{ return ToString(first) + ", " + ToString(args...) ; } diff --git a/projects/clr/hipamd/tests/src/CMakeLists.txt b/projects/clr/hipamd/tests/src/CMakeLists.txt index 4288624ccf..84a19815be 100644 --- a/projects/clr/hipamd/tests/src/CMakeLists.txt +++ b/projects/clr/hipamd/tests/src/CMakeLists.txt @@ -62,7 +62,7 @@ endif() set (HIPCC ${HIP_PATH}/bin/hipcc) set (CMAKE_CXX_COMPILER ${HIPCC}) -set (CMAKE_CXX_FLAGS --hipcc_explicit_lib) +#set (CMAKE_CXX_FLAGS --hipcc_explicit_lib) add_library(test_common OBJECT test_common.cpp ) diff --git a/projects/clr/hipamd/tests/src/buildHIPC.sh b/projects/clr/hipamd/tests/src/buildHIPC.sh new file mode 100755 index 0000000000..b2c9ce2a07 --- /dev/null +++ b/projects/clr/hipamd/tests/src/buildHIPC.sh @@ -0,0 +1,3 @@ +#!/bin/bash + +$HCC_HOME/bin/hcc -I$HCC_HOME/include -I$HSA_PATH/include -I$HIP_PATH/include -std=c11 -c hipC.c diff --git a/projects/clr/hipamd/tests/src/hipC.c b/projects/clr/hipamd/tests/src/hipC.c new file mode 100644 index 0000000000..73526ecfe0 --- /dev/null +++ b/projects/clr/hipamd/tests/src/hipC.c @@ -0,0 +1,22 @@ +#include"hip_runtime.h" +#include + +#define ITER 1<<20 +#define SIZE 1024*1024*sizeof(int) + +__global__ void Iter(hipLaunchParm lp, int *Ad){ + int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + if(tx == 0){ + for(int i=0;i