Enable ROCm/HIP GPU acceleration

Signed-off-by: Gavin Zhao <[email protected]>

Author: GZGavinZhao

CMakeLists.txt

@@ -39,6 +39,7 @@ option(Other_Optimization_For_Other "Enable other optimizatio flag for other if
 option(Enable_IPO "Enable link time optimization, may cause unknown bugs even slow performance" OFF)
 option(Maximum_Optimization "Enable maximum optimizations for the current platform" OFF)
 option(Enable_CUDA "Enable CUDA module" OFF)
+option(Enable_HIP "Enable HIP module" OFF)
 option(Enable_OpenCL "Enable OpenCL module" ON)
 option(Enable_NCNN "Enable ncnn module" OFF)
 option(Enable_Fast_Math "Enable fast math" OFF)
@@ -89,6 +90,8 @@ if(Enable_CUDA)
         set(CMAKE_CUDA_ARCHITECTURES ${CUDA_Minimum_CC}-virtual ${CUDA_CC}-real)
     endif()
     enable_language(CUDA)
+elseif(Enable_HIP)
+    enable_language(HIP)
 endif()
 
 if(Use_Eigen3)
@@ -308,6 +311,7 @@ message(STATUS
 "   Enable IPO ${Enable_IPO}\n"
 "   Enable OpenCL ${Enable_OpenCL}\n"
 "   Enable CUDA ${Enable_CUDA}\n"
+"   Enable HIP ${Enable_HIP}\n"
 "   Enable NCNN ${Enable_NCNN}\n"
 "   Enable OpenCV DNN ${Enable_OpenCV_DNN}\n"
 "   Enable video ${Enable_Video}\n"

core/CMakeLists.txt

@@ -8,11 +8,13 @@ if(Build_Static_Core)
     add_library(${PROJECT_NAME} STATIC 
         ${INCLUDE} ${SOURCE} 
         $<$<BOOL:${Enable_CUDA}>:$<TARGET_OBJECTS:CUDA_Module>>
+        $<$<BOOL:${Enable_HIP}>:$<TARGET_OBJECTS:HIP_Module>>
     )
 else()
     add_library(${PROJECT_NAME} SHARED 
         ${INCLUDE} ${SOURCE} 
         $<$<BOOL:${Enable_CUDA}>:$<TARGET_OBJECTS:CUDA_Module>>
+        $<$<BOOL:${Enable_HIP}>:$<TARGET_OBJECTS:HIP_Module>>
     )
 endif()
 
@@ -21,6 +23,9 @@ target_compile_definitions(
     PUBLIC 
         $<$<BOOL:${Enable_OpenCL}>:ENABLE_OPENCL>
         $<$<BOOL:${Enable_CUDA}>:ENABLE_CUDA>
+        $<$<BOOL:${Enable_HIP}>:ENABLE_HIP>
+        # This allows HIP to emulate as CUDA
+        $<$<BOOL:${Enable_HIP}>:ENABLE_CUDA>
         $<$<BOOL:${Enable_NCNN}>:ENABLE_NCNN>
         $<$<BOOL:${Enable_Video}>:ENABLE_VIDEO>
         $<$<BOOL:${Enable_Preview_GUI}>:ENABLE_PREVIEW_GUI>

hip/CMakeLists.txt

@@ -0,0 +1,33 @@
+if(Enable_HIP)
+	project(HIP_Module LANGUAGES CXX HIP)
+	
+	file(GLOB SOURCES src/*.hip)
+	file(GLOB INCLUDE include/*.hpp include/*.hip)
+	
+	source_group("hip_include" FILES ${INCLUDE}) 
+	source_group("hip_source" FILES ${SOURCE}) 
+	
+	message(STATUS "CMAKE_HIP_COMPILER:         ${CMAKE_HIP_COMPILER}")
+	message(STATUS "CMAKE_HIP_COMPILER_ID:      ${CMAKE_HIP_COMPILER_ID}")
+	message(STATUS "CMAKE_HIP_COMPILER_VERSION: ${CMAKE_HIP_COMPILER_VERSION}")
+	message(STATUS "HIP SOURCES:                ${SOURCES}")
+	message(STATUS "HIP INCLUDE:                ${INCLUDE}")
+	
+	add_library(${PROJECT_NAME} OBJECT ${INCLUDE} ${SOURCES})
+	
+	target_include_directories(${PROJECT_NAME} PRIVATE ${TOP_DIR}/core/include include)
+	
+	target_compile_definitions(${PROJECT_NAME} PRIVATE ENABLE_CUDA ENABLE_HIP)
+	
+	set_target_properties(
+	    ${PROJECT_NAME} PROPERTIES 
+	        HIP_STANDARD 17
+	)
+
+    if((NOT Build_Static_Core) OR (Build_Static_Core AND Build_Static_Core_PIC))
+        set_target_properties(
+            ${PROJECT_NAME} PROPERTIES 
+                POSITION_INDEPENDENT_CODE True
+        )
+    endif()
+endif()

hip/include/ACNetCommon.hpp

@@ -0,0 +1,139 @@
+#ifndef ANIME4KCPP_CUDA_CUDA_HELPER_CUH
+#define ANIME4KCPP_CUDA_CUDA_HELPER_CUH
+
+#include <type_traits>
+#include <cstdint>
+
+#include <hip/hip_runtime.h>
+#include <hip/hip_fp16.h>
+#include <hip/device_functions.h>
+#include <hip/amd_detail/amd_surface_functions.h>
+
+#include "ACException.hpp"
+
+#define CheckCudaErr(err)   \
+    if (err != hipSuccess) \
+    throw Anime4KCPP::ACException<Anime4KCPP::ExceptionType::GPU, true>(hipGetErrorString(err), std::string(__FILE__), __LINE__)
+
+using uchar = std::uint8_t;
+using ushort = std::uint16_t;
+
+extern int currCudaDeviceID;
+
+template <typename T>
+struct PixelValue;
+
+template <>
+struct PixelValue<uchar>
+{
+    __device__ static constexpr uchar max()
+    {
+        return 255;
+    }
+
+    __device__ static constexpr uchar min()
+    {
+        return 0;
+    }
+};
+
+template <>
+struct PixelValue<ushort>
+{
+    __device__ static constexpr ushort max()
+    {
+        return 65535;
+    }
+
+    __device__ static constexpr ushort min()
+    {
+        return 0;
+    }
+};
+
+template <>
+struct PixelValue<float>
+{
+    __device__ static constexpr float max()
+    {
+        return 1.0f;
+    }
+
+    __device__ static constexpr float min()
+    {
+        return 0.0f;
+    }
+};
+
+template <typename T, int dim>
+struct Vec;
+
+template <>
+struct Vec<uchar, 2>
+{
+    using type = uchar2;
+};
+
+template <>
+struct Vec<uchar, 4>
+{
+    using type = uchar4;
+};
+
+template <>
+struct Vec<ushort, 2>
+{
+    using type = ushort2;
+};
+
+template <>
+struct Vec<ushort, 4>
+{
+    using type = ushort4;
+};
+
+template <>
+struct Vec<float, 2>
+{
+    using type = float2;
+};
+
+template <>
+struct Vec<float, 4>
+{
+    using type = float4;
+};
+
+template <typename T>
+using Vec2 = Vec<T, 2>;
+
+template <typename T>
+using Vec4 = Vec<T, 4>;
+
+template <typename T>
+inline __device__ typename Vec4<T>::type makeVec4(T x, T y, T z, T w);
+
+template <>
+inline __device__ typename Vec4<uchar>::type makeVec4(uchar x, uchar y, uchar z, uchar w)
+{
+    return make_uchar4(x, y, z, w);
+}
+
+template <>
+inline __device__ typename Vec4<ushort>::type makeVec4(ushort x, ushort y, ushort z, ushort w)
+{
+    return make_ushort4(x, y, z, w);
+}
+
+template <>
+inline __device__ typename Vec4<float>::type makeVec4(float x, float y, float z, float w)
+{
+    return make_float4(x, y, z, w);
+}
+
+inline __device__ float clamp(float f, float a, float b)
+{
+    return fmaxf(a, fminf(f, b));
+}
+
+#endif // !ANIME4KCPP_CUDA_CUDA_HELPER_CUH