saymrwulf/onnxruntime
1
0
Fork 0
mirror of https://github.com/saymrwulf/onnxruntime.git synced 2026-07-24 19:43:35 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

Make De/QuantizeLinear support half (#3531)

Browse source 
* Make QuantizeLinear support half

* remove unnessary type constraint

* refine kernel definition

* add fp16 support for dequantizelinear

* diable QuantizeLinear_per_tensor_half_int8 for tensorrt

* refine unit test and fix saturate issue for MSDomain QuantizeLinear

* fix build break

* include tensorrt for half_uint8 test
This commit is contained in:
Yufeng Li 2020-04-17 12:17:48 -07:00 committed by GitHub
parent c7b6fab29d
commit f822a54860
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
11 changed files with 365 additions and 133 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

17
onnxruntime/contrib_ops/cpu/quantize_ops.cc
View file

@ -12,8 +12,8 @@ ONNX_CPU_OPERATOR_TYPED_MS_KERNEL(
1,
uint8_t,
KernelDefBuilder()
.TypeConstraint("T1", DataTypeImpl::GetTensorType<float>())
.TypeConstraint("T2", DataTypeImpl::GetTensorType<uint8_t>()),
.TypeConstraint("T1", DataTypeImpl::GetTensorType<uint8_t>())
.TypeConstraint("T2", DataTypeImpl::GetTensorType<float>()),
DequantizeLinear<uint8_t>);
ONNX_CPU_OPERATOR_TYPED_MS_KERNEL(
@ -21,8 +21,8 @@ ONNX_CPU_OPERATOR_TYPED_MS_KERNEL(
1,
int8_t,
KernelDefBuilder()
.TypeConstraint("T1", DataTypeImpl::GetTensorType<float>())
.TypeConstraint("T2", DataTypeImpl::GetTensorType<int8_t>()),
.TypeConstraint("T1", DataTypeImpl::GetTensorType<int8_t>())
.TypeConstraint("T2", DataTypeImpl::GetTensorType<float>()),
DequantizeLinear<int8_t>);
ONNX_CPU_OPERATOR_TYPED_MS_KERNEL(
@ -34,5 +34,14 @@ ONNX_CPU_OPERATOR_TYPED_MS_KERNEL(
.TypeConstraint("T2", DataTypeImpl::GetTensorType<uint8_t>()),
QuantizeLinear<uint8_t>);
ONNX_CPU_OPERATOR_TYPED_MS_KERNEL(
QuantizeLinear,
1,
int8_t,
KernelDefBuilder()
.TypeConstraint("T1", DataTypeImpl::GetTensorType<float>())
.TypeConstraint("T2", DataTypeImpl::GetTensorType<int8_t>()),
QuantizeLinear<int8_t>);
} // namespace contrib
} // namespace onnxruntime

2
onnxruntime/contrib_ops/cpu_contrib_kernels.cc
View file

@ -30,6 +30,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1,
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, uint8_t, DequantizeLinear);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, int8_t, DequantizeLinear);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, uint8_t, QuantizeLinear);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, int8_t, QuantizeLinear);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, float, CDist);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, double, CDist);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, Gelu);
@ -104,6 +105,7 @@ Status RegisterCpuContribKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, uint8_t, DequantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, int8_t, DequantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, uint8_t, QuantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, int8_t, QuantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, float, CDist)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, double, CDist)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, BiasGelu)>,

45
onnxruntime/contrib_ops/cuda/quantize_ops.cc Normal file
View file

@ -0,0 +1,45 @@
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/tensor/quantize_linear.h"
namespace onnxruntime {
namespace contrib {
namespace cuda {
using namespace onnxruntime::cuda;
#define REGISTER_Q_KERNEL_TYPED(T, U) \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
QuantizeLinear, \
kMSDomain, \
1, \
T##_##U, \
kCudaExecutionProvider, \
KernelDefBuilder() \
.TypeConstraint("T1", DataTypeImpl::GetTensorType<U>()) \
.TypeConstraint("T2", DataTypeImpl::GetTensorType<T>()), \
QuantizeLinear<T, U>);
REGISTER_Q_KERNEL_TYPED(int8_t, MLFloat16)
REGISTER_Q_KERNEL_TYPED(uint8_t, MLFloat16)
#define REGISTER_DQ_KERNEL_TYPED(T, U) \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
DequantizeLinear, \
kMSDomain, \
1, \
T##_##U, \
kCudaExecutionProvider, \
KernelDefBuilder() \
.TypeConstraint("T1", DataTypeImpl::GetTensorType<T>()) \
.TypeConstraint("T2", DataTypeImpl::GetTensorType<U>()), \
DequantizeLinear<T, U>);
REGISTER_DQ_KERNEL_TYPED(int8_t, MLFloat16)
REGISTER_DQ_KERNEL_TYPED(uint8_t, MLFloat16)
} // namespace cuda
} // namespace contrib
} // namespace onnxruntime

10
onnxruntime/contrib_ops/cuda_contrib_kernels.cc
View file

@ -56,6 +56,10 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain,
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float_float, LayerNormalization);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double_float, LayerNormalization);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16_float, LayerNormalization);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, int8_t_MLFloat16, QuantizeLinear);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, uint8_t_MLFloat16, QuantizeLinear);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, int8_t_MLFloat16, DequantizeLinear);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, uint8_t_MLFloat16, DequantizeLinear);
Status RegisterCudaContribKernels(KernelRegistry& kernel_registry) {
static const BuildKernelCreateInfoFn function_table[] = {
@ -105,7 +109,11 @@ Status RegisterCudaContribKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16, ThresholdedRelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float_float, LayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double_float, LayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16_float, LayerNormalization)>};
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16_float, LayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, int8_t_MLFloat16, QuantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, uint8_t_MLFloat16, QuantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, int8_t_MLFloat16, DequantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, uint8_t_MLFloat16, DequantizeLinear)>};
for (auto& function_table_entry : function_table) {
ORT_RETURN_IF_ERROR(kernel_registry.Register(function_table_entry()));

32
onnxruntime/core/graph/contrib_ops/contrib_defs.cc
View file

@ -1248,9 +1248,10 @@ activation and leaky_relu_alpha.)DOC")
ONNX_CONTRIB_OPERATOR_SCHEMA_ELSEWHERE(Range, RegisterRangeOpSchema);
static const char* QuantizeLinear_ver1_doc = R"DOC(
The linear quantization operator. It consumes a full precision data, a scale, a zero point and computes the quantized data.
The quantization formula is y = (x / y_scale) + y_zero_point. For (x / y_scale), it computes the nearest integer value to arg (in floating-point format),
rounding halfway cases away from zero. Scale and zero point must have same shape. They must be either scalar (per tensor) or 1-D tensor (per 'axis').)DOC";
The linear quantization operator. It consumes a full precision data, a scale, a zero point to compute the low precision / quantized tensor.
The quantization formula is y = saturate ((x / y_scale) + y_zero_point).For saturation, it saturates to [0, 255] if it's uint8, or [-128, 127] if it's int8.
For (x / y_scale), it's rounding to nearest ties to even. Refer to https://en.wikipedia.org/wiki/Rounding for details.
Scale and zero point must have same shape. They must be either scalar (per tensor) or 1-D tensor (per 'axis').)DOC";
ONNX_CONTRIB_OPERATOR_SCHEMA(QuantizeLinear)
.SetDomain(kMSDomain)
@ -1287,7 +1288,7 @@ The quantization formula is y = (x / y_scale) + y_zero_point. For (x / y_scale),
"T2")
.TypeConstraint(
"T1",
{"tensor(float)"},
{"tensor(float16)", "tensor(float)"},
"Constrain 'x', 'y_scale' to float tensors.")
.TypeConstraint(
"T2",
@ -1318,35 +1319,36 @@ Scale and zero point must have same shape. They must be either scalar (per tenso
"If it's specified, it means per 'axis' quantization and input 'x_scale' and 'x_zero_point' must be 1-D tensors.",
AttributeProto::INT,
false)
.Input(0,
"x",
"N-D quantized Input tensor to be de-quantized.",
"T2")
.Input(
0,
"x",
"N-D quantized Input tensor to be de-quantized.",
"T1")
.Input(
1,
"x_scale",
"Scale for input 'x'. It could be a scalar or a 1-D tensor, which means a per-tensor or per-axis quantization."
"If it's a 1-D tensor, its number of elements should be equal to the dimension value of 'axis' dimension of input 'x'.",
"T1")
"T2")
.Input(
2,
"x_zero_point",
"Zero point for input 'x'. It could be a scalar or a 1-D tensor, which means a per-tensor or per-axis quantization."
"If it's a 1-D tensor, its number of elements should be equal to the dimension value of 'axis' dimension of input 'x'.",
"T2")
"T1")
.Output(
0,
"y",
"N-D full precision output tensor. It has same shape as input 'x'.",
"T1")
"T2")
.TypeConstraint(
"T1",
{"tensor(float)"},
"Constrain 'y', 'x_scale' to float tensors.")
{"tensor(int8)", "tensor(uint8)"},
"Constrain 'x' and 'x_zero_point' to 8-bit integer tensors.")
.TypeConstraint(
"T2",
{"tensor(int8)", "tensor(uint8)"},
"Constrain 'x_zero_point' and 'x' to 8-bit integer tensors.")
{"tensor(float16)", "tensor(float)"},
"Constrain 'y', 'x_scale' to float tensors.")
.SetDoc(DequantizeLinear_ver1_doc)
.TypeAndShapeInferenceFunction([](ONNX_NAMESPACE::InferenceContext& ctx) {
auto y_type = ctx.getOutputType(0);

24
onnxruntime/core/providers/cpu/tensor/quantize_linear.cc
View file

@ -129,25 +129,11 @@ Status QuantizeLinear<T>::Compute(OpKernelContext* ctx) const {
const float* input = x.template Data<float>();
T* output = y.template MutableData<T>();
// Schema of QuantizeLinearOp changed when it was promoted to onnx domain. In order to maintain backward compatiblity
// both the versions need to be supported.
if (ctx->GetOpDomain() != kMSDomain) {
MlasQuantizeLinear(input, output, static_cast<size_t>(block_size), *scale, *zero_point);
} else {
const float qmax = std::numeric_limits<T>::max();
const float qmin_default = std::numeric_limits<T>::min();
// adjust qmin for int8 inputs. This is required to keep zero point as zero
const float qmin = qmin_default == -128 ? -127 : qmin_default;
for (size_t n = 0; n < static_cast<size_t>(N); n++) {
for (size_t bd = 0; bd < static_cast<size_t>(broadcast_dim); bd++) {
float zp = static_cast<float>(zero_point[bd]);
auto sc = scale[bd];
for (size_t bs = 0; bs < static_cast<size_t>(block_size); bs++) {
*output++ = static_cast<T>(clamp(std::round(static_cast<float>(*input++) / sc) + zp, qmin, qmax));
}
}
for (size_t n = 0; n < static_cast<size_t>(N); n++) {
for (size_t bd = 0; bd < static_cast<size_t>(broadcast_dim); bd++) {
MlasQuantizeLinear(input, output, static_cast<size_t>(block_size), scale[bd], zero_point[bd]);
input += block_size;
output += block_size;
}
}

104
onnxruntime/core/providers/cuda/tensor/quantize_linear.cc
View file

@ -9,28 +9,10 @@
namespace onnxruntime {
namespace cuda {
ONNX_OPERATOR_TYPED_KERNEL_EX(QuantizeLinear,
kOnnxDomain,
10,
uint8_t,
kCudaExecutionProvider,
KernelDefBuilder()
.TypeConstraint("T1", DataTypeImpl::GetTensorType<float>())
.TypeConstraint("T2", DataTypeImpl::GetTensorType<uint8_t>()),
QuantizeLinear<uint8_t>);
template <class T, class U>
Status QuantizeLinear<T, U>::ComputeInternal(OpKernelContext* ctx) const {
typedef typename ToCudaType<U>::MappedType CudaU;
ONNX_OPERATOR_TYPED_KERNEL_EX(QuantizeLinear,
kOnnxDomain,
10,
int8_t,
kCudaExecutionProvider,
KernelDefBuilder()
.TypeConstraint("T1", DataTypeImpl::GetTensorType<float>())
.TypeConstraint("T2", DataTypeImpl::GetTensorType<int8_t>()),
QuantizeLinear<int8_t>);
template <class T>
Status QuantizeLinear<T>::ComputeInternal(OpKernelContext* ctx) const {
auto x = ctx->Input<Tensor>(0);
auto y_scale = ctx->Input<Tensor>(1);
auto y_zero_point = ctx->Input<Tensor>(2);
@ -42,14 +24,15 @@ Status QuantizeLinear<T>::ComputeInternal(OpKernelContext* ctx) const {
const auto& x_shape = x->Shape();
const float* input = x->template Data<float>();
const CudaU* input = reinterpret_cast<const CudaU*>(x->template Data<U>());
T* output = y->template MutableData<T>();
ORT_ENFORCE(IsScalarOr1ElementVector(y_scale), "x_scale must be a scalar or 1D tensor of size 1.");
ORT_ENFORCE(IsScalarOr1ElementVector(y_zero_point), "x_zero_point must be a scalar or 1D tensor of size 1.");
// TO DO: support per-channel
ORT_ENFORCE(IsScalarOr1ElementVector(y_scale), "y_scale must be a scalar or 1D tensor of size 1.");
ORT_ENFORCE(IsScalarOr1ElementVector(y_zero_point), "y_zero_point must be a scalar or 1D tensor of size 1.");
const T* zero_point = y_zero_point->template Data<T>();
const float* scale = y_scale->template Data<float>();
const CudaU* scale = reinterpret_cast<const CudaU*>(y_scale->template Data<U>());
const auto num_of_elements = x_shape.Size();
CudaQuantizeLinear(input, output, scale, zero_point, num_of_elements);
@ -57,26 +40,10 @@ Status QuantizeLinear<T>::ComputeInternal(OpKernelContext* ctx) const {
return Status::OK();
}
ONNX_OPERATOR_TYPED_KERNEL_EX(DequantizeLinear,
kOnnxDomain,
10,
uint8_t,
kCudaExecutionProvider,
KernelDefBuilder()
.TypeConstraint("T", DataTypeImpl::GetTensorType<uint8_t>()),
DequantizeLinear<uint8_t>);
template <class T, class U>
Status DequantizeLinear<T, U>::ComputeInternal(OpKernelContext* ctx) const {
typedef typename ToCudaType<U>::MappedType CudaU;
ONNX_OPERATOR_TYPED_KERNEL_EX(DequantizeLinear,
kOnnxDomain,
10,
int8_t,
kCudaExecutionProvider,
KernelDefBuilder()
.TypeConstraint("T", DataTypeImpl::GetTensorType<int8_t>()),
DequantizeLinear<int8_t>);
template <class T>
Status DequantizeLinear<T>::ComputeInternal(OpKernelContext* ctx) const {
auto x = ctx->Input<Tensor>(0);
auto y_scale = ctx->Input<Tensor>(1);
auto y_zero_point = ctx->Input<Tensor>(2);
@ -90,13 +57,13 @@ Status DequantizeLinear<T>::ComputeInternal(OpKernelContext* ctx) const {
ORT_ENFORCE(y != nullptr);
const T* input = x->template Data<T>();
float* output = y->template MutableData<float>();
CudaU* output = reinterpret_cast<CudaU*>(y->template MutableData<U>());
ORT_ENFORCE(IsScalarOr1ElementVector(y_scale), "x_scale must be a scalar or 1D tensor of size 1.");
ORT_ENFORCE(IsScalarOr1ElementVector(y_zero_point), "x_zero_point must be a scalar or 1D tensor of size 1.");
ORT_ENFORCE(IsScalarOr1ElementVector(y_scale), "y_scale must be a scalar or 1D tensor of size 1.");
ORT_ENFORCE(IsScalarOr1ElementVector(y_zero_point), "y_zero_point must be a scalar or 1D tensor of size 1.");
const T* zero_point = y_zero_point->template Data<T>();
const float* scale = y_scale->template Data<float>();
const CudaU* scale = reinterpret_cast<const CudaU*>(y_scale->template Data<U>());
const auto num_of_elements = x_shape.Size();
CudaDequantizeLinear(input, output, scale, zero_point, num_of_elements);
@ -104,5 +71,46 @@ Status DequantizeLinear<T>::ComputeInternal(OpKernelContext* ctx) const {
return Status::OK();
}
// register QuantizeLinear kernels
#define REGISTER_Q_KERNEL_TYPED(T) \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
QuantizeLinear, \
kOnnxDomain, \
10, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder() \
.TypeConstraint("T1", DataTypeImpl::GetTensorType<float>()) \
.TypeConstraint("T2", DataTypeImpl::GetTensorType<T>()), \
QuantizeLinear<T, float>);
REGISTER_Q_KERNEL_TYPED(int8_t)
REGISTER_Q_KERNEL_TYPED(uint8_t)
// register DequantizeLinear kernels
#define REGISTER_DQ_KERNEL_TYPED(T) \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
DequantizeLinear, \
kOnnxDomain, \
10, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder() \
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
DequantizeLinear<T, float>);
REGISTER_DQ_KERNEL_TYPED(int8_t)
REGISTER_DQ_KERNEL_TYPED(uint8_t)
// specialize QuantizeLinear::ComputeInternal and DequantizeLinear::ComputeInternal
#define SPECIALIZED_QDQ_COMPUTE(T, U) \
template Status QuantizeLinear<T, U>::ComputeInternal(OpKernelContext* ctx) const; \
template Status DequantizeLinear<T, U>::ComputeInternal(OpKernelContext* ctx) const;
SPECIALIZED_QDQ_COMPUTE(int8_t, float)
SPECIALIZED_QDQ_COMPUTE(uint8_t, float)
SPECIALIZED_QDQ_COMPUTE(int8_t, MLFloat16)
SPECIALIZED_QDQ_COMPUTE(uint8_t, MLFloat16)
} // namespace cuda
} // namespace onnxruntime

94
onnxruntime/core/providers/cuda/tensor/quantize_linear.cu
View file

@ -8,6 +8,20 @@
namespace onnxruntime {
namespace cuda {
template <int NumThreadsPerBlock, int NumElementsPerThread>
__global__ void QuantizeLinearKernel(const half* input, int8_t* output, const half* scale, const int8_t* zero_point, CUDA_LONG N) {
CUDA_LONG id = NumElementsPerThread * NumThreadsPerBlock * blockIdx.x + threadIdx.x;
#pragma unroll
for (int i = 0; i < NumElementsPerThread; i++) {
if (id < N) {
int value = __half2int_rn(input[id] / (*scale)) + *zero_point;
output[id] = static_cast<int8_t>(max(-128, min(127, value)));
id += NumThreadsPerBlock;
}
}
}
template <int NumThreadsPerBlock, int NumElementsPerThread>
__global__ void QuantizeLinearKernel(const float* input, int8_t* output, const float* scale, const int8_t* zero_point, CUDA_LONG N) {
CUDA_LONG id = NumElementsPerThread * NumThreadsPerBlock * blockIdx.x + threadIdx.x;
@ -36,56 +50,72 @@ __global__ void QuantizeLinearKernel(const float* input, uint8_t* output, const
}
}
template <class T>
Status CudaQuantizeLinear(const float* input, T* output, const float* scale, const T* zero_point, size_t num_of_element) {
if (num_of_element <= 0)
return Status::OK();
int blocksPerGrid = static_cast<int>(CeilDiv(num_of_element, GridDim::maxThreadsPerBlock * GridDim::maxElementsPerThread));
QuantizeLinearKernel<GridDim::maxThreadsPerBlock, GridDim::maxElementsPerThread>
<<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0>>>(
input,
output,
scale,
zero_point,
num_of_element);
return Status::OK();
}
template <class T, int NumThreadsPerBlock, int NumElementsPerThread>
__global__ void DequantizeLinearKernel(const T* input, float* output, const float* scale, const T* zero_point, CUDA_LONG N) {
template <int NumThreadsPerBlock, int NumElementsPerThread>
__global__ void QuantizeLinearKernel(const half* input, uint8_t* output, const half* scale, const uint8_t* zero_point, CUDA_LONG N) {
CUDA_LONG id = NumElementsPerThread * NumThreadsPerBlock * blockIdx.x + threadIdx.x;
#pragma unroll
for (int i = 0; i < NumElementsPerThread; i++) {
if (id < N) {
output[id] = (input[id] - *zero_point) * (*scale);
int value = __half2int_rn(input[id] / (*scale)) + *zero_point;
output[id] = static_cast<uint8_t>(max(0, min(255, value)));
id += NumThreadsPerBlock;
}
}
}
template <class T>
Status CudaDequantizeLinear(const T* input, float* output, const float* scale, const T* zero_point, size_t num_of_element) {
template <class T, class U>
Status CudaQuantizeLinear(const U* input, T* output, const U* scale, const T* zero_point, size_t num_of_element) {
if (num_of_element <= 0)
return Status::OK();
int blocksPerGrid = static_cast<int>(CeilDiv(num_of_element, GridDim::maxThreadsPerBlock * GridDim::maxElementsPerThread));
DequantizeLinearKernel<T, GridDim::maxThreadsPerBlock, GridDim::maxElementsPerThread>
<<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0>>>(
input,
output,
scale,
zero_point,
num_of_element);
QuantizeLinearKernel<GridDim::maxThreadsPerBlock, GridDim::maxElementsPerThread><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0>>>(
input,
output,
scale,
zero_point,
num_of_element);
return Status::OK();
}
template Status CudaQuantizeLinear<int8_t>(const float* input, int8_t* output, const float* scale, const int8_t* zero_point, size_t num_of_element);
template Status CudaQuantizeLinear<uint8_t>(const float* input, uint8_t* output, const float* scale, const uint8_t* zero_point, size_t num_of_element);
template <class T, class U, int NumThreadsPerBlock, int NumElementsPerThread>
__global__ void DequantizeLinearKernel(const T* input, U* output, const U* scale, const T* zero_point, CUDA_LONG N) {
CUDA_LONG id = NumElementsPerThread * NumThreadsPerBlock * blockIdx.x + threadIdx.x;
template Status CudaDequantizeLinear<int8_t>(const int8_t* input, float* output, const float* scale, const int8_t* zero_point, size_t num_of_element);
template Status CudaDequantizeLinear<uint8_t>(const uint8_t* input, float* output, const float* scale, const uint8_t* zero_point, size_t num_of_element);
#pragma unroll
for (int i = 0; i < NumElementsPerThread; i++) {
if (id < N) {
output[id] = static_cast<U>((input[id] - *zero_point)) * (*scale);
id += NumThreadsPerBlock;
}
}
}
template <class T, class U>
Status CudaDequantizeLinear(const T* input, U* output, const U* scale, const T* zero_point, size_t num_of_element) {
if (num_of_element <= 0)
return Status::OK();
int blocksPerGrid = static_cast<int>(CeilDiv(num_of_element, GridDim::maxThreadsPerBlock * GridDim::maxElementsPerThread));
DequantizeLinearKernel<T, U, GridDim::maxThreadsPerBlock, GridDim::maxElementsPerThread><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0>>>(
input,
output,
scale,
zero_point,
num_of_element);
return Status::OK();
}
template Status CudaQuantizeLinear<int8_t, float>(const float* input, int8_t* output, const float* scale, const int8_t* zero_point, size_t num_of_element);
template Status CudaQuantizeLinear<uint8_t, float>(const float* input, uint8_t* output, const float* scale, const uint8_t* zero_point, size_t num_of_element);
template Status CudaQuantizeLinear<int8_t, half>(const half* input, int8_t* output, const half* scale, const int8_t* zero_point, size_t num_of_element);
template Status CudaQuantizeLinear<uint8_t, half>(const half* input, uint8_t* output, const half* scale, const uint8_t* zero_point, size_t num_of_element);
template Status CudaDequantizeLinear<int8_t, float>(const int8_t* input, float* output, const float* scale, const int8_t* zero_point, size_t num_of_element);
template Status CudaDequantizeLinear<uint8_t, float>(const uint8_t* input, float* output, const float* scale, const uint8_t* zero_point, size_t num_of_element);
template Status CudaDequantizeLinear<int8_t, half>(const int8_t* input, half* output, const half* scale, const int8_t* zero_point, size_t num_of_element);
template Status CudaDequantizeLinear<uint8_t, half>(const uint8_t* input, half* output, const half* scale, const uint8_t* zero_point, size_t num_of_element);
} // namespace cuda
} // namespace onnxruntime

8
onnxruntime/core/providers/cuda/tensor/quantize_linear.cuh
View file

@ -11,11 +11,11 @@
namespace onnxruntime {
namespace cuda {
template <class T>
Status CudaQuantizeLinear(const float* input, T* output, const float* scale, const T* zero_point, size_t num_of_element);
template <class T, class U>
Status CudaQuantizeLinear(const U* input, T* output, const U* scale, const T* zero_point, size_t num_of_element);
template <class T>
Status CudaDequantizeLinear(const T* input, float* output, const float* scale, const T* zero_point, size_t num_of_element);
template <class T, class U>
Status CudaDequantizeLinear(const T* input, U* output, const U* scale, const T* zero_point, size_t num_of_element);
} // namespace cuda
} // namespace onnxruntime

4
onnxruntime/core/providers/cuda/tensor/quantize_linear.h
View file

@ -11,7 +11,7 @@
namespace onnxruntime {
namespace cuda {
template <class T>
template <class T, class U = float>
class QuantizeLinear final : public CudaKernel {
public:
QuantizeLinear(const OpKernelInfo& info) : CudaKernel(info) {}
@ -19,7 +19,7 @@ class QuantizeLinear final : public CudaKernel {
Status ComputeInternal(OpKernelContext* p_op_kernel_context) const override;
};
template <class T>
template <class T, class U = float>
class DequantizeLinear final : public CudaKernel {
public:
DequantizeLinear(const OpKernelInfo& info) : CudaKernel(info) {}

158
onnxruntime/test/contrib_ops/quantize_ops_test.cc
View file

@ -2,11 +2,57 @@
// Licensed under the MIT License.
#include "gtest/gtest.h"
#include "test/common/tensor_op_test_utils.h"
#include "test/providers/provider_test_utils.h"
namespace onnxruntime {
namespace test {
// scalar zero & scale with uint8
TEST(DequantizeLinearOpTest, DequantizeLinear_per_tensor_float_uint8) {
OpTester test("DequantizeLinear", 1, onnxruntime::kMSDomain);
std::vector<int64_t> dims{4};
test.AddInput<uint8_t>("x", dims, {0, 3, 128, 255});
test.AddInput<float>("x_scale", {}, {2.0f});
test.AddInput<uint8_t>("x_zero_point", {}, {128});
test.AddOutput<float>("y", dims, {-256.0f, -250.0f, 0.0f, 254.0f});
test.Run();
}
// scalar zero & scale with int8
TEST(DequantizeLinearOpTest, DequantizeLinear_per_tensor_float_int8) {
OpTester test("DequantizeLinear", 1, onnxruntime::kMSDomain);
std::vector<int64_t> dims{4};
test.AddInput<int8_t>("x", dims, {-30, -3, 100, 127});
test.AddInput<float>("x_scale", {}, {2.0f});
test.AddInput<int8_t>("x_zero_point", {}, {-10});
test.AddOutput<float>("y", dims, {-40.0f, 14.0f, 220.0f, 274.0f});
test.Run();
}
#ifdef USE_CUDA
TEST(DequantizeLinearOpTest, DequantizeLinear_per_tensor_half_uint8) {
OpTester test("DequantizeLinear", 1, onnxruntime::kMSDomain);
std::vector<int64_t> dims{4};
test.AddInput<uint8_t>("x", dims, {0, 3, 128, 255});
test.AddInput<MLFloat16>("x_scale", {}, ToFloat16({2.0f}));
test.AddInput<uint8_t>("x_zero_point", {}, {128});
test.AddOutput<MLFloat16>("y", dims, ToFloat16({-256.0f, -250.0f, 0.0f, 254.0f}));
test.Run();
}
// scalar zero & scale with int8
TEST(DequantizeLinearOpTest, DequantizeLinear_per_tensor_half_int8) {
OpTester test("DequantizeLinear", 1, onnxruntime::kMSDomain);
std::vector<int64_t> dims{4};
test.AddInput<int8_t>("x", dims, {-30, -3, 100, 127});
test.AddInput<MLFloat16>("x_scale", {}, ToFloat16({2.0f}));
test.AddInput<int8_t>("x_zero_point", {}, {-10});
test.AddOutput<MLFloat16>("y", dims, ToFloat16({-40.0f, 14.0f, 220.0f, 274.0f}));
test.Run();
}
#endif
// 1d zero & scale with uint8 broadcast axis 0
TEST(DequantizeLinearContribOpTest, DequantizeLinear_0) {
OpTester test("DequantizeLinear", 1, onnxruntime::kMSDomain);
@ -120,18 +166,114 @@ TEST(DequantizeLinearContribOpTest, DequantizeLinear_3) {
}
// quantize with scalar zero point and scale
TEST(QuantizeLinearContribOpTest, QuantizeLinear_0) {
TEST(QuantizeLinearContribOpTest, QuantizeLinear_per_tensor_float_uint8) {
OpTester test("QuantizeLinear", 1, onnxruntime::kMSDomain);
std::vector<int64_t> dims{6};
test.AddInput<float>("x", dims, {0, 2, 3, 1000, -254, -1000});
std::vector<int64_t> dims{16};
test.AddInput<float>("x", dims, {
0.f, 2.f,
3.f, -3.f, // rounding half to even
2.9f, -2.9f, // low case
3.1f, -3.1f, // up case
254.f, -256.f, // critical point
255.f, -257.f, // critical point
256.f, -258.f, // critical point
1000.f, -1000.f // saturate case
});
test.AddInput<float>("y_scale", {}, {2.0f});
test.AddInput<uint8_t>("y_zero_point", {}, {128});
test.AddOutput<uint8_t>("y", dims, {128, 129, 130, 255, 1, 0});
test.AddOutput<uint8_t>("y", dims, {128, 129,
130, 126,
129, 127,
130, 126,
255, 0,
255, 0,
255, 0,
255, 0});
test.Run();
}
TEST(QuantizeLinearContribOpTest, QuantizeLinear_per_tensor_float_int8) {
OpTester test("QuantizeLinear", 1, onnxruntime::kMSDomain);
std::vector<int64_t> dims{16};
test.AddInput<float>("x", dims, {
0.f, 2.f,
3.f, -3.f, // rounding half to even
2.9f, -2.9f, // low case
3.1f, -3.1f, // up case
254.f, -256.f, // critical point
255.f, -257.f, // critical point
256.f, -258.f, // critical point
1000.f, -1000.f // saturate case
});
test.AddInput<float>("y_scale", {}, {2.0f});
test.AddInput<int8_t>("y_zero_point", {}, {1});
test.AddOutput<int8_t>("y", dims, {1, 2,
3, -1,
2, 0,
3, -1,
127, -127,
127, -127,
127, -128,
127, -128});
test.Run();
}
#ifdef USE_CUDA
TEST(QuantizeLinearContribOpTest, QuantizeLinear_per_tensor_half_uint8) {
OpTester test("QuantizeLinear", 1, onnxruntime::kMSDomain);
std::vector<int64_t> dims{16};
test.AddInput<MLFloat16>("x", dims, ToFloat16({
0.f, 2.f,
3.f, -3.f, // rounding half to even
2.9f, -2.9f, // low case
3.1f, -3.1f, // up case
254.f, -256.f, // critical point
255.f, -257.f, // critical point
256.f, -258.f, // critical point
1000.f, -1000.f // saturate case
}));
test.AddInput<MLFloat16>("y_scale", {}, ToFloat16({2.0f}));
test.AddInput<uint8_t>("y_zero_point", {}, {128});
test.AddOutput<uint8_t>("y", dims, {128, 129,
130, 126,
129, 127,
130, 126,
255, 0,
255, 0,
255, 0,
255, 0});
test.Run();
}
TEST(QuantizeLinearContribOpTest, QuantizeLinear_per_tensor_half_int8) {
OpTester test("QuantizeLinear", 1, onnxruntime::kMSDomain);
std::vector<int64_t> dims{16};
test.AddInput<MLFloat16>("x", dims, ToFloat16({
0.f, 2.f,
3.f, -3.f, // rounding half to even
2.9f, -2.9f, // low case
3.1f, -3.1f, // up case
254.f, -256.f, // critical point
255.f, -257.f, // critical point
256.f, -258.f, // critical point
1000.f, -1000.f // saturate case
}));
test.AddInput<MLFloat16>("y_scale", {}, ToFloat16({2.0f}));
test.AddInput<int8_t>("y_zero_point", {}, {1});
test.AddOutput<int8_t>("y", dims, {1, 2,
3, -1,
2, 0,
3, -1,
127, -127,
127, -127,
127, -128,
127, -128});
test.Run();
}
#endif
// quantize with broadcasting
TEST(QuantizeLinearContribOpTest, QuantizeLinear_1) {
TEST(QuantizeLinearContribOpTest, QuantizeLinear_per_channel) {
OpTester test("QuantizeLinear", 1, onnxruntime::kMSDomain);
std::vector<int64_t> dims{3, 4};
test.AddInput<float>("X", dims,
@ -144,12 +286,12 @@ TEST(QuantizeLinearContribOpTest, QuantizeLinear_1) {
test.AddOutput<uint8_t>("Y", dims,
{0, 2, 3, 255,
0, 1, 2, 255,
0, 1, 1, 250});
0, 0, 1, 250});
test.Run();
}
// quantize with broadcasting and negative axis (-2 resolves to axis 0)
TEST(QuantizeLinearContribOpTest, QuantizeLinear_2) {
TEST(QuantizeLinearContribOpTest, QuantizeLinear_per_channel_negative_axis) {
OpTester test("QuantizeLinear", 1, onnxruntime::kMSDomain);
std::vector<int64_t> dims{3, 4};
test.AddInput<float>("X", dims,
@ -162,7 +304,7 @@ TEST(QuantizeLinearContribOpTest, QuantizeLinear_2) {
test.AddOutput<uint8_t>("Y", dims,
{0, 2, 3, 255,
0, 1, 2, 255,
0, 1, 1, 250});
0, 0, 1, 250});
test.Run();
}
} // namespace test