diff --git a/onnxruntime/contrib_ops/cpu/quantize_ops.cc b/onnxruntime/contrib_ops/cpu/quantize_ops.cc index bd85b89985..46587733c4 100644 --- a/onnxruntime/contrib_ops/cpu/quantize_ops.cc +++ b/onnxruntime/contrib_ops/cpu/quantize_ops.cc @@ -12,8 +12,8 @@ ONNX_CPU_OPERATOR_TYPED_MS_KERNEL( 1, uint8_t, KernelDefBuilder() - .TypeConstraint("T1", DataTypeImpl::GetTensorType()) - .TypeConstraint("T2", DataTypeImpl::GetTensorType()), + .TypeConstraint("T1", DataTypeImpl::GetTensorType()) + .TypeConstraint("T2", DataTypeImpl::GetTensorType()), DequantizeLinear); ONNX_CPU_OPERATOR_TYPED_MS_KERNEL( @@ -21,8 +21,8 @@ ONNX_CPU_OPERATOR_TYPED_MS_KERNEL( 1, int8_t, KernelDefBuilder() - .TypeConstraint("T1", DataTypeImpl::GetTensorType()) - .TypeConstraint("T2", DataTypeImpl::GetTensorType()), + .TypeConstraint("T1", DataTypeImpl::GetTensorType()) + .TypeConstraint("T2", DataTypeImpl::GetTensorType()), DequantizeLinear); ONNX_CPU_OPERATOR_TYPED_MS_KERNEL( @@ -34,5 +34,14 @@ ONNX_CPU_OPERATOR_TYPED_MS_KERNEL( .TypeConstraint("T2", DataTypeImpl::GetTensorType()), QuantizeLinear); +ONNX_CPU_OPERATOR_TYPED_MS_KERNEL( + QuantizeLinear, + 1, + int8_t, + KernelDefBuilder() + .TypeConstraint("T1", DataTypeImpl::GetTensorType()) + .TypeConstraint("T2", DataTypeImpl::GetTensorType()), + QuantizeLinear); + } // namespace contrib } // namespace onnxruntime diff --git a/onnxruntime/contrib_ops/cpu_contrib_kernels.cc b/onnxruntime/contrib_ops/cpu_contrib_kernels.cc index 315e4857c5..923cd6908a 100644 --- a/onnxruntime/contrib_ops/cpu_contrib_kernels.cc +++ b/onnxruntime/contrib_ops/cpu_contrib_kernels.cc @@ -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, BuildKernelCreateInfo, BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, diff --git a/onnxruntime/contrib_ops/cuda/quantize_ops.cc b/onnxruntime/contrib_ops/cuda/quantize_ops.cc new file mode 100644 index 0000000000..f0afad58a5 --- /dev/null +++ b/onnxruntime/contrib_ops/cuda/quantize_ops.cc @@ -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()) \ + .TypeConstraint("T2", DataTypeImpl::GetTensorType()), \ + QuantizeLinear); + +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()) \ + .TypeConstraint("T2", DataTypeImpl::GetTensorType()), \ + DequantizeLinear); + +REGISTER_DQ_KERNEL_TYPED(int8_t, MLFloat16) +REGISTER_DQ_KERNEL_TYPED(uint8_t, MLFloat16) + +} // namespace cuda +} // namespace contrib +} // namespace onnxruntime diff --git a/onnxruntime/contrib_ops/cuda_contrib_kernels.cc b/onnxruntime/contrib_ops/cuda_contrib_kernels.cc index 3c1c576969..9a139d45e9 100644 --- a/onnxruntime/contrib_ops/cuda_contrib_kernels.cc +++ b/onnxruntime/contrib_ops/cuda_contrib_kernels.cc @@ -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, BuildKernelCreateInfo, BuildKernelCreateInfo, - BuildKernelCreateInfo}; + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo}; for (auto& function_table_entry : function_table) { ORT_RETURN_IF_ERROR(kernel_registry.Register(function_table_entry())); diff --git a/onnxruntime/core/graph/contrib_ops/contrib_defs.cc b/onnxruntime/core/graph/contrib_ops/contrib_defs.cc index b2c61d5eee..b7fa331655 100644 --- a/onnxruntime/core/graph/contrib_ops/contrib_defs.cc +++ b/onnxruntime/core/graph/contrib_ops/contrib_defs.cc @@ -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); diff --git a/onnxruntime/core/providers/cpu/tensor/quantize_linear.cc b/onnxruntime/core/providers/cpu/tensor/quantize_linear.cc index bf924dc69f..3ba3ba6837 100644 --- a/onnxruntime/core/providers/cpu/tensor/quantize_linear.cc +++ b/onnxruntime/core/providers/cpu/tensor/quantize_linear.cc @@ -129,25 +129,11 @@ Status QuantizeLinear::Compute(OpKernelContext* ctx) const { const float* input = x.template Data(); T* output = y.template MutableData(); - // 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(block_size), *scale, *zero_point); - } else { - const float qmax = std::numeric_limits::max(); - const float qmin_default = std::numeric_limits::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(N); n++) { - for (size_t bd = 0; bd < static_cast(broadcast_dim); bd++) { - float zp = static_cast(zero_point[bd]); - auto sc = scale[bd]; - - for (size_t bs = 0; bs < static_cast(block_size); bs++) { - *output++ = static_cast(clamp(std::round(static_cast(*input++) / sc) + zp, qmin, qmax)); - } - } + for (size_t n = 0; n < static_cast(N); n++) { + for (size_t bd = 0; bd < static_cast(broadcast_dim); bd++) { + MlasQuantizeLinear(input, output, static_cast(block_size), scale[bd], zero_point[bd]); + input += block_size; + output += block_size; } } diff --git a/onnxruntime/core/providers/cuda/tensor/quantize_linear.cc b/onnxruntime/core/providers/cuda/tensor/quantize_linear.cc index 626313096e..6cc9b6e57d 100644 --- a/onnxruntime/core/providers/cuda/tensor/quantize_linear.cc +++ b/onnxruntime/core/providers/cuda/tensor/quantize_linear.cc @@ -9,28 +9,10 @@ namespace onnxruntime { namespace cuda { -ONNX_OPERATOR_TYPED_KERNEL_EX(QuantizeLinear, - kOnnxDomain, - 10, - uint8_t, - kCudaExecutionProvider, - KernelDefBuilder() - .TypeConstraint("T1", DataTypeImpl::GetTensorType()) - .TypeConstraint("T2", DataTypeImpl::GetTensorType()), - QuantizeLinear); +template +Status QuantizeLinear::ComputeInternal(OpKernelContext* ctx) const { + typedef typename ToCudaType::MappedType CudaU; -ONNX_OPERATOR_TYPED_KERNEL_EX(QuantizeLinear, - kOnnxDomain, - 10, - int8_t, - kCudaExecutionProvider, - KernelDefBuilder() - .TypeConstraint("T1", DataTypeImpl::GetTensorType()) - .TypeConstraint("T2", DataTypeImpl::GetTensorType()), - QuantizeLinear); - -template -Status QuantizeLinear::ComputeInternal(OpKernelContext* ctx) const { auto x = ctx->Input(0); auto y_scale = ctx->Input(1); auto y_zero_point = ctx->Input(2); @@ -42,14 +24,15 @@ Status QuantizeLinear::ComputeInternal(OpKernelContext* ctx) const { const auto& x_shape = x->Shape(); - const float* input = x->template Data(); + const CudaU* input = reinterpret_cast(x->template Data()); T* output = y->template MutableData(); - 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(); - const float* scale = y_scale->template Data(); + const CudaU* scale = reinterpret_cast(y_scale->template Data()); const auto num_of_elements = x_shape.Size(); CudaQuantizeLinear(input, output, scale, zero_point, num_of_elements); @@ -57,26 +40,10 @@ Status QuantizeLinear::ComputeInternal(OpKernelContext* ctx) const { return Status::OK(); } -ONNX_OPERATOR_TYPED_KERNEL_EX(DequantizeLinear, - kOnnxDomain, - 10, - uint8_t, - kCudaExecutionProvider, - KernelDefBuilder() - .TypeConstraint("T", DataTypeImpl::GetTensorType()), - DequantizeLinear); +template +Status DequantizeLinear::ComputeInternal(OpKernelContext* ctx) const { + typedef typename ToCudaType::MappedType CudaU; -ONNX_OPERATOR_TYPED_KERNEL_EX(DequantizeLinear, - kOnnxDomain, - 10, - int8_t, - kCudaExecutionProvider, - KernelDefBuilder() - .TypeConstraint("T", DataTypeImpl::GetTensorType()), - DequantizeLinear); - -template -Status DequantizeLinear::ComputeInternal(OpKernelContext* ctx) const { auto x = ctx->Input(0); auto y_scale = ctx->Input(1); auto y_zero_point = ctx->Input(2); @@ -90,13 +57,13 @@ Status DequantizeLinear::ComputeInternal(OpKernelContext* ctx) const { ORT_ENFORCE(y != nullptr); const T* input = x->template Data(); - float* output = y->template MutableData(); + CudaU* output = reinterpret_cast(y->template MutableData()); - 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(); - const float* scale = y_scale->template Data(); + const CudaU* scale = reinterpret_cast(y_scale->template Data()); const auto num_of_elements = x_shape.Size(); CudaDequantizeLinear(input, output, scale, zero_point, num_of_elements); @@ -104,5 +71,46 @@ Status DequantizeLinear::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()) \ + .TypeConstraint("T2", DataTypeImpl::GetTensorType()), \ + QuantizeLinear); + +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()), \ + DequantizeLinear); + +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::ComputeInternal(OpKernelContext* ctx) const; \ + template Status DequantizeLinear::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 diff --git a/onnxruntime/core/providers/cuda/tensor/quantize_linear.cu b/onnxruntime/core/providers/cuda/tensor/quantize_linear.cu index ec49023ced..b5f97f8e99 100644 --- a/onnxruntime/core/providers/cuda/tensor/quantize_linear.cu +++ b/onnxruntime/core/providers/cuda/tensor/quantize_linear.cu @@ -8,6 +8,20 @@ namespace onnxruntime { namespace cuda { +template +__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(max(-128, min(127, value))); + id += NumThreadsPerBlock; + } + } +} + template __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 -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(CeilDiv(num_of_element, GridDim::maxThreadsPerBlock * GridDim::maxElementsPerThread)); - QuantizeLinearKernel - <<>>( - input, - output, - scale, - zero_point, - num_of_element); - return Status::OK(); -} - -template -__global__ void DequantizeLinearKernel(const T* input, float* output, const float* scale, const T* zero_point, CUDA_LONG N) { +template +__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(max(0, min(255, value))); id += NumThreadsPerBlock; } } } -template -Status CudaDequantizeLinear(const T* input, float* output, const float* scale, const T* zero_point, size_t num_of_element) { +template +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(CeilDiv(num_of_element, GridDim::maxThreadsPerBlock * GridDim::maxElementsPerThread)); - DequantizeLinearKernel - <<>>( - input, - output, - scale, - zero_point, - num_of_element); + QuantizeLinearKernel<<>>( + input, + output, + scale, + zero_point, + num_of_element); return Status::OK(); } -template Status CudaQuantizeLinear(const float* input, int8_t* output, const float* scale, const int8_t* zero_point, size_t num_of_element); -template Status CudaQuantizeLinear(const float* input, uint8_t* output, const float* scale, const uint8_t* zero_point, size_t num_of_element); +template +__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(const int8_t* input, float* output, const float* scale, const int8_t* zero_point, size_t num_of_element); -template Status CudaDequantizeLinear(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((input[id] - *zero_point)) * (*scale); + id += NumThreadsPerBlock; + } + } +} + +template +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(CeilDiv(num_of_element, GridDim::maxThreadsPerBlock * GridDim::maxElementsPerThread)); + DequantizeLinearKernel<<>>( + input, + output, + scale, + zero_point, + num_of_element); + return Status::OK(); +} + +template Status CudaQuantizeLinear(const float* input, int8_t* output, const float* scale, const int8_t* zero_point, size_t num_of_element); +template Status CudaQuantizeLinear(const float* input, uint8_t* output, const float* scale, const uint8_t* zero_point, size_t num_of_element); +template Status CudaQuantizeLinear(const half* input, int8_t* output, const half* scale, const int8_t* zero_point, size_t num_of_element); +template Status CudaQuantizeLinear(const half* input, uint8_t* output, const half* scale, const uint8_t* zero_point, size_t num_of_element); + +template Status CudaDequantizeLinear(const int8_t* input, float* output, const float* scale, const int8_t* zero_point, size_t num_of_element); +template Status CudaDequantizeLinear(const uint8_t* input, float* output, const float* scale, const uint8_t* zero_point, size_t num_of_element); +template Status CudaDequantizeLinear(const int8_t* input, half* output, const half* scale, const int8_t* zero_point, size_t num_of_element); +template Status CudaDequantizeLinear(const uint8_t* input, half* output, const half* scale, const uint8_t* zero_point, size_t num_of_element); } // namespace cuda } // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/tensor/quantize_linear.cuh b/onnxruntime/core/providers/cuda/tensor/quantize_linear.cuh index a5d42d1180..5d140981d6 100644 --- a/onnxruntime/core/providers/cuda/tensor/quantize_linear.cuh +++ b/onnxruntime/core/providers/cuda/tensor/quantize_linear.cuh @@ -11,11 +11,11 @@ namespace onnxruntime { namespace cuda { -template -Status CudaQuantizeLinear(const float* input, T* output, const float* scale, const T* zero_point, size_t num_of_element); +template +Status CudaQuantizeLinear(const U* input, T* output, const U* scale, const T* zero_point, size_t num_of_element); -template -Status CudaDequantizeLinear(const T* input, float* output, const float* scale, const T* zero_point, size_t num_of_element); +template +Status CudaDequantizeLinear(const T* input, U* output, const U* scale, const T* zero_point, size_t num_of_element); } // namespace cuda } // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/tensor/quantize_linear.h b/onnxruntime/core/providers/cuda/tensor/quantize_linear.h index 7c964dc4ee..9a36c71639 100644 --- a/onnxruntime/core/providers/cuda/tensor/quantize_linear.h +++ b/onnxruntime/core/providers/cuda/tensor/quantize_linear.h @@ -11,7 +11,7 @@ namespace onnxruntime { namespace cuda { -template +template 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 +template class DequantizeLinear final : public CudaKernel { public: DequantizeLinear(const OpKernelInfo& info) : CudaKernel(info) {} diff --git a/onnxruntime/test/contrib_ops/quantize_ops_test.cc b/onnxruntime/test/contrib_ops/quantize_ops_test.cc index 4da03df8b1..9baae4b1f4 100644 --- a/onnxruntime/test/contrib_ops/quantize_ops_test.cc +++ b/onnxruntime/test/contrib_ops/quantize_ops_test.cc @@ -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 dims{4}; + test.AddInput("x", dims, {0, 3, 128, 255}); + test.AddInput("x_scale", {}, {2.0f}); + test.AddInput("x_zero_point", {}, {128}); + test.AddOutput("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 dims{4}; + test.AddInput("x", dims, {-30, -3, 100, 127}); + test.AddInput("x_scale", {}, {2.0f}); + test.AddInput("x_zero_point", {}, {-10}); + test.AddOutput("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 dims{4}; + test.AddInput("x", dims, {0, 3, 128, 255}); + test.AddInput("x_scale", {}, ToFloat16({2.0f})); + test.AddInput("x_zero_point", {}, {128}); + test.AddOutput("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 dims{4}; + test.AddInput("x", dims, {-30, -3, 100, 127}); + test.AddInput("x_scale", {}, ToFloat16({2.0f})); + test.AddInput("x_zero_point", {}, {-10}); + test.AddOutput("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 dims{6}; - test.AddInput("x", dims, {0, 2, 3, 1000, -254, -1000}); + std::vector dims{16}; + test.AddInput("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("y_scale", {}, {2.0f}); test.AddInput("y_zero_point", {}, {128}); - test.AddOutput("y", dims, {128, 129, 130, 255, 1, 0}); + test.AddOutput("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 dims{16}; + test.AddInput("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("y_scale", {}, {2.0f}); + test.AddInput("y_zero_point", {}, {1}); + test.AddOutput("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 dims{16}; + test.AddInput("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("y_scale", {}, ToFloat16({2.0f})); + test.AddInput("y_zero_point", {}, {128}); + test.AddOutput("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 dims{16}; + test.AddInput("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("y_scale", {}, ToFloat16({2.0f})); + test.AddInput("y_zero_point", {}, {1}); + test.AddOutput("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 dims{3, 4}; test.AddInput("X", dims, @@ -144,12 +286,12 @@ TEST(QuantizeLinearContribOpTest, QuantizeLinear_1) { test.AddOutput("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 dims{3, 4}; test.AddInput("X", dims, @@ -162,7 +304,7 @@ TEST(QuantizeLinearContribOpTest, QuantizeLinear_2) { test.AddOutput("Y", dims, {0, 2, 3, 255, 0, 1, 2, 255, - 0, 1, 1, 250}); + 0, 0, 1, 250}); test.Run(); } } // namespace test