diff --git a/onnxruntime/contrib_ops/cuda/bert/relative_attn_bias.cc b/onnxruntime/contrib_ops/cuda/bert/relative_attn_bias.cc
index 88ebbfe831..af13efe0e2 100644
--- a/onnxruntime/contrib_ops/cuda/bert/relative_attn_bias.cc
+++ b/onnxruntime/contrib_ops/cuda/bert/relative_attn_bias.cc
@@ -57,6 +57,7 @@ Status RelPosAttnBias<T>::ComputeInternal(OpKernelContext* context) const {
 
   typedef typename ToCudaType<T>::MappedType CudaT;
 
+  auto& device_prop = GetDeviceProp();
   return LaunchRelPosAttnBiasKernel<CudaT>(Stream(context),
                                            reinterpret_cast<CudaT*>(output->template MutableData<T>()),
                                            reinterpret_cast<const CudaT*>(bias_table->template Data<T>()),
@@ -64,7 +65,8 @@ Status RelPosAttnBias<T>::ComputeInternal(OpKernelContext* context) const {
                                            static_cast<int>(query_len),
                                            static_cast<int>(num_buckets),
                                            max_distance_,
-                                           is_bidirectional_);
+                                           is_bidirectional_,
+                                           device_prop.maxThreadsPerBlock);
 }
 
 }  // namespace cuda
diff --git a/onnxruntime/contrib_ops/cuda/bert/relative_attn_bias_impl.cu b/onnxruntime/contrib_ops/cuda/bert/relative_attn_bias_impl.cu
index a30fe5b3db..e333152cb5 100644
--- a/onnxruntime/contrib_ops/cuda/bert/relative_attn_bias_impl.cu
+++ b/onnxruntime/contrib_ops/cuda/bert/relative_attn_bias_impl.cu
@@ -36,7 +36,7 @@ __global__ void buildRelativeAttentionBias(T* relative_attention_bias,
                                            const bool is_bidirectional,
                                            const int max_distance) {
   const int head_id = blockIdx.x;
-  for (int seq_id = threadIdx.x; seq_id < seq_len * seq_len; seq_id += blockDim.x) {
+  for (int seq_id = threadIdx.x; seq_id < seq_len * seq_len; seq_id += blockDim.x * gridDim.y) {
     int row_id = seq_id / seq_len;
     int col_id = seq_id % seq_len;
 
@@ -86,20 +86,47 @@ Status LaunchRelPosAttnBiasKernel(
   const int seq_len,
   const int num_bucket,
   const int max_distance,
-  const bool is_bidirectional)
+  const bool is_bidirectional,
+  const int max_threads_per_block)
 {
-  dim3 grid(num_heads);
-  dim3 block(256);
+  const int squared_sq_len = seq_len * seq_len;
+  if (squared_sq_len <= max_threads_per_block) {
+    dim3 grid(num_heads);
+    dim3 block(squared_sq_len);
+    buildRelativeAttentionBias<<<grid, block, 0, stream>>>(output,
+                                                           bias_table,
+                                                           num_heads,
+                                                           seq_len,
+                                                           num_bucket,
+                                                           is_bidirectional,
+                                                           max_distance);
+    return CUDA_CALL(cudaGetLastError());
+  } else if (seq_len >= 128 && seq_len <= 384) {
+    dim3 grid(num_heads, seq_len);
+    dim3 block(seq_len);
+    buildRelativeAttentionBias<<<grid, block, 0, stream>>>(output,
+                                                           bias_table,
+                                                           num_heads,
+                                                           seq_len,
+                                                           num_bucket,
+                                                           is_bidirectional,
+                                                           max_distance);
+    return CUDA_CALL(cudaGetLastError());
+  } else {
+    int blockSize = max_threads_per_block;
+    const int grid_y_Size = (squared_sq_len + blockSize - 1) / blockSize;
+    dim3 grid(num_heads, grid_y_Size);
+    dim3 block(blockSize);
+    buildRelativeAttentionBias<<<grid, block, 0, stream>>>(output,
+                                                           bias_table,
+                                                           num_heads,
+                                                           seq_len,
+                                                           num_bucket,
+                                                           is_bidirectional,
+                                                           max_distance);
 
-  buildRelativeAttentionBias<<<grid, block, 0, stream>>>(output,
-                                                         bias_table,
-                                                         num_heads,
-                                                         seq_len,
-                                                         num_bucket,
-                                                         is_bidirectional,
-                                                         max_distance);
-
-  return CUDA_CALL(cudaGetLastError());
+    return CUDA_CALL(cudaGetLastError());
+  }
 }
 
 template Status LaunchRelPosAttnBiasKernel<float>(cudaStream_t stream,
@@ -109,7 +136,8 @@ template Status LaunchRelPosAttnBiasKernel<float>(cudaStream_t stream,
                                                   const int seq_len,
                                                   const int num_bucket,
                                                   const int max_distance,
-                                                  const bool is_bidirectional);
+                                                  const bool is_bidirectional,
+                                                  const int max_threads_per_block);
 
 template Status LaunchRelPosAttnBiasKernel<half>(cudaStream_t stream,
                                                  half* output,
@@ -118,7 +146,8 @@ template Status LaunchRelPosAttnBiasKernel<half>(cudaStream_t stream,
                                                  const int seq_len,
                                                  const int num_bucket,
                                                  const int max_distance,
-                                                 const bool is_bidirectional);
+                                                 const bool is_bidirectional,
+                                                 const int max_threads_per_block);
 
 }  // namespace cuda
 }  // namespace contrib
diff --git a/onnxruntime/contrib_ops/cuda/bert/relative_attn_bias_impl.h b/onnxruntime/contrib_ops/cuda/bert/relative_attn_bias_impl.h
index a1efd2755f..5a1a229ab6 100644
--- a/onnxruntime/contrib_ops/cuda/bert/relative_attn_bias_impl.h
+++ b/onnxruntime/contrib_ops/cuda/bert/relative_attn_bias_impl.h
@@ -18,7 +18,8 @@ Status LaunchRelPosAttnBiasKernel(
     const int seq_len,
     const int num_bucket,
     const int max_distance,
-    const bool is_bidirectional
+    const bool is_bidirectional,
+    const int max_threads_per_block
 );
 
 }  // namespace cuda