Skip to content

Fix BFloat16 logits returned as garbage in Logits::Get() #2203

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
kunal-vaishnavi merged 3 commits into from
Jun 18, 2026
Merged

Fix BFloat16 logits returned as garbage in Logits::Get() #2203

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
7c912eb
Fix BFloat16 logits returned as garbage in Logits::Get()
justinchuby Jun 9, 2026
6822c1e
Apply suggestions from code review
kunal-vaishnavi Jun 9, 2026
ecb6086
Merge branch 'microsoft:main' into justinchu/bf16-logits-cast
justinchuby Jun 12, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions src/cuda/interface.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
1 change: 1 addition & 0 deletions src/cuda/kernels.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);

Expand Down
14 changes: 14 additions & 0 deletions src/cuda/model_kernels.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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>
Expand Down Expand Up @@ -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) {
int block_size = 256;
Comment thread
kunal-vaishnavi marked this conversation as resolved.
Outdated
int num_blocks = (count + block_size - 1) / block_size;
Comment thread
kunal-vaishnavi marked this conversation as resolved.
Outdated
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) {
Expand Down
6 changes: 3 additions & 3 deletions src/models/logits.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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_);
Comment thread
justinchuby marked this conversation as resolved.

logits_of_last_token = output_last_tokens_.get();
Expand All @@ -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();
}
Expand Down
Loading