Maxpool

src/cudamatrix/cu-kernels.cu

@@ -773,6 +773,61 @@ static void _add_mat_blocks_trans(Real alpha, const Real* src,
     }
 }
 
+template<typename Real>
+__global__
+static void _max_mat_blocks(Real alpha, const Real* src,
+                            int32_cuda num_row_blocks,
+                            int32_cuda num_col_blocks, Real* dst, MatrixDim d,
+                            int src_stride) {
+  int32_cuda i = blockIdx.x * blockDim.x + threadIdx.x;
+  int32_cuda j = blockIdx.y * blockDim.y + threadIdx.y;
+  int32_cuda index = i + j * d.stride;
+  int32_cuda index_src = i + j * src_stride;
+  if (i < d.cols && j < d.rows)
+    for (int32_cuda p = 0; p < num_row_blocks; p++) {
+      for (int32_cuda q = 0; q < num_col_blocks; q++) {
+        dst[index] = fmax(
+            src[index_src + p * src_stride * d.rows + q * d.cols],
+            dst[index]);
+      }
+    }
+}
+
+template<typename Real>
+__global__
+static void _max_mat_blocks_trans(Real alpha, const Real* src,
+                                  int32_cuda num_row_blocks,
+                                  int32_cuda num_col_blocks, Real* dst,
+                                  MatrixDim d, int src_stride) {
+  int32_cuda i = blockIdx.x * blockDim.x + threadIdx.x;
+  int32_cuda j = blockIdx.y * blockDim.y + threadIdx.y;
+  int32_cuda index = i + j * d.stride;
+  int32_cuda index_src = j + i * src_stride;
+  if (i < d.cols && j < d.rows)
+    for (int32_cuda p = 0; p < num_row_blocks; p++) {
+      for (int32_cuda q = 0; q < num_col_blocks; q++) {
+        dst[index] = fmax(
+            src[index_src + p * src_stride * d.cols + q * d.rows],
+            dst[index]);
+      }
+    }
+}
+
+template<typename Real>
+__global__
+static void _max_mat_repeated(Real alpha, const Real* src,
+                              MatrixDim src_dim, Real* dst,
+                              MatrixDim dst_dim) {
+  int32_cuda i = blockIdx.x * blockDim.x + threadIdx.x;
+  int32_cuda j = blockIdx.y * blockDim.y + threadIdx.y;
+  int32_cuda src_i = i % src_dim.cols,
+      src_j = j % src_dim.rows,
+      dst_index = i + j * dst_dim.stride,
+      src_index = src_i + src_j * src_dim.stride;
+  if (i < dst_dim.cols && j < dst_dim.rows)
+    dst[dst_index] += alpha * src[src_index];
+}
+
 template<typename Real>
 __global__
 static void _set_mat_mat_div_mat(const Real* A, const Real* B, const Real* C,
@@ -3952,11 +4007,28 @@ void cudaF_add_mat_blocks(dim3 Gr, dim3 Bl, float alpha, const float* src,
   }
 }
 
+void cudaF_max_mat_blocks(dim3 Gr, dim3 Bl, float alpha, const float* src,
+                          int32_cuda num_row_blocks, int32_cuda num_col_blocks,
+                          float* dst, MatrixDim d, int src_stride,
+                          int A_trans) {
+  if (A_trans) {
+    _max_mat_blocks_trans<<<Gr,Bl>>>(alpha, src, num_row_blocks, num_col_blocks,
+        dst, d, src_stride);
+  } else {
+    _max_mat_blocks<<<Gr,Bl>>>(alpha, src, num_row_blocks, num_col_blocks, dst,
+        d, src_stride);
+  }
+}
+
 void cudaF_add_mat_repeated(dim3 Gr, dim3 Bl, float alpha, const float* src,
                             MatrixDim src_dim, float *dst, MatrixDim dst_dim) {
   _add_mat_repeated<<<Gr,Bl>>>(alpha, src, src_dim, dst, dst_dim);
 }
 
+void cudaF_max_mat_repeated(dim3 Gr, dim3 Bl, float alpha, const float* src,
+                            MatrixDim src_dim, float *dst, MatrixDim dst_dim) {
+  _max_mat_repeated<<<Gr,Bl>>>(alpha, src, src_dim, dst, dst_dim);
+}
 
 void cudaF_set_mat_mat_div_mat(dim3 Gr, dim3 Bl, const float *A, const float *B,
                                const float *C, float *dst, MatrixDim d,
@@ -4656,11 +4728,29 @@ void cudaD_add_mat_blocks(dim3 Gr, dim3 Bl, double alpha, const double* src,
   }
 }
 
+void cudaD_max_mat_blocks(dim3 Gr, dim3 Bl, double alpha, const double* src,
+                          int32_cuda num_row_blocks, int32_cuda num_col_blocks,
+                          double* dst, MatrixDim d, int src_stride,
+                          int A_trans) {
+  if (A_trans) {
+    _max_mat_blocks_trans<<<Gr,Bl>>>(alpha, src, num_row_blocks, num_col_blocks,
+        dst, d, src_stride);
+  } else {
+    _max_mat_blocks<<<Gr,Bl>>>(alpha, src, num_row_blocks, num_col_blocks, dst,
+        d, src_stride);
+  }
+}
+
 void cudaD_add_mat_repeated(dim3 Gr, dim3 Bl, double alpha, const double* src,
                             MatrixDim src_dim, double *dst, MatrixDim dst_dim) {
   _add_mat_repeated<<<Gr,Bl>>>(alpha, src, src_dim, dst, dst_dim);
 }
 
+void cudaD_max_mat_repeated(dim3 Gr, dim3 Bl, double alpha, const double* src,
+                            MatrixDim src_dim, double *dst, MatrixDim dst_dim) {
+  _max_mat_repeated<<<Gr,Bl>>>(alpha, src, src_dim, dst, dst_dim);
+}
+
 void cudaD_set_mat_mat_div_mat(dim3 Gr, dim3 Bl, const double *A,
                                const double *B, const double *C, double *dst,
                                MatrixDim d, int stride_a, int stride_b,

src/cudamatrix/cu-kernels.h

@@ -168,6 +168,30 @@ inline void cuda_add_mat_repeated(dim3 Gr, dim3 Bl, float alpha,
                                   float *dst, MatrixDim dst_dim) {
   cudaF_add_mat_repeated(Gr, Bl, alpha, src, src_dim, dst, dst_dim);
 }
+inline void cuda_max_mat_blocks(dim3 Gr, dim3 Bl, double alpha,
+                                const double *src, int32_cuda num_row_blocks,
+                                int32_cuda num_col_blocks, double *dst,
+                                MatrixDim d, int src_stride, int A_trans) {
+  cudaD_max_mat_blocks(Gr, Bl, alpha, src, num_row_blocks, num_col_blocks, dst,
+                       d, src_stride, A_trans);
+}
+inline void cuda_max_mat_blocks(dim3 Gr, dim3 Bl, float alpha, const float *src,
+                                int32_cuda num_row_blocks,
+                                int32_cuda num_col_blocks, float *dst,
+                                MatrixDim d, int src_stride, int A_trans) {
+  cudaF_max_mat_blocks(Gr, Bl, alpha, src, num_row_blocks, num_col_blocks, dst,
+                       d, src_stride, A_trans);
+}
+inline void cuda_max_mat_repeated(dim3 Gr, dim3 Bl, double alpha,
+                                  const double *src, MatrixDim src_dim,
+                                  double *dst, MatrixDim dst_dim) {
+  cudaD_max_mat_repeated(Gr, Bl, alpha, src, src_dim, dst, dst_dim);
+}
+inline void cuda_max_mat_repeated(dim3 Gr, dim3 Bl, float alpha,
+                                  const float *src, MatrixDim src_dim,
+                                  float *dst, MatrixDim dst_dim) {
+  cudaF_max_mat_repeated(Gr, Bl, alpha, src, src_dim, dst, dst_dim);
+}
 inline void cuda_add_mat_diag_vec(dim3 Gr, dim3 Bl, double alpha, double *mat,
                                   MatrixDim mat_dim, const double *mat2,
                                   int mat2_row_stride, int mat2_col_stride,

src/cudamatrix/cu-matrix.cc

@@ -1187,6 +1187,92 @@ void CuMatrixBase<Real>::AddMatBlocks(Real alpha, const CuMatrixBase<Real> &A,
   }
 }
 
+
+template<typename Real>
+void CuMatrixBase<Real>::MaxMatBlocks(Real alpha, const CuMatrixBase<Real> &A,
+                                      MatrixTransposeType transA) {
+  if (num_rows_ == 0 || num_cols_ == 0) return;
+
+  if (A.NumRows() >= (transA == kNoTrans ? num_rows_ : num_cols_) &&
+      A.NumCols() >= (transA == kNoTrans ? num_cols_ : num_rows_)) {
+    // This is the "summing", not broadcasting, version of AddMatBlocks.
+    // It supports both regular and transposed operation.
+    int32 num_row_blocks, num_col_blocks;
+    if (transA == kNoTrans) {
+      KALDI_ASSERT(A.NumRows() % num_rows_ == 0 && A.NumCols() % num_cols_ == 0);
+      num_row_blocks = A.Mat().NumRows() / num_rows_;
+      num_col_blocks = A.Mat().NumCols() / num_cols_;
+    } else {
+      KALDI_ASSERT(A.NumRows() % num_cols_ == 0 && A.NumCols() % num_rows_ == 0);
+      num_row_blocks = A.Mat().NumRows() / num_cols_;
+      num_col_blocks = A.Mat().NumCols() / num_rows_;
+    }
+#if HAVE_CUDA == 1
+    if (CuDevice::Instantiate().Enabled()) {
+      CuTimer tim;
+      dim3 dimGrid, dimBlock;
+      GetBlockSizesForSimpleMatrixOperation(NumRows(), NumCols(),
+                                            &dimGrid, &dimBlock);
+      cuda_max_mat_blocks(dimGrid, dimBlock, alpha, A.data_, num_row_blocks,
+                          num_col_blocks, data_, Dim(), A.Stride(),
+                          (transA == kTrans ? 1 : 0));
+      CU_SAFE_CALL(cudaGetLastError());
+
+      CuDevice::Instantiate().AccuProfile(__func__, tim);
+    } //else
+#endif
+    // {
+    //   int32 nr, nc;
+    //   if (transA == kNoTrans) {
+    //     nr = num_rows_;
+    //     nc = num_cols_;
+    //   } else {
+    //     nr = num_cols_;
+    //     nc = num_rows_;
+    //   }
+    //   for (int32 i = 0; i < num_row_blocks; i++) {
+    //     for (int32 j = 0; j < num_col_blocks; j++) {
+    //       Mat().AddMat(alpha, SubMatrix<Real>(A.Mat(), i * nr, nr, j * nc, nc),
+    //                    transA);
+    //     }
+    //   }
+    // }
+  } else {
+    // This is the "broadcasting" version of MaxMatBlocks, where
+    // *this is larger than src.
+    if (transA != kNoTrans)
+      KALDI_ERR << "Transposed operation not supported currently.";
+    if (!(num_rows_ % A.NumRows() == 0 && num_cols_ % A.NumCols() == 0))
+      KALDI_ERR << "Invalid sizes of arguments";
+#if HAVE_CUDA == 1
+    if (CuDevice::Instantiate().Enabled()) {
+      CuTimer tim;
+      dim3 dimGrid, dimBlock;
+      GetBlockSizesForSimpleMatrixOperation(NumRows(), NumCols(),
+                                            &dimGrid, &dimBlock);
+      cuda_max_mat_repeated(dimGrid, dimBlock, alpha,
+                            A.data_, A.Dim(), data_, Dim());
+      CU_SAFE_CALL(cudaGetLastError());
+      CuDevice::Instantiate().AccuProfile(__func__, tim);
+    } //else
+#endif
+//     {
+//       const MatrixBase<Real> &src_mat = A.Mat(),
+//           &this_mat = this->Mat();
+//       for (int32 row_offset = 0; row_offset < NumRows();
+//            row_offset += src_mat.NumRows()) {
+//         for (int32 col_offset = 0; col_offset < NumCols();
+//              col_offset += src_mat.NumCols()) {
+//           SubMatrix<Real> this_part(this_mat,
+//                                     row_offset, src_mat.NumRows(),
+//                                     col_offset, src_mat.NumCols());
+//           this_part.AddMat(alpha, src_mat);
+//         }
+//       }
+//     }
+  }
+}
+
 /// dst = a * b / c (by element; when c = 0, dst = a)
 /// dst can be an alias of a, b or c safely and get expected result.
 template<typename Real>

src/cudamatrix/cu-matrix.h

@@ -509,6 +509,9 @@ class CuMatrixBase {
   void AddMatBlocks(Real alpha, const CuMatrixBase<Real> &A,
                     MatrixTransposeType trans = kNoTrans);
 
+  void MaxMatBlocks(Real alpha, const CuMatrixBase<Real> &A,
+                    MatrixTransposeType trans = kNoTrans);
+
   /// (for each column c of *this), c = alpha * col + beta * c
   void AddVecToCols(Real alpha, const CuVectorBase<Real> &col, Real beta = 1.0);
   /// (for each row r of *this), r = alpha * row + beta * r

src/nnet3/nnet-simple-component.cc

@@ -5860,6 +5860,77 @@ void SumBlockComponent::Backprop(
   }
 }
 
+MaxPoolingOverBlock::MaxPoolingOverBlock(const MaxPoolingOverBlock &other):
+    input_dim_(other.input_dim_), output_dim_(other.output_dim_),
+    scale_(other.scale_) { }
+
+void MaxPoolingOverBlock::InitFromConfig(ConfigLine *cfl) {
+  scale_ = 1.0;
+  bool ok = cfl->GetValue("input-dim", &input_dim_) &&
+      cfl->GetValue("output-dim", &output_dim_);
+  if (!ok)
+    KALDI_ERR << "input-dim and output-dim must both be provided.";
+  if (input_dim_ <= 0 || input_dim_ % output_dim_ != 0)
+    KALDI_ERR << "Invalid values input-dim=" << input_dim_
+              << " output-dim=" << output_dim_;
+  cfl->GetValue("scale", &scale_);
+  if (cfl->HasUnusedValues())
+    KALDI_ERR << "Could not process these elements in initializer: "
+              << cfl->UnusedValues();
+}
+
+void MaxPoolingOverBlock::Read(std::istream &is, bool binary) {
+  ExpectOneOrTwoTokens(is, binary, "<MaxPoolingOverBlock>", "<InputDim>");
+  ReadBasicType(is, binary, &input_dim_);
+  ExpectToken(is, binary, "<OutputDim>");
+  ReadBasicType(is, binary, &output_dim_);
+  ExpectToken(is, binary, "<Scale>");
+  ReadBasicType(is, binary, &scale_);
+  ExpectToken(is, binary, "</MaxPoolingOverBlock>");
+}
+
+void MaxPoolingOverBlock::Write(std::ostream &os, bool binary) const {
+  WriteToken(os, binary, "<MaxPoolingOverBlock>");
+  WriteToken(os, binary, "<InputDim>");
+  WriteBasicType(os, binary, input_dim_);
+  WriteToken(os, binary, "<OutputDim>");
+  WriteBasicType(os, binary, output_dim_);
+  WriteToken(os, binary, "<Scale>");
+  WriteBasicType(os, binary, scale_);
+  WriteToken(os, binary, "</MaxPoolingOverBlock>");
+}
+
+std::string MaxPoolingOverBlock::Info() const {
+  std::ostringstream stream;
+  stream << Type() << ", input-dim=" << input_dim_
+         << ", output-dim=" << output_dim_
+         << ", scale=" << scale_;
+  return stream.str();
+}
+
+void* MaxPoolingOverBlock::Propagate(const ComponentPrecomputedIndexes *indexes,
+                                   const CuMatrixBase<BaseFloat> &in,
+                                   CuMatrixBase<BaseFloat> *out) const {
+  KALDI_ASSERT(out->NumRows() == in.NumRows() &&
+               out->NumCols() == output_dim_ &&
+               in.NumCols() == input_dim_);
+  out->MaxMatBlocks(scale_, in, kNoTrans);
+  return NULL;
+}
+
+void MaxPoolingOverBlock::Backprop(
+    const std::string &debug_info,
+    const ComponentPrecomputedIndexes *indexes,
+    const CuMatrixBase<BaseFloat> &, //in_value
+    const CuMatrixBase<BaseFloat> &, // out_value,
+    const CuMatrixBase<BaseFloat> &out_deriv,
+    void *memo,
+    Component *to_update,
+    CuMatrixBase<BaseFloat> *in_deriv) const {
+  if (in_deriv) {
+    in_deriv->MaxMatBlocks(scale_, out_deriv, kNoTrans);
+  }
+}
 
 } // namespace nnet3
 } // namespace kaldi

src/nnet3/nnet-simple-component.h

@@ -1220,6 +1220,49 @@ class SumBlockComponent: public Component {
   SumBlockComponent &operator = (const SumBlockComponent &other); // Disallow.
 };
 
+/**  MaxPoolingOverBlock gets maximum value over blocks of its input: for instance, if
+     you create one with the config "input-dim=400 output-dim=100",
+     its output will be the maximum value over the 4 100-dimensional blocks of
+     the input.
+
+     Accepted values on its config-file line are:
+        input-dim  The input dimension.  Required.
+        output-dim  The block dimension.  Required.  Must divide input-dim.
+        scale      A scaling factor on the output.  Defaults to 1.0.
+ */
+class MaxPoolingOverBlock: public Component {
+ public:
+  explicit MaxPoolingOverBlock(const MaxPoolingOverBlock &other);
+  MaxPoolingOverBlock() { }
+  virtual std::string Type() const { return "MaxPoolingOverBlock"; }
+  virtual int32 Properties() const {
+    return kSimpleComponent|kPropagateAdds|kBackpropAdds;
+  }
+  virtual void (ConfigLine *cfl);
+  virtual int32 InputDim() const { return input_dim_; }
+  virtual int32 OutputDim() const { return output_dim_; }
+  virtual void Read(std::istream &is, bool binary);
+  virtual void Write(std::ostream &os, bool binary) const;
+  virtual std::string Info() const;
+  virtual Component* Copy() const { return new MaxPoolingOverBlock(*this); }
+  virtual void* Propagate(const ComponentPrecomputedIndexes *indexes,
+                          const CuMatrixBase<BaseFloat> &in,
+                          CuMatrixBase<BaseFloat> *out) const;
+  virtual void Backprop(const std::string &debug_info,
+                        const ComponentPrecomputedIndexes *indexes,
+                        const CuMatrixBase<BaseFloat> &, //in_value
+                        const CuMatrixBase<BaseFloat> &, // out_value,
+                        const CuMatrixBase<BaseFloat> &out_deriv,
+                        void *memo,
+                        Component *to_update,
+                        CuMatrixBase<BaseFloat> *in_deriv) const;
+ private:
+  int32 input_dim_;
+  int32 output_dim_;
+  BaseFloat scale_;
+  MaxPoolingOverBlock &operator = (const MaxPoolingOverBlock &other); // Disallow.
+};
+
 
 /*
  ClipGradientComponent just duplicates its input, but clips gradients