a.cu

#define DELEGATE_FAST_GELU_GRADIENT_CUDA_KERNEL(T, FMAFunc, TanhFunc)    \
  template <>                                                            \
  __global__ void FastGeluGradientCUDAKernel<T>(                         \
      const int N, const T* dY, const T* X, T* dX) {                     \
    constexpr T kAlpha = M_2_SQRTPI * M_SQRT1_2;                         \
    constexpr T kBeta = kAlpha * gelu_utils::kFastCoeff * T(3);          \
    const int index = blockIdx.x * CAFFE_CUDA_NUM_THREADS + threadIdx.x; \
    if (index < N) {                                                     \
      const T y = TanhFunc(                                              \
          kAlpha *                                                       \
          FMAFunc(                                                       \
              gelu_utils::kFastCoeff,                                    \
              math::utils::Cube<T>(X[index]),                            \
              X[index]));                                                \
      dX[index] = FMAFunc(                                               \
                      FMAFunc(-X[index], y * y, X[index]),               \
                      FMAFunc(kBeta, X[index] * X[index], kAlpha),       \
                      T(1) + y) *                                        \
          dY[index] * static_cast<T>(0.5);                               \
    }                                                                    \
  }