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
add example_reduce.png
  • Loading branch information
AnnaTrainingG committed Sep 30, 2021
commit 16f2cffe3ee6fd6c86744abdc0bcc7e3a2d12da6
6 changes: 4 additions & 2 deletions docs/guides/11_kernel_primitives_api/reduce_example_cn.md
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,9 @@ struct AddFunctor {
```
### kernel 实现说明

对最高维进行规约操作,将不需要进行规约的维度进行合并,将 blockIdx.x 映射到不需要进行规约的维度,保证访问存储效率最大。线程间数据没有依赖,只需要进行线程内规约操作。当num < blockDim.x时需要将 IsBounary 设置为 true,表明需要进行访存边界判断,避免访问存储越界。
对最高维进行规约操作,将不需要进行规约的维度进行合并,将 blockIdx.x 映射到不需要进行规约的维度,保证访问存储效率最大。线程间数据没有依赖,只需要进行线程内规约操作。当num < blockDim.x时需要将 IsBounary 设置为 true,表明需要进行访存边界判断,避免访问存储越界。</br>
ReduceSum 数据处理过程如下:</br>
![Reduce](./images/example_reduce.png)

### kernel 代码

Expand All @@ -35,7 +37,7 @@ __device__ void HigherDimImp(const Tx* x, Ty* y, ReduceOp reducer,
TransformOp transform, MPType init,
int reduce_num, int left_num,
int block_num) {
const int NY = 1;
const int NY = 2;
int idx = blockIdx.x * blockDim.x;
int idy = blockIdx.y * block_num; // block_offset of rows
Tx reduce_input[NY];
Expand Down
7 changes: 5 additions & 2 deletions docs/guides/11_kernel_primitives_api/reduce_example_en.md
Original file line number Diff line number Diff line change
Expand Up @@ -23,7 +23,10 @@ struct AddFunctor {
```
### Kernel Description

Perform the reduction operation on the highest dimension, merge the dimensions that do not need to be reduced, and map blockIdx.x to the dimensions that do not need to be reduced to ensure maximum access and storage efficiency. There is no dependency on data between threads, only intra-thread protocol operations are required. When num < blockDim.x, IsBounary needs to be set to true, indicating that the memory access boundary judgment needs to be performed to avoid access to memory out of range.
Perform the reduction operation on the highest dimension, merge the dimensions that do not need to be reduced, and map blockIdx.x to the dimensions that do not need to be reduced to ensure maximum access and storage efficiency. There is no dependency on data between threads, only intra-thread protocol operations are required. When num < blockDim.x, IsBounary needs to be set to true, indicating that the memory access boundary judgment needs to be performed to avoid access to memory out of range.</br>

The data processing process of ReduceSum is as follows:</br>
![Reduce](./images/example_reduce.png)

### Code

Expand All @@ -33,7 +36,7 @@ __device__ void HigherDimImp(const Tx* x, Ty* y, ReduceOp reducer,
TransformOp transform, MPType init,
int reduce_num, int left_num,
int block_num) {
const int NY = 1;
const int NY = 2;
int idx = blockIdx.x * blockDim.x;
int idy = blockIdx.y * block_num; // block_offset of rows
Tx reduce_input[NY];
Expand Down