提交 4e228021 编写于 作者: Z zchen0211

Merge branch 'develop' of https://github.com/PaddlePaddle/Paddle into develop

# 构建iOS平台上的PaddlePaddle库
交叉编译iOS平台上适用的PaddlePaddle库,需要在MacOS系统上进行。本文的将介绍在MacOS上,从源码交叉编译iOS平台上适用的PaddlePaddle库。
## 准备交叉编译环境
Apple官方为iOS开发提供了完整的交叉编译工具和集成开发环境,用户从App Store下载安装Xcode即可。也可自行前往官网下载,[Xcode](https://developer.apple.com/cn/xcode/)。安装完成之后,可在命令行执行`xcodebuild -version`,判断是否安装成功。
```bash
$ xcodebuild -version
Xcode 9.0
Build version 9A235
```
## 配置交叉编译参数
PaddlePaddle为交叉编译提供了工具链配置文档[cmake/cross_compiling/ios.cmake](https://github.com/PaddlePaddle/Paddle/blob/develop/cmake/cross_compiling/ios.cmake),以提供一些默认的编译器和编译参数配置。
交叉编译iOS版本的PaddlePaddle库时,有一些必须配置的参数:
- `CMAKE_SYSTEM_NAME`,CMake编译的目标平台,必须设置为`iOS`。在设置`CMAKE_SYSTEM_NAME=iOS`后,PaddlePaddle的CMake系统会自动编译所有的第三方依赖库,并且强制设置一些PaddlePaddle参数的值(`WITH_C_API=ON``WITH_GPU=OFF``WITH_AVX=OFF``WITH_PYTHON=OFF``WITH_RDMA=OFF`)。
- `WITH_C_API`,是否编译C-API预测库,必须设置为ON。在iOS平台上只支持使用C-API来预测。
- `WITH_SWIG_PY`,必须设置为ON。在iOS平台上不支持通过swig调用来训练或者预测。
iOS平台可选配置参数:
- `IOS_PLATFORM`,可设置为`OS/SIMULATOR`,默认值为`OS`
- `OS`,构建目标为`arm`架构的iPhone或者iPad等物理设备。
- `SIMULATOR`,构建目标为`x86`架构的模拟器平台。
- `IOS_ARCH`,目标架构。针对不同的`IOS_PLATFORM`,可设置的目标架构如下表所示:
| IOS_PLATFORM | IOS_ARCH |
|--------------|----------------------|
| OS | armv7, armv7s, arm64 (默认) |
| SIMULATOR | i386, x86_64 (默认) |
- `IOS_DEPLOYMENT_TARGET`,最小的iOS部署版本,默认值为`7.0`
- `IOS_ENABLE_BITCODE`,是否使能[Bitcode](https://developer.apple.com/library/content/documentation/IDEs/Conceptual/AppDistributionGuide/AppThinning/AppThinning.html#//apple_ref/doc/uid/TP40012582-CH35-SW3),可设置`ON/OFF`,默认值为`ON`
- `IOS_USE_VECLIB_FOR_BLAS`,是否使用[vecLib](https://developer.apple.com/documentation/accelerate/veclib)框架进行BLAS矩阵计算,可设置`ON/OFF`,默认值为`OFF`
- `IOS_DEVELOPMENT_ROOT``Developer`目录,可显式指定为`/path/to/platform/Developer`。若未显式指定,PaddlePaddle将会根据`IOS_PLATFORM`自动选择`Xcode`对应`platform``Developer`目录。
- `IOS_SDK_ROOT`,所使用`SDK`的根目录,可显式指定为`/path/to/platform/Developer/SDKs/SDK`。若未显式指定,PaddlePaddle将会自动选择`IOS_DEVELOPMENT_ROOT`目录下最新的`SDK`版本。
其他配置参数:
- `USE_EIGEN_FOR_BLAS`,是否使用Eigen库进行矩阵计算,在`IOS_USE_VECLIB_FOR_BLAS=OFF`时有效。可设置`ON/OFF`,默认值为`OFF`
- `HOST_C/CXX_COMPILER`,宿主机的C/C++编译器。默认值为环境变量`CC/CXX`的值;若环境变量`CC/CXX`未设置,则使用`cc/c++`编译器。
常用的cmake配置如下:
```bash
cmake -DCMAKE_SYSTEM_NAME=iOS \
-DIOS_PLATFORM=OS \
-DIOS_ARCH="arm64" \
-DIOS_ENABLE_BITCODE=ON \
-DIOS_USE_VECLIB_FOR_BLAS=ON \
-DCMAKE_INSTALL_PREFIX=your/path/to/install \
-DWITH_C_API=ON \
-DWITH_TESTING=OFF \
-DWITH_SWIG_PY=OFF \
..
```
```bash
cmake -DCMAKE_SYSTEM_NAME=iOS \
-DIOS_PLATFORM=SIMULATOR \
-DIOS_ARCH="x86_64" \
-DIOS_USE_VECLIB_FOR_BLAS=ON \
-DCMAKE_INSTALL_PREFIX=your/path/to/install \
-DWITH_C_API=ON \
-DWITH_TESTING=OFF \
-DWITH_SWIG_PY=OFF \
..
```
用户还可根据自己的需求设置其他编译参数。比如希望最小化生成库的大小,可以设置`CMAKE_BUILD_TYPE``MinSizeRel`;若希望得到最快的执行速度,则可设置`CMAKE_BUILD_TYPE``Release`。亦可以通过手动设置`CMAKE_C/CXX_FLAGS`来影响PaddlePaddle的编译过程。
**性能TIPS**,为了达到最快的计算速度,在CMake参数配置上,有以下建议:
- 设置`CMAKE_BUILD_TYPE``Release`
- 设置`IOS_USE_VECLIB_FOR_BLAS=ON`,调用`vecLib`框架提供的BLAS函数进行矩阵计算。
## 编译和安装
CMake配置完成后,执行以下命令,PaddlePaddle将自动下载和编译所有第三方依赖库、编译和安装PaddlePaddle预测库。
```
$ make
$ make install
```
注意:如果你曾在源码目录下编译过其他平台的PaddlePaddle库,请先使用`rm -rf`命令删除`third_party`目录和`build`目录,以确保所有的第三方依赖库和PaddlePaddle代码都是针对新的CMake配置重新编译的。
执行完安装命令后,`your/path/to/install`目录中会包含以下内容:
- `include`目录,其中包含所有C-API的头文件
- `lib`目录,其中包含PaddlePaddle的C-API静态库
- `third_party`目录,其中包含所依赖的所有第三方库
注意,不同架构的PaddlePaddle库建议安装到不同的目录下,然后使用`lipo`工具将多个静态库合并成一个支持多个架构的fat库。
自此,PaddlePaddle库已经安装完成,用户可将合成的fat库用于深度学习相关的iOS App中,调用方法见C-API文档。
...@@ -59,4 +59,4 @@ make install ...@@ -59,4 +59,4 @@ make install
注意:如果你曾经在源码目录下编译过其他平台的PaddlePaddle库,请先使用`rm -rf`命令删除`third_party`目录和`build`目录,以确保所有的第三方依赖库和PaddlePaddle代码都是针对新的CMake配置重新编译的。 注意:如果你曾经在源码目录下编译过其他平台的PaddlePaddle库,请先使用`rm -rf`命令删除`third_party`目录和`build`目录,以确保所有的第三方依赖库和PaddlePaddle代码都是针对新的CMake配置重新编译的。
执行完安装命令后,`your/path/to/install`目录中会包含`include``lib`目录,其中`include`中包含C-API的头文件,`lib`中包含一个Raspberry Pi版本的库。 执行完安装命令后,`your/path/to/install`目录中会包含`include``lib`目录,其中`include`中包含C-API的头文件,`lib`中包含一个Raspberry Pi版本的库。
...@@ -44,7 +44,7 @@ cmake -DCMAKE_SYSTEM_NAME=RPi \ ...@@ -44,7 +44,7 @@ cmake -DCMAKE_SYSTEM_NAME=RPi \
.. ..
``` ```
To build the inference library, please set the argument WITH_API to ON: `WITH_C_API=ON`. To build the inference library, please set the argument WITH\_C\_API to ON: `WITH_C_API=ON`.
You can add more arguments. For example, to minimize the size of the generated inference library, you may use `CMAKE_BUILD_TYPE=MinSizeRel`. For performance optimization, you may use `CMAKE_BUILD_TYPE=Release`. You can add more arguments. For example, to minimize the size of the generated inference library, you may use `CMAKE_BUILD_TYPE=MinSizeRel`. For performance optimization, you may use `CMAKE_BUILD_TYPE=Release`.
......
...@@ -36,8 +36,8 @@ TEST(LoDTensor, LoDInGPU) { ...@@ -36,8 +36,8 @@ TEST(LoDTensor, LoDInGPU) {
lod_tensor.mutable_data<float>(place); lod_tensor.mutable_data<float>(place);
lod_tensor.set_lod(src_lod); lod_tensor.set_lod(src_lod);
CHECK_EQ(lod_tensor.lod_element(0, 2).first, 4UL); EXPECT_EQ(lod_tensor.lod_element(0, 2).first, 4UL);
CHECK_EQ(lod_tensor.lod_element(0, 4).first, 8UL); EXPECT_EQ(lod_tensor.lod_element(0, 4).first, 8UL);
auto lod = lod_tensor.lod(); auto lod = lod_tensor.lod();
...@@ -45,6 +45,6 @@ TEST(LoDTensor, LoDInGPU) { ...@@ -45,6 +45,6 @@ TEST(LoDTensor, LoDInGPU) {
cudaDeviceSynchronize(); cudaDeviceSynchronize();
for (size_t i = 0; i < src_lod[0].size(); ++i) { for (size_t i = 0; i < src_lod[0].size(); ++i) {
CHECK_EQ(lod[0].data()[i], src_lod[0].data()[i] * 2); EXPECT_EQ(lod[0].data()[i], src_lod[0].data()[i] * 2);
} }
} }
\ No newline at end of file
...@@ -37,32 +37,32 @@ ExecutionContext::GetEigenDevice<platform::GPUPlace, Eigen::GpuDevice>() const { ...@@ -37,32 +37,32 @@ ExecutionContext::GetEigenDevice<platform::GPUPlace, Eigen::GpuDevice>() const {
std::string OperatorBase::Input(const std::string& name) const { std::string OperatorBase::Input(const std::string& name) const {
auto& ins = Inputs(name); auto& ins = Inputs(name);
PADDLE_ENFORCE_LE(ins.size(), 1UL, PADDLE_ENFORCE_LE(ins.size(), 1UL,
"Op %s input %s should contain only one variable", type_, "Operator %s's input %s should contain only one variable.",
name); type_, name);
return ins.empty() ? kEmptyVarName : ins[0]; return ins.empty() ? kEmptyVarName : ins[0];
} }
const std::vector<std::string>& OperatorBase::Inputs( const std::vector<std::string>& OperatorBase::Inputs(
const std::string& name) const { const std::string& name) const {
auto it = inputs_.find(name); auto it = inputs_.find(name);
PADDLE_ENFORCE(it != inputs_.end(), "Op %s do not have input %s", type_, PADDLE_ENFORCE(it != inputs_.end(), "Operator %s does not have the input %s.",
name); type_, name);
return it->second; return it->second;
} }
std::string OperatorBase::Output(const std::string& name) const { std::string OperatorBase::Output(const std::string& name) const {
auto& outs = Outputs(name); auto& outs = Outputs(name);
PADDLE_ENFORCE_LE(outs.size(), 1UL, PADDLE_ENFORCE_LE(outs.size(), 1UL,
"Op %s output %s should contain only one variable", type_, "Operator %s's output %s should contain only one variable.",
name); type_, name);
return outs.empty() ? kEmptyVarName : outs[0]; return outs.empty() ? kEmptyVarName : outs[0];
} }
const std::vector<std::string>& OperatorBase::Outputs( const std::vector<std::string>& OperatorBase::Outputs(
const std::string& name) const { const std::string& name) const {
auto it = outputs_.find(name); auto it = outputs_.find(name);
PADDLE_ENFORCE(it != outputs_.end(), "Op %s does not have output called %s", PADDLE_ENFORCE(it != outputs_.end(),
type_, name); "Operator %s does not have an output called %s.", type_, name);
return it->second; return it->second;
} }
......
...@@ -427,7 +427,8 @@ class OperatorWithKernel : public OperatorBase { ...@@ -427,7 +427,8 @@ class OperatorWithKernel : public OperatorBase {
int tmp = static_cast<int>(ToDataType(t->type())); int tmp = static_cast<int>(ToDataType(t->type()));
VLOG(3) << "Input " << ipt_name << " with data_type " << tmp; VLOG(3) << "Input " << ipt_name << " with data_type " << tmp;
PADDLE_ENFORCE(tmp == data_type || data_type == -1, PADDLE_ENFORCE(tmp == data_type || data_type == -1,
"DataType of Paddle Op %s must be same.", Type()); "DataType of Paddle Op %s must be the same.",
Type());
data_type = tmp; data_type = tmp;
} }
} }
......
...@@ -118,10 +118,12 @@ class Tensor { ...@@ -118,10 +118,12 @@ class Tensor {
const platform::DeviceContext& ctx); const platform::DeviceContext& ctx);
/** /**
* @brief Return the slice of the tensor. * @brief Return a sub-tensor of the given tensor.
* *
* @param[in] begin_idx The begin index of the slice. * @param[in] begin_idx The index of the start row(inclusive) to slice.
* @param[in] end_idx The end index of the slice. * The index number begins from 0.
* @param[in] end_idx The index of the end row(exclusive) to slice.
* The index number begins from 0.
*/ */
inline Tensor Slice(const int& begin_idx, const int& end_idx) const; inline Tensor Slice(const int& begin_idx, const int& end_idx) const;
......
...@@ -112,9 +112,10 @@ inline void* Tensor::mutable_data(platform::Place place, std::type_index type) { ...@@ -112,9 +112,10 @@ inline void* Tensor::mutable_data(platform::Place place, std::type_index type) {
if (holder_ != nullptr) { if (holder_ != nullptr) {
holder_->set_type(type); holder_->set_type(type);
} }
PADDLE_ENFORCE_GT(numel(), 0, PADDLE_ENFORCE_GT(
"Tensor's numel must be larger than zero to call " numel(), 0,
"Tensor::mutable_data. Call Tensor::set_dim first."); "When calling this method, the Tensor's numel must be larger than zero. "
"Please check Tensor::Resize has been called first.");
int64_t size = numel() * SizeOfType(type); int64_t size = numel() * SizeOfType(type);
/* some versions of boost::variant don't have operator!= */ /* some versions of boost::variant don't have operator!= */
if (holder_ == nullptr || !(holder_->place() == place) || if (holder_ == nullptr || !(holder_->place() == place) ||
...@@ -229,10 +230,12 @@ inline void Tensor::CopyFromVector(const std::vector<T>& src, ...@@ -229,10 +230,12 @@ inline void Tensor::CopyFromVector(const std::vector<T>& src,
inline Tensor Tensor::Slice(const int& begin_idx, const int& end_idx) const { inline Tensor Tensor::Slice(const int& begin_idx, const int& end_idx) const {
check_memory_size(); check_memory_size();
PADDLE_ENFORCE_GE(begin_idx, 0, "Slice begin index is less than zero."); PADDLE_ENFORCE_GE(begin_idx, 0,
PADDLE_ENFORCE_LE(end_idx, dims_[0], "Slice end index is out of bound."); "The start row index must be greater than 0.");
PADDLE_ENFORCE_LT(begin_idx, end_idx, PADDLE_ENFORCE_LE(end_idx, dims_[0], "The end row index is out of bound.");
"Begin index must be less than end index."); PADDLE_ENFORCE_LT(
begin_idx, end_idx,
"The start row index must be lesser than the end row index.");
if (dims_[0] == 1) { if (dims_[0] == 1) {
return *this; return *this;
......
...@@ -101,8 +101,10 @@ void CRFLayer::backward(const UpdateCallback& callback) { ...@@ -101,8 +101,10 @@ void CRFLayer::backward(const UpdateCallback& callback) {
: real(1.0f); : real(1.0f);
instanceWeight *= coeff_; instanceWeight *= coeff_;
MatrixPtr grad = output.grad->subRowMatrix(starts[i], starts[i + 1]); if (output.grad) {
grad->add(*crfs_[i].getXGrad(), real(1.0f), instanceWeight); MatrixPtr grad = output.grad->subRowMatrix(starts[i], starts[i + 1]);
grad->add(*crfs_[i].getXGrad(), real(1.0f), instanceWeight);
}
if (needWGrad) { if (needWGrad) {
weight_->getWGrad()->add( weight_->getWGrad()->add(
*crfs_[i].getWGrad(), real(1.0f), instanceWeight); *crfs_[i].getWGrad(), real(1.0f), instanceWeight);
......
...@@ -102,7 +102,6 @@ real LinearChainCRF::forward(real* x, int* s, int length) { ...@@ -102,7 +102,6 @@ real LinearChainCRF::forward(real* x, int* s, int length) {
} }
void LinearChainCRF::backward(real* x, int* s, int length, bool needWGrad) { void LinearChainCRF::backward(real* x, int* s, int length, bool needWGrad) {
MatrixPtr matX = Matrix::create(x, length, numClasses_);
Matrix::resizeOrCreate(matGrad_, length, numClasses_); Matrix::resizeOrCreate(matGrad_, length, numClasses_);
Matrix::resizeOrCreate(beta_, length, numClasses_); Matrix::resizeOrCreate(beta_, length, numClasses_);
real* b = b_->getData(); real* b = b_->getData();
......
...@@ -70,11 +70,23 @@ void SequenceReshapeLayer::forward(PassType passType) { ...@@ -70,11 +70,23 @@ void SequenceReshapeLayer::forward(PassType passType) {
size_t outDim = getSize(); size_t outDim = getSize();
size_t numSequences = input.getNumSequences(); size_t numSequences = input.getNumSequences();
auto startPositions = input.sequenceStartPositions->getVector(false);
const int* starts = startPositions->getData();
CHECK_EQ(starts[numSequences], input.getBatchSize()); // by default, we assume each instance as a sequence
CHECK_EQ(numSequences, startPositions->getSize() - 1); IVectorPtr seqStarts;
IVector::resizeOrCreate(seqStarts, input.getBatchSize() + 1, false);
int* startsData = seqStarts->getData();
for (int i = 0; i < input.getBatchSize() + 1; i++) {
startsData[i] = i;
}
const int* starts = startsData;
// if there is sequence, then use start positions
if (input.sequenceStartPositions) {
auto startPositions = input.sequenceStartPositions->getVector(false);
starts = startPositions->getData();
CHECK_EQ(starts[numSequences], input.getBatchSize());
CHECK_EQ(numSequences, startPositions->getSize() - 1);
}
for (size_t seqID = 0; seqID < numSequences; seqID++) { for (size_t seqID = 0; seqID < numSequences; seqID++) {
size_t inNumIns = starts[seqID + 1] - starts[seqID]; size_t inNumIns = starts[seqID + 1] - starts[seqID];
......
...@@ -41,7 +41,16 @@ void* CPUAllocator::Alloc(size_t& index, size_t size) { ...@@ -41,7 +41,16 @@ void* CPUAllocator::Alloc(size_t& index, size_t size) {
index = 0; // unlock memory index = 0; // unlock memory
void* p = malloc(size); void* p;
#ifdef PADDLE_USE_MKLDNN
// refer to https://github.com/01org/mkl-dnn/blob/master/include/mkldnn.hpp
// memory alignment
PADDLE_ENFORCE_EQ(posix_memalign(&p, 4096ul, size), 0);
#else
PADDLE_ENFORCE_EQ(posix_memalign(&p, 32ul, size), 0);
#endif
PADDLE_ENFORCE(p, "Fail to allocate CPU memory: size = %d .", size);
if (p != nullptr) { if (p != nullptr) {
if (FLAGS_use_pinned_memory) { if (FLAGS_use_pinned_memory) {
......
...@@ -28,8 +28,9 @@ class CrossEntropyOp : public framework::OperatorWithKernel { ...@@ -28,8 +28,9 @@ class CrossEntropyOp : public framework::OperatorWithKernel {
auto x_dims = ctx->GetInputDim("X"); auto x_dims = ctx->GetInputDim("X");
auto label_dims = ctx->GetInputDim("Label"); auto label_dims = ctx->GetInputDim("Label");
PADDLE_ENFORCE_EQ(x_dims.size(), 2, "Input(X)'s rank should be 2."); PADDLE_ENFORCE_EQ(x_dims.size(), 2UL, "Input(X)'s rank should be 2.");
PADDLE_ENFORCE_EQ(label_dims.size(), 2, "Input(Label)'s rank should be 2."); PADDLE_ENFORCE_EQ(label_dims.size(), 2UL,
"Input(Label)'s rank should be 2.");
PADDLE_ENFORCE_EQ(x_dims[0], label_dims[0], PADDLE_ENFORCE_EQ(x_dims[0], label_dims[0],
"The 1st dimension of Input(X) and Input(Label) should " "The 1st dimension of Input(X) and Input(Label) should "
"be equal."); "be equal.");
...@@ -38,8 +39,8 @@ class CrossEntropyOp : public framework::OperatorWithKernel { ...@@ -38,8 +39,8 @@ class CrossEntropyOp : public framework::OperatorWithKernel {
"If Attr(soft_label) == true, the 2nd dimension of " "If Attr(soft_label) == true, the 2nd dimension of "
"Input(X) and Input(Label) should be equal."); "Input(X) and Input(Label) should be equal.");
} else { } else {
PADDLE_ENFORCE_EQ(label_dims[1], 1, PADDLE_ENFORCE_EQ(label_dims[1], 1UL,
"If Attr(soft_label) == false, the 2nd dimension of " "If Attr(softLabel) == false, the 2nd dimension of "
"Input(Label) should be 1."); "Input(Label) should be 1.");
} }
...@@ -48,7 +49,8 @@ class CrossEntropyOp : public framework::OperatorWithKernel { ...@@ -48,7 +49,8 @@ class CrossEntropyOp : public framework::OperatorWithKernel {
} }
protected: protected:
// CrossEntropy's data type just determined by "X" // Explicitly set that data type of the output of the cross_entropy operator
// is determined by its input "X".
framework::DataType IndicateDataType( framework::DataType IndicateDataType(
const framework::ExecutionContext& ctx) const override { const framework::ExecutionContext& ctx) const override {
return framework::ToDataType(ctx.Input<Tensor>("X")->type()); return framework::ToDataType(ctx.Input<Tensor>("X")->type());
......
/* 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/operators/linear_chain_crf_op.h"
namespace paddle {
namespace operators {
class LinearChainCRFOpMaker : public framework::OpProtoAndCheckerMaker {
public:
LinearChainCRFOpMaker(framework::OpProto* proto,
framework::OpAttrChecker* op_checker)
: OpProtoAndCheckerMaker(proto, op_checker) {
AddInput(
"Emission",
"(LoDTensor, default: LoDTensor<float>). "
"The unscaled emission weight matrix for the linear chain CRF. "
"This input is a LoDTensor with shape [N x D] where N is the size of "
"the mini-batch and D is the total tag number.");
AddInput(
"Transition",
"(Tensor, default: Tensor<float>). A Tensor with shape [(D + 2) x D]. "
"The learnable parameter for the linear_chain_crf operator. "
"See more details in the operator's comments.");
AddInput(
"Label",
"(LoDTensor, default: LoDTensor<int>). The ground truth which is a 2-D "
"LoDTensor with shape [N x 1], where N is the total element number in "
"a mini-batch.");
AddOutput(
"Alpha",
"Tensor, default: Tensor<float>. The forward vectors for the entire "
"batch. A two dimensional tensor with shape [N x D], "
"denoted as \f$\alpha\f$. \f$\alpha$\f is a memo table used to "
"calculate the normalization factor in CRF. \f$\alpha[k, v]$\f stores "
"the unnormalized probabilites of all possible unfinished sequences of "
"tags that end at position \f$k$\f with tag \f$v$\f. For each \f$k$\f, "
"\f$\alpha[k, v]$\f is a vector of length \f$D$\f with a component for "
"each tag value \f$v$\f. This vector is called a forward vecotr and "
"will also be used in backward computations.")
.AsIntermediate();
AddOutput("EmissionExps",
"The exponentials of Input(Emission). This is an intermediate "
"computational result in forward computation, and will be reused "
"in backward computation.")
.AsIntermediate();
AddOutput("TransitionExps",
"The exponentials of Input(Transition). This is an intermediate "
"computational result in forward computation, and will be reused "
"in backward computation.")
.AsIntermediate();
AddOutput(
"LogLikelihood",
"(Tensor, default: Tensor<float>). The logarithm of the conditional "
"likelihood of each training sample in a mini-batch. This is a 2-D "
"tensor with shape [S x 1], where S is the sequence number in a "
"mini-batch. Note: S is equal to the sequence number in a mini-batch. "
"The output is no longer a LoDTensor.");
AddComment(R"DOC(
Conditional Random Field defines an undirected probabilistic graph with nodes
denoting random variables and edges denoting dependencies between these
variables. CRF learns the conditional probability \f$P(Y|X)\f$, where
\f$X = (x_1, x_2, ... , x_n)\f$ are structured inputs and
\f$Y = (y_1, y_2, ... , y_n)\f$ are labels for the inputs.
Linear chain CRF is a special case of CRF that is useful for sequence labeling
task. Sequence labeling tasks do not assume a lot of conditional
independences among inputs. The only constraint they impose is that the input
and output must be linear sequences. Thus, the graph of such a CRF is a simple
chain or a line, which results in the linear chain CRF.
This operator implements the Forward-Backward algorithm for the linear chain
CRF. Please see http://www.cs.columbia.edu/~mcollins/fb.pdf and
http://cseweb.ucsd.edu/~elkan/250Bwinter2012/loglinearCRFs.pdf for reference.
Equation:
- Denote Input(Emission) to this operator as \f$x\f$ here.
- The first D values of Input(Transition) to this operator are for starting
weights, denoted as \f$a\f$ here.
- The next D values of Input(Transition) of this operator are for ending
weights, denoted as \f$b\f$ here.
- The remaning values of Input(Transition) are for transition weights,
denoted as \f$w\f$ here.
- Denote Input(Label) as \f$s\f$ here.
The probability of a sequence \f$s\f$ of length \f$L\f$ is defined as:
\f$P(s) = (1/Z) exp(a_{s_1} + b_{s_L}
+ \sum_{l=1}^L x_{s_l}
+ \sum_{l=2}^L w_{s_{l-1},s_l})\f$
where \f$Z\f$ is a normalization value so that the sum of \f$P(s)\f$ over
all possible sequences is \f$1\f$, and \f$x\f$ is the emission feature weight
to the linear chain CRF.
Finaly, the linear chain CRF operator outputs the logarithm of the conditional
likelihood of each training sample in a mini-batch.
NOTE:
1. The feature function for a CRF is made up of the emission features and the
transition features. The emission feature weights are NOT computed in
this operator. They MUST be computed first before this operator is called.
2. Because this operator performs global normalization over all possible
sequences internally, it expects UNSCALED emission feature weights.
Please do not call this op with the emission feature being output of any
nonlinear activation.
3. The 2nd dimension of Input(Emission) MUST be equal to the tag number.
)DOC");
}
};
class LinearChainCRFOp : public framework::OperatorWithKernel {
public:
using framework::OperatorWithKernel::OperatorWithKernel;
void InferShape(framework::InferShapeContext* ctx) const override {
PADDLE_ENFORCE(ctx->HasInput("Emission"),
"Input(Emission) should be not null.");
PADDLE_ENFORCE(ctx->HasInput("Transition"),
"Input(Transition) should be not null.");
PADDLE_ENFORCE(ctx->HasInput("Label"), "Input(Label) should be not null.");
PADDLE_ENFORCE(ctx->HasOutput("Alpha"),
"Output(Alpha) should be not null.");
PADDLE_ENFORCE(ctx->HasOutput("EmissionExps"),
"Output(EmissionExps) should be not null.");
PADDLE_ENFORCE(ctx->HasOutput("TransitionExps"),
"Output(TransitionExps) should be not null.");
PADDLE_ENFORCE(ctx->HasOutput("LogLikelihood"),
"Output(LogLikelihood) should be not null.");
auto emission_dims = ctx->GetInputDim("Emission");
PADDLE_ENFORCE_EQ(emission_dims.size(), 2UL,
"The Input(Emission) should be a 2-D tensor.");
PADDLE_ENFORCE(emission_dims[0], "An empty mini-batch is not allowed.");
auto transition_dims = ctx->GetInputDim("Transition");
PADDLE_ENFORCE_EQ(transition_dims.size(), 2UL,
"The Input(Transition) should be a 2-D tensor.");
PADDLE_ENFORCE_EQ(
transition_dims[0] - 2, transition_dims[1],
"An invalid dimension for the Input(Transition), which should "
"be a 2-D tensor with shape [(D + 2) x D].");
PADDLE_ENFORCE_EQ(
emission_dims[1], transition_dims[1],
"The 2nd dimension of the Input(Emission) and the Input(Transition) "
"should be equal to the tag number.");
auto label_dims = ctx->GetInputDim("Label");
PADDLE_ENFORCE(label_dims.size() == 2UL && label_dims[1] == 1UL,
"The Input(Label) should be a 2-D tensor with the 2nd "
"dimensions fixed to 1.");
PADDLE_ENFORCE_EQ(
emission_dims[0], label_dims[0],
"The height of Input(Emission) and the height of Input(Label) "
"should be the same.");
ctx->SetOutputDim("Alpha", emission_dims);
ctx->SetOutputDim("EmissionExps", emission_dims);
ctx->SetOutputDim("TransitionExps", transition_dims);
// TODO(caoying) This is tricky. The 1st dimension of Output(LogLikelihood)
// is the sequence number in a mini-batch. The dimension set here should be
// resized to its correct size in the function Compute. Fix this once we can
// get LoD information in the InferShape interface.
ctx->SetOutputDim("LogLikelihood", {emission_dims[0], 1});
}
protected:
// Explicitly set that the data type of output of the linear_chain_crf
// operator is determined by its input "Emission".
framework::DataType IndicateDataType(
const framework::ExecutionContext& ctx) const override {
return framework::ToDataType(ctx.Input<LoDTensor>("Emission")->type());
}
};
class LinearChainCRFGradOp : public framework::OperatorWithKernel {
public:
using framework::OperatorWithKernel::OperatorWithKernel;
void InferShape(framework::InferShapeContext* ctx) const override {
PADDLE_ENFORCE(ctx->HasInput("EmissionExps"),
"Input(EmissionExps) should be not null.");
PADDLE_ENFORCE(ctx->HasInput("TransitionExps"),
"Input(TransitionExps) should be not null.");
PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("LogLikelihood")),
"Input(LogLikelihood@GRAD) shoudl be not null.");
auto emission_exps_dims = ctx->GetInputDim("EmissionExps");
PADDLE_ENFORCE_EQ(emission_exps_dims.size(), 2UL,
"The Input(EmissionExps) should be a 2-D tensor.");
PADDLE_ENFORCE(emission_exps_dims[0],
"An empty mini-batch is not allowed.");
auto transition_exps_dims = ctx->GetInputDim("TransitionExps");
PADDLE_ENFORCE_EQ(transition_exps_dims.size(), 2UL,
"The Input(TransitionExps) should be a 2-D tensor.");
PADDLE_ENFORCE_EQ(
transition_exps_dims[0] - 2, transition_exps_dims[1],
"An invalid dimension for the Input(TransitionExps), which should "
"be a 2-D tensor with shape [(D + 2) x D].");
PADDLE_ENFORCE_EQ(
emission_exps_dims[1], transition_exps_dims[1],
"The 2nd dimension of the Input(EmissionExps) and the "
"Input(TransitionExps) should be equal to the tag number.");
auto label_dims = ctx->GetInputDim("Label");
PADDLE_ENFORCE(label_dims.size() == 2UL && label_dims[1] == 1UL,
"The Input(Label) should be a 2-D tensor with the 2nd "
"dimensions fixed to 1.");
PADDLE_ENFORCE_EQ(
emission_exps_dims[0], label_dims[0],
"The height of Input(EmissionExps) and the height of Input(Label) "
"should be the same.");
if (ctx->HasOutput(framework::GradVarName("Emission"))) {
ctx->SetOutputDim(framework::GradVarName("Emission"), emission_exps_dims);
}
if (ctx->HasOutput(framework::GradVarName("Transition"))) {
ctx->SetOutputDim(framework::GradVarName("Transition"),
transition_exps_dims);
}
}
protected:
// Explicitly set that the data type of output of the linear_chain_crf_grad
// operator is determined by its input: gradients of LogLikelihood.
framework::DataType IndicateDataType(
const framework::ExecutionContext& ctx) const override {
return framework::ToDataType(
ctx.Input<LoDTensor>(framework::GradVarName("LogLikelihood"))->type());
}
};
} // namespace operators
} // namespace paddle
namespace ops = paddle::operators;
REGISTER_OP(linear_chain_crf, ops::LinearChainCRFOp, ops::LinearChainCRFOpMaker,
linear_chain_crf_grad, ops::LinearChainCRFGradOp);
REGISTER_OP_CPU_KERNEL(
linear_chain_crf,
ops::LinearChainCRFOpKernel<paddle::platform::CPUPlace, float>,
ops::LinearChainCRFOpKernel<paddle::platform::CPUPlace, double>);
REGISTER_OP_CPU_KERNEL(
linear_chain_crf_grad,
ops::LinearChainCRFGradOpKernel<paddle::platform::CPUPlace, float>,
ops::LinearChainCRFGradOpKernel<paddle::platform::CPUPlace, double>);
/* 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/operators/linear_chain_crf_op.h"
namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL(
linear_chain_crf,
ops::LinearChainCRFOpKernel<paddle::platform::GPUPlace, float>,
ops::LinearChainCRFOpKernel<paddle::platform::GPUPlace, double>);
REGISTER_OP_GPU_KERNEL(
linear_chain_crf_grad,
ops::LinearChainCRFGradOpKernel<paddle::platform::GPUPlace, float>,
ops::LinearChainCRFGradOpKernel<paddle::platform::GPUPlace, double>);
此差异已折叠。
...@@ -21,7 +21,6 @@ class LSTMOp : public framework::OperatorWithKernel { ...@@ -21,7 +21,6 @@ class LSTMOp : public framework::OperatorWithKernel {
public: public:
using framework::OperatorWithKernel::OperatorWithKernel; using framework::OperatorWithKernel::OperatorWithKernel;
protected:
void InferShape(framework::InferShapeContext* ctx) const override { void InferShape(framework::InferShapeContext* ctx) const override {
PADDLE_ENFORCE(ctx->HasInput("Input"), PADDLE_ENFORCE(ctx->HasInput("Input"),
"Input(Input) of LSTM should not be null."); "Input(Input) of LSTM should not be null.");
...@@ -29,9 +28,13 @@ class LSTMOp : public framework::OperatorWithKernel { ...@@ -29,9 +28,13 @@ class LSTMOp : public framework::OperatorWithKernel {
"Output(Hidden) of LSTM should not be null."); "Output(Hidden) of LSTM should not be null.");
PADDLE_ENFORCE(ctx->HasOutput("Cell"), PADDLE_ENFORCE(ctx->HasOutput("Cell"),
"Output(Cell) of LSTM should not be null."); "Output(Cell) of LSTM should not be null.");
PADDLE_ENFORCE(ctx->HasOutput("BatchGate"),
"Output(BatchGate) of LSTM should not be null.");
PADDLE_ENFORCE(ctx->HasOutput("BatchCellPreAct"),
"Output(BatchGate) of LSTM should not be null.");
auto x_dims = ctx->GetInputDim("Input"); auto in_dims = ctx->GetInputDim("Input");
PADDLE_ENFORCE_EQ(x_dims.size(), 2, "Input(X)'s rank must be 2."); PADDLE_ENFORCE_EQ(in_dims.size(), 2, "Input(X)'s rank must be 2.");
if (ctx->HasInput("H0")) { if (ctx->HasInput("H0")) {
PADDLE_ENFORCE(ctx->HasInput("C0"), PADDLE_ENFORCE(ctx->HasInput("C0"),
...@@ -44,7 +47,7 @@ class LSTMOp : public framework::OperatorWithKernel { ...@@ -44,7 +47,7 @@ class LSTMOp : public framework::OperatorWithKernel {
"should be the same."); "should be the same.");
} }
int frame_size = x_dims[1] / 4; int frame_size = in_dims[1] / 4;
auto w_dims = ctx->GetInputDim("Weight"); auto w_dims = ctx->GetInputDim("Weight");
PADDLE_ENFORCE_EQ(w_dims.size(), 2, PADDLE_ENFORCE_EQ(w_dims.size(), 2,
"The rank of Input(Weight) should be 2."); "The rank of Input(Weight) should be 2.");
...@@ -71,12 +74,21 @@ class LSTMOp : public framework::OperatorWithKernel { ...@@ -71,12 +74,21 @@ class LSTMOp : public framework::OperatorWithKernel {
"4 * %d if disable peepholes connection", "4 * %d if disable peepholes connection",
frame_size); frame_size);
} }
ctx->SetOutputDim("Hidden", {x_dims[0], frame_size}); framework::DDim out_dims({in_dims[0], frame_size});
ctx->SetOutputDim("Cell", {x_dims[0], frame_size}); ctx->SetOutputDim("Hidden", out_dims);
ctx->SetOutputDim("BatchGate", x_dims); ctx->SetOutputDim("Cell", out_dims);
ctx->SetOutputDim("BatchGate", in_dims);
ctx->SetOutputDim("BatchCellPreAct", out_dims);
ctx->ShareLoD("Input", "Hidden"); ctx->ShareLoD("Input", "Hidden");
ctx->ShareLoD("Input", "Cell"); ctx->ShareLoD("Input", "Cell");
} }
protected:
framework::DataType IndicateDataType(
const framework::ExecutionContext& ctx) const override {
return framework::ToDataType(
ctx.Input<framework::LoDTensor>("Input")->type());
}
}; };
class LSTMOpMaker : public framework::OpProtoAndCheckerMaker { class LSTMOpMaker : public framework::OpProtoAndCheckerMaker {
...@@ -86,16 +98,18 @@ class LSTMOpMaker : public framework::OpProtoAndCheckerMaker { ...@@ -86,16 +98,18 @@ class LSTMOpMaker : public framework::OpProtoAndCheckerMaker {
AddInput("Input", AddInput("Input",
"(LoDTensor) the first input is a LodTensor, which support " "(LoDTensor) the first input is a LodTensor, which support "
"variable-time length input sequence. The underlying tensor in " "variable-time length input sequence. The underlying tensor in "
"this LoDTensor is a matrix with shape (T X 4D), where, T is the " "this LoDTensor is a matrix with shape (T X 4D), where T is the "
"total time steps in this mini-batch, D is the hidden size."); "total time steps in this mini-batch, D is the hidden size.");
AddInput("H0", AddInput("H0",
"(Tensor, optional) the initial hidden state is an optional " "(Tensor, optional) the initial hidden state is an optional "
"input. This is a tensor with shape (N x D), where N is the " "input. This is a tensor with shape (N x D), where N is the "
"batch size, D is the hidden size."); "batch size, D is the hidden size.")
.AsDispensable();
AddInput("C0", AddInput("C0",
"(Tensor, optional) the initial cell state is an optional " "(Tensor, optional) the initial cell state is an optional "
"input. This is a tensor with shape (N x D), where N is the " "input. This is a tensor with shape (N x D), where N is the "
"batch size. `H0` and `C0` can be NULL but only at the same time"); "batch size. `H0` and `C0` can be NULL but only at the same time")
.AsDispensable();
AddInput("Weight", AddInput("Weight",
"(Tensor) the learnable hidden-hidden weights." "(Tensor) the learnable hidden-hidden weights."
" - The shape is (D x 4D), where D is the hidden size. " " - The shape is (D x 4D), where D is the hidden size. "
...@@ -109,22 +123,27 @@ class LSTMOpMaker : public framework::OpProtoAndCheckerMaker { ...@@ -109,22 +123,27 @@ class LSTMOpMaker : public framework::OpProtoAndCheckerMaker {
" - Bias = {b_c, b_i, b_f, b_o}." " - Bias = {b_c, b_i, b_f, b_o}."
"2. `usePeepholes = True` " "2. `usePeepholes = True` "
" - The shape is (1 x 7D). " " - The shape is (1 x 7D). "
" - Bias = {b_c, b_i, b_f, b_o, W_ic, W_fc, W_oc}."); " - Bias = {b_c, b_i, b_f, b_o, W_ic, W_fc, W_oc}.")
.AsDispensable();
AddOutput("Hidden",
"(LoDTensor) the hidden state of LSTM operator. "
"The shape is (T x D), and lod is the same with the `Input`.");
AddOutput("Cell",
"(LoDTensor) the cell state of LSTM operator. "
"The shape is (T x D), and lod is the same with the `Input`.");
AddOutput("BatchGate", AddOutput("BatchGate",
"(LoDTensor) This LoDTensor contains input gate, forget gate " "(LoDTensor) This LoDTensor contains input gate, forget gate "
"and output gate after the nonlinear computation. This " "and output gate after the nonlinear computation. This "
"LoDTensor has the same shape with the reorganized input, which " "LoDTensor has the same shape with the reorganized input, which "
"was also be called batch input. The LoD size is 2. The first " "is also be called batch input. The LoD size is 2. The first "
"LoD is the batch offsets and the second LoD contains the " "LoD is the batch offsets and the second LoD contains the "
"indexes, which denote the position of reorganized sequence " "indexes, which denote the position of reorganized sequence "
"in the raw input.") "in the raw input.")
.AsIntermediate(); .AsIntermediate();
AddOutput("Hidden", AddOutput("BatchCellPreAct",
"(LoDTensor) the hidden state lod tensor of LSTM operator. " "(LoDTensor) This LoDTensor is got in the forward and used "
"The shape and lod is the same with the `Input`."); "in the backward.")
AddOutput("Cell", .AsIntermediate();
"(LoDTensor) the cell state lod tensor of LSTM operator. "
"The shape and lod is the same with the `Input`.");
AddAttr<bool>("usePeepholes", AddAttr<bool>("usePeepholes",
"(bool, defalut: True) " "(bool, defalut: True) "
"whether to enable diagonal/peephole connections.") "whether to enable diagonal/peephole connections.")
...@@ -202,15 +221,37 @@ class LSTMGradOp : public framework::OperatorWithKernel { ...@@ -202,15 +221,37 @@ class LSTMGradOp : public framework::OperatorWithKernel {
public: public:
using framework::OperatorWithKernel::OperatorWithKernel; using framework::OperatorWithKernel::OperatorWithKernel;
protected:
void InferShape(framework::InferShapeContext* ctx) const override { void InferShape(framework::InferShapeContext* ctx) const override {
PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Hidden")), PADDLE_ENFORCE(ctx->HasInput("Input"),
"Input(Hidden@GRAD) should not be null"); "Input(Input) of LSTM should not be null.");
PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Cell")), PADDLE_ENFORCE(ctx->HasInput("Hidden"),
"Input(Cell@GRAD) should not be null"); "Input(Hidden) of LSTM should not be null.");
ctx->SetOutputDim(framework::GradVarName("Weight"), PADDLE_ENFORCE(ctx->HasInput("Cell"),
ctx->GetInputDim("Weight")); "Input(Cell) of LSTM should not be null.");
ctx->SetOutputDim(framework::GradVarName("Bias"), ctx->GetInputDim("Bias"));
PADDLE_ENFORCE(ctx->HasInput("BatchGate"),
"Input(BatchGate) of LSTM should not be null.");
PADDLE_ENFORCE(ctx->HasInput("BatchCellPreAct"),
"Input(BatchGate) of LSTM should not be null.");
auto in_g_name = framework::GradVarName("Input");
if (ctx->HasOutput(in_g_name))
ctx->SetOutputDim(in_g_name, ctx->GetInputDim("Input"));
auto w_g_name = framework::GradVarName("Weight");
if (ctx->HasOutput(w_g_name))
ctx->SetOutputDim(w_g_name, ctx->GetInputDim("Weight"));
auto b_g_name = framework::GradVarName("Bias");
if (ctx->HasOutput(b_g_name))
ctx->SetOutputDim(b_g_name, ctx->GetInputDim("Bias"));
}
protected:
framework::DataType IndicateDataType(
const framework::ExecutionContext& ctx) const override {
return framework::ToDataType(
ctx.Input<framework::LoDTensor>("Input")->type());
} }
}; };
......
...@@ -21,8 +21,9 @@ limitations under the License. */ ...@@ -21,8 +21,9 @@ limitations under the License. */
namespace paddle { namespace paddle {
namespace operators { namespace operators {
using framework::LoDTensor; using LoDTensor = framework::LoDTensor;
using framework::Tensor; using Tensor = framework::Tensor;
template <typename T, int MajorType = Eigen::RowMajor, template <typename T, int MajorType = Eigen::RowMajor,
typename IndexType = Eigen::DenseIndex> typename IndexType = Eigen::DenseIndex>
using EigenMatrix = framework::EigenMatrix<T, MajorType, IndexType>; using EigenMatrix = framework::EigenMatrix<T, MajorType, IndexType>;
...@@ -31,15 +32,15 @@ template <typename Place, typename T> ...@@ -31,15 +32,15 @@ template <typename Place, typename T>
class LSTMKernel : public framework::OpKernel<T> { class LSTMKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& ctx) const override { void Compute(const framework::ExecutionContext& ctx) const override {
auto* input = ctx.Input<framework::LoDTensor>("Input"); auto* input = ctx.Input<LoDTensor>("Input");
auto* weight = ctx.Input<framework::Tensor>("Weight"); auto* weight = ctx.Input<Tensor>("Weight");
auto* bias = ctx.Input<framework::Tensor>("Bias"); auto* bias = ctx.Input<Tensor>("Bias");
auto* batch_gate = ctx.Output<framework::LoDTensor>("BatchGate"); auto* batch_gate = ctx.Output<LoDTensor>("BatchGate");
batch_gate->mutable_data<T>(ctx.GetPlace()); batch_gate->mutable_data<T>(ctx.GetPlace());
auto* hidden_out = ctx.Output<framework::LoDTensor>("Hidden"); auto* hidden_out = ctx.Output<LoDTensor>("Hidden");
hidden_out->mutable_data<T>(ctx.GetPlace()); hidden_out->mutable_data<T>(ctx.GetPlace());
auto* cell_out = ctx.Output<framework::LoDTensor>("Cell"); auto* cell_out = ctx.Output<LoDTensor>("Cell");
cell_out->mutable_data<T>(ctx.GetPlace()); cell_out->mutable_data<T>(ctx.GetPlace());
// Now the function ShareLoD in InferShape is not implemented. // Now the function ShareLoD in InferShape is not implemented.
...@@ -49,7 +50,8 @@ class LSTMKernel : public framework::OpKernel<T> { ...@@ -49,7 +50,8 @@ class LSTMKernel : public framework::OpKernel<T> {
bool is_reverse = ctx.Attr<bool>("isReverse"); bool is_reverse = ctx.Attr<bool>("isReverse");
math::LoDTensor2BatchFunctor<Place, T> to_batch; math::LoDTensor2BatchFunctor<Place, T> to_batch;
to_batch(ctx.device_context(), *input, *batch_gate, is_reverse); auto& device_ctx = ctx.device_context();
to_batch(device_ctx, *input, *batch_gate, true, is_reverse);
auto in_dims = input->dims(); auto in_dims = input->dims();
int frame_size = static_cast<int>(in_dims[1] / 4); int frame_size = static_cast<int>(in_dims[1] / 4);
...@@ -69,17 +71,26 @@ class LSTMKernel : public framework::OpKernel<T> { ...@@ -69,17 +71,26 @@ class LSTMKernel : public framework::OpKernel<T> {
} }
math::LstmMetaValue<T> lstm_value; math::LstmMetaValue<T> lstm_value;
T* bias_data = const_cast<T*>(bias->data<T>()); if (bias) {
// the code style in LstmMetaValue will be updated later. T* bias_data = const_cast<T*>(bias->data<T>());
lstm_value.checkIg = bias_data + 4 * frame_size; // the code style in LstmMetaValue will be updated later.
lstm_value.checkFg = lstm_value.checkIg + frame_size;
lstm_value.checkOg = lstm_value.checkFg + frame_size; lstm_value.checkIg = bias_data + 4 * frame_size;
lstm_value.checkFg = lstm_value.checkIg + frame_size;
lstm_value.checkOg = lstm_value.checkFg + frame_size;
} else {
lstm_value.checkIg = nullptr;
lstm_value.checkFg = nullptr;
lstm_value.checkOg = nullptr;
}
lstm_value.prevStateValue = nullptr; lstm_value.prevStateValue = nullptr;
framework::LoDTensor batch_out, batch_cell, batch_cell_pre_act; // Use the local variable as here.
batch_out.mutable_data<T>(dims, ctx.GetPlace()); LoDTensor batch_hidden, batch_cell;
auto* batch_cell_pre_act = ctx.Output<LoDTensor>("BatchCellPreAct");
batch_hidden.mutable_data<T>(dims, ctx.GetPlace());
batch_cell.mutable_data<T>(dims, ctx.GetPlace()); batch_cell.mutable_data<T>(dims, ctx.GetPlace());
batch_cell_pre_act.mutable_data<T>(dims, ctx.GetPlace()); batch_cell_pre_act->mutable_data<T>(dims, ctx.GetPlace());
auto batch_starts = batch_gate->lod()[0]; auto batch_starts = batch_gate->lod()[0];
size_t num_batch = batch_starts.size() - 1; size_t num_batch = batch_starts.size() - 1;
...@@ -92,18 +103,18 @@ class LSTMKernel : public framework::OpKernel<T> { ...@@ -92,18 +103,18 @@ class LSTMKernel : public framework::OpKernel<T> {
int bend = static_cast<int>(batch_starts[n + 1]); int bend = static_cast<int>(batch_starts[n + 1]);
Tensor gate_t = batch_gate->Slice(bstart, bend); Tensor gate_t = batch_gate->Slice(bstart, bend);
Tensor out_t = batch_out.Slice(bstart, bend); Tensor out_t = batch_hidden.Slice(bstart, bend);
Tensor cell_t = batch_cell.Slice(bstart, bend); Tensor cell_t = batch_cell.Slice(bstart, bend);
Tensor cell_pre_act_t = batch_cell_pre_act.Slice(bstart, bend); Tensor cell_pre_act_t = batch_cell_pre_act->Slice(bstart, bend);
int cur_batch_size = bend - bstart; int cur_batch_size = bend - bstart;
if (n != 0) { if (n != 0) {
int pre_h_start = static_cast<int>(batch_starts[n - 1]); int pre_h_start = static_cast<int>(batch_starts[n - 1]);
int pre_h_end = pre_h_start + cur_batch_size; int pre_h_end = pre_h_start + cur_batch_size;
auto pre_hidden_t = batch_out.Slice(pre_h_start, pre_h_end); auto pre_hidden_t = batch_hidden.Slice(pre_h_start, pre_h_end);
math::matmul<Place, T>(ctx.device_context(), pre_hidden_t, false, math::matmul<Place, T>(device_ctx, pre_hidden_t, false, *weight, false,
*weight, false, static_cast<T>(1.0), &gate_t, static_cast<T>(1.0), &gate_t,
static_cast<T>(1.0)); static_cast<T>(1.0));
} }
// else if : FIXME support the initial hidden and cell // else if : FIXME support the initial hidden and cell
...@@ -112,27 +123,186 @@ class LSTMKernel : public framework::OpKernel<T> { ...@@ -112,27 +123,186 @@ class LSTMKernel : public framework::OpKernel<T> {
lstm_value.outputValue = out_t.data<T>(); lstm_value.outputValue = out_t.data<T>();
lstm_value.stateValue = cell_t.data<T>(); lstm_value.stateValue = cell_t.data<T>();
lstm_value.stateActiveValue = cell_pre_act_t.data<T>(); lstm_value.stateActiveValue = cell_pre_act_t.data<T>();
math::LstmUnitFunctor<Place, T>::compute(ctx.device_context(), lstm_value, math::LstmUnitFunctor<Place, T>::compute(device_ctx, lstm_value,
frame_size, cur_batch_size, frame_size, cur_batch_size,
gate_act, cell_act, cand_act); gate_act, cell_act, cand_act);
lstm_value.prevStateValue = lstm_value.stateValue; lstm_value.prevStateValue = lstm_value.stateValue;
} }
math::Batch2LoDTensorFunctor<Place, T> to_seq; math::Batch2LoDTensorFunctor<Place, T> to_seq;
batch_out.set_lod(batch_gate->lod()); batch_hidden.set_lod(batch_gate->lod());
// restore the output hidden in LoDTensor from the batch hidden // restore the output hidden in LoDTensor from the batch hidden
to_seq(ctx.device_context(), batch_out, *hidden_out); to_seq(device_ctx, batch_hidden, *hidden_out);
batch_cell.set_lod(batch_gate->lod()); batch_cell.set_lod(batch_gate->lod());
// restore the output cell state in LoDTensor from the batch cell // restore the output cell state in LoDTensor from the batch cell
to_seq(ctx.device_context(), batch_cell, *cell_out); to_seq(device_ctx, batch_cell, *cell_out);
} }
}; };
template <typename Place, typename T> template <typename Place, typename T>
class LSTMGradKernel : public framework::OpKernel<T> { class LSTMGradKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& ctx) const override {} void Compute(const framework::ExecutionContext& ctx) const override {
auto* input = ctx.Input<LoDTensor>("Input");
auto* weight = ctx.Input<Tensor>("Weight");
auto* bias = ctx.Input<Tensor>("Bias");
auto* hidden_out = ctx.Input<LoDTensor>("Hidden");
auto* cell_out = ctx.Input<LoDTensor>("Cell");
auto* batch_gate = ctx.Input<LoDTensor>("BatchGate");
auto* batch_cell_pre_act = ctx.Input<LoDTensor>("BatchCellPreAct");
auto* hidden_g = ctx.Input<LoDTensor>(framework::GradVarName("Hidden"));
auto* in_g = ctx.Output<LoDTensor>(framework::GradVarName("Input"));
auto* weight_g = ctx.Output<Tensor>(framework::GradVarName("Weight"));
auto* bias_g = ctx.Output<Tensor>(framework::GradVarName("Bias"));
auto& device_ctx = ctx.device_context();
math::SetConstant<Place, T> zero;
if (weight_g) {
weight_g->mutable_data<T>(ctx.GetPlace());
zero(device_ctx, weight_g, static_cast<T>(0.0));
}
auto in_dims = input->dims();
auto out_dims = hidden_g->dims();
int frame_size = static_cast<int>(in_dims[1] / 4);
PADDLE_ENFORCE_EQ(frame_size, out_dims[1]);
math::LstmMetaValue<T> lstm_value;
if (bias) {
T* bias_data = const_cast<T*>(bias->data<T>());
lstm_value.checkIg = bias_data + 4 * frame_size;
lstm_value.checkFg = lstm_value.checkIg + frame_size;
lstm_value.checkOg = lstm_value.checkFg + frame_size;
} else {
lstm_value.checkIg = nullptr;
lstm_value.checkFg = nullptr;
lstm_value.checkOg = nullptr;
}
math::LstmMetaGrad<T> lstm_grad;
if (bias && bias_g) {
T* bias_g_data = const_cast<T*>(bias_g->mutable_data<T>(ctx.GetPlace()));
zero(device_ctx, bias_g, static_cast<T>(0.0));
lstm_grad.checkIgGrad = bias_g_data + 4 * frame_size;
lstm_grad.checkFgGrad = lstm_grad.checkIgGrad + frame_size;
lstm_grad.checkOgGrad = lstm_grad.checkFgGrad + frame_size;
} else {
lstm_grad.checkIgGrad = nullptr;
lstm_grad.checkFgGrad = nullptr;
lstm_grad.checkOgGrad = nullptr;
}
math::LoDTensor2BatchFunctor<Place, T> to_batch;
// use the local variable as here.
LoDTensor batch_hidden;
batch_hidden.mutable_data<T>(out_dims, ctx.GetPlace());
batch_hidden.set_lod(batch_gate->lod());
to_batch(device_ctx, *hidden_out, batch_hidden, false);
LoDTensor batch_hidden_g;
batch_hidden_g.mutable_data<T>(out_dims, ctx.GetPlace());
batch_hidden_g.set_lod(batch_gate->lod());
to_batch(device_ctx, *hidden_g, batch_hidden_g, false);
LoDTensor batch_cell;
batch_cell.mutable_data<T>(out_dims, ctx.GetPlace());
batch_cell.set_lod(batch_gate->lod());
to_batch(device_ctx, *cell_out, batch_cell, false);
LoDTensor batch_cell_g;
batch_cell_g.mutable_data<T>(out_dims, ctx.GetPlace());
batch_cell_g.set_lod(batch_gate->lod());
// TODO(qingqing) support the case output cell has gradient.
// to_batch(device_ctx, *cell_g, batch_cell_g, false);
zero(device_ctx, &batch_cell_g, static_cast<T>(0.0));
LoDTensor batch_gate_g;
batch_gate_g.mutable_data<T>(batch_gate->dims(), ctx.GetPlace());
batch_gate_g.set_lod(batch_gate->lod());
auto gate_act = ctx.Attr<std::string>("gateActivation");
auto cell_act = ctx.Attr<std::string>("cellActivation");
auto cand_act = ctx.Attr<std::string>("candidateActivation");
auto batch_starts = batch_gate->lod()[0];
size_t num_batch = batch_starts.size() - 1;
for (int n = static_cast<int>(num_batch) - 1; n >= 0; n--) {
int bstart = static_cast<int>(batch_starts[n]);
int bend = static_cast<int>(batch_starts[n + 1]);
Tensor gate = batch_gate->Slice(bstart, bend);
Tensor cell = batch_cell.Slice(bstart, bend);
Tensor cell_pre_act = batch_cell_pre_act->Slice(bstart, bend);
lstm_value.gateValue = gate.data<T>();
lstm_value.stateValue = cell.data<T>();
lstm_value.stateActiveValue = cell_pre_act.data<T>();
Tensor out_g = batch_hidden_g.Slice(bstart, bend);
Tensor gate_g = batch_gate_g.Slice(bstart, bend);
Tensor cell_g = batch_cell_g.Slice(bstart, bend);
lstm_grad.stateGrad = cell_g.data<T>();
lstm_grad.gateGrad = gate_g.data<T>();
lstm_grad.outputGrad = out_g.data<T>();
if (n) {
int bstart_pre = static_cast<int>(batch_starts[n - 1]);
Tensor cell_pre = batch_cell.Slice(bstart_pre, bstart);
Tensor cell_pre_g = batch_cell_g.Slice(bstart_pre, bstart);
lstm_value.prevStateValue = cell_pre.data<T>();
lstm_grad.prevStateGrad = cell_pre_g.data<T>();
} else {
lstm_value.prevStateValue = nullptr;
lstm_grad.prevStateGrad = nullptr;
}
int cur_batch_size = bend - bstart;
math::LstmUnitGradFunctor<Place, T>::compute(
device_ctx, lstm_value, lstm_grad, frame_size, cur_batch_size,
gate_act, cell_act, cand_act);
if (n != 0) {
int pre_h_start = static_cast<int>(batch_starts[n - 1]);
int pre_h_end = pre_h_start + cur_batch_size;
auto pre_hidden_g = batch_hidden_g.Slice(pre_h_start, pre_h_end);
math::matmul<Place, T>(device_ctx, gate_g, false, *weight, true,
static_cast<T>(1.0), &pre_hidden_g,
static_cast<T>(1.0));
if (weight_g) {
/* backward weight */
auto pre_hidden = batch_hidden.Slice(pre_h_start, pre_h_end);
math::matmul<Place, T>(device_ctx, pre_hidden, true, gate_g, false,
static_cast<T>(1.0), weight_g,
static_cast<T>(1.0));
}
}
}
math::Batch2LoDTensorFunctor<Place, T> to_seq;
if (in_g) {
/* backward data */
in_g->mutable_data<T>(ctx.GetPlace());
to_seq(device_ctx, batch_gate_g, *in_g);
}
if (bias && bias_g) {
/* backward bias */
int m = static_cast<int>(batch_gate_g.dims()[0]);
int n = static_cast<int>(batch_gate_g.dims()[1]);
Tensor ones;
ones.mutable_data<T>({m}, ctx.GetPlace());
math::SetConstant<Place, T> set;
set(device_ctx, &ones, static_cast<T>(1.0));
math::gemv<Place, T>(device_ctx, true, m, n, 1., batch_gate_g.data<T>(),
ones.data<T>(), 0., bias_g->data<T>());
}
}
}; };
} // namespace operators } // namespace operators
......
...@@ -26,10 +26,7 @@ namespace detail { ...@@ -26,10 +26,7 @@ namespace detail {
template <class T, class Op> template <class T, class Op>
void naive_lstm_forward_one_sequence(Op op, LstmMetaValue<T> value, void naive_lstm_forward_one_sequence(Op op, LstmMetaValue<T> value,
int frameSize, int frameSize) {
activation_mode_t active_node,
activation_mode_t active_gate,
activation_mode_t active_state) {
T rValueIn; T rValueIn;
T rValueIg; T rValueIg;
T rValueFg; T rValueFg;
...@@ -60,10 +57,8 @@ void naive_lstm_forward_one_sequence(Op op, LstmMetaValue<T> value, ...@@ -60,10 +57,8 @@ void naive_lstm_forward_one_sequence(Op op, LstmMetaValue<T> value,
rPrevState = value.prevStateValue[i]; rPrevState = value.prevStateValue[i];
} }
hppl::cpu::ForwardAct<T> act;
op(rValueIn, rValueIg, rValueFg, rValueOg, rPrevState, rState, rStateAtv, op(rValueIn, rValueIg, rValueFg, rValueOg, rPrevState, rState, rStateAtv,
rOut, rCheckI, rCheckF, rCheckO, act(active_node), act(active_gate), rOut, rCheckI, rCheckF, rCheckO);
act(active_state));
valueIn[i] = rValueIn; valueIn[i] = rValueIn;
valueIg[i] = rValueIg; valueIg[i] = rValueIg;
...@@ -77,10 +72,7 @@ void naive_lstm_forward_one_sequence(Op op, LstmMetaValue<T> value, ...@@ -77,10 +72,7 @@ void naive_lstm_forward_one_sequence(Op op, LstmMetaValue<T> value,
template <class T, class Op> template <class T, class Op>
void naive_lstm_backward_one_sequence(Op op, LstmMetaValue<T> value, void naive_lstm_backward_one_sequence(Op op, LstmMetaValue<T> value,
LstmMetaGrad<T> grad, int frameSize, LstmMetaGrad<T> grad, int frameSize) {
activation_mode_t active_node,
activation_mode_t active_gate,
activation_mode_t active_state) {
T rValueIn; T rValueIn;
T rValueIg; T rValueIg;
T rValueFg; T rValueFg;
...@@ -127,11 +119,10 @@ void naive_lstm_backward_one_sequence(Op op, LstmMetaValue<T> value, ...@@ -127,11 +119,10 @@ void naive_lstm_backward_one_sequence(Op op, LstmMetaValue<T> value,
rPrevState = value.prevStateValue[i]; rPrevState = value.prevStateValue[i];
} }
hppl::cpu::BackwardAct<T> act;
op(rValueIn, rValueIg, rValueFg, rValueOg, rGradIn, rGradIg, rGradFg, op(rValueIn, rValueIg, rValueFg, rValueOg, rGradIn, rGradIg, rGradFg,
rGradOg, rPrevState, rPrevStateGrad, rState, rStateGrad, rStateAtv, rGradOg, rPrevState, rPrevStateGrad, rState, rStateGrad, rStateAtv,
rOutputGrad, rCheckI, rCheckF, rCheckO, rCheckIGrad, rCheckFGrad, rOutputGrad, rCheckI, rCheckF, rCheckO, rCheckIGrad, rCheckFGrad,
rCheckOGrad, act(active_node), act(active_gate), act(active_state)); rCheckOGrad);
gradIn[i] = rGradIn; gradIn[i] = rGradIn;
gradIg[i] = rGradIg; gradIg[i] = rGradIg;
...@@ -283,8 +274,7 @@ void cpu_lstm_forward(Op op, LstmMetaValue<T> value, int frameSize, ...@@ -283,8 +274,7 @@ void cpu_lstm_forward(Op op, LstmMetaValue<T> value, int frameSize,
avx_lstm_forward_one_sequence<T>(op, value, frameSize, active_node, avx_lstm_forward_one_sequence<T>(op, value, frameSize, active_node,
active_gate, active_state); active_gate, active_state);
} else { } else {
naive_lstm_forward_one_sequence<T>(op, value, frameSize, active_node, naive_lstm_forward_one_sequence<T>(op, value, frameSize);
active_gate, active_state);
} }
} }
...@@ -297,8 +287,7 @@ void cpu_lstm_backward(Op op, LstmMetaValue<T> value, LstmMetaGrad<T> grad, ...@@ -297,8 +287,7 @@ void cpu_lstm_backward(Op op, LstmMetaValue<T> value, LstmMetaGrad<T> grad,
avx_lstm_backward_one_sequence<T>(op, value, grad, frameSize, active_node, avx_lstm_backward_one_sequence<T>(op, value, grad, frameSize, active_node,
active_gate, active_state); active_gate, active_state);
} else { } else {
naive_lstm_backward_one_sequence<T>(op, value, grad, frameSize, active_node, naive_lstm_backward_one_sequence<T>(op, value, grad, frameSize);
active_gate, active_state);
} }
} }
......
...@@ -32,9 +32,7 @@ namespace detail { ...@@ -32,9 +32,7 @@ namespace detail {
*/ */
template <class T, class Op, bool isBatch> template <class T, class Op, bool isBatch>
__global__ void KeLstmForward(Op op, LstmMetaValue<T> value, int frameSize, __global__ void KeLstmForward(Op op, LstmMetaValue<T> value, int frameSize,
int batchSize, activation_mode_t active_node, int batchSize) {
activation_mode_t active_gate,
activation_mode_t active_state) {
const int frameIdx = blockIdx.x * blockDim.x + threadIdx.x; const int frameIdx = blockIdx.x * blockDim.x + threadIdx.x;
if (frameIdx >= frameSize) return; if (frameIdx >= frameSize) return;
...@@ -70,10 +68,8 @@ __global__ void KeLstmForward(Op op, LstmMetaValue<T> value, int frameSize, ...@@ -70,10 +68,8 @@ __global__ void KeLstmForward(Op op, LstmMetaValue<T> value, int frameSize,
rPrevState = value.prevStateValue[frameIdx]; rPrevState = value.prevStateValue[frameIdx];
} }
hppl::gpu::ForwardAct<T> act;
op(rValueIn, rValueIg, rValueFg, rValueOg, rPrevState, rState, rStateAtv, op(rValueIn, rValueIg, rValueFg, rValueOg, rPrevState, rState, rStateAtv,
rOut, rCheckI, rCheckF, rCheckO, act(active_node), act(active_gate), rOut, rCheckI, rCheckF, rCheckO);
act(active_state));
value.gateValue[frameIdx] = rValueIn; value.gateValue[frameIdx] = rValueIn;
value.gateValue[frameIdx + frameSize] = rValueIg; value.gateValue[frameIdx + frameSize] = rValueIg;
...@@ -92,9 +88,7 @@ __global__ void KeLstmForward(Op op, LstmMetaValue<T> value, int frameSize, ...@@ -92,9 +88,7 @@ __global__ void KeLstmForward(Op op, LstmMetaValue<T> value, int frameSize,
template <class T, class Op, bool isBatch> template <class T, class Op, bool isBatch>
__global__ void KeLstmBackward(Op op, LstmMetaValue<T> value, __global__ void KeLstmBackward(Op op, LstmMetaValue<T> value,
LstmMetaGrad<T> grad, int frameSize, LstmMetaGrad<T> grad, int frameSize,
int batchSize, activation_mode_t active_node, int batchSize) {
activation_mode_t active_gate,
activation_mode_t active_state) {
const int frameIdx = blockIdx.x * blockDim.x + threadIdx.x; const int frameIdx = blockIdx.x * blockDim.x + threadIdx.x;
if (frameIdx >= frameSize) return; if (frameIdx >= frameSize) return;
...@@ -145,11 +139,9 @@ __global__ void KeLstmBackward(Op op, LstmMetaValue<T> value, ...@@ -145,11 +139,9 @@ __global__ void KeLstmBackward(Op op, LstmMetaValue<T> value,
rPrevState = value.prevStateValue[frameIdx]; rPrevState = value.prevStateValue[frameIdx];
} }
hppl::gpu::BackwardAct<T> act;
op(rValueIn, rValueIg, rValueFg, rValueOg, rGradIn, rGradIg, rGradFg, rGradOg, op(rValueIn, rValueIg, rValueFg, rValueOg, rGradIn, rGradIg, rGradFg, rGradOg,
rPrevState, rPrevStateGrad, rState, rStateGrad, rStateAtv, rOutputGrad, rPrevState, rPrevStateGrad, rState, rStateGrad, rStateAtv, rOutputGrad,
rCheckI, rCheckF, rCheckO, rCheckIGrad, rCheckFGrad, rCheckOGrad, rCheckI, rCheckF, rCheckO, rCheckIGrad, rCheckFGrad, rCheckOGrad);
act(active_node), act(active_gate), act(active_state));
grad.gateGrad[frameIdx] = rGradIn; grad.gateGrad[frameIdx] = rGradIn;
grad.gateGrad[frameIdx + frameSize] = rGradIg; grad.gateGrad[frameIdx + frameSize] = rGradIg;
...@@ -205,13 +197,11 @@ void gpu_lstm_forward(const platform::DeviceContext& context, Op op, ...@@ -205,13 +197,11 @@ void gpu_lstm_forward(const platform::DeviceContext& context, Op op,
if (batchSize == 1) { if (batchSize == 1) {
KeLstmForward<T, Op, KeLstmForward<T, Op,
/* isBatch= */ false><<<grid, threads, 0, stream>>>( /* isBatch= */ false><<<grid, threads, 0, stream>>>(
op, value, frameSize, batchSize, active_node, active_gate, op, value, frameSize, batchSize);
active_state);
} else { } else {
KeLstmForward<T, Op, KeLstmForward<T, Op,
/* isBatch= */ true><<<grid, threads, 0, stream>>>( /* isBatch= */ true><<<grid, threads, 0, stream>>>(
op, value, frameSize, batchSize, active_node, active_gate, op, value, frameSize, batchSize);
active_state);
} }
} }
...@@ -240,13 +230,11 @@ void gpu_lstm_backward(const platform::DeviceContext& context, Op op, ...@@ -240,13 +230,11 @@ void gpu_lstm_backward(const platform::DeviceContext& context, Op op,
if (batchSize == 1) { if (batchSize == 1) {
KeLstmBackward<T, Op, KeLstmBackward<T, Op,
/* isBatch= */ false><<<grid, threads, 0, stream>>>( /* isBatch= */ false><<<grid, threads, 0, stream>>>(
op, value, grad, frameSize, batchSize, active_node, active_gate, op, value, grad, frameSize, batchSize);
active_state);
} else { } else {
KeLstmBackward<T, Op, KeLstmBackward<T, Op,
/* isBatch= */ true><<<grid, threads, 0, stream>>>( /* isBatch= */ true><<<grid, threads, 0, stream>>>(
op, value, grad, frameSize, batchSize, active_node, active_gate, op, value, grad, frameSize, batchSize);
active_state);
} }
} }
......
...@@ -24,15 +24,29 @@ namespace detail { ...@@ -24,15 +24,29 @@ namespace detail {
namespace forward { namespace forward {
template <typename T>
DEVICE inline T sigmoid(const T a) {
const T min = SIGMOID_THRESHOLD_MIN;
const T max = SIGMOID_THRESHOLD_MAX;
T tmp = (a < min) ? min : ((a > max) ? max : a);
return static_cast<T>(1.0) / (static_cast<T>(1.0) + exp(-tmp));
}
template <typename T>
DEVICE inline T tanh(const T a) {
T tmp = -2.0 * a;
tmp = (tmp > EXP_MAX_INPUT) ? EXP_MAX_INPUT : tmp;
return (2.0 / (1.0 + exp(tmp))) - 1.0;
}
template <class T> template <class T>
class lstm { class lstm {
public: public:
HOSTDEVICE void operator()(T &valueIn, T &valueIg, T &valueFg, T &valueOg, HOSTDEVICE void operator()(T &valueIn, T &valueIg, T &valueFg, T &valueOg,
T &prevState, T &state, T &stateAtv, T &output, T &prevState, T &state, T &stateAtv, T &output,
T &checkI, T &checkF, T &checkO, T &checkI, T &checkF, T &checkO) {
typename hppl::ForwardActType<T>::type actInput, #if 0
typename hppl::ForwardActType<T>::type actGate, // TODO(qingqing) support to activation speficed by users
typename hppl::ForwardActType<T>::type actState) {
valueIn = actInput(valueIn); valueIn = actInput(valueIn);
valueIg = actGate(valueIg + prevState * checkI); valueIg = actGate(valueIg + prevState * checkI);
valueFg = actGate(valueFg + prevState * checkF); valueFg = actGate(valueFg + prevState * checkF);
...@@ -40,6 +54,15 @@ class lstm { ...@@ -40,6 +54,15 @@ class lstm {
valueOg = actGate(valueOg + state * checkO); valueOg = actGate(valueOg + state * checkO);
stateAtv = actState(state); stateAtv = actState(state);
output = valueOg * stateAtv; output = valueOg * stateAtv;
#else
valueIn = tanh<T>(valueIn);
valueIg = sigmoid<T>(valueIg + prevState * checkI);
valueFg = sigmoid<T>(valueFg + prevState * checkF);
state = valueIn * valueIg + prevState * valueFg;
valueOg = sigmoid<T>(valueOg + state * checkO);
stateAtv = tanh<T>(state);
output = valueOg * stateAtv;
#endif
} }
#ifndef __NVCC__ #ifndef __NVCC__
#ifndef __AVX__ // If not compiled with AVX instructs. Disable AVX by default #ifndef __AVX__ // If not compiled with AVX instructs. Disable AVX by default
...@@ -72,6 +95,16 @@ class lstm { ...@@ -72,6 +95,16 @@ class lstm {
namespace backward { namespace backward {
template <typename T>
DEVICE inline T sigmoid(const T a, const T b) {
return a * b * (1.0 - b);
}
template <typename T>
DEVICE inline T tanh(const T a, const T b) {
return a * (1.0 - b * b);
}
template <class T> template <class T>
class lstm { class lstm {
public: public:
...@@ -80,10 +113,9 @@ class lstm { ...@@ -80,10 +113,9 @@ class lstm {
T &prevState, T &prevStateGrad, T &state, T &prevState, T &prevStateGrad, T &state,
T &stateGrad, T &stateAtv, T &outputGrad, T &stateGrad, T &stateAtv, T &outputGrad,
T &checkI, T &checkF, T &checkO, T &checkIGrad, T &checkI, T &checkF, T &checkO, T &checkIGrad,
T &checkFGrad, T &checkOGrad, T &checkFGrad, T &checkOGrad) {
typename hppl::BackwardActType<T>::type actInput, #if 0
typename hppl::BackwardActType<T>::type actGate, // TODO(qingqing) support to activation speficed by users
typename hppl::BackwardActType<T>::type actState) {
gradOg = actGate(outputGrad * stateAtv, valueOg); gradOg = actGate(outputGrad * stateAtv, valueOg);
stateGrad += actState(outputGrad * valueOg, stateAtv) + gradOg * checkO; stateGrad += actState(outputGrad * valueOg, stateAtv) + gradOg * checkO;
gradIn = actInput(stateGrad * valueIg, valueIn); gradIn = actInput(stateGrad * valueIg, valueIn);
...@@ -93,6 +125,17 @@ class lstm { ...@@ -93,6 +125,17 @@ class lstm {
checkIGrad = gradIg * prevState; checkIGrad = gradIg * prevState;
checkFGrad = gradFg * prevState; checkFGrad = gradFg * prevState;
checkOGrad = gradOg * state; checkOGrad = gradOg * state;
#else
gradOg = sigmoid<T>(outputGrad * stateAtv, valueOg);
stateGrad += tanh<T>(outputGrad * valueOg, stateAtv) + gradOg * checkO;
gradIn = tanh<T>(stateGrad * valueIg, valueIn);
gradIg = sigmoid<T>(stateGrad * valueIn, valueIg);
gradFg = sigmoid<T>(stateGrad * prevState, valueFg);
prevStateGrad = gradIg * checkI + gradFg * checkF + stateGrad * valueFg;
checkIGrad = gradIg * prevState;
checkFGrad = gradFg * prevState;
checkOGrad = gradOg * state;
#endif
} }
#ifndef __NVCC__ #ifndef __NVCC__
#ifndef __AVX__ // If not compiled with AVX instructs. Disable AVX by default #ifndef __AVX__ // If not compiled with AVX instructs. Disable AVX by default
......
...@@ -211,6 +211,26 @@ void batched_gemm<platform::CPUPlace, double>( ...@@ -211,6 +211,26 @@ void batched_gemm<platform::CPUPlace, double>(
} }
#endif #endif
template <>
void gemv<platform::CPUPlace, float>(const platform::DeviceContext& context,
const bool trans_a, const int M,
const int N, const float alpha,
const float* A, const float* B,
const float beta, float* C) {
CBLAS_TRANSPOSE transA = (trans_a == false) ? CblasNoTrans : CblasTrans;
cblas_sgemv(CblasRowMajor, transA, M, N, alpha, A, N, B, 1, beta, C, 1);
}
template <>
void gemv<platform::CPUPlace, double>(const platform::DeviceContext& context,
const bool trans_a, const int M,
const int N, const double alpha,
const double* A, const double* B,
const double beta, double* C) {
CBLAS_TRANSPOSE transA = (trans_a == false) ? CblasNoTrans : CblasTrans;
cblas_dgemv(CblasRowMajor, transA, M, N, alpha, A, N, B, 1, beta, C, 1);
}
template struct SetConstant<platform::CPUPlace, float>; template struct SetConstant<platform::CPUPlace, float>;
} // namespace math } // namespace math
......
...@@ -203,6 +203,33 @@ void batched_gemm<platform::GPUPlace, double>( ...@@ -203,6 +203,33 @@ void batched_gemm<platform::GPUPlace, double>(
&beta, C, ldc, strideC, batchCount)); &beta, C, ldc, strideC, batchCount));
} }
template <>
void gemv<platform::GPUPlace, float>(const platform::DeviceContext& context,
const bool trans_a, const int M,
const int N, const float alpha,
const float* A, const float* B,
const float beta, float* C) {
cublasOperation_t cuTransA = (trans_a == false) ? CUBLAS_OP_T : CUBLAS_OP_N;
PADDLE_ENFORCE(platform::dynload::cublasSgemv(
reinterpret_cast<const platform::CUDADeviceContext&>(context)
.cublas_handle(),
cuTransA, N, M, &alpha, A, N, B, 1, &beta, C, 1));
}
template <>
void gemv<platform::GPUPlace, double>(const platform::DeviceContext& context,
const bool trans_a, const int M,
const int N, const double alpha,
const double* A, const double* B,
const double beta, double* C) {
cublasOperation_t cuTransA = (trans_a == false) ? CUBLAS_OP_T : CUBLAS_OP_N;
PADDLE_ENFORCE(platform::dynload::cublasDgemv(
reinterpret_cast<const platform::CUDADeviceContext&>(context)
.cublas_handle(),
cuTransA, N, M, &alpha, A, N, B, 1, &beta, C, 1));
}
template struct SetConstant<platform::GPUPlace, float>; template struct SetConstant<platform::GPUPlace, float>;
} // namespace math } // namespace math
......
...@@ -93,6 +93,11 @@ void batched_gemm(const platform::DeviceContext& context, ...@@ -93,6 +93,11 @@ void batched_gemm(const platform::DeviceContext& context,
const T* A, const T* B, const T beta, T* C, const T* A, const T* B, const T beta, T* C,
const int batchCount, const int strideA, const int strideB); const int batchCount, const int strideA, const int strideB);
template <typename Place, typename T>
void gemv(const platform::DeviceContext& context, const bool trans_a,
const int M, const int N, const T alpha, const T* A, const T* B,
const T beta, T* C);
template <typename Place, typename T> template <typename Place, typename T>
struct SetConstant { struct SetConstant {
void operator()(const platform::DeviceContext& context, void operator()(const platform::DeviceContext& context,
......
...@@ -89,3 +89,53 @@ TEST(math_function, zero) { ...@@ -89,3 +89,53 @@ TEST(math_function, zero) {
EXPECT_EQ(t[2], 1); EXPECT_EQ(t[2], 1);
EXPECT_EQ(t[3], 1); EXPECT_EQ(t[3], 1);
} }
template <typename T>
void GemvTest(int m, int n, bool trans) {
paddle::framework::Tensor mat_a;
paddle::framework::Tensor vec_b;
paddle::framework::Tensor vec_c;
auto* cpu_place = new paddle::platform::CPUPlace();
int b_num = trans ? m : n;
int c_num = trans ? n : m;
T* data_a = mat_a.mutable_data<T>({m, n}, *cpu_place);
T* data_b = vec_b.mutable_data<T>({b_num}, *cpu_place);
T* data_c = vec_c.mutable_data<T>({c_num}, *cpu_place);
for (int i = 0; i < mat_a.numel(); ++i) {
data_a[i] = static_cast<T>(i);
}
for (int i = 0; i < vec_b.numel(); ++i) {
data_b[i] = static_cast<T>(i);
}
paddle::platform::CPUDeviceContext context(*cpu_place);
paddle::operators::math::gemv<paddle::platform::CPUPlace, T>(
context, trans, static_cast<int>(m), static_cast<int>(n), 1., data_a,
data_b, 0., data_c);
if (!trans) {
for (int i = 0; i < m; ++i) {
T sum = 0.0;
for (int j = 0; j < n; ++j) {
sum += data_a[i * n + j] * data_b[j];
}
ASSERT_FLOAT_EQ(data_c[i], sum);
}
} else {
for (int i = 0; i < n; ++i) {
T sum = 0.0;
for (int j = 0; j < m; ++j) {
sum += data_a[j * n + i] * data_b[j];
}
ASSERT_FLOAT_EQ(data_c[i], sum);
}
}
}
TEST(math_function, gemv) {
GemvTest<float>(3, 13, false);
GemvTest<double>(4, 5, false);
GemvTest<float>(12, 7, true);
GemvTest<double>(7, 9, true);
}
...@@ -177,3 +177,65 @@ TEST(math_function, gemm_trans_cublas) { ...@@ -177,3 +177,65 @@ TEST(math_function, gemm_trans_cublas) {
EXPECT_EQ(input3_ptr[7], 99); EXPECT_EQ(input3_ptr[7], 99);
delete gpu_place; delete gpu_place;
} }
template <typename T>
void GemvTest(int m, int n, bool trans) {
paddle::framework::Tensor mat_a;
paddle::framework::Tensor vec_b;
paddle::framework::Tensor vec_c;
auto* cpu_place = new paddle::platform::CPUPlace();
T* data_a = mat_a.mutable_data<T>({m, n}, *cpu_place);
T* data_b = vec_b.mutable_data<T>({trans ? m : n}, *cpu_place);
T* data_c = vec_c.mutable_data<T>({trans ? n : m}, *cpu_place);
auto* gpu_place = new paddle::platform::GPUPlace(0);
paddle::framework::Tensor g_mat_a;
paddle::framework::Tensor g_vec_b;
paddle::framework::Tensor g_vec_c;
T* g_data_a = g_mat_a.mutable_data<T>(mat_a.dims(), *gpu_place);
T* g_data_b = g_vec_b.mutable_data<T>(vec_b.dims(), *gpu_place);
T* g_data_c = g_vec_c.mutable_data<T>(vec_c.dims(), *gpu_place);
for (int i = 0; i < mat_a.numel(); ++i) {
data_a[i] = static_cast<T>(i);
}
for (int i = 0; i < vec_b.numel(); ++i) {
data_b[i] = static_cast<T>(i);
}
paddle::platform::CUDADeviceContext context(*gpu_place);
g_mat_a.CopyFrom(mat_a, *gpu_place, context);
g_vec_b.CopyFrom(vec_b, *gpu_place, context);
paddle::operators::math::gemv<paddle::platform::GPUPlace, T>(
context, trans, static_cast<int>(m), static_cast<int>(n), 1., g_data_a,
g_data_b, 0., g_data_c);
vec_c.CopyFrom(g_vec_c, paddle::platform::CPUPlace(), context);
if (!trans) {
for (int i = 0; i < m; ++i) {
T sum = 0.0;
for (int j = 0; j < n; ++j) {
sum += data_a[i * n + j] * data_b[j];
}
ASSERT_FLOAT_EQ(data_c[i], sum);
}
} else {
for (int i = 0; i < n; ++i) {
T sum = 0.0;
for (int j = 0; j < m; ++j) {
sum += data_a[j * n + i] * data_b[j];
}
ASSERT_FLOAT_EQ(data_c[i], sum);
}
}
}
TEST(math_function, gemv) {
GemvTest<float>(3, 13, false);
GemvTest<double>(3, 13, false);
GemvTest<float>(3, 13, true);
GemvTest<double>(3, 13, true);
}
...@@ -53,7 +53,18 @@ class LoDTensor2BatchFunctor { ...@@ -53,7 +53,18 @@ class LoDTensor2BatchFunctor {
public: public:
void operator()(const platform::DeviceContext& context, void operator()(const platform::DeviceContext& context,
const framework::LoDTensor& lod_tensor, const framework::LoDTensor& lod_tensor,
framework::LoDTensor& batch, bool is_reverse) const { framework::LoDTensor& batch, bool is_cal_batch_lod,
bool is_reverse = false) const {
if (!is_cal_batch_lod) {
auto lods = batch.lod();
PADDLE_ENFORCE_EQ(lods.size(), 2UL);
PADDLE_ENFORCE_EQ(lods[1].size(),
static_cast<size_t>(lod_tensor.dims()[0]));
CopyMatrixRowsFunctor<Place, T> to_batch;
to_batch(context, lod_tensor, lods[1].data(), batch, true);
return;
}
auto lods = lod_tensor.lod(); auto lods = lod_tensor.lod();
PADDLE_ENFORCE_EQ(lods.size(), 1UL, "Only support one level sequence now."); PADDLE_ENFORCE_EQ(lods.size(), 1UL, "Only support one level sequence now.");
auto lod = lods[0]; auto lod = lods[0];
...@@ -101,10 +112,10 @@ class LoDTensor2BatchFunctor { ...@@ -101,10 +112,10 @@ class LoDTensor2BatchFunctor {
size_t* batch_starts = batch_lods[0].data(); size_t* batch_starts = batch_lods[0].data();
size_t* seq2batch_idx = batch_lods[1].data(); size_t* seq2batch_idx = batch_lods[1].data();
batch_starts[0] = 0; batch_starts[0] = 0;
for (size_t n = 0; n < num_batch; n++) { for (int n = 0; n < num_batch; n++) {
auto batch_id = static_cast<int>(batch_starts[n]); auto batch_id = static_cast<int>(batch_starts[n]);
for (size_t i = 0; i < seq_info.size(); ++i) { for (size_t i = 0; i < seq_info.size(); ++i) {
size_t seq_len = seq_info[i].length; int seq_len = seq_info[i].length;
int start = seq_info[i].start; int start = seq_info[i].start;
if (n < seq_len) { if (n < seq_len) {
seq2batch_idx[batch_id] = seq2batch_idx[batch_id] =
...@@ -132,11 +143,8 @@ class Batch2LoDTensorFunctor { ...@@ -132,11 +143,8 @@ class Batch2LoDTensorFunctor {
auto in_lod = batch.lod(); auto in_lod = batch.lod();
PADDLE_ENFORCE_EQ(in_lod.size(), 2UL, PADDLE_ENFORCE_EQ(in_lod.size(), 2UL,
"The LoD size of input `batch` should be 2."); "The LoD size of input `batch` should be 2.");
auto out_lod = lod_tensor.lod()[0]; PADDLE_ENFORCE_EQ(in_lod[1].size(),
auto num = out_lod[out_lod.size() - 1]; static_cast<size_t>(lod_tensor.dims()[0]));
PADDLE_ENFORCE_EQ(num, lod_tensor.dims()[0]);
PADDLE_ENFORCE_EQ(num, in_lod[1].size());
PADDLE_ENFORCE_EQ(num, batch.dims()[0]);
CopyMatrixRowsFunctor<Place, T> to_seq; CopyMatrixRowsFunctor<Place, T> to_seq;
size_t* index = in_lod[1].data(); size_t* index = in_lod[1].data();
to_seq(context, batch, index, lod_tensor, false); to_seq(context, batch, index, lod_tensor, false);
......
...@@ -185,7 +185,7 @@ TEST_F(NCCLTester, ncclAllReduceOp) { ...@@ -185,7 +185,7 @@ TEST_F(NCCLTester, ncclAllReduceOp) {
recv_tensor.numel() * sizeof(float), recv_tensor.numel() * sizeof(float),
static_cast<p::CUDADeviceContext *>(dev_ctxs[i])->stream()); static_cast<p::CUDADeviceContext *>(dev_ctxs[i])->stream());
for (size_t j = 0; j < f::product(kDims); ++j) { for (int64_t j = 0; j < f::product(kDims); ++j) {
ASSERT_NEAR(ct[j], result, 1e-5); ASSERT_NEAR(ct[j], result, 1e-5);
} }
} }
...@@ -234,7 +234,7 @@ TEST_F(NCCLTester, ncclReduceOp) { ...@@ -234,7 +234,7 @@ TEST_F(NCCLTester, ncclReduceOp) {
recv_tensor.numel() * sizeof(float), recv_tensor.numel() * sizeof(float),
static_cast<p::CUDADeviceContext *>(dev_ctxs[kRoot])->stream()); static_cast<p::CUDADeviceContext *>(dev_ctxs[kRoot])->stream());
for (int j = 0; j < f::product(kDims); ++j) { for (int64_t j = 0; j < f::product(kDims); ++j) {
ASSERT_NEAR(ct[j], result, 1e-5); ASSERT_NEAR(ct[j], result, 1e-5);
} }
} }
...@@ -282,7 +282,7 @@ TEST_F(NCCLTester, ncclBcastOp) { ...@@ -282,7 +282,7 @@ TEST_F(NCCLTester, ncclBcastOp) {
recv_tensor.numel() * sizeof(float), recv_tensor.numel() * sizeof(float),
static_cast<p::CUDADeviceContext *>(dev_ctxs[idx])->stream()); static_cast<p::CUDADeviceContext *>(dev_ctxs[idx])->stream());
for (size_t j = 0; j < f::product(kDims); ++j) { for (int64_t j = 0; j < f::product(kDims); ++j) {
ASSERT_NEAR(ct[j], result, 1e-5); ASSERT_NEAR(ct[j], result, 1e-5);
} }
} }
......
...@@ -36,7 +36,7 @@ class ReshapeOp : public framework::OperatorWithKernel { ...@@ -36,7 +36,7 @@ class ReshapeOp : public framework::OperatorWithKernel {
PADDLE_ENFORCE(shape.size() > 0, "Attr(shape) shouldn't be empty."); PADDLE_ENFORCE(shape.size() > 0, "Attr(shape) shouldn't be empty.");
auto x_dims = ctx->GetInputDim("X"); auto x_dims = ctx->GetInputDim("X");
// TODO(qiao) change batch_size // TODO(qiao) change batch_size
for (int i = 1; i < shape.size(); ++i) { for (size_t i = 1; i < shape.size(); ++i) {
PADDLE_ENFORCE(shape[i] > 0, PADDLE_ENFORCE(shape[i] > 0,
"Each dimension of shape " "Each dimension of shape "
"must be positiv except the first."); "must be positiv except the first.");
......
...@@ -34,7 +34,7 @@ TEST(SaveLoadOp, CPU) { ...@@ -34,7 +34,7 @@ TEST(SaveLoadOp, CPU) {
tensor->set_lod(expect_lod); tensor->set_lod(expect_lod);
int* expect = tensor->mutable_data<int>(place); int* expect = tensor->mutable_data<int>(place);
for (size_t i = 0; i < paddle::framework::product(tensor->dims()); ++i) { for (int64_t i = 0; i < tensor->numel(); ++i) {
expect[i] = static_cast<int>(i); expect[i] = static_cast<int>(i);
} }
paddle::framework::AttributeMap attrs; paddle::framework::AttributeMap attrs;
...@@ -50,7 +50,7 @@ TEST(SaveLoadOp, CPU) { ...@@ -50,7 +50,7 @@ TEST(SaveLoadOp, CPU) {
"load", {}, {{"Out", {"out_var"}}}, attrs); "load", {}, {{"Out", {"out_var"}}}, attrs);
load_op->Run(scope, ctx); load_op->Run(scope, ctx);
int* actual = target->data<int>(); int* actual = target->data<int>();
for (size_t i = 0; i < paddle::framework::product(tensor->dims()); ++i) { for (int64_t i = 0; i < tensor->numel(); ++i) {
EXPECT_EQ(expect[i], actual[i]); EXPECT_EQ(expect[i], actual[i]);
} }
auto& actual_lod = target->lod(); auto& actual_lod = target->lod();
...@@ -60,4 +60,4 @@ TEST(SaveLoadOp, CPU) { ...@@ -60,4 +60,4 @@ TEST(SaveLoadOp, CPU) {
EXPECT_EQ(expect_lod[i][j], actual_lod[i][j]); EXPECT_EQ(expect_lod[i][j], actual_lod[i][j]);
} }
} }
} }
\ No newline at end of file
...@@ -39,15 +39,14 @@ class SequencePoolOpMaker : public framework::OpProtoAndCheckerMaker { ...@@ -39,15 +39,14 @@ class SequencePoolOpMaker : public framework::OpProtoAndCheckerMaker {
AddOutput("Out", AddOutput("Out",
"(Tensor), output of SequencePoolOp, which does not contain LoD " "(Tensor), output of SequencePoolOp, which does not contain LoD "
"infomation."); "infomation.");
AddAttr<int>( AddAttr<std::string>(
"strategy", "pooltype",
"(int, default AVERAGE) the pooling strategy of SequencePoolOp.") "(int, default AVERAGE) the pooling pooltype of SequencePoolOp.")
.SetDefault(AVERAGE) .SetDefault("AVERAGE");
.InEnum({AVERAGE, SUM, SQRT, MAX, LAST, FIRST});
AddComment(R"DOC( AddComment(R"DOC(
SequencePoolOp pools features of all time-steps of each instance. SequencePoolOp pools features of all time-steps of each instance.
It supports six pooling strategy: It supports six pooling pooltype:
- AVERAGE: Out[i] = average_{for each instance in i-th sequence}{X[i]} - AVERAGE: Out[i] = average_{for each instance in i-th sequence}{X[i]}
- SUM: Out[i] = sum_{for each instance in i-th sequence}{X[i]} - SUM: Out[i] = sum_{for each instance in i-th sequence}{X[i]}
- SQRT: Out[i] = sum_{for each instance in i-th sequence}{X[i]} - SQRT: Out[i] = sum_{for each instance in i-th sequence}{X[i]}
...@@ -63,7 +62,7 @@ class SequencePoolOpMaker : public framework::OpProtoAndCheckerMaker { ...@@ -63,7 +62,7 @@ class SequencePoolOpMaker : public framework::OpProtoAndCheckerMaker {
and the value of X = [[1, 3], [2, 4, 6], [5, 1]]. and the value of X = [[1, 3], [2, 4, 6], [5, 1]].
Thus, Out is a [3,1,1] Tensor without LoD infomation. Thus, Out is a [3,1,1] Tensor without LoD infomation.
And for different strategy, the value of Out is as follows: And for different pooltype, the value of Out is as follows:
- AVERAGE: [2, 4, 3], where 2=(1+3)/2, 4=(2+4+6)/3, 3=(5+1)/2 - AVERAGE: [2, 4, 3], where 2=(1+3)/2, 4=(2+4+6)/3, 3=(5+1)/2
- SUM: [4, 12, 6], where 4=1+3, 12=2+4+6, 6=5+1 - SUM: [4, 12, 6], where 4=1+3, 12=2+4+6, 6=5+1
......
...@@ -29,22 +29,13 @@ template <typename T, int MajorType = Eigen::RowMajor, ...@@ -29,22 +29,13 @@ template <typename T, int MajorType = Eigen::RowMajor,
typename IndexType = Eigen::DenseIndex> typename IndexType = Eigen::DenseIndex>
using EigenMatrix = framework::EigenMatrix<T, MajorType, IndexType>; using EigenMatrix = framework::EigenMatrix<T, MajorType, IndexType>;
enum SeqPoolType {
AVERAGE = 0,
SUM = 1,
SQRT = 2, // square_root_n
MAX = 3,
LAST = 4,
FIRST = 5
};
template <typename Place, typename T> template <typename Place, typename T>
class SequencePoolKernel : public framework::OpKernel<T> { class SequencePoolKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& context) const override { void Compute(const framework::ExecutionContext& context) const override {
auto* in = context.Input<LoDTensor>("X"); auto* in = context.Input<LoDTensor>("X");
auto* out = context.Output<LoDTensor>("Out"); auto* out = context.Output<LoDTensor>("Out");
int strategy = context.Attr<int>("strategy"); std::string pooltype = context.Attr<std::string>("pooltype");
auto dims = in->dims(); auto dims = in->dims();
auto lod = in->lod(); auto lod = in->lod();
...@@ -71,28 +62,21 @@ class SequencePoolKernel : public framework::OpKernel<T> { ...@@ -71,28 +62,21 @@ class SequencePoolKernel : public framework::OpKernel<T> {
auto in_e = EigenMatrix<T>::From(in_t, framework::make_ddim({h, w})); auto in_e = EigenMatrix<T>::From(in_t, framework::make_ddim({h, w}));
auto out_e = EigenVector<T>::Flatten(out_t); auto out_e = EigenVector<T>::Flatten(out_t);
switch (strategy) { if (pooltype == "AVERAGE") {
case AVERAGE: out_e.device(place) = in_e.mean(Eigen::array<int, 1>({{0}}));
out_e.device(place) = in_e.mean(Eigen::array<int, 1>({{0}})); } else if (pooltype == "SUM") {
break; out_e.device(place) = in_e.sum(Eigen::array<int, 1>({{0}}));
case SUM: } else if (pooltype == "SQRT") {
out_e.device(place) = in_e.sum(Eigen::array<int, 1>({{0}})); out_e.device(place) = in_e.sum(Eigen::array<int, 1>({{0}})) /
break; std::sqrt(static_cast<T>(h));
case SQRT: } else if (pooltype == "MAX") {
out_e.device(place) = in_e.sum(Eigen::array<int, 1>({{0}})) / out_e.device(place) = in_e.maximum(Eigen::array<int, 1>({{0}}));
std::sqrt(static_cast<T>(h)); } else if (pooltype == "LAST") {
break; out_e.device(place) = in_e.chip(h - 1, 0);
case MAX: } else if (pooltype == "FIRST") {
out_e.device(place) = in_e.maximum(Eigen::array<int, 1>({{0}})); out_e.device(place) = in_e.chip(0, 0);
break; } else {
case LAST: PADDLE_THROW("unsupported pooling pooltype");
out_e.device(place) = in_e.chip(h - 1, 0);
break;
case FIRST:
out_e.device(place) = in_e.chip(0, 0);
break;
default:
PADDLE_THROW("unsupported pooling strategy");
} }
} }
} }
...@@ -105,15 +89,15 @@ class SequencePoolGradKernel : public framework::OpKernel<T> { ...@@ -105,15 +89,15 @@ class SequencePoolGradKernel : public framework::OpKernel<T> {
auto* in = context.Input<LoDTensor>("X"); auto* in = context.Input<LoDTensor>("X");
auto* in_g = context.Output<LoDTensor>(framework::GradVarName("X")); auto* in_g = context.Output<LoDTensor>(framework::GradVarName("X"));
auto* out_g = context.Input<LoDTensor>(framework::GradVarName("Out")); auto* out_g = context.Input<LoDTensor>(framework::GradVarName("Out"));
int strategy = context.Attr<int>("strategy"); std::string pooltype = context.Attr<std::string>("pooltype");
auto dims = in->dims(); auto dims = in->dims();
auto lod = in->lod()[0]; auto lod = in->lod()[0];
int64_t w = in->numel() / dims[0]; int64_t w = in->numel() / dims[0];
in_g->mutable_data<T>(context.GetPlace()); in_g->mutable_data<T>(context.GetPlace());
if (strategy == LAST || strategy == FIRST) { if (pooltype == "LAST" || pooltype == "FIRST") {
// set X@Grad be zero at first when strategy is LAST/FIRST // set X@Grad be zero at first when pooltype is LAST/FIRST
math::SetConstant<Place, T> functor; math::SetConstant<Place, T> functor;
functor(context.device_context(), in_g, 0); functor(context.device_context(), in_g, 0);
} }
...@@ -127,41 +111,33 @@ class SequencePoolGradKernel : public framework::OpKernel<T> { ...@@ -127,41 +111,33 @@ class SequencePoolGradKernel : public framework::OpKernel<T> {
auto out_g_e = EigenMatrix<T>::From(out_g_t, {1, w}); auto out_g_e = EigenMatrix<T>::From(out_g_t, {1, w});
Eigen::DSizes<int, 2> bcast(h, 1); Eigen::DSizes<int, 2> bcast(h, 1);
switch (strategy) { if (pooltype == "AVERAGE") {
case AVERAGE: in_g_e.device(place) = (out_g_e / static_cast<T>(h)).broadcast(bcast);
in_g_e.device(place) = (out_g_e / static_cast<T>(h)).broadcast(bcast); } else if (pooltype == "SUM") {
break; in_g_e.device(place) = (out_g_e).broadcast(bcast);
case SUM: } else if (pooltype == "SQRT") {
in_g_e.device(place) = (out_g_e).broadcast(bcast); in_g_e.device(place) =
break; (out_g_e / std::sqrt(static_cast<T>(h))).broadcast(bcast);
case SQRT: } else if (pooltype == "MAX") {
in_g_e.device(place) = auto in_t =
(out_g_e / std::sqrt(static_cast<T>(h))).broadcast(bcast); in->Slice(static_cast<int>(lod[i]), static_cast<int>(lod[i + 1]));
break; Eigen::Map<const Eigen::Matrix<T, Eigen::Dynamic, Eigen::Dynamic>>
case MAX: { in_t_map(in_t.data<T>(), h, w);
auto in_t = int row_id;
in->Slice(static_cast<int>(lod[i]), static_cast<int>(lod[i + 1])); Eigen::array<int, 2> extents{{1, 1}};
Eigen::Map<const Eigen::Matrix<T, Eigen::Dynamic, Eigen::Dynamic>> for (int col_id = 0; col_id < w; col_id++) {
in_t_map(in_t.data<T>(), h, w); in_t_map.col(col_id).maxCoeff(&row_id);
int row_id; Eigen::array<int, 2> in_offsets{{row_id, col_id}};
Eigen::array<int, 2> extents{{1, 1}}; Eigen::array<int, 2> out_offsets{{0, col_id}};
for (int col_id = 0; col_id < w; col_id++) { in_g_e.slice(in_offsets, extents).device(place) =
in_t_map.col(col_id).maxCoeff(&row_id); out_g_e.slice(out_offsets, extents);
Eigen::array<int, 2> in_offsets{{row_id, col_id}};
Eigen::array<int, 2> out_offsets{{0, col_id}};
in_g_e.slice(in_offsets, extents).device(place) =
out_g_e.slice(out_offsets, extents);
}
break;
} }
case LAST: } else if (pooltype == "LAST") {
in_g_e.chip(h - 1, 0).device(place) = out_g_e; in_g_e.chip(h - 1, 0).device(place) = out_g_e;
break; } else if (pooltype == "FIRST") {
case FIRST: in_g_e.chip(0, 0).device(place) = out_g_e;
in_g_e.chip(0, 0).device(place) = out_g_e; } else {
break; PADDLE_THROW("unsupported pooling pooltype");
default:
PADDLE_THROW("unsupported pooling strategy");
} }
} }
} }
......
...@@ -32,9 +32,9 @@ class SoftmaxWithCrossEntropyOpMaker ...@@ -32,9 +32,9 @@ class SoftmaxWithCrossEntropyOpMaker
AddInput("Label", AddInput("Label",
"(Tensor, default: Tensor<int>), The ground truth which is a 2-D " "(Tensor, default: Tensor<int>), The ground truth which is a 2-D "
"tensor. " "tensor. "
"If softLable is set to 0, Label is a Tensor<int> with shape [N x " "If softLabel is set to false, Label is a Tensor<int> with shape "
"1]. " "[N x 1]."
"If softLable is set to 1, Label is a Tensor<float/double> " "If softLabel is set to true, Label is a Tensor<float/double> "
"with shape [N x K]."); "with shape [N x K].");
AddOutput( AddOutput(
"Softmax", "Softmax",
...@@ -60,19 +60,23 @@ Because this operators performs a softmax on logits internally, it expects ...@@ -60,19 +60,23 @@ Because this operators performs a softmax on logits internally, it expects
unscaled logits. Please do not call this op with the output of softmax operator, unscaled logits. Please do not call this op with the output of softmax operator,
which will produce incorrect results. which will produce incorrect results.
This operators expects mutually exclusive hard labels, each sample in a batch When the attribute softLabel is set false, this operators expects mutually
is in exactly one class with probabilities 1. Each sample in the batch with one exclusive hard labels, each sample in a batch is in exactly one class with
and only one label. probabilities 1. Each sample in the batch with one and only one label.
Equation: Equation:
1) hard label (one-hot label) 1) hard label (one-hot label)
Loss_j = -\text{Logit}_{Label_j} + \log\left(\sum_{i=0}^{K}\exp(\text{Logit}_i)\right), j = 1, ..., K Loss_j = \f$ -\text{Logit}_{Label_j} +
\log\left(\sum_{i=0}^{K}\exp(\text{Logit}_i)\right),
j = 1, ..., K $\f
2) soft label (a distribution over all classes) 2) soft label (a distribution over all classes)
Loss_j = -\sum_{i=0}^{K}\text{Label}_i\left(\text{Logit}_i-\log\left(\sum_{i=0}^{K}\exp(\text{Logit}_i)\right)\right), j = 1,...,K Loss_j = \f$ -\sum_{i=0}^{K}\text{Label}_i\left(\text{Logit}_i -
\log\left(\sum_{i=0}^{K}\exp(\text{Logit}_i)\right)\right),
j = 1,...,K $\f
)DOC"); )DOC");
} }
......
...@@ -4,6 +4,10 @@ set -xe ...@@ -4,6 +4,10 @@ set -xe
if [ $ANDROID_ABI == "arm64-v8a" ]; then if [ $ANDROID_ABI == "arm64-v8a" ]; then
ANDROID_ARCH=arm64 ANDROID_ARCH=arm64
if [ $ANDROID_API -lt 21 ]; then
echo "Warning: arm64-v8a requires ANDROID_API >= 21."
ANDROID_API=21
fi
else # armeabi, armeabi-v7a else # armeabi, armeabi-v7a
ANDROID_ARCH=arm ANDROID_ARCH=arm
fi fi
......
...@@ -27,6 +27,13 @@ using namespace paddle; // NOLINT ...@@ -27,6 +27,13 @@ using namespace paddle; // NOLINT
using namespace std; // NOLINT using namespace std; // NOLINT
int main(int argc, char** argv) { int main(int argc, char** argv) {
if (FLAGS_model_dir.empty() || FLAGS_config_file.empty() ||
FLAGS_model_file.empty()) {
LOG(INFO) << "Usage: ./paddle_merge_model --model_dir=pass-00000 "
"--config_file=config.py --model_file=out.paddle";
return 0;
}
initMain(argc, argv); initMain(argc, argv);
initPython(argc, argv); initPython(argc, argv);
......
...@@ -351,32 +351,21 @@ def conv2d(input, ...@@ -351,32 +351,21 @@ def conv2d(input,
return helper.append_activation(pre_act) return helper.append_activation(pre_act)
def sequence_pool(input, pool_type, program=None, init_program=None): def sequence_pool(input, pool_type, **kwargs):
# FIXME(dzh) : want to unify the argument of python layer ENUM_POOL_TYPE = set(["MAX", "AVG", "SQRT", "LAST", "FIRST"])
# function. So we ignore some unecessary attributes
ENUM_POOL_TYPE = dict({
"AVERAGE": 0,
"SUM": 1,
"SQRT": 2,
"MAX": 3,
"LAST": 4,
"FIRST": 5
})
if pool_type.upper() not in ENUM_POOL_TYPE: if pool_type.upper() not in ENUM_POOL_TYPE:
raise ValueError("Unknown pool_type: '%s'. It can only be %s.", raise ValueError("Unknown pool_type: '%s'. It can only be %s.",
str(pool_type), " ".join(ENUM_POOL_TYPE.keys())) str(pool_type), " ".join(ENUM_POOL_TYPE))
helper = LayerHelper('sequence_pool', **locals()) helper = LayerHelper('sequence_pool', **kwargs)
dtype = helper.input_dtype() dtype = helper.input_dtype()
pool_out = helper.create_tmp_variable(dtype) pool_out = helper.create_tmp_variable(dtype)
# FIXME(dzh): strategy
helper.append_op( helper.append_op(
type="sequence_pool", type="sequence_pool",
inputs={"X": [input]}, inputs={"X": [input]},
outputs={"Out": [pool_out]}, outputs={"Out": [pool_out]},
attrs={"strategy": ENUM_POOL_TYPE[pool_type.upper()]}) attrs={"pooltype": pool_type.upper()})
return pool_out return pool_out
......
import unittest
import random
import numpy as np
from op_test import OpTest
class LinearChainCrfForward(object):
def __init__(self, seq_start_positions, emission_weights, emission_row_max,
emission_exps, transition_weights, transition_exps, labels):
self.tag_num = emission_weights.shape[1]
self.seq_num = len(seq_start_positions) - 1
self.seq_start_positions = seq_start_positions
self.labels = labels
self.x = emission_weights
self.x_row_max = emission_row_max
self.x_exps = emission_exps
# unnormalized logits of the transition weights for the start mark.
self.a = transition_weights[0, :]
self.a_exps = transition_exps[0, :]
# unnormalized logits of the transition weights for the end mark.
self.b = transition_weights[1, :]
self.b_exps = transition_exps[1, :]
# unnormalized logits of the transition weights for all the other tags.
self.w = transition_weights[2:, :]
self.w_exps = transition_exps[2:, :]
# The output of linear chain crf operator.
# alpha is a memo table in dynamic programming to caculate
# nomalization factor.
self.alpha = np.zeros(
(seq_start_positions[-1], self.tag_num), dtype="float64")
self.log_likelihood = np.zeros((self.seq_num, 1))
def _l1_norm(self, x):
s = np.sum(x)
x /= s
return s
def _forward_a_sequence(self, x, x_row_max, x_exps, label, alpha):
seq_len = x_row_max.shape[0]
log_likelihood = 0.
for i in range(self.tag_num):
alpha[0, i] = self.a_exps[i] * x_exps[0, i]
log_likelihood = -x_row_max[0] - np.log(self._l1_norm(alpha[0, :]))
# calculate the unnormalized logits of the normalization factor.
for k in range(1, seq_len):
for i in range(self.tag_num):
s = 0.
for j in range(self.tag_num):
s += alpha[k - 1, j] * self.w_exps[j, i]
alpha[k, i] = x_exps[k, i] * s
log_likelihood -= x_row_max[k] + np.log(self._l1_norm(alpha[k, :]))
s = 0.
for i in range(self.tag_num):
s += alpha[-1, i] * self.b_exps[i]
log_likelihood -= np.log(s)
# calculate the nominator part.
log_likelihood += (
self.a[label[0]] + x[0, label[0]] + self.b[label[-1]])
for k in range(1, seq_len):
log_likelihood += (x[k, label[k]] + self.w[label[k - 1], label[k]])
return -log_likelihood
def crf_forward_compute(self):
for i in range(self.seq_num):
start = self.seq_start_positions[i]
end = self.seq_start_positions[i + 1]
self.log_likelihood[i] = self._forward_a_sequence(
self.x[start:end, :], self.x_row_max[start:end, :],
self.x_exps[start:end, :], self.labels[start:end, :],
self.alpha[start:end, :])
return self.alpha, self.log_likelihood
class TestLinearChainCrfOp(OpTest):
def set_test_data(self):
# TODO(caoying) Fix the unittest by: add the boundary cases when
# sequence lengths are 1, 2, and 3.
SEQ_NUM = 3
TAG_NUM = 17
MAX_SEQ_LEN = 5
# the linear_chain_crf operator only supports sequence (LoD level = 1)
lod = [[0]]
for i in range(SEQ_NUM):
lod[-1].append(lod[-1][-1] + random.randint(1, MAX_SEQ_LEN))
emission = np.random.uniform(-1, 1,
[lod[-1][-1], TAG_NUM]).astype("float64")
emission_row_max = np.amax(emission, axis=1, keepdims=True)
emission_exps = np.exp(emission - emission_row_max)
transition = np.random.uniform(-0.5, 0.5,
[TAG_NUM + 2, TAG_NUM]).astype("float64")
transition_exps = np.exp(transition)
labels = np.random.randint(
low=0, high=TAG_NUM, size=(lod[-1][-1], 1), dtype="int32")
self.inputs = {
"Emission": (emission, lod),
"Transition": transition,
"Label": (labels, lod)
}
crf = LinearChainCrfForward(lod[0], emission, emission_row_max,
emission_exps, transition, transition_exps,
labels)
alpha, log_likelihood = crf.crf_forward_compute()
self.outputs = {
"Alpha": alpha,
"EmissionExps": emission_exps,
"TransitionExps": transition_exps,
"LogLikelihood": log_likelihood
}
def setUp(self):
self.op_type = "linear_chain_crf"
self.set_test_data()
def test_check_output(self):
self.check_output()
def test_check_grad(self):
self.check_grad(["Emission", "Transition"], "LogLikelihood")
def test_check_grad_ignore_transition(self):
self.check_grad(
["Emission"], "LogLikelihood", no_grad_set=set("Transition"))
if __name__ == "__main__":
unittest.main()
...@@ -52,7 +52,7 @@ def lstm( ...@@ -52,7 +52,7 @@ def lstm(
g = np.dot(h_pre, w_h) # 1 x 4D g = np.dot(h_pre, w_h) # 1 x 4D
g = g + x g = g + x
g = np.reshape(g, (1, g.size)) g = np.reshape(g, (1, g.size))
c_tmp, g_i, g_f, g_o = np.split(g, 4, axis=1) c, g_i, g_f, g_o = np.split(g, 4, axis=1)
if w_c is None: if w_c is None:
g_i = act_gate(g_i) # 1 x D g_i = act_gate(g_i) # 1 x D
g_f = act_gate(g_f) # 1 x D g_f = act_gate(g_f) # 1 x D
...@@ -60,7 +60,7 @@ def lstm( ...@@ -60,7 +60,7 @@ def lstm(
w_ic, w_fc, w_oc = np.split(w_c, 3, axis=1) w_ic, w_fc, w_oc = np.split(w_c, 3, axis=1)
g_i = act_gate(g_i + w_ic * c_pre) # 1 x D g_i = act_gate(g_i + w_ic * c_pre) # 1 x D
g_f = act_gate(g_f + w_fc * c_pre) # 1 x D g_f = act_gate(g_f + w_fc * c_pre) # 1 x D
c = g_f * c_pre + g_i * act_cand(c_tmp) # 1 x D c = g_f * c_pre + g_i * act_cand(c) # 1 x D
if w_c is None: if w_c is None:
g_o = act_gate(g_o) # 1 x D g_o = act_gate(g_o) # 1 x D
...@@ -68,8 +68,7 @@ def lstm( ...@@ -68,8 +68,7 @@ def lstm(
_, _, w_oc = np.split(w_c, 3, axis=1) _, _, w_oc = np.split(w_c, 3, axis=1)
g_o = act_gate(g_o + w_oc * c) # 1 x D g_o = act_gate(g_o + w_oc * c) # 1 x D
h = g_o * act_cell(c) h = g_o * act_cell(c)
bg = np.concatenate((act_cand(c_tmp), g_i, g_f, g_o), axis=1) return h, c
return h, c, bg
def _reverse(x, lod): def _reverse(x, lod):
y = np.zeros_like(x) y = np.zeros_like(x)
...@@ -82,7 +81,6 @@ def lstm( ...@@ -82,7 +81,6 @@ def lstm(
batch_size = len(offset) - 1 batch_size = len(offset) - 1
hidden = [] hidden = []
cell = [] cell = []
gate = []
input = _reverse(input, offset) if is_reverse else input input = _reverse(input, offset) if is_reverse else input
if w_b is not None: if w_b is not None:
input = input + np.tile(w_b, (offset[-1], 1)) input = input + np.tile(w_b, (offset[-1], 1))
...@@ -94,96 +92,109 @@ def lstm( ...@@ -94,96 +92,109 @@ def lstm(
c_pre = c0[i] # 1 x D c_pre = c0[i] # 1 x D
for j in range(seq_len): for j in range(seq_len):
# compute one step # compute one step
h_pre, c_pre, g_pre = _step(x[j], w_h, w_c, h_pre, c_pre, act_gate, h_pre, c_pre = _step(x[j], w_h, w_c, h_pre, c_pre, act_gate,
act_cell, act_cand) act_cell, act_cand)
hidden.append(h_pre.flatten()) hidden.append(h_pre.flatten())
cell.append(c_pre.flatten()) cell.append(c_pre.flatten())
gate.append(g_pre.flatten())
hidden = np.array(hidden).astype("float64") hidden = np.array(hidden).astype('float64')
cell = np.array(cell).astype("float64") cell = np.array(cell).astype('float64')
gate = np.array(gate).astype("float64")
hidden = _reverse(hidden, offset) if is_reverse else hidden hidden = _reverse(hidden, offset) if is_reverse else hidden
cell = _reverse(cell, offset) if is_reverse else cell cell = _reverse(cell, offset) if is_reverse else cell
assert gate.shape == input.shape
assert hidden.shape == (input.shape[0], input.shape[1] / 4) assert hidden.shape == (input.shape[0], input.shape[1] / 4)
assert cell.shape == (input.shape[0], input.shape[1] / 4) assert cell.shape == (input.shape[0], input.shape[1] / 4)
return hidden, cell, gate return hidden, cell
class TestLstmOp(OpTest): class TestLstmOp(OpTest):
def set_data(self): def set_argument(self):
self.lod = [[0, 2, 6, 9]] self.lod = [[0, 2, 5, 7]]
self.D = 64 self.D = 16
self.sort_idx = [2, 6, 0, 3, 7, 1, 4, 8, 5]
self.act_gate = "sigmoid" self.act_gate = 'sigmoid'
self.act_cell = "tanh" self.act_cell = 'tanh'
self.act_cand = "tanh" self.act_cand = 'tanh'
self.has_initial_state = True
self.is_reverse = False self.is_reverse = False
def setUp(self): def setUp(self):
self.set_data() self.set_argument()
self.op_type = "lstm" self.op_type = 'lstm'
T = self.lod[0][-1] T = self.lod[0][-1]
N = len(self.lod[0]) - 1 N = len(self.lod[0]) - 1
x = np.random.normal(size=(T, 4 * self.D)).astype("float64") x = np.random.normal(size=(T, 4 * self.D)).astype('float64')
h0 = np.zeros((N, self.D)).astype("float64") h0 = np.zeros((N, self.D)).astype('float64')
c0 = np.zeros((N, self.D)).astype("float64") c0 = np.zeros((N, self.D)).astype('float64')
w = np.random.normal(size=(self.D, 4 * self.D)).astype("float64") w = np.random.normal(size=(self.D, 4 * self.D)).astype('float64')
b = np.random.normal(size=(1, 7 * self.D)).astype("float64") b = np.random.normal(size=(1, 7 * self.D)).astype('float64')
w_b = b[:, 0:4 * self.D] w_b = b[:, 0:4 * self.D]
w_c = b[:, 4 * self.D:] w_c = b[:, 4 * self.D:]
h, c, g = lstm(x, self.lod, h0, c0, w, w_b, w_c, self.is_reverse, h, c = lstm(x, self.lod, h0, c0, w, w_b, w_c, self.is_reverse,
ACTVATION[self.act_gate], ACTVATION[self.act_cell], ACTVATION[self.act_gate], ACTVATION[self.act_cell],
ACTVATION[self.act_cand]) ACTVATION[self.act_cand])
g_sort = np.zeros_like(x) self.inputs = {'Input': (x, self.lod), 'Weight': w, 'Bias': b}
for i, j in enumerate(self.sort_idx): if self.has_initial_state:
g_sort[i, :] = g[j, :] self.inputs['H0'] = h0
self.inputs['C0'] = c0
self.inputs = {
'Input': (x, self.lod),
'H0': h0,
'C0': c0,
'Weight': w,
'Bias': b
}
self.outputs = { self.outputs = {
'Hidden': (h, self.lod), 'Hidden': (h, self.lod),
'Cell': (c, self.lod), 'Cell': (c, self.lod),
'BatchGate': g_sort
} }
self.attrs = { self.attrs = {
'usePeepholes': True, 'usePeepholes': True,
'isReverse': self.is_reverse, 'isReverse': self.is_reverse,
'gateActivation': 'sigmoid', 'gateActivation': self.act_gate,
'cellActivation': 'tanh', 'cellActivation': self.act_cell,
'candidateActivation': 'tanh' 'candidateActivation': self.act_cand
} }
def test_check_output(self): def test_check_output(self):
self.check_output() self.check_output(atol=1e-8)
#TODO(qingqing) add more unit testing case
def test_check_grad(self):
# TODO(qingqing) remove folowing lines after the check_grad is refined.
N = len(self.lod[0]) - 1
self.outputs['BatchGate'] = np.zeros((N, 4 * self.D)).astype('float64')
self.outputs['BatchCellPreAct'] = np.zeros(
(N, self.D)).astype('float64')
self.check_grad(
['Input', 'Weight', 'Bias'], ['Hidden'], max_relative_error=5e-4)
class TestLstmOpHasNoInitial(TestLstmOp):
def set_argument(self):
self.lod = [[0, 2, 5, 7]]
self.D = 16
self.act_gate = 'sigmoid'
self.act_cell = 'tanh'
self.act_cand = 'tanh'
self.has_initial_state = False
self.is_reverse = True
class TestLstmOpRerverse(TestLstmOp): class TestLstmOpRerverse(TestLstmOp):
def set_data(self): def set_argument(self):
self.lod = [[0, 2, 6, 9]] self.lod = [[0, 2, 5, 7]]
self.D = 64 self.D = 16
self.sort_idx = [2, 6, 0, 3, 7, 1, 4, 8, 5]
self.act_gate = "sigmoid" self.act_gate = 'sigmoid'
self.act_cell = "tanh" self.act_cell = 'tanh'
self.act_cand = "tanh" self.act_cand = 'tanh'
self.has_initial_state = True
self.is_reverse = True self.is_reverse = True
if __name__ == "__main__": if __name__ == '__main__':
unittest.main() unittest.main()
...@@ -3,15 +3,6 @@ import numpy as np ...@@ -3,15 +3,6 @@ import numpy as np
from op_test import OpTest from op_test import OpTest
class SeqPoolType(OpTest):
AVERAGE = 0
SUM = 1
SQRT = 2
MAX = 3
LAST = 4
FIRST = 5
class TestSeqAvgPool(OpTest): class TestSeqAvgPool(OpTest):
def set_data(self): def set_data(self):
self.op_type = 'sequence_pool' self.op_type = 'sequence_pool'
...@@ -25,7 +16,7 @@ class TestSeqAvgPool(OpTest): ...@@ -25,7 +16,7 @@ class TestSeqAvgPool(OpTest):
return x, lod, out return x, lod, out
def compute(self, x, lod, out): def compute(self, x, lod, out):
self.attrs = {'strategy': SeqPoolType.AVERAGE} self.attrs = {'pooltype': "AVERAGE"}
for i in range(4): for i in range(4):
sub_x = x[lod[0][i]:lod[0][i + 1], :] sub_x = x[lod[0][i]:lod[0][i + 1], :]
out[i] = sub_x.mean(axis=0) out[i] = sub_x.mean(axis=0)
...@@ -54,7 +45,7 @@ class TestSeqAvgPool2D(TestSeqAvgPool): ...@@ -54,7 +45,7 @@ class TestSeqAvgPool2D(TestSeqAvgPool):
return x, lod, out return x, lod, out
def compute(self, x, lod, out): def compute(self, x, lod, out):
self.attrs = {'strategy': SeqPoolType.AVERAGE} self.attrs = {'pooltype': "AVERAGE"}
for i in range(4): for i in range(4):
sub_x = np.reshape(x[lod[0][i]:lod[0][i + 1], :], (-1, 3 * 17)) sub_x = np.reshape(x[lod[0][i]:lod[0][i + 1], :], (-1, 3 * 17))
out[i] = np.reshape(sub_x.mean(axis=0), (3, 17)) out[i] = np.reshape(sub_x.mean(axis=0), (3, 17))
...@@ -62,7 +53,7 @@ class TestSeqAvgPool2D(TestSeqAvgPool): ...@@ -62,7 +53,7 @@ class TestSeqAvgPool2D(TestSeqAvgPool):
class TestSeqSumPool(TestSeqAvgPool): class TestSeqSumPool(TestSeqAvgPool):
def compute(self, x, lod, out): def compute(self, x, lod, out):
self.attrs = {'strategy': SeqPoolType.SUM} self.attrs = {'pooltype': "SUM"}
for i in range(4): for i in range(4):
sub_x = x[lod[0][i]:lod[0][i + 1], :] sub_x = x[lod[0][i]:lod[0][i + 1], :]
out[i] = sub_x.sum(axis=0) out[i] = sub_x.sum(axis=0)
...@@ -70,7 +61,7 @@ class TestSeqSumPool(TestSeqAvgPool): ...@@ -70,7 +61,7 @@ class TestSeqSumPool(TestSeqAvgPool):
class TestSeqSumPool2D(TestSeqAvgPool2D): class TestSeqSumPool2D(TestSeqAvgPool2D):
def compute(self, x, lod, out): def compute(self, x, lod, out):
self.attrs = {'strategy': SeqPoolType.SUM} self.attrs = {'pooltype': "SUM"}
for i in range(4): for i in range(4):
sub_x = np.reshape(x[lod[0][i]:lod[0][i + 1], :], (-1, 3 * 17)) sub_x = np.reshape(x[lod[0][i]:lod[0][i + 1], :], (-1, 3 * 17))
out[i] = np.reshape(sub_x.sum(axis=0), (3, 17)) out[i] = np.reshape(sub_x.sum(axis=0), (3, 17))
...@@ -78,7 +69,7 @@ class TestSeqSumPool2D(TestSeqAvgPool2D): ...@@ -78,7 +69,7 @@ class TestSeqSumPool2D(TestSeqAvgPool2D):
class TestSeqSqrtPool(TestSeqAvgPool): class TestSeqSqrtPool(TestSeqAvgPool):
def compute(self, x, lod, out): def compute(self, x, lod, out):
self.attrs = {'strategy': SeqPoolType.SQRT} self.attrs = {'pooltype': "SQRT"}
for i in range(4): for i in range(4):
sub_x = x[lod[0][i]:lod[0][i + 1], :] sub_x = x[lod[0][i]:lod[0][i + 1], :]
len = lod[0][i + 1] - lod[0][i] len = lod[0][i + 1] - lod[0][i]
...@@ -87,7 +78,7 @@ class TestSeqSqrtPool(TestSeqAvgPool): ...@@ -87,7 +78,7 @@ class TestSeqSqrtPool(TestSeqAvgPool):
class TestSeqSqrtPool2D(TestSeqAvgPool2D): class TestSeqSqrtPool2D(TestSeqAvgPool2D):
def compute(self, x, lod, out): def compute(self, x, lod, out):
self.attrs = {'strategy': SeqPoolType.SQRT} self.attrs = {'pooltype': "SQRT"}
for i in range(4): for i in range(4):
sub_x = np.reshape(x[lod[0][i]:lod[0][i + 1], :], (-1, 3 * 17)) sub_x = np.reshape(x[lod[0][i]:lod[0][i + 1], :], (-1, 3 * 17))
len = lod[0][i + 1] - lod[0][i] len = lod[0][i + 1] - lod[0][i]
...@@ -99,7 +90,7 @@ class TestSeqSqrtPool2D(TestSeqAvgPool2D): ...@@ -99,7 +90,7 @@ class TestSeqSqrtPool2D(TestSeqAvgPool2D):
class TestSeqMaxPool(TestSeqAvgPool): class TestSeqMaxPool(TestSeqAvgPool):
def compute(self, x, lod, out): def compute(self, x, lod, out):
self.attrs = {'strategy': SeqPoolType.MAX} self.attrs = {'pooltype': "MAX"}
for i in range(4): for i in range(4):
sub_x = x[lod[0][i]:lod[0][i + 1], :] sub_x = x[lod[0][i]:lod[0][i + 1], :]
out[i] = np.amax(sub_x, axis=0) out[i] = np.amax(sub_x, axis=0)
...@@ -111,7 +102,7 @@ class TestSeqMaxPool(TestSeqAvgPool): ...@@ -111,7 +102,7 @@ class TestSeqMaxPool(TestSeqAvgPool):
class TestSeqMaxPool2D(TestSeqAvgPool2D): class TestSeqMaxPool2D(TestSeqAvgPool2D):
def compute(self, x, lod, out): def compute(self, x, lod, out):
self.attrs = {'strategy': SeqPoolType.MAX} self.attrs = {'pooltype': "MAX"}
for i in range(4): for i in range(4):
sub_x = np.reshape(x[lod[0][i]:lod[0][i + 1], :], (-1, 3 * 17)) sub_x = np.reshape(x[lod[0][i]:lod[0][i + 1], :], (-1, 3 * 17))
out[i] = np.reshape(np.amax(sub_x, axis=0), (3, 17)) out[i] = np.reshape(np.amax(sub_x, axis=0), (3, 17))
...@@ -123,7 +114,7 @@ class TestSeqMaxPool2D(TestSeqAvgPool2D): ...@@ -123,7 +114,7 @@ class TestSeqMaxPool2D(TestSeqAvgPool2D):
class TestSeqLastPool(TestSeqAvgPool): class TestSeqLastPool(TestSeqAvgPool):
def compute(self, x, lod, out): def compute(self, x, lod, out):
self.attrs = {'strategy': SeqPoolType.LAST} self.attrs = {'pooltype': "LAST"}
for i in range(4): for i in range(4):
sub_x = x[lod[0][i]:lod[0][i + 1], :] sub_x = x[lod[0][i]:lod[0][i + 1], :]
out[i] = sub_x[-1, :] out[i] = sub_x[-1, :]
...@@ -131,7 +122,7 @@ class TestSeqLastPool(TestSeqAvgPool): ...@@ -131,7 +122,7 @@ class TestSeqLastPool(TestSeqAvgPool):
class TestSeqLastPool2D(TestSeqAvgPool2D): class TestSeqLastPool2D(TestSeqAvgPool2D):
def compute(self, x, lod, out): def compute(self, x, lod, out):
self.attrs = {'strategy': SeqPoolType.LAST} self.attrs = {'pooltype': "LAST"}
for i in range(4): for i in range(4):
sub_x = np.reshape(x[lod[0][i]:lod[0][i + 1], :], (-1, 3 * 17)) sub_x = np.reshape(x[lod[0][i]:lod[0][i + 1], :], (-1, 3 * 17))
out[i] = np.reshape(sub_x[-1, :], (3, 17)) out[i] = np.reshape(sub_x[-1, :], (3, 17))
...@@ -139,7 +130,7 @@ class TestSeqLastPool2D(TestSeqAvgPool2D): ...@@ -139,7 +130,7 @@ class TestSeqLastPool2D(TestSeqAvgPool2D):
class TestSeqFirstPool(TestSeqAvgPool): class TestSeqFirstPool(TestSeqAvgPool):
def compute(self, x, lod, out): def compute(self, x, lod, out):
self.attrs = {'strategy': SeqPoolType.FIRST} self.attrs = {'pooltype': "FIRST"}
for i in range(4): for i in range(4):
sub_x = x[lod[0][i]:lod[0][i + 1], :] sub_x = x[lod[0][i]:lod[0][i + 1], :]
out[i] = sub_x[0, :] out[i] = sub_x[0, :]
...@@ -147,7 +138,7 @@ class TestSeqFirstPool(TestSeqAvgPool): ...@@ -147,7 +138,7 @@ class TestSeqFirstPool(TestSeqAvgPool):
class TestSeqFirstPool2D(TestSeqAvgPool2D): class TestSeqFirstPool2D(TestSeqAvgPool2D):
def compute(self, x, lod, out): def compute(self, x, lod, out):
self.attrs = {'strategy': SeqPoolType.FIRST} self.attrs = {'pooltype': "FIRST"}
for i in range(4): for i in range(4):
sub_x = np.reshape(x[lod[0][i]:lod[0][i + 1], :], (-1, 3 * 17)) sub_x = np.reshape(x[lod[0][i]:lod[0][i + 1], :], (-1, 3 * 17))
out[i] = np.reshape(sub_x[0, :], (3, 17)) out[i] = np.reshape(sub_x[0, :], (3, 17))
......
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册