@@ -511,7 +511,7 @@ __global__ void KernelMaxPool2DGradCompatible(
511511 T* input_grad,
512512 FastDivModForPooling<IndexT> divmods,
513513 bool channel_last = false ) {
514- using MPType = typename MPTypeTrait<T>::Type;
514+ using MT = typename MPTypeTrait<T>::Type;
515515
516516 CUDA_KERNEL_LOOP (index, input_height * input_width) {
517517 IndexT h = index / input_width;
@@ -523,19 +523,19 @@ __global__ void KernelMaxPool2DGradCompatible(
523523 T input_data_value = input_data[h * input_width + w];
524524 for (IndexT n = blockIdx .y ; n < batch_size; n += gridDim .y ) {
525525 for (IndexT c = blockIdx .z ; c < channels; c += gridDim .z ) {
526- MPType gradient = static_cast <MPType >(0 .0f );
526+ MT gradient = static_cast <MT >(0 .0f );
527527 IndexT offset = (n * channels + c) * output_height * output_width;
528528 for (int ph = phstart; ph < phend; ++ph) {
529529 for (int pw = pwstart; pw < pwend; ++pw) {
530530 T output_data_value = output_data[ph * output_width + pw + offset];
531531 if (output_data_value == input_data_value) {
532- gradient += static_cast <MPType>(
533- output_grad[ph * output_width + pw + offset]);
532+ gradient +=
533+ static_cast <MT>( output_grad[ph * output_width + pw + offset]);
534534 }
535535 }
536536 }
537537 input_grad[(n * channels + c) * input_height * input_width + index] =
538- static_cast <MPType >(gradient);
538+ static_cast <MT >(gradient);
539539 }
540540 }
541541 }
0 commit comments