     const int total = batch_size * hidden_size;
     const dim3 blocks( (total - 1) / threads + 1 );
 
-    AT_DISPATCH_FLOATING_TYPES(U.type(), "sru_forward_cuda", ([&] {
+    AT_DISPATCH_FLOATING_TYPES_AND_HALF(U.type(), "sru_forward_cuda", ([&] {
         cuda_forward_kernel<scalar_t><<<blocks, threads>>>(
             h.data<scalar_t>(),
             c.data<scalar_t>(),

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

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

template <typename scalar_t>

    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_forward_cuda", ([&] {
        cuda_bi_forward_kernel<scalar_t><<<blocks, threads>>>(
            h.data<scalar_t>(),
            c.data<scalar_t>(),
