diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt index 187b32045b..8b9a603ad8 100644 --- a/projects/hip/CMakeLists.txt +++ b/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 diff --git a/projects/hip/include/hip/hcc_detail/channel_descriptor.h b/projects/hip/include/hip/hcc_detail/channel_descriptor.h index de290fafc1..38acff9951 100644 --- a/projects/hip/include/hip/hcc_detail/channel_descriptor.h +++ b/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 #include #include #ifdef __cplusplus +HIP_PUBLIC_API hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f); static inline hipChannelFormatDesc hipCreateChannelDescHalf() { diff --git a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp b/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp index f502fddf42..883470b1fb 100644 --- a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/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 -#include #include #include -#include -#include -#include #include -#include #include #include -#include #include -#include + +#pragma GCC visibility push(hidden) namespace hip_impl { template {}>::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 diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index e71ac8d8fb..cec9eb9c94 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/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 - #include - #include - #include - #include - #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 /** *------------------------------------------------------------------------------------------------- diff --git a/projects/hip/include/hip/hcc_detail/program_state.hpp b/projects/hip/include/hip/hcc_detail/program_state.hpp index 4e0fbb76a4..fca88f8094 100644 --- a/projects/hip/include/hip/hcc_detail/program_state.hpp +++ b/projects/hip/include/hip/hcc_detail/program_state.hpp @@ -31,11 +31,19 @@ THE SOFTWARE. #include #include +#include + 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() { diff --git a/projects/hip/include/hip/hip_common.h b/projects/hip/include/hip/hip_common.h index 71285fc8dc..79c787be4e 100644 --- a/projects/hip/include/hip/hip_common.h +++ b/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) diff --git a/projects/hip/include/hip/hip_hcc.h b/projects/hip/include/hip/hip_hcc.h index 3935f4903b..a82a8c622b 100644 --- a/projects/hip/include/hip/hip_hcc.h +++ b/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, diff --git a/projects/hip/src/functional_grid_launch.inl b/projects/hip/src/functional_grid_launch.inl index 3c611ed468..a99a15db28 100644 --- a/projects/hip/src/functional_grid_launch.inl +++ b/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( diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index dcfdfb6137..da0b38ed4f 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/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"