   output[idx] = input[idx];
 }
 
-template <typename T, typename Tin>
+template <typename T, typename Tin, int NumElementsPerThread>
 __global__ void _GatherGradImpl(
     const Tin* input,
     const Tin* indices,

      }
      idx++;
    } while (idx < numel && input[idx] == input[idx - 1]);
  }
}


        }

#pragma unroll
        for (int ii = 0; ii < NumElementsPerThread; ii++) {
          weight[ii] += gradient[ii];
        }

#pragma unroll
        for (int ii = 0; ii < NumElementsPerThread; ii++) {
          int feature_dim = start_feature + ii * GPU_WARP_SIZE;
          if (feature_dim < stride) {
            grad_weight[weight_row + feature_dim] = static_cast<T>(weight[ii]);
