diff --git a/hipamd/CMakeLists.txt b/hipamd/CMakeLists.txt index b6c4adaa09..f64ca45559 100644 --- a/hipamd/CMakeLists.txt +++ b/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 diff --git a/hipamd/src/hiprtc.cpp b/hipamd/src/hiprtc.cpp index 5198bf0cbb..8dcb944f72 100644 --- a/hipamd/src/hiprtc.cpp +++ b/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 args{hipcc, "-shared"}; + vector args{hipcc, "-fPIC -shared"}; if (n) args.insert(args.cend(), o, o + n); handleTarget(args); diff --git a/hipamd/tests/src/hiprtc/saxpy.cpp b/hipamd/tests/src/hiprtc/saxpy.cpp index 156a44afe4..d063578757 100755 --- a/hipamd/tests/src/hiprtc/saxpy.cpp +++ b/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);