diff --git a/docs/ContribOperators.md b/docs/ContribOperators.md
index d8630076c1..ac2130f36c 100644
--- a/docs/ContribOperators.md
+++ b/docs/ContribOperators.md
@@ -486,7 +486,7 @@ This version of the operator has been available since version 1 of the 'com.micr
#### Type Constraints
-- T : tensor(float16), tensor(float), tensor(double)
+- T : tensor(float16), tensor(float), tensor(double), tensor(bfloat16)
- Constrain input and output types to float tensors.
diff --git a/docs/OperatorKernels.md b/docs/OperatorKernels.md
index c3e2ff94d3..7e311e15c0 100644
--- a/docs/OperatorKernels.md
+++ b/docs/OperatorKernels.md
@@ -729,7 +729,7 @@ Do not modify directly.*
|**Operator Domain:** *com.microsoft*||||
|Attention|*in* input:**T**
*in* weight:**T**
*in* bias:**T**
*in* mask_index:**M**
*in* past:**T**
*in* extra_add:**T**
*out* output:**T**
*out* present:**T**|1+|**T** = tensor(float), tensor(float16)|
|BiasDropout|*in* data:**T**
*in* bias:**T**
*in* residual:**T**
*in* ratio:**T1**
*in* training_mode:**T2**
*out* output:**T**
*out* mask:**T2**|1+|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16)
**T1** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16)
**T2** = tensor(bool)|
-|BiasGelu|*in* A:**T**
*in* B:**T**
*out* C:**T**|1+|**T** = tensor(double), tensor(float), tensor(float16)|
+|BiasGelu|*in* A:**T**
*in* B:**T**
*out* C:**T**|1+|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16)|
|BiasSoftmax|*in* data:**T**
*in* bias:**T**
*out* output:**T**|1+|**T** = tensor(double), tensor(float), tensor(float16)|
|ComplexMul|*in* A:**T**
*in* B:**T**
*out* C:**T**|1+|**T** = tensor(float), tensor(float16)|
|ComplexMulConj|*in* A:**T**
*in* B:**T**
*out* C:**T**|1+|**T** = tensor(float), tensor(float16)|
diff --git a/include/onnxruntime/core/framework/float16.h b/include/onnxruntime/core/framework/float16.h
index 5613b66c48..8de851d9ae 100644
--- a/include/onnxruntime/core/framework/float16.h
+++ b/include/onnxruntime/core/framework/float16.h
@@ -3,9 +3,19 @@
#pragma once
#include "endian.h"
+#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
+#include "cuda_bf16.h"
+#endif
-namespace onnxruntime
-{
+#include "core/common/common.h"
+
+namespace onnxruntime {
+
+#if defined(__CUDACC__) || defined(__HIPCC__)
+#define ORT_HOST_DEVICE __host__ __device__
+#else
+#define ORT_HOST_DEVICE
+#endif
// MLFloat16
struct MLFloat16 {
@@ -17,53 +27,64 @@ struct MLFloat16 {
float ToFloat() const;
- operator float() const {
- return ToFloat();
- }
+ operator float() const { return ToFloat(); }
};
-inline bool operator==(const MLFloat16& left, const MLFloat16& right) {
- return left.val == right.val;
-}
+inline bool operator==(const MLFloat16& left, const MLFloat16& right) { return left.val == right.val; }
+inline bool operator!=(const MLFloat16& left, const MLFloat16& right) { return left.val != right.val; }
+inline bool operator<(const MLFloat16& left, const MLFloat16& right) { return left.val < right.val; }
-inline bool operator!=(const MLFloat16& left, const MLFloat16& right) {
- return left.val != right.val;
-}
-
-inline bool operator<(const MLFloat16& left, const MLFloat16& right) {
- return left.val < right.val;
-}
-
-//BFloat16
+// BFloat16
struct BFloat16 {
uint16_t val{0};
- explicit BFloat16() = default;
- explicit BFloat16(uint16_t v) : val(v) {}
- explicit BFloat16(float v) {
+#if defined(USE_ROCM)
+ ORT_HOST_DEVICE BFloat16() = default;
+#else
+ BFloat16() = default;
+#endif
+
+ struct FromBitsT {};
+ static constexpr ORT_HOST_DEVICE FromBitsT FromBits() { return FromBitsT(); }
+ constexpr ORT_HOST_DEVICE BFloat16(unsigned short bits, FromBitsT) : val(bits){};
+
+ inline ORT_HOST_DEVICE BFloat16(float v) {
+#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000 && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
+ val = __bfloat16_as_ushort(__float2bfloat16(v));
+#else
ORT_IF_CONSTEXPR(endian::native == endian::little) {
std::memcpy(&val, reinterpret_cast(&v) + sizeof(uint16_t), sizeof(uint16_t));
- } else {
+ }
+ else {
std::memcpy(&val, &v, sizeof(uint16_t));
}
+#endif
}
- float ToFloat() const {
+ inline ORT_HOST_DEVICE float ToFloat() const {
+#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
+ return __bfloat162float(*reinterpret_cast(&val));
+#else
float result;
char* const first = reinterpret_cast(&result);
char* const second = first + sizeof(uint16_t);
ORT_IF_CONSTEXPR(endian::native == endian::little) {
std::memset(first, 0, sizeof(uint16_t));
std::memcpy(second, &val, sizeof(uint16_t));
- } else {
+ }
+ else {
std::memcpy(first, &val, sizeof(uint16_t));
std::memset(second, 0, sizeof(uint16_t));
}
return result;
+#endif
}
- operator float() const {
- return ToFloat();
- }
+ inline ORT_HOST_DEVICE operator float() const { return ToFloat(); }
+
+#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
+ ORT_HOST_DEVICE BFloat16(const __nv_bfloat16& value) { val = *reinterpret_cast(&value); }
+ explicit ORT_HOST_DEVICE operator __nv_bfloat16() const { return *reinterpret_cast(&val); }
+#endif
};
inline void BFloat16ToFloat(const BFloat16* blf, float* flt, size_t size) {
@@ -82,16 +103,4 @@ inline void FloatToBFloat16(const float* flt, BFloat16* blf, size_t size) {
}
}
-inline bool operator==(const BFloat16& left, const BFloat16& right) {
- return left.val == right.val;
-}
-
-inline bool operator!=(const BFloat16& left, const BFloat16& right) {
- return left.val != right.val;
-}
-
-inline bool operator<(const BFloat16& left, const BFloat16& right) {
- return left.val < right.val;
-}
-
-}
\ No newline at end of file
+} // namespace onnxruntime
diff --git a/onnxruntime/contrib_ops/cuda/bert/fast_gelu.cc b/onnxruntime/contrib_ops/cuda/bert/fast_gelu.cc
index 864c614b99..29f823d994 100644
--- a/onnxruntime/contrib_ops/cuda/bert/fast_gelu.cc
+++ b/onnxruntime/contrib_ops/cuda/bert/fast_gelu.cc
@@ -25,9 +25,7 @@ namespace cuda {
REGISTER_KERNEL_TYPED(float)
REGISTER_KERNEL_TYPED(MLFloat16)
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
REGISTER_KERNEL_TYPED(BFloat16)
-#endif
using namespace ONNX_NAMESPACE;
diff --git a/onnxruntime/contrib_ops/cuda/bert/fast_gelu_impl.cu b/onnxruntime/contrib_ops/cuda/bert/fast_gelu_impl.cu
index 4e9863edd2..61879024c2 100644
--- a/onnxruntime/contrib_ops/cuda/bert/fast_gelu_impl.cu
+++ b/onnxruntime/contrib_ops/cuda/bert/fast_gelu_impl.cu
@@ -94,11 +94,10 @@ bool LaunchFastGeluKernel(const cudaDeviceProp& prop, cudaStream_t stream, int i
#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
template
-__global__ void FastGeluKernel2(const nv_bfloat162 a, const nv_bfloat162 b, const nv_bfloat162 c,
- int input_length, int bias_length,
- const nv_bfloat162* input, const nv_bfloat162* bias, nv_bfloat162* output) {
+__global__ void FastGeluKernel2(const nv_bfloat162 a, const nv_bfloat162 b, const nv_bfloat162 c, int input_length,
+ int bias_length, const nv_bfloat162* input, const nv_bfloat162* bias,
+ nv_bfloat162* output) {
const int idx = blockIdx.x * TPB + threadIdx.x;
-
if (idx < input_length) {
const nv_bfloat162 x = input[idx];
const nv_bfloat162 in = (bias == nullptr) ? x : (x + bias[idx % bias_length]);
@@ -106,11 +105,13 @@ __global__ void FastGeluKernel2(const nv_bfloat162 a, const nv_bfloat162 b, cons
output[idx] = in * cdf;
}
}
+#endif
template <>
-bool LaunchFastGeluKernel(const cudaDeviceProp& prop, cudaStream_t stream, int input_length, int bias_length, const nv_bfloat16* input, const nv_bfloat16* bias, nv_bfloat16* output, bool /*use_half2*/) {
+bool LaunchFastGeluKernel(const cudaDeviceProp& prop, cudaStream_t stream, int input_length, int bias_length,
+ const BFloat16* input, const BFloat16* bias, BFloat16* output, bool /*use_half2*/) {
constexpr int blockSize = 256;
-
+#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
if (0 == (bias_length & 1) && prop.major >= 7) {
const int n = input_length / 2;
const int gridSize = (n + blockSize - 1) / blockSize;
@@ -120,15 +121,18 @@ bool LaunchFastGeluKernel(const cudaDeviceProp& prop, cudaStream_t stream, int i
const nv_bfloat162* input2 = reinterpret_cast(input);
const nv_bfloat162* bias2 = reinterpret_cast(bias);
nv_bfloat162* output2 = reinterpret_cast(output);
- FastGeluKernel2<<>>(A2, B2, C2, n, bias_length / 2, input2, bias2, output2);
+ FastGeluKernel2
+ <<>>(A2, B2, C2, n, bias_length / 2, input2, bias2, output2);
} else {
+#endif
const int gridSize = (input_length + blockSize - 1) / blockSize;
- FastGeluKernel<<>>(A, B, C, input_length, bias_length, input, bias, output);
+ FastGeluKernel
+ <<>>(A, B, C, input_length, bias_length, input, bias, output);
+#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
}
-
+#endif
return CUDA_CALL(cudaPeekAtLastError());
}
-#endif
} // namespace cuda
} // namespace contrib
diff --git a/onnxruntime/contrib_ops/cuda/cuda_contrib_kernels.cc b/onnxruntime/contrib_ops/cuda/cuda_contrib_kernels.cc
index 19b226858a..442138d3ad 100644
--- a/onnxruntime/contrib_ops/cuda/cuda_contrib_kernels.cc
+++ b/onnxruntime/contrib_ops/cuda/cuda_contrib_kernels.cc
@@ -18,6 +18,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, BiasGelu);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, BiasGelu);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, BiasGelu);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, BFloat16, BiasGelu);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, TransposeMatMul); // backward compatibility
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, TransposeMatMul); // backward compatibility
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, TransposeMatMul); // backward compatibility
@@ -86,13 +87,10 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float_int8_t, QAttention);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16_int8_t, QAttention);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, FusedConv);
-
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, BFloat16, FastGelu);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, BFloat16, TransposeMatMul); // backward compatibility
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, BFloat16, FusedMatMul);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, BFloat16_float, LayerNormalization);
-#endif
template <>
KernelCreateInfo BuildKernelCreateInfo() {
@@ -112,6 +110,7 @@ Status RegisterCudaContribKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
BuildKernelCreateInfo, // backward compatibility
BuildKernelCreateInfo, // backward compatibility
BuildKernelCreateInfo, // backward compatibility
@@ -180,14 +179,11 @@ Status RegisterCudaContribKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
-
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
BuildKernelCreateInfo,
// TransposedMatMul is still here for backward compatibility
BuildKernelCreateInfo, // backward compatibility
BuildKernelCreateInfo,
BuildKernelCreateInfo,
-#endif
BuildKernelCreateInfo,
};
diff --git a/onnxruntime/contrib_ops/cuda/layer_norm.cc b/onnxruntime/contrib_ops/cuda/layer_norm.cc
index 3095ebf437..3f7360bd77 100644
--- a/onnxruntime/contrib_ops/cuda/layer_norm.cc
+++ b/onnxruntime/contrib_ops/cuda/layer_norm.cc
@@ -35,9 +35,7 @@ namespace cuda {
REGISTER_KERNEL_TYPED(float, float)
REGISTER_KERNEL_TYPED(double, double)
REGISTER_KERNEL_TYPED(MLFloat16, float)
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
REGISTER_KERNEL_TYPED(BFloat16, float)
-#endif
template
LayerNorm::LayerNorm(const OpKernelInfo& op_kernel_info) : CudaKernel(op_kernel_info) {
diff --git a/onnxruntime/contrib_ops/cuda/layer_norm_impl.cu b/onnxruntime/contrib_ops/cuda/layer_norm_impl.cu
index 740aae4d45..d476ca3e25 100644
--- a/onnxruntime/contrib_ops/cuda/layer_norm_impl.cu
+++ b/onnxruntime/contrib_ops/cuda/layer_norm_impl.cu
@@ -394,10 +394,8 @@ LAYERNORM_LINEAR_IMPL(half, float, false)
LAYERNORM_LINEAR_IMPL(double, double, false)
//LAYERNORM_LINEAR_IMPL(half, half)
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-LAYERNORM_LINEAR_IMPL(nv_bfloat16, float, true)
-LAYERNORM_LINEAR_IMPL(nv_bfloat16, float, false)
-#endif
+LAYERNORM_LINEAR_IMPL(BFloat16, float, true)
+LAYERNORM_LINEAR_IMPL(BFloat16, float, false)
} // namespace cuda
} // namespace contrib
diff --git a/onnxruntime/contrib_ops/cuda/math/bias_dropout.cc b/onnxruntime/contrib_ops/cuda/math/bias_dropout.cc
index 192c28c86d..a16a1b88ef 100644
--- a/onnxruntime/contrib_ops/cuda/math/bias_dropout.cc
+++ b/onnxruntime/contrib_ops/cuda/math/bias_dropout.cc
@@ -17,8 +17,8 @@ ONNX_OPERATOR_KERNEL_EX(
1,
kCudaExecutionProvider,
(*KernelDefBuilder::Create())
- .TypeConstraint("T", ALL_IEEE_FLOAT_TENSOR_TYPES)
- .TypeConstraint("T1", ALL_IEEE_FLOAT_TENSOR_TYPES)
+ .TypeConstraint("T", BuildKernelDefConstraints())
+ .TypeConstraint("T1", BuildKernelDefConstraints())
.TypeConstraint("T2", DataTypeImpl::GetTensorType())
.InputMemoryType(OrtMemTypeCPUInput, 3)
.InputMemoryType(OrtMemTypeCPUInput, 4),
@@ -96,7 +96,7 @@ Status BiasDropout::ComputeInternal(OpKernelContext* context) const {
float ratio_data = default_ratio_;
auto ratio = context->Input(3);
if (ratio) {
- utils::MLTypeCallDispatcher t_disp(ratio->GetElementType());
+ utils::MLTypeCallDispatcher t_disp(ratio->GetElementType());
t_disp.Invoke(ratio, ratio_data);
}
@@ -117,7 +117,7 @@ Status BiasDropout::ComputeInternal(OpKernelContext* context) const {
const fast_divmod fdm_dim(gsl::narrow_cast(dim));
PhiloxGenerator& generator = generator_ ? *generator_ : PhiloxGenerator::Default();
- utils::MLTypeCallDispatcher t_disp(X->GetElementType());
+ utils::MLTypeCallDispatcher t_disp(X->GetElementType());
return t_disp.InvokeRet(
GetDeviceProp(), Stream(), N, fdm_dim, ratio_data, generator, *X, *bias, residual, *Y, mask_data, has_same_shape_bias);
}
diff --git a/onnxruntime/contrib_ops/cuda/math/bias_dropout_impl.cu b/onnxruntime/contrib_ops/cuda/math/bias_dropout_impl.cu
index ad2ff05a58..727d4f4878 100644
--- a/onnxruntime/contrib_ops/cuda/math/bias_dropout_impl.cu
+++ b/onnxruntime/contrib_ops/cuda/math/bias_dropout_impl.cu
@@ -238,9 +238,7 @@ void BiasDropoutKernelImpl(
SPECIALIZED_BIAS_DROPOUT_IMPL(float)
SPECIALIZED_BIAS_DROPOUT_IMPL(double)
SPECIALIZED_BIAS_DROPOUT_IMPL(half)
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-SPECIALIZED_BIAS_DROPOUT_IMPL(nv_bfloat16)
-#endif
+SPECIALIZED_BIAS_DROPOUT_IMPL(BFloat16)
} // namespace cuda
} // namespace contrib {
diff --git a/onnxruntime/contrib_ops/cuda/math/binary_elementwise_ops.cc b/onnxruntime/contrib_ops/cuda/math/binary_elementwise_ops.cc
index 0df46fb763..6ce2446e44 100644
--- a/onnxruntime/contrib_ops/cuda/math/binary_elementwise_ops.cc
+++ b/onnxruntime/contrib_ops/cuda/math/binary_elementwise_ops.cc
@@ -61,7 +61,8 @@ namespace cuda {
#define CONTRIB_BINARY_OP_HFD(name, ver) \
CONTRIB_BINARY_OP_TYPED(name, ver, MLFloat16) \
CONTRIB_BINARY_OP_TYPED(name, ver, float) \
- CONTRIB_BINARY_OP_TYPED(name, ver, double)
+ CONTRIB_BINARY_OP_TYPED(name, ver, double) \
+ CONTRIB_BINARY_OP_TYPED(name, ver, BFloat16)
CONTRIB_BINARY_OP_HFD(BiasGelu, 1)
diff --git a/onnxruntime/contrib_ops/cuda/math/binary_elementwise_ops_impl.cu b/onnxruntime/contrib_ops/cuda/math/binary_elementwise_ops_impl.cu
index 01791ed94c..bb522fe9e2 100644
--- a/onnxruntime/contrib_ops/cuda/math/binary_elementwise_ops_impl.cu
+++ b/onnxruntime/contrib_ops/cuda/math/binary_elementwise_ops_impl.cu
@@ -63,7 +63,8 @@ namespace cuda {
#define CONTRIB_SPECIALIZED_BINARY_ELEMENTWISE_IMPL_HFD(x) \
CONTRIB_SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, half) \
CONTRIB_SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, float) \
- CONTRIB_SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, double)
+ CONTRIB_SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, double) \
+ CONTRIB_SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, BFloat16)
// create declarations for op and impl
#define CONTRIB_BINARY_OP_NAME_EXPR(name, expr) \
diff --git a/onnxruntime/contrib_ops/cuda/math/fused_matmul.cc b/onnxruntime/contrib_ops/cuda/math/fused_matmul.cc
index 86ce8d3efa..a2fad0ac4d 100644
--- a/onnxruntime/contrib_ops/cuda/math/fused_matmul.cc
+++ b/onnxruntime/contrib_ops/cuda/math/fused_matmul.cc
@@ -22,15 +22,12 @@ namespace cuda {
REGISTER_KERNEL_TYPED(TransposeMatMul, float)
REGISTER_KERNEL_TYPED(TransposeMatMul, double)
REGISTER_KERNEL_TYPED(TransposeMatMul, MLFloat16)
+REGISTER_KERNEL_TYPED(TransposeMatMul, BFloat16)
REGISTER_KERNEL_TYPED(FusedMatMul, float)
REGISTER_KERNEL_TYPED(FusedMatMul, double)
REGISTER_KERNEL_TYPED(FusedMatMul, MLFloat16)
-
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
-REGISTER_KERNEL_TYPED(TransposeMatMul, BFloat16)
REGISTER_KERNEL_TYPED(FusedMatMul, BFloat16)
-#endif
} // namespace cuda
} // namespace contrib
diff --git a/onnxruntime/contrib_ops/cuda/math/isfinite.cuh b/onnxruntime/contrib_ops/cuda/math/isfinite.cuh
index 1116228a9d..d27a6631c1 100644
--- a/onnxruntime/contrib_ops/cuda/math/isfinite.cuh
+++ b/onnxruntime/contrib_ops/cuda/math/isfinite.cuh
@@ -5,10 +5,6 @@
#include "core/providers/cuda/cu_inc/common.cuh"
#include "contrib_ops/cuda/math/isfinite.h"
-#if CUDA_VERSION >= 11000
-#include "cuda_bf16.h"
-#endif
-
namespace onnxruntime {
namespace cuda {
@@ -54,22 +50,20 @@ __device__ __forceinline__ bool IsNaNScalar(const half value) {
#endif
}
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
template <>
-__device__ __forceinline__ bool IsFiniteScalar(const nv_bfloat16 value) {
- return !__hisinf(value) && !__hisnan(value);
+__device__ __forceinline__ bool IsFiniteScalar(const BFloat16 value) {
+ return isfinite(static_cast(value));
}
template <>
-__device__ __forceinline__ bool IsInfScalar(const nv_bfloat16 value) {
- return __hisinf(value);
+__device__ __forceinline__ bool IsInfScalar(const BFloat16 value) {
+ return isinf(static_cast(value));
}
template <>
-__device__ __forceinline__ bool IsNaNScalar(const nv_bfloat16 value) {
- return __hisnan(value);
+__device__ __forceinline__ bool IsNaNScalar(const BFloat16 value) {
+ return isnan(static_cast(value));
}
-#endif
} // namespace cuda
} // namespace onnxruntime
\ No newline at end of file
diff --git a/onnxruntime/core/framework/tensorprotoutils.cc b/onnxruntime/core/framework/tensorprotoutils.cc
index 2b6d2bbed1..08e86d815f 100644
--- a/onnxruntime/core/framework/tensorprotoutils.cc
+++ b/onnxruntime/core/framework/tensorprotoutils.cc
@@ -389,7 +389,7 @@ Status UnpackTensor(const ONNX_NAMESPACE::TensorProto& tensor, const void* raw_d
if (v < 0 || v > max_value) {
return Status(common::ONNXRUNTIME, common::INVALID_ARGUMENT, "data overflow");
}
- p_data[i] = BFloat16(static_cast(v));
+ p_data[i] = BFloat16(static_cast(v), BFloat16::FromBits());
}
return Status::OK();
diff --git a/onnxruntime/core/graph/contrib_ops/contrib_defs.cc b/onnxruntime/core/graph/contrib_ops/contrib_defs.cc
index 502244fa4b..fcc219799b 100644
--- a/onnxruntime/core/graph/contrib_ops/contrib_defs.cc
+++ b/onnxruntime/core/graph/contrib_ops/contrib_defs.cc
@@ -2874,7 +2874,7 @@ It's an extension of Gelu. It takes the sum of input A and bias input B as the i
.Output(0, "C", "The output.", "T")
.TypeConstraint(
"T",
- {"tensor(float16)", "tensor(float)", "tensor(double)"},
+ {"tensor(float16)", "tensor(float)", "tensor(double)", "tensor(bfloat16)"},
"Constrain input and output types to float tensors.")
.TypeAndShapeInferenceFunction(ONNX_NAMESPACE::propagateShapeAndTypeFromFirstInput);
diff --git a/onnxruntime/core/providers/cuda/activation/activations.cc b/onnxruntime/core/providers/cuda/activation/activations.cc
index 08e0f86106..df02e333b3 100644
--- a/onnxruntime/core/providers/cuda/activation/activations.cc
+++ b/onnxruntime/core/providers/cuda/activation/activations.cc
@@ -58,25 +58,17 @@ namespace cuda {
REGISTER_ACTIVATION_KERNEL(name, ver, T) \
UNARY_ACTIVATION_COMPUTE(name, T)
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
-#define UNARY_ACTIVATION_OP_TYPED_BF16(name, ver) UNARY_ACTIVATION_OP_TYPED(name, ver, BFloat16)
-#define UNARY_ACTIVATION_OP_VERSIONED_TYPED_BF16(name, startver, endver) UNARY_ACTIVATION_OP_VERSIONED_TYPED(name, startver, endver, BFloat16)
-#else
-#define UNARY_ACTIVATION_OP_TYPED_BF16(name, ver)
-#define UNARY_ACTIVATION_OP_VERSIONED_TYPED_BF16(name, startver, endver)
-#endif
-
#define UNARY_ACTIVATION_OP_VERSIONED_HFD_WITH_BF16(name, startver, endver) \
UNARY_ACTIVATION_OP_VERSIONED_TYPED(name, startver, endver, MLFloat16) \
UNARY_ACTIVATION_OP_VERSIONED_TYPED(name, startver, endver, float) \
UNARY_ACTIVATION_OP_VERSIONED_TYPED(name, startver, endver, double) \
- UNARY_ACTIVATION_OP_VERSIONED_TYPED_BF16(name, startver, endver)
+ UNARY_ACTIVATION_OP_VERSIONED_TYPED(name, startver, endver, BFloat16)
#define UNARY_ACTIVATION_OP_HFD(name, ver) \
UNARY_ACTIVATION_OP_TYPED(name, ver, MLFloat16) \
- UNARY_ACTIVATION_OP_TYPED_BF16(name, ver) \
UNARY_ACTIVATION_OP_TYPED(name, ver, float) \
- UNARY_ACTIVATION_OP_TYPED(name, ver, double)
+ UNARY_ACTIVATION_OP_TYPED(name, ver, double) \
+ UNARY_ACTIVATION_OP_TYPED(name, ver, BFloat16)
UNARY_ACTIVATION_OP_HFD(Elu, 6);
UNARY_ACTIVATION_OP_HFD(HardSigmoid, 6);
diff --git a/onnxruntime/core/providers/cuda/activation/activations_impl.cu b/onnxruntime/core/providers/cuda/activation/activations_impl.cu
index 2ff5a4748f..1a4dc3cf79 100644
--- a/onnxruntime/core/providers/cuda/activation/activations_impl.cu
+++ b/onnxruntime/core/providers/cuda/activation/activations_impl.cu
@@ -91,20 +91,15 @@ struct OP_ThresholdedRelu : public CtxThresholdedRelu {
count); \
}
-#define SPECIALIZED_UNARY_ACTIVATION_IMPL(name, T) \
- template void Impl_##name(cudaStream_t stream, const T* input_data, T* output_data, const Ctx##name* func_ctx, size_t count);
-
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-#define SPECIALIZED_UNARY_ACTIVATION_IMPL_BF16(name) SPECIALIZED_UNARY_ACTIVATION_IMPL(name, nv_bfloat16)
-#else
-#define SPECIALIZED_UNARY_ACTIVATION_IMPL_BF16(name)
-#endif
+#define SPECIALIZED_UNARY_ACTIVATION_IMPL(name, T) \
+ template void Impl_##name(cudaStream_t stream, const T* input_data, T* output_data, const Ctx##name* func_ctx, \
+ size_t count);
#define SPECIALIZED_UNARY_ACTIVATIONL_HFD(name) \
SPECIALIZED_UNARY_ACTIVATION_IMPL(name, half) \
- SPECIALIZED_UNARY_ACTIVATION_IMPL_BF16(name) \
SPECIALIZED_UNARY_ACTIVATION_IMPL(name, float) \
- SPECIALIZED_UNARY_ACTIVATION_IMPL(name, double)
+ SPECIALIZED_UNARY_ACTIVATION_IMPL(name, double) \
+ SPECIALIZED_UNARY_ACTIVATION_IMPL(name, BFloat16)
#define UNARY_ACTIVATION_OP_NAME(name) \
UNARY_ACTIVATION_IMPL(name); \
diff --git a/onnxruntime/core/providers/cuda/atomic/common.cuh b/onnxruntime/core/providers/cuda/atomic/common.cuh
index 098f650986..d6751f9bd0 100644
--- a/onnxruntime/core/providers/cuda/atomic/common.cuh
+++ b/onnxruntime/core/providers/cuda/atomic/common.cuh
@@ -20,10 +20,7 @@
#include "cuda.h"
#include "cuda_fp16.h"
#include "cuda_runtime.h"
-
-#if CUDA_VERSION >= 11000
-#include "cuda_bf16.h"
-#endif
+#include "core/framework/float16.h"
namespace onnxruntime {
namespace cuda {
@@ -72,22 +69,20 @@ __device__ __forceinline__ void atomic_add(half *address, half value) {
#endif
}
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-__device__ __forceinline__ void atomic_add(nv_bfloat16 *address, nv_bfloat16 value) {
- unsigned int * base_address = reinterpret_cast(reinterpret_cast(address) - (reinterpret_cast(address) & 2));
+__device__ __forceinline__ void atomic_add(BFloat16* address, BFloat16 value) {
+ unsigned int* base_address =
+ reinterpret_cast(reinterpret_cast(address) - (reinterpret_cast(address) & 2));
unsigned int old = *base_address;
unsigned int assumed;
- unsigned short x;
-
+ BFloat16 bsum;
do {
assumed = old;
- x = reinterpret_cast(address) & 2 ? (old >> 16) : (old & 0xffff);
- x = __bfloat16_as_short(__float2bfloat16(__bfloat162float(*reinterpret_cast(&x)) + __bfloat162float(value)));
- old = reinterpret_cast(address) & 2 ? (old & 0xffff) | (x << 16) : (old & 0xffff0000) | x;
+ bsum.val = reinterpret_cast(address) & 2 ? (old >> 16) : (old & 0xffff);
+ bsum = bsum + value;
+ old = reinterpret_cast(address) & 2 ? (old & 0xffff) | (bsum.val << 16) : (old & 0xffff0000) | bsum.val;
old = atomicCAS(base_address, assumed, old);
} while (assumed != old);
}
-#endif
} // namespace cuda
} // namespace onnxruntime
\ No newline at end of file
diff --git a/onnxruntime/core/providers/cuda/cu_inc/common.cuh b/onnxruntime/core/providers/cuda/cu_inc/common.cuh
index 7b283731e9..ea9ddf0450 100644
--- a/onnxruntime/core/providers/cuda/cu_inc/common.cuh
+++ b/onnxruntime/core/providers/cuda/cu_inc/common.cuh
@@ -81,6 +81,82 @@ __device__ __forceinline__ half2 operator*(const half2& lh, const half2& rh) { h
__device__ __forceinline__ half2 operator/(const half2& lh, const half2& rh) { half2 r; r.x = lh.x / rh.x; r.y = lh.y / rh.y; return r; }
#endif
+/// Arithmetic for BFloat16
+
+__device__ __forceinline__ BFloat16 operator+(const BFloat16& a, const BFloat16& b) {
+ return static_cast(a) + static_cast(b);
+}
+
+__device__ __forceinline__ BFloat16 operator-(const BFloat16& a, const BFloat16& b) {
+ return static_cast(a) - static_cast(b);
+}
+
+__device__ __forceinline__ BFloat16 operator*(const BFloat16& a, const BFloat16& b) {
+ return static_cast(a) * static_cast(b);
+}
+
+__device__ __forceinline__ BFloat16 operator/(const BFloat16& a, const BFloat16& b) {
+ return static_cast(a) / static_cast(b);
+}
+
+__device__ __forceinline__ BFloat16 operator-(const BFloat16& a) { return -static_cast(a); }
+
+__device__ __forceinline__ BFloat16& operator+=(BFloat16& a, const BFloat16& b) {
+ a = a + b;
+ return a;
+}
+
+__device__ __forceinline__ BFloat16& operator-=(BFloat16& a, const BFloat16& b) {
+ a = a - b;
+ return a;
+}
+
+__device__ __forceinline__ BFloat16& operator*=(BFloat16& a, const BFloat16& b) {
+ a = a * b;
+ return a;
+}
+
+__device__ __forceinline__ BFloat16& operator/=(BFloat16& a, const BFloat16& b) {
+ a = a / b;
+ return a;
+}
+
+/// Arithmetic with floats
+
+__device__ __forceinline__ float operator+(BFloat16 a, float b) { return static_cast(a) + b; }
+__device__ __forceinline__ float operator-(BFloat16 a, float b) { return static_cast(a) - b; }
+__device__ __forceinline__ float operator*(BFloat16 a, float b) { return static_cast(a) * b; }
+__device__ __forceinline__ float operator/(BFloat16 a, float b) { return static_cast(a) / b; }
+
+__device__ __forceinline__ float operator+(float a, BFloat16 b) { return a + static_cast(b); }
+__device__ __forceinline__ float operator-(float a, BFloat16 b) { return a - static_cast(b); }
+__device__ __forceinline__ float operator*(float a, BFloat16 b) { return a * static_cast(b); }
+__device__ __forceinline__ float operator/(float a, BFloat16 b) { return a / static_cast(b); }
+
+__device__ __forceinline__ float& operator+=(float& a, const BFloat16& b) { return a += static_cast(b); }
+__device__ __forceinline__ float& operator-=(float& a, const BFloat16& b) { return a -= static_cast(b); }
+__device__ __forceinline__ float& operator*=(float& a, const BFloat16& b) { return a *= static_cast(b); }
+__device__ __forceinline__ float& operator/=(float& a, const BFloat16& b) { return a /= static_cast(b); }
+
+/// Arithmetic with doubles
+
+__device__ __forceinline__ double operator+(BFloat16 a, double b) { return static_cast(a) + b; }
+__device__ __forceinline__ double operator-(BFloat16 a, double b) { return static_cast(a) - b; }
+__device__ __forceinline__ double operator*(BFloat16 a, double b) { return static_cast(a) * b; }
+__device__ __forceinline__ double operator/(BFloat16 a, double b) { return static_cast(a) / b; }
+
+__device__ __forceinline__ double operator+(double a, BFloat16 b) { return a + static_cast(b); }
+__device__ __forceinline__ double operator-(double a, BFloat16 b) { return a - static_cast(b); }
+__device__ __forceinline__ double operator*(double a, BFloat16 b) { return a * static_cast(b); }
+__device__ __forceinline__ double operator/(double a, BFloat16 b) { return a / static_cast(b); }
+
+// Overloading < and > operators
+
+__device__ __forceinline__ bool operator==(BFloat16& lhs, BFloat16& rhs) { return float(lhs) == float(rhs); }
+__device__ __forceinline__ bool operator!=(BFloat16& lhs, BFloat16& rhs) { return float(lhs) != float(rhs); }
+__device__ __forceinline__ bool operator>(BFloat16& lhs, BFloat16& rhs) { return float(lhs) > float(rhs); }
+__device__ __forceinline__ bool operator<(BFloat16& lhs, BFloat16& rhs) { return float(lhs) < float(rhs); }
+
template
__device__ __inline__ T _Ceil(T a);
@@ -263,24 +339,19 @@ __device__ __inline__ double _Normcdf(double a) { return normcdf(a); }
template <>
__device__ __inline__ half _Normcdf(half a) { return half(normcdff((float)a)); }
-template
-__device__ __inline__ T _Gelu(T a) {
- return a * _Normcdf(a);
-}
+template <>
+__device__ __inline__ BFloat16 _Sqrt(BFloat16 a) { return sqrtf(static_cast(a)); }
+
+template <>
+__device__ __inline__ BFloat16 _Exp(BFloat16 a) { return expf(static_cast(a)); }
+
+template <>
+__device__ __inline__ BFloat16 _Log(BFloat16 a) { return logf(static_cast(a)); }
+
+template <>
+__device__ __inline__ BFloat16 _Tanh(BFloat16 a) { return tanhf(static_cast(a)); }
#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-template <>
-__device__ __inline__ nv_bfloat16 _Sqrt(nv_bfloat16 a) { return nv_bfloat16(sqrtf(static_cast(a))); }
-
-template <>
-__device__ __inline__ nv_bfloat16 _Exp(nv_bfloat16 a) { return nv_bfloat16(expf(static_cast(a))); }
-
-template <>
-__device__ __inline__ nv_bfloat16 _Log(nv_bfloat16 a) { return nv_bfloat16(logf(static_cast(a))); }
-
-template <>
-__device__ __inline__ nv_bfloat16 _Tanh(nv_bfloat16 a) { return nv_bfloat16(tanhf(static_cast(a))); }
-
template <>
__device__ __inline__ nv_bfloat162 _Tanh(nv_bfloat162 a) {
float2 tmp = (__bfloat1622float2(a));
@@ -288,10 +359,15 @@ __device__ __inline__ nv_bfloat162 _Tanh(nv_bfloat162 a) {
tmp.y = tanhf(tmp.y);
return __float22bfloat162_rn(tmp);
}
+#endif
template <>
-__device__ __inline__ nv_bfloat16 _Normcdf(nv_bfloat16 a) { return nv_bfloat16(normcdff(static_cast(a))); }
-#endif
+__device__ __inline__ BFloat16 _Normcdf(BFloat16 a) { return normcdff(static_cast(a)); }
+
+template
+__device__ __inline__ T _Gelu(T a) {
+ return a * _Normcdf(a);
+}
// We would like to use 64-bit integer to support large matrices. However, CUDA seems to support only 32-bit integer
// For now, use int32_t to ensure that both Linux and Windows see this as 32 bit integer type.
diff --git a/onnxruntime/core/providers/cuda/cuda_common.h b/onnxruntime/core/providers/cuda/cuda_common.h
index 94f56a346c..1d41b47536 100644
--- a/onnxruntime/core/providers/cuda/cuda_common.h
+++ b/onnxruntime/core/providers/cuda/cuda_common.h
@@ -69,18 +69,6 @@ class ToCudaType {
}
};
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
-template <>
-class ToCudaType {
- public:
- typedef nv_bfloat16 MappedType;
- static MappedType FromFloat(float f) {
- uint16_t h = BFloat16(f).val;
- return *reinterpret_cast(&h);
- }
-};
-#endif
-
inline bool CalculateFdmStrides(gsl::span p, const std::vector& dims) {
int stride = 1;
if (dims.empty() || p.size() < dims.size())
diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc
index cd5e7ad6b9..5ebdb8dd1f 100755
--- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc
+++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc
@@ -1116,8 +1116,6 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain,
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, bool, Pad);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, SpaceToDepth);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, DepthToSpace);
-
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 13, BFloat16, Add);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 13, BFloat16, Sub);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 13, BFloat16, Mul);
@@ -1130,7 +1128,6 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain,
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, BFloat16, Tanh);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, BFloat16, Gemm);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, BFloat16, ReduceSum);
-#endif
// OpSet 14
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 14, CumSum);
@@ -1187,14 +1184,11 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain,
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 14, uint8_t, ReduceMin);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 14, int64_t, ReduceMin);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 14, Trilu);
-
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 14, BFloat16, Add);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 14, BFloat16, Sub);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 14, BFloat16, Mul);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 14, BFloat16, Div);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 14, BFloat16, Relu);
-#endif
// OpSet 15
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 15, Pow);
@@ -1966,8 +1960,6 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
-
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
@@ -1980,7 +1972,7 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
-#endif
+
// OpSet 14
BuildKernelCreateInfo,
BuildKernelCreateInfo,
@@ -2035,13 +2027,11 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
-#endif
BuildKernelCreateInfo,
// OpSet 15
diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.h b/onnxruntime/core/providers/cuda/cuda_execution_provider.h
index 67120de11c..7e5326ddc7 100644
--- a/onnxruntime/core/providers/cuda/cuda_execution_provider.h
+++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.h
@@ -143,13 +143,11 @@ class CUDAExecutionProvider : public IExecutionProvider {
constant_ones_half_ = cuda::CreateConstantOnes();
}
return reinterpret_cast(constant_ones_half_->GetBuffer(stream_, count));
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
- } else if (std::is_same::value) {
+ } else if (std::is_same::value) {
if (!constant_ones_bfloat16_) {
- constant_ones_bfloat16_ = cuda::CreateConstantOnes();
+ constant_ones_bfloat16_ = cuda::CreateConstantOnes();
}
return reinterpret_cast(constant_ones_bfloat16_->GetBuffer(stream_, count));
-#endif
} else {
return nullptr;
}
@@ -172,9 +170,7 @@ class CUDAExecutionProvider : public IExecutionProvider {
std::unique_ptr> constant_ones_float_;
std::unique_ptr> constant_ones_double_;
std::unique_ptr> constant_ones_half_;
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
- std::unique_ptr> constant_ones_bfloat16_;
-#endif
+ std::unique_ptr> constant_ones_bfloat16_;
AllocatorPtr allocator_;
};
diff --git a/onnxruntime/core/providers/cuda/cuda_utils.cu b/onnxruntime/core/providers/cuda/cuda_utils.cu
index 2296b127ae..cc9eaf9ae3 100644
--- a/onnxruntime/core/providers/cuda/cuda_utils.cu
+++ b/onnxruntime/core/providers/cuda/cuda_utils.cu
@@ -71,9 +71,7 @@ std::unique_ptr> CreateConstantOnes() {
template std::unique_ptr> CreateConstantOnes();
template std::unique_ptr> CreateConstantOnes();
template std::unique_ptr> CreateConstantOnes();
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-template std::unique_ptr> CreateConstantOnes();
-#endif
+template std::unique_ptr> CreateConstantOnes();
#define SPECIALIZED_FILL(T) \
template void Fill(cudaStream_t stream, T * output, T value, int64_t count);
@@ -85,9 +83,7 @@ SPECIALIZED_FILL(int64_t)
SPECIALIZED_FILL(float)
SPECIALIZED_FILL(double)
SPECIALIZED_FILL(__half)
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-SPECIALIZED_FILL(nv_bfloat16)
-#endif
+SPECIALIZED_FILL(BFloat16)
} // namespace cuda
} // namespace onnxruntime
diff --git a/onnxruntime/core/providers/cuda/cudnn_common.cc b/onnxruntime/core/providers/cuda/cudnn_common.cc
index 52889ff8c9..d85ca2d7cf 100644
--- a/onnxruntime/core/providers/cuda/cudnn_common.cc
+++ b/onnxruntime/core/providers/cuda/cudnn_common.cc
@@ -161,10 +161,8 @@ const float Consts::Zero = 0;
const float Consts::One = 1;
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
-const float Consts::Zero = 0;
-const float Consts::One = 1;
-#endif
+const float Consts::Zero = 0;
+const float Consts::One = 1;
template <>
const int8_t Consts::Zero = 0;
diff --git a/onnxruntime/core/providers/cuda/cudnn_common.h b/onnxruntime/core/providers/cuda/cudnn_common.h
index 6526fc0aae..f8cee8fa89 100644
--- a/onnxruntime/core/providers/cuda/cudnn_common.h
+++ b/onnxruntime/core/providers/cuda/cudnn_common.h
@@ -124,13 +124,11 @@ struct Consts {
static const float One;
};
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
template <>
-struct Consts {
+struct Consts {
static const float Zero;
static const float One;
};
-#endif
inline double ClampCudnnBatchNormEpsilon(double epsilon) {
if (epsilon < CUDNN_BN_MIN_EPSILON) {
diff --git a/onnxruntime/core/providers/cuda/fpgeneric.cu b/onnxruntime/core/providers/cuda/fpgeneric.cu
index 695c6038c0..3ae0b9c8eb 100644
--- a/onnxruntime/core/providers/cuda/fpgeneric.cu
+++ b/onnxruntime/core/providers/cuda/fpgeneric.cu
@@ -55,13 +55,12 @@ __global__ void CopyVectorHalf(const half* x, int incx, half* y, int incy, int n
y[id * incy] = x[id * incx];
}
-#if CUDA_VERSION >= 11000
-__global__ void CopyVectorBFloat16(const nv_bfloat16* x, int incx, nv_bfloat16* y, int incy, int n) {
+__global__ void CopyVectorBFloat16(const onnxruntime::BFloat16* x, int incx, onnxruntime::BFloat16* y, int incy,
+ int n) {
int id = blockIdx.x * blockDim.x + threadIdx.x;
if (id >= n) return;
y[id * incy] = x[id * incx];
}
-#endif
} // namespace
@@ -84,13 +83,10 @@ cublasStatus_t cublasCopyHelper(cudaStream_t stream, cublasHandle_t, int n, cons
return CUBLAS_STATUS_SUCCESS;
}
-#if CUDA_VERSION >= 11000
-cublasStatus_t cublasCopyHelper(cudaStream_t stream, cublasHandle_t, int n, const nv_bfloat16* x, int incx, nv_bfloat16* y, int incy) {
+cublasStatus_t cublasCopyHelper(cudaStream_t stream, cublasHandle_t, int n, const onnxruntime::BFloat16* x, int incx,
+ onnxruntime::BFloat16* y, int incy) {
dim3 dimGrid((unsigned int)(n + COPY_BLOCK_DIM - 1) / COPY_BLOCK_DIM, 1, 1);
dim3 dimBlock(COPY_BLOCK_DIM, 1, 1);
CopyVectorBFloat16<<>>(x, incx, y, incy, n);
return CUBLAS_STATUS_SUCCESS;
}
-
-
-#endif
diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc
index fd669a72a5..e98e639d67 100644
--- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc
+++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc
@@ -189,22 +189,6 @@ Status BinaryElementwise::Prepare(OpKernelContext* context, Bin
// D: double
// O: bool
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
-#define BINARY_OP_TYPED_BF16(name, ver) BINARY_OP_TYPED(name, ver, BFloat16)
-#define BINARY_OP_VERSIONED_TYPED_BF16(name, startver, endver) BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED(name, startver, endver, BFloat16)
-#define BINARY_ELEMENTWISE_REGISTER_KERNEL_TYPED_BF16(name, ver) BINARY_ELEMENTWISE_REGISTER_KERNEL_TYPED(name, ver, BFloat16)
-#define BINARY_ELEMENTWISE_LOGICALOP_REGISTER_KERNEL_TYPED_BF16(name, ver) BINARY_ELEMENTWISE_LOGICALOP_REGISTER_KERNEL_TYPED(name, ver, BFloat16)
-#define BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED_BF16(name, startver, endver) BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED(name, startver, endver, BFloat16)
-#define BINARY_OP_TYPED_VERSIONED_V_BF16(name, class_name, startver, endver) BINARY_OP_TYPED_VERSIONED_V(name, class_name, startver, endver, BFloat16)
-#else
-#define BINARY_OP_TYPED_BF16(name, ver)
-#define BINARY_OP_VERSIONED_TYPED_BF16(name, startver, endver)
-#define BINARY_ELEMENTWISE_REGISTER_KERNEL_TYPED_BF16(name, ver)
-#define BINARY_ELEMENTWISE_LOGICALOP_REGISTER_KERNEL_TYPED_BF16(name, ver)
-#define BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED_BF16(name, startver, endver)
-#define BINARY_OP_TYPED_VERSIONED_V_BF16(name, class_name, startver, endver)
-#endif
-
#define BINARY_OP_VERSIONED_HFD(name, startver, endver) \
BINARY_OP_VERSIONED_TYPED(name, startver, endver, MLFloat16) \
BINARY_OP_VERSIONED_TYPED(name, startver, endver, float) \
@@ -222,14 +206,14 @@ Status BinaryElementwise::Prepare(OpKernelContext* context, Bin
BINARY_OP_VERSIONED_TYPED(name, startver, endver, uint64_t) \
BINARY_OP_VERSIONED_TYPED(name, startver, endver, int32_t) \
BINARY_OP_VERSIONED_TYPED(name, startver, endver, int64_t) \
- BINARY_OP_VERSIONED_TYPED_BF16(name, startver, endver) \
- BINARY_OP_VERSIONED_HFD(name, startver, endver)
+ BINARY_OP_VERSIONED_HFD(name, startver, endver) \
+ BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED(name, startver, endver, BFloat16)
#define BINARY_OP_HFD(name, ver) \
BINARY_OP_TYPED(name, ver, MLFloat16) \
- BINARY_OP_TYPED_BF16(name, ver) \
BINARY_OP_TYPED(name, ver, float) \
- BINARY_OP_TYPED(name, ver, double)
+ BINARY_OP_TYPED(name, ver, double) \
+ BINARY_OP_TYPED(name, ver, BFloat16)
#define BINARY_OP_UZILHFD(name, ver) \
BINARY_OP_TYPED(name, ver, uint32_t) \
@@ -250,9 +234,9 @@ Status BinaryElementwise::Prepare(OpKernelContext* context, Bin
#define BINARY_OP_REGISTER_HFD(name, ver) \
BINARY_ELEMENTWISE_REGISTER_KERNEL_TYPED(name, ver, MLFloat16) \
- BINARY_ELEMENTWISE_REGISTER_KERNEL_TYPED_BF16(name, ver) \
BINARY_ELEMENTWISE_REGISTER_KERNEL_TYPED(name, ver, float) \
- BINARY_ELEMENTWISE_REGISTER_KERNEL_TYPED(name, ver, double)
+ BINARY_ELEMENTWISE_REGISTER_KERNEL_TYPED(name, ver, double) \
+ BINARY_ELEMENTWISE_REGISTER_KERNEL_TYPED(name, ver, BFloat16)
#define BINARY_OP_REGISTER_UZILHFD(name, ver) \
BINARY_ELEMENTWISE_REGISTER_KERNEL_TYPED(name, ver, uint32_t) \
@@ -267,21 +251,21 @@ Status BinaryElementwise::Prepare(OpKernelContext* context, Bin
BINARY_ELEMENTWISE_LOGICALOP_REGISTER_KERNEL_TYPED(name, ver, int32_t) \
BINARY_ELEMENTWISE_LOGICALOP_REGISTER_KERNEL_TYPED(name, ver, int64_t) \
BINARY_ELEMENTWISE_LOGICALOP_REGISTER_KERNEL_TYPED(name, ver, MLFloat16) \
- BINARY_ELEMENTWISE_LOGICALOP_REGISTER_KERNEL_TYPED_BF16(name, ver) \
BINARY_ELEMENTWISE_LOGICALOP_REGISTER_KERNEL_TYPED(name, ver, float) \
- BINARY_ELEMENTWISE_LOGICALOP_REGISTER_KERNEL_TYPED(name, ver, double)
+ BINARY_ELEMENTWISE_LOGICALOP_REGISTER_KERNEL_TYPED(name, ver, double) \
+ BINARY_ELEMENTWISE_LOGICALOP_REGISTER_KERNEL_TYPED(name, ver, BFloat16)
#define BINARY_OP_REGISTER_VERSIONED_HFD(name, startver, endver) \
BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED(name, startver, endver, MLFloat16) \
- BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED_BF16(name, startver, endver) \
BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED(name, startver, endver, float) \
- BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED(name, startver, endver, double)
+ BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED(name, startver, endver, double) \
+ BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED(name, startver, endver, BFloat16)
#define BINARY_OP_REGISTER_VERSIONED_CLASS_HFD(name, class_name, startver, endver) \
BINARY_OP_TYPED_VERSIONED_V(name, class_name, startver, endver, MLFloat16) \
- BINARY_OP_TYPED_VERSIONED_V_BF16(name, class_name, startver, endver) \
BINARY_OP_TYPED_VERSIONED_V(name, class_name, startver, endver, float) \
- BINARY_OP_TYPED_VERSIONED_V(name, class_name, startver, endver, double)
+ BINARY_OP_TYPED_VERSIONED_V(name, class_name, startver, endver, double) \
+ BINARY_OP_TYPED_VERSIONED_V(name, class_name, startver, endver, BFloat16)
#define BINARY_OP_REGISTER_VERSIONED_UZILHFD(name, startver, endver) \
BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED(name, startver, endver, uint32_t) \
diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu
index 39ae0acfd8..e408cbbc04 100644
--- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu
+++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu
@@ -79,23 +79,15 @@ namespace cuda {
const TArray* rhs_padded_strides, const T2* rhs_data, \
const TArray* fdm_output_strides, const fast_divmod& fdm_H, const fast_divmod& fdm_C, T* output_data, size_t count);
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-#define SPECIALIZED_BINARY_ELEMENTWISE_IMPL_BF16(x) SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, nv_bfloat16)
-#define SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T2_BF16(name) SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T2(name, bool, nv_bfloat16, nv_bfloat16)
-#else
-#define SPECIALIZED_BINARY_ELEMENTWISE_IMPL_BF16(x)
-#define SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T2_BF16(name)
-#endif
-
#define SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(x) \
SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, uint32_t) \
SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, uint64_t) \
SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, int32_t) \
SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, int64_t) \
SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, half) \
- SPECIALIZED_BINARY_ELEMENTWISE_IMPL_BF16(x) \
SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, float) \
- SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, double)
+ SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, double) \
+ SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, BFloat16)
#define SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1_ILHFD(x, T) \
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(x, T, int32_t) \
@@ -111,9 +103,9 @@ namespace cuda {
#define SPECIALIZED_BINARY_ELEMENTWISE_IMPL_HFD(x) \
SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, half) \
- SPECIALIZED_BINARY_ELEMENTWISE_IMPL_BF16(x) \
SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, float) \
- SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, double)
+ SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, double) \
+ SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, BFloat16)
// create declarations for impl
#define BINARY_OP_NAME_EXPR(name, expr) \
@@ -172,9 +164,9 @@ BINARY_OPS2()
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T2(name, bool, int32_t, int32_t) \
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T2(name, bool, int64_t, int64_t) \
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T2(name, bool, half, half) \
- SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T2_BF16(name) \
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T2(name, bool, float, float) \
- SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T2(name, bool, double, double)
+ SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T2(name, bool, double, double) \
+ SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T2(name, bool, BFloat16, BFloat16)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD2(Greater)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD2(Equal)
diff --git a/onnxruntime/core/providers/cuda/math/gemm.cc b/onnxruntime/core/providers/cuda/math/gemm.cc
index def63e4db2..5e61bf4e08 100644
--- a/onnxruntime/core/providers/cuda/math/gemm.cc
+++ b/onnxruntime/core/providers/cuda/math/gemm.cc
@@ -53,9 +53,7 @@ namespace cuda {
REGISTER_KERNEL_TYPED(float)
REGISTER_KERNEL_TYPED(double)
REGISTER_KERNEL_TYPED(MLFloat16)
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
REGISTER_KERNEL_TYPED(BFloat16)
-#endif
template
Status Gemm::ComputeInternal(OpKernelContext* ctx) const {
diff --git a/onnxruntime/core/providers/cuda/math/matmul.cc b/onnxruntime/core/providers/cuda/math/matmul.cc
index c7632dcdc5..abdef98751 100644
--- a/onnxruntime/core/providers/cuda/math/matmul.cc
+++ b/onnxruntime/core/providers/cuda/math/matmul.cc
@@ -41,9 +41,7 @@ namespace cuda {
REGISTER_KERNEL_TYPED(float)
REGISTER_KERNEL_TYPED(double)
REGISTER_KERNEL_TYPED(MLFloat16)
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
REGISTER_KERNEL_TYPED(BFloat16)
-#endif
// StridedBatchedGemm can be used for the following GEMM computation
// C[pnm] = A[pnk]*B[km] or C[pnm] = A[pnk]*B[pkm]
diff --git a/onnxruntime/core/providers/cuda/math/softmax.cc b/onnxruntime/core/providers/cuda/math/softmax.cc
index bddf60947d..9116178dda 100644
--- a/onnxruntime/core/providers/cuda/math/softmax.cc
+++ b/onnxruntime/core/providers/cuda/math/softmax.cc
@@ -44,29 +44,23 @@ SPECIALIZED_SOFTMAX_HELPER_IMPL(float)
SPECIALIZED_SOFTMAX_HELPER_IMPL(double)
SPECIALIZED_SOFTMAX_HELPER_IMPL(MLFloat16)
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
// cudnnSoftmaxForward/Backward doesn't support BFloat16.
-#define SPECIALIZED_SOFTMAX_HELPER_IMPL_BFloat16(is_log_softmax) \
- template <> \
- Status SoftMaxComputeHelper( \
- cudaStream_t stream, \
- const BFloat16* X, \
- const TensorShape& input_shape, \
- BFloat16* Y, \
- int64_t axis) { \
- typedef typename ToCudaType::MappedType CudaT; \
- int64_t N = input_shape.SizeToDimension(axis); \
- int64_t D = input_shape.SizeFromDimension(axis); \
- auto Y_data = reinterpret_cast(Y); \
- auto X_data = reinterpret_cast(X); \
- dispatch_warpwise_softmax_forward, is_log_softmax>( \
- stream, Y_data, X_data, gsl::narrow_cast(D), gsl::narrow_cast(D), gsl::narrow_cast(N)); \
- return Status::OK(); \
+#define SPECIALIZED_SOFTMAX_HELPER_IMPL_BFloat16(is_log_softmax) \
+ template <> \
+ Status SoftMaxComputeHelper(cudaStream_t stream, const BFloat16* X, \
+ const TensorShape& input_shape, BFloat16* Y, int64_t axis) { \
+ typedef typename ToCudaType::MappedType CudaT; \
+ int64_t N = input_shape.SizeToDimension(axis); \
+ int64_t D = input_shape.SizeFromDimension(axis); \
+ auto Y_data = reinterpret_cast(Y); \
+ auto X_data = reinterpret_cast(X); \
+ dispatch_warpwise_softmax_forward, is_log_softmax>( \
+ stream, Y_data, X_data, gsl::narrow_cast(D), gsl::narrow_cast(D), gsl::narrow_cast(N)); \
+ return Status::OK(); \
}
SPECIALIZED_SOFTMAX_HELPER_IMPL_BFloat16(true)
- SPECIALIZED_SOFTMAX_HELPER_IMPL_BFloat16(false)
-#endif
+SPECIALIZED_SOFTMAX_HELPER_IMPL_BFloat16(false)
#define REGISTER_KERNEL_TYPED(T) \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
@@ -225,9 +219,7 @@ SPECIALIZED_SOFTMAX_HELPER_IMPL_BFloat16(true)
SPECIALIZED_COMPUTE(float)
SPECIALIZED_COMPUTE(double)
SPECIALIZED_COMPUTE(MLFloat16)
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
SPECIALIZED_COMPUTE(BFloat16)
-#endif
} // namespace cuda
} // namespace onnxruntime
diff --git a/onnxruntime/core/providers/cuda/math/softmax_impl.cu b/onnxruntime/core/providers/cuda/math/softmax_impl.cu
index 8a3396a448..4ef7f6674b 100644
--- a/onnxruntime/core/providers/cuda/math/softmax_impl.cu
+++ b/onnxruntime/core/providers/cuda/math/softmax_impl.cu
@@ -109,9 +109,7 @@ template void dispatch_warpwise_softmax_forward(
SPECIALIZED_WRAPWISE_SOFTMAX_IMPL(float, float, float)
SPECIALIZED_WRAPWISE_SOFTMAX_IMPL(half, half, float)
SPECIALIZED_WRAPWISE_SOFTMAX_IMPL(double, double, double)
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-SPECIALIZED_WRAPWISE_SOFTMAX_IMPL(nv_bfloat16, nv_bfloat16, float)
-#endif
+SPECIALIZED_WRAPWISE_SOFTMAX_IMPL(BFloat16, BFloat16, float)
template
void dispatch_blockwise_softmax_forward(cudaStream_t stream, output_t* output, const input_t* input, int softmax_elements, int softmax_elements_stride, int batch_count) {
@@ -127,17 +125,17 @@ void dispatch_blockwise_softmax_forward(cudaStream_t stream, output_t* output, c
}
}
-#define SPECIALIZED_BLOCKWISE_SOFTMAX_IMPL(input_t, output_t, acc_t) \
-template void dispatch_blockwise_softmax_forward(cudaStream_t stream, output_t* output, const input_t* src, int softmax_elements, int softmax_elements_stride, int batch_count); \
-template void dispatch_blockwise_softmax_forward(cudaStream_t stream, output_t* output, const input_t* src, int softmax_elements, int softmax_elements_stride, int batch_count);
+#define SPECIALIZED_BLOCKWISE_SOFTMAX_IMPL(input_t, output_t, acc_t) \
+ template void dispatch_blockwise_softmax_forward( \
+ cudaStream_t stream, output_t * output, const input_t* src, int softmax_elements, int softmax_elements_stride, \
+ int batch_count); \
+ template void dispatch_blockwise_softmax_forward( \
+ cudaStream_t stream, output_t * output, const input_t* src, int softmax_elements, int softmax_elements_stride, \
+ int batch_count);
SPECIALIZED_BLOCKWISE_SOFTMAX_IMPL(float, float, float)
SPECIALIZED_BLOCKWISE_SOFTMAX_IMPL(half, half, float)
SPECIALIZED_BLOCKWISE_SOFTMAX_IMPL(double, double, double)
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-SPECIALIZED_BLOCKWISE_SOFTMAX_IMPL(nv_bfloat16, nv_bfloat16, float)
-#endif
-
-
+SPECIALIZED_BLOCKWISE_SOFTMAX_IMPL(BFloat16, BFloat16, float)
}
}
diff --git a/onnxruntime/core/providers/cuda/math/unary_elementwise_ops_impl.cu b/onnxruntime/core/providers/cuda/math/unary_elementwise_ops_impl.cu
index 66d50dafa5..9bb5a2d855 100644
--- a/onnxruntime/core/providers/cuda/math/unary_elementwise_ops_impl.cu
+++ b/onnxruntime/core/providers/cuda/math/unary_elementwise_ops_impl.cu
@@ -94,24 +94,18 @@ struct ViaTypeMap {
typedef float ViaT;
};
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
template <>
-struct ViaTypeMap {
+struct ViaTypeMap {
typedef float ViaT;
};
-#endif
template
struct OP_Cast {
__device__ __inline__ OutT operator()(const InT& a) const {
const bool any_float16 = std::is_same::value || std::is_same::value;
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
- const bool any_bf16 = std::is_same::value || std::is_same::value;
- typedef typename std::conditional::type T1;
+ const bool any_bf16 = std::is_same::value || std::is_same::value;
+ typedef typename std::conditional::type T1;
typedef typename std::conditional::type T;
-#else
- typedef typename std::conditional::type T;
-#endif
typedef typename ViaTypeMap::ViaT ViaT;
return (OutT)((ViaT)a);
}
@@ -133,15 +127,8 @@ void Impl_Cast(
#define SPECIALIZED_CAST_IMPL2(InT, OutT) \
template void Impl_Cast(cudaStream_t stream, const InT* input_data, OutT* output_data, size_t count);
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-#define SPECIALIZED_CAST_IMPL2_BF16(T) SPECIALIZED_CAST_IMPL2(T, nv_bfloat16)
-#else
-#define SPECIALIZED_CAST_IMPL2_BF16(T)
-#endif
-
#define SPECIALIZED_CAST_FROM(T) \
SPECIALIZED_CAST_IMPL2(T, half) \
- SPECIALIZED_CAST_IMPL2_BF16(T) \
SPECIALIZED_CAST_IMPL2(T, float) \
SPECIALIZED_CAST_IMPL2(T, double) \
SPECIALIZED_CAST_IMPL2(T, int8_t) \
@@ -152,7 +139,8 @@ void Impl_Cast(
SPECIALIZED_CAST_IMPL2(T, uint16_t) \
SPECIALIZED_CAST_IMPL2(T, uint32_t) \
SPECIALIZED_CAST_IMPL2(T, uint64_t) \
- SPECIALIZED_CAST_IMPL2(T, bool)
+ SPECIALIZED_CAST_IMPL2(T, bool) \
+ SPECIALIZED_CAST_IMPL2(T, BFloat16)
SPECIALIZED_CAST_FROM(half)
SPECIALIZED_CAST_FROM(float)
@@ -166,9 +154,7 @@ SPECIALIZED_CAST_FROM(uint16_t)
SPECIALIZED_CAST_FROM(uint32_t)
SPECIALIZED_CAST_FROM(uint64_t)
SPECIALIZED_CAST_FROM(bool)
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-SPECIALIZED_CAST_FROM(nv_bfloat16)
-#endif
+SPECIALIZED_CAST_FROM(BFloat16)
} // namespace cuda
} // namespace onnxruntime
diff --git a/onnxruntime/core/providers/cuda/math/variadic_elementwise_ops.cc b/onnxruntime/core/providers/cuda/math/variadic_elementwise_ops.cc
index dc4db5a830..81d9e15ce3 100644
--- a/onnxruntime/core/providers/cuda/math/variadic_elementwise_ops.cc
+++ b/onnxruntime/core/providers/cuda/math/variadic_elementwise_ops.cc
@@ -228,26 +228,18 @@ Status VariadicElementwiseOp
namespace {
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
-#define ALL_IEEE_FLOAT_DATA_TYPES MLFloat16, float, double, BFloat16
-#else
-#define ALL_IEEE_FLOAT_DATA_TYPES MLFloat16, float, double
-#endif
+using SumOp = VariadicElementwiseOp;
-using SumOp = VariadicElementwiseOp<
- variadic_elementwise_ops::Sum,
- ALL_IEEE_FLOAT_DATA_TYPES>;
+using MinOp = VariadicElementwiseOp;
-using MinOp = VariadicElementwiseOp<
- variadic_elementwise_ops::Min,
- uint32_t, uint64_t, int32_t, int64_t, ALL_IEEE_FLOAT_DATA_TYPES>;
+using MaxOp = VariadicElementwiseOp;
-using MaxOp = VariadicElementwiseOp<
- variadic_elementwise_ops::Max,
- uint32_t, uint64_t, int32_t, int64_t, ALL_IEEE_FLOAT_DATA_TYPES>;
-
-const DeleteOnUnloadPtr> k_uzilhfd_datatypes = new std::vector(BuildKernelDefConstraints());
-const DeleteOnUnloadPtr> k_hfd_datatypes = new std::vector(BuildKernelDefConstraints());
+const DeleteOnUnloadPtr> k_uzilhfd_datatypes = new std::vector(
+ BuildKernelDefConstraints());
+const DeleteOnUnloadPtr> k_hfd_datatypes =
+ new std::vector(BuildKernelDefConstraints());
} // namespace
diff --git a/onnxruntime/core/providers/cuda/math/variadic_elementwise_ops_impl.cu b/onnxruntime/core/providers/cuda/math/variadic_elementwise_ops_impl.cu
index da1f228e5b..a17dc0df13 100644
--- a/onnxruntime/core/providers/cuda/math/variadic_elementwise_ops_impl.cu
+++ b/onnxruntime/core/providers/cuda/math/variadic_elementwise_ops_impl.cu
@@ -129,17 +129,11 @@ void Impl_NoBroadcastInputBatch(
// D: double
// O: bool
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-#define SPECIALIZE_IMPL_BF16(VariadicElementwiseOpTag) SPECIALIZE_IMPL(nv_bfloat16, VariadicElementwiseOpTag)
-#else
-#define SPECIALIZE_IMPL_BF16(VariadicElementwiseOpTag)
-#endif
-
#define SPECIALIZE_IMPL_HFD(VariadicElementwiseOpTag) \
SPECIALIZE_IMPL(half, VariadicElementwiseOpTag) \
- SPECIALIZE_IMPL_BF16(VariadicElementwiseOpTag) \
SPECIALIZE_IMPL(float, VariadicElementwiseOpTag) \
- SPECIALIZE_IMPL(double, VariadicElementwiseOpTag)
+ SPECIALIZE_IMPL(double, VariadicElementwiseOpTag) \
+ SPECIALIZE_IMPL(BFloat16, VariadicElementwiseOpTag)
#define SPECIALIZE_IMPL_UZILHFD(VariadicElementwiseOpTag) \
SPECIALIZE_IMPL(uint32_t, VariadicElementwiseOpTag) \
diff --git a/onnxruntime/core/providers/cuda/nn/dropout.cc b/onnxruntime/core/providers/cuda/nn/dropout.cc
index 91d9560582..40c4826b56 100644
--- a/onnxruntime/core/providers/cuda/nn/dropout.cc
+++ b/onnxruntime/core/providers/cuda/nn/dropout.cc
@@ -25,8 +25,8 @@ ONNX_OPERATOR_KERNEL_EX(
13,
kCudaExecutionProvider,
(*KernelDefBuilder::Create())
- .TypeConstraint("T", ALL_IEEE_FLOAT_TENSOR_TYPES)
- .TypeConstraint("T1", ALL_IEEE_FLOAT_TENSOR_TYPES)
+ .TypeConstraint("T", BuildKernelDefConstraints())
+ .TypeConstraint("T1", BuildKernelDefConstraints())
.TypeConstraint("T2", DataTypeImpl::GetTensorType())
.InputMemoryType(OrtMemTypeCPUInput, 1)
.InputMemoryType(OrtMemTypeCPUInput, 2),
@@ -50,7 +50,7 @@ Status Dropout::ComputeInternal(OpKernelContext* context) const {
float ratio_data = default_ratio_;
auto ratio = context->Input(1);
if (ratio) {
- utils::MLTypeCallDispatcher t_disp(ratio->GetElementType());
+ utils::MLTypeCallDispatcher t_disp(ratio->GetElementType());
t_disp.Invoke(ratio, ratio_data);
}
@@ -80,7 +80,7 @@ Status Dropout::ComputeInternal(OpKernelContext* context) const {
PhiloxGenerator& generator = generator_ ? *generator_ : PhiloxGenerator::Default();
- utils::MLTypeCallDispatcher t_disp(X->GetElementType());
+ utils::MLTypeCallDispatcher t_disp(X->GetElementType());
t_disp.Invoke(GetDeviceProp(), Stream(), N, ratio_data, generator, *X, *Y, mask_data);
return Status::OK();
diff --git a/onnxruntime/core/providers/cuda/nn/dropout.h b/onnxruntime/core/providers/cuda/nn/dropout.h
index af04a36b40..78ef4e3237 100644
--- a/onnxruntime/core/providers/cuda/nn/dropout.h
+++ b/onnxruntime/core/providers/cuda/nn/dropout.h
@@ -10,18 +10,6 @@
namespace onnxruntime {
namespace cuda {
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
-#define ALL_IEEE_FLOAT_TENSOR_TYPES \
- { DataTypeImpl::GetTensorType(), \
- DataTypeImpl::GetTensorType(), \
- DataTypeImpl::GetTensorType(), \
- DataTypeImpl::GetTensorType() }
-#define ALL_IEEE_FLOAT_DATA_TYPES float, MLFloat16, double, BFloat16
-#else
-#define ALL_IEEE_FLOAT_TENSOR_TYPES DataTypeImpl::AllIEEEFloatTensorTypes()
-#define ALL_IEEE_FLOAT_DATA_TYPES float, MLFloat16, double
-#endif
-
template
struct GetRatioDataImpl {
void operator()(const Tensor* ratio, float& ratio_data) const {
diff --git a/onnxruntime/core/providers/cuda/nn/dropout_impl.cu b/onnxruntime/core/providers/cuda/nn/dropout_impl.cu
index 27e1145431..a7418a17d0 100644
--- a/onnxruntime/core/providers/cuda/nn/dropout_impl.cu
+++ b/onnxruntime/core/providers/cuda/nn/dropout_impl.cu
@@ -158,9 +158,7 @@ void DropoutKernelImpl(
SPECIALIZED_DROPOUT_IMPL(float)
SPECIALIZED_DROPOUT_IMPL(double)
SPECIALIZED_DROPOUT_IMPL(half)
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-SPECIALIZED_DROPOUT_IMPL(nv_bfloat16)
-#endif
+SPECIALIZED_DROPOUT_IMPL(BFloat16)
} // namespace cuda
} // namespace onnxruntime
diff --git a/onnxruntime/core/providers/cuda/reduction/reduction_functions.cu b/onnxruntime/core/providers/cuda/reduction/reduction_functions.cu
index e48ad4f4f5..9488454e5f 100644
--- a/onnxruntime/core/providers/cuda/reduction/reduction_functions.cu
+++ b/onnxruntime/core/providers/cuda/reduction/reduction_functions.cu
@@ -348,6 +348,8 @@ INSTANTIATE_REDUCE_SUM(half, half);
INSTANTIATE_REDUCE_SUM(half, float);
INSTANTIATE_REDUCE_SUM(float, float);
INSTANTIATE_REDUCE_SUM(double, double);
+INSTANTIATE_REDUCE_SUM(BFloat16, BFloat16);
+INSTANTIATE_REDUCE_SUM(BFloat16, float);
#undef INSTANTIATE_REDUCE_SUM
#define INSTANTIATE_REDUCE_SQUARE_SUM(TIn, TOut) \
@@ -355,9 +357,7 @@ INSTANTIATE_REDUCE_SUM(double, double);
INSTANTIATE_REDUCE_SQUARE_SUM(half, float);
INSTANTIATE_REDUCE_SQUARE_SUM(float, float);
INSTANTIATE_REDUCE_SQUARE_SUM(double, double);
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-INSTANTIATE_REDUCE_SQUARE_SUM(nv_bfloat16, float);
-#endif
+INSTANTIATE_REDUCE_SQUARE_SUM(BFloat16, float);
#undef INSTANTIATE_REDUCE_SQUARE_SUM
#define INSTANTIATE_REDUCE_L2_NORM(TIn, TOut) \
@@ -480,9 +480,7 @@ void UnaryDiv(cudaStream_t stream, const T* input, T* output, T denominator, siz
INSTANTIATE_UNARY_DIV(half);
INSTANTIATE_UNARY_DIV(float);
INSTANTIATE_UNARY_DIV(double);
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-INSTANTIATE_UNARY_DIV(nv_bfloat16);
-#endif
+INSTANTIATE_UNARY_DIV(BFloat16);
#undef INSTANTIATE_UNARY_DIV
template
@@ -496,9 +494,7 @@ Status reduce_matrix_rows(cudaStream_t stream, const TIn* input, TOut* output, i
INSTANTIATE_REDUCE_MATRIX_ROWS(half);
INSTANTIATE_REDUCE_MATRIX_ROWS(float);
INSTANTIATE_REDUCE_MATRIX_ROWS(double);
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-INSTANTIATE_REDUCE_MATRIX_ROWS(nv_bfloat16);
-#endif
+INSTANTIATE_REDUCE_MATRIX_ROWS(BFloat16);
#undef INSTANTIATE_REDUCE_MATRIX_ROWS
template
@@ -512,9 +508,7 @@ Status reduce_matrix_columns(cudaStream_t stream, const TIn* input, TOut* output
INSTANTIATE_REDUCE_MATRIX_COLUMNS(half);
INSTANTIATE_REDUCE_MATRIX_COLUMNS(float);
INSTANTIATE_REDUCE_MATRIX_COLUMNS(double);
-#if CUDA_VERSION >= 11000 && (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-INSTANTIATE_REDUCE_MATRIX_COLUMNS(nv_bfloat16);
-#endif
+INSTANTIATE_REDUCE_MATRIX_COLUMNS(BFloat16);
#undef INSTANTIATE_REDUCE_MATRIX_COLUMNS
} // namespace cuda
diff --git a/onnxruntime/core/providers/cuda/reduction/reduction_ops.cc b/onnxruntime/core/providers/cuda/reduction/reduction_ops.cc
index 88e3af5ff6..8ce2c268ae 100644
--- a/onnxruntime/core/providers/cuda/reduction/reduction_ops.cc
+++ b/onnxruntime/core/providers/cuda/reduction/reduction_ops.cc
@@ -835,7 +835,6 @@ SPECIALIZED_REDUCEKERNEL_COMPUTEIMPL(int64_t)
SPECIALIZED_REDUCEKERNEL_COMPUTEIMPL(int8_t)
SPECIALIZED_REDUCEKERNEL_COMPUTEIMPL(uint8_t)
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
template <>
template <>
Status ReduceKernel::ComputeImpl(
@@ -940,7 +939,6 @@ Status ReduceKernel::ComputeImpl
return Status::OK();
}
-#endif
namespace ReductionOps {
@@ -997,17 +995,11 @@ template std::unique_ptr ReduceCompute= 11000
-#define REGISTER_KERNEL_TYPED_BF16(name) REGISTER_KERNEL_TYPED(name, BFloat16)
-#else
-#define REGISTER_KERNEL_TYPED_BF16(name)
-#endif
-
#define REGISTER_KERNEL_HFD(name) \
REGISTER_KERNEL_TYPED(name, MLFloat16) \
- REGISTER_KERNEL_TYPED_BF16(name) \
REGISTER_KERNEL_TYPED(name, float) \
- REGISTER_KERNEL_TYPED(name, double)
+ REGISTER_KERNEL_TYPED(name, double) \
+ REGISTER_KERNEL_TYPED(name, BFloat16)
#define REGISTER_KERNEL_HFD_11(name) \
REGISTER_KERNEL_TYPED_11(name, MLFloat16) \
@@ -1052,9 +1044,7 @@ REGISTER_KERNEL_TYPED_13(ReduceSum, float)
REGISTER_KERNEL_TYPED_13(ReduceSum, double)
REGISTER_KERNEL_TYPED_13(ReduceSum, int32_t)
REGISTER_KERNEL_TYPED_13(ReduceSum, int64_t)
-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
REGISTER_KERNEL_TYPED_13(ReduceSum, BFloat16)
-#endif
REGISTER_KERNEL_HFD(ReduceLogSum)
REGISTER_KERNEL_HFD(ReduceSumSquare)
diff --git a/onnxruntime/core/providers/cuda/shared_inc/accumulation_type.h b/onnxruntime/core/providers/cuda/shared_inc/accumulation_type.h
index d0eafc9d23..3aa7e477e4 100644
--- a/onnxruntime/core/providers/cuda/shared_inc/accumulation_type.h
+++ b/onnxruntime/core/providers/cuda/shared_inc/accumulation_type.h
@@ -4,6 +4,7 @@
#pragma once
#include
+#include "core/framework/float16.h"
namespace onnxruntime {
namespace cuda {
@@ -17,10 +18,8 @@ template <>
struct AccumulationType { using type = float; };
template <>
struct AccumulationType