【Hackathon 5th No.102】 move fused_embedding_eltwise_layernorm/fusion_transpose_flatten_concat/fused_fc_elementwise_layernorm to phi #57865

zeroRains · 2023-10-03T06:19:12Z

PR types

Others

PR changes

Others

Description

move fused_embedding_eltwise_layernorm/fusion_transpose_flatten_concat/fused_fc_elementwise_layernorm to phi
#57262

paddle-bot · 2023-10-03T06:19:17Z

你的PR提交成功，感谢你对开源项目的贡献!
请关注后续CI自动化测试结果，详情请参考Paddle-CI手册。
Your PR has been submitted. Thanks for your contribution!
Please wait for the result of CI firstly. See Paddle CI Manual for details.

CLAassistant · 2023-10-03T10:24:33Z

All committers have signed the CLA.

…mbedding_eltwise_op also have the bug in new IR. because the wrong memory accesss

…ss the test with new IR

yuanlehome · 2023-10-09T10:59:17Z

paddle/phi/infermeta/fusion.cc

+  auto dim_output = phi::make_ddim({batch, seq_len, hidden});
+  out->set_dims(dim_output);
+  // out->share_lod(ids);
+  // context->ShareLoD("Ids", /*->*/ "Out");


后面可以单独提个PR 这里注释给删掉，out->share_lod(ids); 这行为什么给注释掉？

因为在phi 中的share_lod需要一个tensor类型的变量作为参数，但是在这个算子里面的ids是一个vector<DesneTensor*>的类型，暂时不知道怎么处理，所以先留着。这种情况应该怎么处理呢@yuanlehome

因为在phi 中的share_lod需要一个tensor类型的变量作为参数，但是在这个算子里面的ids是一个vector<DesneTensor*>的类型，暂时不知道怎么处理，所以先留着。这种情况应该怎么处理呢@yuanlehome

share第0个，out->share_lod(*ids[0]);

Done，等内存访问异常BUG解决后一起合并

yuanlehome

LGTM

ZzSean

LGTM for OP-Benchmark

yuanlehome · 2023-10-10T06:45:58Z

问下，这个 “开启FLAGS_enable_new_ir_in_executor=1，单测也可以运行成功” 验证了吗？

zeroRains · 2023-10-10T10:38:58Z

问下，这个 “开启FLAGS_enable_new_ir_in_executor=1，单测也可以运行成功” 验证了吗？

额，test_ir_embedding_eltwise_layernorm_fuse_pass这个单测有一个内存访问异常的报错，但是我没看出来哪里有问题。可以麻烦您看一下吗？@yuanlehome

其他两个算子的验证都通过了。

yuanlehome · 2023-10-10T10:43:38Z

问下，这个 “开启FLAGS_enable_new_ir_in_executor=1，单测也可以运行成功” 验证了吗？

额，test_ir_embedding_eltwise_layernorm_fuse_pass这个单测有一个内存访问异常的报错，但是我没看出来哪里有问题。可以麻烦您看一下吗？@yuanlehome

其他两个算子的验证都通过了。

可以 set GLOG_v=5 看下日志，定位下出错位置

zeroRains · 2023-10-11T01:55:46Z

问下，这个 “开启FLAGS_enable_new_ir_in_executor=1，单测也可以运行成功” 验证了吗？

额，test_ir_embedding_eltwise_layernorm_fuse_pass这个单测有一个内存访问异常的报错，但是我没看出来哪里有问题。可以麻烦您看一下吗？@yuanlehome
其他两个算子的验证都通过了。

可以 set GLOG_v=5 看下日志，定位下出错位置

我在kernel中使用了VLOG的方式进行输出定位，但是在我使用命令GLOG_v=5 ctest -R embedding_eltwise_layernorm_fuse_pass -VV运行单测时，却没有显示我在kernel中设置的VLOG输出，请问这是为什么呢？
@yuanlehome

yuanlehome · 2023-10-11T03:50:09Z

问下，这个 “开启FLAGS_enable_new_ir_in_executor=1，单测也可以运行成功” 验证了吗？

额，test_ir_embedding_eltwise_layernorm_fuse_pass这个单测有一个内存访问异常的报错，但是我没看出来哪里有问题。可以麻烦您看一下吗？@yuanlehome
其他两个算子的验证都通过了。

可以 set GLOG_v=5 看下日志，定位下出错位置

我在kernel中使用了VLOG的方式进行输出定位，但是在我使用命令GLOG_v=5 ctest -R embedding_eltwise_layernorm_fuse_pass -VV运行单测时，却没有显示我在kernel中设置的VLOG输出，请问这是为什么呢？ @yuanlehome

对应的实现是这个kernel paddle/phi/kernels/fusion/gpu/fused_embedding_eltwise_layernorm_kernel.cu

zeroRains · 2023-10-11T04:26:08Z

fused_embedding_eltwise_layernorm_kernel

好吧，我看错了。。。sorry

yuanlehome · 2023-10-11T05:10:30Z

fused_embedding_eltwise_layernorm_kernel

好吧，我看错了。。。sorry

我这边定位到在kernel最后copy out，会挂掉，你可以继续看看是不是输入有问题，导致结果计算出错

yuanlehome · 2023-10-11T07:10:35Z

fused_embedding_eltwise_layernorm_kernel

好吧，我看错了。。。sorry

我这边定位到在kernel最后copy out，会挂掉，你可以继续看看是不是输入有问题，导致结果计算出错

加一个同步，PADDLE_ENFORCE_GPU_SUCCESS(cudaDeviceSynchronize());

发现PADDLE_ENFORCE_GPU_SUCCESS(cudaDeviceSynchronize())这一行就挂掉了，说明前面计算有，kernel实现有问题

zeroRains · 2023-10-14T06:36:10Z

fused_embedding_eltwise_layernorm_kernel

好吧，我看错了。。。sorry

我这边定位到在kernel最后copy out，会挂掉，你可以继续看看是不是输入有问题，导致结果计算出错

加一个同步，PADDLE_ENFORCE_GPU_SUCCESS(cudaDeviceSynchronize()); 发现PADDLE_ENFORCE_GPU_SUCCESS(cudaDeviceSynchronize())这一行就挂掉了，说明前面计算有，kernel实现有问题

应该是这个函数EmbEltwiseLayernormKernel出了问题，这个是从paddle/fluid/operators/math/bert_encoder_functor.cu中复制过来的（没怎么改过）应该不会有问题才对，感觉是哪一步索引出了问题。如果从输入数据的角度考虑的话，会不会是这个异步复制数据的操作还没复制完成就执行了kernel导致CUDA地址索引异常呢？

zeroRains · 2023-10-14T07:05:56Z

fused_embedding_eltwise_layernorm_kernel

好吧，我看错了。。。sorry

我这边定位到在kernel最后copy out，会挂掉，你可以继续看看是不是输入有问题，导致结果计算出错

加一个同步，PADDLE_ENFORCE_GPU_SUCCESS(cudaDeviceSynchronize()); 发现PADDLE_ENFORCE_GPU_SUCCESS(cudaDeviceSynchronize())这一行就挂掉了，说明前面计算有，kernel实现有问题

应该是这个函数EmbEltwiseLayernormKernel出了问题，这个是从paddle/fluid/operators/math/bert_encoder_functor.cu中复制过来的（没怎么改过）应该不会有问题才对，感觉是哪一步索引出了问题。如果从输入数据的角度考虑的话，会不会是这个异步复制数据的操作还没复制完成就执行了kernel导致CUDA地址索引异常呢？

尝试使用cudaStreamSynchronize()对复制进行同步，也没能解决问题，(:з」∠)

…transpose_flatten_concat/fused_fc_elementwise_layernorm to phi (PaddlePaddle#57865) * transplant fused_embedding_elt_wise_layer_norm_kernel * fix the error * fix some bug * move the transpose to phi but new IR have a bug in output==nullptr. embedding_eltwise_op also have the bug in new IR. because the wrong memory accesss * remove some useless code * move fused_fc_elementwise_layernorm to phi, but have a bug in making * fix the bug in build the fused_fc_elementwise_layernorm_kernel and pass the test with new IR * try to fix the bug

yuanlehome · 2023-10-16T05:30:30Z

fused_embedding_eltwise_layernorm_kernel

好吧，我看错了。。。sorry

我这边定位到在kernel最后copy out，会挂掉，你可以继续看看是不是输入有问题，导致结果计算出错

加一个同步，PADDLE_ENFORCE_GPU_SUCCESS(cudaDeviceSynchronize()); 发现PADDLE_ENFORCE_GPU_SUCCESS(cudaDeviceSynchronize())这一行就挂掉了，说明前面计算有，kernel实现有问题

应该是这个函数EmbEltwiseLayernormKernel出了问题，这个是从paddle/fluid/operators/math/bert_encoder_functor.cu中复制过来的（没怎么改过）应该不会有问题才对，感觉是哪一步索引出了问题。如果从输入数据的角度考虑的话，会不会是这个异步复制数据的操作还没复制完成就执行了kernel导致CUDA地址索引异常呢？

尝试使用cudaStreamSynchronize()对复制进行同步，也没能解决问题，(:з」∠)

这个我看一下，你先不用管了哈

zeroRains · 2023-10-16T05:58:37Z

fused_embedding_eltwise_layernorm_kernel

好吧，我看错了。。。sorry

我这边定位到在kernel最后copy out，会挂掉，你可以继续看看是不是输入有问题，导致结果计算出错

加一个同步，PADDLE_ENFORCE_GPU_SUCCESS(cudaDeviceSynchronize()); 发现PADDLE_ENFORCE_GPU_SUCCESS(cudaDeviceSynchronize())这一行就挂掉了，说明前面计算有，kernel实现有问题

应该是这个函数EmbEltwiseLayernormKernel出了问题，这个是从paddle/fluid/operators/math/bert_encoder_functor.cu中复制过来的（没怎么改过）应该不会有问题才对，感觉是哪一步索引出了问题。如果从输入数据的角度考虑的话，会不会是这个异步复制数据的操作还没复制完成就执行了kernel导致CUDA地址索引异常呢？

尝试使用cudaStreamSynchronize()对复制进行同步，也没能解决问题，(:з」∠)

这个我看一下，你先不用管了哈

好的，麻烦你了

…transpose_flatten_concat/fused_fc_elementwise_layernorm to phi (PaddlePaddle#57865) * transplant fused_embedding_elt_wise_layer_norm_kernel * fix the error * fix some bug * move the transpose to phi but new IR have a bug in output==nullptr. embedding_eltwise_op also have the bug in new IR. because the wrong memory accesss * remove some useless code * move fused_fc_elementwise_layernorm to phi, but have a bug in making * fix the bug in build the fused_fc_elementwise_layernorm_kernel and pass the test with new IR * try to fix the bug

yuanlehome · 2023-10-16T07:34:17Z

这个PR #58115 修复了 test/ir/test_ir_embedding_eltwise_layernorm_fuse_pass.py 单测在 export FLAGS_enable_new_ir_in_executor=1 时的运行报错。

…transpose_flatten_concat/fused_fc_elementwise_layernorm to phi (PaddlePaddle#57865) * transplant fused_embedding_elt_wise_layer_norm_kernel * fix the error * fix some bug * move the transpose to phi but new IR have a bug in output==nullptr. embedding_eltwise_op also have the bug in new IR. because the wrong memory accesss * remove some useless code * move fused_fc_elementwise_layernorm to phi, but have a bug in making * fix the bug in build the fused_fc_elementwise_layernorm_kernel and pass the test with new IR * try to fix the bug

transplant fused_embedding_elt_wise_layer_norm_kernel

526b4bd

paddle-bot bot added the contributor External developers label Oct 3, 2023

fix the error

3b648d4

zeroRains added 6 commits October 4, 2023 01:59

fix some bug

e5d31f4

move the transpose to phi but new IR have a bug in output==nullptr. e…

a5f9b18

…mbedding_eltwise_op also have the bug in new IR. because the wrong memory accesss

remove some useless code

ac3de17

move fused_fc_elementwise_layernorm to phi, but have a bug in making

c7859b9

fix the bug in build the fused_fc_elementwise_layernorm_kernel and pa…

3949aee

…ss the test with new IR

try to fix the bug

e2ad002

Ligoml mentioned this pull request Oct 7, 2023

【PaddlePaddle Hackathon 5th】开源贡献个人挑战赛 #57262

Open

luotao1 added the PaddlePaddle Hackathon label Oct 9, 2023

luotao1 assigned luotao1 and yuanlehome Oct 9, 2023

yuanlehome reviewed Oct 9, 2023

View reviewed changes

yuanlehome approved these changes Oct 9, 2023

View reviewed changes

ZzSean approved these changes Oct 10, 2023

View reviewed changes

luotao1 merged commit 5ae0019 into PaddlePaddle:develop Oct 10, 2023
27 checks passed

zeroRains deleted the hei branch October 25, 2023 12:04

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

【Hackathon 5th No.102】 move fused_embedding_eltwise_layernorm/fusion_transpose_flatten_concat/fused_fc_elementwise_layernorm to phi #57865

【Hackathon 5th No.102】 move fused_embedding_eltwise_layernorm/fusion_transpose_flatten_concat/fused_fc_elementwise_layernorm to phi #57865

zeroRains commented Oct 3, 2023 •

edited

paddle-bot bot commented Oct 3, 2023

CLAassistant commented Oct 3, 2023 •

edited

yuanlehome Oct 9, 2023

zeroRains Oct 9, 2023 •

edited

yuanlehome Oct 9, 2023

zeroRains Oct 10, 2023

yuanlehome left a comment

ZzSean left a comment

yuanlehome commented Oct 10, 2023

zeroRains commented Oct 10, 2023 •

edited

yuanlehome commented Oct 10, 2023 •

edited

zeroRains commented Oct 11, 2023

yuanlehome commented Oct 11, 2023

zeroRains commented Oct 11, 2023

yuanlehome commented Oct 11, 2023

yuanlehome commented Oct 11, 2023

zeroRains commented Oct 14, 2023

zeroRains commented Oct 14, 2023

yuanlehome commented Oct 16, 2023

zeroRains commented Oct 16, 2023

yuanlehome commented Oct 16, 2023

【Hackathon 5th No.102】 move fused_embedding_eltwise_layernorm/fusion_transpose_flatten_concat/fused_fc_elementwise_layernorm to phi #57865

【Hackathon 5th No.102】 move fused_embedding_eltwise_layernorm/fusion_transpose_flatten_concat/fused_fc_elementwise_layernorm to phi #57865

Conversation

zeroRains commented Oct 3, 2023 • edited

PR types

PR changes

Description

paddle-bot bot commented Oct 3, 2023

CLAassistant commented Oct 3, 2023 • edited

yuanlehome Oct 9, 2023

Choose a reason for hiding this comment

zeroRains Oct 9, 2023 • edited

Choose a reason for hiding this comment

yuanlehome Oct 9, 2023

Choose a reason for hiding this comment

zeroRains Oct 10, 2023

Choose a reason for hiding this comment

yuanlehome left a comment

Choose a reason for hiding this comment

ZzSean left a comment

Choose a reason for hiding this comment

yuanlehome commented Oct 10, 2023

zeroRains commented Oct 10, 2023 • edited

yuanlehome commented Oct 10, 2023 • edited

zeroRains commented Oct 11, 2023

yuanlehome commented Oct 11, 2023

zeroRains commented Oct 11, 2023

yuanlehome commented Oct 11, 2023

yuanlehome commented Oct 11, 2023

zeroRains commented Oct 14, 2023

zeroRains commented Oct 14, 2023

yuanlehome commented Oct 16, 2023

zeroRains commented Oct 16, 2023

yuanlehome commented Oct 16, 2023

zeroRains commented Oct 3, 2023 •

edited

CLAassistant commented Oct 3, 2023 •

edited

zeroRains Oct 9, 2023 •

edited

zeroRains commented Oct 10, 2023 •

edited

yuanlehome commented Oct 10, 2023 •

edited