Skip to content

Commit

Permalink
Neural Network Extension - add CMake support for HIP GPU backend
Browse files Browse the repository at this point in the history
- This PR also adds initial HIP kernel support for the gather layer.
- The support for executing the gather layer with HIP GPU backend will be added in the next PR.
  • Loading branch information
AryanSalmanpour committed Jul 20, 2021
1 parent 2788b7a commit 9c62123
Show file tree
Hide file tree
Showing 5 changed files with 203 additions and 16 deletions.
21 changes: 13 additions & 8 deletions amd_openvx_extensions/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -27,12 +27,12 @@ find_package(FFmpeg QUIET)
find_package(AMDRPP QUIET)

if(GPU_SUPPORT)
find_package(miopengemm PATHS ${ROCM_PATH} QUIET)
find_package(miopen PATHS ${ROCM_PATH} QUIET)
if("${BACKEND}" STREQUAL "OPENCL")
find_package(OpenCL QUIET)
if(OpenCL_FOUND)
SET(BUILD_OPENCL True)
find_package(miopengemm PATHS ${ROCM_PATH} QUIET)
find_package(miopen PATHS ${ROCM_PATH} QUIET)
else()
set(GPU_SUPPORT OFF)
set(BUILD_OPENCL False)
Expand Down Expand Up @@ -85,13 +85,18 @@ if(GPU_SUPPORT)
endif(GPU_SUPPORT)

if (NEURAL_NET)
if(GPU_SUPPORT AND OpenCL_FOUND AND BUILD_OPENCL AND miopengemm_FOUND AND miopen_FOUND)
add_subdirectory(amd_nn)
message("-- ${Green}AMD OpenVX Neural Network Extension -- amd_nn module added${ColourReset}")
elseif("${BACKEND}" STREQUAL "HIP" OR "${BACKEND}" STREQUAL "CPU")
message("-- ${Yellow}AMD OpenVX Neural Network Extension -- amd_nn module excluded with HIP/CPU Backend${ColourReset}")
if(GPU_SUPPORT AND miopengemm_FOUND AND miopen_FOUND)
if (OpenCL_FOUND AND BUILD_OPENCL)
add_subdirectory(amd_nn)
message("-- ${Green}AMD OpenVX Neural Network Extension -- amd_nn module added${ColourReset}")
elseif("${BACKEND}" STREQUAL "HIP" AND HIP_FOUND)
add_subdirectory(amd_nn)
message("-- ${Green}AMD OpenVX Neural Network Extension -- amd_nn module added${ColourReset}")
else()
message("-- ${Red}WARNING:GPU support with OpenCL/HIP Not Found -- amd_nn module excluded${ColourReset}")
endif()
else()
message("-- ${Red}WARNING:OpenCL/MIOpen/MIOpenGEMM Not Found -- amd_nn module excluded${ColourReset}")
message("-- ${Red}WARNING:GPU_SUPPORT/MIOpen/MIOpenGEMM Not Found -- amd_nn module excluded${ColourReset}")
endif()
else()
message("-- ${Cyan}Neural Net Modules turned OFF by user option -D NEURAL_NET=OFF ${ColourReset}")
Expand Down
56 changes: 48 additions & 8 deletions amd_openvx_extensions/amd_nn/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -24,25 +24,41 @@ project(amd_nn)
list(APPEND CMAKE_MODULE_PATH ${PROJECT_SOURCE_DIR}/../../amd_openvx/cmake)
set(CMAKE_CXX_STANDARD 11)

find_package(OpenCL REQUIRED)
find_package(miopengemm PATHS ${ROCM_PATH} REQUIRED)
find_package(miopen PATHS ${ROCM_PATH} REQUIRED)
find_package(Protobuf QUIET)

list(APPEND PACKAGE_DEPENDS PACKAGE OpenCL)
if(GPU_SUPPORT AND "${BACKEND}" STREQUAL "OPENCL")
find_package(OpenCL REQUIRED)
list(APPEND PACKAGE_DEPENDS PACKAGE OpenCL)
elseif(GPU_SUPPORT AND "${BACKEND}" STREQUAL "HIP")
set(OpenCL_FOUND FALSE)
if(NOT DEFINED ENV{HSA_PATH})
SET(HSA_PATH ${ROCM_PATH}/hsa)
else()
SET(HSA_PATH $ENV{HSA_PATH})
endif()
find_package(HIP QUIET REQUIRED)
if(HIP_COMPILER STREQUAL clang)
set(HIP_LIBRARY amdhip64)
else()
message(FATAL_ERROR "Unsupported HIP compiler")
endif()
list(APPEND PACKAGE_DEPENDS PACKAGE HIP)
endif()

if(Protobuf_FOUND)
add_subdirectory(../../utilities/inference_generator ${CMAKE_CURRENT_BINARY_DIR}/bin)
message("-- ${Green}Inference Generator -- utilities/inference_generator module added${ColourReset}")
add_subdirectory(../../utilities/inference_generator ${CMAKE_CURRENT_BINARY_DIR}/bin)
message("-- ${Green}Inference Generator -- utilities/inference_generator module added${ColourReset}")
else(Protobuf_FOUND)
message("-- ${Red}WARNING:Protobuf Not Found -- utilities/inference_generator module excluded${ColourReset}")
message("-- ${Red}WARNING:Protobuf Not Found -- utilities/inference_generator module excluded${ColourReset}")
endif(Protobuf_FOUND)

include_directories(../../amd_openvx/openvx/include
${OpenCL_INCLUDE_DIRS}
src
include
../../utilities/inference_generator
../../utilities/inference_generator
)

list(APPEND SOURCES
Expand Down Expand Up @@ -88,8 +104,32 @@ list(APPEND SOURCES
src/profiler.cpp
)

add_library(vx_nn SHARED ${SOURCES})
target_link_libraries(vx_nn openvx MIOpen miopengemm)

if(GPU_SUPPORT AND "${BACKEND}" STREQUAL "OPENCL" AND OpenCL_FOUND)
message("-- ${Green}AMD OpenVX Neural Network Extension -- BUILDING WITH OPENCL BACKEND${ColourReset}")
set(ENABLE_OPENCL 1)
set(ENABLE_HIP 0)
add_definitions(-DENABLE_OPENCL=${ENABLE_OPENCL} -DENABLE_HIP=${ENABLE_HIP})
include_directories(${OpenCL_INCLUDE_DIRS} ${OpenCL_INCLUDE_DIRS}/Headers)
add_library(vx_nn SHARED ${SOURCES})
target_link_libraries(vx_nn openvx MIOpen miopengemm)
elseif (GPU_SUPPORT AND "${BACKEND}" STREQUAL "HIP" AND HIP_FOUND)
message("-- ${Green}AMD OpenVX Neural Network Extension -- BUILDING WITH HIP BACKEND${ColourReset}")
set(ENABLE_OPENCL 0)
set(ENABLE_HIP 1)
add_definitions(-DENABLE_OPENCL=${ENABLE_OPENCL} -DENABLE_HIP=${ENABLE_HIP} -D__HIP_PLATFORM_HCC__)
include_directories(${HIP_PATH}/include ${HSA_PATH}/include)
link_directories(${HIP_PATH}/lib)
add_subdirectory(nn_hip)
add_library(vx_nn SHARED ${SOURCES})
set_target_properties(openvx PROPERTIES LINKER_LANGUAGE CXX)
set_target_properties(openvx PROPERTIES POSITION_INDEPENDENT_CODE ON)
target_link_libraries(vx_nn openvx openvx_hip nn_hip MIOpen miopengemm ${HIP_LIBRARY})
else()
message("-- ${Red}AMD OpenVX Neural Network Extension -- OPENCL/HIP NOT FOUND${ColourReset}")
message("-- ${Red}AMD OpenVX Neural Network Extension module excluded {ColourReset}")
endif()


install(TARGETS vx_nn DESTINATION lib)
install(FILES include/vx_amd_nn.h DESTINATION include)
Expand Down
39 changes: 39 additions & 0 deletions amd_openvx_extensions/amd_nn/nn_hip/CMakeLists.txt
Original file line number Diff line number Diff line change
@@ -0,0 +1,39 @@
# Copyright (c) 2015 - 2020 Advanced Micro Devices, Inc. All rights reserved.
#
# Permission is hereby granted, free of charge, to any person obtaining a copy
# of this software and associated documentation files (the "Software"), to deal
# in the Software without restriction, including without limitation the rights
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
# copies of the Software, and to permit persons to whom the Software is
# furnished to do so, subject to the following conditions:
#
# The above copyright notice and this permission notice shall be included in
# all copies or substantial portions of the Software.
#
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
# THE SOFTWARE.

cmake_minimum_required(VERSION 3.5)
project(nn_hip CXX)

set(NN_HIP_SOURCES
nn_hip_kernels.cpp
)

set(hip_library_name amdhip64)
set_source_files_properties(${NN_HIP_SOURCES} PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1)
set(HIP_CXX_FLAGS -std=c++14)

if(CMAKE_BUILD_TYPE MATCHES Debug)
set(HIP_CXX_FLAGS ${HIP_CXX_FLAGS} -g)
ENDIF()

hip_add_library(nn_hip ${NN_HIP_SOURCES} HIPCC_OPTIONS "${HIP_CXX_FLAGS}" SHARED)
target_compile_definitions(nn_hip PRIVATE __HIP_PLATFORM_HCC__)
install(TARGETS nn_hip DESTINATION lib)

32 changes: 32 additions & 0 deletions amd_openvx_extensions/amd_nn/nn_hip/nn_hip_host_decls.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,32 @@
/*
Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/


#ifndef NN_HIP_HOST_DECLS_H
#define NN_HIP_HOST_DECLS_H
#include "hip/hip_runtime.h"
#include <VX/vx.h>

int HipExec_Gather_layer(hipStream_t stream, dim3 globalThreads, dim3 localThreads, vx_enum type, unsigned char* in, uint in_offset,
uint4 in_stride, unsigned char* ind, uint ind_offset, uint4 ind_stride, unsigned char* out, uint out_offset,
uint4 out_stride, uint axis);

#endif //NN_HIP_HOST_DECLS_H
71 changes: 71 additions & 0 deletions amd_openvx_extensions/amd_nn/nn_hip/nn_hip_kernels.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,71 @@
/*
Copyright (c) 2015 - 2020 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/

#include "../../../amd_openvx/openvx/hipvx/hip_common_funcs.h"
#include "nn_hip_host_decls.h"
#include "hip/hip_fp16.h"

// ----------------------------------------------------------------------------
// Neural Network kernels for hip backend
// ----------------------------------------------------------------------------

template <typename T>
__global__ void __attribute__((visibility("default")))
Hip_Gather_layer(uchar * in, uint in_offset, uint4 in_stride, uchar * ind, uint ind_offset, uint4 ind_stride,
uchar *out, uint out_offset, uint4 out_stride, uint axis) {

uint x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
uint y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
uint c = hipBlockDim_z * hipBlockIdx_z + hipThreadIdx_z;

int indices = *(int*)&ind[ind_offset + y * ind_stride.x];
T value;
uint offset;
if (axis == 0) {
value = *(T*)&in[in_offset + x * in_stride.x + indices * in_stride.y + c * in_stride.z];
offset = out_offset + x * out_stride.x + y * out_stride.y + c * out_stride.z;
} else if (axis == 1) {
value = *(T*)&in[in_offset + indices * in_stride.x + c * in_stride.y];
offset = out_offset + y * out_stride.x + c * out_stride.y;
} else if (axis == 2) {
value = *(T*)&in[in_offset + c * in_stride.x];
offset = out_offset + c * out_stride.x;
}
out += offset;
*(T *)&out[0] = value;
}

int HipExec_Gather_layer(hipStream_t stream, dim3 globalThreads, dim3 localThreads, vx_enum type, uchar* in, uint in_offset, uint4 in_stride, uchar* ind, uint ind_offset,
uint4 ind_stride, uchar* out, uint out_offset, uint4 out_stride, uint axis) {

if (type == VX_TYPE_FLOAT32) {
hipLaunchKernelGGL(Hip_Gather_layer<float>,
dim3(ceil((float)globalThreads.x/localThreads.x), ceil((float)globalThreads.y/localThreads.y), ceil((float)globalThreads.z/localThreads.z)),
dim3(localThreads.x, localThreads.y, localThreads.z), 0, stream, in, in_offset, in_stride, ind, ind_offset, ind_stride, out, out_offset, out_stride, axis);
} else {
hipLaunchKernelGGL(Hip_Gather_layer<__half>,
dim3(ceil((float)globalThreads.x/localThreads.x), ceil((float)globalThreads.y/localThreads.y), ceil((float)globalThreads.z/localThreads.z)),
dim3(localThreads.x, localThreads.y, localThreads.z), 0, stream, in, in_offset, in_stride, ind, ind_offset, ind_stride, out, out_offset, out_stride, axis);
}

return VX_SUCCESS;
}

0 comments on commit 9c62123

Please sign in to comment.