Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

add kernel primitive api #3890

Merged
merged 75 commits into from
Oct 13, 2021
Merged
Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
75 commits
Select commit Hold shift + click to select a range
fafbb84
add kernel primitive api
AnnaTrainingG Sep 17, 2021
997a971
Merge branch 'develop' of https://github.com/PaddlePaddle/docs into P…
AnnaTrainingG Sep 17, 2021
7f0812f
modfied index_cn in guides
AnnaTrainingG Sep 18, 2021
c83cb7b
update
AnnaTrainingG Sep 18, 2021
210ab74
update
AnnaTrainingG Sep 18, 2021
319ba2f
update
AnnaTrainingG Sep 18, 2021
c95af9d
update
AnnaTrainingG Sep 18, 2021
3a6b52a
update
AnnaTrainingG Sep 18, 2021
6012225
update
AnnaTrainingG Sep 18, 2021
21bc6f3
update
AnnaTrainingG Sep 18, 2021
98a2eb7
update
AnnaTrainingG Sep 18, 2021
12cf54f
yes
AnnaTrainingG Sep 22, 2021
66621b9
update
AnnaTrainingG Sep 22, 2021
cbcf3b8
update
AnnaTrainingG Sep 22, 2021
b54eb97
update
AnnaTrainingG Sep 22, 2021
84b6984
add case
AnnaTrainingG Sep 22, 2021
cce555f
update
AnnaTrainingG Sep 22, 2021
8e8ab02
all in io
AnnaTrainingG Sep 22, 2021
da7593c
temp
AnnaTrainingG Sep 22, 2021
e18ec8f
temp update
AnnaTrainingG Sep 22, 2021
d674dd7
temp update
AnnaTrainingG Sep 22, 2021
c21109b
update temp
AnnaTrainingG Sep 22, 2021
cdf48c1
update
AnnaTrainingG Sep 22, 2021
f6d915c
add index_en
AnnaTrainingG Sep 23, 2021
4b4d40d
add en
AnnaTrainingG Sep 23, 2021
01943ef
update en
AnnaTrainingG Sep 23, 2021
4b37f31
update
AnnaTrainingG Sep 23, 2021
57d114d
update
AnnaTrainingG Sep 24, 2021
f7f36e9
update
AnnaTrainingG Sep 24, 2021
b48082c
update
AnnaTrainingG Sep 24, 2021
e2d1e5f
update
AnnaTrainingG Sep 24, 2021
cc30464
upate
AnnaTrainingG Sep 24, 2021
fa50b68
Update index_en.rst
AnnaTrainingG Sep 24, 2021
4d62970
update
AnnaTrainingG Sep 24, 2021
3924dd2
Merge branch 'Primitive_API_31094' of https://github.com/niuliling123…
AnnaTrainingG Sep 24, 2021
84986c6
update
AnnaTrainingG Sep 24, 2021
19d0f5b
update
AnnaTrainingG Sep 24, 2021
cbccb13
update
AnnaTrainingG Sep 24, 2021
caab6d2
update
AnnaTrainingG Sep 24, 2021
fec8e48
add url
AnnaTrainingG Sep 26, 2021
fd8552b
update from ELe to Ele
AnnaTrainingG Sep 27, 2021
2f82c40
update
AnnaTrainingG Sep 27, 2021
4e23256
update Block
AnnaTrainingG Sep 27, 2021
8ca099b
update
AnnaTrainingG Sep 28, 2021
94123a6
add functor
AnnaTrainingG Sep 29, 2021
d6c2ec2
update
AnnaTrainingG Sep 29, 2021
8fbf2cb
update
AnnaTrainingG Sep 29, 2021
50f114b
add images
AnnaTrainingG Sep 29, 2021
6a0433c
update
AnnaTrainingG Sep 29, 2021
b3611e2
temp
AnnaTrainingG Sep 29, 2021
315ab1f
add en
AnnaTrainingG Sep 29, 2021
22ca7e0
update
AnnaTrainingG Sep 29, 2021
898b611
add static_cast
AnnaTrainingG Sep 29, 2021
f782de7
add static_cast
AnnaTrainingG Sep 29, 2021
a185ff0
update
AnnaTrainingG Sep 30, 2021
6f806d9
add functor_en
AnnaTrainingG Sep 30, 2021
32b04d6
update image
AnnaTrainingG Sep 30, 2021
2c788b5
update
AnnaTrainingG Sep 30, 2021
ffbc1a5
update
AnnaTrainingG Sep 30, 2021
16f2cff
add example_reduce.png
AnnaTrainingG Sep 30, 2021
757594e
add example_reduce.png
AnnaTrainingG Sep 30, 2021
0c12dcb
update and add example_add png
AnnaTrainingG Oct 8, 2021
a531027
update
AnnaTrainingG Oct 8, 2021
f65400c
update en
AnnaTrainingG Oct 9, 2021
e6a3d6d
update index and functor
AnnaTrainingG Oct 9, 2021
253bf26
update
AnnaTrainingG Oct 11, 2021
bc68805
update all
AnnaTrainingG Oct 11, 2021
b438f08
update paddlepaddle
AnnaTrainingG Oct 11, 2021
3b9322c
add url for functor
AnnaTrainingG Oct 11, 2021
df442c7
update for merge
AnnaTrainingG Oct 12, 2021
187147f
update functor_en url
AnnaTrainingG Oct 12, 2021
4dfed57
add init in io
AnnaTrainingG Oct 12, 2021
a23984a
update stride_nx in image
AnnaTrainingG Oct 12, 2021
978688d
update ;
AnnaTrainingG Oct 12, 2021
14856dd
update image and notes
AnnaTrainingG Oct 12, 2021
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
Next Next commit
update
  • Loading branch information
AnnaTrainingG committed Sep 30, 2021
commit 2c788b52473a549b155754c83a8d934bb3ef99d3
28 changes: 15 additions & 13 deletions docs/guides/11_kernel_primitives_api/functor_api_cn.md
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
# API 介绍 - Functor
介绍 Kernel Primitive API 定义的计算 Functor,当前一共有 13 个 Functor 可以直接使用。
介绍 Kernel Primitive API 定义的 Functor,当前一共有 13 个 Functor 可以直接使用。

## Unary Functor

Expand All @@ -14,12 +14,12 @@ struct kps::ExpFunctor<Tx, Ty>();
对 Tx 类型的输入数做 Exp 操作,并将结果转成 Ty 类型返回。

#### 模板参数
> Tx : 输入数据的类型。
> Ty : 返回类型。
> Tx : 输入数据的类型。</br>
> Ty : 返回类型。</br>

#### 使用示例
```
auto functor = kps::ExpFunctor<float>(1);
auto functor = kps::ExpFunctor<float>();
float input = 0;
float out = functor(input);

Expand All @@ -39,8 +39,8 @@ struct kps::IdentityFunctor<Tx, Ty>();
将 Tx 类型的输入数据转成 Ty 类型返回。

#### 模板参数
> Tx : 输入数据的类型。
> Ty : 返回类型。
> Tx : 输入数据的类型。</br>
> Ty : 返回类型。</br>

#### 使用示例
```
Expand All @@ -62,8 +62,8 @@ struct kps::DivideFunctor<Tx, Ty>(num);
将 Tx 类型的输入数据除以 num,并将结果转成 Ty 类型返回。

#### 模板参数
> Tx : 输入数据的类型。
> Ty : 返回类型。
> Tx : 输入数据的类型。</br>
> Ty : 返回类型。</br>

#### 使用示例
```
Expand All @@ -86,8 +86,8 @@ struct kps::SquareFunctor<Tx, Ty>();
对 Tx 类型的输入数做 Square 操作,并将结果转成 Ty 类型返回。

#### 模板参数
> Tx : 输入数据的类型。
> Ty : 返回类型。
> Tx : 输入数据的类型。</br>
> Ty : 返回类型。</br>
#### 使用示例
```
auto functor = kps::SquareFunctor<float>();
Expand Down Expand Up @@ -254,6 +254,7 @@ struct kps::SubFunctor<T>();

#### 模板参数
> T : 数据类型。

#### 使用示例
```
auto functor = kps::SubFunctor<float>();
Expand All @@ -277,6 +278,7 @@ struct kps::DivFunctor<T>();

#### 模板参数
> T : 数据类型。

#### 使用示例
```
auto functor = kps::DivFunctor<float>();
Expand Down Expand Up @@ -313,10 +315,10 @@ float out = functor(input1, input2);
// out = 0
```
## Functor 定义规则
ElementwiseAny 支持指针数组的形式传递,对于其他计算函数仅仅支持普通参数传递
当前计算函数中仅 ElementwiseAny 支持 Functor 参数设置为指针,其他计算函数的 Functor 仅能设置为普通参数

### 指针数组传递
目前指针数组传递仅仅在 ElementwiseAny 中支持,因此在进行 ElementwiseAny 的 Functor 定义时候需要保证重载的 operate() 函数的参数是指针数组。例如要实现功能: (a + b) * c + d, 则可以结合 ElementwiseAny 与 Functor 完成对应计算。
### 指针传递
在进行 ElementwiseAny 的 Functor 定义时候需要保证重载的 operate() 函数的参数是数组指针。例如要实现功能: (a + b) * c + d, 则可以结合 ElementwiseAny 与 Functor 完成对应计算。

ExampleFunctor1 定义:
```
Expand Down
19 changes: 10 additions & 9 deletions docs/guides/11_kernel_primitives_api/functor_api_en.md
Original file line number Diff line number Diff line change
Expand Up @@ -14,12 +14,12 @@ struct kps::ExpFunctor<Tx, Ty>();
Exp operation is performed on the input number of Tx type, and the result is converted into Ty type and returned.

#### Template Parameters
> Tx : Data type of input.
> Ty : The type of return.
> Tx : Data type of input.</br>
> Ty : The type of return.</br>

#### Example
```
auto functor = kps::ExpFunctor<float>(1);
auto functor = kps::ExpFunctor<float>();
float input = 0;
float out = functor(input);

Expand Down Expand Up @@ -62,8 +62,8 @@ struct kps::DivideFunctor<Tx, Ty>(num);
Divide the input data of Tx type by num, and convert the result into Ty type to return.

#### Template Parameters
> Tx : Data type of input.
> Ty : The type of return.
> Tx : Data type of input.</br>
> Ty : The type of return.</br>

#### Example
```
Expand All @@ -86,8 +86,8 @@ struct kps::SquareFunctor<Tx, Ty>();
Perform Square operation on the input number of Tx type, and convert the result into Ty type to return.

#### Template Parameters
> Tx : Data type of input.
> Ty : The type of return.
> Tx : Data type of input.</br>
> Ty : The type of return.</br>

#### Example
```
Expand Down Expand Up @@ -138,6 +138,7 @@ Returns the maximum value of the two inputs. MaxFunctor provides the initial() f

#### Template Parameters
> T : The type of data.

#### Example
```
auto functor = kps::MaxFunctor<float>();
Expand Down Expand Up @@ -318,10 +319,10 @@ float out = functor(input1, input2);
// out = 0
```
## Functor Definition Rules
ElementwiseAny supports the transfer of pointer arrays, and only supports ordinary parameter transfer for other calculation functions.
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.

### Pointer Array
Pointer array transfer is only supported in ElementwiseAny, so when defining ElementwiseAny's Functor, you need to ensure that the parameter of the overloaded operate() function is an array of pointers. For example, to realize the function: (a + b) * c + d, you can combine ElementwiseAny and Functor to complete the corresponding calculation.
When defining ElementwiseAny's Functor, you need to ensure that the parameter of the overloaded operate() function is a pointer. For example, to realize the function: (a + b) * c + d, you can combine ElementwiseAny and Functor to complete the calculation.

ExampleFunctor1:
```
Expand Down
7 changes: 3 additions & 4 deletions docs/guides/11_kernel_primitives_api/index_cn.rst
Original file line number Diff line number Diff line change
Expand Up @@ -55,9 +55,8 @@ API 列表
| | 根据当前 Block 的数据偏移和原始输入指针计算输入坐标, |
| | 并将输入数据读取到寄存器中。 |
+--------------------------------------+-------------------------------------------------------+
| ReadDataReduce | IO, |
| | 根据当前 Block 的数据偏移和原始输入指针计算输入坐标, |
| | 并将输入数据读取到寄存器中。 |
| ReadDataReduce | IO,Reduce 形式的读取数据, |
| | 将需要进行规约的数据从全局内存读取到寄存器中。 |
+--------------------------------------+-------------------------------------------------------+
| WriteData | IO,将数据从寄存器写入全局内存中。 |
+--------------------------------------+-------------------------------------------------------+
Expand All @@ -82,7 +81,7 @@ API 列表
API 介绍
##############

- `Functor <./functor_api_cn.html>`_ : 介绍Kernel Primitive API 提供的 Functor。
- `Functor <./functor_api_cn.html>`_ : 介绍 Kernel Primitive API 提供的 Functor。
- `IO API <./io_api_cn.html>`_ : 介绍 IO 类 API 的定义和功能。
- `Compute API <./compute_api_cn.html>`_ : 介绍 Compute 类 API 的定义和功能。

Expand Down
13 changes: 6 additions & 7 deletions docs/guides/11_kernel_primitives_api/index_en.rst
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@ Functor List
+--------------------------------------+-------------------------------------------------------+
| Functor name | Descriptions |
+======================================+=======================================================+
| ExpFunctor | Unary Functor, perform Exp operation. |
| ExpFunctor | Unary Functor, performs Exp operation. |
+--------------------------------------+-------------------------------------------------------+
| IdentityFunctor | Unary Functor, which performs type conversion on input|
| | data. |
Expand All @@ -41,13 +41,13 @@ Functor List
| LogicalOrFunctor | Binary Functor, which returns the logical or result of|
| | the inputs. |
+--------------------------------------+-------------------------------------------------------+
| LogicalAndFunctor | Binary Functor, which returns the logical AND result |
| LogicalAndFunctor | Binary Functor, which returns the logical and result |
| | of the inputs. |
+--------------------------------------+-------------------------------------------------------+
| DivFunctor | Binary Functor, returns the result of division of two |
| DivFunctor | Binary Functor, returns the result of division of the |
| | inputs. |
+--------------------------------------+-------------------------------------------------------+
| FloorDivFunctor | Binary Functor, returns the result of division of two |
| FloorDivFunctor | Binary Functor, returns the result of division of the |
| | inputs. |
+--------------------------------------+-------------------------------------------------------+

Expand All @@ -65,9 +65,8 @@ API List
| | block and the original input pointer, and read the input |
| | data into the register. |
+--------------------------------------+-------------------------------------------------------------+
| ReadDataReduce | IO, calculate the input coordinates according to the data |
| | offset of the current block and the original input pointer, |
| | and read the input data into the register. |
| ReadDataReduce | IO, read data in Reduce form, read the input data from |
| | global memory into register for Reduce. |
+--------------------------------------+-------------------------------------------------------------+
| WriteData | IO, write data from registers to global memory. |
+--------------------------------------+-------------------------------------------------------------+
Expand Down
2 changes: 1 addition & 1 deletion docs/guides/11_kernel_primitives_api/io_api_cn.md
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
# API 介绍 - IO
介绍目前 Kernel Primitive API 提供的用于全局内存和寄存器进行数据交换的 API。当前实现的 IO 类 API 均是 Block 级别的多线程 API。函数内部以 blockDim.x 或 blockDim.y 进行线程索引,因此请详细阅读 API 使用规则。
介绍目前 Kernel Primitive API 提供的用于全局内存和寄存器进行数据交换的 API。当前实现的 IO 类 API 均是 Block 级别的多线程 API。函数内部以 blockDim.x 或 blockDim.y 进行线程索引,请详细阅读 API 使用规则。
## [ReadData](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/fluid/operators/kernel_primitives/datamover_primitives.h#L121)
### 函数定义

Expand Down