sequence2batch.cu 2.9 KB
Newer Older
D
dangqingqing 已提交
1 2 3 4 5 6 7 8 9 10 11 12 13 14
/* 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. */

15
#define EIGEN_USE_GPU
D
dangqingqing 已提交
16 17 18 19 20 21
#include "paddle/operators/math/sequence2batch.h"

namespace paddle {
namespace operators {
namespace math {

D
dangqingqing 已提交
22
template <typename T, int BlockDimX, int BlockDimY, int GridDimX>
23 24
__global__ void CopyMatrixRowsKernel(const T* src, T* dst, const size_t* index,
                                     int64_t height, int64_t width,
Y
Yu Yang 已提交
25
                                     bool is_src_index) {
D
dangqingqing 已提交
26 27 28 29 30 31
  int idx = threadIdx.x;
  int idy = threadIdx.y;
  int id = blockIdx.x + idy * GridDimX;
  while (id < height) {
    int src_idx = is_src_index ? index[id] : id;
    int dst_idx = is_src_index ? id : index[id];
32
    const T* src_data = src + src_idx * width;
D
dangqingqing 已提交
33 34 35 36 37 38 39 40 41 42 43 44
    T* dst_data = dst + dst_idx * width;
    for (int i = idx; i < width; i += BlockDimX) {
      dst_data[i] = src_data[i];
    }
    id += BlockDimY * GridDimX;
  }
}

template <typename T>
class CopyMatrixRowsFunctor<platform::GPUPlace, T> {
 public:
  void operator()(const platform::DeviceContext& context,
45 46
                  const framework::Tensor& src, const size_t* index,
                  framework::Tensor& dst, bool is_src_index) {
D
dangqingqing 已提交
47 48
    auto src_dims = src.dims();
    auto dst_dims = dst.dims();
49 50 51 52
    PADDLE_ENFORCE_EQ(src_dims.size(), 2,
                      "The src must be matrix with rank 2.");
    PADDLE_ENFORCE_EQ(dst_dims.size(), 2,
                      "The dst must be matrix with rank 2.");
D
dangqingqing 已提交
53 54 55 56 57 58 59 60 61
    PADDLE_ENFORCE_EQ(src_dims[1], dst_dims[1],
                      "The width of src and dst must be same.");
    auto height = dst_dims[0];
    auto width = dst_dims[1];
    auto* src_data = src.data<T>();
    auto* dst_data = dst.data<T>();

    dim3 threads(128, 8);
    dim3 grid(8, 1);
62 63
    auto stream =
        reinterpret_cast<const platform::CUDADeviceContext&>(context).stream();
D
dangqingqing 已提交
64
    CopyMatrixRowsKernel<T, 128, 8, 8><<<grid, threads, 0, stream>>>(
65
        src_data, dst_data, index, height, width, is_src_index);
D
dangqingqing 已提交
66 67 68 69 70 71
  }
};

template class CopyMatrixRowsFunctor<platform::GPUPlace, float>;
template class CopyMatrixRowsFunctor<platform::GPUPlace, double>;

D
dangqingqing 已提交
72
template class LoDTensor2BatchFunctor<platform::GPUPlace, float>;
73 74 75
template class LoDTensor2BatchFunctor<platform::GPUPlace, double>;
template class Batch2LoDTensorFunctor<platform::GPUPlace, float>;
template class Batch2LoDTensorFunctor<platform::GPUPlace, double>;
D
dangqingqing 已提交
76 77 78 79

}  // namespace math
}  // namespace operators
}  // namespace paddle
新手
引导
客服 返回
顶部