Merge branch 'privatestaging' of https://github.com/AMDComputeLibraries/HIP-privatestaging into privatestaging

Conflicts:
	include/hcc_detail/trace_helper.h


[ROCm/clr commit: ba114a238f]

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)
 

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.

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 <iostream>
 #include <iomanip>
@@ -35,7 +35,7 @@ THE SOFTWARE.
 
 // Building block functions:
 template <typename T>
-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 <typename T>
-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 <typename T, typename... Args> 
-std::string ToString(T first, Args... args) {
+inline std::string ToString(T first, Args... args) 
+{
     return ToString(first) + ", " + ToString(args...) ;
 }

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 )

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

projects/clr/hipamd/tests/src/hipC.c

@@ -0,0 +1,22 @@
+#include"hip_runtime.h"
+#include<stdio.h>
+
+#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<ITER;i++){
+            Ad[tx] += 1;
+        }
+    }
+}
+
+int main(){
+    int A=0, *Ad;
+    hipMalloc((void**)&Ad, SIZE);
+    hipMemcpy(Ad, &A, SIZE, hipMemcpyHostToDevice);
+    hipLaunchKernel(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, 0, Ad);
+    hipMemcpy(&A, Ad, SIZE, hipMemcpyDeviceToHost);
+}