From c51b500ca73dd8b9b8ec332735ac609e8dc6269d Mon Sep 17 00:00:00 2001
From: Tianlei Wu <tlwu@microsoft.com>
Date: Mon, 28 Feb 2022 23:11:51 -0800
Subject: [PATCH] replace std::numeric_limits<T> by cub::FpLimits<T> (#10703)

---
 .../contrib_ops/cuda/transformers/beam_search_impl.cu    | 9 +++++----
 1 file changed, 5 insertions(+), 4 deletions(-)

diff --git a/onnxruntime/contrib_ops/cuda/transformers/beam_search_impl.cu b/onnxruntime/contrib_ops/cuda/transformers/beam_search_impl.cu
index 5b681457c0..4e171c0c46 100644
--- a/onnxruntime/contrib_ops/cuda/transformers/beam_search_impl.cu
+++ b/onnxruntime/contrib_ops/cuda/transformers/beam_search_impl.cu
@@ -4,6 +4,7 @@
 #include "beam_search_impl.h"
 #include "core/providers/cuda/cuda_common.h"
 #include "core/providers/cuda/cu_inc/common.cuh"
+#include "cub/util_type.cuh"
 
 namespace onnxruntime {
 namespace contrib {
@@ -109,27 +110,27 @@ __global__ void LogitsProcessKernel(
       }
 
       if (found) {
-        next_token_scores[index] = std::numeric_limits<T>::lowest();
+        next_token_scores[index] = cub::FpLimits<T>::Lowest();
         return;
       }
     }
 
     // VocabMaskLogitsProcessor
     if (vocab_mask != nullptr && vocab_mask[word_id] == 0) {
-      next_token_scores[index] = std::numeric_limits<T>::lowest();
+      next_token_scores[index] = cub::FpLimits<T>::Lowest();
       return;
     }
 
     // PrefixVocabMaskLogitsProcessor
     int batch_id = batch_beam_index / num_beams;
     if (prefix_vocab_mask != nullptr && prefix_vocab_mask[batch_id * vocab_size + word_id] == 0) {
-      next_token_scores[index] = std::numeric_limits<T>::lowest();
+      next_token_scores[index] = cub::FpLimits<T>::Lowest();
       return;
     }
 
     // MinLengthLogitsProcessor
     if (word_id == demote_token_id) {
-      next_token_scores[index] = std::numeric_limits<T>::lowest();
+      next_token_scores[index] = cub::FpLimits<T>::Lowest();
     }
   }
 }