                                           int shift_size,
                                           int window_size,
                                           const float eps) {
+#if CUDA_ARCH_FP16_SUPPORTED(__CUDA_ARCH__)
   const int batch_offset = blockIdx.z * gridDim.y * gridDim.x;
   const int bid = batch_offset + blockIdx.y * gridDim.x + blockIdx.x;
   const int shifted_H_idx =

        (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>

  for (int i = 0; i < kITE; i++) {
    int col_id = i * blockDim.x + tid;
    if (col_id < n) {
#if CUDA_ARCH_FP16_SUPPORTED(__CUDA_ARCH__)
      local_out[i] = static_cast<float>(__ldg(input + offset + col_id));
#else
      local_out[i] = static_cast<float>(input[offset + col_id]);
#endif
      sum += local_out[i];
    }
  }
