conv shift op: change to CamelCase & fix bug #5558
Conversation
mkliegl
force-pushed
the
conv_shift_fix_camel_case
branch
2 times, most recently
from
dffdfa7
to
77b400b
Compare
|
@chengduoZH The tests pass now after rebasing. Could you please take a look? |
paddle/operators/conv_shift_op.cu
Outdated
| int batch_size) { | ||
| __global__ void ConvShiftForward(const T *x, const T *y, T *out, int x_width, | ||
| int y_width, int y_half_width, | ||
| int batch_size) { | ||
| extern __shared__ T mem[]; | ||
|
|
||
| int tx = threadIdx.x; |
There was a problem hiding this comment.
Although the problem mentioned in issue has been solved, there is a problem with the ConvShiftForward function.
Please recheck line 62~68, you should not use return before __syncthreads();
There was a problem hiding this comment.
Good catch. Let me fix this now quickly.
|
@mkliegl |
|
If some functions are used in Op, the functions are generally written in the form of functor, such as LSTM. |
mkliegl
changed the title
|
@chengduoZH Thank you for your review! I added a fix for the issue you pointed out now rather than making another PR. Could you please check one more time? Regarding making functors: It doesn't seem likely that functors for the conv shift forward and backward kernels are going to be needed anywhere else, so I'm not sure it is worth making this change now. But if you still prefer this for consistency or other reasons, let me know and I can do it. Or else please feel free to just make the changes yourself. |
paddle/operators/conv_shift_op.cu
Outdated
| int batch_size) { | ||
| __global__ void ConvShiftForward(const T *x, const T *y, T *out, int x_width, | ||
| int y_width, int y_half_width, | ||
| int batch_size) { |
There was a problem hiding this comment.
Please reorder parameter.
Function_Parameter_Ordering
The following function parameters are the same.
paddle/operators/conv_shift_op.cu
Outdated
| @@ -160,20 +160,20 @@ class ConvShiftGradKernel<platform::GPUPlace, T> | |||
| auto stream = context.cuda_device_context().stream(); | |||
|
|
|||
| const int x_per_block = 256; | |||
| int num_x_blocks = div_up(x_width, x_per_block); | |||
| int num_x_blocks = DivUp(x_width, x_per_block); | |||
| dim3 grid_dim(num_x_blocks, y_width, batch_size); | |||
|
|
|||
| if (dX) { | |||
| T *dx_data = dX->mutable_data<T>(context.GetPlace()); | |||
| cudaMemsetAsync(dx_data, 0, dX->numel() * sizeof(T), stream); | |||
There was a problem hiding this comment.
You can clean memory in this way.
|
@mkliegl |
mkliegl
force-pushed
the
conv_shift_fix_camel_case
branch
from
6d218bb
to
372c365
Compare
mkliegl
force-pushed
the
conv_shift_fix_camel_case
branch
from
372c365
to
d0b601c
Compare
Addressing part of issue #5549 .