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
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
temp
  • Loading branch information
AnnaTrainingG committed Sep 22, 2021
commit da7593c82e2b5b29c1ed51c6de8683328a553c7b
60 changes: 20 additions & 40 deletions docs/guides/11_kernel_primitives_api/io_api_cn.md
Original file line number Diff line number Diff line change
Expand Up @@ -13,25 +13,21 @@ __device__ void ReadData(Ty* dst, const Tx* src, int size_nx, int size_ny, int s

### 模板参数

```
Tx :数据存储在全局内存中的数据类型。
Ty :数据存储到寄存器上的类型。
NX :每个线程读取 NX 列数据。
NY :每个线程读取 NY 行数据。
BlockSize :设备属性,标识当前设备线程索引方式。对于 GPU,threadIdx.x 用作线程索引,对于 XPU,core_id() 用作线程索引。
IsBoundary :标识是否进行访存边界判断。当block处理的数据总数小于 NX x NY x blockDim.x 时,需要进行边界判断以避免访存越界。
```
> Tx :数据存储在全局内存中的数据类型。
> Ty :数据存储到寄存器上的类型。
> NX :每个线程读取 NX 列数据。
> NY :每个线程读取 NY 行数据。
> BlockSize :设备属性,标识当前设备线程索引方式。对于 GPU,threadIdx.x 用作线程索引,对于 XPU,core_id() 用作线程索引。
> IsBoundary :标识是否进行访存边界判断。当block处理的数据总数小于 NX x NY x blockDim.x 时,需要进行边界判断以避免访存越界。

### 函数参数

```
dst :输出寄存器指针,数据类型为Ty, 大小为 NX x NY。
src :当前 block 的输入数据指针,数据类型为 Tx,指针计算方式通常为 input + blockIdx.x x blockDim.x x NX。
size_nx :block 需要读取 size_nx 列数据,参数仅在 IsBoundary=true 时使用。
size_ny :block 需要读取 size_ny 行数据,参数仅在 IsBoundary=true 时使用。
stride_nx :每读取 1 列数据需要偏移 stride_nx 列。
stride_ny :每读取 NX 列需要偏移 stride_nx 行。
```
> dst :输出寄存器指针,数据类型为Ty, 大小为 NX x NY。
> src :当前 block 的输入数据指针,数据类型为 Tx,指针计算方式通常为 input + blockIdx.x x blockDim.x x NX。
> size_nx :block 需要读取 size_nx 列数据,参数仅在 IsBoundary=true 时使用。
> size_ny :block 需要读取 size_ny 行数据,参数仅在 IsBoundary=true 时使用。
> stride_nx :每读取 1 列数据需要偏移 stride_nx 列。
> stride_ny :每读取 NX 列需要偏移 stride_nx 行。

------------------

Expand All @@ -51,21 +47,17 @@ __device__ void ReadData(T* dst, const T* src, int num);

### 模板参数

```
T :元素类型。
NX :每个线程读取 NX 列数据。
NY :每个线程读取 NY 行数据,当前仅支持为NY = 1。
BlockSize :设备属性,标识当前设备线程索引方式。对于 GPU,threadIdx.x 用作线程索引,对于 XPU,core_id() 用作线程索引。
IsBoundary :标识是否进行访存边界判断。当block处理的数据总数小于 NX x NY x blockDim.x 时,需要进行边界判断以避免访存越界。
```
> T :元素类型。
> NX :每个线程读取 NX 列数据。
> NY :每个线程读取 NY 行数据,当前仅支持为NY = 1。
> BlockSize :设备属性,标识当前设备线程索引方式。对于 GPU,threadIdx.x 用作线程索引,对于 XPU,core_id() 用作线程索引。
> IsBoundary :标识是否进行访存边界判断。当block处理的数据总数小于 NX x NY x blockDim.x 时,需要进行边界判断以避免访存越界。

### 函数参数

```
dst : 输出寄存器指针,大小为 NX x NY。
src : 当前 block 的输入数据指针,通常为 input + blockIdx.x x blockDim.x x NX。
num : 当前 block 最多读取 num 个元素,参数仅在 IsBoundary = true 时使用。
```
> dst : 输出寄存器指针,大小为 NX x NY。
> src : 当前 block 的输入数据指针,通常为 input + blockIdx.x x blockDim.x x NX。
> num : 当前 block 最多读取 num 个元素,参数仅在 IsBoundary = true 时使用。

------------------

Expand All @@ -89,26 +81,22 @@ __device__ void ReadDataBc(T* dst, const T* src,

### 模板参数

```
T :元素类型。
NX :每个线程读取 NX 列数据。
NY :每个线程读取 NY 行数据。
BlockSize :设备属性,标识当前设备线程索引方式。对于 GPU,threadIdx.x 用作线程索引,而对于XPU,core_id() 用作线程索引。
Rank :原始输出数据的维度。
IsBoundary :标识是否进行访存边界判断。当block处理的数据总数小于 NX x NY x blockDim.x 时,需要进行边界判断以避免访存越界。
```

### 函数参数

```
dst :输出寄存器指针,大小为 NX x NY。
src :原始输入数据指针。
block_offset :当前block的数据偏移,通常为 blockIdx.x * blockDim.x * NX。
config :输入输出坐标映射函数,可通过 BroadcastConfig(const std::vector<int64_t>& out_dims, const std::vector<int64_t>& in_dims, int dim_size) 进行定义。
total_num_output :原始输出的总数据个数,避免访存越界,参数仅在 IsBoundary = true 时使用。
stride_nx :每读取 1 列数据需要偏移 stride_nx 列。
stride_ny :每读取 NX 列需要偏移 stride_nx 行。
```


------------------
Expand Down Expand Up @@ -136,7 +124,6 @@ __device__ void ReadDataReduce(T* dst,

### 模板参数

```
T :元素类型
NX :每个线程读取 NX 列数据。
NY :每个线程读取 NY 行数据。
Expand All @@ -150,11 +137,9 @@ IndexCal :输入输出坐标映射规则。定义方式如下:
};
IsBoundary :标识是否进行访存边界判断。当block处理的数据总数小于 NX x NY x blockDim 时,需要进行边界判断以避免访存越界。

```

### 函数参数

```
dst :输出寄存器指针,大小为 NX x NY。
src :原始输入数据指针。
block_offset : 当前block的数据偏移,通常为 blockIdx.x * blockDim.x * NX。
Expand All @@ -164,7 +149,6 @@ size_ny : block 需要读取 size_ny 行数据,参数仅在 IsBoundary = true
stride_nx : 每读取 1 列数据需要偏移 stride_nx 列。
stride_ny : 每读取 NX 列需要偏移 stride_nx 行。
reduce_last_dim:原始输入数据的最低维是否进行reduce,当reduce_last_dim = true 按照 threadIdx.x 进行索引,否则使用 threadIdx.y。
```

------------------

Expand All @@ -184,18 +168,14 @@ __device__ void WriteData(T* dst, T* src, int num);

### 模板参数

```
T :元素类型。
NX :每个线程读取 NX 列数据。
NY :每个线程读取 NY 行数据, 当前仅支持为NY = 1。
BlockSize :设备属性,标识当前设备线程索引方式。对于 GPU,threadIdx.x 用作线程索引,而对于XPU,core_id() 用作线程索引。
IsBoundary :标识是否进行访存边界判断。当block处理的数据总数小于 NX x NY x blockDim 时,需要进行边界判断以避免访存越界。
```

### 函数参数

```
dst : 当前 block 的输出数据指针,通常为 input + blockIdx.x x blockDim.x x NX。
src : 寄存器指针,大小为 NX x NY。,通常为 input + blockIdx.x * blockDim.x * NX。
num : 当前 block 对多读取 num 个元素,参数仅在 IsBoundary = true 时使用。
```