   // set effective_element_count as the num of elements when warps do effective
   // work
   // set effective_element_count as 0, when warps do ineffective work
-  int effective_element_count = (effective_warp_id <= 0) ? 0 : element_count;
+  int effective_element_count = (batch_id < batch_size) ? element_count : 0;
   for (int it = 0; it < warp_iter; ++it) {
     int element_index = thread_in_warp_idx + it * kernel_warp_size;
     if (element_index < effective_element_count) {

  }
};

}  // operators
}  // paddle


  constexpr int warp_iter = near_greater_power_of_two / kernel_warp_size;
  int batch_id = blockDim.y * blockIdx.x + threadIdx.y;

  int thread_in_warp_idx = threadIdx.x;

  // 1.read data from global memory to registers
