API Description - OpFunc¶

Introduce the calculation functors defined by the Kernel Primitive API. There are currently 13 functors that can be used directly.

Unary Functor¶

ExpFunctor¶

Definition¶

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

Description¶

Exp operation is performed on the input of Tx type, and the result is converted into Ty type and returned.

Template Parameters¶

Tx : The type of input.
Ty : The type of return.

Example¶

```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¶

Definition¶

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

Description¶

Convert the input data of Tx type to Ty type and return.

Template Parameters¶

Tx : The type of input.
Ty : The type of return.

Example¶

```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¶

Definition¶

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

Description¶

Divide the input data of Tx type by num, and convert the result into Ty type to return.

Template Parameters¶

Tx : The type of input.
Ty : The type of return.

Example¶

```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¶

Definition¶

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

Description¶

Perform Square operation on the input number of Tx type, and convert the result into Ty type to return.

Template Parameters¶

Tx : The type of input.
Ty : The type of return.

Example¶

```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¶

Definition¶

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

Description¶

Returns the minimum of the two inputs. MinFunctor provides the initial() function for data initialization and returns the maximum value represented by the T type.

Template Parameters¶

T : The type of data.

Example¶

```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¶

Definition¶

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

Description¶

Returns the maximum value of the two inputs. MaxFunctor provides the initial() function for data initialization, and returns the minimum value represented by the T type.

Template Parameters¶

T : The type of data.

Example¶

```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
```

Definition¶

```namespace kps = paddle::operators::kernel_primitives;
template <typename T>
```

Description¶

Returns the sum of the two inputs. AddFunctor provides the initial() function for data initialization, and returns the data 0 represented by the T type.

Template Parameters¶

T : The type of data.

Example¶

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

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

MulFunctor¶

Definition¶

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

Description¶

Returns the product of two inputs. MulFunctor provides the initial() function for data initialization, and returns data 1 represented by the T type.

Template Parameters¶

T : The type of data.

Example¶

```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¶

Definition¶

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

Description¶

Returns the result of the logical or operation of two inputs. LogicalOrFunctor provides the initial() function for data initialization, and returns false for the data represented by the T type.

Template Parameters¶

T : The type of data.

Example¶

```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¶

Definition¶

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

Description¶

Returns the result of the logical and operation of the two inputs. LogicalAndFunctor provides the initial() function for data initialization, and returns true for the data represented by the T type.

Template Parameters¶

T : The type of data.

Example¶

```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¶

Definition¶

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

Description¶

Two inputs are subtracted. SubFunctor provides the initial() function for data initialization, and returns the data 0 represented by the T type.

Template Parameters¶

T : The type of data.

Example¶

```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¶

Definition¶

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

Description¶

Two inputs are divided. DivFunctor provides the initial() function for data initialization, and returns the data 1 represented by the T type.

Template Parameters¶

T : The type of data.

Example¶

```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¶

Definition¶

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

Description¶

Divide the two inputs and return the integer part. FloorDivFunctor provides the initial() function for data initialization, and returns data 1 represented by the T type.

Template Parameters¶

T : The type of data.

Example¶

```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 Definition Rules¶

In the current calculation function, only ElementwiseAny supports setting the functor parameter as a pointer, and the functor of other calculation functions can only be set as a normal parameter.

Common parameters¶

Except ElementwiseAny API, other calculation functions only support ordinary parameter transfer. For example, to realize (a + b) * c, Functor can be defined as follows:

ExampleFunctor2:

```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);
}
};
```

Example¶

```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];

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);

...

}
```

Pointer¶

When defining the Functor of ElementwiseAny, you need to ensure that the parameter of the operate() function is an array pointer. For example, to realize the function: (a + b) * c + d, you can combine ElementwiseAny and Functor to complete the calculation.

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]); }
};
```

example¶

```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];