Skip to content

[Cuda] Updated bfloat16 math defs. #10258

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
masahi merged 1 commit into from
Feb 17, 2022
Merged

[Cuda] Updated bfloat16 math defs. #10258

Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
4 changes: 2 additions & 2 deletions src/target/source/literal/cuda_half_t.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -338,15 +338,15 @@ __pack_nv_bfloat162(const nv_bfloat16 x, const nv_bfloat16 y) {
// so we define them here to make sure the generated CUDA code
// is valid.
#define CUDA_UNSUPPORTED_HALF_MATH_BINARY(HALF_MATH_NAME, FP32_MATH_NAME) \
static inline __device__ __host__ half HALF_MATH_NAME(half x, half y) { \
static inline __device__ __host__ nv_bfloat16 HALF_MATH_NAME(nv_bfloat16 x, nv_bfloat16 y) { \
float tmp_x = __bfloat162float(x); \
float tmp_y = __bfloat162float(y); \
float result = FP32_MATH_NAME(tmp_x, tmp_y); \
return __float2bfloat16(result); \
}

#define CUDA_UNSUPPORTED_HALF_MATH_UNARY(HALF_MATH_NAME, FP32_MATH_NAME) \
static inline __device__ __host__ half HALF_MATH_NAME(half x) { \
static inline __device__ __host__ nv_bfloat16 HALF_MATH_NAME(nv_bfloat16 x) { \
float tmp_x = __bfloat162float(x); \
float result = FP32_MATH_NAME(tmp_x); \
return __float2bfloat16(result); \
Expand Down