-
Notifications
You must be signed in to change notification settings - Fork 5.6k
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
Sparse Conv3d gpu backward #40143
Conversation
Thanks for your contribution! |
… conv3d_gpu_backward
… conv3d_gpu_backward
… into conv3d_gpu_backward
… conv3d_gpu_backward
… conv3d_gpu_backward
… conv3d_gpu_backward
@@ -157,68 +150,6 @@ __global__ void ProductRuleBookKernel(const int* x_indices, | |||
} | |||
} | |||
|
|||
// TODO(zhangkaihuo): After the GatherCUDAKernel is migrated to phi, replace |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
使用了不同的方法吗 ?
There was a problem hiding this comment.
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); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
这里不能使用ResizeAndAlloc吗?
There was a problem hiding this comment.
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()); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
同上
There was a problem hiding this comment.
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( |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
这个上面不是已经计算过了吗?为什么不直接使用
There was a problem hiding this comment.
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( |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
可以创建一个meta 然后传到这三个Empty中
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
下个PR优化。
PR types
New features
PR changes
OPs
Describe