Skip to content

P
Paddle

PaddlePaddle / Paddle
1 年多 前同步成功

通知 2305
Star 20932
Fork 5423
P
Paddle
提交 4bc1dc26 编写于 作者: T tensor-tang
浏览文件
操作

Merge remote-tracking branch 'upstream/develop' into merge-doc

上级 6405438b 91b674fc
展开全部 显示空白变更内容
内联 并排
Showing with 5866 addition and 4152 deletion
.pre-commit-config.yaml
浏览文件 @ 4bc1dc26
...@@ -24,7 +24,7 @@ ...@@ -24,7 +24,7 @@
description: Format files with ClangFormat. description: Format files with ClangFormat.
entry: clang-format -i entry: clang-format -i
language: system language: system
files: \.(c|cc|cxx|cpp|h|hpp|hxx)$ files: \.(c|cc|cxx|cpp|cu|h|hpp|hxx|proto)$
- repo: https://github.com/PaddlePaddle/pre-commit-golang - repo: https://github.com/PaddlePaddle/pre-commit-golang
sha: 8337620115c25ff8333f1b1a493bd031049bd7c0 sha: 8337620115c25ff8333f1b1a493bd031049bd7c0
hooks: hooks:
......
CMakeLists.txt
浏览文件 @ 4bc1dc26
...@@ -36,8 +36,8 @@ include(simd) ...@@ -36,8 +36,8 @@ include(simd)
################################ Configurations ####################################### ################################ Configurations #######################################
option(WITH_GPU "Compile PaddlePaddle with NVIDIA GPU" ${CUDA_FOUND}) option(WITH_GPU "Compile PaddlePaddle with NVIDIA GPU" ${CUDA_FOUND})
option(WITH_AVX "Compile PaddlePaddle with AVX intrinsics" ${AVX_FOUND}) option(WITH_AVX "Compile PaddlePaddle with AVX intrinsics" ${AVX_FOUND})
option(WITH_MKLDNN "Compile PaddlePaddle with mkl-dnn support." ${AVX_FOUND}) option(WITH_MKLDNN "Compile PaddlePaddle with mkl-dnn support." OFF)
option(WITH_MKLML "Compile PaddlePaddle with mklml package." ${AVX_FOUND}) option(WITH_MKLML "Compile PaddlePaddle with mklml package." OFF)
option(WITH_DSO "Compile PaddlePaddle with dynamic linked CUDA" ON) option(WITH_DSO "Compile PaddlePaddle with dynamic linked CUDA" ON)
option(WITH_TESTING "Compile PaddlePaddle with unit testing" ON) option(WITH_TESTING "Compile PaddlePaddle with unit testing" ON)
option(WITH_SWIG_PY "Compile PaddlePaddle with inference api" ON) option(WITH_SWIG_PY "Compile PaddlePaddle with inference api" ON)
......
Dockerfile
浏览文件 @ 4bc1dc26
...@@ -27,13 +27,16 @@ RUN apt-get update && \ ...@@ -27,13 +27,16 @@ RUN apt-get update && \
git python-pip python-dev openssh-server bison \ git python-pip python-dev openssh-server bison \
wget unzip unrar tar xz-utils bzip2 gzip coreutils ntp \ wget unzip unrar tar xz-utils bzip2 gzip coreutils ntp \
curl sed grep graphviz libjpeg-dev zlib1g-dev \ curl sed grep graphviz libjpeg-dev zlib1g-dev \
python-numpy python-matplotlib gcc-4.8 g++-4.8 \ python-matplotlib gcc-4.8 g++-4.8 \
automake locales clang-format-3.8 swig doxygen cmake \ automake locales clang-format-3.8 swig doxygen cmake \
liblapack-dev liblapacke-dev libboost-dev \ liblapack-dev liblapacke-dev libboost-dev \
clang-3.8 llvm-3.8 libclang-3.8-dev \ clang-3.8 llvm-3.8 libclang-3.8-dev \
net-tools && \ net-tools && \
apt-get clean -y apt-get clean -y
# paddle is using numpy.flip, which is introduced since 1.12.0
RUN pip --no-cache-dir install 'numpy>=1.12.0'
# Install Go and glide # Install Go and glide
RUN wget -O go.tgz https://storage.googleapis.com/golang/go1.8.1.linux-amd64.tar.gz && \ RUN wget -O go.tgz https://storage.googleapis.com/golang/go1.8.1.linux-amd64.tar.gz && \
tar -C /usr/local -xzf go.tgz && \ tar -C /usr/local -xzf go.tgz && \
......
cmake/cpplint.cmake
浏览文件 @ 4bc1dc26
...@@ -56,11 +56,14 @@ macro(add_style_check_target TARGET_NAME) ...@@ -56,11 +56,14 @@ macro(add_style_check_target TARGET_NAME)
# cpplint code style # cpplint code style
get_filename_component(base_filename ${filename} NAME) get_filename_component(base_filename ${filename} NAME)
set(CUR_GEN ${CMAKE_CURRENT_BINARY_DIR}/${base_filename}.cpplint) set(CUR_GEN ${CMAKE_CURRENT_BINARY_DIR}/${base_filename}.cpplint)
add_custom_command(TARGET ${TARGET_NAME} PRE_BUILD add_custom_command(OUTPUT ${CUR_GEN} PRE_BUILD
COMMAND "${PYTHON_EXECUTABLE}" "${PROJ_ROOT}/paddle/scripts/cpplint.py" COMMAND "${PYTHON_EXECUTABLE}" "${PROJ_ROOT}/paddle/scripts/cpplint.py"
"--filter=${STYLE_FILTER}" "--filter=${STYLE_FILTER}"
"--write-success=${CUR_GEN}" ${filename} "--write-success=${CUR_GEN}" ${filename}
DEPENDS ${filename} ${PROJ_ROOT}/paddle/scripts/cpplint.py
WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}) WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR})
add_custom_target(${base_filename}.cpplint DEPENDS ${CUR_GEN})
add_dependencies(${TARGET_NAME} ${base_filename}.cpplint)
endif() endif()
endforeach() endforeach()
endif() endif()
......
cmake/external/gflags.cmake
浏览文件 @ 4bc1dc26
...@@ -28,7 +28,14 @@ INCLUDE_DIRECTORIES(${GFLAGS_INCLUDE_DIR}) ...@@ -28,7 +28,14 @@ INCLUDE_DIRECTORIES(${GFLAGS_INCLUDE_DIR})
ExternalProject_Add( ExternalProject_Add(
extern_gflags extern_gflags
${EXTERNAL_PROJECT_LOG_ARGS} ${EXTERNAL_PROJECT_LOG_ARGS}
GIT_REPOSITORY "https://github.com/gflags/gflags.git" # TODO(yiwang): The annoying warnings mentioned in
# https://github.com/PaddlePaddle/Paddle/issues/3277 are caused by
# gflags. I fired a PR https://github.com/gflags/gflags/pull/230
# to fix it. Before it gets accepted by the gflags team, we use
# my personal fork, which contains above fix, temporarily. Let's
# change this back to the official Github repo once my PR is
# merged.
GIT_REPOSITORY "https://github.com/wangkuiyi/gflags.git"
PREFIX ${GFLAGS_SOURCES_DIR} PREFIX ${GFLAGS_SOURCES_DIR}
UPDATE_COMMAND "" UPDATE_COMMAND ""
CMAKE_ARGS -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER} CMAKE_ARGS -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER}
......
cmake/external/python.cmake
浏览文件 @ 4bc1dc26
...@@ -24,7 +24,6 @@ IF(WITH_PYTHON) ...@@ -24,7 +24,6 @@ IF(WITH_PYTHON)
ENDIF(WITH_PYTHON) ENDIF(WITH_PYTHON)
SET(py_env "") SET(py_env "")
SET(USE_VIRTUALENV_FOR_TEST 1)
IF(PYTHONINTERP_FOUND) IF(PYTHONINTERP_FOUND)
find_python_module(pip REQUIRED) find_python_module(pip REQUIRED)
find_python_module(numpy REQUIRED) find_python_module(numpy REQUIRED)
......
cmake/generic.cmake
浏览文件 @ 4bc1dc26
...@@ -187,7 +187,13 @@ function(cc_library TARGET_NAME) ...@@ -187,7 +187,13 @@ function(cc_library TARGET_NAME)
endif() endif()
# cpplint code style # cpplint code style
add_style_check_target(${TARGET_NAME} ${cc_library_SRCS}) foreach(source_file ${cc_library_SRCS})
string(REGEX REPLACE "\\.[^.]*$" "" source ${source_file})
if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/${source}.h)
list(APPEND cc_library_HEADERS ${CMAKE_CURRENT_SOURCE_DIR}/${source}.h)
endif()
endforeach()
add_style_check_target(${TARGET_NAME} ${cc_library_SRCS} ${cc_library_HEADERS})
else(cc_library_SRCS) else(cc_library_SRCS)
if (cc_library_DEPS) if (cc_library_DEPS)
...@@ -239,6 +245,14 @@ function(nv_library TARGET_NAME) ...@@ -239,6 +245,14 @@ function(nv_library TARGET_NAME)
add_dependencies(${TARGET_NAME} ${nv_library_DEPS}) add_dependencies(${TARGET_NAME} ${nv_library_DEPS})
target_link_libraries(${TARGET_NAME} ${nv_library_DEPS}) target_link_libraries(${TARGET_NAME} ${nv_library_DEPS})
endif() endif()
# cpplint code style
foreach(source_file ${nv_library_SRCS})
string(REGEX REPLACE "\\.[^.]*$" "" source ${source_file})
if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/${source}.h)
list(APPEND cc_library_HEADERS ${CMAKE_CURRENT_SOURCE_DIR}/${source}.h)
endif()
endforeach()
add_style_check_target(${TARGET_NAME} ${nv_library_SRCS} ${nv_library_HEADERS})
else(nv_library_SRCS) else(nv_library_SRCS)
if (nv_library_DEPS) if (nv_library_DEPS)
merge_static_libs(${TARGET_NAME} ${nv_library_DEPS}) merge_static_libs(${TARGET_NAME} ${nv_library_DEPS})
......
cmake/util.cmake
浏览文件 @ 4bc1dc26
...@@ -118,7 +118,6 @@ endfunction() ...@@ -118,7 +118,6 @@ endfunction()
macro(add_unittest_without_exec TARGET_NAME) macro(add_unittest_without_exec TARGET_NAME)
add_executable(${TARGET_NAME} ${ARGN}) add_executable(${TARGET_NAME} ${ARGN})
link_paddle_test(${TARGET_NAME}) link_paddle_test(${TARGET_NAME})
add_style_check_target(${TARGET_NAME} ${ARGN})
endmacro() endmacro()
# add_unittest # add_unittest
...@@ -150,9 +149,12 @@ endfunction() ...@@ -150,9 +149,12 @@ endfunction()
# Create a python unittest using run_python_tests.sh, # Create a python unittest using run_python_tests.sh,
# which takes care of making correct running environment # which takes care of making correct running environment
function(add_python_test TEST_NAME) function(add_python_test TEST_NAME)
add_test(NAME ${TEST_NAME} foreach(arg ${ARGN})
COMMAND env PADDLE_PACKAGE_DIR=${PADDLE_PYTHON_PACKAGE_DIR} get_filename_component(py_fn ${arg} NAME_WE)
bash ${PROJ_ROOT}/paddle/scripts/run_python_tests.sh set(TRG_NAME ${TEST_NAME}_${py_fn})
${USE_VIRTUALENV_FOR_TEST} ${PYTHON_EXECUTABLE} ${ARGN} add_test(NAME ${TRG_NAME}
COMMAND env PYTHONPATH=${PADDLE_PYTHON_PACKAGE_DIR}
python2 ${arg}
WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}) WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR})
endforeach()
endfunction() endfunction()
paddle/.set_python_path.sh
浏览文件 @ 4bc1dc26
...@@ -21,22 +21,15 @@ ...@@ -21,22 +21,15 @@
# #
# It same as PYTHONPATH=${YOUR_PYTHON_PATH}:$PYTHONPATH {exec...} # It same as PYTHONPATH=${YOUR_PYTHON_PATH}:$PYTHONPATH {exec...}
# #
PYPATH=""
if ! python -c "import paddle" >/dev/null 2>/dev/null; then set -x
PYPATH="" while getopts "d:" opt; do
set -x
while getopts "d:" opt; do
case $opt in case $opt in
d) d)
PYPATH=$OPTARG PYPATH=$OPTARG
;; ;;
esac esac
done done
shift $(($OPTIND - 1)) shift $(($OPTIND - 1))
export PYTHONPATH=$PYPATH:$PYTHONPATH export PYTHONPATH=$PYPATH:$PYTHONPATH
$@ $@
else
echo "paddle package is already in your PYTHONPATH. But unittest need a clean environment."
echo "Please uninstall paddle package before start unittest. Try to 'pip uninstall paddle'"
exit 1
fi
paddle/CMakeLists.txt
浏览文件 @ 4bc1dc26
...@@ -15,7 +15,6 @@ if(Boost_FOUND) ...@@ -15,7 +15,6 @@ if(Boost_FOUND)
add_subdirectory(platform) add_subdirectory(platform)
add_subdirectory(framework) add_subdirectory(framework)
add_subdirectory(operators) add_subdirectory(operators)
add_subdirectory(pybind)
endif() endif()
if(WITH_C_API) if(WITH_C_API)
......
paddle/cuda/src/hl_batch_transpose.cu
浏览文件 @ 4bc1dc26
...@@ -12,17 +12,15 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,17 +12,15 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "hl_batch_transpose.h"
#include "hl_base.h" #include "hl_base.h"
#include "hl_batch_transpose.h"
const int TILE_DIM = 64; const int TILE_DIM = 64;
const int BLOCK_ROWS = 16; const int BLOCK_ROWS = 16;
// No bank-conflict transpose for a batch of data. // No bank-conflict transpose for a batch of data.
__global__ void batchTransposeNoBankConflicts(real* odata, __global__ void batchTransposeNoBankConflicts(
const real* idata, real* odata, const real* idata, int numSamples, int width, int height) {
int numSamples, int width,
int height) {
__shared__ float tile[TILE_DIM][TILE_DIM + 1]; __shared__ float tile[TILE_DIM][TILE_DIM + 1];
const int x = blockIdx.x * TILE_DIM + threadIdx.x; const int x = blockIdx.x * TILE_DIM + threadIdx.x;
...@@ -50,12 +48,12 @@ __global__ void batchTransposeNoBankConflicts(real* odata, ...@@ -50,12 +48,12 @@ __global__ void batchTransposeNoBankConflicts(real* odata,
newX] = tile[threadIdx.x][j]; newX] = tile[threadIdx.x][j];
} }
void batchTranspose(const real* input, real* output, int width, int height, void batchTranspose(
int batchSize) { const real* input, real* output, int width, int height, int batchSize) {
dim3 dimBlock(TILE_DIM, BLOCK_ROWS, 1); dim3 dimBlock(TILE_DIM, BLOCK_ROWS, 1);
dim3 dimGrid(DIVUP(width, TILE_DIM), DIVUP(height, TILE_DIM), batchSize); dim3 dimGrid(DIVUP(width, TILE_DIM), DIVUP(height, TILE_DIM), batchSize);
batchTransposeNoBankConflicts<<<dimGrid, dimBlock, 0, STREAM_DEFAULT>>> batchTransposeNoBankConflicts<<<dimGrid, dimBlock, 0, STREAM_DEFAULT>>>(
(output, input, batchSize, width, height); output, input, batchSize, width, height);
CHECK_SYNC("batchTranspose failed!"); CHECK_SYNC("batchTranspose failed!");
} }
paddle/cuda/src/hl_cuda_aggregate.cu
浏览文件 @ 4bc1dc26
...@@ -12,27 +12,23 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,27 +12,23 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "hl_aggregate.h"
#include "hl_base.h" #include "hl_base.h"
#include "hl_cuda.h" #include "hl_cuda.h"
#include "hl_cuda.ph" #include "hl_cuda.ph"
#include "hl_aggregate.h"
#include "hl_thread.ph"
#include "hl_matrix_base.cuh" #include "hl_matrix_base.cuh"
#include "hl_thread.ph"
#include "paddle/utils/Logging.h" #include "paddle/utils/Logging.h"
/** /**
* @brief matrix row operator. * @brief matrix row operator.
*/ */
template<class Agg, int blockSize> template <class Agg, int blockSize>
__global__ void KeMatrixRowOp(Agg agg, __global__ void KeMatrixRowOp(Agg agg, real *E, real *Sum, int dimN) {
real *E,
real *Sum,
int dimN) {
__shared__ real sum_s[blockSize]; __shared__ real sum_s[blockSize];
int cnt = (dimN + blockSize -1) / blockSize; int cnt = (dimN + blockSize - 1) / blockSize;
int rowId = blockIdx.x + blockIdx.y*gridDim.x; int rowId = blockIdx.x + blockIdx.y * gridDim.x;
int index = rowId*dimN; int index = rowId * dimN;
int tid = threadIdx.x; int tid = threadIdx.x;
int lmt = tid; int lmt = tid;
...@@ -44,7 +40,7 @@ __global__ void KeMatrixRowOp(Agg agg, ...@@ -44,7 +40,7 @@ __global__ void KeMatrixRowOp(Agg agg,
sum_s[tid] = tmp; sum_s[tid] = tmp;
__syncthreads(); __syncthreads();
for (int stride = blockSize/2; stride > 0; stride = stride/2) { for (int stride = blockSize / 2; stride > 0; stride = stride / 2) {
if (tid < stride) { if (tid < stride) {
sum_s[tid] = agg(sum_s[tid], sum_s[tid + stride]); sum_s[tid] = agg(sum_s[tid], sum_s[tid + stride]);
} }
...@@ -58,29 +54,21 @@ __global__ void KeMatrixRowOp(Agg agg, ...@@ -58,29 +54,21 @@ __global__ void KeMatrixRowOp(Agg agg,
} }
template <class Agg> template <class Agg>
void hl_matrix_row_op(Agg agg, void hl_matrix_row_op(Agg agg, real *A_d, real *C_d, int dimM, int dimN) {
real *A_d,
real *C_d,
int dimM,
int dimN) {
int blocksX = dimM; int blocksX = dimM;
int blocksY = 1; int blocksY = 1;
dim3 threads(128, 1); dim3 threads(128, 1);
dim3 grid(blocksX, blocksY); dim3 grid(blocksX, blocksY);
KeMatrixRowOp<Agg, 128><<< grid, threads, 0, STREAM_DEFAULT >>> KeMatrixRowOp<Agg, 128><<<grid, threads, 0, STREAM_DEFAULT>>>(
(agg, A_d, C_d, dimN); agg, A_d, C_d, dimN);
} }
void hl_matrix_row_sum(real *A_d, real *C_d, int dimM, int dimN) { void hl_matrix_row_sum(real *A_d, real *C_d, int dimM, int dimN) {
CHECK_NOTNULL(A_d); CHECK_NOTNULL(A_d);
CHECK_NOTNULL(C_d); CHECK_NOTNULL(C_d);
hl_matrix_row_op(aggregate::sum(), hl_matrix_row_op(aggregate::sum(), A_d, C_d, dimM, dimN);
A_d,
C_d,
dimM,
dimN);
CHECK_SYNC("hl_matrix_row_sum failed"); CHECK_SYNC("hl_matrix_row_sum failed");
} }
...@@ -88,11 +76,7 @@ void hl_matrix_row_max(real *A_d, real *C_d, int dimM, int dimN) { ...@@ -88,11 +76,7 @@ void hl_matrix_row_max(real *A_d, real *C_d, int dimM, int dimN) {
CHECK_NOTNULL(A_d); CHECK_NOTNULL(A_d);
CHECK_NOTNULL(C_d); CHECK_NOTNULL(C_d);
hl_matrix_row_op(aggregate::max(), hl_matrix_row_op(aggregate::max(), A_d, C_d, dimM, dimN);
A_d,
C_d,
dimM,
dimN);
CHECK_SYNC("hl_matrix_row_max failed"); CHECK_SYNC("hl_matrix_row_max failed");
} }
...@@ -100,23 +84,16 @@ void hl_matrix_row_min(real *A_d, real *C_d, int dimM, int dimN) { ...@@ -100,23 +84,16 @@ void hl_matrix_row_min(real *A_d, real *C_d, int dimM, int dimN) {
CHECK_NOTNULL(A_d); CHECK_NOTNULL(A_d);
CHECK_NOTNULL(C_d); CHECK_NOTNULL(C_d);
hl_matrix_row_op(aggregate::min(), hl_matrix_row_op(aggregate::min(), A_d, C_d, dimM, dimN);
A_d,
C_d,
dimM,
dimN);
CHECK_SYNC("hl_matrix_row_min failed"); CHECK_SYNC("hl_matrix_row_min failed");
} }
/** /**
* @brief matrix column operator. * @brief matrix column operator.
*/ */
template<class Agg> template <class Agg>
__global__ void KeMatrixColumnOp(Agg agg, __global__ void KeMatrixColumnOp(
real *E, Agg agg, real *E, real *Sum, int dimM, int dimN) {
real *Sum,
int dimM,
int dimN) {
int rowIdx = blockIdx.x * blockDim.x + threadIdx.x; int rowIdx = blockIdx.x * blockDim.x + threadIdx.x;
real tmp = agg.init(); real tmp = agg.init();
if (rowIdx < dimN) { if (rowIdx < dimN) {
...@@ -127,13 +104,10 @@ __global__ void KeMatrixColumnOp(Agg agg, ...@@ -127,13 +104,10 @@ __global__ void KeMatrixColumnOp(Agg agg,
} }
} }
template<class Agg, int blockDimX, int blockDimY> template <class Agg, int blockDimX, int blockDimY>
__global__ void KeMatrixColumnOp_S(Agg agg, __global__ void KeMatrixColumnOp_S(
real *E, Agg agg, real *E, real *Sum, int dimM, int dimN) {
real *Sum, __shared__ real _sum[blockDimX * blockDimY];
int dimM,
int dimN) {
__shared__ real _sum[blockDimX*blockDimY];
int rowIdx = blockIdx.x * blockDim.x + threadIdx.x; int rowIdx = blockIdx.x * blockDim.x + threadIdx.x;
int index = threadIdx.y; int index = threadIdx.y;
...@@ -144,14 +118,14 @@ __global__ void KeMatrixColumnOp_S(Agg agg, ...@@ -144,14 +118,14 @@ __global__ void KeMatrixColumnOp_S(Agg agg,
index += blockDimY; index += blockDimY;
} }
} }
_sum[threadIdx.x + threadIdx.y*blockDimX] = tmp; _sum[threadIdx.x + threadIdx.y * blockDimX] = tmp;
__syncthreads(); __syncthreads();
if (rowIdx < dimN) { if (rowIdx < dimN) {
if (threadIdx.y ==0) { if (threadIdx.y == 0) {
real tmp = agg.init(); real tmp = agg.init();
for (int i=0; i < blockDimY; i++) { for (int i = 0; i < blockDimY; i++) {
tmp = agg(tmp, _sum[threadIdx.x + i*blockDimX]); tmp = agg(tmp, _sum[threadIdx.x + i * blockDimX]);
} }
Sum[rowIdx] = tmp; Sum[rowIdx] = tmp;
} }
...@@ -159,25 +133,21 @@ __global__ void KeMatrixColumnOp_S(Agg agg, ...@@ -159,25 +133,21 @@ __global__ void KeMatrixColumnOp_S(Agg agg,
} }
template <class Agg> template <class Agg>
void hl_matrix_column_op(Agg agg, void hl_matrix_column_op(Agg agg, real *A_d, real *C_d, int dimM, int dimN) {
real *A_d,
real *C_d,
int dimM,
int dimN) {
if (dimN >= 8192) { if (dimN >= 8192) {
int blocksX = (dimN + 128 -1) / 128; int blocksX = (dimN + 128 - 1) / 128;
int blocksY = 1; int blocksY = 1;
dim3 threads(128, 1); dim3 threads(128, 1);
dim3 grid(blocksX, blocksY); dim3 grid(blocksX, blocksY);
KeMatrixColumnOp<Agg><<< grid, threads, 0, STREAM_DEFAULT >>> KeMatrixColumnOp<Agg><<<grid, threads, 0, STREAM_DEFAULT>>>(
(agg, A_d, C_d, dimM, dimN); agg, A_d, C_d, dimM, dimN);
} else { } else {
int blocksX = (dimN + 32 -1) / 32; int blocksX = (dimN + 32 - 1) / 32;
int blocksY = 1; int blocksY = 1;
dim3 threads(32, 32); dim3 threads(32, 32);
dim3 grid(blocksX, blocksY); dim3 grid(blocksX, blocksY);
KeMatrixColumnOp_S<Agg, 32, 32><<< grid, threads, 0, STREAM_DEFAULT>>> KeMatrixColumnOp_S<Agg, 32, 32><<<grid, threads, 0, STREAM_DEFAULT>>>(
(agg, A_d, C_d, dimM, dimN); agg, A_d, C_d, dimM, dimN);
} }
return; return;
...@@ -187,11 +157,7 @@ void hl_matrix_column_sum(real *A_d, real *C_d, int dimM, int dimN) { ...@@ -187,11 +157,7 @@ void hl_matrix_column_sum(real *A_d, real *C_d, int dimM, int dimN) {
CHECK_NOTNULL(A_d); CHECK_NOTNULL(A_d);
CHECK_NOTNULL(C_d); CHECK_NOTNULL(C_d);
hl_matrix_column_op(aggregate::sum(), hl_matrix_column_op(aggregate::sum(), A_d, C_d, dimM, dimN);
A_d,
C_d,
dimM,
dimN);
CHECK_SYNC("hl_matrix_column_sum failed"); CHECK_SYNC("hl_matrix_column_sum failed");
} }
...@@ -200,11 +166,7 @@ void hl_matrix_column_max(real *A_d, real *C_d, int dimM, int dimN) { ...@@ -200,11 +166,7 @@ void hl_matrix_column_max(real *A_d, real *C_d, int dimM, int dimN) {
CHECK_NOTNULL(A_d); CHECK_NOTNULL(A_d);
CHECK_NOTNULL(C_d); CHECK_NOTNULL(C_d);
hl_matrix_column_op(aggregate::max(), hl_matrix_column_op(aggregate::max(), A_d, C_d, dimM, dimN);
A_d,
C_d,
dimM,
dimN);
CHECK_SYNC("hl_matrix_column_max failed"); CHECK_SYNC("hl_matrix_column_max failed");
} }
...@@ -213,11 +175,7 @@ void hl_matrix_column_min(real *A_d, real *C_d, int dimM, int dimN) { ...@@ -213,11 +175,7 @@ void hl_matrix_column_min(real *A_d, real *C_d, int dimM, int dimN) {
CHECK_NOTNULL(A_d); CHECK_NOTNULL(A_d);
CHECK_NOTNULL(C_d); CHECK_NOTNULL(C_d);
hl_matrix_column_op(aggregate::min(), hl_matrix_column_op(aggregate::min(), A_d, C_d, dimM, dimN);
A_d,
C_d,
dimM,
dimN);
CHECK_SYNC("hl_matrix_column_min failed"); CHECK_SYNC("hl_matrix_column_min failed");
} }
...@@ -226,16 +184,16 @@ template <int blockSize> ...@@ -226,16 +184,16 @@ template <int blockSize>
__global__ void KeVectorSum(real *E, real *Sum, int dimM) { __global__ void KeVectorSum(real *E, real *Sum, int dimM) {
__shared__ double sum_s[blockSize]; __shared__ double sum_s[blockSize];
int tid = threadIdx.x; int tid = threadIdx.x;
int index = blockIdx.y*blockDim.x+threadIdx.x; int index = blockIdx.y * blockDim.x + threadIdx.x;
sum_s[tid] = 0.0f; sum_s[tid] = 0.0f;
while (index < dimM) { while (index < dimM) {
sum_s[tid] += E[index]; sum_s[tid] += E[index];
index += blockDim.x*gridDim.y; index += blockDim.x * gridDim.y;
} }
__syncthreads(); __syncthreads();
for (int stride = blockSize/2; stride > 0; stride = stride/2) { for (int stride = blockSize / 2; stride > 0; stride = stride / 2) {
if (tid < stride) { if (tid < stride) {
sum_s[tid] += sum_s[tid + stride]; sum_s[tid] += sum_s[tid + stride];
} }
...@@ -261,36 +219,37 @@ void hl_vector_sum(real *A_d, real *C_h, int dimM) { ...@@ -261,36 +219,37 @@ void hl_vector_sum(real *A_d, real *C_h, int dimM) {
struct _hl_event_st hl_event_st = {.cu_event = t_resource.event}; struct _hl_event_st hl_event_st = {.cu_event = t_resource.event};
hl_event_t hl_event = &hl_event_st; hl_event_t hl_event = &hl_event_st;
while (!hl_cuda_event_is_ready(hl_event)) {} while (!hl_cuda_event_is_ready(hl_event)) {
}
KeVectorSum<128><<< grid, threads, 0, STREAM_DEFAULT >>> KeVectorSum<128><<<grid, threads, 0, STREAM_DEFAULT>>>(
(A_d, t_resource.gpu_mem, dimM); A_d, t_resource.gpu_mem, dimM);
KeVectorSum<128><<< 1, threads, 0, STREAM_DEFAULT >>> KeVectorSum<128><<<1, threads, 0, STREAM_DEFAULT>>>(
(t_resource.gpu_mem, t_resource.cpu_mem, 128); t_resource.gpu_mem, t_resource.cpu_mem, 128);
hl_memcpy_async(C_h, t_resource.cpu_mem, sizeof(real), HPPL_STREAM_DEFAULT); hl_memcpy_async(C_h, t_resource.cpu_mem, sizeof(real), HPPL_STREAM_DEFAULT);
hl_stream_record_event(HPPL_STREAM_DEFAULT, hl_event); hl_stream_record_event(HPPL_STREAM_DEFAULT, hl_event);
hl_stream_synchronize(HPPL_STREAM_DEFAULT); hl_stream_synchronize(HPPL_STREAM_DEFAULT);
cudaError_t err = (cudaError_t)hl_get_device_last_error(); cudaError_t err = (cudaError_t)hl_get_device_last_error();
CHECK_EQ(cudaSuccess, err) CHECK_EQ(cudaSuccess, err) << "CUDA error: "
<< "CUDA error: " << hl_get_device_error_string((size_t)err); << hl_get_device_error_string((size_t)err);
} }
template <int blockSize> template <int blockSize>
__global__ void KeVectorAbsSum(real *E, real *Sum, int dimM) { __global__ void KeVectorAbsSum(real *E, real *Sum, int dimM) {
__shared__ double sum_s[blockSize]; __shared__ double sum_s[blockSize];
int tid = threadIdx.x; int tid = threadIdx.x;
int index = blockIdx.y*blockDim.x+threadIdx.x; int index = blockIdx.y * blockDim.x + threadIdx.x;
sum_s[tid] = 0.0f; sum_s[tid] = 0.0f;
while (index < dimM) { while (index < dimM) {
sum_s[tid] += abs(E[index]); sum_s[tid] += abs(E[index]);
index += blockDim.x*gridDim.y; index += blockDim.x * gridDim.y;
} }
__syncthreads(); __syncthreads();
for (int stride = blockSize/2; stride > 0; stride = stride/2) { for (int stride = blockSize / 2; stride > 0; stride = stride / 2) {
if (tid < stride) { if (tid < stride) {
sum_s[tid] += sum_s[tid + stride]; sum_s[tid] += sum_s[tid + stride];
} }
...@@ -316,18 +275,19 @@ void hl_vector_abs_sum(real *A_d, real *C_h, int dimM) { ...@@ -316,18 +275,19 @@ void hl_vector_abs_sum(real *A_d, real *C_h, int dimM) {
struct _hl_event_st hl_event_st = {.cu_event = t_resource.event}; struct _hl_event_st hl_event_st = {.cu_event = t_resource.event};
hl_event_t hl_event = &hl_event_st; hl_event_t hl_event = &hl_event_st;
while (!hl_cuda_event_is_ready(hl_event)) {} while (!hl_cuda_event_is_ready(hl_event)) {
}
KeVectorAbsSum<128><<< grid, threads, 0, STREAM_DEFAULT >>> KeVectorAbsSum<128><<<grid, threads, 0, STREAM_DEFAULT>>>(
(A_d, t_resource.gpu_mem, dimM); A_d, t_resource.gpu_mem, dimM);
KeVectorAbsSum<128><<< 1, threads, 0, STREAM_DEFAULT >>> KeVectorAbsSum<128><<<1, threads, 0, STREAM_DEFAULT>>>(
(t_resource.gpu_mem, t_resource.cpu_mem, 128); t_resource.gpu_mem, t_resource.cpu_mem, 128);
hl_memcpy_async(C_h, t_resource.cpu_mem, sizeof(real), HPPL_STREAM_DEFAULT); hl_memcpy_async(C_h, t_resource.cpu_mem, sizeof(real), HPPL_STREAM_DEFAULT);
hl_stream_record_event(HPPL_STREAM_DEFAULT, hl_event); hl_stream_record_event(HPPL_STREAM_DEFAULT, hl_event);
hl_stream_synchronize(HPPL_STREAM_DEFAULT); hl_stream_synchronize(HPPL_STREAM_DEFAULT);
cudaError_t err = (cudaError_t)hl_get_device_last_error(); cudaError_t err = (cudaError_t)hl_get_device_last_error();
CHECK_EQ(cudaSuccess, err) CHECK_EQ(cudaSuccess, err) << "CUDA error: "
<< "CUDA error: " << hl_get_device_error_string((size_t)err); << hl_get_device_error_string((size_t)err);
} }
paddle/cuda/src/hl_cuda_cnn.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/cuda/src/hl_cuda_lstm.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/cuda/src/hl_cuda_matrix.cu
浏览文件 @ 4bc1dc26
...@@ -12,22 +12,21 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,22 +12,21 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "hl_base.h" #include "hl_base.h"
#include "hl_device_functions.cuh"
#include "hl_gpu_matrix_kernel.cuh"
#include "hl_matrix.h" #include "hl_matrix.h"
#include "hl_matrix_ops.cuh"
#include "hl_matrix_apply.cuh" #include "hl_matrix_apply.cuh"
#include "hl_matrix_ops.cuh"
#include "hl_sequence.h" #include "hl_sequence.h"
#include "hl_sparse.ph" #include "hl_sparse.ph"
#include "paddle/utils/Logging.h" #include "paddle/utils/Logging.h"
#include "hl_device_functions.cuh"
#include "hl_gpu_matrix_kernel.cuh"
DEFINE_MATRIX_UNARY_OP(Zero, a = 0); DEFINE_MATRIX_UNARY_OP(Zero, a = 0);
DEFINE_MATRIX_TERNARY_PARAMETER_OP(_add, TWO_PARAMETER, c = p1*a + p2*b); DEFINE_MATRIX_TERNARY_PARAMETER_OP(_add, TWO_PARAMETER, c = p1 * a + p2 * b);
void hl_matrix_add(real *A_d, void hl_matrix_add(real* A_d,
real *B_d, real* B_d,
real *C_d, real* C_d,
int dimM, int dimM,
int dimN, int dimN,
real alpha, real alpha,
...@@ -36,8 +35,8 @@ void hl_matrix_add(real *A_d, ...@@ -36,8 +35,8 @@ void hl_matrix_add(real *A_d,
CHECK_NOTNULL(B_d); CHECK_NOTNULL(B_d);
CHECK_NOTNULL(C_d); CHECK_NOTNULL(C_d);
hl_gpu_apply_ternary_op hl_gpu_apply_ternary_op<real, ternary::_add<real>, 0, 0>(
<real, ternary::_add<real>, 0, 0>(ternary::_add<real>(alpha, beta), ternary::_add<real>(alpha, beta),
A_d, A_d,
B_d, B_d,
C_d, C_d,
...@@ -50,12 +49,11 @@ void hl_matrix_add(real *A_d, ...@@ -50,12 +49,11 @@ void hl_matrix_add(real *A_d,
} }
#ifdef PADDLE_TYPE_DOUBLE #ifdef PADDLE_TYPE_DOUBLE
#define THRESHOLD 128 #define THRESHOLD 128
#else #else
#define THRESHOLD 64 #define THRESHOLD 64
#endif #endif
__device__ __forceinline__ __device__ __forceinline__ void findMax(real* I,
void findMax(real* I,
real* dfMax_s, real* dfMax_s,
int blockSize, int blockSize,
int base, int base,
...@@ -89,8 +87,7 @@ void findMax(real* I, ...@@ -89,8 +87,7 @@ void findMax(real* I,
__syncthreads(); __syncthreads();
} }
__device__ __forceinline__ __device__ __forceinline__ void subMaxAndExp(real* I,
void subMaxAndExp(real* I,
real* O, real* O,
int curIdx, int curIdx,
int nextIdx, int nextIdx,
...@@ -115,8 +112,7 @@ void subMaxAndExp(real* I, ...@@ -115,8 +112,7 @@ void subMaxAndExp(real* I,
__syncthreads(); __syncthreads();
} }
__device__ __forceinline__ __device__ __forceinline__ void valueSum(real* O,
void valueSum(real* O,
real* dfMax_s, real* dfMax_s,
int blockSize, int blockSize,
int base, int base,
...@@ -141,13 +137,8 @@ void valueSum(real* O, ...@@ -141,13 +137,8 @@ void valueSum(real* O,
__syncthreads(); __syncthreads();
} }
__device__ __forceinline__ __device__ __forceinline__ void divSum(
void divSum(real* O, real* O, real sum, int curIdx, int nextIdx, int blockSize, int dimN) {
real sum,
int curIdx,
int nextIdx,
int blockSize,
int dimN) {
while (curIdx < dimN) { while (curIdx < dimN) {
O[nextIdx] /= sum; O[nextIdx] /= sum;
nextIdx += blockSize; nextIdx += blockSize;
...@@ -155,8 +146,7 @@ void divSum(real* O, ...@@ -155,8 +146,7 @@ void divSum(real* O,
} }
} }
__device__ __forceinline__ __device__ __forceinline__ void softmax(real* I,
void softmax(real* I,
real* O, real* O,
real* dfMax_s, real* dfMax_s,
int blockSize, int blockSize,
...@@ -167,8 +157,7 @@ void softmax(real* I, ...@@ -167,8 +157,7 @@ void softmax(real* I,
__shared__ real max; __shared__ real max;
// find the max number // find the max number
findMax(I, dfMax_s, blockSize, base, curIdx, findMax(I, dfMax_s, blockSize, base, curIdx, nextIdx, dimN, &max);
nextIdx, dimN, &max);
// sub max Value and do Exp operation // sub max Value and do Exp operation
subMaxAndExp(I, O, base, nextIdx, blockSize, dimN, max); subMaxAndExp(I, O, base, nextIdx, blockSize, dimN, max);
...@@ -181,8 +170,8 @@ void softmax(real* I, ...@@ -181,8 +170,8 @@ void softmax(real* I,
divSum(O, dfMax_s[0], curIdx, nextIdx, blockSize, dimN); divSum(O, dfMax_s[0], curIdx, nextIdx, blockSize, dimN);
} }
template<int blockSize> template <int blockSize>
__global__ void KeMatrixSoftMax(real *O, real *I, int dimN) { __global__ void KeMatrixSoftMax(real* O, real* I, int dimN) {
int base = threadIdx.x; int base = threadIdx.x;
__shared__ real dfMax_s[blockSize]; __shared__ real dfMax_s[blockSize];
int nextIdx = blockIdx.x * dimN + base; int nextIdx = blockIdx.x * dimN + base;
...@@ -191,19 +180,18 @@ __global__ void KeMatrixSoftMax(real *O, real *I, int dimN) { ...@@ -191,19 +180,18 @@ __global__ void KeMatrixSoftMax(real *O, real *I, int dimN) {
softmax(I, O, dfMax_s, blockSize, base, curIdx, nextIdx, dimN); softmax(I, O, dfMax_s, blockSize, base, curIdx, nextIdx, dimN);
} }
void hl_matrix_softmax(real *A_d, real *C_d, int dimM, int dimN) { void hl_matrix_softmax(real* A_d, real* C_d, int dimM, int dimN) {
CHECK_NOTNULL(A_d); CHECK_NOTNULL(A_d);
CHECK_NOTNULL(C_d); CHECK_NOTNULL(C_d);
dim3 block(512, 1); dim3 block(512, 1);
dim3 grid(dimM, 1); dim3 grid(dimM, 1);
KeMatrixSoftMax<512> KeMatrixSoftMax<512><<<grid, block, 0, STREAM_DEFAULT>>>(C_d, A_d, dimN);
<<<grid, block, 0, STREAM_DEFAULT>>>(C_d, A_d, dimN);
CHECK_SYNC("hl_matrix_softmax failed"); CHECK_SYNC("hl_matrix_softmax failed");
} }
template<int blockSize> template <int blockSize>
__global__ void KeSequenceSoftMax(real *O, real *I, const int* index) { __global__ void KeSequenceSoftMax(real* O, real* I, const int* index) {
int base = threadIdx.x; int base = threadIdx.x;
int bid = blockIdx.x; int bid = blockIdx.x;
__shared__ real dfMax_s[blockSize]; __shared__ real dfMax_s[blockSize];
...@@ -217,8 +205,8 @@ __global__ void KeSequenceSoftMax(real *O, real *I, const int* index) { ...@@ -217,8 +205,8 @@ __global__ void KeSequenceSoftMax(real *O, real *I, const int* index) {
softmax(I, O, dfMax_s, blockSize, base, curIdx, nextIdx, dimN); softmax(I, O, dfMax_s, blockSize, base, curIdx, nextIdx, dimN);
} }
void hl_sequence_softmax_forward(real *A_d, void hl_sequence_softmax_forward(real* A_d,
real *C_d, real* C_d,
const int* index, const int* index,
int numSequence) { int numSequence) {
CHECK_NOTNULL(A_d); CHECK_NOTNULL(A_d);
...@@ -226,59 +214,48 @@ void hl_sequence_softmax_forward(real *A_d, ...@@ -226,59 +214,48 @@ void hl_sequence_softmax_forward(real *A_d,
dim3 block(512, 1); dim3 block(512, 1);
dim3 grid(numSequence, 1); dim3 grid(numSequence, 1);
KeSequenceSoftMax<512> KeSequenceSoftMax<512><<<grid, block, 0, STREAM_DEFAULT>>>(C_d, A_d, index);
<<<grid, block, 0, STREAM_DEFAULT>>>(C_d, A_d, index);
CHECK_SYNC("hl_sequence_softmax_forward failed"); CHECK_SYNC("hl_sequence_softmax_forward failed");
} }
__global__ void KeMatrixDerivative(real *grad_d, __global__ void KeMatrixDerivative(
real *output_d, real* grad_d, real* output_d, real* sftmaxSum_d, int dimM, int dimN) {
real *sftmaxSum_d, int rowIdx = blockIdx.x * blockDim.x + threadIdx.x;
int dimM, int colIdx = blockIdx.y * blockDim.y + threadIdx.y;
int dimN) {
int rowIdx = blockIdx.x*blockDim.x + threadIdx.x;
int colIdx = blockIdx.y*blockDim.y + threadIdx.y;
int index; int index;
if (rowIdx < dimM && colIdx < dimN) { if (rowIdx < dimM && colIdx < dimN) {
index = rowIdx*dimN + colIdx; index = rowIdx * dimN + colIdx;
grad_d[index] = output_d[index] * (grad_d[index] - sftmaxSum_d[rowIdx]); grad_d[index] = output_d[index] * (grad_d[index] - sftmaxSum_d[rowIdx]);
} }
} }
void hl_matrix_softmax_derivative(real *grad_d, void hl_matrix_softmax_derivative(
real *output_d, real* grad_d, real* output_d, real* sftmaxSum_d, int dimM, int dimN) {
real *sftmaxSum_d,
int dimM,
int dimN) {
CHECK_NOTNULL(grad_d); CHECK_NOTNULL(grad_d);
CHECK_NOTNULL(output_d); CHECK_NOTNULL(output_d);
CHECK_NOTNULL(sftmaxSum_d); CHECK_NOTNULL(sftmaxSum_d);
int blocksX = (dimM + 0) / 1; int blocksX = (dimM + 0) / 1;
int blocksY = (dimN + 1024 -1) / 1024; int blocksY = (dimN + 1024 - 1) / 1024;
dim3 threads(1, 1024); dim3 threads(1, 1024);
dim3 grid(blocksX, blocksY); dim3 grid(blocksX, blocksY);
KeMatrixDerivative<<< grid, threads, 0, STREAM_DEFAULT >>> KeMatrixDerivative<<<grid, threads, 0, STREAM_DEFAULT>>>(
(grad_d, output_d, sftmaxSum_d, dimM, dimN); grad_d, output_d, sftmaxSum_d, dimM, dimN);
CHECK_SYNC("hl_matrix_softmax_derivative failed"); CHECK_SYNC("hl_matrix_softmax_derivative failed");
} }
__global__ void KeMatrixMultiBinaryCrossEntropy(real* output, __global__ void KeMatrixMultiBinaryCrossEntropy(
real* entropy, real* output, real* entropy, int* row, int* col, int dimM, int dimN) {
int* row,
int* col,
int dimM,
int dimN) {
int index = blockIdx.x * blockDim.x + threadIdx.x; int index = blockIdx.x * blockDim.x + threadIdx.x;
if (index < dimM) { if (index < dimM) {
for (int i = 0; i < dimN; i ++) { for (int i = 0; i < dimN; i++) {
entropy[index] -= log(1 - output[index * dimN + i]); entropy[index] -= log(1 - output[index * dimN + i]);
} }
int *row_col = col + row[index]; int* row_col = col + row[index];
int col_num = row[index + 1] - row[index]; int col_num = row[index + 1] - row[index];
for (int i = 0; i < col_num; i ++) { for (int i = 0; i < col_num; i++) {
real o = output[index * dimN + row_col[i]]; real o = output[index * dimN + row_col[i]];
entropy[index] -= log(o / (1 - o)); entropy[index] -= log(o / (1 - o));
} }
...@@ -299,37 +276,30 @@ void hl_matrix_multi_binary_cross_entropy(real* output, ...@@ -299,37 +276,30 @@ void hl_matrix_multi_binary_cross_entropy(real* output,
dim3 threads(n_threads); dim3 threads(n_threads);
dim3 grid(blocks); dim3 grid(blocks);
hl_csr_matrix mat = (hl_csr_matrix)(csr_mat->matrix); hl_csr_matrix mat = (hl_csr_matrix)(csr_mat->matrix);
KeMatrixMultiBinaryCrossEntropy<<< grid, threads, 0, STREAM_DEFAULT >>> KeMatrixMultiBinaryCrossEntropy<<<grid, threads, 0, STREAM_DEFAULT>>>(
(output, entropy, mat->csr_row, mat->csr_col, dimM, dimN); output, entropy, mat->csr_row, mat->csr_col, dimM, dimN);
CHECK_SYNC("hl_matrix_multi_binary_cross_entropy failed"); CHECK_SYNC("hl_matrix_multi_binary_cross_entropy failed");
} }
__global__ void KeMatrixMultiBinaryCrossEntropyBp(real* output, __global__ void KeMatrixMultiBinaryCrossEntropyBp(
real* grad, real* output, real* grad, int* row, int* col, int dimM, int dimN) {
int* row,
int* col,
int dimM,
int dimN) {
int row_idx = blockIdx.x * blockDim.x + threadIdx.x; int row_idx = blockIdx.x * blockDim.x + threadIdx.x;
if (row_idx < dimM) { if (row_idx < dimM) {
for (int i = 0; i < dimN; i ++) { for (int i = 0; i < dimN; i++) {
int index = row_idx * dimN + i; int index = row_idx * dimN + i;
grad[index] += 1.0 / (1 - output[index]); grad[index] += 1.0 / (1 - output[index]);
} }
int col_num = row[row_idx + 1] - row[row_idx]; int col_num = row[row_idx + 1] - row[row_idx];
int *row_col = col + row[row_idx]; int* row_col = col + row[row_idx];
for (int i = 0; i < col_num; i ++) { for (int i = 0; i < col_num; i++) {
int index = row_idx * dimN + row_col[i]; int index = row_idx * dimN + row_col[i];
grad[index] -= 1.0 / (output[index] * (1 - output[index])); grad[index] -= 1.0 / (output[index] * (1 - output[index]));
} }
} }
} }
void hl_matrix_multi_binary_cross_entropy_bp(real* output, void hl_matrix_multi_binary_cross_entropy_bp(
real* grad, real* output, real* grad, hl_sparse_matrix_s csr_mat, int dimM, int dimN) {
hl_sparse_matrix_s csr_mat,
int dimM,
int dimN) {
CHECK_NOTNULL(output); CHECK_NOTNULL(output);
CHECK_NOTNULL(grad); CHECK_NOTNULL(grad);
CHECK_NOTNULL(csr_mat); CHECK_NOTNULL(csr_mat);
...@@ -339,16 +309,13 @@ void hl_matrix_multi_binary_cross_entropy_bp(real* output, ...@@ -339,16 +309,13 @@ void hl_matrix_multi_binary_cross_entropy_bp(real* output,
dim3 threads(n_threads); dim3 threads(n_threads);
dim3 grid(blocks); dim3 grid(blocks);
hl_csr_matrix mat = (hl_csr_matrix)(csr_mat->matrix); hl_csr_matrix mat = (hl_csr_matrix)(csr_mat->matrix);
KeMatrixMultiBinaryCrossEntropyBp<<< grid, threads, 0, STREAM_DEFAULT >>> KeMatrixMultiBinaryCrossEntropyBp<<<grid, threads, 0, STREAM_DEFAULT>>>(
(output, grad, mat->csr_row, mat->csr_col, dimM, dimN); output, grad, mat->csr_row, mat->csr_col, dimM, dimN);
CHECK_SYNC("hl_matrix_multi_binary_cross_entropy_bp failed"); CHECK_SYNC("hl_matrix_multi_binary_cross_entropy_bp failed");
} }
__global__ void KeMatrixCrossEntropy(real* O, __global__ void KeMatrixCrossEntropy(
real* E, real* O, real* E, int* label, int dimM, int dimN) {
int* label,
int dimM,
int dimN) {
int index = blockIdx.x * blockDim.x + threadIdx.x; int index = blockIdx.x * blockDim.x + threadIdx.x;
int newBase; int newBase;
if (index < dimM) { if (index < dimM) {
...@@ -358,59 +325,49 @@ __global__ void KeMatrixCrossEntropy(real* O, ...@@ -358,59 +325,49 @@ __global__ void KeMatrixCrossEntropy(real* O,
} }
} }
void hl_matrix_cross_entropy(real* A_d, void hl_matrix_cross_entropy(
real* C_d, real* A_d, real* C_d, int* label_d, int dimM, int dimN) {
int* label_d,
int dimM,
int dimN) {
CHECK_NOTNULL(A_d); CHECK_NOTNULL(A_d);
CHECK_NOTNULL(C_d); CHECK_NOTNULL(C_d);
int blocks = (dimM + 1024 - 1) / 1024; int blocks = (dimM + 1024 - 1) / 1024;
dim3 threads(1024, 1); dim3 threads(1024, 1);
dim3 grid(blocks, 1); dim3 grid(blocks, 1);
KeMatrixCrossEntropy<<< grid, threads, 0, STREAM_DEFAULT >>> KeMatrixCrossEntropy<<<grid, threads, 0, STREAM_DEFAULT>>>(
(A_d, C_d, label_d, dimM, dimN); A_d, C_d, label_d, dimM, dimN);
CHECK_SYNC("hl_matrix_cross_entropy failed"); CHECK_SYNC("hl_matrix_cross_entropy failed");
} }
__global__ void KeMatrixCrossEntropyBp(real* grad_d, __global__ void KeMatrixCrossEntropyBp(
real* output_d, real* grad_d, real* output_d, int* label_d, int dimM, int dimN) {
int* label_d, int rowIdx = blockIdx.x * blockDim.x + threadIdx.x;
int dimM, int colIdx = blockIdx.y * blockDim.y + threadIdx.y;
int dimN) {
int rowIdx = blockIdx.x*blockDim.x + threadIdx.x;
int colIdx = blockIdx.y*blockDim.y + threadIdx.y;
int index; int index;
if (rowIdx < dimM && colIdx < dimN) { if (rowIdx < dimM && colIdx < dimN) {
index = rowIdx*dimN + colIdx; index = rowIdx * dimN + colIdx;
if (label_d[rowIdx] == colIdx) { if (label_d[rowIdx] == colIdx) {
grad_d[index] -= 1.0f / output_d[index]; grad_d[index] -= 1.0f / output_d[index];
} }
} }
} }
void hl_matrix_cross_entropy_bp(real* grad_d, void hl_matrix_cross_entropy_bp(
real* output_d, real* grad_d, real* output_d, int* label_d, int dimM, int dimN) {
int* label_d,
int dimM,
int dimN) {
CHECK_NOTNULL(grad_d); CHECK_NOTNULL(grad_d);
CHECK_NOTNULL(output_d); CHECK_NOTNULL(output_d);
CHECK_NOTNULL(label_d); CHECK_NOTNULL(label_d);
int blocksX = (dimM + 0)/1; int blocksX = (dimM + 0) / 1;
int blocksY = (dimN + 1024 -1) / 1024; int blocksY = (dimN + 1024 - 1) / 1024;
dim3 threads(1, 1024); dim3 threads(1, 1024);
dim3 grid(blocksX, blocksY); dim3 grid(blocksX, blocksY);
KeMatrixCrossEntropyBp<<< grid, threads, 0, STREAM_DEFAULT >>> KeMatrixCrossEntropyBp<<<grid, threads, 0, STREAM_DEFAULT>>>(
(grad_d, output_d, label_d, dimM, dimN); grad_d, output_d, label_d, dimM, dimN);
CHECK_SYNC("hl_matrix_cross_entropy_bp failed"); CHECK_SYNC("hl_matrix_cross_entropy_bp failed");
} }
void hl_matrix_zero_mem(real* data, int num) { void hl_matrix_zero_mem(real* data, int num) {
hl_gpu_apply_unary_op( hl_gpu_apply_unary_op(unary::Zero<real>(), data, 1, num, num);
unary::Zero<real>(), data, 1, num, num);
} }
__global__ void KeParamReluForward(real* output, __global__ void KeParamReluForward(real* output,
...@@ -423,8 +380,8 @@ __global__ void KeParamReluForward(real* output, ...@@ -423,8 +380,8 @@ __global__ void KeParamReluForward(real* output,
int ty = blockIdx.y * blockDim.y + threadIdx.y; int ty = blockIdx.y * blockDim.y + threadIdx.y;
if (tx < width && ty < height) { if (tx < width && ty < height) {
int index = ty * width + tx; int index = ty * width + tx;
output[index] = input[index] > 0 ? input[index] : output[index] =
input[index] * w[tx / partial_sum]; input[index] > 0 ? input[index] : input[index] * w[tx / partial_sum];
} }
} }
...@@ -439,14 +396,14 @@ void hl_param_relu_forward(real* output, ...@@ -439,14 +396,14 @@ void hl_param_relu_forward(real* output,
CHECK_NOTNULL(w); CHECK_NOTNULL(w);
dim3 threads(16, 16); dim3 threads(16, 16);
int blockX = (width + 16 - 1) / 16; int blockX = (width + 16 - 1) / 16;
int blockY = (height + 16 -1) / 16; int blockY = (height + 16 - 1) / 16;
dim3 grid(blockX, blockY); dim3 grid(blockX, blockY);
KeParamReluForward<<<grid, threads, 0, STREAM_DEFAULT>>> KeParamReluForward<<<grid, threads, 0, STREAM_DEFAULT>>>(
(output, input, w, width, height, partial_sum); output, input, w, width, height, partial_sum);
CHECK_SYNC("hl_param_relu_forward failed"); CHECK_SYNC("hl_param_relu_forward failed");
} }
template<int blockSize> template <int blockSize>
__global__ void KeParamReluBackWardW(real* grad_w, __global__ void KeParamReluBackWardW(real* grad_w,
real* grad_o, real* grad_o,
real* input, real* input,
...@@ -491,8 +448,8 @@ void hl_param_relu_backward_w(real* grad_w, ...@@ -491,8 +448,8 @@ void hl_param_relu_backward_w(real* grad_w,
int grid_num = width / partial_sum; int grid_num = width / partial_sum;
dim3 threads(blockSize, 1); dim3 threads(blockSize, 1);
dim3 grid(grid_num, 1); dim3 grid(grid_num, 1);
KeParamReluBackWardW<blockSize><<<grid, threads, 0, STREAM_DEFAULT>>> KeParamReluBackWardW<blockSize><<<grid, threads, 0, STREAM_DEFAULT>>>(
(grad_w, grad_o, input, width, height, partial_sum); grad_w, grad_o, input, width, height, partial_sum);
CHECK_SYNC("hl_param_relu_backward_w failed"); CHECK_SYNC("hl_param_relu_backward_w failed");
} }
...@@ -524,19 +481,15 @@ void hl_param_relu_backward_diff(real* grad_o, ...@@ -524,19 +481,15 @@ void hl_param_relu_backward_diff(real* grad_o,
CHECK_NOTNULL(diff); CHECK_NOTNULL(diff);
dim3 threads(16, 16); dim3 threads(16, 16);
int blockX = (width + 16 - 1) / 16; int blockX = (width + 16 - 1) / 16;
int blockY = (height + 16 -1) / 16; int blockY = (height + 16 - 1) / 16;
dim3 grid(blockX, blockY); dim3 grid(blockX, blockY);
KeParamReluBackwardDiff<<<grid, threads, 0, STREAM_DEFAULT>>> KeParamReluBackwardDiff<<<grid, threads, 0, STREAM_DEFAULT>>>(
(grad_o, data, w, diff, width, height, partial_sum); grad_o, data, w, diff, width, height, partial_sum);
CHECK_SYNC("hl_param_relu_backward_diff failed"); CHECK_SYNC("hl_param_relu_backward_diff failed");
} }
__global__ void KeMatrixAddSharedBias(real* A, __global__ void KeMatrixAddSharedBias(
real* B, real* A, real* B, const int channel, const int M, const int N, real scale) {
const int channel,
const int M,
const int N,
real scale) {
int index = blockIdx.x * blockDim.x + threadIdx.x; int index = blockIdx.x * blockDim.x + threadIdx.x;
int dim = N / channel; int dim = N / channel;
if (index < M * N) { if (index < M * N) {
...@@ -554,15 +507,14 @@ void hl_matrix_add_shared_bias(real* A_d, ...@@ -554,15 +507,14 @@ void hl_matrix_add_shared_bias(real* A_d,
real scale) { real scale) {
const int blocks = 512; const int blocks = 512;
const int grids = DIVUP(dimM * dimN, blocks); const int grids = DIVUP(dimM * dimN, blocks);
KeMatrixAddSharedBias<<<grids, blocks, 0, STREAM_DEFAULT>>> KeMatrixAddSharedBias<<<grids, blocks, 0, STREAM_DEFAULT>>>(
(A_d, B_d, channel, dimM, dimN, scale); A_d, B_d, channel, dimM, dimN, scale);
CHECK_SYNC("hl_matrix_add_shared_bias failed"); CHECK_SYNC("hl_matrix_add_shared_bias failed");
} }
template <int blockSize> template <int blockSize>
__global__ void KeMatrixCollectSharedBias(real *B, __global__ void KeMatrixCollectSharedBias(real* B,
real *A, real* A,
const int channel, const int channel,
const int M, const int M,
const int N, const int N,
...@@ -611,14 +563,13 @@ void hl_matrix_collect_shared_bias(real* B_d, ...@@ -611,14 +563,13 @@ void hl_matrix_collect_shared_bias(real* B_d,
const int limit = 64; const int limit = 64;
int grids = (dimM * dim) < limit ? DIVUP(channel, blocks) : channel; int grids = (dimM * dim) < limit ? DIVUP(channel, blocks) : channel;
KeMatrixCollectSharedBias<blocks> KeMatrixCollectSharedBias<blocks><<<grids, blocks, 0, STREAM_DEFAULT>>>(
<<< grids, blocks, 0, STREAM_DEFAULT>>> B_d, A_d, channel, dimM, dimN, dim, limit, scale);
(B_d, A_d, channel, dimM, dimN, dim, limit, scale);
CHECK_SYNC("hl_matrix_collect_shared_bias failed"); CHECK_SYNC("hl_matrix_collect_shared_bias failed");
} }
__global__ void keMatrixRotate(real* mat, real* matRot, __global__ void keMatrixRotate(
int dimM, int dimN, bool clockWise) { real* mat, real* matRot, int dimM, int dimN, bool clockWise) {
int idx = blockIdx.x * blockDim.x + threadIdx.x; int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < dimM * dimN) { if (idx < dimM * dimN) {
int i = idx / dimN; int i = idx / dimN;
...@@ -631,13 +582,13 @@ __global__ void keMatrixRotate(real* mat, real* matRot, ...@@ -631,13 +582,13 @@ __global__ void keMatrixRotate(real* mat, real* matRot,
} }
} }
void hl_matrix_rotate(real *mat, real* matRot, void hl_matrix_rotate(
int dimM, int dimN, bool clockWise) { real* mat, real* matRot, int dimM, int dimN, bool clockWise) {
CHECK_NOTNULL(mat); CHECK_NOTNULL(mat);
CHECK_NOTNULL(matRot); CHECK_NOTNULL(matRot);
const int threads = 512; const int threads = 512;
const int blocks = DIVUP(dimM * dimN, threads); const int blocks = DIVUP(dimM * dimN, threads);
keMatrixRotate<<< blocks, threads, 0, STREAM_DEFAULT >>> keMatrixRotate<<<blocks, threads, 0, STREAM_DEFAULT>>>(
(mat, matRot, dimM, dimN, clockWise); mat, matRot, dimM, dimN, clockWise);
CHECK_SYNC("hl_matrix_rotate failed"); CHECK_SYNC("hl_matrix_rotate failed");
} }
paddle/cuda/src/hl_cuda_sequence.cu
浏览文件 @ 4bc1dc26
...@@ -16,36 +16,36 @@ limitations under the License. */ ...@@ -16,36 +16,36 @@ limitations under the License. */
#include "hl_device_functions.cuh" #include "hl_device_functions.cuh"
#include "paddle/utils/Logging.h" #include "paddle/utils/Logging.h"
__global__ void KeMaxSequenceForward(real *input, __global__ void KeMaxSequenceForward(real* input,
const int *sequence, const int* sequence,
real* output, real* output,
int *index, int* index,
int numSequences, int numSequences,
int dim) { int dim) {
int dimIdx = threadIdx.x; int dimIdx = threadIdx.x;
int sequenceId = blockIdx.x; int sequenceId = blockIdx.x;
if (sequenceId >= numSequences) return; if (sequenceId >= numSequences) return;
int start = sequence[sequenceId]; int start = sequence[sequenceId];
int end = sequence[sequenceId+1]; int end = sequence[sequenceId + 1];
for (int i = dimIdx; i < dim; i += blockDim.x) { for (int i = dimIdx; i < dim; i += blockDim.x) {
real tmp = -HL_FLOAT_MAX; real tmp = -HL_FLOAT_MAX;
int tmpId = -1; int tmpId = -1;
for (int insId = start; insId < end; insId++) { for (int insId = start; insId < end; insId++) {
if (tmp < input[insId*dim + i]) { if (tmp < input[insId * dim + i]) {
tmp = input[insId*dim + i]; tmp = input[insId * dim + i];
tmpId = insId; tmpId = insId;
} }
} }
output[sequenceId*dim + i] = tmp; output[sequenceId * dim + i] = tmp;
index[sequenceId*dim + i] = tmpId; index[sequenceId * dim + i] = tmpId;
} }
} }
void hl_max_sequence_forward(real* input, void hl_max_sequence_forward(real* input,
const int* sequence, const int* sequence,
real* output, real* output,
int *index, int* index,
int numSequences, int numSequences,
int dim) { int dim) {
CHECK_NOTNULL(input); CHECK_NOTNULL(input);
...@@ -55,29 +55,23 @@ void hl_max_sequence_forward(real* input, ...@@ -55,29 +55,23 @@ void hl_max_sequence_forward(real* input,
dim3 threads(256, 1); dim3 threads(256, 1);
dim3 grid(numSequences, 1); dim3 grid(numSequences, 1);
KeMaxSequenceForward<<< grid, threads, 0, STREAM_DEFAULT >>> KeMaxSequenceForward<<<grid, threads, 0, STREAM_DEFAULT>>>(
(input, sequence, output, index, numSequences, dim); input, sequence, output, index, numSequences, dim);
CHECK_SYNC("hl_max_sequence_forward failed"); CHECK_SYNC("hl_max_sequence_forward failed");
} }
__global__ void KeMaxSequenceBackward(real *outputGrad, __global__ void KeMaxSequenceBackward(
int *index, real* outputGrad, int* index, real* inputGrad, int numSequences, int dim) {
real* inputGrad,
int numSequences,
int dim) {
int idx = threadIdx.x + blockIdx.x * blockDim.x; int idx = threadIdx.x + blockIdx.x * blockDim.x;
int colIdx = idx % dim; int colIdx = idx % dim;
if (idx < numSequences*dim) { if (idx < numSequences * dim) {
int insId = index[idx]; int insId = index[idx];
inputGrad[insId * dim + colIdx] += outputGrad[idx]; inputGrad[insId * dim + colIdx] += outputGrad[idx];
} }
} }
void hl_max_sequence_backward(real* outputGrad, void hl_max_sequence_backward(
int *index, real* outputGrad, int* index, real* inputGrad, int numSequences, int dim) {
real* inputGrad,
int numSequences,
int dim) {
CHECK_NOTNULL(outputGrad); CHECK_NOTNULL(outputGrad);
CHECK_NOTNULL(index); CHECK_NOTNULL(index);
CHECK_NOTNULL(inputGrad); CHECK_NOTNULL(inputGrad);
...@@ -85,12 +79,12 @@ void hl_max_sequence_backward(real* outputGrad, ...@@ -85,12 +79,12 @@ void hl_max_sequence_backward(real* outputGrad,
unsigned int blocks = (numSequences * dim + 128 - 1) / 128; unsigned int blocks = (numSequences * dim + 128 - 1) / 128;
dim3 threads(128, 1); dim3 threads(128, 1);
dim3 grid(blocks, 1); dim3 grid(blocks, 1);
KeMaxSequenceBackward<<< grid, threads, 0, STREAM_DEFAULT >>> KeMaxSequenceBackward<<<grid, threads, 0, STREAM_DEFAULT>>>(
(outputGrad, index, inputGrad, numSequences, dim); outputGrad, index, inputGrad, numSequences, dim);
CHECK_SYNC("hl_max_sequence_backward failed"); CHECK_SYNC("hl_max_sequence_backward failed");
} }
template<int blockDimX, int blockDimY, int gridDimX, bool AddRow> template <int blockDimX, int blockDimY, int gridDimX, bool AddRow>
__global__ void KeMatrixAddRows(real* output, __global__ void KeMatrixAddRows(real* output,
real* table, real* table,
int* ids, int* ids,
...@@ -104,8 +98,8 @@ __global__ void KeMatrixAddRows(real* output, ...@@ -104,8 +98,8 @@ __global__ void KeMatrixAddRows(real* output,
while (sampleId < numSamples) { while (sampleId < numSamples) {
int tableId = ids[sampleId]; int tableId = ids[sampleId];
if ((0 <= tableId) && (tableId < tableSize)) { if ((0 <= tableId) && (tableId < tableSize)) {
real *outputData = output + sampleId * dim; real* outputData = output + sampleId * dim;
real *tableData = table + tableId * dim; real* tableData = table + tableId * dim;
for (int i = idx; i < dim; i += blockDimX) { for (int i = idx; i < dim; i += blockDimX) {
if (AddRow == 0) { if (AddRow == 0) {
outputData[i] += tableData[i]; outputData[i] += tableData[i];
...@@ -114,15 +108,18 @@ __global__ void KeMatrixAddRows(real* output, ...@@ -114,15 +108,18 @@ __global__ void KeMatrixAddRows(real* output,
} }
} }
} }
sampleId += blockDimY*gridDimX; sampleId += blockDimY * gridDimX;
} }
} }
template<int blockDimX, int blockDimY, int gridDimX, bool seq2batch, bool isAdd> template <int blockDimX,
__global__ int blockDimY,
void KeSequence2Batch(real *batch, int gridDimX,
real *sequence, bool seq2batch,
const int *batchIndex, bool isAdd>
__global__ void KeSequence2Batch(real* batch,
real* sequence,
const int* batchIndex,
int seqWidth, int seqWidth,
int batchCount) { int batchCount) {
int idx = threadIdx.x; int idx = threadIdx.x;
...@@ -130,8 +127,8 @@ void KeSequence2Batch(real *batch, ...@@ -130,8 +127,8 @@ void KeSequence2Batch(real *batch,
int id = blockIdx.x + idy * gridDimX; int id = blockIdx.x + idy * gridDimX;
while (id < batchCount) { while (id < batchCount) {
int seqId = batchIndex[id]; int seqId = batchIndex[id];
real* batchData = batch + id*seqWidth; real* batchData = batch + id * seqWidth;
real* seqData = sequence + seqId*seqWidth; real* seqData = sequence + seqId * seqWidth;
for (int i = idx; i < seqWidth; i += blockDimX) { for (int i = idx; i < seqWidth; i += blockDimX) {
if (seq2batch) { if (seq2batch) {
if (isAdd) { if (isAdd) {
...@@ -147,13 +144,13 @@ void KeSequence2Batch(real *batch, ...@@ -147,13 +144,13 @@ void KeSequence2Batch(real *batch,
} }
} }
} }
id += blockDimY*gridDimX; id += blockDimY * gridDimX;
} }
} }
void hl_sequence2batch_copy(real *batch, void hl_sequence2batch_copy(real* batch,
real *sequence, real* sequence,
const int *batchIndex, const int* batchIndex,
int seqWidth, int seqWidth,
int batchCount, int batchCount,
bool seq2batch) { bool seq2batch) {
...@@ -164,18 +161,18 @@ void hl_sequence2batch_copy(real *batch, ...@@ -164,18 +161,18 @@ void hl_sequence2batch_copy(real *batch,
dim3 threads(128, 8); dim3 threads(128, 8);
dim3 grid(8, 1); dim3 grid(8, 1);
if (seq2batch) { if (seq2batch) {
KeSequence2Batch<128, 8, 8, 1, 0><<< grid, threads, 0, STREAM_DEFAULT >>> KeSequence2Batch<128, 8, 8, 1, 0><<<grid, threads, 0, STREAM_DEFAULT>>>(
(batch, sequence, batchIndex, seqWidth, batchCount); batch, sequence, batchIndex, seqWidth, batchCount);
} else { } else {
KeSequence2Batch<128, 8, 8, 0, 0><<< grid, threads, 0, STREAM_DEFAULT >>> KeSequence2Batch<128, 8, 8, 0, 0><<<grid, threads, 0, STREAM_DEFAULT>>>(
(batch, sequence, batchIndex, seqWidth, batchCount); batch, sequence, batchIndex, seqWidth, batchCount);
} }
CHECK_SYNC("hl_sequence2batch_copy failed"); CHECK_SYNC("hl_sequence2batch_copy failed");
} }
void hl_sequence2batch_add(real *batch, void hl_sequence2batch_add(real* batch,
real *sequence, real* sequence,
int *batchIndex, int* batchIndex,
int seqWidth, int seqWidth,
int batchCount, int batchCount,
bool seq2batch) { bool seq2batch) {
...@@ -186,18 +183,17 @@ void hl_sequence2batch_add(real *batch, ...@@ -186,18 +183,17 @@ void hl_sequence2batch_add(real *batch,
dim3 threads(128, 8); dim3 threads(128, 8);
dim3 grid(8, 1); dim3 grid(8, 1);
if (seq2batch) { if (seq2batch) {
KeSequence2Batch<128, 8, 8, 1, 1><<< grid, threads, 0, STREAM_DEFAULT >>> KeSequence2Batch<128, 8, 8, 1, 1><<<grid, threads, 0, STREAM_DEFAULT>>>(
(batch, sequence, batchIndex, seqWidth, batchCount); batch, sequence, batchIndex, seqWidth, batchCount);
} else { } else {
KeSequence2Batch<128, 8, 8, 0, 1><<< grid, threads, 0, STREAM_DEFAULT >>> KeSequence2Batch<128, 8, 8, 0, 1><<<grid, threads, 0, STREAM_DEFAULT>>>(
(batch, sequence, batchIndex, seqWidth, batchCount); batch, sequence, batchIndex, seqWidth, batchCount);
} }
CHECK_SYNC("hl_sequence2batch_add failed"); CHECK_SYNC("hl_sequence2batch_add failed");
} }
template<bool normByTimes, bool seq2batch> template <bool normByTimes, bool seq2batch>
__global__ __global__ void KeSequence2BatchPadding(real* batch,
void KeSequence2BatchPadding(real* batch,
real* sequence, real* sequence,
const int* sequenceStartPositions, const int* sequenceStartPositions,
const size_t sequenceWidth, const size_t sequenceWidth,
...@@ -276,37 +272,49 @@ void hl_sequence2batch_copy_padding(real* batch, ...@@ -276,37 +272,49 @@ void hl_sequence2batch_copy_padding(real* batch,
if (seq2batch) { if (seq2batch) {
/* sequence -> batch */ /* sequence -> batch */
if (normByTimes) { if (normByTimes) {
KeSequence2BatchPadding<1, 1><<< grid, threads, 0, STREAM_DEFAULT >>>( KeSequence2BatchPadding<1, 1><<<grid, threads, 0, STREAM_DEFAULT>>>(
batch, sequence, sequenceStartPositions, batch,
sequenceWidth, maxSequenceLength, numSequences); sequence,
sequenceStartPositions,
sequenceWidth,
maxSequenceLength,
numSequences);
} else { } else {
KeSequence2BatchPadding<0, 1><<< grid, threads, 0, STREAM_DEFAULT >>>( KeSequence2BatchPadding<0, 1><<<grid, threads, 0, STREAM_DEFAULT>>>(
batch, sequence, sequenceStartPositions, batch,
sequenceWidth, maxSequenceLength, numSequences); sequence,
sequenceStartPositions,
sequenceWidth,
maxSequenceLength,
numSequences);
} }
} else { } else {
/* batch -> sequence */ /* batch -> sequence */
if (normByTimes) { if (normByTimes) {
KeSequence2BatchPadding<1, 0><<< grid, threads, 0, STREAM_DEFAULT >>>( KeSequence2BatchPadding<1, 0><<<grid, threads, 0, STREAM_DEFAULT>>>(
batch, sequence, sequenceStartPositions, batch,
sequenceWidth, maxSequenceLength, numSequences); sequence,
sequenceStartPositions,
sequenceWidth,
maxSequenceLength,
numSequences);
} else { } else {
KeSequence2BatchPadding<0, 0><<< grid, threads, 0, STREAM_DEFAULT >>>( KeSequence2BatchPadding<0, 0><<<grid, threads, 0, STREAM_DEFAULT>>>(
batch, sequence, sequenceStartPositions, batch,
sequenceWidth, maxSequenceLength, numSequences); sequence,
sequenceStartPositions,
sequenceWidth,
maxSequenceLength,
numSequences);
} }
} }
CHECK_SYNC("hl_sequence2batch_copy_padding failed"); CHECK_SYNC("hl_sequence2batch_copy_padding failed");
} }
__device__ inline float my_rsqrt(float x) { __device__ inline float my_rsqrt(float x) { return rsqrtf(x); }
return rsqrtf(x);
}
__device__ inline double my_rsqrt(double x) { __device__ inline double my_rsqrt(double x) { return rsqrt(x); }
return rsqrt(x);
}
__global__ void KeSequenceAvgForward(real* dst, __global__ void KeSequenceAvgForward(real* dst,
real* src, real* src,
...@@ -327,8 +335,8 @@ __global__ void KeSequenceAvgForward(real* dst, ...@@ -327,8 +335,8 @@ __global__ void KeSequenceAvgForward(real* dst,
for (int i = start; i < end; i++) { for (int i = start; i < end; i++) {
sum += src[i * width + col]; sum += src[i * width + col];
} }
sum = mode == 1 ? sum : sum = mode == 1 ? sum : (mode == 0 ? sum / seqLength
(mode == 0 ? sum / seqLength : sum * my_rsqrt((real)seqLength)); : sum * my_rsqrt((real)seqLength));
dst[gid] += sum; dst[gid] += sum;
} }
} }
...@@ -349,8 +357,8 @@ void hl_sequence_avg_forward(real* dst, ...@@ -349,8 +357,8 @@ void hl_sequence_avg_forward(real* dst,
CHECK(mode == 0 || mode == 1 || mode == 2) CHECK(mode == 0 || mode == 1 || mode == 2)
<< "mode error in hl_sequence_avg_forward!"; << "mode error in hl_sequence_avg_forward!";
KeSequenceAvgForward<<< grid, block, 0, STREAM_DEFAULT >>> KeSequenceAvgForward<<<grid, block, 0, STREAM_DEFAULT>>>(
(dst, src, starts, height, width, mode); dst, src, starts, height, width, mode);
CHECK_SYNC("hl_sequence_avg_forward failed"); CHECK_SYNC("hl_sequence_avg_forward failed");
} }
...@@ -370,8 +378,8 @@ __global__ void KeSequenceAvgBackward(real* dst, ...@@ -370,8 +378,8 @@ __global__ void KeSequenceAvgBackward(real* dst,
int seqLength = end - start; int seqLength = end - start;
if (seqLength == 0) return; if (seqLength == 0) return;
real grad = src[gid]; real grad = src[gid];
grad = mode == 1 ? grad : grad = mode == 1 ? grad : (mode == 0 ? grad / seqLength
(mode == 0 ? grad / seqLength : grad * my_rsqrt((real)seqLength)); : grad * my_rsqrt((real)seqLength));
for (int i = start; i < end; i++) { for (int i = start; i < end; i++) {
dst[i * width + col] += grad; dst[i * width + col] += grad;
} }
...@@ -394,7 +402,7 @@ void hl_sequence_avg_backward(real* dst, ...@@ -394,7 +402,7 @@ void hl_sequence_avg_backward(real* dst,
CHECK(mode == 0 || mode == 1 || mode == 2) CHECK(mode == 0 || mode == 1 || mode == 2)
<< "mode error in hl_sequence_avg_backward!"; << "mode error in hl_sequence_avg_backward!";
KeSequenceAvgBackward<<< grid, block, 0, STREAM_DEFAULT >>> KeSequenceAvgBackward<<<grid, block, 0, STREAM_DEFAULT>>>(
(dst, src, starts, height, width, mode); dst, src, starts, height, width, mode);
CHECK_SYNC("hl_sequence_avg_backward failed"); CHECK_SYNC("hl_sequence_avg_backward failed");
} }
paddle/cuda/src/hl_cuda_sparse.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/cuda/src/hl_perturbation_util.cu
浏览文件 @ 4bc1dc26
...@@ -12,13 +12,12 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,13 +12,12 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include <cmath>
#include <stdlib.h> #include <stdlib.h>
#include "hl_cuda.h" #include <cmath>
#include "hl_time.h"
#include "hl_base.h" #include "hl_base.h"
#include "hl_cuda.h"
#include "hl_perturbation_util.cuh" #include "hl_perturbation_util.cuh"
#include "hl_time.h"
#define _USE_MATH_DEFINES #define _USE_MATH_DEFINES
...@@ -30,10 +29,16 @@ limitations under the License. */ ...@@ -30,10 +29,16 @@ limitations under the License. */
* centerX, centerY: translation. * centerX, centerY: translation.
* sourceX, sourceY: output coordinates in the original image. * sourceX, sourceY: output coordinates in the original image.
*/ */
__device__ void getTranformCoord(int x, int y, real theta, real scale, __device__ void getTranformCoord(int x,
real tgtCenter, real imgCenter, int y,
real centerR, real centerC, real theta,
int* sourceX, int* sourceY) { real scale,
real tgtCenter,
real imgCenter,
real centerR,
real centerC,
int* sourceX,
int* sourceY) {
real H[4] = {cosf(-theta), -sinf(-theta), sinf(-theta), cosf(-theta)}; real H[4] = {cosf(-theta), -sinf(-theta), sinf(-theta), cosf(-theta)};
// compute coornidates in the rotated and scaled image // compute coornidates in the rotated and scaled image
...@@ -57,11 +62,17 @@ __device__ void getTranformCoord(int x, int y, real theta, real scale, ...@@ -57,11 +62,17 @@ __device__ void getTranformCoord(int x, int y, real theta, real scale,
* created by Wei Xu (genome), converted by Jiang Wang * created by Wei Xu (genome), converted by Jiang Wang
*/ */
__global__ void kSamplingPatches(const real* imgs, real* targets, __global__ void kSamplingPatches(const real* imgs,
int imgSize, int tgtSize, const int channels, real* targets,
int samplingRate, const real* thetas, int imgSize,
const real* scales, const int* centerRs, int tgtSize,
const int* centerCs, const real padValue, const int channels,
int samplingRate,
const real* thetas,
const real* scales,
const int* centerRs,
const int* centerCs,
const real padValue,
const int numImages) { const int numImages) {
const int caseIdx = blockIdx.x * 4 + threadIdx.x; const int caseIdx = blockIdx.x * 4 + threadIdx.x;
const int pxIdx = blockIdx.y * 128 + threadIdx.y; const int pxIdx = blockIdx.y * 128 + threadIdx.y;
...@@ -80,8 +91,15 @@ __global__ void kSamplingPatches(const real* imgs, real* targets, ...@@ -80,8 +91,15 @@ __global__ void kSamplingPatches(const real* imgs, real* targets,
const int pxY = pxIdx / tgtSize; const int pxY = pxIdx / tgtSize;
int srcPxX, srcPxY; int srcPxX, srcPxY;
getTranformCoord(pxX, pxY, thetas[imgIdx], scales[imgIdx], tgtCenter, getTranformCoord(pxX,
imgCenter, centerCs[caseIdx], centerRs[caseIdx], &srcPxX, pxY,
thetas[imgIdx],
scales[imgIdx],
tgtCenter,
imgCenter,
centerCs[caseIdx],
centerRs[caseIdx],
&srcPxX,
&srcPxY); &srcPxY);
imgs += (imgIdx * imgPixels + srcPxY * imgSize + srcPxX) * channels; imgs += (imgIdx * imgPixels + srcPxY * imgSize + srcPxX) * channels;
...@@ -100,10 +118,15 @@ __global__ void kSamplingPatches(const real* imgs, real* targets, ...@@ -100,10 +118,15 @@ __global__ void kSamplingPatches(const real* imgs, real* targets,
* *
* created by Wei Xu * created by Wei Xu
*/ */
void hl_generate_disturb_params(real*& gpuAngle, real*& gpuScaleRatio, void hl_generate_disturb_params(real*& gpuAngle,
int*& gpuCenterR, int*& gpuCenterC, real*& gpuScaleRatio,
int numImages, int imgSize, real rotateAngle, int*& gpuCenterR,
real scaleRatio, int samplingRate, int*& gpuCenterC,
int numImages,
int imgSize,
real rotateAngle,
real scaleRatio,
int samplingRate,
bool isTrain) { bool isTrain) {
// The number of output samples. // The number of output samples.
int numPatches = numImages * samplingRate; int numPatches = numImages * samplingRate;
...@@ -123,7 +146,8 @@ void hl_generate_disturb_params(real*& gpuAngle, real*& gpuScaleRatio, ...@@ -123,7 +146,8 @@ void hl_generate_disturb_params(real*& gpuAngle, real*& gpuScaleRatio,
for (int i = 0; i < numImages; i++) { for (int i = 0; i < numImages; i++) {
r_angle[i] = r_angle[i] =
(rotateAngle * M_PI / 180.0) * (rand() / (RAND_MAX + 1.0) // NOLINT (rotateAngle * M_PI / 180.0) * (rand() / (RAND_MAX + 1.0) // NOLINT
- 0.5); -
0.5);
s_ratio[i] = s_ratio[i] =
1 + (rand() / (RAND_MAX + 1.0) - 0.5) * scaleRatio; // NOLINT 1 + (rand() / (RAND_MAX + 1.0) - 0.5) * scaleRatio; // NOLINT
} }
...@@ -140,8 +164,10 @@ void hl_generate_disturb_params(real*& gpuAngle, real*& gpuScaleRatio, ...@@ -140,8 +164,10 @@ void hl_generate_disturb_params(real*& gpuAngle, real*& gpuScaleRatio,
int pxY = int pxY =
(int)(real(imgSize - 1) * rand() / (RAND_MAX + 1.0)); // NOLINT (int)(real(imgSize - 1) * rand() / (RAND_MAX + 1.0)); // NOLINT
const real H[4] = {cos(-r_angle[i]), -sin(-r_angle[i]), const real H[4] = {cos(-r_angle[i]),
sin(-r_angle[i]), cos(-r_angle[i])}; -sin(-r_angle[i]),
sin(-r_angle[i]),
cos(-r_angle[i])};
real x = pxX - imgCenter; real x = pxX - imgCenter;
real y = pxY - imgCenter; real y = pxY - imgCenter;
real xx = H[0] * x + H[1] * y; real xx = H[0] * x + H[1] * y;
...@@ -185,9 +211,12 @@ void hl_generate_disturb_params(real*& gpuAngle, real*& gpuScaleRatio, ...@@ -185,9 +211,12 @@ void hl_generate_disturb_params(real*& gpuAngle, real*& gpuScaleRatio,
delete[] center_c; delete[] center_c;
} }
void hl_conv_random_disturb_with_params(const real* images, int imgSize, void hl_conv_random_disturb_with_params(const real* images,
int tgtSize, int channels, int imgSize,
int numImages, int samplingRate, int tgtSize,
int channels,
int numImages,
int samplingRate,
const real* gpuRotationAngle, const real* gpuRotationAngle,
const real* gpuScaleRatio, const real* gpuScaleRatio,
const int* gpuCenterR, const int* gpuCenterR,
...@@ -202,29 +231,59 @@ void hl_conv_random_disturb_with_params(const real* images, int imgSize, ...@@ -202,29 +231,59 @@ void hl_conv_random_disturb_with_params(const real* images, int imgSize,
dim3 threadsPerBlock(4, 128); dim3 threadsPerBlock(4, 128);
dim3 numBlocks(DIVUP(numPatches, 4), DIVUP(targetSize, 128)); dim3 numBlocks(DIVUP(numPatches, 4), DIVUP(targetSize, 128));
kSamplingPatches <<<numBlocks, threadsPerBlock>>> kSamplingPatches<<<numBlocks, threadsPerBlock>>>(images,
(images, target, imgSize, tgtSize, channels, samplingRate, target,
gpuRotationAngle, gpuScaleRatio, gpuCenterR, gpuCenterC, imgSize,
paddingValue, numImages); tgtSize,
channels,
samplingRate,
gpuRotationAngle,
gpuScaleRatio,
gpuCenterR,
gpuCenterC,
paddingValue,
numImages);
hl_device_synchronize(); hl_device_synchronize();
} }
void hl_conv_random_disturb(const real* images, int imgSize, void hl_conv_random_disturb(const real* images,
int tgtSize, int channels, int numImages, int imgSize,
real scaleRatio, real rotateAngle, int tgtSize,
int samplingRate, real* gpu_r_angle, int channels,
real* gpu_s_ratio, int* gpu_center_r, int numImages,
int* gpu_center_c, int paddingValue, real scaleRatio,
bool isTrain, real* targets) { real rotateAngle,
int samplingRate,
real* gpu_r_angle,
real* gpu_s_ratio,
int* gpu_center_r,
int* gpu_center_c,
int paddingValue,
bool isTrain,
real* targets) {
// generate the random disturbance sequence and the sampling locations // generate the random disturbance sequence and the sampling locations
hl_generate_disturb_params(gpu_r_angle, gpu_s_ratio, gpu_center_r, hl_generate_disturb_params(gpu_r_angle,
gpu_center_c, numImages, imgSize, rotateAngle, gpu_s_ratio,
scaleRatio, samplingRate, isTrain); gpu_center_r,
gpu_center_c,
hl_conv_random_disturb_with_params( numImages,
images, imgSize, tgtSize, channels, numImages, imgSize,
samplingRate, gpu_r_angle, gpu_s_ratio, rotateAngle,
gpu_center_r, gpu_center_r, paddingValue, scaleRatio,
samplingRate,
isTrain);
hl_conv_random_disturb_with_params(images,
imgSize,
tgtSize,
channels,
numImages,
samplingRate,
gpu_r_angle,
gpu_s_ratio,
gpu_center_r,
gpu_center_r,
paddingValue,
targets); targets);
} }
paddle/cuda/src/hl_table_apply.cu
浏览文件 @ 4bc1dc26
...@@ -12,15 +12,16 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,15 +12,16 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "hl_base.h" #include "hl_base.h"
#include "hl_device_functions.cuh"
#include "hl_cuda.h" #include "hl_cuda.h"
#include "hl_device_functions.cuh"
#include "paddle/utils/Logging.h" #include "paddle/utils/Logging.h"
template<int blockDimX, int blockDimY, int gridDimX, bool AddRow> template <int blockDimX, int blockDimY, int gridDimX, bool AddRow>
__global__ void KeMatrixAddRows(real* output, int ldo, __global__ void KeMatrixAddRows(real* output,
real* table, int ldt, int ldo,
real* table,
int ldt,
int* ids, int* ids,
int numSamples, int numSamples,
int tableSize, int tableSize,
...@@ -31,8 +32,8 @@ __global__ void KeMatrixAddRows(real* output, int ldo, ...@@ -31,8 +32,8 @@ __global__ void KeMatrixAddRows(real* output, int ldo,
while (idy < numSamples) { while (idy < numSamples) {
int tableId = ids[idy]; int tableId = ids[idy];
if ((0 <= tableId) && (tableId < tableSize)) { if ((0 <= tableId) && (tableId < tableSize)) {
real *out = output + idy * ldo; real* out = output + idy * ldo;
real *tab = table + tableId * ldt; real* tab = table + tableId * ldt;
for (int i = idx; i < dim; i += blockDimX) { for (int i = idx; i < dim; i += blockDimX) {
if (AddRow) { if (AddRow) {
paddle::paddleAtomicAdd(&tab[i], out[i]); paddle::paddleAtomicAdd(&tab[i], out[i]);
...@@ -45,8 +46,10 @@ __global__ void KeMatrixAddRows(real* output, int ldo, ...@@ -45,8 +46,10 @@ __global__ void KeMatrixAddRows(real* output, int ldo,
} }
} }
void hl_matrix_select_rows(real* output, int ldo, void hl_matrix_select_rows(real* output,
real* table, int ldt, int ldo,
real* table,
int ldt,
int* ids, int* ids,
int numSamples, int numSamples,
int tableSize, int tableSize,
...@@ -57,14 +60,16 @@ void hl_matrix_select_rows(real* output, int ldo, ...@@ -57,14 +60,16 @@ void hl_matrix_select_rows(real* output, int ldo,
dim3 threads(128, 8); dim3 threads(128, 8);
dim3 grid(8, 1); dim3 grid(8, 1);
KeMatrixAddRows<128, 8, 8, 0><<< grid, threads, 0, STREAM_DEFAULT >>> KeMatrixAddRows<128, 8, 8, 0><<<grid, threads, 0, STREAM_DEFAULT>>>(
(output, ldo, table, ldt, ids, numSamples, tableSize, dim); output, ldo, table, ldt, ids, numSamples, tableSize, dim);
CHECK_SYNC("hl_matrix_select_rows failed"); CHECK_SYNC("hl_matrix_select_rows failed");
} }
void hl_matrix_add_to_rows(real* table, int ldt, void hl_matrix_add_to_rows(real* table,
real* input, int ldi, int ldt,
real* input,
int ldi,
int* ids, int* ids,
int numSamples, int numSamples,
int tableSize, int tableSize,
...@@ -75,16 +80,15 @@ void hl_matrix_add_to_rows(real* table, int ldt, ...@@ -75,16 +80,15 @@ void hl_matrix_add_to_rows(real* table, int ldt,
dim3 threads(128, 8); dim3 threads(128, 8);
dim3 grid(8, 1); dim3 grid(8, 1);
KeMatrixAddRows<128, 8, 8, 1><<< grid, threads, 0, STREAM_DEFAULT >>> KeMatrixAddRows<128, 8, 8, 1><<<grid, threads, 0, STREAM_DEFAULT>>>(
(input, ldi, table, ldt, ids, numSamples, tableSize, dim); input, ldi, table, ldt, ids, numSamples, tableSize, dim);
CHECK_SYNC("hl_matrix_add_to_rows failed"); CHECK_SYNC("hl_matrix_add_to_rows failed");
} }
template<class T, int blockDimX, int gridDimX> template <class T, int blockDimX, int gridDimX>
__global__ void KeVectorSelect(T* dst, int sized, __global__ void KeVectorSelect(
const T* src, int sizes, T* dst, int sized, const T* src, int sizes, const int* ids, int sizei) {
const int* ids, int sizei) {
int idx = threadIdx.x + blockDimX * blockIdx.x; int idx = threadIdx.x + blockDimX * blockIdx.x;
while (idx < sizei) { while (idx < sizei) {
int index = ids[idx]; int index = ids[idx];
...@@ -95,9 +99,8 @@ __global__ void KeVectorSelect(T* dst, int sized, ...@@ -95,9 +99,8 @@ __global__ void KeVectorSelect(T* dst, int sized,
} }
template <class T> template <class T>
void hl_vector_select_from(T* dst, int sized, void hl_vector_select_from(
const T* src, int sizes, T* dst, int sized, const T* src, int sizes, const int* ids, int sizei) {
const int* ids, int sizei) {
CHECK_NOTNULL(dst); CHECK_NOTNULL(dst);
CHECK_NOTNULL(src); CHECK_NOTNULL(src);
CHECK_NOTNULL(ids); CHECK_NOTNULL(ids);
...@@ -105,18 +108,17 @@ void hl_vector_select_from(T* dst, int sized, ...@@ -105,18 +108,17 @@ void hl_vector_select_from(T* dst, int sized,
dim3 threads(512, 1); dim3 threads(512, 1);
dim3 grid(8, 1); dim3 grid(8, 1);
KeVectorSelect<T, 512, 8><<< grid, threads, 0, STREAM_DEFAULT >>> KeVectorSelect<T, 512, 8><<<grid, threads, 0, STREAM_DEFAULT>>>(
(dst, sized, src, sizes, ids, sizei); dst, sized, src, sizes, ids, sizei);
CHECK_SYNC("hl_vector_select_from failed"); CHECK_SYNC("hl_vector_select_from failed");
} }
template template void hl_vector_select_from(real* dst,
void hl_vector_select_from(real* dst, int sized, int sized,
const real* src, int sizes, const real* src,
const int* ids, int sizei); int sizes,
template const int* ids,
void hl_vector_select_from(int* dst, int sized, int sizei);
const int* src, int sizes, template void hl_vector_select_from(
const int* ids, int sizei); int* dst, int sized, const int* src, int sizes, const int* ids, int sizei);
paddle/cuda/src/hl_top_k.cu
浏览文件 @ 4bc1dc26
...@@ -12,45 +12,37 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,45 +12,37 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "hl_base.h" #include "hl_base.h"
#include "hl_top_k.h"
#include "hl_sparse.ph" #include "hl_sparse.ph"
#include "hl_top_k.h"
#include "paddle/utils/Logging.h" #include "paddle/utils/Logging.h"
// using namespace hppl; // using namespace hppl;
struct Pair { struct Pair {
__device__ __forceinline__ __device__ __forceinline__ Pair() {}
Pair() {}
__device__ __forceinline__ __device__ __forceinline__ Pair(real value, int id) : v_(value), id_(id) {}
Pair(real value, int id) : v_(value), id_(id) {}
__device__ __forceinline__ __device__ __forceinline__ void set(real value, int id) {
void set(real value, int id) {
v_ = value; v_ = value;
id_ = id; id_ = id;
} }
__device__ __forceinline__ __device__ __forceinline__ void operator=(const Pair& in) {
void operator=(const Pair& in) {
v_ = in.v_; v_ = in.v_;
id_ = in.id_; id_ = in.id_;
} }
__device__ __forceinline__ __device__ __forceinline__ bool operator<(const real value) const {
bool operator<(const real value) const {
return (v_ < value); return (v_ < value);
} }
__device__ __forceinline__ __device__ __forceinline__ bool operator<(const Pair& in) const {
bool operator<(const Pair& in) const {
return (v_ < in.v_) || ((v_ == in.v_) && (id_ > in.id_)); return (v_ < in.v_) || ((v_ == in.v_) && (id_ > in.id_));
} }
__device__ __forceinline__ __device__ __forceinline__ bool operator>(const Pair& in) const {
bool operator>(const Pair& in) const {
return (v_ > in.v_) || ((v_ == in.v_) && (id_ < in.id_)); return (v_ > in.v_) || ((v_ == in.v_) && (id_ < in.id_));
} }
...@@ -58,8 +50,9 @@ struct Pair { ...@@ -58,8 +50,9 @@ struct Pair {
int id_; int id_;
}; };
__device__ __forceinline__ __device__ __forceinline__ void addTo(Pair topK[],
void addTo(Pair topK[], const Pair &p, int beamSize) { const Pair& p,
int beamSize) {
for (int k = beamSize - 2; k >= 0; k--) { for (int k = beamSize - 2; k >= 0; k--) {
if (topK[k] < p) { if (topK[k] < p) {
topK[k + 1] = topK[k]; topK[k + 1] = topK[k];
...@@ -71,9 +64,8 @@ void addTo(Pair topK[], const Pair &p, int beamSize) { ...@@ -71,9 +64,8 @@ void addTo(Pair topK[], const Pair &p, int beamSize) {
topK[0] = p; topK[0] = p;
} }
template<int beamSize> template <int beamSize>
__device__ __forceinline__ __device__ __forceinline__ void addTo(Pair topK[], const Pair& p) {
void addTo(Pair topK[], const Pair &p) {
for (int k = beamSize - 2; k >= 0; k--) { for (int k = beamSize - 2; k >= 0; k--) {
if (topK[k] < p) { if (topK[k] < p) {
topK[k + 1] = topK[k]; topK[k + 1] = topK[k];
...@@ -85,9 +77,9 @@ void addTo(Pair topK[], const Pair &p) { ...@@ -85,9 +77,9 @@ void addTo(Pair topK[], const Pair &p) {
topK[0] = p; topK[0] = p;
} }
template<int blockSize> template <int blockSize>
__device__ __forceinline__ __device__ __forceinline__ void getTopK(
void getTopK(Pair topK[], real *src, int idx, int dim, int beamSize) { Pair topK[], real* src, int idx, int dim, int beamSize) {
while (idx < dim) { while (idx < dim) {
if (topK[beamSize - 1] < src[idx]) { if (topK[beamSize - 1] < src[idx]) {
Pair tmp(src[idx], idx); Pair tmp(src[idx], idx);
...@@ -97,10 +89,9 @@ void getTopK(Pair topK[], real *src, int idx, int dim, int beamSize) { ...@@ -97,10 +89,9 @@ void getTopK(Pair topK[], real *src, int idx, int dim, int beamSize) {
} }
} }
template<int blockSize> template <int blockSize>
__device__ __forceinline__ __device__ __forceinline__ void getTopK(
void getTopK(Pair topK[], real *src, int idx, int dim, Pair topK[], real* src, int idx, int dim, const Pair& max, int beamSize) {
const Pair& max, int beamSize) {
while (idx < dim) { while (idx < dim) {
if (topK[beamSize - 1] < src[idx]) { if (topK[beamSize - 1] < src[idx]) {
Pair tmp(src[idx], idx); Pair tmp(src[idx], idx);
...@@ -112,10 +103,9 @@ void getTopK(Pair topK[], real *src, int idx, int dim, ...@@ -112,10 +103,9 @@ void getTopK(Pair topK[], real *src, int idx, int dim,
} }
} }
template<int blockSize> template <int blockSize>
__device__ __forceinline__ __device__ __forceinline__ void getTopK(
void getTopK(Pair topK[], real *val, int *col, Pair topK[], real* val, int* col, int idx, int dim, int beamSize) {
int idx, int dim, int beamSize) {
while (idx < dim) { while (idx < dim) {
if (topK[beamSize - 1] < val[idx]) { if (topK[beamSize - 1] < val[idx]) {
Pair tmp(val[idx], col[idx]); Pair tmp(val[idx], col[idx]);
...@@ -125,10 +115,14 @@ void getTopK(Pair topK[], real *val, int *col, ...@@ -125,10 +115,14 @@ void getTopK(Pair topK[], real *val, int *col,
} }
} }
template<int blockSize> template <int blockSize>
__device__ __forceinline__ __device__ __forceinline__ void getTopK(Pair topK[],
void getTopK(Pair topK[], real *val, int *col, int idx, int dim, real* val,
const Pair& max, int beamSize) { int* col,
int idx,
int dim,
const Pair& max,
int beamSize) {
while (idx < dim) { while (idx < dim) {
if (topK[beamSize - 1] < val[idx]) { if (topK[beamSize - 1] < val[idx]) {
Pair tmp(val[idx], col[idx]); Pair tmp(val[idx], col[idx]);
...@@ -140,12 +134,16 @@ void getTopK(Pair topK[], real *val, int *col, int idx, int dim, ...@@ -140,12 +134,16 @@ void getTopK(Pair topK[], real *val, int *col, int idx, int dim,
} }
} }
template<int maxLength, int blockSize> template <int maxLength, int blockSize>
__device__ __forceinline__ __device__ __forceinline__ void threadGetTopK(Pair topK[],
void threadGetTopK(Pair topK[], int& beam, int beamSize, int& beam,
int beamSize,
real* src, real* src,
bool& firstStep, bool& isEmpty, Pair& max, bool& firstStep,
int dim, const int tid) { bool& isEmpty,
Pair& max,
int dim,
const int tid) {
if (beam > 0) { if (beam > 0) {
int length = beam < beamSize ? beam : beamSize; int length = beam < beamSize ? beam : beamSize;
if (firstStep) { if (firstStep) {
...@@ -160,8 +158,7 @@ void threadGetTopK(Pair topK[], int& beam, int beamSize, ...@@ -160,8 +158,7 @@ void threadGetTopK(Pair topK[], int& beam, int beamSize,
} }
} }
if (!isEmpty) { if (!isEmpty) {
getTopK<blockSize>(topK + maxLength - beam, src, tid, dim, getTopK<blockSize>(topK + maxLength - beam, src, tid, dim, max, length);
max, length);
} }
} }
...@@ -171,12 +168,17 @@ void threadGetTopK(Pair topK[], int& beam, int beamSize, ...@@ -171,12 +168,17 @@ void threadGetTopK(Pair topK[], int& beam, int beamSize,
} }
} }
template<int maxLength, int blockSize> template <int maxLength, int blockSize>
__device__ __forceinline__ __device__ __forceinline__ void threadGetTopK(Pair topK[],
void threadGetTopK(Pair topK[], int& beam, int beamSize, int& beam,
real* val, int* col, int beamSize,
bool& firstStep, bool& isEmpty, Pair& max, real* val,
int dim, const int tid) { int* col,
bool& firstStep,
bool& isEmpty,
Pair& max,
int dim,
const int tid) {
if (beam > 0) { if (beam > 0) {
int length = beam < beamSize ? beam : beamSize; int length = beam < beamSize ? beam : beamSize;
if (firstStep) { if (firstStep) {
...@@ -191,8 +193,8 @@ void threadGetTopK(Pair topK[], int& beam, int beamSize, ...@@ -191,8 +193,8 @@ void threadGetTopK(Pair topK[], int& beam, int beamSize,
} }
} }
if (!isEmpty) { if (!isEmpty) {
getTopK<blockSize>(topK + maxLength - beam, val, col, tid, dim, getTopK<blockSize>(
max, length); topK + maxLength - beam, val, col, tid, dim, max, length);
} }
} }
...@@ -202,12 +204,16 @@ void threadGetTopK(Pair topK[], int& beam, int beamSize, ...@@ -202,12 +204,16 @@ void threadGetTopK(Pair topK[], int& beam, int beamSize,
} }
} }
template<int maxLength, int blockSize> template <int maxLength, int blockSize>
__device__ __forceinline__ __device__ __forceinline__ void blockReduce(Pair* shTopK,
void blockReduce(Pair* shTopK, int* maxId, Pair topK[], int* maxId,
real** topVal, int** topIds, Pair topK[],
int& beam, int& beamSize, real** topVal,
const int tid, const int warp) { int** topIds,
int& beam,
int& beamSize,
const int tid,
const int warp) {
while (true) { while (true) {
__syncthreads(); __syncthreads();
if (tid < blockSize / 2) { if (tid < blockSize / 2) {
...@@ -218,7 +224,7 @@ void blockReduce(Pair* shTopK, int* maxId, Pair topK[], ...@@ -218,7 +224,7 @@ void blockReduce(Pair* shTopK, int* maxId, Pair topK[],
} }
} }
__syncthreads(); __syncthreads();
for (int stride = blockSize / 4; stride > 0; stride = stride/2) { for (int stride = blockSize / 4; stride > 0; stride = stride / 2) {
if (tid < stride) { if (tid < stride) {
if (shTopK[maxId[tid]] < shTopK[maxId[tid + stride]]) { if (shTopK[maxId[tid]] < shTopK[maxId[tid + stride]]) {
maxId[tid] = maxId[tid + stride]; maxId[tid] = maxId[tid + stride];
...@@ -257,10 +263,12 @@ void blockReduce(Pair* shTopK, int* maxId, Pair topK[], ...@@ -257,10 +263,12 @@ void blockReduce(Pair* shTopK, int* maxId, Pair topK[],
* 3. go to the second setp, until one thread's topK value is null; * 3. go to the second setp, until one thread's topK value is null;
* 4. go to the first setp, until get the topK value. * 4. go to the first setp, until get the topK value.
*/ */
template<int maxLength, int blockSize> template <int maxLength, int blockSize>
__global__ void KeMatrixTopK(real* topVal, int ldv, __global__ void KeMatrixTopK(real* topVal,
int * topIds, int ldv,
real* src, int lds, int* topIds,
real* src,
int lds,
int dim, int dim,
int beamSize) { int beamSize) {
__shared__ Pair shTopK[blockSize]; __shared__ Pair shTopK[blockSize];
...@@ -281,18 +289,19 @@ __global__ void KeMatrixTopK(real* topVal, int ldv, ...@@ -281,18 +289,19 @@ __global__ void KeMatrixTopK(real* topVal, int ldv,
topK[k].set(-HL_FLOAT_MAX, -1); topK[k].set(-HL_FLOAT_MAX, -1);
} }
while (beamSize) { while (beamSize) {
threadGetTopK<maxLength, blockSize> threadGetTopK<maxLength, blockSize>(
(topK, beam, beamSize, src, firstStep, isEmpty, max, dim, tid); topK, beam, beamSize, src, firstStep, isEmpty, max, dim, tid);
shTopK[tid] = topK[0]; shTopK[tid] = topK[0];
blockReduce<maxLength, blockSize> blockReduce<maxLength, blockSize>(
(shTopK, maxId, topK, &topVal, &topIds, beam, beamSize, tid, warp); shTopK, maxId, topK, &topVal, &topIds, beam, beamSize, tid, warp);
} }
} }
template<int maxLength, int blockSize> template <int maxLength, int blockSize>
__global__ void KeSMatrixTopK(real* topVal, int ldv, __global__ void KeSMatrixTopK(real* topVal,
int * topIds, int ldv,
int* topIds,
real* val, real* val,
int* row, int* row,
int* col, int* col,
...@@ -330,18 +339,20 @@ __global__ void KeSMatrixTopK(real* topVal, int ldv, ...@@ -330,18 +339,20 @@ __global__ void KeSMatrixTopK(real* topVal, int ldv,
topK[k].set(-HL_FLOAT_MAX, -1); topK[k].set(-HL_FLOAT_MAX, -1);
} }
while (beamSize) { while (beamSize) {
threadGetTopK<maxLength, blockSize> threadGetTopK<maxLength, blockSize>(
(topK, beam, beamSize, val, col, firstStep, isEmpty, max, dim, tid); topK, beam, beamSize, val, col, firstStep, isEmpty, max, dim, tid);
shTopK[tid] = topK[0]; shTopK[tid] = topK[0];
blockReduce<maxLength, blockSize> blockReduce<maxLength, blockSize>(
(shTopK, maxId, topK, &topVal, &topIds, beam, beamSize, tid, warp); shTopK, maxId, topK, &topVal, &topIds, beam, beamSize, tid, warp);
} }
} }
void hl_matrix_top_k(real* topVal, int ldv, void hl_matrix_top_k(real* topVal,
int * topIds, int ldv,
real* src, int lds, int* topIds,
real* src,
int lds,
int dim, int dim,
int beamSize, int beamSize,
int numSamples) { int numSamples) {
...@@ -353,33 +364,32 @@ void hl_matrix_top_k(real* topVal, int ldv, ...@@ -353,33 +364,32 @@ void hl_matrix_top_k(real* topVal, int ldv,
dim3 threads(256, 1); dim3 threads(256, 1);
dim3 grid(numSamples, 1); dim3 grid(numSamples, 1);
KeMatrixTopK<5, 256><<< grid, threads, 0, STREAM_DEFAULT >>> KeMatrixTopK<5, 256><<<grid, threads, 0, STREAM_DEFAULT>>>(
(topVal, ldv, topIds, src, lds, dim, beamSize); topVal, ldv, topIds, src, lds, dim, beamSize);
CHECK_SYNC("hl_matrix_top_k failed"); CHECK_SYNC("hl_matrix_top_k failed");
} }
void hl_sparse_matrix_top_k(real* topVal, int ldv, void hl_sparse_matrix_top_k(real* topVal,
int * topIds, int ldv,
int* topIds,
hl_sparse_matrix_s src, hl_sparse_matrix_s src,
int beamSize, int beamSize,
int numSamples) { int numSamples) {
CHECK_NOTNULL(topVal); CHECK_NOTNULL(topVal);
CHECK_NOTNULL(topIds); CHECK_NOTNULL(topIds);
CHECK_NOTNULL(src); CHECK_NOTNULL(src);
CHECK_EQ(src->format, HL_SPARSE_CSR) CHECK_EQ(src->format, HL_SPARSE_CSR) << "sparse matrix format error!";
<<"sparse matrix format error!";
hl_csr_matrix csr = (hl_csr_matrix)src->matrix; hl_csr_matrix csr = (hl_csr_matrix)src->matrix;
if (csr->csr_val == NULL || csr->csr_row == NULL || if (csr->csr_val == NULL || csr->csr_row == NULL || csr->csr_col == NULL) {
csr->csr_col == NULL) {
LOG(FATAL) << "parameter src is null!"; LOG(FATAL) << "parameter src is null!";
} }
dim3 threads(256, 1); dim3 threads(256, 1);
dim3 grid(numSamples, 1); dim3 grid(numSamples, 1);
KeSMatrixTopK<5, 256><<< grid, threads, 0, STREAM_DEFAULT >>> KeSMatrixTopK<5, 256><<<grid, threads, 0, STREAM_DEFAULT>>>(
(topVal, ldv, topIds, csr->csr_val, csr->csr_row, csr->csr_col, beamSize); topVal, ldv, topIds, csr->csr_val, csr->csr_row, csr->csr_col, beamSize);
CHECK_SYNC("hl_sparse_matrix_top_k failed"); CHECK_SYNC("hl_sparse_matrix_top_k failed");
} }
...@@ -392,10 +402,12 @@ void hl_sparse_matrix_top_k(real* topVal, int ldv, ...@@ -392,10 +402,12 @@ void hl_sparse_matrix_top_k(real* topVal, int ldv,
* 3. go to the second setp, until one thread's topK value is null; * 3. go to the second setp, until one thread's topK value is null;
* 4. go to the first setp, until get the topK value. * 4. go to the first setp, until get the topK value.
*/ */
template<int maxLength, int blockSize> template <int maxLength, int blockSize>
__global__ void KeMatrixTopKClassificationError(real* topVal, int ldv, __global__ void KeMatrixTopKClassificationError(real* topVal,
int * topIds, int ldv,
real* src, int lds, int* topIds,
real* src,
int lds,
int dim, int dim,
int beamSize, int beamSize,
int* label, int* label,
...@@ -420,12 +432,12 @@ __global__ void KeMatrixTopKClassificationError(real* topVal, int ldv, ...@@ -420,12 +432,12 @@ __global__ void KeMatrixTopKClassificationError(real* topVal, int ldv,
} }
while (beamSize) { while (beamSize) {
threadGetTopK<maxLength, blockSize> threadGetTopK<maxLength, blockSize>(
(topK, beam, beamSize, src, firstStep, isEmpty, max, dim, tid); topK, beam, beamSize, src, firstStep, isEmpty, max, dim, tid);
shTopK[tid] = topK[0]; shTopK[tid] = topK[0];
blockReduce<maxLength, blockSize> blockReduce<maxLength, blockSize>(
(shTopK, maxId, topK, &topVal, &topIds, beam, beamSize, tid, warp); shTopK, maxId, topK, &topVal, &topIds, beam, beamSize, tid, warp);
} }
__syncthreads(); __syncthreads();
...@@ -440,9 +452,11 @@ __global__ void KeMatrixTopKClassificationError(real* topVal, int ldv, ...@@ -440,9 +452,11 @@ __global__ void KeMatrixTopKClassificationError(real* topVal, int ldv,
} }
} }
void hl_matrix_classification_error(real* topVal, int ldv, void hl_matrix_classification_error(real* topVal,
int ldv,
int* topIds, int* topIds,
real* src, int lds, real* src,
int lds,
int dim, int dim,
int topkSize, int topkSize,
int numSamples, int numSamples,
...@@ -456,9 +470,8 @@ void hl_matrix_classification_error(real* topVal, int ldv, ...@@ -456,9 +470,8 @@ void hl_matrix_classification_error(real* topVal, int ldv,
dim3 threads(256, 1); dim3 threads(256, 1);
dim3 grid(numSamples, 1); dim3 grid(numSamples, 1);
KeMatrixTopKClassificationError<5, 256> KeMatrixTopKClassificationError<5, 256><<<grid, threads, 0, STREAM_DEFAULT>>>(
<<< grid, threads, 0, STREAM_DEFAULT >>> topVal, ldv, topIds, src, lds, dim, topkSize, label, recResult);
(topVal, ldv, topIds, src, lds, dim, topkSize, label, recResult);
CHECK_SYNC("hl_matrix_top_k classification error failed"); CHECK_SYNC("hl_matrix_top_k classification error failed");
} }
paddle/framework/CMakeLists.txt
浏览文件 @ 4bc1dc26
...@@ -12,13 +12,15 @@ cc_test(variable_test SRCS variable_test.cc) ...@@ -12,13 +12,15 @@ cc_test(variable_test SRCS variable_test.cc)
cc_library(scope SRCS scope.cc) cc_library(scope SRCS scope.cc)
cc_test(scope_test SRCS scope_test.cc DEPS scope) cc_test(scope_test SRCS scope_test.cc DEPS scope)
proto_library(attr_type SRCS attr_type.proto) proto_library(attribute_proto SRCS attribute.proto)
proto_library(op_proto SRCS op_proto.proto DEPS attr_type) proto_library(op_proto SRCS op_proto.proto DEPS attribute_proto)
proto_library(op_desc SRCS op_desc.proto DEPS attr_type) proto_library(op_desc SRCS op_desc.proto DEPS attribute_proto)
cc_test(op_proto_test SRCS op_proto_test.cc DEPS op_proto protobuf) cc_test(op_proto_test SRCS op_proto_test.cc DEPS op_proto protobuf)
cc_test(op_desc_test SRCS op_desc_test.cc DEPS op_desc protobuf) cc_test(op_desc_test SRCS op_desc_test.cc DEPS op_desc protobuf)
cc_library(operator SRCS operator.cc DEPS op_desc device_context tensor scope) cc_library(attribute SRCS attribute.cc DEPS op_desc op_proto)
cc_library(operator SRCS operator.cc DEPS op_desc device_context tensor scope attribute)
cc_test(operator_test SRCS operator_test.cc DEPS operator op_registry) cc_test(operator_test SRCS operator_test.cc DEPS operator op_registry)
cc_library(grad_op_builder SRCS grad_op_builder.cc DEPS op_proto operator) cc_library(grad_op_builder SRCS grad_op_builder.cc DEPS op_proto operator)
...@@ -26,13 +28,19 @@ cc_library(op_registry SRCS op_registry.cc DEPS op_desc grad_op_builder) ...@@ -26,13 +28,19 @@ cc_library(op_registry SRCS op_registry.cc DEPS op_desc grad_op_builder)
cc_test(op_registry_test SRCS op_registry_test.cc DEPS op_registry) cc_test(op_registry_test SRCS op_registry_test.cc DEPS op_registry)
cc_test(grad_op_builder_test SRCS grad_op_builder_test.cc DEPS grad_op_builder op_registry add_op) cc_test(grad_op_builder_test SRCS grad_op_builder_test.cc DEPS grad_op_builder op_registry add_op)
py_proto_compile(framework_py_proto SRCS attr_type.proto op_proto.proto op_desc.proto) py_proto_compile(framework_py_proto SRCS attribute.proto op_proto.proto op_desc.proto)
# Generate an empty __init__.py to make framework_py_proto as a valid python module. # Generate an empty __init__.py to make framework_py_proto as a valid python module.
add_custom_target(framework_py_proto_init ALL COMMAND ${CMAKE_COMMAND} -E touch __init__.py) add_custom_target(framework_py_proto_init ALL COMMAND ${CMAKE_COMMAND} -E touch __init__.py)
add_dependencies(framework_py_proto framework_py_proto_init) add_dependencies(framework_py_proto framework_py_proto_init)
cc_library(net SRCS net.cc DEPS op_registry) cc_library(backward SRCS backward.cc DEPS net_op)
cc_test(net_op_test SRCS net_op_test.cc DEPS net)
cc_library(backward SRCS backward.cc DEPS net)
cc_test(backward_test SRCS backward_test.cc DEPS backward) cc_test(backward_test SRCS backward_test.cc DEPS backward)
cc_library(paddle_pybind SHARED
SRCS pybind.cc
DEPS pybind python backward
fc_op
sgd_op
add_op
mean_op
cross_entropy_op
recurrent_op)
paddle/framework/attribute.cc 0 → 100644
浏览文件 @ 4bc1dc26
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
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 "paddle/framework/attribute.h"
#include <vector>
namespace paddle {
namespace framework {
template <>
AttrType AttrTypeID<int>() {
return INT;
}
template <>
AttrType AttrTypeID<float>() {
return FLOAT;
}
template <>
AttrType AttrTypeID<std::string>() {
return STRING;
}
template <>
AttrType AttrTypeID<std::vector<int>>() {
return INTS;
}
template <>
AttrType AttrTypeID<std::vector<float>>() {
return FLOATS;
}
template <>
AttrType AttrTypeID<std::vector<std::string>>() {
return STRINGS;
}
Attribute GetAttrValue(const AttrDesc& attr_desc) {
switch (attr_desc.type()) {
case paddle::framework::AttrType::INT: {
return attr_desc.i();
}
case paddle::framework::AttrType::FLOAT: {
return attr_desc.f();
}
case paddle::framework::AttrType::STRING: {
return attr_desc.s();
}
case paddle::framework::AttrType::INTS: {
std::vector<int> val(attr_desc.ints_size());
for (int i = 0; i < attr_desc.ints_size(); ++i) {
val[i] = attr_desc.ints(i);
}
return val;
}
case paddle::framework::AttrType::FLOATS: {
std::vector<float> val(attr_desc.floats_size());
for (int i = 0; i < attr_desc.floats_size(); ++i) {
val[i] = attr_desc.floats(i);
}
return val;
}
case paddle::framework::AttrType::STRINGS: {
std::vector<std::string> val(attr_desc.strings_size());
for (int i = 0; i < attr_desc.strings_size(); ++i) {
val[i] = attr_desc.strings(i);
}
return val;
}
}
PADDLE_ENFORCE(false, "Unknown OpDesc::AttrDesc::type !");
return boost::blank();
}
} // namespace framework
} // namespace paddle
paddle/framework/attr_checker.h → paddle/framework/attribute.h
浏览文件 @ 4bc1dc26
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
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. */
#pragma once #pragma once
#include <boost/variant.hpp> #include <boost/variant.hpp>
...@@ -6,6 +20,9 @@ ...@@ -6,6 +20,9 @@
#include <unordered_map> #include <unordered_map>
#include <unordered_set> #include <unordered_set>
#include <vector> #include <vector>
#include "paddle/framework/attribute.pb.h"
#include "paddle/framework/op_desc.pb.h"
#include "paddle/platform/enforce.h" #include "paddle/platform/enforce.h"
namespace paddle { namespace paddle {
...@@ -14,13 +31,19 @@ namespace framework { ...@@ -14,13 +31,19 @@ namespace framework {
typedef boost::variant<boost::blank, int, float, std::string, std::vector<int>, typedef boost::variant<boost::blank, int, float, std::string, std::vector<int>,
std::vector<float>, std::vector<std::string>> std::vector<float>, std::vector<std::string>>
Attribute; Attribute;
typedef std::unordered_map<std::string, Attribute> AttributeMap; typedef std::unordered_map<std::string, Attribute> AttributeMap;
template <typename T>
AttrType AttrTypeID();
Attribute GetAttrValue(const AttrDesc& attr_desc);
// check whether a value(attribute) fit a certain limit // check whether a value(attribute) fit a certain limit
template <typename T> template <typename T>
class LargerThanChecker { class LargerThanChecker {
public: public:
LargerThanChecker(T lower_bound) : lower_bound_(lower_bound) {} explicit LargerThanChecker(T lower_bound) : lower_bound_(lower_bound) {}
void operator()(T& value) const { void operator()(T& value) const {
PADDLE_ENFORCE(value > lower_bound_, "larger_than check fail"); PADDLE_ENFORCE(value > lower_bound_, "larger_than check fail");
} }
...@@ -35,7 +58,8 @@ class LargerThanChecker { ...@@ -35,7 +58,8 @@ class LargerThanChecker {
template <typename T> template <typename T>
class DefaultValueSetter { class DefaultValueSetter {
public: public:
DefaultValueSetter(T default_value) : default_value_(default_value) {} explicit DefaultValueSetter(T default_value)
: default_value_(default_value) {}
void operator()(T& value) const { value = default_value_; } void operator()(T& value) const { value = default_value_; }
private: private:
...@@ -78,7 +102,8 @@ class TypedAttrChecker { ...@@ -78,7 +102,8 @@ class TypedAttrChecker {
typedef std::function<void(T&)> ValueChecker; typedef std::function<void(T&)> ValueChecker;
public: public:
TypedAttrChecker(const std::string& attr_name) : attr_name_(attr_name) {} explicit TypedAttrChecker(const std::string& attr_name)
: attr_name_(attr_name) {}
TypedAttrChecker& InEnum(const std::unordered_set<T>& range) { TypedAttrChecker& InEnum(const std::unordered_set<T>& range) {
value_checkers_.push_back(EnumInContainer<T>(range)); value_checkers_.push_back(EnumInContainer<T>(range));
......
paddle/framework/attr_type.proto → paddle/framework/attribute.proto
浏览文件 @ 4bc1dc26
...@@ -12,7 +12,7 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,7 +12,7 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
syntax="proto2"; syntax = "proto2";
package paddle.framework; package paddle.framework;
// Attribute Type for paddle's Op. // Attribute Type for paddle's Op.
......
paddle/framework/backward.cc
浏览文件 @ 4bc1dc26
...@@ -14,8 +14,8 @@ ...@@ -14,8 +14,8 @@
#include "paddle/framework/backward.h" #include "paddle/framework/backward.h"
#include <list> #include <list>
#include "paddle/framework/net.h"
#include "paddle/framework/op_registry.h" #include "paddle/framework/op_registry.h"
#include "paddle/operators/net_op.h"
namespace paddle { namespace paddle {
namespace framework { namespace framework {
...@@ -32,7 +32,7 @@ static bool AllInSet(const std::vector<std::string>& names, ...@@ -32,7 +32,7 @@ static bool AllInSet(const std::vector<std::string>& names,
} }
static std::shared_ptr<OperatorBase> NOP() { static std::shared_ptr<OperatorBase> NOP() {
auto net_op = std::make_shared<NetOp>(); auto net_op = std::make_shared<operators::NetOp>();
net_op->type_ = "@NOP@"; net_op->type_ = "@NOP@";
net_op->CompleteAddOp(); net_op->CompleteAddOp();
return net_op; return net_op;
...@@ -42,9 +42,9 @@ static std::shared_ptr<OperatorBase> NOP() { ...@@ -42,9 +42,9 @@ static std::shared_ptr<OperatorBase> NOP() {
// //
// no_grad_names the gradient variable names without gradient calculating. // no_grad_names the gradient variable names without gradient calculating.
// //
// uniq_id is a unique index used inside recursively calling BackwardRecursive. // uniq_id is a unique index used inside recursively calling
// use `uid = uniq_id++;` to get the unique index, and pass `uniq_id` through // BackwardRecursive. use `uid = uniq_id++;` to get the unique index, and
// recursive calling. // pass `uniq_id` through recursive calling.
// //
// returns The backward operator. For simple situation, it is a simple // returns The backward operator. For simple situation, it is a simple
// operator. For complex situation, it is a NetOp. // operator. For complex situation, it is a NetOp.
...@@ -59,32 +59,30 @@ std::shared_ptr<OperatorBase> BackwardRecursive( ...@@ -59,32 +59,30 @@ std::shared_ptr<OperatorBase> BackwardRecursive(
// If all input gradients of forwarding operator do not need to calculate, // If all input gradients of forwarding operator do not need to calculate,
// just return an NOP. Not return null ptr because NOP does not take // just return an NOP. Not return null ptr because NOP does not take
// too much time for calculation, but it is useful for simplifying logic. // too much time for calculation, but it is useful for simplifying logic.
if (AllInSet(forwardOp.inputs_, OperatorBase::GRAD_VAR_SUFFIX(), if (AllInSet(forwardOp.inputs_, kGradVarSuffix, no_grad_names)) {
no_grad_names)) {
return NOP(); return NOP();
} }
// All output gradients of forwarding operator do not need to calculate. Then // All output gradients of forwarding operator do not need to calculate.
// all input gradients cannot be computed at all, and we put them into // Then all input gradients cannot be computed at all, and we put them into
// `no_grad_names` set. Return an NOP. // `no_grad_names` set. Return an NOP.
if (AllInSet(forwardOp.outputs_, OperatorBase::GRAD_VAR_SUFFIX(), if (AllInSet(forwardOp.outputs_, kGradVarSuffix, no_grad_names)) {
no_grad_names)) {
for (auto& name : forwardOp.inputs_) { for (auto& name : forwardOp.inputs_) {
// Mark all input is not need // Mark all input is not need
no_grad_names.insert(name + OperatorBase::GRAD_VAR_SUFFIX()); no_grad_names.insert(name + kGradVarSuffix);
} }
return NOP(); return NOP();
} }
// Returned gradient network // Returned gradient network
auto net = std::make_shared<NetOp>(); auto net = std::make_shared<operators::NetOp>();
if (forwardOp.IsNetOp()) { if (forwardOp.IsNetOp()) {
// Because forwardOp is a net op, it can static_cast. // Because forwardOp is a net op, it can static_cast.
auto& forwardNet = static_cast<const NetOp&>(forwardOp); auto& forwardNet = static_cast<const operators::NetOp&>(forwardOp);
// Map from output gradient variable name to operator's indices in backward // Map from output gradient variable name to operator's indices in
// net. That operator generates that variable. // backward net. That operator generates that variable.
std::unordered_map<std::string, std::vector<size_t>> dup_output_ops; std::unordered_map<std::string, std::vector<size_t>> dup_output_ops;
size_t local_op_id = 0; size_t local_op_id = 0;
...@@ -134,9 +132,9 @@ std::shared_ptr<OperatorBase> BackwardRecursive( ...@@ -134,9 +132,9 @@ std::shared_ptr<OperatorBase> BackwardRecursive(
std::shared_ptr<OperatorBase> grad_op = OpRegistry::CreateGradOp(forwardOp); std::shared_ptr<OperatorBase> grad_op = OpRegistry::CreateGradOp(forwardOp);
for (std::string& grad_input : grad_op->inputs_) { for (std::string& grad_input : grad_op->inputs_) {
if (no_grad_names.count(grad_input)) { if (no_grad_names.count(grad_input)) {
std::string prefix = grad_input.substr( std::string prefix =
0, grad_input.size() - OperatorBase::GRAD_VAR_SUFFIX().size()); grad_input.substr(0, grad_input.size() - kGradVarSuffix.size());
grad_input = prefix + OperatorBase::ZERO_VAR_SUFFIX(); grad_input = prefix + kZeroVarSuffix;
// If part of input gradient of that operator is not calculated, fill // If part of input gradient of that operator is not calculated, fill
// zero variables to that input gradient. // zero variables to that input gradient.
...@@ -147,7 +145,7 @@ std::shared_ptr<OperatorBase> BackwardRecursive( ...@@ -147,7 +145,7 @@ std::shared_ptr<OperatorBase> BackwardRecursive(
for (std::string& grad_output : grad_op->outputs_) { for (std::string& grad_output : grad_op->outputs_) {
if (no_grad_names.count(grad_output)) { if (no_grad_names.count(grad_output)) {
grad_output = OperatorBase::EMPTY_VAR_NAME(); grad_output = kEmptyVarName;
} }
} }
...@@ -168,14 +166,14 @@ std::shared_ptr<OperatorBase> Backward( ...@@ -168,14 +166,14 @@ std::shared_ptr<OperatorBase> Backward(
std::unordered_set<std::string> no_grad_names; std::unordered_set<std::string> no_grad_names;
no_grad_names.reserve(no_grad_vars.size()); no_grad_names.reserve(no_grad_vars.size());
no_grad_names.insert(OperatorBase::EMPTY_VAR_NAME() + no_grad_names.insert(kEmptyVarName + kGradVarSuffix);
OperatorBase::GRAD_VAR_SUFFIX());
for (auto& name : no_grad_vars) { for (auto& name : no_grad_vars) {
no_grad_names.insert(name + OperatorBase::GRAD_VAR_SUFFIX()); no_grad_names.insert(name + kGradVarSuffix);
} }
size_t uid = 0; size_t uid = 0;
return BackwardRecursive(forwardOp, no_grad_names, uid); return BackwardRecursive(forwardOp, no_grad_names, uid);
} }
} // namespace framework } // namespace framework
} // namespace paddle } // namespace paddle
paddle/framework/backward_test.cc
浏览文件 @ 4bc1dc26
...@@ -15,8 +15,9 @@ ...@@ -15,8 +15,9 @@
#include "paddle/framework/backward.h" #include "paddle/framework/backward.h"
#include <gtest/gtest.h> #include <gtest/gtest.h>
#include "paddle/framework/net.h"
#include "paddle/framework/op_registry.h" #include "paddle/framework/op_registry.h"
#include "paddle/operators/net_op.h"
#include "paddle/operators/type_alias.h"
namespace paddle { namespace paddle {
namespace framework { namespace framework {
...@@ -70,21 +71,21 @@ class NoGradOpMaker : public OpProtoAndCheckerMaker { ...@@ -70,21 +71,21 @@ class NoGradOpMaker : public OpProtoAndCheckerMaker {
} }
}; };
class FcOp : public NetOp { class FcOp : public ops::NetOp {
public: public:
void Init() override { void Init() override {
AddOp(OpRegistry::CreateOp("mul", {Input("X"), Input("W")}, AddOp(OpRegistry::CreateOp("mul", {Input("X"), Input("W")},
{Output("mul_result")}, {})); {Output("mul_result")}, {}));
auto b_name = Input("b"); auto b_name = Input("b");
std::string before_act = "mul_result"; std::string before_act = "mul_result";
if (b_name != EMPTY_VAR_NAME()) { if (b_name != kEmptyVarName) {
AddOp(OpRegistry::CreateOp("rowwise_add", {Output("mul_result"), b_name}, AddOp(OpRegistry::CreateOp("rowwise_add", {Output("mul_result"), b_name},
{Output("add_result")}, {})); {Output("add_result")}, {}));
before_act = "add_result"; before_act = "add_result";
} else { } else {
auto out_varname = Output("add_result"); auto out_varname = Output("add_result");
if (out_varname != EMPTY_VAR_NAME()) { if (out_varname != kEmptyVarName) {
this->Rename(out_varname, EMPTY_VAR_NAME()); this->Rename(out_varname, kEmptyVarName);
} }
} }
...@@ -161,14 +162,13 @@ TEST(Backward, simple_op_grad) { ...@@ -161,14 +162,13 @@ TEST(Backward, simple_op_grad) {
auto fwd = f::OpRegistry::CreateOp("rowwise_add", {"X", "b"}, {"Out"}, {}); auto fwd = f::OpRegistry::CreateOp("rowwise_add", {"X", "b"}, {"Out"}, {});
ASSERT_NE(fwd, nullptr); ASSERT_NE(fwd, nullptr);
auto gop = f::OpRegistry::CreateGradOp(*fwd); auto gop = f::OpRegistry::CreateGradOp(*fwd);
ASSERT_EQ(1UL, gop->inputs_.size()); ASSERT_EQ(4UL, gop->inputs_.size());
ASSERT_EQ("Out" + f::OperatorBase::GRAD_VAR_SUFFIX(), gop->inputs_[0]); ASSERT_EQ(f::kEmptyVarName, gop->inputs_[0]);
ASSERT_EQ("rowwise_add_grad", gop->type_); ASSERT_EQ("rowwise_add_grad", gop->type_);
ASSERT_EQ("X" + f::OperatorBase::GRAD_VAR_SUFFIX(), gop->outputs_[0]); ASSERT_EQ("X" + f::kGradVarSuffix, gop->outputs_[0]);
ASSERT_EQ("b" + f::OperatorBase::GRAD_VAR_SUFFIX(), gop->outputs_[1]); ASSERT_EQ("b" + f::kGradVarSuffix, gop->outputs_[1]);
ASSERT_EQ("X" + f::OperatorBase::GRAD_VAR_SUFFIX(), ASSERT_EQ("X" + f::kGradVarSuffix, gop->Output("X" + f::kGradVarSuffix));
gop->Output("X" + f::OperatorBase::GRAD_VAR_SUFFIX()));
} }
TEST(Backward, simple_op_not_need_grad) { TEST(Backward, simple_op_not_need_grad) {
...@@ -176,13 +176,14 @@ TEST(Backward, simple_op_not_need_grad) { ...@@ -176,13 +176,14 @@ TEST(Backward, simple_op_not_need_grad) {
ASSERT_NE(fwd, nullptr); ASSERT_NE(fwd, nullptr);
auto gop = f::Backward(*fwd, {"X"}); auto gop = f::Backward(*fwd, {"X"});
ASSERT_EQ(std::find(gop->outputs_.begin(), gop->outputs_.end(), ASSERT_EQ(std::find(gop->outputs_.begin(), gop->outputs_.end(),
"X" + f::OperatorBase::GRAD_VAR_SUFFIX()), "X" + f::kGradVarSuffix),
gop->outputs_.end()); gop->outputs_.end());
auto no_input_gop = f::Backward(*fwd, {"X", "b"}); auto no_input_gop = f::Backward(*fwd, {"X", "b"});
ASSERT_NE(no_input_gop, nullptr); ASSERT_NE(no_input_gop, nullptr);
ASSERT_TRUE(no_input_gop->IsNetOp()); ASSERT_TRUE(no_input_gop->IsNetOp());
ASSERT_EQ(0UL, std::static_pointer_cast<f::NetOp>(no_input_gop)->ops_.size()); ASSERT_EQ(0UL,
std::static_pointer_cast<ops::NetOp>(no_input_gop)->ops_.size());
} }
TEST(Backward, net_fc_backward_normal) { TEST(Backward, net_fc_backward_normal) {
...@@ -191,7 +192,7 @@ TEST(Backward, net_fc_backward_normal) { ...@@ -191,7 +192,7 @@ TEST(Backward, net_fc_backward_normal) {
ASSERT_NE(fwd, nullptr); ASSERT_NE(fwd, nullptr);
std::shared_ptr<f::OperatorBase> gop = f::Backward(*fwd, {}); std::shared_ptr<f::OperatorBase> gop = f::Backward(*fwd, {});
ASSERT_TRUE(gop->IsNetOp()); ASSERT_TRUE(gop->IsNetOp());
auto net = static_cast<f::NetOp *>(gop.get()); auto net = static_cast<ops::NetOp *>(gop.get());
ASSERT_NO_THROW(net->DebugString()); ASSERT_NO_THROW(net->DebugString());
...@@ -208,13 +209,13 @@ TEST(Backward, net_fc_backward_normal) { ...@@ -208,13 +209,13 @@ TEST(Backward, net_fc_backward_normal) {
} }
TEST(Backward, net_fc_backward_not_have_b) { TEST(Backward, net_fc_backward_not_have_b) {
std::shared_ptr<f::OperatorBase> fwd = f::OpRegistry::CreateOp( std::shared_ptr<f::OperatorBase> fwd =
"fc", {"X", "w", f::OperatorBase::EMPTY_VAR_NAME()}, f::OpRegistry::CreateOp("fc", {"X", "w", f::kEmptyVarName},
{"mul_result", "add_result", "tmp"}, {}); {"mul_result", "add_result", "tmp"}, {});
ASSERT_NE(fwd, nullptr); ASSERT_NE(fwd, nullptr);
std::shared_ptr<f::OperatorBase> gop = f::Backward(*fwd, {}); std::shared_ptr<f::OperatorBase> gop = f::Backward(*fwd, {});
ASSERT_TRUE(gop->IsNetOp()); ASSERT_TRUE(gop->IsNetOp());
auto net = static_cast<f::NetOp *>(gop.get()); auto net = static_cast<ops::NetOp *>(gop.get());
ASSERT_NO_THROW(net->DebugString()); ASSERT_NO_THROW(net->DebugString());
...@@ -228,7 +229,7 @@ TEST(Backward, net_fc_backward_not_have_b) { ...@@ -228,7 +229,7 @@ TEST(Backward, net_fc_backward_not_have_b) {
} }
TEST(Backward, net_input_of_network_not_need_grad) { TEST(Backward, net_input_of_network_not_need_grad) {
f::NetOp net; ops::NetOp net;
net.AddOp(f::OpRegistry::CreateOp("fc", {"X", "W1", "b1"}, net.AddOp(f::OpRegistry::CreateOp("fc", {"X", "W1", "b1"},
{"mul_tmp_0", "add_tmp_0", "hidden0"}, {})); {"mul_tmp_0", "add_tmp_0", "hidden0"}, {}));
net.AddOp(f::OpRegistry::CreateOp("fc", {"hidden0", "W2", "b2"}, net.AddOp(f::OpRegistry::CreateOp("fc", {"hidden0", "W2", "b2"},
...@@ -236,39 +237,36 @@ TEST(Backward, net_input_of_network_not_need_grad) { ...@@ -236,39 +237,36 @@ TEST(Backward, net_input_of_network_not_need_grad) {
net.CompleteAddOp(); net.CompleteAddOp();
auto bwd = Backward(net, {"X"}); // X@GRAD is not need. auto bwd = Backward(net, {"X"}); // X@GRAD is not need.
ASSERT_TRUE(bwd->IsNetOp()); ASSERT_TRUE(bwd->IsNetOp());
auto bwd_net = static_cast<f::NetOp *>(bwd.get()); auto bwd_net = static_cast<ops::NetOp *>(bwd.get());
std::unordered_set<std::string> all_output = std::unordered_set<std::string>( std::unordered_set<std::string> all_output = std::unordered_set<std::string>(
bwd_net->outputs_.begin(), bwd_net->outputs_.end()); bwd_net->outputs_.begin(), bwd_net->outputs_.end());
all_output.erase(f::OperatorBase::EMPTY_VAR_NAME()); all_output.erase(f::kEmptyVarName);
for (auto &out : {"W1", "b1", "hidden0", "W2", "b2"}) { for (auto &out : {"W1", "b1", "hidden0", "W2", "b2"}) {
ASSERT_NE(all_output.find(out + f::OperatorBase::GRAD_VAR_SUFFIX()), ASSERT_NE(all_output.find(out + f::kGradVarSuffix), all_output.end());
all_output.end());
} }
// Not Generated X // Not Generated X
ASSERT_EQ(all_output.find("X" + f::OperatorBase::GRAD_VAR_SUFFIX()), ASSERT_EQ(all_output.find("X" + f::kGradVarSuffix), all_output.end());
all_output.end());
ASSERT_EQ(2UL, bwd_net->ops_.size()); ASSERT_EQ(2UL, bwd_net->ops_.size());
ASSERT_TRUE(bwd_net->ops_[1]->IsNetOp()); ASSERT_TRUE(bwd_net->ops_[1]->IsNetOp());
auto first_fc_grad = static_cast<f::NetOp *>(bwd_net->ops_[1].get()); auto first_fc_grad = static_cast<ops::NetOp *>(bwd_net->ops_[1].get());
ASSERT_EQ(3UL, first_fc_grad->ops_.size()); ASSERT_EQ(3UL, first_fc_grad->ops_.size());
ASSERT_EQ( ASSERT_EQ(f::kEmptyVarName,
f::OperatorBase::EMPTY_VAR_NAME(), first_fc_grad->ops_[2]->Output("A" + f::kGradVarSuffix));
first_fc_grad->ops_[2]->Output("A" + f::OperatorBase::GRAD_VAR_SUFFIX()));
} }
TEST(Backward, net_shared_weight) { TEST(Backward, net_shared_weight) {
f::NetOp net; ops::NetOp net;
net.AddOp(f::OpRegistry::CreateOp("mul", {"X", "W"}, {"Out"}, {})); net.AddOp(f::OpRegistry::CreateOp("mul", {"X", "W"}, {"Out"}, {}));
net.AddOp(f::OpRegistry::CreateOp("mul", {"Out", "W"}, {"FinalOut"}, {})); net.AddOp(f::OpRegistry::CreateOp("mul", {"Out", "W"}, {"FinalOut"}, {}));
net.CompleteAddOp(); net.CompleteAddOp();
auto bwd = f::Backward(net, {}); auto bwd = f::Backward(net, {});
ASSERT_TRUE(bwd->IsNetOp()); ASSERT_TRUE(bwd->IsNetOp());
auto bwd_net = static_cast<f::NetOp *>(bwd.get()); auto bwd_net = static_cast<ops::NetOp *>(bwd.get());
ASSERT_EQ(3UL, bwd_net->ops_.size()); ASSERT_EQ(3UL, bwd_net->ops_.size());
ASSERT_EQ("add", bwd_net->ops_[2]->type_); ASSERT_EQ("add", bwd_net->ops_[2]->type_);
} }
...@@ -285,7 +283,7 @@ TEST(Backward, op_all_input_are_not_need) { ...@@ -285,7 +283,7 @@ TEST(Backward, op_all_input_are_not_need) {
auto fwd = f::OpRegistry::CreateOp("rowwise_add", {"X", "b"}, {"Out"}, {}); auto fwd = f::OpRegistry::CreateOp("rowwise_add", {"X", "b"}, {"Out"}, {});
auto backward = f::Backward(*fwd, {"X", "b"}); auto backward = f::Backward(*fwd, {"X", "b"});
ASSERT_TRUE(backward->IsNetOp()); ASSERT_TRUE(backward->IsNetOp());
auto net = static_cast<f::NetOp *>(backward.get()); auto net = static_cast<ops::NetOp *>(backward.get());
ASSERT_TRUE(net->ops_.empty()); ASSERT_TRUE(net->ops_.empty());
} }
...@@ -293,7 +291,7 @@ TEST(Backward, op_all_output_are_not_need) { ...@@ -293,7 +291,7 @@ TEST(Backward, op_all_output_are_not_need) {
auto fwd = f::OpRegistry::CreateOp("rowwise_add", {"X", "b"}, {"Out"}, {}); auto fwd = f::OpRegistry::CreateOp("rowwise_add", {"X", "b"}, {"Out"}, {});
auto backward = f::Backward(*fwd, {"Out"}); auto backward = f::Backward(*fwd, {"Out"});
ASSERT_TRUE(backward->IsNetOp()); ASSERT_TRUE(backward->IsNetOp());
auto net = static_cast<f::NetOp *>(backward.get()); auto net = static_cast<ops::NetOp *>(backward.get());
ASSERT_TRUE(net->ops_.empty()); ASSERT_TRUE(net->ops_.empty());
} }
...@@ -301,7 +299,7 @@ TEST(Backward, op_part_of_output_are_not_need) { ...@@ -301,7 +299,7 @@ TEST(Backward, op_part_of_output_are_not_need) {
auto fwd = f::OpRegistry::CreateOp("many_output_op", {"X"}, {"Y", "Z"}, {}); auto fwd = f::OpRegistry::CreateOp("many_output_op", {"X"}, {"Y", "Z"}, {});
auto backward = f::Backward(*fwd, {"Z"}); auto backward = f::Backward(*fwd, {"Z"});
ASSERT_TRUE(backward->IsNetOp()); ASSERT_TRUE(backward->IsNetOp());
auto net = static_cast<f::NetOp *>(backward.get()); auto net = static_cast<ops::NetOp *>(backward.get());
ASSERT_EQ(net->ops_.size(), 2UL); ASSERT_EQ(net->ops_.size(), 2UL);
auto &fill_zero = *net->ops_[0]; auto &fill_zero = *net->ops_[0];
...@@ -309,17 +307,15 @@ TEST(Backward, op_part_of_output_are_not_need) { ...@@ -309,17 +307,15 @@ TEST(Backward, op_part_of_output_are_not_need) {
ASSERT_EQ(1UL, fill_zero.inputs_.size()); ASSERT_EQ(1UL, fill_zero.inputs_.size());
ASSERT_EQ("Z", fill_zero.inputs_[0]); ASSERT_EQ("Z", fill_zero.inputs_[0]);
ASSERT_EQ(1UL, fill_zero.outputs_.size()); ASSERT_EQ(1UL, fill_zero.outputs_.size());
ASSERT_EQ("Z" + f::OperatorBase::ZERO_VAR_SUFFIX(), fill_zero.outputs_[0]); ASSERT_EQ("Z" + f::kZeroVarSuffix, fill_zero.outputs_[0]);
auto &d_many_out = *net->ops_[1]; auto &d_many_out = *net->ops_[1];
ASSERT_EQ("many_output_op_grad", d_many_out.type_); ASSERT_EQ("many_output_op_grad", d_many_out.type_);
ASSERT_EQ(1UL + 2UL + 2UL, d_many_out.inputs_.size()); // I/O/OG ASSERT_EQ(1UL + 2UL + 2UL, d_many_out.inputs_.size()); // I/O/OG
ASSERT_EQ("Z" + f::OperatorBase::ZERO_VAR_SUFFIX(), ASSERT_EQ("Z" + f::kZeroVarSuffix, d_many_out.Input("z" + f::kGradVarSuffix));
d_many_out.Input("z" + f::OperatorBase::GRAD_VAR_SUFFIX())); ASSERT_EQ("Y" + f::kGradVarSuffix, d_many_out.Input("y" + f::kGradVarSuffix));
ASSERT_EQ("Y" + f::OperatorBase::GRAD_VAR_SUFFIX(), ASSERT_EQ("X" + f::kGradVarSuffix,
d_many_out.Input("y" + f::OperatorBase::GRAD_VAR_SUFFIX())); d_many_out.Output("x" + f::kGradVarSuffix));
ASSERT_EQ("X" + f::OperatorBase::GRAD_VAR_SUFFIX(),
d_many_out.Output("x" + f::OperatorBase::GRAD_VAR_SUFFIX()));
} }
TEST(Backward, op_part_of_input_are_not_need) { TEST(Backward, op_part_of_input_are_not_need) {
...@@ -329,19 +325,17 @@ TEST(Backward, op_part_of_input_are_not_need) { ...@@ -329,19 +325,17 @@ TEST(Backward, op_part_of_input_are_not_need) {
ASSERT_EQ(grad_mul.type_, "mul_grad"); ASSERT_EQ(grad_mul.type_, "mul_grad");
ASSERT_EQ(grad_mul.inputs_.size(), 2UL + 1UL + 1UL); ASSERT_EQ(grad_mul.inputs_.size(), 2UL + 1UL + 1UL);
ASSERT_EQ(grad_mul.outputs_.size(), 2UL); ASSERT_EQ(grad_mul.outputs_.size(), 2UL);
ASSERT_EQ(grad_mul.Output("A" + f::OperatorBase::GRAD_VAR_SUFFIX()), ASSERT_EQ(grad_mul.Output("A" + f::kGradVarSuffix), f::kEmptyVarName);
f::OperatorBase::EMPTY_VAR_NAME()); ASSERT_EQ(grad_mul.Output("B" + f::kGradVarSuffix), "b" + f::kGradVarSuffix);
ASSERT_EQ(grad_mul.Output("B" + f::OperatorBase::GRAD_VAR_SUFFIX()), ASSERT_EQ(grad_mul.Input("Out" + f::kGradVarSuffix),
"b" + f::OperatorBase::GRAD_VAR_SUFFIX()); "out" + f::kGradVarSuffix);
ASSERT_EQ(grad_mul.Input("Out" + f::OperatorBase::GRAD_VAR_SUFFIX()),
"out" + f::OperatorBase::GRAD_VAR_SUFFIX());
ASSERT_EQ(grad_mul.Input("A"), "a"); ASSERT_EQ(grad_mul.Input("A"), "a");
ASSERT_EQ(grad_mul.Input("B"), "b"); ASSERT_EQ(grad_mul.Input("B"), "b");
ASSERT_EQ(grad_mul.Input("Out"), "out"); ASSERT_EQ(grad_mul.Input("Out"), "out");
} }
TEST(Backward, linear_net_intermediate_variable_has_no_grad) { TEST(Backward, linear_net_intermediate_variable_has_no_grad) {
f::NetOp net; ops::NetOp net;
net.AddOp(f::OpRegistry::CreateOp("fc", {"x1", "w1", "b1"}, net.AddOp(f::OpRegistry::CreateOp("fc", {"x1", "w1", "b1"},
{"mul_out1", "add_out1", "out1"}, {})); {"mul_out1", "add_out1", "out1"}, {}));
net.AddOp(f::OpRegistry::CreateOp("fc", {"out1", "w2", "b2"}, net.AddOp(f::OpRegistry::CreateOp("fc", {"out1", "w2", "b2"},
...@@ -351,14 +345,13 @@ TEST(Backward, linear_net_intermediate_variable_has_no_grad) { ...@@ -351,14 +345,13 @@ TEST(Backward, linear_net_intermediate_variable_has_no_grad) {
net.CompleteAddOp(); net.CompleteAddOp();
auto backward = f::Backward(net, {"mul_out2", "tmp_out2", "out2"}); auto backward = f::Backward(net, {"mul_out2", "tmp_out2", "out2"});
ASSERT_TRUE(backward->IsNetOp()); ASSERT_TRUE(backward->IsNetOp());
auto bwd_net = static_cast<f::NetOp *>(backward.get()); auto bwd_net = static_cast<ops::NetOp *>(backward.get());
ASSERT_EQ(bwd_net->ops_.size(), 3UL); ASSERT_EQ(bwd_net->ops_.size(), 3UL);
auto &grad_fc = *bwd_net->ops_[0]; auto &grad_fc = *bwd_net->ops_[0];
EXPECT_EQ(grad_fc.inputs_.size(), EXPECT_EQ(grad_fc.inputs_.size(),
3UL /* external input number */ 3UL /* external input number */
+ 1UL /* external output number*/ + 1UL /* external output number*/
+ 1UL /* number of gradient of external output*/ + 1UL /* number of gradient of external output*/
- 1UL /*ignoreGradient varable number*/
+ 2U /* internal variable number*/); + 2U /* internal variable number*/);
EXPECT_EQ(grad_fc.outputs_.size(), 2UL /* input number of mul*/ EXPECT_EQ(grad_fc.outputs_.size(), 2UL /* input number of mul*/
+ 2UL /* input number of rowwise_add */ + 2UL /* input number of rowwise_add */
...@@ -367,23 +360,4 @@ TEST(Backward, linear_net_intermediate_variable_has_no_grad) { ...@@ -367,23 +360,4 @@ TEST(Backward, linear_net_intermediate_variable_has_no_grad) {
EXPECT_EQ(bwd_net->ops_[1]->outputs_.size(), 0UL); EXPECT_EQ(bwd_net->ops_[1]->outputs_.size(), 0UL);
EXPECT_EQ(bwd_net->ops_[2]->inputs_.size(), 0UL); EXPECT_EQ(bwd_net->ops_[2]->inputs_.size(), 0UL);
EXPECT_EQ(bwd_net->ops_[2]->outputs_.size(), 0UL); EXPECT_EQ(bwd_net->ops_[2]->outputs_.size(), 0UL);
/*
EXPECT_EQ(grad_fc.Output("X" + f::OperatorBase::GRAD_VAR_SUFFIX()),
f::OperatorBase::EMPTY_VAR_NAME());
EXPECT_EQ(grad_fc.Output("W" + f::OperatorBase::GRAD_VAR_SUFFIX()),
"w3" + f::OperatorBase::GRAD_VAR_SUFFIX());
EXPECT_EQ(grad_fc.Output("b" + f::OperatorBase::GRAD_VAR_SUFFIX()),
"b3" + f::OperatorBase::GRAD_VAR_SUFFIX());
EXPECT_EQ(grad_fc.Output("mul_result" + f::OperatorBase::GRAD_VAR_SUFFIX()),
"mul_out3" + f::OperatorBase::GRAD_VAR_SUFFIX());
EXPECT_EQ(grad_fc.Input("Out" + f::OperatorBase::GRAD_VAR_SUFFIX()),
"out3" + f::OperatorBase::GRAD_VAR_SUFFIX());
EXPECT_EQ(grad_fc.Input("X"), "out2");
EXPECT_EQ(grad_fc.Input("W"), "w3");
EXPECT_EQ(grad_fc.Input("mul_result"), "mul_out3");
EXPECT_EQ(grad_fc.Input("add_result"), "tmp_out3");
EXPECT_EQ(grad_fc.Input("Out"), "out3");
*/
} }
paddle/framework/ddim.h
浏览文件 @ 4bc1dc26
...@@ -25,18 +25,15 @@ limitations under the License. */ ...@@ -25,18 +25,15 @@ limitations under the License. */
namespace paddle { namespace paddle {
namespace framework { namespace framework {
namespace {
typedef boost::variant<Dim<1>, Dim<2>, Dim<3>, Dim<4>, Dim<5>, Dim<6>, Dim<7>,
Dim<8>, Dim<9>>
DDimVar;
}
/** /**
* \brief A dynamically sized dimension. * \brief A dynamically sized dimension.
* *
* The number of dimensions must be between [1, 9]. * The number of dimensions must be between [1, 9].
*/ */
struct DDim { struct DDim {
typedef boost::variant<Dim<1>, Dim<2>, Dim<3>, Dim<4>, Dim<5>, Dim<6>, Dim<7>,
Dim<8>, Dim<9>>
DDimVar;
DDimVar var; DDimVar var;
DDim() : var(Dim<1>()) {} DDim() : var(Dim<1>()) {}
......
paddle/framework/grad_op_builder.cc
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/framework/grad_op_builder.h
浏览文件 @ 4bc1dc26
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
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. */
#pragma once #pragma once
#include "paddle/framework/op_proto.pb.h"
#include "paddle/framework/operator.h" #include "paddle/framework/operator.h"
namespace paddle { namespace paddle {
namespace framework { namespace framework {
class OpRegistry;
OperatorBase* BuildGradOp(const OperatorBase* op);
enum InOutType { IN, OUT };
struct OpInOutArg {
OpInOutArg(const std::string& proto_name, const InOutType& type,
bool needed_in_grad, size_t begin_idx, size_t end_idx)
: proto_name_(proto_name),
type_(type),
needed_in_grad_(needed_in_grad),
begin_idx_(begin_idx),
end_idx_(end_idx) {}
std::string proto_name_;
InOutType type_;
bool needed_in_grad_;
size_t begin_idx_;
size_t end_idx_;
};
class GradOpBuilder {
using VarIndexMap = std::unordered_map<std::string, int>;
public:
GradOpBuilder(const OperatorBase& op) : op_(op) {}
OperatorBase* Build();
private:
OpInOutArg* BuildArg(const VarProto& var, const VarIndexMap& var_map,
const std::vector<int>& format, InOutType type);
void BuildOpInOutArgList();
void AddArgIntoGradOp(const OpInOutArg* arg, std::vector<std::string>& in_out,
std::vector<int>& format, VarIndexMap* varmap, int& idx,
bool is_grad) const;
void CompleteGradOp(OperatorBase* grad_op) const;
const OperatorBase& op_;
std::vector<std::shared_ptr<OpInOutArg>> arg_list_;
};
} // namespace framework } // namespace framework
} // namespace paddle } // namespace paddle
paddle/framework/op_desc.proto
浏览文件 @ 4bc1dc26
...@@ -12,10 +12,10 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ...@@ -12,10 +12,10 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
syntax="proto2"; syntax = "proto2";
package paddle.framework; package paddle.framework;
import "attr_type.proto"; import "attribute.proto";
// AttrDesc is used to describe Attributes of an Operator. It contain's // AttrDesc is used to describe Attributes of an Operator. It contain's
// name, type, and value of Attribute. // name, type, and value of Attribute.
......
paddle/framework/op_proto.proto
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/framework/op_registry.cc
浏览文件 @ 4bc1dc26
...@@ -14,37 +14,8 @@ limitations under the License. */ ...@@ -14,37 +14,8 @@ limitations under the License. */
#include <paddle/framework/op_registry.h> #include <paddle/framework/op_registry.h>
namespace paddle { #include <vector>
namespace framework {
template <>
void AttrTypeHelper::SetAttrType<int>(AttrProto* attr) {
attr->set_type(paddle::framework::AttrType::INT);
}
template <>
void AttrTypeHelper::SetAttrType<float>(AttrProto* attr) {
attr->set_type(paddle::framework::AttrType::FLOAT);
}
template <>
void AttrTypeHelper::SetAttrType<std::string>(AttrProto* attr) {
attr->set_type(paddle::framework::AttrType::STRING);
}
template <> namespace paddle {
void AttrTypeHelper::SetAttrType<std::vector<int>>(AttrProto* attr) { namespace framework {} // namespace framework
attr->set_type(paddle::framework::AttrType::INTS);
}
template <>
void AttrTypeHelper::SetAttrType<std::vector<float>>(AttrProto* attr) {
attr->set_type(paddle::framework::AttrType::FLOATS);
}
template <>
void AttrTypeHelper::SetAttrType<std::vector<std::string>>(AttrProto* attr) {
attr->set_type(paddle::framework::AttrType::STRINGS);
}
} // namespace framework
} // namespace paddle } // namespace paddle
paddle/framework/op_registry.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/framework/operator.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/framework/operator_test.cc
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/framework/tensor.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/function/BufferArgTest.cpp
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/function/CosSimOpGpu.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/function/CropOpGpu.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/function/FunctionTest.cpp
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/function/Im2ColOpGpu.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/function/MulOpGpu.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/function/PadOpGpu.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/function/RowConvOpGpu.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/function/TensorShapeTest.cpp
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/function/TensorTypeTest.cpp
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/gserver/layers/GruCompute.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/gserver/tests/CMakeLists.txt
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/math/BaseMatrix.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/math/MathUtils.cpp
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/math/TrainingAlgorithmOp.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/math/tests/test_Tensor.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/memory/detail/meta_cache.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/memory/memory.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/.clang-format 0 → 100644
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/CMakeLists.txt
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/add_op.cc
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/add_op.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/add_op.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/cross_entropy_op.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/fc_op.cc
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/mean_op.cc
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/mean_op.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/mean_op.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/mul_op.cc
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/mul_op.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/recurrent_op.cc
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/recurrent_op.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/rnn/recurrent_op_utils.cc 0 → 100644
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/rnn/recurrent_op_utils.h 0 → 100644
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/rowwise_add_op.cc
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/rowwise_add_op.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/rowwise_add_op.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/sgd_op.cc
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/sgd_op.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/sgd_op.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/sigmoid_op.cc
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/sigmoid_op.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/sigmoid_op.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/softmax_op.cc
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/softmax_op.cu
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/softmax_op.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/operators/type_alias.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/platform/device_context.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/platform/dynload/cublas.cc
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/platform/dynload/cudnn.cc
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/platform/dynload/curand.cc
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/platform/enforce.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/platform/enforce_test.cc
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/platform/place.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/scripts/docker/build.sh
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/scripts/run_python_tests.sh 已删除 100755 → 0
浏览文件 @ 6405438b
此差异已折叠。
paddle/setup.py.in
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/string/piece.h
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/trainer/tests/compare_sparse_data 0 → 100644
浏览文件 @ 4bc1dc26
此差异已折叠。
paddle/trainer/tests/train_sparse.list 0 → 100644
浏览文件 @ 4bc1dc26
此差异已折叠。
proto/DataConfig.proto
浏览文件 @ 4bc1dc26
此差异已折叠。
proto/DataFormat.proto
浏览文件 @ 4bc1dc26
此差异已折叠。
proto/ModelConfig.proto
浏览文件 @ 4bc1dc26
此差异已折叠。
proto/OptimizerConfig.proto
浏览文件 @ 4bc1dc26
此差异已折叠。
proto/ParameterConfig.proto
浏览文件 @ 4bc1dc26
此差异已折叠。
proto/ParameterServerConfig.proto
浏览文件 @ 4bc1dc26
此差异已折叠。
proto/ParameterService.proto
浏览文件 @ 4bc1dc26
此差异已折叠。
proto/TrainerConfig.proto
浏览文件 @ 4bc1dc26
此差异已折叠。
python/paddle/v2/dataset/cifar.py
浏览文件 @ 4bc1dc26
此差异已折叠。
python/paddle/v2/dataset/common.py
浏览文件 @ 4bc1dc26
此差异已折叠。
python/paddle/v2/dataset/conll05.py
浏览文件 @ 4bc1dc26
此差异已折叠。
python/paddle/v2/dataset/imdb.py
浏览文件 @ 4bc1dc26
此差异已折叠。
python/paddle/v2/dataset/mnist.py
浏览文件 @ 4bc1dc26
此差异已折叠。
python/paddle/v2/dataset/wmt14.py
浏览文件 @ 4bc1dc26
此差异已折叠。
python/paddle/v2/framework/tests/gradient_checker.py 0 → 100644
浏览文件 @ 4bc1dc26
此差异已折叠。
python/setup.py.in
浏览文件 @ 4bc1dc26
此差异已折叠。
Markdown is supported
0% 或 .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑！
想要评论请 注册 或