Nigel
/
mindspore
forked from mindspore-Ecosystem/mindspore
1
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

!30120 Package the cuda operators as a dynamic link library 

Merge pull request !30120 from jinjiali-kali/cuda_ops_3
This commit is contained in:
i-robot 2022-02-18 01:28:24 +00:00 committed by Gitee
parent 7a864db171 cd7db0f57f
commit 8e21ffae65
No known key found for this signature in database
GPG Key ID: 173E9B9CA92EEF8F
772 changed files with 11356 additions and 10384 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
cmake/package.cmake
View File

@ -212,6 +212,11 @@ if(ENABLE_GPU)
DESTINATION ${INSTALL_LIB_DIR}
COMPONENT mindspore
)
install(
TARGETS cuda_ops
DESTINATION ${INSTALL_LIB_DIR}
COMPONENT mindspore
)
endif()
if(ENABLE_D)

6
mindspore/ccsrc/CMakeLists.txt
View File

@ -114,6 +114,8 @@ if(ENABLE_GPU)
"plugin/device/gpu/kernel/*.cu"
)
list(REMOVE_ITEM GPU_SRC_LIST "plugin/device/gpu/kernel/cuda_impl/cuda_ops/*.cu")
list(APPEND CUDA_NVCC_FLAGS -arch=sm_53 --expt-relaxed-constexpr)
list(REMOVE_ITEM GPU_SRC_LIST "plugin/device/gpu/hal/device/blocking_queue.cc"
"plugin/device/gpu/hal/device/gpu_buffer_mgr.cc")
@ -145,6 +147,8 @@ if(ENABLE_GPU)
cuda_add_library(gpu_cuda_lib STATIC ${GPU_SRC_LIST})
set(CMAKE_CXX_FLAGS ${NVCC_TMP_CMAKE_CXX_FLAGS})
add_compile_definitions(ENABLE_GPU)
add_subdirectory(plugin/device/gpu/kernel/cuda_impl/cuda_ops)
endif()
@ -430,7 +434,7 @@ endif()
if(ENABLE_GPU)
message("add gpu lib to c_expression")
target_link_libraries(_c_expression PRIVATE gpu_cuda_lib gpu_queue cublas
target_link_libraries(_c_expression PRIVATE gpu_cuda_lib gpu_queue cublas cuda_ops
${CUDA_PATH}/lib64/libcurand.so
${CUDNN_LIBRARY_PATH}
${CUDA_PATH}/lib64/libcudart.so

2
mindspore/ccsrc/cxx_api/CMakeLists.txt
View File

@ -140,7 +140,7 @@ if(ENABLE_D)
endif()
if(ENABLE_GPU)
target_link_libraries(mindspore_shared_lib PRIVATE gpu_cuda_lib gpu_queue cublas
target_link_libraries(mindspore_shared_lib PRIVATE gpu_cuda_lib gpu_queue cublas cuda_ops
${CUDA_PATH}/lib64/libcurand.so
${CUDNN_LIBRARY_PATH}
${CUDA_PATH}/lib64/libcudart.so

2
mindspore/ccsrc/plugin/device/gpu/hal/device/kernel_info_setter.cc
View File

@ -26,7 +26,7 @@
#include "kernel/oplib/oplib.h"
#include "backend/common/session/anf_runtime_algorithm.h"
#include "plugin/device/gpu/kernel/custom/custom_aot_gpu_kernel.h"
#include "plugin/device/gpu/hal/device/cuda_common.h"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cuda_common.h"
#include "utils/ms_context.h"
#include "utils/ms_utils.h"
#include "utils/utils.h"

2
mindspore/ccsrc/plugin/device/gpu/hal/hardware/gpu_device_context.cc
View File

@ -28,7 +28,7 @@
#include "plugin/device/gpu/hal/device/gpu_buffer_mgr.h"
#include "kernel/common_utils.h"
#include "plugin/device/gpu/hal/device/gpu_common.h"
#include "plugin/device/gpu/hal/device/cuda_common.h"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cuda_common.h"
#include "plugin/device/gpu/hal/hardware/optimizer.h"
#include "utils/ms_device_shape_transfer.h"
#include "utils/context/graph_kernel_flags.h"

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/argmax_gpu_kernel.h
View File

@ -21,7 +21,7 @@
#include <string>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/argmax_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/argmax_impl.cuh"
namespace mindspore {
namespace kernel {
template <typename T, typename S>

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/argmaxandminwithvalue_gpu_kernel.h
View File

@ -22,7 +22,7 @@
#include <map>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/general_reduction_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/general_reduction_impl.cuh"
namespace mindspore {
namespace kernel {
template <typename T, typename S>

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/batchtospace_gpu_kernel.h
View File

@ -21,7 +21,7 @@
#include <string>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/batchtospace_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/batchtospace_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/broadcast_to_gpu_kernel.h
View File

@ -21,7 +21,7 @@
#include <string>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/broadcast_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/broadcast_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/cast_gpu_kernel.h
View File

@ -21,7 +21,7 @@
#include <string>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/cast_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cast_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/concatv2_gpu_kernel.h
View File

@ -22,7 +22,7 @@
#include <memory>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/concatv2_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/concatv2_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/crop_and_resize_gpu_kernel.h
View File

@ -21,7 +21,7 @@
#include <string>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/crop_and_resize_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/crop_and_resize_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/depthtospace_gpu_kernel.h
View File

@ -20,7 +20,7 @@
#include <vector>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/depthtospace_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/depthtospace_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/dynamic_range_gpu_kernel.h
View File

@ -21,7 +21,7 @@
#include <vector>
#include "plugin/device/gpu/kernel/cuda_impl/dynamic_range_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/dynamic_range_impl.cuh"
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/embedding_lookup_gpu_kernel.h
View File

@ -21,7 +21,7 @@
#include <algorithm>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/embedding_lookup_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/embedding_lookup_impl.cuh"
namespace mindspore {
namespace kernel {

4
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/extract_image_patches_gpu_kernel.h
View File

@ -23,8 +23,8 @@
#include <vector>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/transpose_impl_opt.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/extract_image_patches_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/transpose_impl_opt.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/extract_image_patches_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/gather_gpu_kernel.h
View File

@ -20,7 +20,7 @@
#include <vector>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/gather.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/gather.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/gather_grad_gpu_kernel.h
View File

@ -20,7 +20,7 @@
#include <vector>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/gather_grad.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/gather_grad.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/gathernd_gpu_kernel.h
View File

@ -20,7 +20,7 @@
#include <vector>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/gathernd.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/gathernd.cuh"
#include "backend/common/session/anf_runtime_algorithm.h"
namespace mindspore {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/gatherv2_gpu_kernel.h
View File

@ -21,7 +21,7 @@
#include <algorithm>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/gatherv2.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/gatherv2.cuh"
#include "backend/common/session/anf_runtime_algorithm.h"
namespace mindspore {

6
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/in_top_k_gpu_kernel.h
View File

@ -22,9 +22,9 @@
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/cast_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/in_top_k_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/topk_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cast_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/in_top_k_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/topk_impl.cuh"
namespace mindspore {
namespace kernel {

6
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/matrix_band_part_gpu_kernel.h
View File

@ -23,9 +23,9 @@
#include <cuda_runtime.h>
#include <vector>
#include <algorithm>
#include "utils/complex.h"
#include "plugin/device/gpu/kernel/cuda_impl/matrix_band_part_impl.cuh"
#include "plugin/device/gpu/hal/device/cuda_common.h"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/complex.h"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/matrix_band_part_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cuda_common.h"
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/kernel_constants.h"

6
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/matrix_diag_part_gpu_kernel.h
View File

@ -25,9 +25,9 @@
#include <string>
#include <utility>
#include <algorithm>
#include "utils/complex.h"
#include "plugin/device/gpu/kernel/cuda_impl/matrix_diag_part_impl.cuh"
#include "plugin/device/gpu/hal/device/cuda_common.h"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/complex.h"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/matrix_diag_part_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cuda_common.h"
#include "kernel/common_utils.h"
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/matrix_set_diag_gpu_kernel.h
View File

@ -26,7 +26,7 @@
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "kernel/common_utils.h"
#include "plugin/device/gpu/kernel/kernel_constants.h"
#include "plugin/device/gpu/kernel/cuda_impl/matrix_set_diag_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/matrix_set_diag_impl.cuh"
namespace mindspore {
namespace kernel {
template <typename T>

4
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/meshgrid_gpu_kernel.h
View File

@ -22,8 +22,8 @@
#include <utility>
#include <vector>
#include "plugin/device/gpu/kernel/cuda_impl/broadcast_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/oneslike_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/broadcast_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/oneslike_impl.cuh"
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/math/broadcast_gpu_kernel.h"

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/one_hot_gpu_kernel.h
View File

@ -20,7 +20,7 @@
#include <vector>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/one_hot_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/one_hot_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/oneslike_gpu_kernel.h
View File

@ -19,7 +19,7 @@
#include <vector>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/oneslike_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops//oneslike_impl.cuh"
namespace mindspore {
namespace kernel {
template <typename T>

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/pack_gpu_kernel.h
View File

@ -22,7 +22,7 @@
#include <memory>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/pack.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/pack.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/range_gpu_kernel.h
View File

@ -19,7 +19,7 @@
#include <vector>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/range_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/range_impl.cuh"
namespace mindspore {
namespace kernel {
constexpr float kStartDefault = 0.;

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/resize_nearest_neighbor_gpu_kernel.h
View File

@ -20,7 +20,7 @@
#include <vector>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/resize_nearest_neighbor_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/resize_nearest_neighbor_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/resize_nearest_neighbor_grad_gpu_kernel.h
View File

@ -20,7 +20,7 @@
#include <vector>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/resize_nearest_neighbor_grad_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/resize_nearest_neighbor_grad_impl.cuh"
namespace mindspore {
namespace kernel {

4
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/reverse_sequence_gpu_kernel.h
View File

@ -22,8 +22,8 @@
#include <iostream>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/hal/device/cuda_common.h"
#include "plugin/device/gpu/kernel/cuda_impl/reverse_sequence_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cuda_common.h"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/reverse_sequence_impl.cuh"
#include "plugin/device/gpu/kernel/kernel_constants.h"
namespace mindspore {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/reverse_v2_gpu_kernel.h
View File

@ -22,7 +22,7 @@
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/reverse_v2_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/reverse_v2_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/scatter_functor_gpu_kernel.h
View File

@ -22,7 +22,7 @@
#include <map>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/scatter_functor_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/scatter_functor_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/scatter_nd_functor_gpu_kernel.h
View File

@ -22,7 +22,7 @@
#include <map>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/scatter_nd_functor_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/scatter_nd_functor_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/scatter_nd_gpu_kernel.h
View File

@ -19,7 +19,7 @@
#include <vector>
#include <algorithm>
#include "plugin/device/gpu/kernel/cuda_impl/scatter_nd.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/scatter_nd.cuh"
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/select_gpu_kernel.h
View File

@ -21,7 +21,7 @@
#include <string>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/select_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/select_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/slice_gpu_kernel.h
View File

@ -23,7 +23,7 @@
#include <algorithm>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/slice_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/slice_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/slice_grad_gpu_kernel.h
View File

@ -23,7 +23,7 @@
#include <utility>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/slice_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/slice_impl.cuh"
namespace mindspore {
namespace kernel {

6
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/sort_gpu_kernel.h
View File

@ -24,9 +24,9 @@
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/topk_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/transpose_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/unary_op_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/topk_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/transpose_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/unary_op_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/spacetobatch_gpu_kernel.h
View File

@ -21,7 +21,7 @@
#include <string>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/spacetobatch_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/spacetobatch_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/spacetodepth_gpu_kernel.h
View File

@ -21,7 +21,7 @@
#include <vector>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/spacetodepth_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/spacetodepth_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/split_gpu_kernel.h
View File

@ -22,7 +22,7 @@
#include <memory>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/split_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/split_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/strided_slice_gpu_kernel.h
View File

@ -23,7 +23,7 @@
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/arrays/strided_slice_gpu_common.h"
#include "plugin/device/gpu/kernel/cuda_impl/slice_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/slice_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/strided_slice_grad_gpu_kernel.h
View File

@ -23,7 +23,7 @@
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/arrays/strided_slice_gpu_common.h"
#include "plugin/device/gpu/kernel/cuda_impl/slice_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/slice_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/tensor_copy_slices_gpu_kernel.h
View File

@ -25,7 +25,7 @@
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "kernel/common_utils.h"
#include "plugin/device/gpu/kernel/cuda_impl/slice_copy_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/slice_copy_impl.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/tensor_scatter_add_gpu_kernel.h
View File

@ -20,7 +20,7 @@
#include <vector>
#include <string>
#include <algorithm>
#include "plugin/device/gpu/kernel/cuda_impl/tensor_scatter_add.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/tensor_scatter_add.cuh"
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/tensor_scatter_max_gpu_kernel.h
View File

@ -20,7 +20,7 @@
#include <vector>
#include <algorithm>
#include "plugin/device/gpu/kernel/cuda_impl/tensor_scatter_max.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/tensor_scatter_max.cuh"
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/tensor_scatter_min_gpu_kernel.h
View File

@ -19,7 +19,7 @@
#include <vector>
#include <algorithm>
#include "plugin/device/gpu/kernel/cuda_impl/tensor_scatter_min.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/tensor_scatter_min.cuh"
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/tensor_scatter_sub_gpu_kernel.h
View File

@ -20,7 +20,7 @@
#include <vector>
#include <algorithm>
#include "plugin/device/gpu/kernel/cuda_impl/tensor_scatter_sub.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/tensor_scatter_sub.cuh"
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/tensor_scatter_update_gpu_kernel.h
View File

@ -19,7 +19,7 @@
#include <vector>
#include <algorithm>
#include "plugin/device/gpu/kernel/cuda_impl/tensor_scatter_update.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/tensor_scatter_update.cuh"
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/tile_gpu_kernel.h
View File

@ -20,7 +20,7 @@
#include <vector>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/tile_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/tile_impl.cuh"
namespace mindspore {
namespace kernel {

4
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/topk_gpu_kernel.h
View File

@ -21,8 +21,8 @@
#include <vector>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/cast_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/topk_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cast_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/topk_impl.cuh"
namespace mindspore {
namespace kernel {

4
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/transpose_gpu_kernel.h
View File

@ -21,8 +21,8 @@
#include <algorithm>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/transpose_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/transpose_impl_opt.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/transpose_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/transpose_impl_opt.cuh"
namespace mindspore {
namespace kernel {
constexpr size_t kDimSize4 = 4;

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/unique_gpu_kernel.h
View File

@ -20,7 +20,7 @@
#include <vector>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/unique_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/unique_impl.cuh"
namespace mindspore {
namespace kernel {
template <typename T, typename S>

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/unpack_gpu_kernel.h
View File

@ -22,7 +22,7 @@
#include <memory>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/unpack.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/unpack.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/unsorted_segment_max_gpu_kernel.h
View File

@ -21,7 +21,7 @@
#include <limits>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/unsorted_segment_max.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/unsorted_segment_max.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/unsorted_segment_min_gpu_kernel.h
View File

@ -21,7 +21,7 @@
#include <limits>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/unsorted_segment_min.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/unsorted_segment_min.cuh"
namespace mindspore {
namespace kernel {

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/unsorted_segment_sum_gpu_kernel.h
View File

@ -20,7 +20,7 @@
#include <vector>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "plugin/device/gpu/kernel/cuda_impl/unsorted_segment_sum.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/unsorted_segment_sum.cuh"
namespace mindspore {
namespace kernel {

30
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/adagrad_impl.cuh
View File

@ -1,30 +0,0 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_ADAGRAD_IMPL_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_ADAGRAD_IMPL_H_
#include "plugin/device/gpu/hal/device/cuda_common.h"
template <typename T, typename S, typename G>
void ApplyAdagrad(const size_t size,
const bool update_slots,
const S *learning_rate,
const G *gradient,
T *variable,
T *accumulation,
cudaStream_t stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_ADAGRAD_IMPL_H_

29
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/adam_impl.cuh
View File

@ -1,29 +0,0 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_ADAM_IMPL_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_ADAM_IMPL_H_
#include "plugin/device/gpu/hal/device/cuda_common.h"
template <typename T>
void ApplyAdam(const size_t size, const T *gradient, const T *beta1_power, const T *beta2_power, const T *learning_rate,
const T *beta1, const T *beta2, const T *epsilon, T *variable, T *m, T *v, cudaStream_t cuda_stream);
template <typename T>
void AdamWeightDecayOp(const size_t size, const T *gradient, const float *learning_rate, const float *beta1,
const float *beta2, const float *epsilon, const float *decay, T *variable, T *m, T *v,
cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_ADAM_IMPL_H_

24
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/adam_weight_decay_impl.cuh
View File

@ -1,24 +0,0 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_ADAM_WEIGHT_DECAY_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_ADAM_WEIGHT_DECAY_H_
template <typename T>
void AdamWeightDecay(const int &element_num_, const bool &need_decay, const float *beta1, const float *one_sub_beta1,
const float *beta2, const float *one_sub_beta2, const float *epsilon, const float *lr,
const float *weight_decay, T *m, T *v, T *param, T *gradient, cudaStream_t stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_ADAM_WEIGHT_DECAY_H_

26
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/adaptive_avg_pool2d_grad_impl.cuh
View File

@ -1,26 +0,0 @@
/**
* Copyright 2021 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_ADAPTIVE_AVGPOOL2D_GRAD_IMPL_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_ADAPTIVE_AVGPOOL2D_GRAD_IMPL_H_
#include "plugin/device/gpu/hal/device/cuda_common.h"
template <typename T>
void ApplyAdaptiveAvgPool2DGrad(const uint size, const uint input_height, const uint input_width,
const uint output_height, const uint output_width, T *input_data,
T *output_data, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_ADAPTIVE_AVGPOOL2D_GRAD_IMPL_H_

25
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/adaptive_avg_pool2d_impl.cuh
View File

@ -1,25 +0,0 @@
/**
* Copyright 2021 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_ADAPTIVEAVGPOOL2D_IMPL_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_ADAPTIVEAVGPOOL2D_IMPL_H_
#include "plugin/device/gpu/hal/device/cuda_common.h"
template <typename T>
void ApplyAdaptiveAvgPool2D(const uint size, const uint input_height, const uint input_width, const uint output_height,
const uint output_width, T *input_data, T *output_data, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_ADAPTIVEAVGPOOL2D_IMPL_H_

27
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/add_relu_v2_impl.cuh
View File

@ -1,27 +0,0 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_ADD_RELU_V2_IMPL_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_ADD_RELU_V2_IMPL_H_
#include "plugin/device/gpu/hal/device/cuda_common.h"
template <typename T>
void AddReluV2(const size_t num, const T *x1, const T *x2, T *y, uint32_t *mask, cudaStream_t cuda_stream);
template <typename T>
void AddReluGradV2(const size_t size, const T *x1, const T *x2, const uint32_t *mask, T *dx, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_ADD_RELU_IMPL_H_

23
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/argmax_impl.cuh
View File

@ -1,23 +0,0 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_CUDA_IMPL_ARGMAX_IMPL_CUH_
#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_CUDA_IMPL_ARGMAX_IMPL_CUH_
template <typename T, typename S>
void CalArgmax(const T *input, const S bound, const size_t outer_size, const size_t inner_size, S *output,
cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_CUDA_IMPL_ARGMAX_IMPL_CUH_

40
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/batchnorm_fold2_impl.cuh
View File

@ -1,40 +0,0 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BATCHNORMFOLD2_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BATCHNORMFOLD2_H_
#include "plugin/device/gpu/hal/device/cuda_common.h"
template <typename T>
void BatchNormFold2Forward(const T *x, const T *beta, const T *gamma, const T *batch_std, const T *batch_mean,
const T *running_std, const T *running_mean, const int *global_step, T *y, int freeze_bn,
size_t N, size_t C, size_t H, size_t W, cudaStream_t cuda_stream);
template <typename T>
void CalBatchNormFold2GradNotFreeze(const T *d_beta, const T *reduce_x, const T *batch_mean, const T *batch_std,
const T *running_mean, const T *running_std, const T *gamma, T *d_gamma,
T *d_batch_mean, T *d_batch_std, size_t C, cudaStream_t cuda_stream);
template <typename T>
void CalBatchNormFold2GradFreeze(const T *d_beta, const T *reduce_x, const T *batch_mean, const T *batch_std,
const T *running_mean, const T *running_std, const T *gamma, T *d_gamma,
T *d_batch_mean, T *d_batch_std, size_t C, cudaStream_t cuda_stream);
template <typename T>
void BatchNormFold2GradReduce(const T *dout, const T *x, T *d_beta, T *tmp, T *reduce_x, T *tmp2, T *tmp_x, size_t N,
size_t C, size_t H, size_t W, cudaStream_t cuda_stream);
template <typename T>
void CalBatchNormFold2GradNotFreezeDxMul(const T *batch_std, const T *running_std, T *d_x, size_t N, size_t C, size_t H,
size_t W, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BATCHNORMFOLD2_H_

32
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/batchnorm_fold_impl.cuh
View File

@ -1,32 +0,0 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BATCHNORM_FOLD_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BATCHNORM_FOLD_H_
template <typename T>
void CalUpdateRunningStd(int channel_size, double epsilon, T* running_std, cudaStream_t cuda_stream);
template <typename T>
void CalUpdateBatchStd(int channel_size, T* batch_std, cudaStream_t cuda_stream);
template <typename T>
void CalBatchNormFoldGrad(const T* d_batch_mean, const T* d_batch_std, const T* x, const T* batch_mean,
const T* batch_std, int batch_size, int channel_size, int height, int width, T* dx,
cudaStream_t cuda_stream);
template <typename T>
void ThrustFillWith(T* array, int size, T tofill, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_BATCHNORM_FOLD_H_

133
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/batchtospace_impl.cu
View File

@ -1,133 +0,0 @@
/**
* Copyright 2021 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include <cuda_runtime.h>
#include "batchtospace_impl.cuh"
#include "plugin/device/gpu/hal/device/cuda_common.h"
template <typename T>
__global__ void BatchToSpace(const size_t size, const T *input, const size_t in,
const size_t ih, const size_t iw, const size_t ic,
const size_t on, const size_t oh, const size_t ow,
const size_t oc, const size_t crop_up, const size_t crop_dn,
const size_t crop_lft, const size_t crop_rht, const size_t block_num,
T *output) {
size_t temp_stride = 0;
size_t temp_pos = 0;
size_t idx_on = 0;
size_t idx_oc = 0;
size_t idx_oh = 0;
size_t idx_ow = 0;
size_t idx_in = 0;
size_t input_pos = 0;
for (size_t pos = blockIdx.x * blockDim.x + threadIdx.x; pos < size;
pos += blockDim.x * gridDim.x) {
temp_stride = oc * oh * ow;
idx_on = pos / temp_stride;
temp_pos = pos % temp_stride;
temp_stride /= oc;
idx_oc = temp_pos / temp_stride;
temp_pos = pos % temp_stride;
temp_stride /= oh;
idx_oh = temp_pos / temp_stride;
temp_pos = pos % temp_stride;
temp_stride /= ow;
idx_ow = temp_pos / temp_stride;
idx_in = (((idx_oh + crop_up) % block_num) * block_num + ((idx_ow + crop_lft) % block_num)) * on + idx_on;
input_pos = idx_in * ic;
input_pos = (input_pos + idx_oc) * ih;
input_pos = (input_pos + ((idx_oh + crop_up) - (idx_in / (on * block_num))) / block_num) * iw;
input_pos = (input_pos + ((idx_ow + crop_lft) - ((idx_in / on) % block_num)) / block_num);
output[pos] = input[input_pos];
}
return;
}
template <typename T>
void CalBatchToSpace(const size_t size, const T *input, const size_t in,
const size_t ih, const size_t iw, const size_t ic,
const size_t on, const size_t oh, const size_t ow,
const size_t oc, const size_t crop_up, const size_t crop_dn,
const size_t crop_lft, const size_t crop_rht, const size_t block_num,
T *output, cudaStream_t cuda_stream) {
BatchToSpace<<<GET_BLOCKS(size), GET_THREADS, 0, cuda_stream>>>(
size, input, in, ih, iw, ic, on, oh, ow, oc, crop_up, crop_dn, crop_lft, crop_rht, block_num, output);
return;
}
template void CalBatchToSpace<float>(const size_t size, const float *input, const size_t in,
const size_t ih, const size_t iw, const size_t ic,
const size_t on, const size_t oh, const size_t ow,
const size_t oc, const size_t crop_up, const size_t crop_dn,
const size_t crop_lft, const size_t crop_rht, const size_t block_num,
float *output, cudaStream_t cuda_stream);
template void CalBatchToSpace<half>(const size_t size, const half *input, const size_t in,
const size_t ih, const size_t iw, const size_t ic,
const size_t on, const size_t oh, const size_t ow,
const size_t oc, const size_t crop_up, const size_t crop_dn,
const size_t crop_lft, const size_t crop_rht, const size_t block_num,
half *output, cudaStream_t cuda_stream);
template void CalBatchToSpace<int>(const size_t size, const int *input, const size_t in,
const size_t ih, const size_t iw, const size_t ic,
const size_t on, const size_t oh, const size_t ow,
const size_t oc, const size_t crop_up, const size_t crop_dn,
const size_t crop_lft, const size_t crop_rht, const size_t block_num,
int *output, cudaStream_t cuda_stream);
template void CalBatchToSpace<int64_t>(const size_t size, const int64_t *input, const size_t in,
const size_t ih, const size_t iw, const size_t ic,
const size_t on, const size_t oh, const size_t ow,
const size_t oc, const size_t crop_up, const size_t crop_dn,
const size_t crop_lft, const size_t crop_rht, const size_t block_num,
int64_t *output, cudaStream_t cuda_stream);
template void CalBatchToSpace<int16_t>(const size_t size, const int16_t *input, const size_t in,
const size_t ih, const size_t iw, const size_t ic,
const size_t on, const size_t oh, const size_t ow,
const size_t oc, const size_t crop_up, const size_t crop_dn,
const size_t crop_lft, const size_t crop_rht, const size_t block_num,
int16_t *output, cudaStream_t cuda_stream);
template void CalBatchToSpace<int8_t>(const size_t size, const int8_t *input, const size_t in,
const size_t ih, const size_t iw, const size_t ic,
const size_t on, const size_t oh, const size_t ow,
const size_t oc, const size_t crop_up, const size_t crop_dn,
const size_t crop_lft, const size_t crop_rht, const size_t block_num,
int8_t *output, cudaStream_t cuda_stream);
template void CalBatchToSpace<uint8_t>(const size_t size, const uint8_t *input, const size_t in,
const size_t ih, const size_t iw, const size_t ic,
const size_t on, const size_t oh, const size_t ow,
const size_t oc, const size_t crop_up, const size_t crop_dn,
const size_t crop_lft, const size_t crop_rht, const size_t block_num,
uint8_t *output, cudaStream_t cuda_stream);
template void CalBatchToSpace<uint16_t>(const size_t size, const uint16_t *input, const size_t in,
const size_t ih, const size_t iw, const size_t ic,
const size_t on, const size_t oh, const size_t ow,
const size_t oc, const size_t crop_up, const size_t crop_dn,
const size_t crop_lft, const size_t crop_rht, const size_t block_num,
uint16_t *output, cudaStream_t cuda_stream);
template void CalBatchToSpace<uint32_t>(const size_t size, const uint32_t *input, const size_t in,
const size_t ih, const size_t iw, const size_t ic,
const size_t on, const size_t oh, const size_t ow,
const size_t oc, const size_t crop_up, const size_t crop_dn,
const size_t crop_lft, const size_t crop_rht, const size_t block_num,
uint32_t *output, cudaStream_t cuda_stream);
template void CalBatchToSpace<uint64_t>(const size_t size, const uint64_t *input, const size_t in,
const size_t ih, const size_t iw, const size_t ic,
const size_t on, const size_t oh, const size_t ow,
const size_t oc, const size_t crop_up, const size_t crop_dn,
const size_t crop_lft, const size_t crop_rht, const size_t block_num,
uint64_t *output, cudaStream_t cuda_stream);

27
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/batchtospace_impl.cuh
View File

@ -1,27 +0,0 @@
/**
* Copyright 2021 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BATCHTOSPACE_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BATCHTOSPACE_H_
template <typename T>
void CalBatchToSpace(const size_t size, const T *input, const size_t in,
const size_t ih, const size_t iw, const size_t ic,
const size_t on, const size_t oh, const size_t ow,
const size_t oc, const size_t crop_up, const size_t crop_dn,
const size_t crop_lft, const size_t crop_rht, const size_t block_num,
T *output, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BATCHTOSPACE_H_

30
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/bce_with_logits_loss_impl.cuh
View File

@ -1,30 +0,0 @@
/**
* Copyright 2021 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_CUDA_IMPL_BCE_WITH_LOGITS_LOSS_IMPL_CUH_
#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_CUDA_IMPL_BCE_WITH_LOGITS_LOSS_IMPL_CUH_
#define MAX_LOGITS_DIMENSION 8
#include "plugin/device/gpu/hal/device/cuda_common.h"
template <typename T>
void CalBCEWithLogitsLoss(const size_t input_size, const T *predict, const T *target, const size_t *input_shape,
const size_t shape_size, const T *weight, const size_t *weight_shape,
const bool weight_need_broadcast, const T *pos_weight, const size_t *pos_weight_shape,
const bool pos_weight_need_broadcast, T *shape_broadcasted, T *output,
cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_CUDA_IMPL_BCE_WITH_LOGITS_LOSS_IMPL_CUH_

27
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/bias_add_grad_impl.cuh
View File

@ -1,27 +0,0 @@
/**
* Copyright 2021 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BIASADDGRAD_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BIASADDGRAD_H_
#include "plugin/device/gpu/hal/device/cuda_common.h"
template <typename T>
void CalBiasAddGradNHWC(const size_t size, const size_t bias_size,
const T* dy, T* db, cudaStream_t cuda_stream);
template <typename T>
void CalBiasAddGradNCHW(const size_t size, const size_t bias_size, const int height, const int width,
const T* dy, T* db, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BIASADDGRAD_H_

27
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/boundingbox_decode_impl.cuh
View File

@ -1,27 +0,0 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_BOUNDINGBOX_DECODE_IMPL_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_BOUNDINGBOX_DECODE_IMPL_H_
#include "plugin/device/gpu/hal/device/cuda_common.h"
template <typename T>
void BoundingBoxDecode(const size_t size, const T *rois, const T *deltas, T *bboxes, const float &m1, const float &m2,
const float &m3, const float &m4, const float &s1, const float &s2, const float &s3,
const float &s4, const int &max_height, const int &max_width, const float &ratio_clip,
cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_BOUNDINGBOX_DECODE_IMPL_H_

26
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/boundingbox_encode_impl.cuh
View File

@ -1,26 +0,0 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_BOUNDINGBOX_ENCODE_IMPL_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_BOUNDINGBOX_ENCODE_IMPL_H_
#include "plugin/device/gpu/hal/device/cuda_common.h"
template <typename T>
void BoundingBoxEncode(const size_t size, const T *anchor_box, const T *groundtruth_box, T *deltas, const float &m1,
const float &m2, const float &m3, const float &m4, const float &s1, const float &s2,
const float &s3, const float &s4, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_BOUNDINGBOX_ENCODE_IMPL_H_

38
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/broadcast_grad_impl.cuh
View File

@ -1,38 +0,0 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BROADCAST_GRAD_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BROADCAST_GRAD_H_
#include "plugin/device/gpu/hal/device/cuda_common.h"
enum BroadcastGradOpType {
BROADCAST_GRAD_TYPE_MAXIMUM = 0,
BROADCAST_GRAD_TYPE_MINIMUM = 1,
BROADCAST_GRAD_TYPE_INVALID = 0xffffffff,
};
template <typename T>
void BroadcastGrad(const int &l0, const int &l1, const int &l2, const int &l3, const int &r0, const int &r1,
const int &r2, const int &r3, const int &d0, const int &d1, const int &d2, const int &d3,
const bool &grad_x1, const bool &grad_x2, enum BroadcastGradOpType op, const T *x1, const T *x2,
const T *dy, T *dx1, T *dx2, cudaStream_t stream);
template <typename T>
void NoBroadcastGrad(const int &nums, const bool &grad_x1, const bool &grad_x2, enum BroadcastGradOpType op,
const T *x1, const T *x2, const T *dy, T *dx1, T *dx2, cudaStream_t stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BROADCAST_GRAD_H_

89
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/broadcast_impl.cuh
View File

@ -1,89 +0,0 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BROADCAST_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BROADCAST_H_
#include <vector>
#include "plugin/device/gpu/hal/device/cuda_common.h"
#include "utils/complex.h"
const float kFloatEplison = 1e-37;
enum BroadcastOpType {
BROADCAST_TYPE_GREATER = 0,
BROADCAST_TYPE_LESS = 1,
BROADCAST_TYPE_MAXIMUM = 2,
BROADCAST_TYPE_MINIMUM = 3,
BROADCAST_TYPE_POWER = 4,
BROADCAST_TYPE_REALDIV = 5,
BROADCAST_TYPE_MUL = 6,
BROADCAST_TYPE_SUB = 7,
BROADCAST_TYPE_ADD = 8,
BROADCAST_TYPE_FLOORDIV = 9,
BROADCAST_TYPE_ABSGRAD = 10,
BROADCAST_TYPE_DIV = 11,
BROADCAST_TYPE_DIVNONAN = 12,
BROADCAST_TYPE_EQUAL = 13,
BROADCAST_TYPE_SQUARED_DIFFERENCE = 14,
BROADCAST_TYPE_MOD = 15,
BROADCAST_TYPE_FLOORMOD = 16,
BROADCAST_TYPE_ATAN2 = 17,
BROADCAST_TYPE_GREATER_EQUAL = 18,
BROADCAST_TYPE_LESS_EQUAL = 19,
BROADCAST_TYPE_NOT_EQUAL = 20,
BROADCAST_TYPE_LOGICAL_AND = 21,
BROADCAST_TYPE_LOGICAL_OR = 22,
BROADCAST_TYPE_TRUNCATEDIV = 23,
BROADCAST_TYPE_TRUNCATEMOD = 24,
BROADCAST_TYPE_COMPLEX = 25,
BROADCAST_TYPE_INVALID = 0xffffffff,
};
template <typename T>
void ElewiseCmp(const int &nums, enum BroadcastOpType op, const T *x0, const T *x1, bool *y, cudaStream_t stream);
template <typename T>
void ElewiseArith(const int &nums, enum BroadcastOpType op, const T *x0, const T *x1, T *y, cudaStream_t stream);
template <typename T1, typename T2, typename T3>
void ElewiseComplexArith(const int &nums, enum BroadcastOpType op, const T1 *x0, const T2 *x1,
Complex<T3> *y, cudaStream_t stream);
template <typename T>
void BroadcastCmp(const std::vector<size_t> &x0_dims, const std::vector<size_t> &x1_dims,
const std::vector<size_t> &y_dims, enum BroadcastOpType op, const T *x0, const T *x1, bool *y,
cudaStream_t stream);
template <typename T>
void BroadcastArith(const std::vector<size_t> &x0_dims, const std::vector<size_t> &x1_dims,
const std::vector<size_t> &y_dims, enum BroadcastOpType op, const T *x0, const T *x1, T *y,
cudaStream_t stream);
template <typename T1, typename T2, typename T3>
void BroadcastComplexArith(const std::vector<size_t> &x0_dims, const std::vector<size_t> &x1_dims,
const std::vector<size_t> &y_dims, enum BroadcastOpType op, const T1 *x0, const T2 *x1,
Complex<T3> *y, cudaStream_t stream);
template <typename T>
void BroadcastComplexArith(const std::vector<size_t> &x0_dims, const std::vector<size_t> &x1_dims,
const std::vector<size_t> &y_dims, enum BroadcastOpType op, const T *x0, const T *x1,
Complex<T> *y, cudaStream_t stream);
template <typename T>
void BroadcastTo(const size_t &i0, const size_t &i1, const size_t &i2, const size_t &i3, const size_t &o0,
const size_t &o1, const size_t &o2, const size_t &o3, const T *input_addr, T *output_addr,
cudaStream_t stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_BROADCAST_H_

318
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cast_impl.cu
View File

@ -1,318 +0,0 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include <vector>
#include <iostream>
#include "plugin/device/gpu/kernel/cuda_impl/cast_impl.cuh"
#include "plugin/device/gpu/hal/device/cuda_common.h"
// Generic cast
template <typename S, typename T>
__device__ __forceinline__ void CastBase(const S *input_addr, T *output_addr) {
*output_addr = static_cast<T>((*input_addr));
}
// half --> integer
__device__ __forceinline__ void CastBase(const half *input_addr, uint64_t *output_addr) {
*output_addr = __half2ull_rz((*input_addr));
}
__device__ __forceinline__ void CastBase(const half *input_addr, int64_t *output_addr) {
*output_addr = __half2ll_rz((*input_addr));
}
__device__ __forceinline__ void CastBase(const half *input_addr, uint32_t *output_addr) {
*output_addr = __half2uint_rz((*input_addr));
}
__device__ __forceinline__ void CastBase(const half *input_addr, int32_t *output_addr) {
*output_addr = __half2int_rz((*input_addr));
}
__device__ __forceinline__ void CastBase(const half *input_addr, uint16_t *output_addr) {
*output_addr = __half2ushort_rz((*input_addr));
}
__device__ __forceinline__ void CastBase(const half *input_addr, int16_t *output_addr) {
*output_addr = __half2short_rz((*input_addr));
}
__device__ __forceinline__ void CastBase(const half *input_addr, uint8_t *output_addr) {
*output_addr = static_cast<uint8_t>(__half2ushort_rz((*input_addr)));
}
__device__ __forceinline__ void CastBase(const half *input_addr, int8_t *output_addr) {
*output_addr = static_cast<int8_t>(__half2short_rz((*input_addr)));
}
// integer --> half
__device__ __forceinline__ void CastBase(const uint64_t *input_addr, half *output_addr) {
*output_addr = __ull2half_rn((*input_addr));
}
__device__ __forceinline__ void CastBase(const int64_t *input_addr, half *output_addr) {
*output_addr = __ll2half_rn((*input_addr));
}
__device__ __forceinline__ void CastBase(const uint32_t *input_addr, half *output_addr) {
*output_addr = __uint2half_rn((*input_addr));
}
__device__ __forceinline__ void CastBase(const int32_t *input_addr, half *output_addr) {
*output_addr = __int2half_rn((*input_addr));
}
__device__ __forceinline__ void CastBase(const uint16_t *input_addr, half *output_addr) {
*output_addr = __ushort2half_rn((*input_addr));
}
__device__ __forceinline__ void CastBase(const int16_t *input_addr, half *output_addr) {
*output_addr = __short2half_rn((*input_addr));
}
__device__ __forceinline__ void CastBase(const uint8_t *input_addr, half *output_addr) {
*output_addr = __ushort2half_rn(static_cast<uint16_t>(*input_addr));
}
__device__ __forceinline__ void CastBase(const int8_t *input_addr, half *output_addr) {
*output_addr = __short2half_rn(static_cast<int16_t>(*input_addr));
}
// Cast
template <typename S, typename T>
__global__ void CastKernel(const int input_size, const S *input_addr, T *output_addr) {
for (size_t pos = blockIdx.x * blockDim.x + threadIdx.x; pos < input_size; pos += blockDim.x * gridDim.x) {
CastBase(input_addr + pos, output_addr + pos);
}
}
template <typename S, typename T>
void Cast(const int input_size, const S *input_addr, T *output_addr, cudaStream_t stream) {
CastKernel<<<GET_BLOCKS(input_size), GET_THREADS, 0, stream>>>(input_size, input_addr, output_addr);
}
template void Cast(const int input_size, const int8_t *input_addr, int8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int8_t *input_addr, int16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int8_t *input_addr, int32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int8_t *input_addr, int64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int8_t *input_addr, uint8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int8_t *input_addr, uint16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int8_t *input_addr, uint32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int8_t *input_addr, uint64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int8_t *input_addr, float *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int8_t *input_addr, double *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int8_t *input_addr, half *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int8_t *input_addr, bool *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int8_t *input_addr, Complex<float> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int8_t *input_addr, Complex<double> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int16_t *input_addr, int8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int16_t *input_addr, int16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int16_t *input_addr, int32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int16_t *input_addr, int64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int16_t *input_addr, uint8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int16_t *input_addr, uint16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int16_t *input_addr, uint32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int16_t *input_addr, uint64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int16_t *input_addr, float *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int16_t *input_addr, double *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int16_t *input_addr, half *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int16_t *input_addr, bool *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int16_t *input_addr, Complex<float> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int16_t *input_addr, Complex<double> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int32_t *input_addr, int8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int32_t *input_addr, int16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int32_t *input_addr, int32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int32_t *input_addr, int64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int32_t *input_addr, uint8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int32_t *input_addr, uint16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int32_t *input_addr, uint32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int32_t *input_addr, uint64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int32_t *input_addr, float *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int32_t *input_addr, double *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int32_t *input_addr, half *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int32_t *input_addr, bool *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int32_t *input_addr, Complex<float> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int32_t *input_addr, Complex<double> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int64_t *input_addr, int8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int64_t *input_addr, int16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int64_t *input_addr, int32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int64_t *input_addr, int64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int64_t *input_addr, uint8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int64_t *input_addr, uint16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int64_t *input_addr, uint32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int64_t *input_addr, uint64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int64_t *input_addr, float *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int64_t *input_addr, double *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int64_t *input_addr, half *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int64_t *input_addr, bool *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int64_t *input_addr, Complex<float> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const int64_t *input_addr, Complex<double> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint8_t *input_addr, int8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint8_t *input_addr, int16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint8_t *input_addr, int32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint8_t *input_addr, int64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint8_t *input_addr, uint8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint8_t *input_addr, uint16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint8_t *input_addr, uint32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint8_t *input_addr, uint64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint8_t *input_addr, float *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint8_t *input_addr, double *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint8_t *input_addr, half *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint8_t *input_addr, bool *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint8_t *input_addr, Complex<float> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint8_t *input_addr, Complex<double> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint16_t *input_addr, int8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint16_t *input_addr, int16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint16_t *input_addr, int32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint16_t *input_addr, int64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint16_t *input_addr, uint8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint16_t *input_addr, uint16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint16_t *input_addr, uint32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint16_t *input_addr, uint64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint16_t *input_addr, float *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint16_t *input_addr, double *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint16_t *input_addr, half *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint16_t *input_addr, bool *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint16_t *input_addr, Complex<float> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint16_t *input_addr, Complex<double> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint32_t *input_addr, int8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint32_t *input_addr, int16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint32_t *input_addr, int32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint32_t *input_addr, int64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint32_t *input_addr, uint8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint32_t *input_addr, uint16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint32_t *input_addr, uint32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint32_t *input_addr, uint64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint32_t *input_addr, float *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint32_t *input_addr, double *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint32_t *input_addr, half *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint32_t *input_addr, bool *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint32_t *input_addr, Complex<float> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint32_t *input_addr, Complex<double> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint64_t *input_addr, int8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint64_t *input_addr, int16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint64_t *input_addr, int32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint64_t *input_addr, int64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint64_t *input_addr, uint8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint64_t *input_addr, uint16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint64_t *input_addr, uint32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint64_t *input_addr, uint64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint64_t *input_addr, float *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint64_t *input_addr, double *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint64_t *input_addr, half *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint64_t *input_addr, bool *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint64_t *input_addr, Complex<float> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const uint64_t *input_addr, Complex<double> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const half *input_addr, int8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const half *input_addr, int16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const half *input_addr, int32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const half *input_addr, int64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const half *input_addr, uint8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const half *input_addr, uint16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const half *input_addr, uint32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const half *input_addr, uint64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const half *input_addr, float *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const half *input_addr, double *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const half *input_addr, half *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const half *input_addr, bool *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const half *input_addr, Complex<float> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const half *input_addr, Complex<double> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const float *input_addr, int8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const float *input_addr, int16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const float *input_addr, int32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const float *input_addr, int64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const float *input_addr, uint8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const float *input_addr, uint16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const float *input_addr, uint32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const float *input_addr, uint64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const float *input_addr, float *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const float *input_addr, double *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const float *input_addr, half *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const float *input_addr, bool *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const float *input_addr, Complex<float> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const float *input_addr, Complex<double> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const double *input_addr, int8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const double *input_addr, int16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const double *input_addr, int32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const double *input_addr, int64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const double *input_addr, uint8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const double *input_addr, uint16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const double *input_addr, uint32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const double *input_addr, uint64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const double *input_addr, float *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const double *input_addr, double *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const double *input_addr, half *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const double *input_addr, bool *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const double *input_addr, Complex<float> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const double *input_addr, Complex<double> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const bool *input_addr, int8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const bool *input_addr, int16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const bool *input_addr, int32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const bool *input_addr, int64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const bool *input_addr, uint8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const bool *input_addr, uint16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const bool *input_addr, uint32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const bool *input_addr, uint64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const bool *input_addr, float *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const bool *input_addr, double *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const bool *input_addr, half *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const bool *input_addr, bool *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const bool *input_addr, Complex<float> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const bool *input_addr, Complex<double> *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<float> *input_addr, int8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<float> *input_addr, int16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<float> *input_addr, int32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<float> *input_addr, int64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<float> *input_addr, uint8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<float> *input_addr, uint16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<float> *input_addr, uint32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<float> *input_addr, uint64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<float> *input_addr, float *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<float> *input_addr, double *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<float> *input_addr, half *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<float> *input_addr, bool *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<float> *input_addr, Complex<double> *output_addr,
cudaStream_t stream);
template void Cast(const int input_size, const Complex<double> *input_addr, int8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<double> *input_addr, int16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<double> *input_addr, int32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<double> *input_addr, int64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<double> *input_addr, uint8_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<double> *input_addr, uint16_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<double> *input_addr, uint32_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<double> *input_addr, uint64_t *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<double> *input_addr, float *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<double> *input_addr, double *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<double> *input_addr, half *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<double> *input_addr, bool *output_addr, cudaStream_t stream);
template void Cast(const int input_size, const Complex<double> *input_addr, Complex<float> *output_addr,
cudaStream_t stream);

29
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/clip_grad_norm_impl.cuh
View File

@ -1,29 +0,0 @@
/**
* Copyright 2021 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_CLIP_GRAD_NORM_IMPL_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_CLIP_GRAD_NORM_IMPL_H_
#include "plugin/device/gpu/hal/device/cuda_common.h"
template <typename T>
void ScalingGradOp(const size_t size, const T *x, const float *scaling_factor, float *scaling_out_addr,
cudaStream_t cuda_stream);
template <typename T>
void ClipGradNormOp(const size_t size, const float *x, const T *clip_norm, const float *reduce_sum_value,
float *output_addr, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_CLIP_GRAD_NORM_IMPL_H_

92
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/concatv2_impl.cu
View File

@ -1,92 +0,0 @@
/**
* Copyright 2019-2021 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include <stdio.h>
#include <stdint.h>
#include <cuda_runtime.h>
#include "plugin/device/gpu/kernel/cuda_impl/concatv2_impl.cuh"
template <typename T>
__global__ void Concat(const size_t size, const int input_num, const int all_size_before_axis, const int all_size_axis,
int *len_axis, T **inputs, T *output) {
for (int pos = blockIdx.x * blockDim.x + threadIdx.x; pos < (size); pos += blockDim.x * gridDim.x) {
int num = pos % all_size_before_axis / all_size_axis;
int block = -1;
int axis_inc = 0;
int block_len = 0;
for (int i = 0; i < input_num; i++) {
if (axis_inc <= num) {
block++;
axis_inc += len_axis[i];
} else {
break;
}
}
block_len = len_axis[block];
axis_inc -= len_axis[block];
int block_pos =
pos / all_size_before_axis * block_len * all_size_axis + (num - axis_inc) * all_size_axis + pos % all_size_axis;
output[pos] = inputs[block][block_pos];
}
return;
}
template <typename T>
void ConcatKernel(const size_t size, const int input_num, const int all_size_before_axis, const int all_size_axis,
int *len_axis, T **inputs, T *output, cudaStream_t cuda_stream) {
Concat<<<GET_BLOCKS(size), GET_THREADS, 0, cuda_stream>>>(size, input_num, all_size_before_axis, all_size_axis,
len_axis, inputs, output);
return;
}
template void ConcatKernel(const size_t size, const int input_num, const int all_size_before_axis,
const int all_size_axis, int *len_axis, double **inputs, double *output,
cudaStream_t cuda_stream);
template void ConcatKernel(const size_t size, const int input_num, const int all_size_before_axis,
const int all_size_axis, int *len_axis, float **inputs, float *output,
cudaStream_t cuda_stream);
template void ConcatKernel(const size_t size, const int input_num, const int all_size_before_axis,
const int all_size_axis, int *len_axis, half **inputs, half *output,
cudaStream_t cuda_stream);
template void ConcatKernel(const size_t size, const int input_num, const int all_size_before_axis,
const int all_size_axis, int *len_axis, int64_t **inputs, int64_t *output,
cudaStream_t cuda_stream);
template void ConcatKernel(const size_t size, const int input_num, const int all_size_before_axis,
const int all_size_axis, int *len_axis, int **inputs, int *output,
cudaStream_t cuda_stream);
template void ConcatKernel(const size_t size, const int input_num, const int all_size_before_axis,
const int all_size_axis, int *len_axis, short **inputs, short *output, // NOLINT
cudaStream_t cuda_stream);
template void ConcatKernel(const size_t size, const int input_num, const int all_size_before_axis,
const int all_size_axis, int *len_axis, char **inputs, char *output,
cudaStream_t cuda_stream);
template void ConcatKernel(const size_t size, const int input_num, const int all_size_before_axis,
const int all_size_axis, int *len_axis, uint64_t **inputs, uint64_t *output,
cudaStream_t cuda_stream);
template void ConcatKernel(const size_t size, const int input_num, const int all_size_before_axis,
const int all_size_axis, int *len_axis, uint32_t **inputs, uint32_t *output,
cudaStream_t cuda_stream);
template void ConcatKernel(const size_t size, const int input_num, const int all_size_before_axis,
const int all_size_axis, int *len_axis, uint16_t **inputs, uint16_t *output,
cudaStream_t cuda_stream);
template void ConcatKernel(const size_t size, const int input_num, const int all_size_before_axis,
const int all_size_axis, int *len_axis, unsigned char **inputs, unsigned char *output,
cudaStream_t cuda_stream);
template void ConcatKernel(const size_t size, const int input_num, const int all_size_before_axis,
const int all_size_axis, int *len_axis, bool **inputs, bool *output,
cudaStream_t cuda_stream);

34
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/convert_gradient_impl.cuh
View File

@ -1,34 +0,0 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_CONVERTGRADIENT_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_CONVERTGRADIENT_H_
#include "plugin/device/gpu/hal/device/cuda_common.h"
template <typename T>
void ConvertGradient(const size_t size, const size_t height_h, const size_t height_w, const size_t batchwidth,
const size_t width, T *input_addr, T *outt_addr, cudaStream_t cuda_stream);
template <typename T>
void ConvertGradientBack(const size_t size, const size_t height_h, const size_t height_w, const size_t batchwidth,
const size_t width, T *input_addr, T *output_addr, cudaStream_t cuda_stream);
template <typename T>
void ConvertGradientBack(const size_t size, const size_t height_h, const size_t height_w, const size_t ori_h,
const size_t ori_w, const size_t batchwidth, const size_t width, T *input_addr, T *output_addr,
cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_CONVERTGRADIENT_H_

27
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/correction_mul_impl.cuh
View File

@ -1,27 +0,0 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_CORRECTIONMUL_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_CORRECTIONMUL_H_
template <typename T>
void CalCorrectionMul(const T* weight, const T* gamma, const T* running_std, int batch_size, int channel_size,
int height, int width, T* output, cudaStream_t cuda_stream);
template <typename T>
void CalCorrectionMulGrad(const T* d_out, const T* weight, const T* running_std, int batch_size, int channel_size,
int height, int width, T* d_gamma, T* tmp, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_CORRECTIONMUL_H_

25
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/crop_and_resize_impl.cuh
View File

@ -1,25 +0,0 @@
/**
* Copyright 2021 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_CUDA_IMPL_CROP_AND_RESIZE_IMPL_H_
#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_CUDA_IMPL_CROP_AND_RESIZE_IMPL_H_
#include <cuda_runtime.h>
#include "plugin/device/gpu/hal/device/cuda_common.h"
template <typename T>
void CalCropAndResize(const size_t size, const T *input_image, float *input_boxes, int *input_box_index, int batch,
int input_height, int input_width, int final_height, int final_width, int channel,
int method, float extrapol_val, float *output, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_CUDA_IMPL_CROP_AND_RESIZE_IMPL_H_

36
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cross_entropy_impl.cuh
View File

@ -1,36 +0,0 @@
/**
* Copyright 2019 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_CROSSENTROPY_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_CROSSENTROPY_H_
#include "plugin/device/gpu/hal/device/cuda_common.h"
// The batch size limit to judge whether to use multiple threads.
constexpr int kLargeBatchLowLimit = 32768;
template <typename T, typename S>
void CrossEntropyWithSparse(const T *logits, const S *labels, const size_t batch_size, const size_t class_num, T *loss,
cudaStream_t cuda_stream);
template <typename T, typename S>
void CrossEntropyGradWithSparse(const T *logits, const S *labels, const size_t batch_size, const size_t class_num,
T *grad, cudaStream_t cuda_stream);
template <typename T, typename S>
void CrossEntropy(const T *logits, const S *labels, const size_t batch_size, const size_t class_num, T *losses,
T *dlogits, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_CROSSENTROPY_H_

51
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/ctcloss_impl.cuh
View File

@ -1,51 +0,0 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_IMPL_CUH
#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_IMPL_CUH
template <typename T>
void CalculateFwdVar(T *log_alpha_b, int *label_value_with_blank, T *softmax_probs, const int *sequence_length,
bool ctc_merge_repeated, int batch, int SOffSet, int maxtime, int blank, int *label_squence_length,
int *cum_labels_length, bool ignore_longer_outputs_than_inputs, cudaStream_t stream);
template <typename T>
void CalculateBwdVar(T *log_beta_b, int *label_value_with_blank, T *softmax_probs, const int *sequence_length,
bool ctc_merge_repeated, int batch, int SOffSet, int maxtime, int blank, int *label_squence_length,
int *cum_labels_length, bool ignore_longer_outputs_than_inputs, cudaStream_t stream);
template <typename T>
void InnerSoftMax(const T *probs, T *softmax_cost, const int *sequence_length, int max_time, int batch, int numclass,
cudaStream_t stream);
void GenLabelValuePCR(int *label_value_sp, int *label_value_pcr, int *label_squence_length, int *cum_labels_length,
int *max_labels_length, int batch, cudaStream_t stream);
void GenLabelWithBlank(int *label_value, int *label_value_with_blank, int *label_squence_length,
int *precum_labels_length, int *cum_labels_length, int batch, int blank, cudaStream_t stream);
void GenLabelValue(int *label_value_sp, const int64_t *label_indices, const int *label_values,
int *label_squence_length, int *cum_labels_length, int *max_labels_length, int size, int blank,
int batch, cudaStream_t stream);
void CalculatePreLength(int *label_squence_length, int *precum_labels_length, int *cum_labels_length,
int *max_labels_length, const int64_t *label_indices, int batch, int size, cudaStream_t stream);
void CalculateMaxSequence(const int *sequence_length, int *max_labels_length, int batch, cudaStream_t stream);
template <typename T>
void CTCLoss(T *log_alpha_b, T *log_beta_b, T *softmax_probs, int *label_value_with_blank, int batch, int SOffSet,
int maxtime, int numclass, const int *sequence_length, int *label_squence_length, int *cum_labels_length,
T *cost, T *grads, T *prob_num, bool ignore_longer_outputs_than_inputs, cudaStream_t stream);
#endif // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_IMPL_CUH

27
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/CMakeLists.txt Normal file
View File

@ -0,0 +1,27 @@
file(GLOB_RECURSE CUDA_OPS_SRC_LIST RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} "*.cu")
if(CMAKE_SYSTEM_NAME MATCHES "Darwin")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-delete-non-abstract-non-virtual-dtor -Wno-overloaded-virtual")
endif()
if(${CUDA_VERSION} VERSION_LESS 11.0)
string(REPLACE "-std=c++17" "-std=c++11" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}")
else()
string(REPLACE "-std=c++17" "-std=c++14" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}")
endif()
set_property(SOURCE ${CUDA_OPS_SRC_LIST} PROPERTY COMPILE_DEFINITIONS SUBMODULE_ID=mindspore::SubModuleId::SM_KERNEL)
if(ENABLE_GPU)
add_library(cuda_common_obj OBJECT cuda_common.cc)
target_compile_options(cuda_common_obj PRIVATE "-std=c++17")
cuda_add_library(cuda_ops SHARED ${CUDA_OPS_SRC_LIST} $<TARGET_OBJECTS:cuda_common_obj>)
message("add gpu lib to cuda_ops")
target_link_libraries(cuda_ops mindspore_core
${CUDA_PATH}/lib64/libcurand.so
${CUDNN_LIBRARY_PATH}
${CUDA_PATH}/lib64/libcudart.so
${CUDA_PATH}/lib64/stubs/libcuda.so
${CUDA_PATH}/lib64/libcusolver.so
${CUDA_PATH}/lib64/libcufft.so
${CUDA_PATH}/lib64/libcublas.so)
endif()

87
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/adagrad_impl.cu → mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/adagrad_impl.cu
View File

@ -14,7 +14,8 @@
* limitations under the License.
*/
#include "plugin/device/gpu/kernel/cuda_impl/adagrad_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/adagrad_impl.cuh"
#include "include/cuda_fp16.h"
template <typename T>
__device__ __forceinline__ T SqrtFunc(T input) {
@ -113,50 +114,50 @@ void ApplyAdagrad(const size_t size,
size, update_slots, learning_rate, gradient, variable, accumulation);
}
template void ApplyAdagrad<float, float, float>(const size_t size,
const bool update_slots,
const float *learning_rate,
const float *gradient,
float *variable,
float *accumulation,
cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void ApplyAdagrad<float, float, float>(const size_t size,
const bool update_slots,
const float *learning_rate,
const float *gradient,
float *variable,
float *accumulation,
cudaStream_t cuda_stream);
template void ApplyAdagrad<half, half, half>(const size_t size,
const bool update_slots,
const half *learning_rate,
const half *gradient,
half *variable,
half *accumulation,
cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void ApplyAdagrad<half, half, half>(const size_t size,
const bool update_slots,
const half *learning_rate,
const half *gradient,
half *variable,
half *accumulation,
cudaStream_t cuda_stream);
template void ApplyAdagrad<half, float, half>(const size_t size,
const bool update_slots,
const float *learning_rate,
const half *gradient,
half *variable,
half *accumulation,
cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void ApplyAdagrad<half, float, half>(const size_t size,
const bool update_slots,
const float *learning_rate,
const half *gradient,
half *variable,
half *accumulation,
cudaStream_t cuda_stream);
template void ApplyAdagrad<float, float, half>(const size_t size,
const bool update_slots,
const float *learning_rate,
const half *gradient,
float *variable,
float *accumulation,
cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void ApplyAdagrad<float, float, half>(const size_t size,
const bool update_slots,
const float *learning_rate,
const half *gradient,
float *variable,
float *accumulation,
cudaStream_t cuda_stream);
template void ApplyAdagrad<float, half, float>(const size_t size,
const bool update_slots,
const half *learning_rate,
const float *gradient,
float *variable,
float *accumulation,
cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void ApplyAdagrad<float, half, float>(const size_t size,
const bool update_slots,
const half *learning_rate,
const float *gradient,
float *variable,
float *accumulation,
cudaStream_t cuda_stream);
template void ApplyAdagrad<half, float, float>(const size_t size,
const bool update_slots,
const float *learning_rate,
const float *gradient,
half *variable,
half *accumulation,
cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void ApplyAdagrad<half, float, float>(const size_t size,
const bool update_slots,
const float *learning_rate,
const float *gradient,
half *variable,
half *accumulation,
cudaStream_t cuda_stream);

29
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/adagrad_impl.cuh Normal file
View File

@ -0,0 +1,29 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADAGRAD_IMPL_CUH_
#define MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADAGRAD_IMPL_CUH_
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cuda_common.h"
template <typename T, typename S, typename G>
CUDA_LIB_EXPORT void ApplyAdagrad(const size_t size,
const bool update_slots,
const S *learning_rate,
const G *gradient,
T *variable,
T *accumulation,
cudaStream_t stream);
#endif // MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADAGRAD_IMPL_CUH_

32
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/adam_impl.cu → mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/adam_impl.cu
View File

@ -14,7 +14,8 @@
* limitations under the License.
*/
#include "plugin/device/gpu/kernel/cuda_impl/adam_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/adam_impl.cuh"
#include "include/cuda_fp16.h"
template <typename T>
__device__ __forceinline__ T SqrtFunc(T input) {
@ -82,16 +83,19 @@ void AdamWeightDecayOp(const size_t size, const T *gradient, const float *learni
epsilon, decay, variable, m, v);
}
template void ApplyAdam<float>(const size_t size, const float *gradient, const float *beta1_power,
const float *beta2_power, const float *learning_rate, const float *beta1,
const float *beta2, const float *epsilon, float *variable, float *m, float *v,
cudaStream_t cuda_stream);
template void ApplyAdam<half>(const size_t size, const half *gradient, const half *beta1_power, const half *beta2_power,
const half *learning_rate, const half *beta1, const half *beta2, const half *epsilon,
half *variable, half *m, half *v, cudaStream_t cuda_stream);
template void AdamWeightDecayOp<float>(const size_t size, const float *gradient, const float *learning_rate,
const float *beta1, const float *beta2, const float *epsilon, const float *decay,
float *variable, float *m, float *v, cudaStream_t cuda_stream);
template void AdamWeightDecayOp<half>(const size_t size, const half *gradient, const float *learning_rate,
const float *beta1, const float *beta2, const float *epsilon, const float *decay,
half *variable, half *m, half *v, cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void ApplyAdam<float>(const size_t size, const float *gradient, const float *beta1_power,
const float *beta2_power, const float *learning_rate, const float *beta1,
const float *beta2, const float *epsilon, float *variable, float *m,
float *v, cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void ApplyAdam<half>(const size_t size, const half *gradient, const half *beta1_power,
const half *beta2_power, const half *learning_rate, const half *beta1,
const half *beta2, const half *epsilon, half *variable, half *m, half *v,
cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void AdamWeightDecayOp<float>(const size_t size, const float *gradient,
const float *learning_rate, const float *beta1,
const float *beta2, const float *epsilon, const float *decay,
float *variable, float *m, float *v, cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void AdamWeightDecayOp<half>(const size_t size, const half *gradient,
const float *learning_rate, const float *beta1,
const float *beta2, const float *epsilon, const float *decay,
half *variable, half *m, half *v, cudaStream_t cuda_stream);

29
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/adam_impl.cuh Normal file
View File

@ -0,0 +1,29 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADAM_IMPL_CUH_
#define MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADAM_IMPL_CUH_
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cuda_common.h"
template <typename T>
CUDA_LIB_EXPORT void ApplyAdam(const size_t size, const T *gradient, const T *beta1_power, const T *beta2_power,
const T *learning_rate, const T *beta1, const T *beta2, const T *epsilon, T *variable,
T *m, T *v, cudaStream_t cuda_stream);
template <typename T>
CUDA_LIB_EXPORT void AdamWeightDecayOp(const size_t size, const T *gradient, const float *learning_rate,
const float *beta1, const float *beta2, const float *epsilon, const float *decay,
T *variable, T *m, T *v, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADAM_IMPL_CUH_

10
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/adam_weight_decay_impl.cu → mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/adam_weight_decay_impl.cu
View File

@ -15,7 +15,6 @@
*/
#include "adam_weight_decay_impl.cuh"
#include "plugin/device/gpu/hal/device/cuda_common.h"
template <typename T>
__global__ void AdamWeightDecayKernel(const int element_num_, const bool need_decay, const float *beta1,
@ -44,7 +43,8 @@ void AdamWeightDecay(const int &element_num_, const bool &need_decay, const floa
gradient);
}
template void AdamWeightDecay(const int &element_num_, const bool &need_decay, const float *beta1,
const float *one_sub_beta1, const float *beta2, const float *one_sub_beta2,
const float *epsilon, const float *lr, const float *weight_decay, float *m, float *v,
float *param, float *gradient, cudaStream_t stream);
template CUDA_LIB_EXPORT void AdamWeightDecay(const int &element_num_, const bool &need_decay, const float *beta1,
const float *one_sub_beta1, const float *beta2,
const float *one_sub_beta2, const float *epsilon, const float *lr,
const float *weight_decay, float *m, float *v, float *param,
float *gradient, cudaStream_t stream);

26
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/adam_weight_decay_impl.cuh Normal file
View File

@ -0,0 +1,26 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADAM_WEIGHT_DECAY_IMPL_CUH_
#define MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADAM_WEIGHT_DECAY_IMPL_CUH_
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cuda_common.h"
template <typename T>
CUDA_LIB_EXPORT void AdamWeightDecay(const int &element_num_, const bool &need_decay, const float *beta1,
const float *one_sub_beta1, const float *beta2, const float *one_sub_beta2,
const float *epsilon, const float *lr, const float *weight_decay, T *m, T *v,
T *param, T *gradient, cudaStream_t stream);
#endif // MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADAM_WEIGHT_DECAY_IMPL_CUH_

24
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/adaptive_avg_pool2d_grad_impl.cu → mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/adaptive_avg_pool2d_grad_impl.cu
View File

@ -14,7 +14,8 @@
* limitations under the License.
*/
#include "plugin/device/gpu/kernel/cuda_impl/adaptive_avg_pool2d_grad_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/adaptive_avg_pool2d_grad_impl.cuh"
#include "include/cuda_fp16.h"
__device__ inline uint start_index(uint a, uint b, uint c) {
return floorf(__uint2float_rn(a * c) / __uint2float_rn(b));
@ -168,14 +169,17 @@ void ApplyAdaptiveAvgPool2DGrad(const uint size, const uint input_height, const
size, input_height, input_width, output_height, output_width, input_data, output_data);
}
template void ApplyAdaptiveAvgPool2DGrad<float>(const uint size, const uint input_height, const uint input_width,
const uint output_height, const uint output_width, float *input_data,
float *output_data, cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void ApplyAdaptiveAvgPool2DGrad<float>(const uint size, const uint input_height,
const uint input_width, const uint output_height,
const uint output_width, float *input_data,
float *output_data, cudaStream_t cuda_stream);
template void ApplyAdaptiveAvgPool2DGrad<half>(const uint size, const uint input_height, const uint input_width,
const uint output_height, const uint output_width, half *input_data,
half *output_data, cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void ApplyAdaptiveAvgPool2DGrad<half>(const uint size, const uint input_height,
const uint input_width, const uint output_height,
const uint output_width, half *input_data,
half *output_data, cudaStream_t cuda_stream);
template void ApplyAdaptiveAvgPool2DGrad<double>(const uint size, const uint input_height, const uint input_width,
const uint output_height, const uint output_width, double *input_data,
double *output_data, cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void ApplyAdaptiveAvgPool2DGrad<double>(const uint size, const uint input_height,
const uint input_width, const uint output_height,
const uint output_width, double *input_data,
double *output_data, cudaStream_t cuda_stream);

25
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/adaptive_avg_pool2d_grad_impl.cuh Normal file
View File

@ -0,0 +1,25 @@
/**
* Copyright 2021 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADAPTIVE_AVGPOOL2D_GRAD_IMPL_CUH_
#define MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADAPTIVE_AVGPOOL2D_GRAD_IMPL_CUH_
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cuda_common.h"
template <typename T>
CUDA_LIB_EXPORT void ApplyAdaptiveAvgPool2DGrad(const uint size, const uint input_height, const uint input_width,
const uint output_height, const uint output_width, T *input_data,
T *output_data, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADAPTIVE_AVGPOOL2D_GRAD_IMPL_CUH_

24
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/adaptive_avg_pool2d_impl.cu → mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/adaptive_avg_pool2d_impl.cu
View File

@ -14,7 +14,8 @@
* limitations under the License.
*/
#include "plugin/device/gpu/kernel/cuda_impl/adaptive_avg_pool2d_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/adaptive_avg_pool2d_impl.cuh"
#include "include/cuda_fp16.h"
__device__ inline uint start_index(uint a, uint b, uint c) {
return floorf(__uint2float_rn(a * c) / __uint2float_rn(b));
@ -155,14 +156,17 @@ void ApplyAdaptiveAvgPool2D(const uint size, const uint input_height, const uint
size, input_height, input_width, output_height, output_width, input_data, output_data);
}
template void ApplyAdaptiveAvgPool2D<float>(const uint size, const uint input_height, const uint input_width,
const uint output_height, const uint output_width, float *input_data,
float *output_data, cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void ApplyAdaptiveAvgPool2D<float>(const uint size, const uint input_height,
const uint input_width, const uint output_height,
const uint output_width, float *input_data,
float *output_data, cudaStream_t cuda_stream);
template void ApplyAdaptiveAvgPool2D<half>(const uint size, const uint input_height, const uint input_width,
const uint output_height, const uint output_width, half *input_data,
half *output_data, cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void ApplyAdaptiveAvgPool2D<half>(const uint size, const uint input_height,
const uint input_width, const uint output_height,
const uint output_width, half *input_data,
half *output_data, cudaStream_t cuda_stream);
template void ApplyAdaptiveAvgPool2D<double>(const uint size, const uint input_height, const uint input_width,
const uint output_height, const uint output_width, double *input_data,
double *output_data, cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void ApplyAdaptiveAvgPool2D<double>(const uint size, const uint input_height,
const uint input_width, const uint output_height,
const uint output_width, double *input_data,
double *output_data, cudaStream_t cuda_stream);

25
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/adaptive_avg_pool2d_impl.cuh Normal file
View File

@ -0,0 +1,25 @@
/**
* Copyright 2021 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADAPTIVE_AVGPOOL2D_IMPL_CUH_
#define MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADAPTIVE_AVGPOOL2D_IMPL_CUH_
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cuda_common.h"
template <typename T>
CUDA_LIB_EXPORT void ApplyAdaptiveAvgPool2D(const uint size, const uint input_height, const uint input_width,
const uint output_height, const uint output_width, T *input_data,
T *output_data, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADAPTIVE_AVGPOOL2D_IMPL_CUH_

15
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/sigmoid_cross_entropy_with_logits_impl.cuh → mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/add_relu_impl.cuh
View File

@ -14,12 +14,15 @@
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_SIGMOID_CROSS_ENTROPY_WITH_LOGITS_IMPL_H_
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_SIGMOID_CROSS_ENTROPY_WITH_LOGITS_IMPL_H_
#ifndef MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADD_RELU_IMPL_CUH_
#define MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADD_RELU_IMPL_CUH_
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cuda_common.h"
template <typename T>
CUDA_LIB_EXPORT void AddReluV2(const size_t num, const T *x1, const T *x2, T *y, uint32_t *mask,
cudaStream_t cuda_stream);
#include "plugin/device/gpu/hal/device/cuda_common.h"
template <typename T, typename S>
void SigmoidCrossEntropyWithLogits(const size_t size, const T *logits, const S *labels, T *outputs,
template <typename T>
CUDA_LIB_EXPORT void AddReluGradV2(const size_t size, const T *x1, const T *x2, const uint32_t *mask, T *dx,
cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMP_SIGMOID_CROSS_ENTROPY_WITH_LOGITS_IMPL_H_
#endif // MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADD_RELU_IMPL_CUH_

36
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/add_relu_v2_impl.cu → mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/add_relu_v2_impl.cu
View File

@ -14,8 +14,8 @@
* limitations under the License.
*/
#include "plugin/device/gpu/kernel/cuda_impl/add_relu_v2_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/util.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/add_relu_v2_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/util.cuh"
template <typename T>
__global__ void AddReluV2Kernel(const size_t num, const T *x1, const T *x2, T *y, uint32_t *mask) {
@ -49,20 +49,20 @@ void AddReluGradV2(const size_t num, const T *x1, const T *x2, const uint32_t *m
AddReluGradV2Kernel<<<kBlocksPerGrid(num), kThreadsPerBlock, 0, cuda_stream>>>(num, x1, x2, mask, dx);
}
template void AddReluV2(const size_t num, const float *x1, const float *x2, float *y, uint32_t *mask,
cudaStream_t cuda_stream);
template void AddReluV2(const size_t num, const half *x1, const half *x2, half *y, uint32_t *mask,
cudaStream_t cuda_stream);
template void AddReluV2(const size_t num, const int32_t *x1, const int32_t *x2, int32_t *y, uint32_t *mask,
cudaStream_t cuda_stream);
template void AddReluV2(const size_t num, const int64_t *x1, const int64_t *x2, int64_t *y, uint32_t *mask,
cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void AddReluV2(const size_t num, const float *x1, const float *x2, float *y, uint32_t *mask,
cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void AddReluV2(const size_t num, const half *x1, const half *x2, half *y, uint32_t *mask,
cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void AddReluV2(const size_t num, const int32_t *x1, const int32_t *x2, int32_t *y,
uint32_t *mask, cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void AddReluV2(const size_t num, const int64_t *x1, const int64_t *x2, int64_t *y,
uint32_t *mask, cudaStream_t cuda_stream);
template void AddReluGradV2(const size_t num, const float *x1, const float *x2, const uint32_t *mask, float *dx,
cudaStream_t cuda_stream);
template void AddReluGradV2(const size_t num, const half *x1, const half *x2, const uint32_t *mask, half *dx,
cudaStream_t cuda_stream);
template void AddReluGradV2(const size_t num, const int32_t *x1, const int32_t *x2, const uint32_t *mask, int32_t *dx,
cudaStream_t cuda_stream);
template void AddReluGradV2(const size_t num, const int64_t *x1, const int64_t *x2, const uint32_t *mask, int64_t *dx,
cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void AddReluGradV2(const size_t num, const float *x1, const float *x2, const uint32_t *mask,
float *dx, cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void AddReluGradV2(const size_t num, const half *x1, const half *x2, const uint32_t *mask,
half *dx, cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void AddReluGradV2(const size_t num, const int32_t *x1, const int32_t *x2,
const uint32_t *mask, int32_t *dx, cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void AddReluGradV2(const size_t num, const int64_t *x1, const int64_t *x2,
const uint32_t *mask, int64_t *dx, cudaStream_t cuda_stream);

28
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/add_relu_v2_impl.cuh Normal file
View File

@ -0,0 +1,28 @@
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADD_RELU_V2_IMPL_CUH_
#define MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADD_RELU_V2_IMPL_CUH_
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cuda_common.h"
template <typename T>
CUDA_LIB_EXPORT void AddReluV2(const size_t num, const T *x1, const T *x2, T *y, uint32_t *mask,
cudaStream_t cuda_stream);
template <typename T>
CUDA_LIB_EXPORT void AddReluGradV2(const size_t size, const T *x1, const T *x2, const uint32_t *mask, T *dx,
cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_ADD_RELU_V2_IMPL_CUH_

12
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/apply_gradient_descent_impl.cu → mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/apply_gradient_descent_impl.cu
View File

@ -14,7 +14,8 @@
* limitations under the License.
*/
#include "plugin/device/gpu/kernel/cuda_impl/apply_gradient_descent_impl.cuh"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/apply_gradient_descent_impl.cuh"
#include "include/cuda_fp16.h"
template <typename T>
__global__ void ApplyGradientDescent(const size_t size, T *var, const T *alpha, const T *delta, T *output) {
@ -31,7 +32,8 @@ void CalApplyGradientDescent(const size_t &size, T *var, const T *alpha, const T
ApplyGradientDescent<<<GET_BLOCKS(size), GET_THREADS, 0, cuda_stream>>>(size, var, alpha, delta, output);
}
template void CalApplyGradientDescent<float>(const size_t &size, float *var, const float *alpha, const float *delta,
float *output, cudaStream_t cuda_stream);
template void CalApplyGradientDescent<half>(const size_t &size, half *var, const half *alpha, const half *delta,
half *output, cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void CalApplyGradientDescent<float>(const size_t &size, float *var, const float *alpha,
const float *delta, float *output,
cudaStream_t cuda_stream);
template CUDA_LIB_EXPORT void CalApplyGradientDescent<half>(const size_t &size, half *var, const half *alpha,
const half *delta, half *output, cudaStream_t cuda_stream);

15
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/hsigmoid_impl.cuh → mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/apply_gradient_descent_impl.cuh
View File

@ -14,16 +14,13 @@
* limitations under the License.
*/
#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_CUDA_IMPL_HSIGMOID_IMPL_CUH_
#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_CUDA_IMPL_HSIGMOID_IMPL_CUH_
#ifndef MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_APPLY_GRADIENT_DESCENT_IMPL_CUH_
#define MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_APPLY_GRADIENT_DESCENT_IMPL_CUH_
#include <cuda_runtime.h>
#include "plugin/device/gpu/hal/device/cuda_common.h"
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/cuda_common.h"
template <typename T>
void CalHSigmoid(const size_t &size, const T *input, T *output, cudaStream_t cuda_stream);
CUDA_LIB_EXPORT void CalApplyGradientDescent(const size_t &size, T *var, const T *alpha, const T *delta, T *output,
cudaStream_t cuda_stream);
template <typename T>
void CalHSigmoidGrad(const size_t &size, const T *dout, const T *x, T *output, cudaStream_t cuda_stream);
#endif // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_CUDA_IMPL_HSIGMOID_IMPL_CUH_
#endif // MINDSPORE_CCSRC_PLUGIN_DEVICE_GPU_KERNEL_CUDA_IMPL_CUDA_OPS_APPLY_GRADIENT_DESCENT_IMPL_CUH_

Some files were not shown because too many files have changed in this diff Show More