saymrwulf/onnxruntime
1
0
Fork 0
mirror of https://github.com/saymrwulf/onnxruntime.git synced 2026-07-12 17:57:38 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

Revert Gather Grad optimization in PR 6381 targeted for Rocm (#6880)

Browse source 
* revert gather_grad_impl.cu

* put stream changes back in

* restrict changes to commenting launch of optimized version
This commit is contained in:
Suffian Khan 2021-03-04 10:21:49 -08:00 committed by GitHub
parent 54cdb6af71
commit 7915b6709a
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
1 changed files with 14 additions and 12 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

26
orttraining/orttraining/training_ops/rocm/tensor/gather_grad_impl.cu
View file

@ -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<T, Tin, GridDim::maxElementsPerThread>), 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<T, Tin, GridDim::maxElementsPerThread>), 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<T, Tin, GridDim::maxElementsPerThread>), 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) \