API 介绍 - OpFunc

介绍 Kernel Primitive API 定义的 Functor,当前一共有 13 个 Functor 可以直接使用。

Unary Functor

ExpFunctor

定义

namespace kps = paddle::operators::kernel_primitives;
template <typename Tx, typename Ty = Tx>
struct kps::ExpFunctor<Tx, Ty>();

功能介绍

对 Tx 类型的输入数据做 Exp 操作,并将结果转成 Ty 类型返回。

模板参数

Tx : 输入数据的类型。
Ty : 返回类型。

使用示例

const int VecSize = 1;
float data[VecSize];
float out[VecSize];

kps::ElementwiseUnary<float, float, VecSize, 1, 1, kps::ExpFunctor<float>>(out, data, kps::ExpFunctor<float>());

// data[0] = 0;
// out[0] = exp(data);
// out[0] = 1

IdentityFunctor

定义

namespace kps = paddle::operators::kernel_primitives;
template <typename Tx, typename Ty = Tx>
struct kps::IdentityFunctor<Tx, Ty>();

功能介绍

将 Tx 类型的输入数据转成 Ty 类型返回。

模板参数

Tx : 输入数据的类型。
Ty : 返回类型。

使用示例

const int VecSize = 1;
float data[VecSize];
int out[VecSize];

kps::ElementwiseUnary<float, int, VecSize, 1, 1, kps::IdentityFunctor<float, int>>(out, data, kps::IdentityFunctor<float, int>());

// data[0] = 1.3;
// out[0] = 1

DivideFunctor

定义

namespace kps = paddle::operators::kernel_primitives;
template <typename Tx, typename Ty = Tx>
struct kps::DivideFunctor<Tx, Ty>(num);

功能介绍

将 Tx 类型的输入数据除以 num,并将结果转成 Ty 类型返回。

模板参数

Tx : 输入数据的类型。
Ty : 返回类型。

使用示例

const int VecSize = 1;
float data[VecSize];
float out[VecSize];
float num = 10.0;

kps::ElementwiseUnary<float, float, VecSize, 1, 1, kps::DivideFunctor<float>>(out, data, kps::DivideFunctor<float>(num));

// data[0] = 3.0
// out[0] = (3.0 / 10.0)
// out[0] = 0.3

SquareFunctor

定义

namespace kps = paddle::operators::kernel_primitives;
template <typename Tx, typename Ty = Tx>
struct kps::SquareFunctor<Tx, Ty>();

功能介绍

对 Tx 类型的输入数做 Square 操作,并将结果转成 Ty 类型返回。

模板参数

Tx : 输入数据的类型。
Ty : 返回类型。

使用示例

const int VecSize = 1;
float data[VecSize];
float out[VecSize];

kps::ElementwiseUnary<float, float, VecSize, 1, 1, kps::SquareFunctor<float>>(out, data, kps::SquareFunctor<float>());

// data[0] = 3.0
// out[0] = (3.0 * 3.0)
// out[0] = 9.0

Binary Functor

MinFunctor

定义

namespace kps = paddle::operators::kernel_primitives;
template <typename T>
struct kps::MinFunctor<T>();

功能介绍

返回两个输入中的最小值。MinFunctor 提供了用于数据初始化的 initial() 函数,返回 T 类型表示的最大值。

模板参数

T : 数据类型。

使用示例

const int VecSize = 1;
float input1[VecSize];
float input2[VecSize];
float out[VecSize];

kps::ElementwiseBinary<float, float, VecSize, 1, 1, kps::MinFunctor<float>>(out, input1, input2, kps::MinFunctor<float>());

// input1[0] = 3.0
// input2[0] = 1.0
// out = input1[0] < input2[0] ? input1[0] : input2[0]
// out[0] = 1.0

MaxFunctor

定义

namespace kps = paddle::operators::kernel_primitives;
template <typename T>
struct kps::MaxFunctor<T>();

功能介绍

返回两个输入中的最大值。MaxFunctor 提供了用于数据初始化的 initial() 函数,返回 T 类型表示的最小值。

模板参数

T : 数据类型。

使用示例

const int VecSize = 1;
float input1[VecSize];
float input2[VecSize];
float out[VecSize];

kps::ElementwiseBinary<float, float, VecSize, 1, 1, kps::MaxFunctor<float>>(out, input1, input2, kps::MaxFunctor<float>());

// input1[0] = 3.0
// input2[0] = 1.0
// out = input1[0] > input2[0] ? input1[0] : input2[0]
// out[0] = 3.0

AddFunctor

定义

namespace kps = paddle::operators::kernel_primitives;
template <typename T>
struct kps::AddFunctor<T>();

功能介绍

返回两个输入之和。AddFunctor 提供了用于数据初始化的 initial() 函数,返回 T 类型表示的数据 0。

模板参数

T : 数据类型。

使用示例

const int VecSize = 1;
float input1[VecSize];
float input2[VecSize];
float out[VecSize];

kps::ElementwiseBinary<float, float, VecSize, 1, 1, kps::AddFunctor<float>>(out, input1, input2, kps::AddFunctor<float>());

// input1[0] = 3.0
// input2[0] = 1.0
// out = input1[0] + input2[0]
// out[0] = 4.0

MulFunctor

定义

namespace kps = paddle::operators::kernel_primitives;
template <typename T>
struct kps::MulFunctor<T>();

功能介绍

返回两个输入的乘积。MulFunctor 提供了用于数据初始化的 initial() 函数,返回 T 类型表示的数据 1。

模板参数

T : 数据类型。

使用示例

const int VecSize = 1;
float input1[VecSize];
float input2[VecSize];
float out[VecSize];

kps::ElementwiseBinary<float, float, VecSize, 1, 1, kps::MulFunctor<float>>(out, input1, input2, kps::MulFunctor<float>());

// input1[0] = 3.0
// input2[0] = 1.0
// out = input1[0] * input2[0]
// out[0] = 3.0

LogicalOrFunctor

定义

namespace kps = paddle::operators::kernel_primitives;
template <typename T>
struct kps::LogicalOrFunctor<T>();

功能介绍

返回两个输入元素逻辑或操作后的结果。LogicalOrFunctor 提供了用于数据初始化的 initial() 函数,返回 T 类型表示的数据 false。

模板参数

T : 数据类型。

使用示例

const int VecSize = 1;
bool input1[VecSize];
bool input2[VecSize];
bool out[VecSize];

kps::ElementwiseBinary<bool, bool, VecSize, 1, 1, kps::LogicalOrFunctor<bool>>(out, input1, input2, kps::LogicalOrFunctor<bool>());

// input1[0] = false
// input2[0] = true
// out = input1[0] || input2[0]
// out[0] = true

LogicalAndFunctor

定义

namespace kps = paddle::operators::kernel_primitives;
template <typename T>
struct kps::LogicalAndFunctor<T>();

功能介绍

返回两个输入元素逻辑与操作后的结果。LogicalAndFunctor 提供了用于数据初始化的 initial() 函数,返回 T 类型表示的数据 true。

模板参数

T : 数据类型。

使用示例

const int VecSize = 1;
bool input1[VecSize];
bool input2[VecSize];
bool out[VecSize];

kps::ElementwiseBinary<bool, bool, VecSize, 1, 1, kps::LogicalAndFunctor<bool>>(out, input1, input2, kps::LogicalAndFunctor<bool>());

// input1[0] = false
// input2[0] = true
// out = input1[0] && input2[0]
// out[0] = false

SubFunctor

定义

namespace kps = paddle::operators::kernel_primitives;
template <typename T>
struct kps::SubFunctor<T>();

功能介绍

两个输入进行减法操作。SubFunctor 提供了用于数据初始化的 initial() 函数,返回 T 类型表示的数据 0。

模板参数

T : 数据类型。

使用示例

const int VecSize = 1;
float input1[VecSize];
float input2[VecSize];
float out[VecSize];

kps::ElementwiseBinary<float, float, VecSize, 1, 1, kps::SubFunctor<float>>(out, input1, input2, kps::SubFunctor<float>());

// input1[0] = 3.0
// input2[0] = 1.0
// out = input1[0] - input2[0]
// out[0] = 2.0

DivFunctor

定义

namespace kps = paddle::operators::kernel_primitives;
template <typename T>
struct kps::DivFunctor<T>();

功能介绍

两个输入进行除法操作。DivFunctor 提供了用于数据初始化的 initial() 函数,返回 T 类型表示的数据 1。

模板参数

T : 数据类型。

使用示例

const int VecSize = 1;
float input1[VecSize];
float input2[VecSize];
float out[VecSize];

kps::ElementwiseBinary<float, float, VecSize, 1, 1, kps::DivideFunctor<float>>(out, input1, input2, kps::DivideFunctor<float>());

// input1[0] = 3.0
// input2[0] = 1.0
// out = input1[0] / input2[0]
// out[0] = 3.0

FloorDivFunctor

定义

namespace kps = paddle::operators::kernel_primitives;
template <typename T>
struct kps::FloorDivFunctor<T>();

功能介绍

两个输入进行除法操作,返回整数部分。FloorDivFunctor 提供了用于数据初始化的 initial() 函数,返回 T 类型表示的数据 1。

模板参数

T : 数据类型。

使用示例

const int VecSize = 1;
float input1[VecSize];
float input2[VecSize];
float out[VecSize];

kps::ElementwiseBinary<float, float, VecSize, 1, 1, kps::FloorDivFunctor<float>>(out, input1, input2, kps::FloorDivFunctor<float>());

// input1[0] = 3.0
// input2[0] = 2.0
// out = input1[0] / input2[0]
// out[0] = 1.0

Functor 定义规则

当前计算函数中仅 ElementwiseAny 支持 Functor 参数设置为指针,其他计算函数的 Functor 仅能设置为普通参数。

普通参数传递

除 ElementwiseAny API 外其他计算函数仅支持普通参数传递。例如需要实现 (a + b) * c 可将 Functor 定义如下:

ExampleTernaryFunctor:

namespace kps = paddle::operators::kernel_primitives;
template <typename T>
struct ExampleTernaryFunctor {
   inline HOSTDEVICE T operator()(const T &input1, const T &input2, const T &input3) const {
       return ((input1 + input2) * input3);
   }
};

示例

template<int VecSize, typename InT, typename OutT>
__global__ void ElementwiseTernaryKernel(InT *input0, InT *input1, InT * input2, OutT *output, int size) {
// global memory input pointer input0, input1, input2
auto functor = ExampleTernaryFunctor<InT>();

const int NX = 4;
const int NY = 1;
const int BlockSize = 1;
const bool IsBoundary = false;
int data_offset = NX * blockIdx.x * blockDim.x;
int num = size - data_offset;
// if  num < NX * blockDim.x set IsBoundary = true

InT in0[NX * NY];
InT in1[NX * NY];
InT in2[NX * NY];
OutT out[NX * NY];

// each thread reads NX data continuously, and each block reads num data continuously
kps::ReadData<InT, NX, NY, BlockSize, IsBoundary>(in0, input0 + data_offset, num);
kps::ReadData<InT, NX, NY, BlockSize, IsBoundary>(in1, input1 + data_offset, num);
kps::ReadData<InT, NX, NY, BlockSize, IsBoundary>(in2, input2 + data_offset, num);
kps::ElementwiseTernary<InT, OutT, NX, NY, BlockSize, ExampleTernaryFunctor<InT>>(out, in0, in1, in2, functor);

...

}

指针传递

在进行 ElementwiseAny 的 Functor 定义时,需要保证 operate() 函数的参数是数组指针。例如要实现功能: (a + b) * c + d, 则可以结合 ElementwiseAny 与 Functor 完成对应计算。

ExampleAnyFunctor 定义:

namespace kps = paddle::operators::kernel_primitives;
template <typename T>
struct ExampleAnyFunctor {
   inline HOSTDEVICE T operator()(const T * args) const { return ((arg[0] + arg[1]) * arg[2] + arg[3]); }
};

示例

template<int VecSize, typename InT, typename OutT>
__global__ void ElementwiseAnyKernel(InT *input0, InT *input1, InT * input2, InT * input3, OutT *output, int size) {
// global memory input pointer input0, input1, input2, input3
auto functor = ExampleAnyFunctor<InT>();

const int NX = 4;
const int NY = 1;
const int BlockSize = 1;
const bool IsBoundary = false;
const int Arity = 4; // the pointers of inputs
int data_offset = NX * blockIdx.x * blockDim.x;
int num = size - data_offset;
// if  num < NX * blockDim.x set IsBoundary = true

InT inputs[Arity][NX * NY];
OutT out[NX * NY];

// each thread reads NX data continuously, and each block reads num data continuously
kps::ReadData<InT, NX, NY, BlockSize, IsBoundary>(inputs[0], input0 + data_offset, num);
kps::ReadData<InT, NX, NY, BlockSize, IsBoundary>(inputs[1], input1 + data_offset, num);
kps::ReadData<InT, NX, NY, BlockSize, IsBoundary>(inputs[2], input2 + data_offset, num);
kps::ReadData<InT, NX, NY, BlockSize, IsBoundary>(inputs[3], input3 + data_offset, num);
kps::ElementwiseAny<InT, OutT, NX, NY, BlockSize, Arity, ExampleAnyFunctor<InT>>(out, inputs, functor);

...

}