math_function.cu 6.7 KB
Newer Older
1
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserved.
Q
qijun 已提交
2 3 4 5 6 7 8 9 10 11 12 13

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. */
Y
Yu Yang 已提交
14
#include <vector>
Y
Yi Wang 已提交
15
#include "paddle/fluid/framework/data_type.h"
Y
Yu Yang 已提交
16
#include "paddle/fluid/operators/math/blas.h"
Y
Yi Wang 已提交
17 18
#include "paddle/fluid/operators/math/math_function.h"
#include "paddle/fluid/operators/math/math_function_impl.h"
19
#include "paddle/fluid/platform/float16.h"
Q
qijun 已提交
20

Q
qijun 已提交
21 22 23 24
namespace paddle {
namespace operators {
namespace math {

25 26
using float16 = paddle::platform::float16;

K
Kexin Zhao 已提交
27
template struct SetConstant<platform::CUDADeviceContext, platform::float16>;
Q
QI JUN 已提交
28 29 30 31 32
template struct SetConstant<platform::CUDADeviceContext, float>;
template struct SetConstant<platform::CUDADeviceContext, double>;
template struct SetConstant<platform::CUDADeviceContext, int>;
template struct SetConstant<platform::CUDADeviceContext, int64_t>;
template struct SetConstant<platform::CUDADeviceContext, bool>;
33

Q
qingqing01 已提交
34 35 36 37
#define DEFINE_GPU_TRANS(RANK)                                           \
  template struct Transpose<platform::CUDADeviceContext, float, RANK>;   \
  template struct Transpose<platform::CUDADeviceContext, double, RANK>;  \
  template struct Transpose<platform::CUDADeviceContext, float16, RANK>; \
38 39 40
  template struct Transpose<platform::CUDADeviceContext, int8_t, RANK>;  \
  template struct Transpose<platform::CUDADeviceContext, int32_t, RANK>; \
  template struct Transpose<platform::CUDADeviceContext, int64_t, RANK>;
41 42 43 44 45 46 47

DEFINE_GPU_TRANS(1);
DEFINE_GPU_TRANS(2);
DEFINE_GPU_TRANS(3);
DEFINE_GPU_TRANS(4);
DEFINE_GPU_TRANS(5);
DEFINE_GPU_TRANS(6);
Q
qijun 已提交
48

49 50
struct TensorSetConstantGPU {
  TensorSetConstantGPU(const platform::DeviceContext& context,
D
dangqingqing 已提交
51
                       framework::Tensor* tensor, float value)
52 53 54
      : context_(context), tensor_(tensor), value_(value) {}

  template <typename T>
D
dzhwinter 已提交
55
  void apply() const {
Q
QI JUN 已提交
56 57 58
    SetConstant<platform::CUDADeviceContext, T> functor;
    functor(reinterpret_cast<const platform::CUDADeviceContext&>(context_),
            tensor_, static_cast<T>(value_));
59 60 61 62 63 64 65 66
  }

  const platform::DeviceContext& context_;
  framework::Tensor* tensor_;
  float value_;
};

template <>
D
dzhwinter 已提交
67
void set_constant_with_place<platform::CUDAPlace>(
68 69
    const platform::DeviceContext& context, framework::Tensor* tensor,
    float value) {
Y
Yu Yang 已提交
70
  framework::VisitDataType(tensor->type(),
71
                           TensorSetConstantGPU(context, tensor, value));
72 73
}

Q
qingqing01 已提交
74
template <typename T>
Q
qingqing01 已提交
75 76 77
__global__ void RowwiseAddKernel(const T* a, const T* b, T* c, int width,
                                 int num) {
  T tmp = 1.0 / width;
78
  CUDA_KERNEL_LOOP(i, num) {
Q
qingqing01 已提交
79 80 81
    int h = i * tmp;
    int w = i - h * width;
    c[i] = a[i] + b[w];
Q
qingqing01 已提交
82 83 84 85 86 87 88 89 90
  }
}

template <typename T>
struct RowwiseAdd<platform::CUDADeviceContext, T> {
  void operator()(const platform::CUDADeviceContext& context,
                  const framework::Tensor& input,
                  const framework::Tensor& vector, framework::Tensor* output) {
    auto in_dims = input.dims();
Q
qingqing01 已提交
91 92 93
    auto size = input.numel() / in_dims[0];
    PADDLE_ENFORCE_EQ(vector.numel(), size);
    PADDLE_ENFORCE_EQ(output->dims(), in_dims);
Q
qingqing01 已提交
94 95 96
    int blocks = 512;
    int grids = (input.numel() + blocks - 1) / blocks;
    RowwiseAddKernel<T><<<grids, blocks, 0, context.stream()>>>(
Q
qingqing01 已提交
97 98
        input.data<T>(), vector.data<T>(), output->data<T>(),
        static_cast<int>(in_dims[1]), static_cast<int>(input.numel()));
Q
qingqing01 已提交
99 100 101
  }
};

Q
QI JUN 已提交
102 103 104
template struct RowwiseAdd<platform::CUDADeviceContext, float>;
template struct RowwiseAdd<platform::CUDADeviceContext, double>;
template struct ColwiseSum<platform::CUDADeviceContext, float>;
Y
yangyaming 已提交
105 106
template struct ColwiseSum<platform::CUDADeviceContext, int>;
template struct ColwiseSum<platform::CUDADeviceContext, int64_t>;
Q
QI JUN 已提交
107 108
// template struct ColwiseSum<platform::CUDADeviceContext, double>;
// The ColwiseSum<platform::CUDADeviceContext, double> failed in debug mode,
109 110
// and only failed for this case. So reimplemented it.
template <>
Q
QI JUN 已提交
111 112
void ColwiseSum<platform::CUDADeviceContext, double>::operator()(
    const platform::CUDADeviceContext& context, const framework::Tensor& input,
113 114 115 116 117 118
    framework::Tensor* vector) {
  auto in_dims = input.dims();
  auto size = input.numel() / in_dims[0];
  PADDLE_ENFORCE_EQ(vector->numel(), size);
  framework::Tensor one;
  one.mutable_data<double>({in_dims[0]}, context.GetPlace());
Q
QI JUN 已提交
119
  SetConstant<platform::CUDADeviceContext, double> set;
120
  set(context, &one, static_cast<double>(1.0));
Y
Yu Yang 已提交
121 122 123
  GetBlas<platform::CUDADeviceContext, double>(context).GEMV(
      true, static_cast<int>(in_dims[0]), static_cast<int>(in_dims[1]), 1.0,
      input.data<double>(), one.data<double>(), 0.0, vector->data<double>());
124
}
125

C
chengduoZH 已提交
126 127 128 129 130 131 132 133 134 135 136 137 138 139 140 141
template struct RowwiseSum<platform::CUDADeviceContext, float>;
// template struct RowwiseSum<platform::CUDADeviceContext, double>;
// TODO(zcd): Following ColwiseSum format, need to confirm.
// The RowwiseSum<platform::CUDADeviceContext, double> failed in debug mode,
// and only failed for this case. So reimplemented it.
template <>
void RowwiseSum<platform::CUDADeviceContext, double>::operator()(
    const platform::CUDADeviceContext& context, const framework::Tensor& input,
    framework::Tensor* vector) {
  auto in_dims = input.dims();
  auto size = input.numel() / in_dims[0];
  PADDLE_ENFORCE_EQ(vector->numel(), in_dims[0]);
  framework::Tensor one;
  one.mutable_data<double>({size}, context.GetPlace());
  SetConstant<platform::CUDADeviceContext, double> set;
  set(context, &one, static_cast<double>(1.0));
Y
Yu Yang 已提交
142 143 144
  GetBlas<platform::CUDADeviceContext, double>(context).GEMV(
      true, static_cast<int>(in_dims[1]), static_cast<int>(in_dims[0]), 1.0,
      one.data<double>(), input.data<double>(), 0.0, vector->data<double>());
C
chengduoZH 已提交
145 146 147 148 149
}

template struct RowwiseMean<platform::CUDADeviceContext, float>;
template struct RowwiseMean<platform::CUDADeviceContext, double>;

150 151 152 153 154 155 156 157 158 159 160 161 162
template <typename T>
struct ElementwiseAddTo<platform::CUDADeviceContext, T> {
  void operator()(platform::CUDADeviceContext* ctx,
                  const framework::Tensor& src, framework::Tensor* dst) {
    auto in = framework::EigenVector<T>::Flatten(src);
    auto out = framework::EigenVector<T>::Flatten(*dst);
    auto& place = *(ctx->eigen_device());
    out.device(place) = out + in;
  }
};

template struct ElementwiseAddTo<platform::CUDADeviceContext,
                                 platform::float16>;
Q
qijun 已提交
163 164 165
}  // namespace math
}  // namespace operators
}  // namespace paddle