From fa2cc448de45a052cefc52bd891df5ee7adcfd8c Mon Sep 17 00:00:00 2001
From: Jatin Chaudhary <Jatin.Chaudhary@amd.com>
Date: Fri, 21 Feb 2020 21:52:19 +0530
Subject: [PATCH] Generating hiprtc lib with hcc+hip-clang

Review comments - generate hiprtc lib everytime when HIP_PLATFORM is hcc

Changes for hip-clang

Removing pre processor directive to simplify

Change-Id: Id38ab368362b58ee0458baeb8051fea709ae6bba


[ROCm/hip commit: 54447268cc1115fc84305c0a5327bc0ac5a12c5c]
---
 projects/hip/CMakeLists.txt             | 24 +++++++++---------------
 projects/hip/src/hiprtc.cpp             |  4 ++--
 projects/hip/tests/src/hiprtc/saxpy.cpp |  2 +-
 3 files changed, 12 insertions(+), 18 deletions(-)

diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt
index b6c4adaa09..f64ca45559 100644
--- a/projects/hip/CMakeLists.txt
+++ b/projects/hip/CMakeLists.txt
@@ -373,18 +373,17 @@ if(HIP_PLATFORM STREQUAL "hcc")
     set_property ( TARGET hip_hcc PROPERTY VERSION "${HIP_LIB_VERSION_STRING}" )
     set_property ( TARGET hip_hcc PROPERTY SOVERSION "${HIP_LIB_VERSION_MAJOR}" )
 
-    if(HIP_COMPILER STREQUAL "hcc")
-        target_link_libraries(hip_hcc PRIVATE hc_am)
-        target_link_libraries(hip_hcc_static PRIVATE hc_am)
+    target_link_libraries(hip_hcc PRIVATE hc_am)
+    target_link_libraries(hip_hcc_static PRIVATE hc_am)
 
-        add_library(hiprtc SHARED src/hiprtc.cpp src/code_object_bundle.cpp)
-        set_property ( TARGET hiprtc PROPERTY VERSION "${HIP_LIB_VERSION_STRING}" )
-        set_property ( TARGET hiprtc PROPERTY SOVERSION "${HIP_LIB_VERSION_MAJOR}" )
+    add_library(hiprtc SHARED src/hiprtc.cpp src/code_object_bundle.cpp)
+    set_property ( TARGET hiprtc PROPERTY VERSION "${HIP_LIB_VERSION_STRING}" )
+    set_property ( TARGET hiprtc PROPERTY SOVERSION "${HIP_LIB_VERSION_MAJOR}" )
+
+    target_include_directories(
+        hiprtc SYSTEM
+            PRIVATE ${PROJECT_SOURCE_DIR}/include ${HSA_PATH}/include)
 
-        target_include_directories(
-            hiprtc SYSTEM
-                PRIVATE ${PROJECT_SOURCE_DIR}/include ${HSA_PATH}/include)
-endif()
     set_target_properties(hip_hcc PROPERTIES CXX_VISIBILITY_PRESET hidden)
     set_target_properties(hip_hcc PROPERTIES VISIBILITY_INLINES_HIDDEN 1)
 
@@ -448,12 +447,7 @@ endif()
 #############################
 # Install hip_hcc if platform is hcc
 if(HIP_PLATFORM STREQUAL "hcc")
-    if(HIP_COMPILER STREQUAL "hcc")
       install(TARGETS hip_hcc_static hip_hcc hiprtc DESTINATION lib)
-    else()
-      install(TARGETS hip_hcc_static hip_hcc DESTINATION lib)
-    endif()
-
 endif()
 
 # Install .hipInfo
diff --git a/projects/hip/src/hiprtc.cpp b/projects/hip/src/hiprtc.cpp
index 5198bf0cbb..8dcb944f72 100644
--- a/projects/hip/src/hiprtc.cpp
+++ b/projects/hip/src/hiprtc.cpp
@@ -235,7 +235,7 @@ struct _hiprtcProgram {
 
         const auto it{find_if(reader.sections.begin(), reader.sections.end(),
                               [](const section* x) {
-            return x->get_name() == ".kernel";
+            return (x->get_name() == ".hip_fatbin") || (x->get_name() == ".kernel");
         })};
 
         if (it == reader.sections.end()) return false;
@@ -513,7 +513,7 @@ hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o)
 
     const auto src{p->writeTemporaryFiles(tmp.path())};
 
-    vector<string> args{hipcc, "-shared"};
+    vector<string> args{hipcc, "-fPIC -shared"};
     if (n) args.insert(args.cend(), o, o + n);
 
     handleTarget(args);
diff --git a/projects/hip/tests/src/hiprtc/saxpy.cpp b/projects/hip/tests/src/hiprtc/saxpy.cpp
index 156a44afe4..d063578757 100755
--- a/projects/hip/tests/src/hiprtc/saxpy.cpp
+++ b/projects/hip/tests/src/hiprtc/saxpy.cpp
@@ -143,7 +143,7 @@ int main()
     hipMemcpyDtoH(hOut.get(), dOut, bufferSize);
 
     for (size_t i = 0; i < n; ++i) {
-       if (a * hX[i] + hY[i] != hOut[i]) { failed("Validation failed."); }
+        if (fabs(a * hX[i] + hY[i] - hOut[i]) > fabs(hOut[i])* 1e-6) { failed("Validation failed."); }
     }
 
     hipFree(dX);