Add maxout operator.

paddle/operators/CMakeLists.txt

@@ -162,6 +162,7 @@ set(DEPS_OPS
     softmax_with_cross_entropy_op
     sum_op
     pool_op
+    maxout_op
     pool_with_index_op
     conv_op
     lstm_op
@@ -182,6 +183,7 @@ op_library(softmax_with_cross_entropy_op DEPS cross_entropy softmax)
 op_library(conv_op DEPS vol2col)
 op_library(sum_op DEPS net_op selected_rows_functor)
 op_library(pool_op DEPS pooling)
+op_library(maxout_op DEPS maxouting)
 op_library(pool_with_index_op DEPS pooling)
 op_library(lod_rank_table_op SRCS lod_rank_table_op.cc DEPS lod_rank_table)
 op_library(lod_tensor_to_array_op SRCS lod_tensor_to_array_op.cc DEPS lod_rank_table_op)

paddle/operators/math/CMakeLists.txt

@@ -8,6 +8,7 @@ if(WITH_GPU)
     nv_library(softmax SRCS softmax.cc softmax.cu DEPS operator)
     nv_library(cross_entropy SRCS cross_entropy.cc cross_entropy.cu DEPS operator)
     nv_library(pooling SRCS pooling.cc pooling.cu DEPS device_context)
+    nv_library(maxouting SRCS maxouting.cc maxouting.cu DEPS device_context)
     nv_library(sequence_pooling SRCS sequence_pooling.cc sequence_pooling.cu DEPS device_context math_function)
     nv_library(vol2col SRCS vol2col.cc vol2col.cu DEPS device_context)
     nv_library(context_project SRCS context_project.cc context_project.cu DEPS device_context)
@@ -20,6 +21,7 @@ else()
     cc_library(softmax SRCS softmax.cc DEPS operator)
     cc_library(cross_entropy SRCS cross_entropy.cc DEPS operator)
     cc_library(pooling SRCS pooling.cc DEPS device_context)
+    cc_library(maxouting SRCS maxouting.cc DEPS device_context)
     cc_library(sequence_pooling SRCS sequence_pooling.cc DEPS device_context math_function)
     cc_library(vol2col SRCS vol2col.cc DEPS device_context)
     cc_library(context_project SRCS context_project.cc DEPS device_context)

paddle/operators/math/maxouting.cc

@@ -0,0 +1,117 @@
+/* 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/math/maxouting.h"
+
+namespace paddle {
+namespace operators {
+namespace math {
+
+/*
+ * All tensors are in NCHW format.
+ * Ksize, strides, paddings are two elements. These two elements represent
+ * height and width, respectively.
+ */
+template <typename MaxOutProcess, typename T>
+class MaxOutFunctor<platform::CPUPlace, MaxOutProcess, T> {
+ public:
+  void operator()(const platform::DeviceContext& context,
+                  const framework::Tensor& input, framework::Tensor& output,
+                  int groups, int num_channels, MaxOutProcess maxout_process) {
+    const int batch_size = input.dims()[0];
+    const int input_height = input.dims()[2];
+    const int input_width = input.dims()[3];
+    const int output_channels = num_channels/groups;
+
+    int fea_size = input_height * input_width;
+    int c_size = fea_size * output_channels;
+
+    const T* input_data = input.data<T>();
+    T* output_data = output.mutable_data<T>(context.GetPlace());
+
+    for (int i = 0; i < batch_size; i++) {
+      int new_bindex =  c_size * i;
+      for (int c = 0; c < output_channels; ++c) {
+        int new_cindex = fea_size * c;
+        for (int f = 0; f < fea_size; f++) {
+          T ele = maxout_process.initial();
+          for (int ph = 0; ph < groups; ++ph) {
+            maxout_process.compute(ele,
+              input_data[(new_bindex+new_cindex) * groups+ph*fea_size+f]);
+          }
+          maxout_process.finalize(ele, (static_cast<T>(groups)));
+          output_data[(new_bindex+new_cindex+f)] = ele;
+        }
+      }
+    }
+  }
+};
+
+
+
+template <class T>
+class MaxOutGradFunctor<platform::CPUPlace, T> {
+public:
+  void operator()(const platform::DeviceContext& context,
+                  const framework::Tensor& input,
+                  framework::Tensor& input_grad,
+                  const framework::Tensor& output,
+                  const framework::Tensor& output_grad,
+                  int groups, int num_channels) {
+    const int batch_size = input.dims()[0];
+    const int input_height = input.dims()[2];
+    const int input_width = input.dims()[3];
+    const int output_channels = num_channels / groups;
+
+    int fea_size = input_height * input_width;
+
+    const T* input_data = input.data<T>();
+    const T* output_data = output.data<T>();
+    const T* output_grad_data = output_grad.data<T>();
+    T* input_grad_data = input_grad.mutable_data<T>(context.GetPlace());
+
+    for (int i = 0; i < batch_size; i++) {
+      int blen = fea_size * output_channels * i;
+      for (int c = 0; c < output_channels; ++c) {
+        int clen = fea_size * c;
+        for (int f = 0; f < fea_size; f++) {
+          int input_idx = 0;
+          bool stop = false;
+          int output_idx = blen + clen + f;
+          for (int g = 0; g < groups && !stop; g++) {
+              input_idx = (blen + clen) * groups + fea_size * g + f;
+              input_grad_data[input_idx] = 0;
+              if (input_data[input_idx] == output_data[output_idx]) {
+                input_grad_data[input_idx] += output_grad_data[output_idx];
+                stop = true;
+              } else {
+                input_grad_data[input_idx] = 0;
+              }
+          }
+        }
+      }
+    }
+  }
+};
+
+template class MaxOutGradFunctor<platform::CPUPlace, float>;
+template class MaxOutGradFunctor<platform::CPUPlace, double>;
+template class MaxOutFunctor<platform::CPUPlace,
+                             paddle::operators::math::MaxOut<float>, float>;
+template class MaxOutFunctor<platform::CPUPlace,
+                             paddle::operators::math::MaxOut<double>, double>;
+
+}  // namespace math
+}  // namespace operators
+}  // namespace paddle

paddle/operators/math/maxouting.cu

@@ -0,0 +1,161 @@
+/* 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/math/maxouting.h"
+#include "paddle/platform/cuda_helper.h"
+
+namespace paddle {
+namespace operators {
+namespace math {
+
+template <typename MaxOutProcess, typename T>
+__global__ void KernelMaxOut(const int nthreads, const T* input_data,
+                             T* output_data, const int channels,
+                             const int input_height, const int input_width,
+                             int groups, MaxOutProcess maxout_process) {
+  int size = input_height * input_width * channels / groups;
+  int featLen = input_height * input_width;
+  for (int index = blockIdx.x * blockDim.x + threadIdx.x; index < nthreads;
+              index += blockDim.x * gridDim.x) {
+    int batch_idx = index / size;
+    int i = index % size;
+    int channel_idx = i / featLen;
+    int feat_idx = i % featLen;
+    int data_idx =
+      (batch_idx * size + channel_idx * featLen) * groups + feat_idx;
+    T ele = maxout_process.initial();
+    for (int g = 0; g < groups; g++) {
+      maxout_process.compute(ele, input_data[data_idx + g * featLen]);
+    }
+    maxout_process.finalize(ele, (static_cast<T>(groups)));
+    output_data[index] = ele;
+  }
+}
+template <typename T>
+__global__ void KernelMaxoutGrad(
+    const int nthreads, const T* input_data, const T* output_data,
+    const T* output_grad, T* input_grad, const int channels,
+    const int input_height, const int input_width, int groups) {
+    int size = input_height * input_width * channels / groups;
+    int featLen = input_height * input_width;
+    for (int index = blockIdx.x * blockDim.x + threadIdx.x; index < nthreads;
+         index += blockDim.x * gridDim.x) {
+      int batch_idx = index / size;
+      int i = index % size;
+      int channel_idx = i / featLen;
+      int feat_idx = i % featLen;
+      int data_idx =
+        (batch_idx * size + channel_idx * featLen) * groups + feat_idx;
+      int maxIndex = -1;
+      bool stop = false;
+      for (int g = 0; g < groups && !stop; g++) {
+        if (input_data[data_idx + g * featLen] == output_data[index]) {
+          maxIndex = data_idx + g * featLen;
+          stop = true;
+        }
+      }
+      if (maxIndex != -1) {
+        // atomic add
+        platform::CudaAtomicAdd(input_grad + maxIndex, output_grad[index]);
+      }
+    }
+}
+/*
+ * All tensors are in NCHW format.
+ * Ksize, strides, paddings are two elements. These two elements represent
+ * height and width, respectively.
+ */
+template <typename MaxOutProcess, typename T>
+class MaxOutFunctor<platform::GPUPlace, MaxOutProcess, T> {
+ public:
+  void operator()(const platform::DeviceContext& context,
+                  const framework::Tensor& input, framework::Tensor& output,
+                  int groups, int num_channels,
+                  MaxOutProcess maxout_process) {
+    const int batch_size = input.dims()[0];
+    const int input_channels = input.dims()[1];
+    const int input_height = input.dims()[2];
+    const int input_width = input.dims()[3];
+    const int output_channels = num_channels / groups;
+    const int output_height = output.dims()[2];
+    const int output_width = output.dims()[3];
+
+    const T* input_data = input.data<T>();
+    T* output_data = output.mutable_data<T>(context.GetPlace());
+
+    int nthreads = batch_size * output_channels * output_height * output_width;
+    int blocks = (nthreads + 1024 - 1) / 1024;
+    dim3 threads(1024, 1);
+    dim3 grid(blocks, 1);
+
+    KernelMaxOut<
+        MaxOutProcess,
+        T><<<grid, threads, 0,
+             reinterpret_cast<const platform::CUDADeviceContext&>(context)
+                 .stream()>>>(nthreads, input_data, output_data, input_channels,
+                              input_height, input_width, groups,
+                              maxout_process);
+  }
+};
+/*
+ * All tensors are in NCHW format.
+ * Ksize, strides, paddings are two elements. These two elements represent
+ * height and width, respectively.
+ */
+template <typename T>
+class MaxOutGradFunctor<platform::GPUPlace, T> {
+ public:
+  void operator()(const platform::DeviceContext& context,
+                  const framework::Tensor& input, framework::Tensor& input_grad,
+                  const framework::Tensor& output,
+                  const framework::Tensor& output_grad,
+                  int groups, int num_channels) {
+    const int batch_size = input.dims()[0];
+    const int input_channels = input.dims()[1];
+    const int input_height = input.dims()[2];
+    const int input_width = input.dims()[3];
+    const int output_channels = output.dims()[1];
+    const int output_height = output.dims()[2];
+    const int output_width = output.dims()[3];
+
+    const T* input_data = input.data<T>();
+    const T* output_data = output.data<T>();
+    const T* output_grad_data = output_grad.data<T>();
+    T* input_grad_data = input_grad.mutable_data<T>(context.GetPlace());
+
+    int nthreads = batch_size * output_channels * output_height * output_width;
+    int blocks = (nthreads + 1024 - 1) / 1024;
+    dim3 threads(1024, 1);
+    dim3 grid(blocks, 1);
+
+    KernelMaxoutGrad<
+        T><<<grid, threads, 0,
+             reinterpret_cast<const platform::CUDADeviceContext&>(context)
+                 .stream()>>>(
+        nthreads, input_data, output_data, output_grad_data, input_grad_data,
+        input_channels, input_height, input_width, groups);
+  }
+};
+
+template class MaxOutGradFunctor<platform::GPUPlace, float>;
+template class MaxOutGradFunctor<platform::GPUPlace, double>;
+
+template class MaxOutFunctor<platform::GPUPlace,
+                             paddle::operators::math::MaxOut<float>, float>;
+template class MaxOutFunctor<platform::GPUPlace,
+                             paddle::operators::math::MaxOut<double>, double>;
+
+}  // namespace math
+}  // namespace operators
+}  // namespace paddle