diff --git a/src/cuda/interface.cpp b/src/cuda/interface.cpp
index 2f273d0173..d0ffde2ec4 100644
--- a/src/cuda/interface.cpp
+++ b/src/cuda/interface.cpp
@@ -122,6 +122,8 @@ struct CudaInterfaceImplBase : DeviceInterface {
       cuda::LaunchFp32ToFp16(reinterpret_cast<const float*>(input_data), reinterpret_cast<uint16_t*>(output_data), static_cast<int>(element_count), GetStream());
     } else if (input_type == Ort::TypeToTensorType<Ort::Float16_t> && output_type == Ort::TypeToTensorType<float>) {
       cuda::LaunchFp16ToFp32(reinterpret_cast<const uint16_t*>(input_data), reinterpret_cast<float*>(output_data), static_cast<int>(element_count), GetStream());
+    } else if (input_type == Ort::TypeToTensorType<Ort::BFloat16_t> && output_type == Ort::TypeToTensorType<float>) {
+      cuda::LaunchBf16ToFp32(reinterpret_cast<const uint16_t*>(input_data), reinterpret_cast<float*>(output_data), static_cast<int>(element_count), GetStream());
     } else if (input_type == Ort::TypeToTensorType<int32_t> && output_type == Ort::TypeToTensorType<int64_t>) {
       cuda::LaunchInt32ToInt64(reinterpret_cast<const int32_t*>(input_data), reinterpret_cast<int64_t*>(output_data), static_cast<int>(element_count), GetStream());
     } else
diff --git a/src/cuda/kernels.h b/src/cuda/kernels.h
index d926563201..19c22f95bd 100644
--- a/src/cuda/kernels.h
+++ b/src/cuda/kernels.h
@@ -13,6 +13,7 @@ void Launch_UpdateAttentionMask(T* mask_data, T* old_data, int batch_beam_size,
 
 void LaunchAddLogitsMask(float* batch_logits, int batch_beam_size, int vocab_size, const uint32_t* logits_mask, cudaStream_t stream);
 void LaunchFp16ToFp32(const uint16_t* fp16, float* fp32, int count, cudaStream_t stream);
+void LaunchBf16ToFp32(const uint16_t* bf16, float* fp32, int count, cudaStream_t stream);
 void LaunchFp32ToFp16(const float* fp32, uint16_t* fp16, int count, cudaStream_t stream);
 void LaunchInt32ToInt64(const int32_t* src, int64_t* dst, int count, cudaStream_t stream);
 
diff --git a/src/cuda/model_kernels.cu b/src/cuda/model_kernels.cu
index 0f96276b8a..f0eccde135 100644
--- a/src/cuda/model_kernels.cu
+++ b/src/cuda/model_kernels.cu
@@ -2,6 +2,7 @@
 // Licensed under the MIT License.
 
 #include <cuda_fp16.h>
+#include <cuda_bf16.h>
 #include <cuda_runtime.h>
 #include <stdint.h>
 #include <limits>
@@ -128,6 +129,19 @@ void LaunchFp32ToFp16(const float* fp32, uint16_t* fp16, int count, cudaStream_t
   CUDA_CHECK_LAUNCH();
 }
 
+__global__ void ConvertBf16ToFp32(const __nv_bfloat16* src, float* dst, int count) {
+  int idx = threadIdx.x + blockIdx.x * blockDim.x;
+  if (idx < count)
+    dst[idx] = __bfloat162float(src[idx]);
+}
+
+void LaunchBf16ToFp32(const uint16_t* bf16, float* fp32, int count, cudaStream_t stream) {
+  constexpr int block_size = 256;
+  const int num_blocks = (count + block_size - 1) / block_size;
+  ConvertBf16ToFp32<<<num_blocks, block_size, 0, stream>>>(reinterpret_cast<const __nv_bfloat16*>(bf16), fp32, count);
+  CUDA_CHECK_LAUNCH();
+}
+
 __global__ void ConvertInt32ToInt64(const int32_t* src, int64_t* dst, int count) {
   int idx = threadIdx.x + blockIdx.x * blockDim.x;
   if (idx < count) {
diff --git a/src/models/logits.cpp b/src/models/logits.cpp
index 477cd38e14..32d1992c43 100644
--- a/src/models/logits.cpp
+++ b/src/models/logits.cpp
@@ -44,7 +44,7 @@ DeviceSpan<float> Logits::Get() {
     // create new OrtValue for logits_of_last_token and use output_last_tokens_ to hold it
     output_last_tokens_ = OrtValue::CreateTensor(model_.p_device_inputs_->GetAllocator(), shape_last, type_);
 
-    if (type_ == Ort::TypeToTensorType<Ort::Float16_t>)
+    if (type_ == Ort::TypeToTensorType<Ort::Float16_t> || type_ == Ort::TypeToTensorType<Ort::BFloat16_t>)
       logits_of_last_token_fp32_ = OrtValue::CreateTensor<float>(model_.p_device_inputs_->GetAllocator(), shape_);
 
     logits_of_last_token = output_last_tokens_.get();
@@ -69,8 +69,8 @@ DeviceSpan<float> Logits::Get() {
     element_count = shape_[0] * shape_[2];  // shape_[1] is now 1, so the element count must be updated
   }
 
-  // Convert from float16 to float32 if necessary
-  if (type_ == Ort::TypeToTensorType<Ort::Float16_t>) {
+  // Convert from float16/bfloat16 to float32 if necessary
+  if (type_ == Ort::TypeToTensorType<Ort::Float16_t> || type_ == Ort::TypeToTensorType<Ort::BFloat16_t>) {
     Cast(*logits_of_last_token, logits_of_last_token_fp32_, *model_.p_device_inputs_, Ort::TypeToTensorType<float>);
     logits_of_last_token = logits_of_last_token_fp32_.get();
   }