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

Fix a bug in ReadData, ReadDataBc and ReadDataReduce when NX != 1 #36373

Merged
merged 32 commits into from
Oct 21, 2021

Conversation

AnnaTrainingG
Copy link
Contributor

@AnnaTrainingG AnnaTrainingG commented Oct 12, 2021

PR types

Others

PR changes

APIs

Describe

  1. Fix a bug in ReadData, ReadDataBc and ReadDataReduce when NX != 1
  2. update ReduceAnyKernel with kernel primitive api

问题描述:
在进行2D数据读取时,线程数据访问起始位置设置成 threadIdx.x * NX, 对于带stride的读取方式,正确的线程数据起始位置应该为 threadIdx.x; 原始情况会出现数据访存越界的情况,如下图所示:
截屏2021-10-21 上午10 43 47

出错的场景: 当使用大Kernel Primitive API 中的 2D 数据读取时, 只要将NX 设置成大于 1的数就会报错。

出错case: 当前实现并未涉及 NX > 1 的case

AnnaTrainingG and others added 22 commits March 25, 2021 16:46
@paddle-bot-old
Copy link

Thanks for your contribution!
Please wait for the result of CI firstly. See Paddle CI Manual for details.

@AnnaTrainingG AnnaTrainingG changed the title Primitive init Fix a bug in ReadData, ReadDataBc and ReadDataReduce when NX != 1 Oct 13, 2021
zkh2016
zkh2016 previously approved these changes Oct 13, 2021
Liu-xiandong
Liu-xiandong previously approved these changes Oct 13, 2021
@limin2021
Copy link
Contributor

LGTM for modifications in attn_bias_add.cu.h.

xingfeng01
xingfeng01 previously approved these changes Oct 18, 2021
zkh2016
zkh2016 previously approved these changes Oct 18, 2021
Liu-xiandong
Liu-xiandong previously approved these changes Oct 18, 2021
@ZzSean
Copy link
Contributor

ZzSean commented Oct 21, 2021

LGTM

@zhangting2020
Copy link
Contributor

zhangting2020 commented Oct 21, 2021

修复bug类的PR,描述里需要写清楚场景,最好贴上报错log、复现的case

@AnnaTrainingG AnnaTrainingG merged commit 921c091 into PaddlePaddle:develop Oct 21, 2021
AnnaTrainingG added a commit to AnnaTrainingG/Paddle that referenced this pull request Oct 21, 2021
…ddlePaddle#36373)

* Fix a bug in ReadData, ReadDataBc and ReadDataReduce when NX != 1
* Update the implement of reduceAnyKernel according to kernel primitive api
lanxianghit pushed a commit that referenced this pull request Oct 22, 2021
…6373) (#36616)

* Fix a bug in ReadData, ReadDataBc and ReadDataReduce when NX != 1
* Update the implement of reduceAnyKernel according to kernel primitive api
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

7 participants