   float mean = 0.0f;
   float variance = 0.0f;
 
+#if CUDA_ARCH_FP16_SUPPORTED(__CUDA_ARCH__)
   float local_out =
       (tid < n) ? static_cast<float>(__ldg(input + bid * n + tid)) : 0.0f;
+#else
+  float local_out = (tid < n) ? static_cast<float>(input[bid * n + tid]) : 0.0f;
+#endif
 
   mean = blockReduceSum<float>(local_out);
   if (threadIdx.x == 0) {

                                          int shift_size,
                                          int window_size,
                                          const float eps) {
  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 =

  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__)
      out[output_offset + col_id] =
          (T)(local_out[i] * s_variance *
                  static_cast<float>(__ldg(&gamma[col_id])) +
              static_cast<float>(__ldg(&beta[col_id])));
#else
      out[output_offset + col_id] =
          (T)(local_out[i] * s_variance * static_cast<float>(gamma[col_id]) +
              static_cast<float>(beta[col_id]));
#endif
    }
  }
}

template <>
__global__ void layernorm_shift_partition_v2(half2 *out_ptr,
                                             const half2 *__restrict input_ptr,
