             const auto u0 = *up;
             const auto u1 = *(up + 1);
             const auto u2 = *(up + 2);
-            const auto x_val = (skip_type) ? (*xp) : 0;
+            const auto x_val = (skip_type) ? (*xp) : (scalar_t)0;
             const auto gh_val = *ghp;
             const auto g1 = sigmoidf(u1 + wc1*prev_c_val + bias1);
             const auto g2 = sigmoidf(u2 + wc2*prev_c_val + bias2);

    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));
    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_AND_HALF(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,
