From 7915b6709a1dff27d31cdafc9528f3ba451f0e76 Mon Sep 17 00:00:00 2001 From: Suffian Khan Date: Thu, 4 Mar 2021 10:21:49 -0800 Subject: [PATCH] Revert Gather Grad optimization in PR 6381 targeted for Rocm (#6880) * revert gather_grad_impl.cu * put stream changes back in * restrict changes to commenting launch of optimized version --- .../rocm/tensor/gather_grad_impl.cu | 26 ++++++++++--------- 1 file changed, 14 insertions(+), 12 deletions(-) diff --git a/orttraining/orttraining/training_ops/rocm/tensor/gather_grad_impl.cu b/orttraining/orttraining/training_ops/rocm/tensor/gather_grad_impl.cu index 7263591b48..bd988d0402 100644 --- a/orttraining/orttraining/training_ops/rocm/tensor/gather_grad_impl.cu +++ b/orttraining/orttraining/training_ops/rocm/tensor/gather_grad_impl.cu @@ -162,17 +162,19 @@ void GatherGradImpl( dim3 block(GPU_WARP_SIZE, 4); dim3 grid(CeilDiv(num_indices, 4), CeilDiv(stride, GridDim::maxElementsPerThread * GPU_WARP_SIZE)); - if (param_itrs == 1) - { - hipLaunchKernelGGL(HIP_KERNEL_NAME(_GatherAxis0GradImpl), dim3(grid), dim3(block), 0, stream, - indices_data_sorted.get(), - original_indices_sorted.get(), - grad_data, - output_data, - num_indices, - num_inputs, - stride); - } else { + +// commented optimization resulted in increase variance of loss for BERT across multiple reasons +//if (param_itrs == 1) +//{ +// hipLaunchKernelGGL(HIP_KERNEL_NAME(_GatherAxis0GradImpl), dim3(grid), dim3(block), 0, stream, +// indices_data_sorted.get(), +// original_indices_sorted.get(), +// grad_data, +// output_data, +// num_indices, +// num_inputs, +// stride); +//} else { hipLaunchKernelGGL(HIP_KERNEL_NAME(_GatherGradImpl), dim3(grid), dim3(block), 0, stream, indices_data_sorted.get(), original_indices_sorted.get(), @@ -182,7 +184,7 @@ void GatherGradImpl( num_inputs, param_itrs, stride); - } +//} } #define SPECIALIZED_GRAD_IMPL2(T) \