                                              int shift_size,
                                              int window_size,
                                              const float eps) {
+#if CUDA_ARCH_FP16_SUPPORTED(__CUDA_ARCH__)
   // constexpr int ite = 4;
   const int tid = threadIdx.x;
   const int batch_offset = blockIdx.z * gridDim.y * gridDim.x;

  __syncthreads();

  if (tid < n) {
    out[output_bid * n + tid] =
        (T)(((local_out - s_mean) * rsqrtf(s_variance)) *
                static_cast<float>(__ldg(&gamma[tid])) +
            static_cast<float>(__ldg(&beta[tid])));
  }
}

#if CUDA_ARCH_FP16_SUPPORTED(__CUDA_ARCH__)
template <>
__global__ void layernorm_shift_partition(half2 *out_ptr,
                                          const half2 *input_ptr,

        (local_out_fp2.y - s_mean) * s_variance * gamma_val.y + beta_val.y;
    out_ptr[output_bid * n + tid] = __float22half2_rn(local_out_fp2);
  }
#endif
}

#define kITE 4
template <typename T>
