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

Support MaskedSelectGrad op with Kernel Primitive API #40617

Merged
merged 6 commits into from
Mar 21, 2022

Conversation

AnnaTrainingG
Copy link
Contributor

PR types

Others

PR changes

Others

Describe

Support MaskedSelectGrad op with Kernel Primitive API

  1. 添加单线程ReadData API
  2. 添加SelectCaller 以支持masked_select_grad,通过MaskData模板参数区分MaskData = 0 where_index; MaskData = 1 masked_select; MaskData = 2 masked_select_grad

@paddle-bot-old
Copy link

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

typename Functor,
int VecSize,
int IsBoundary,
int IsMaskData>
int MaskData>
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

建议在下个PR注释里说明下 maskdata =0,1,2 对应的情况

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

好的

@@ -123,6 +123,15 @@ __device__ __forceinline__ void WriteData(T* dst,
dst[i] = src[i];
}
}

template <typename T>
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

这个和之前的 readdata 可以复用吗 ?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

不可以 这个是线程级别的API

#include "paddle/phi/kernels/masked_select_grad_kernel.h"

#include "paddle/phi/backends/cpu/cpu_context.h"
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

这个头文件应该没用到,可以删掉试试

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

好的下个PR再修改


SelectGradWithPrefixMask<T><<<grid, threads, 0, stream>>>(
mask_prefix_sum_data, mask_data, input_data, out_data, mask_size);
auto* out_data = x_grad->mutable_data<T>(dev_ctx.GetPlace());
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

kernel里分配内存调用新接口:dev_ctx.template Alloc(x_grad)

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

好的 下个PR再修改

@AnnaTrainingG AnnaTrainingG merged commit f53db25 into PaddlePaddle:develop Mar 21, 2022
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.

4 participants