Compile HIP runtime with hidden visibility by default (#1303)

* add default visibility to most APIs in program_state

* remove unwanted C++ headers

* Add symbol visibility pragmas and compiler flags

* Add visibility attribute to APIs in channel_descriptor and hip_hcc

* remove unused headers

* simplify build flags with hcc

* add pragma visibility hidden to functional_grid_launch

* [CMake] add gfx908 back


[ROCm/hip commit: 83af327ef2]

projects/hip/CMakeLists.txt

@@ -221,7 +221,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
     set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -DHIP_VERSION_MAJOR=${HIP_VERSION_MAJOR} -DHIP_VERSION_MINOR=${HIP_VERSION_MINOR} -DHIP_VERSION_PATCH=${HIP_VERSION_PATCH}")
 
     # Add remaining flags
-    execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --cxxflags OUTPUT_VARIABLE HCC_CXX_FLAGS)
+    set(HCC_CXX_FLAGS  "-hc -fno-gpu-rdc --amdgpu-target=gfx803 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 ")
     set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -fPIC ${HCC_CXX_FLAGS} -I${HSA_PATH}/include")
 
     # Set compiler and compiler flags
@@ -252,9 +252,6 @@ if(HIP_PLATFORM STREQUAL "hcc")
         src/env.cpp
         src/h2f.cpp)
 
-    execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS)
-    set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS}")
-    set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx803 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908")
     if(COMPILE_HIP_ATP_MARKER)
         set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L/opt/rocm/profiler/CXLActivityLogger/bin/x86_64 -lCXLActivityLogger")
     endif()
@@ -270,7 +267,8 @@ if(HIP_PLATFORM STREQUAL "hcc")
                 PRIVATE ${PROJECT_SOURCE_DIR}/include ${HSA_PATH}/include)
         target_link_libraries(hiprtc PUBLIC stdc++fs)
     endif()
-
+    set_target_properties(hip_hcc PROPERTIES CXX_VISIBILITY_PRESET hidden)
+    set_target_properties(hip_hcc PROPERTIES VISIBILITY_INLINES_HIDDEN 1)
 
     if(HIP_PLATFORM STREQUAL "hcc")
         find_package(amd_comgr REQUIRED CONFIG

projects/hip/include/hip/hcc_detail/channel_descriptor.h

@@ -23,11 +23,13 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H
 #define HIP_INCLUDE_HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H
 
+#include <hip/hip_common.h>
 #include <hip/hcc_detail/driver_types.h>
 #include <hip/hcc_detail/hip_vector_types.h>
 
 #ifdef __cplusplus
 
+HIP_PUBLIC_API
 hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f);
 
 static inline hipChannelFormatDesc hipCreateChannelDescHalf() {

projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp

@@ -25,25 +25,16 @@ THE SOFTWARE.
 #include "concepts.hpp"
 #include "helpers.hpp"
 #include "program_state.hpp"
+#include "hip_runtime_api.h"
 
-#include "hc.hpp"
-#include "hip/hip_hcc.h"
-#include "hip_runtime.h"
-
-#include <algorithm>
-#include <cstddef>
 #include <cstdint>
 #include <cstring>
-#include <functional>
-#include <iostream>
-#include <mutex>
 #include <stdexcept>
-#include <string>
 #include <tuple>
 #include <type_traits>
-#include <unordered_map>
 #include <utility>
-#include <vector>
+
+#pragma GCC visibility push(hidden)
 
 namespace hip_impl {
 template <typename T, typename std::enable_if<std::is_integral<T>{}>::type* = nullptr>
@@ -114,7 +105,7 @@ inline hip_impl::kernarg make_kernarg(
 }
 
 
-hsa_agent_t target_agent(hipStream_t stream);
+HIP_INTERNAL_EXPORTED_API hsa_agent_t target_agent(hipStream_t stream);
 
 inline
 __attribute__((visibility("hidden")))
@@ -194,3 +185,5 @@ inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBloc
     hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{},
                        std::move(args)...);
 }
+
+#pragma GCC visibility pop

projects/hip/include/hip/hcc_detail/hip_runtime_api.h

@@ -67,19 +67,18 @@ THE SOFTWARE.
 #define HIP_LAUNCH_PARAM_END ((void*)0x03)
 
 #ifdef __cplusplus
-  #include <algorithm>
-  #include <mutex>
-  #include <string>
-  #include <unordered_map>
-  #include <vector>
-
   #define __dparm(x) \
           = x
 #else
   #define __dparm(x)
 #endif
 
+#ifdef __GNUC__
+#pragma GCC visibility push (default)
+#endif
+
 #ifdef __cplusplus
+
 namespace hip_impl {
 hipError_t hip_init();
 }  // namespace hip_impl
@@ -3233,6 +3232,9 @@ hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject);
 
 #endif
 
+#ifdef __GNUC__
+#pragma GCC visibility pop
+#endif
 
 /**
  *-------------------------------------------------------------------------------------------------

projects/hip/include/hip/hcc_detail/program_state.hpp

@@ -31,11 +31,19 @@ THE SOFTWARE.
 #include <cstdint>
 #include <cstdlib>
 
+#include <hip/hip_common.h>
+
 struct ihipModuleSymbol_t;
 using hipFunction_t = ihipModuleSymbol_t*;
 
 namespace hip_impl {
 
+// This section contains internal APIs that
+// needs to be exported
+#ifdef __GNUC__
+#pragma GCC visibility push (default)
+#endif
+
 struct kernarg_impl;
 class kernarg {
 public:
@@ -83,6 +91,10 @@ private:
     friend kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t);
 };
 
+#ifdef __GNUC__
+#pragma GCC visibility pop
+#endif
+
 inline
 __attribute__((visibility("hidden")))
 program_state& get_program_state() {

projects/hip/include/hip/hip_common.h

@@ -45,6 +45,14 @@ THE SOFTWARE.
 #define __HIP_DEVICE_COMPILE__ 1
 #endif
 
+#ifdef __GNUC__
+#define HIP_PUBLIC_API              __attribute__ ((visibility ("default")))
+#define HIP_INTERNAL_EXPORTED_API   __attribute__ ((visibility ("default")))
+#else
+#define HIP_PUBLIC_API
+#define HIP_INTERNAL_EXPORTED_API 
+#endif
+
 #if __HIP_DEVICE_COMPILE__ == 0
 // 32-bit Atomics
 #define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (0)

projects/hip/include/hip/hip_hcc.h

@@ -48,6 +48,7 @@ class accelerator_view;
  * @brief Return hc::accelerator associated with the specified deviceId
  * @return #hipSuccess, #hipErrorInvalidDevice
  */
+HIP_PUBLIC_API
 hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc);
 
 /**
@@ -56,6 +57,7 @@ hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc);
  * If stream is 0, the accelerator_view for the default stream is returned.
  * @return #hipSuccess
  */
+HIP_PUBLIC_API
 hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** av);
 
 
@@ -89,6 +91,7 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** a
 
  * HIP/ROCm actually updates the start event when the associated kernel completes.
  */
+HIP_PUBLIC_API
 hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
                                     uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
                                     uint32_t localWorkSizeX, uint32_t localWorkSizeY,
@@ -98,6 +101,7 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
                                     hipEvent_t stopEvent = nullptr,
                                     uint32_t flags = 0);
 
+HIP_PUBLIC_API
 hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
                                     uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
                                     uint32_t localWorkSizeX, uint32_t localWorkSizeY,

projects/hip/src/functional_grid_launch.inl

@@ -41,7 +41,7 @@ using namespace std;
 
 namespace hip_impl
 {
-    hsa_agent_t target_agent(hipStream_t stream)
+    HIP_INTERNAL_EXPORTED_API hsa_agent_t target_agent(hipStream_t stream)
     {
         if (stream) {
             return *static_cast<hsa_agent_t*>(

projects/hip/src/hip_hcc.cpp

@@ -46,6 +46,7 @@ THE SOFTWARE.
 #include "hsa/hsa_ext_image.h"
 #include "hip/hip_runtime.h"
 #include "hip_hcc_internal.h"
+#include "hip/hip_hcc.h"
 #include "trace_helper.h"
 #include "env.h"