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

Sparse Conv3d gpu backward #40143

Merged
merged 30 commits into from
Mar 9, 2022
Merged

Conversation

zkh2016
Copy link
Contributor

@zkh2016 zkh2016 commented Mar 3, 2022

PR types

New features

PR changes

OPs

Describe

  1. SparseConv3d 反向GPU代码,反向CPU代码见PR
  2. 去除不带meta的Empty调用

@paddle-bot-old
Copy link

paddle-bot-old bot commented Mar 3, 2022

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

@@ -157,68 +150,6 @@ __global__ void ProductRuleBookKernel(const int* x_indices,
}
}

// TODO(zhangkaihuo): After the GatherCUDAKernel is migrated to phi, replace
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.

目前换了个位置,后面会把这个方法替换成已有的gatter kernel.

T* in_features_ptr = in_features.data<T>();
T* d_x_features_ptr = d_x_features.data<T>();
T* out_grad_features_ptr = out_grad_features.data<T>();
kernel_grad->Resize(kernel_dims);
Copy link
Contributor

Choose a reason for hiding this comment

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

这里不能使用ResizeAndAlloc吗?

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优化。


// 4. scatter
x_grad->Resize(x.non_zero_elements().dims());
dev_ctx.Alloc(x_grad, x_grad->dtype(), sizeof(T) * x_grad->numel());
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优化。

&unique_key,
&unique_value);

config = phi::backends::gpu::GetGpuLaunchConfig1D(
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.

用同一个变量,中间被覆盖了。

DenseTensor out_index = phi::Empty<int, Context>(dev_ctx);
DenseTensor unique_key = phi::Empty<int, Context>(dev_ctx);
DenseTensor unique_value = phi::Empty<int, Context>(dev_ctx);
DenseTensor out_index = phi::Empty(
Copy link
Contributor

Choose a reason for hiding this comment

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

可以创建一个meta 然后传到这三个Empty中

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优化。

@zkh2016 zkh2016 merged commit 60b86b2 into PaddlePaddle:develop Mar 9, 2022
@zkh2016 zkh2016 deleted the conv3d_gpu_backward branch March 28, 2022 09:32
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.

3 participants