Merge "Generating hiprtc lib with hcc+hip-clang" into amd-master-next

hipamd/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

hipamd/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);

hipamd/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);