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

hipamd/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() {

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

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

hipamd/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() {

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

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