                                        bias_qk_[threadIdx.x + bias_offset]))
                  : 0.0f;
   float tmp = threadIdx.x < seq_len ? static_cast<float>(qk) : -1e20f;
-  float max_val = blockReduceMax<float>(tmp);
+
+  float max_val = blockReduceMax<float>(tmp, mask);
+
   if (threadIdx.x == 0) s_max = max_val;
   __syncthreads();
 
   float qk_tmp =
       threadIdx.x < seq_len ? __expf(static_cast<float>(tmp - s_max)) : 0.0f;
-  float sum_val = blockReduceSum<float>(qk_tmp);
+  float sum_val = blockReduceSum<float>(qk_tmp, mask);
 
   if (threadIdx.x == 0) {
     s_sum = sum_val + 1e-6f;


/* Calculate the maximum of all elements in a block */
template <typename T>
__inline__ __device__ T blockReduceMax(T val) {
  static __shared__ T shared[WARP_SIZE];
  int lane = threadIdx.x & 0x1f;
  int wid = threadIdx.x >> 5;

  val = warpReduceMax(val);

  if (lane == 0) shared[wid] = val;

  __syncthreads();

  val = (threadIdx.x < (blockDim.x >> 5)) ? shared[lane] : -1e10f;
  val = warpReduceMax(val);

  return val;
}

#define WARP_SIZE 32

template <typename T>
__inline__ __device__ T warpReduceSum(T val, unsigned lane_mask) {
  for (int mask = HALF_WARP; mask > 0; mask >>= 1)
#if __CUDA_ARCH__ >= 350 && CUDA_VERSION >= 9000
    val += __shfl_xor_sync(lane_mask, val, mask, warpSize);
#else
    val += __shfl_xor(val, mask, warpSize);
#endif
