| CARVIEW |
Navigation Menu
-
Notifications
You must be signed in to change notification settings - Fork 5.8k
[sparse] Add backend conv2d support #54707
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
Conversation
|
你的PR提交成功，感谢你对开源项目的贡献! |
| @@ -443,4 +443,5 @@ def test2D(self): | |||
|
|
|||
|
|
|||
| if __name__ == "__main__": | |||
| # paddle.device.set_device("cpu") | |||
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.
done
| const auto& kernel_dims = kernel.dims(); | ||
| DDim out_dims = {1, 1, 1, 1, 1}; | ||
|
|
||
| int count_tmp = is2D == true ? 4 : 5; |
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.
count_tmp改成rank
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.
done
| @@ -164,4 +164,4 @@ void LaunchEmbeddingGradDeterministicKernel(const GPUContext& ctx, | |||
| } | |||
|
|
|||
| } // namespace funcs | |||
| } // namespace phi | |||
| } // namespace phi | |||
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.
这个改动去掉。
| phi::errors::InvalidArgument( | ||
| "the shape of kernel should be (D, H, W, C, OC)")); | ||
| bool is2D = out_dims->size() == 4 ? true : false; | ||
| printf("dim = %d\n", out_dims->size()); |
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.
这个printf去掉
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.
done
paddle/phi/kernels/sparse/cpu/conv.h
Outdated
| const int64_t non_zero_num = x.nnz(); | ||
| const auto& indices = x.indices(); | ||
| const IntT* indices_ptr = indices.data<IntT>(); | ||
| int kernel_size = kernel_sizes[0] * kernel_sizes[1] * kernel_sizes[2]; | ||
| int kernel_size = is2D == true |
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.
is2D==true直接改成is2D
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.
done
paddle/phi/kernels/sparse/cpu/conv.h
Outdated
| ddim0 = 1; | ||
| ddim1 = is2D == true ? dilations[1] : dilations[2]; | ||
| ddim2 = is2D == true ? dilations[0] : dilations[1]; | ||
| ddim3 = is2D == true ? 1 : dilations[0]; |
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.
这些地方的复制可以用宏封装一下。is2D==true直接改成is2D
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.
done
paddle/phi/kernels/sparse/cpu/conv.h
Outdated
| IntT in_y = is2D == true ? indices_ptr[i + non_zero_num] | ||
| : indices_ptr[i + 2 * non_zero_num]; | ||
| IntT in_x = is2D == true ? indices_ptr[i + 2 * non_zero_num] | ||
| : indices_ptr[i + 3 * non_zero_num]; |
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.
is2D==true改成is2D
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.
done
| IntT in_x = is2D == true ? indices_ptr[i + 2 * non_zero_num] | ||
| : indices_ptr[i + 3 * non_zero_num]; | ||
| IntT index = phi::funcs::sparse::PointToIndex<Dims4D>( | ||
| batch, in_x, in_y, in_z, c_x_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.
这里为啥改用Dims4D了
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.
Dims4D在2D和3D的环境下内容相同，不用修改原函数
paddle/phi/kernels/sparse/cpu/conv.h
Outdated
| for (int kx = 0; kx < kernel_sizes[2]; kx++) { | ||
| int zceil = is2D == true ? 1 : kernel_sizes[0]; | ||
| int yceil = is2D == true ? kernel_sizes[0] : kernel_sizes[1]; | ||
| int xceil = is2D == true ? kernel_sizes[1] : kernel_sizes[2]; |
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.
is2D==true改成is2D
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.
done
paddle/phi/kernels/sparse/cpu/conv.h
Outdated
| IntT in_x = is2D == true ? indices_ptr[i + 2 * non_zero_num] | ||
| : indices_ptr[i + 3 * non_zero_num]; | ||
|
|
||
| IntT out_z = is2D == true ? 0 : (in_z + paddings[0] - kz * dilations[0]) / strides[0];; |
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.
is2D==true改成is2D
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.
done
paddle/phi/kernels/sparse/cpu/conv.h
Outdated
| @@ -149,7 +201,8 @@ void UpdateRulebookAndOutIndex(const Context& dev_ctx, | |||
| } | |||
|
|
|||
| int out_non_zero_num = out_indexs.size(); | |||
| const int64_t sparse_dim = 4; | |||
| int tmpindex = is2D == true ? 3 : 4; | |||
| const int64_t sparse_dim = tmpindex; | |||
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.
这里写成const int64_t sparse_dim = is2D?3:4;
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.
done
| dilations[i - 1] * (kernel_sizes[i - 1] - 1) - 1) / | ||
| strides[i - 1] + | ||
| 1; | ||
| bool is2D = out_dims->size() == 4 ? true : false; |
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.
声明时可以加个const
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.
done
| int* counter) { | ||
| int* counter, | ||
| bool is2D) { | ||
| // bool is2D = x_dims.size() == 4 ? true : false; |
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.
done
| @@ -332,7 +332,9 @@ __global__ void ProductRuleBookKernel(const T* x_indices, | |||
| const Dims4D dilations, | |||
| const Dims4D strides, | |||
| T* rulebook, | |||
| int* counter) { | |||
| int* counter, | |||
| bool is2D) { | |||
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.
加个const，is2D这个参数要放rulebook和counter前面
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.
done
| @@ -390,12 +392,16 @@ __global__ void GetOutIndexTable1(const IntT* indices, | |||
| const IntT non_zero_num, | |||
| const Dims4D dims, | |||
| int* index_flags, | |||
| int* out_index_table) { | |||
| int* out_index_table, | |||
| bool is2D) { | |||
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.
加const，输入参数要放输出参数前面。
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.
done
| int* out_index_table) { | ||
| int* out_index_table, | ||
| bool is2D) { | ||
| // bool is2D = dims.size() == 4 ? true : false; |
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.
done
| int* counter) { | ||
| int* counter, | ||
| bool is2D) { | ||
| // bool is2D = x_dims.size() == 4 ? true : false; |
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.
done
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.
部分代码的写法请在后面添加sparse convnd时修改。
PR types
Function optimization
PR changes
APIs
Description
修改con3d后端实现以支持2d形状，从而在前端省去reshape的代价。
Pcard-70459