     const auto wc2 = *(weight_c + (col%d) + d);
     const auto bias1 = *(bias + (col%d));
     const auto bias2 = *(bias + (col%d) + d);
-    const auto  mask = (mask_c == NULL) ? 1.0 : (*(mask_c + col));
+    const auto  mask = (mask_c == NULL) ? (scalar_t)1.0 : (*(mask_c + col));
     auto cur = *(init + col);
     const auto* up = u + (col*k);
     const auto* xp = (skip_type == 0) ? NULL : ((skip_type == 1) ? (x + col) : (up + 3));

    const int total = batch_size * hidden_size * 2;
    const dim3 blocks( (total - 1) / threads + 1 );

    AT_DISPATCH_FLOATING_TYPES(U.type(), "sru_bi_backward_cuda", ([&] {
        cuda_bi_backward_kernel<scalar_t><<<blocks, threads>>>(
            grad_u.data<scalar_t>(),
            grad_x.numel() ? grad_x.data<scalar_t>() : NULL,

template <typename scalar_t>
__forceinline__ __device__ scalar_t calc_activation(int type, scalar_t x)
{
    return type ? (scalar_t)tanh(x) : x;
}

template <typename scalar_t>
__forceinline__ __device__ scalar_t calc_grad_activation(int type, scalar_t x)
{
    return type ? (1 - x * x) : (scalar_t)1;
}

template <typename scalar_t>
