mirrors/opencv
2
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2026-09-25 04:09:57 +03:00
Code Issues Packages Projects Releases Wiki Activity

Merge pull request #29834 from abhishek-gola:dnn-fp8-support

Browse Source 
FP8 model support in DNN - #29834

ONNX coverage after this PR: 76.6%

co-authored by: @SavyaSanchi-Sharma 

### Pull Request Readiness Checklist

See details at https://github.com/opencv/opencv/wiki/How_to_contribute#making-a-good-pull-request

- [x] I agree to contribute to the project under Apache 2 License.
- [x] To the best of my knowledge, the proposed patch is not based on a code under GPL or another license that is incompatible with OpenCV
- [x] The PR is proposed to the proper branch
- [x] There is a reference to the original bug report and related work
- [x] There is accuracy test, performance test and test data in opencv_extra repository, if applicable
      Patch to opencv_extra has the same branch name.
- [x] The feature is well documented and sample code can be built with the project CMake
This commit is contained in:
Abhishek Gola
2026-09-17 16:15:30 +03:00
committed by GitHub
parent f7d88d11a2
commit b2b4f34820
16 changed files with 362 additions and 58 deletions
Download Patch File Download Diff File Expand all files Collapse all files
modules/core/include/opencv2/core/cvdef.h
+14 -9
View File
@@ -985,39 +985,42 @@ protected:
namespace fp8_detail {
// round-half-up (ties away from zero) — deliberately NOT round-to-nearest-even / OCP-ONNX spec
inline unsigned roundHalfUp(unsigned full, int shift)
// IEEE 754 default rounding, same as cv::hfloat and every hardware FP8 convert.
inline unsigned roundNearestEven(unsigned full, int shift)
{
if (shift <= 0) return full << (-shift);
unsigned q = full >> shift, rem = full & ((1u << shift) - 1), half = 1u << (shift - 1);
if (rem >= half) q++;
if (rem > half || (rem == half && (q & 1))) q++;
return q;
}
inline float pow2(int n) { Cv32suf s; s.u = (unsigned)((n + 127) << 23); return s.f; }
// E4M3 encode; no inf, bias/fnuz select E4M3FN vs E4M3FNUZ
inline uchar encodeE4M3(float x, int bias, bool fnuz)
// E4M3 encode; no inf, bias/fnuz select E4M3FN vs E4M3FNUZ.
// saturate (default): overflow and inf clamp to max finite, otherwise they become NaN.
inline uchar encodeE4M3(float x, int bias, bool fnuz, bool saturate = true)
{
Cv32suf in; in.f = x;
unsigned u = in.u, sign = (u >> 31) & 1, e = (u >> 23) & 0xFF, m = u & 0x7FFFFF;
const unsigned sbit = sign << 7;
const unsigned nanc = fnuz ? 0x80u : (sbit | 0x7Fu);
if (e == 0xFF) return (uchar)nanc;
const unsigned maxfin = sbit | (fnuz ? 0x7Fu : 0x7Eu); // 240 (FNUZ) / 448 (FN)
const unsigned ovc = saturate ? maxfin : nanc;
if (e == 0xFF) return (uchar)(m != 0 ? nanc : ovc); // NaN stays NaN, inf saturates
if (e == 0 && m == 0) return (uchar)(fnuz ? 0u : sbit);
int newexp = (int)e - 127 + bias;
const unsigned full = (1u << 23) | m;
if (newexp <= 0) // subnormal
{
const int shift = 20 + (1 - newexp);
unsigned mant = (shift >= 32) ? 0u : roundHalfUp(full, shift);
unsigned mant = (shift >= 32) ? 0u : roundNearestEven(full, shift);
return (uchar)(mant == 0 ? (fnuz ? 0u : sbit) : (sbit | mant));
}
unsigned rounded = roundHalfUp(full, 20);
unsigned rounded = roundNearestEven(full, 20);
if (rounded & 16u) { rounded >>= 1; newexp++; } // carry into exponent
const unsigned mant = rounded & 7u;
const bool ov = fnuz ? ((unsigned)newexp > 15)
: ((unsigned)newexp > 15 || ((unsigned)newexp == 15 && mant == 7));
if (ov) return (uchar)nanc;
if (ov) return (uchar)ovc;
return (uchar)(sbit | ((unsigned)newexp << 3) | mant);
}
@@ -1038,6 +1041,7 @@ struct fp8_t // E4M3FN: bias 7, no inf, max 448
{
fp8_t() : b(0) {}
explicit fp8_t(float x) : b(fp8_detail::encodeE4M3(x, 7, false)) {}
fp8_t(float x, bool saturate) : b(fp8_detail::encodeE4M3(x, 7, false, saturate)) {}
operator float() const { return table()[b]; }
static const float* decodeLUT() { return table(); }
protected:
@@ -1051,6 +1055,7 @@ struct fp8a_t // E4M3FNUZ: bias 8, no inf, single NaN, no -0, max 240
{
fp8a_t() : b(0) {}
explicit fp8a_t(float x) : b(fp8_detail::encodeE4M3(x, 8, true)) {}
fp8a_t(float x, bool saturate) : b(fp8_detail::encodeE4M3(x, 8, true, saturate)) {}
operator float() const { return table()[b]; }
static const float* decodeLUT() { return table(); }
protected:
modules/core/src/convert.hpp
+7 -8
View File
@@ -491,13 +491,16 @@ fp8Prepare(const v_float32& vf, v_uint32& full, v_uint32& sbit, v_int32& newexpR
fallbackMask = v_or(isNan, v_le(newexpRaw, vx_setall_s32(0)));
}
// round-half-up + carry + overflow-to-NaN; valid only when fp8Prepare's fallbackMask is false
// Valid only when fp8Prepare's fallbackMask is false; must stay bit-identical to encodeE4M3.
template<int bias, bool fnuz> static inline v_int32
encodeFp8Finish(v_uint32 full, v_uint32 sbit, v_int32 newexpRaw)
{
v_uint32 one = vx_setall_u32(1u);
v_uint32 half = vx_setall_u32(1u << 19);
v_uint32 q = v_shr<20>(full);
v_uint32 rem = v_and(full, vx_setall_u32((1u << 20) - 1));
v_uint32 inc = v_and(v_ge(rem, vx_setall_u32(1u << 19)), vx_setall_u32(1u));
v_uint32 tie = v_and(v_eq(rem, half), v_ne(v_and(q, one), vx_setzero_u32()));
v_uint32 inc = v_and(v_or(v_gt(rem, half), tie), one);
v_uint32 rounded = v_add(q, inc);
v_uint32 carry = v_ne(v_and(rounded, vx_setall_u32(16u)), vx_setall_u32(0u));
@@ -514,14 +517,10 @@ encodeFp8Finish(v_uint32 full, v_uint32 sbit, v_int32 newexpRaw)
overflow = v_or(gt15, v_and(v_eq(newexp, vx_setall_s32(15)),
v_eq(v_reinterpret_as_s32(mant), vx_setall_s32(7))));
v_uint32 nanCode;
if constexpr (fnuz)
nanCode = vx_setall_u32(0x80u);
else
nanCode = v_or(sbit, vx_setall_u32(0x7Fu));
v_uint32 maxFin = v_or(sbit, vx_setall_u32(fnuz ? 0x7Fu : 0x7Eu));
v_uint32 normal = v_or(v_or(sbit, v_shl<3>(v_reinterpret_as_u32(newexp))), mant);
v_uint32 result = v_select(v_reinterpret_as_u32(overflow), nanCode, normal);
v_uint32 result = v_select(v_reinterpret_as_u32(overflow), maxFin, normal);
return v_reinterpret_as_s32(result);
}
modules/core/test/test_fp8.cpp
+21 -6
View File
@@ -41,6 +41,11 @@ TEST(Core_FP8, scalar_roundtrip_exact)
}
// round-to-nearest-even onto the grid
EXPECT_EQ((float)cv::fp8_t(1.234f), 1.25f); // 3 mantissa bits
// exact ties go to the even code, not away from zero
EXPECT_EQ((float)cv::fp8_t(1.0625f), 1.f); // between 1.0 (even) and 1.125
EXPECT_EQ((float)cv::fp8_t(1.1875f), 1.25f); // between 1.125 and 1.25 (even)
EXPECT_EQ((float)cv::fp8_t(-1.0625f), -1.f);
EXPECT_EQ((float)cv::fp8_t(0.0009765625f), 0.f); // subnormal tie -> even
}
TEST(Core_FP8, format_specific_limits)
@@ -49,14 +54,24 @@ TEST(Core_FP8, format_specific_limits)
EXPECT_EQ((float)cv::fp8_t(448.f), 448.f);
EXPECT_EQ((float)cv::fp8a_t(240.f), 240.f);
// overflow: these formats have no inf -> overflow to NaN
EXPECT_TRUE(cvIsNaN((float)cv::fp8_t(1e6f)));
EXPECT_TRUE(cvIsNaN((float)cv::fp8a_t(1e6f)));
// 448 exceeds the FNUZ E4M3 range (max 240) -> NaN
EXPECT_TRUE(cvIsNaN((float)cv::fp8a_t(448.f)));
// overflow: these formats have no inf, so out-of-range magnitudes clamp to max finite
EXPECT_EQ((float)cv::fp8_t(1e6f), 448.f);
EXPECT_EQ((float)cv::fp8a_t(1e6f), 240.f);
EXPECT_EQ((float)cv::fp8_t(-1e6f), -448.f);
// 448 exceeds the FNUZ E4M3 range (max 240)
EXPECT_EQ((float)cv::fp8a_t(448.f), 240.f);
const float inf = std::numeric_limits<float>::infinity();
EXPECT_EQ((float)cv::fp8_t(inf), 448.f);
EXPECT_EQ((float)cv::fp8a_t(-inf), -240.f);
// NaN propagates
// saturate=false keeps the old overflow-to-NaN behaviour (ONNX saturate=0)
EXPECT_TRUE(cvIsNaN((float)cv::fp8_t(1e6f, false)));
EXPECT_TRUE(cvIsNaN((float)cv::fp8a_t(1e6f, false)));
EXPECT_TRUE(cvIsNaN((float)cv::fp8_t(inf, false)));
// NaN propagates regardless
EXPECT_TRUE(cvIsNaN((float)cv::fp8_t(std::numeric_limits<float>::quiet_NaN())));
EXPECT_TRUE(cvIsNaN((float)cv::fp8_t(std::numeric_limits<float>::quiet_NaN(), false)));
// smallest E4M3FN subnormal is 2^-9
EXPECT_EQ((float)cv::fp8_t(0.001953125f), 0.001953125f);
modules/dnn/include/opencv2/dnn/all_layers.hpp
+1
View File
@@ -810,6 +810,7 @@ CV__DNN_INLINE_NS_BEGIN
int axis;
int block_size;
int output_dtype;
int output_onnx_dtype; // raw ONNX dtype; disambiguates E5M2 vs E5M2FNUZ
bool saturate;
static Ptr<QuantizeLinearLayer> create(const LayerParams& params);
modules/dnn/src/layers/cast2_layer.cpp
+1 -1
View File
@@ -333,7 +333,7 @@ public:
const int ddepth = dst.depth();
if (ddepth == CV_8F_E4M3FN || ddepth == CV_8F_E4M3FNUZ)
{
// Store the ONNX-encoded byte: core's E4M3 encode rounds differently.
// Native depth: store the raw FP8 byte.
uchar* d = dst.ptr<uchar>();
for (size_t i = 0; i < total; i++)
d[i] = onnx_dtype::f32ToFp8(CV_DNN_SRC_F(i), fmt, saturate);
modules/dnn/src/layers/conv2_layer.cpp
+11 -2
View File
@@ -119,7 +119,16 @@ public:
Mat bias_ = bias_arr.getMat();
CV_Assert(!weights_.empty());
int wtype0 = weights_.type();
CV_Assert(wtype0 == CV_32F || wtype0 == CV_16F || wtype0 == CV_16BF);
CV_Assert(wtype0 == CV_32F || wtype0 == CV_16F || wtype0 == CV_16BF ||
wtype0 == CV_8F_E4M3FN || wtype0 == CV_8F_E4M3FNUZ);
// repackConvWeights is FP32-only; FP8 reaches here only from vendor-quantised
// graphs that feed Conv without a DequantizeLinear. Half stays as it was --
// repackDepthwiseConvWeights takes CV_16F/CV_16BF natively.
if (wtype0 == CV_8F_E4M3FN || wtype0 == CV_8F_E4M3FNUZ) {
Mat widened;
weights_.convertTo(widened, CV_32F);
weights_ = widened;
}
CV_Assert(accuracy == -1 || accuracy == CV_32F);
int wtype = accuracy < 0 ? CV_32F : accuracy;
@@ -145,7 +154,7 @@ public:
// >= 256*256 so the reorder cost is amortized).
mlas_packed_B_.release();
mlas_packed_M_ = mlas_packed_K_ = 0;
if (!depthwise && ngroups == 1 && wtype0 == CV_32F &&
if (!depthwise && ngroups == 1 && weights_.type() == CV_32F &&
wtype == CV_32F && mlasAvailable())
{
bool ksize_all_one = wshape0.dims >= 3;
modules/dnn/src/layers/dequantizelinear_layer.cpp
+108 -3
View File
@@ -7,6 +7,7 @@
#include "../precomp.hpp"
#include "layers_common.hpp"
#include "../net_impl.hpp"
#include "../onnx/onnx_dtype_convert.hpp"
#if defined(__x86_64__) || defined(_M_X64)
#include <immintrin.h>
@@ -160,15 +161,65 @@ static void dequantizeLinear(const _InpTp* inp_, const _ScaleTp* scale_,
});
}
// Native FP8 bytes through core's decode table, bit-identical to the ONNX reference.
template <typename _ScaleTp, typename _OutTp>
static void dequantizeLinearFp8Native(const uchar* inp, const _ScaleTp* scale, const uchar* zp,
_OutTp* out, const float* fp8lut,
int64_t nslices, int sz_a, int64_t slice_size)
{
parallel_for_(Range(0, (int)nslices), [&](const Range& r) {
for (int slice = r.start; slice < r.end; slice++) {
size_t base = (size_t)slice * sz_a * slice_size;
for (int a = 0; a < sz_a; a++) {
float sc = (float)scale[a];
float zpv = zp ? fp8lut[zp[a]] : 0.f;
for (int64_t j = 0; j < slice_size; j++) {
size_t idx = base + (size_t)a * slice_size + j;
out[idx] = _OutTp((fp8lut[inp[idx]] - zpv) * sc);
}
}
}
});
}
// E5M2 has no native depth; already decoded to real values upstream, held in
// CV_16F, or in CV_32F once setInput has widened a graph input.
template <typename _InpTp, typename _ScaleTp, typename _OutTp>
static void dequantizeLinearFp8Wide(const _InpTp* inp, const _ScaleTp* scale, const _InpTp* zp,
_OutTp* out, int64_t nslices, int sz_a, int64_t slice_size)
{
parallel_for_(Range(0, (int)nslices), [&](const Range& r) {
for (int slice = r.start; slice < r.end; slice++) {
size_t base = (size_t)slice * sz_a * slice_size;
for (int a = 0; a < sz_a; a++) {
float sc = (float)scale[a];
float zpv = zp ? (float)zp[a] : 0.f;
for (int64_t j = 0; j < slice_size; j++) {
size_t idx = base + (size_t)a * slice_size + j;
out[idx] = _OutTp(((float)inp[idx] - zpv) * sc);
}
}
}
});
}
// Dequantize INT8/UINT8 to FP32/FP16; out must be preallocated
static void dequantizeLinear(const Mat& inp, const Mat& scale_, const Mat& zp,
static void dequantizeLinear(const Mat& inp, const Mat& scale_, const Mat& zp_,
int axis, int block_size, Mat& out)
{
Mat scale = scale_;
Mat zp = zp_;
CV_Assert(inp.isContinuous());
CV_Assert(scale.isContinuous());
CV_Assert(out.isContinuous());
// E5M2 rides in CV_16F, but widenHalfConstants() may already have promoted a constant
// zero_point to CV_32F, so the two can reach here at different widths.
if (!zp.empty() && zp.type() != inp.type() &&
(inp.type() == CV_16F || inp.type() == CV_32F) &&
(zp.type() == CV_16F || zp.type() == CV_32F))
zp_.convertTo(zp, inp.type());
int inptype = inp.type();
int outtype = out.type();
int sctype = scale.type();
@@ -179,9 +230,14 @@ static void dequantizeLinear(const Mat& inp, const Mat& scale_, const Mat& zp,
int i, ndims = inpshape.dims;
int64_t nslices = 1, slice_size = 1;
CV_Assert(inptype == CV_8U || inptype == CV_8S || inptype == CV_32S);
CV_Assert(inptype == CV_8U || inptype == CV_8S || inptype == CV_32S ||
inptype == CV_8F_E4M3FN || inptype == CV_8F_E4M3FNUZ ||
inptype == CV_16F || inptype == CV_32F);
CV_Assert(sctype == CV_32F || sctype == CV_16F);
CV_Assert(outtype == CV_32F || outtype == CV_16F);
if (inptype == CV_8F_E4M3FN || inptype == CV_8F_E4M3FNUZ ||
inptype == CV_16F || inptype == CV_32F)
CV_Assert(block_size == 0); // block-wise FP8 dequantization not yet supported
if (!zp.empty()) {
CV_Assert(zp.isContinuous());
@@ -317,6 +373,52 @@ static void dequantizeLinear(const Mat& inp, const Mat& scale_, const Mat& zp,
reinterpret_cast<const int32_t*>(zp.data),
reinterpret_cast<hfloat*>(out.data),
nslices, sz_a, slice_size, block_size);
else if (inptype == CV_8F_E4M3FN || inptype == CV_8F_E4M3FNUZ) {
const float* fp8lut = inptype == CV_8F_E4M3FN ? fp8_t::decodeLUT() : fp8a_t::decodeLUT();
const uchar* zpdata = zp.empty() ? nullptr : reinterpret_cast<const uchar*>(zp.data);
if (sctype == CV_32F && outtype == CV_32F)
dequantizeLinearFp8Native(reinterpret_cast<const uchar*>(inp.data), reinterpret_cast<const float*>(scale.data),
zpdata, reinterpret_cast<float*>(out.data), fp8lut, nslices, sz_a, slice_size);
else if (sctype == CV_16F && outtype == CV_32F)
dequantizeLinearFp8Native(reinterpret_cast<const uchar*>(inp.data), reinterpret_cast<const hfloat*>(scale.data),
zpdata, reinterpret_cast<float*>(out.data), fp8lut, nslices, sz_a, slice_size);
else if (sctype == CV_32F && outtype == CV_16F)
dequantizeLinearFp8Native(reinterpret_cast<const uchar*>(inp.data), reinterpret_cast<const float*>(scale.data),
zpdata, reinterpret_cast<hfloat*>(out.data), fp8lut, nslices, sz_a, slice_size);
else
dequantizeLinearFp8Native(reinterpret_cast<const uchar*>(inp.data), reinterpret_cast<const hfloat*>(scale.data),
zpdata, reinterpret_cast<hfloat*>(out.data), fp8lut, nslices, sz_a, slice_size);
}
else if (inptype == CV_16F) {
const hfloat* zpdata = zp.empty() ? nullptr : reinterpret_cast<const hfloat*>(zp.data);
if (sctype == CV_32F && outtype == CV_32F)
dequantizeLinearFp8Wide(reinterpret_cast<const hfloat*>(inp.data), reinterpret_cast<const float*>(scale.data),
zpdata, reinterpret_cast<float*>(out.data), nslices, sz_a, slice_size);
else if (sctype == CV_16F && outtype == CV_32F)
dequantizeLinearFp8Wide(reinterpret_cast<const hfloat*>(inp.data), reinterpret_cast<const hfloat*>(scale.data),
zpdata, reinterpret_cast<float*>(out.data), nslices, sz_a, slice_size);
else if (sctype == CV_32F && outtype == CV_16F)
dequantizeLinearFp8Wide(reinterpret_cast<const hfloat*>(inp.data), reinterpret_cast<const float*>(scale.data),
zpdata, reinterpret_cast<hfloat*>(out.data), nslices, sz_a, slice_size);
else
dequantizeLinearFp8Wide(reinterpret_cast<const hfloat*>(inp.data), reinterpret_cast<const hfloat*>(scale.data),
zpdata, reinterpret_cast<hfloat*>(out.data), nslices, sz_a, slice_size);
}
else if (inptype == CV_32F) {
const float* zpdata = zp.empty() ? nullptr : reinterpret_cast<const float*>(zp.data);
if (sctype == CV_32F && outtype == CV_32F)
dequantizeLinearFp8Wide(reinterpret_cast<const float*>(inp.data), reinterpret_cast<const float*>(scale.data),
zpdata, reinterpret_cast<float*>(out.data), nslices, sz_a, slice_size);
else if (sctype == CV_16F && outtype == CV_32F)
dequantizeLinearFp8Wide(reinterpret_cast<const float*>(inp.data), reinterpret_cast<const hfloat*>(scale.data),
zpdata, reinterpret_cast<float*>(out.data), nslices, sz_a, slice_size);
else if (sctype == CV_32F && outtype == CV_16F)
dequantizeLinearFp8Wide(reinterpret_cast<const float*>(inp.data), reinterpret_cast<const float*>(scale.data),
zpdata, reinterpret_cast<hfloat*>(out.data), nslices, sz_a, slice_size);
else
dequantizeLinearFp8Wide(reinterpret_cast<const float*>(inp.data), reinterpret_cast<const hfloat*>(scale.data),
zpdata, reinterpret_cast<hfloat*>(out.data), nslices, sz_a, slice_size);
}
else {
CV_Error_(Error::StsNotImplemented,
("the following combination of types is not supported in "
@@ -371,7 +473,10 @@ public:
size_t ninputs = inputs.size();
CV_Assert(2 <= ninputs && ninputs <= 3);
if (ninputs == 3) {
CV_Assert(inputs[0] == inputs[2]);
// Both widths are accepted on the E5M2 pair; see dequantizeLinear() above.
const bool wideFp8Pair = (inputs[0] == CV_16F || inputs[0] == CV_32F) &&
(inputs[2] == CV_16F || inputs[2] == CV_32F);
CV_Assert(inputs[0] == inputs[2] || wideFp8Pair);
}
outputs.assign(1, getOutType());
}
modules/dnn/src/layers/gather_elements_layer.cpp
+2 -1
View File
@@ -68,7 +68,8 @@ public:
{
CV_CheckEQ(inputs.size(), (size_t)2, "");
CV_CheckType(inputs[0], inputs[0] == CV_32F || inputs[0] == CV_32S || inputs[0] == CV_64S || inputs[0] == CV_16F || inputs[0] == CV_8U || inputs[0] == CV_8S || inputs[0] == CV_Bool ||
inputs[0] == CV_64F || inputs[0] == CV_16U || inputs[0] == CV_16S || inputs[0] == CV_32U || inputs[0] == CV_64U, "");
inputs[0] == CV_64F || inputs[0] == CV_16U || inputs[0] == CV_16S || inputs[0] == CV_32U || inputs[0] == CV_64U ||
inputs[0] == CV_8F_E4M3FN || inputs[0] == CV_8F_E4M3FNUZ, "");
CV_CheckType(inputs[1], inputs[1] == CV_64S || inputs[1] == CV_32S, "");
outputs.assign(1, inputs[0]);
}
modules/dnn/src/layers/gather_layer.cpp
+2 -1
View File
@@ -54,7 +54,8 @@ public:
std::vector<MatType>& internals) const CV_OVERRIDE
{
CV_CheckEQ(inputs.size(), (size_t)2, "");
CV_CheckType(inputs[0], inputs[0] == CV_32F || inputs[0] == CV_32S || inputs[0] == CV_64S || inputs[0] == CV_16F || inputs[0] == CV_8U || inputs[0] == CV_8S || inputs[0] == CV_Bool, "");
CV_CheckType(inputs[0], inputs[0] == CV_32F || inputs[0] == CV_32S || inputs[0] == CV_64S || inputs[0] == CV_16F || inputs[0] == CV_8U || inputs[0] == CV_8S || inputs[0] == CV_Bool ||
inputs[0] == CV_8F_E4M3FN || inputs[0] == CV_8F_E4M3FNUZ, "");
CV_CheckType(inputs[1], inputs[1] == CV_64S || inputs[1] == CV_32S, "");
outputs.assign(1, inputs[0]);
}
modules/dnn/src/layers/gemm_layer.cpp
+2 -1
View File
@@ -80,7 +80,8 @@ public:
real_ndims_C = params.get<int>("real_ndims_C", -1);
for (Mat& blob : blobs) {
if (blob.type() == CV_16F || blob.type() == CV_16BF) {
if (blob.type() == CV_16F || blob.type() == CV_16BF ||
blob.type() == CV_8F_E4M3FN || blob.type() == CV_8F_E4M3FNUZ) {
Mat widened;
blob.convertTo(widened, CV_32F);
blob = widened;
modules/dnn/src/layers/matmul_layer.cpp
+5 -1
View File
@@ -51,8 +51,12 @@ class MatMulLayerImpl CV_FINAL : public MatMulLayer {
real_ndims_C = params.get<int>("real_ndims_C", -1);
// The GEMM kernels are FP32-only, so narrow constant weights are decoded
// once here. FP8 arrives this way from vendor-quantised models that feed
// MatMul directly instead of going through DequantizeLinear.
for (Mat& blob : blobs) {
if (blob.type() == CV_16F || blob.type() == CV_16BF) {
if (blob.type() == CV_16F || blob.type() == CV_16BF ||
blob.type() == CV_8F_E4M3FN || blob.type() == CV_8F_E4M3FNUZ) {
Mat widened;
blob.convertTo(widened, CV_32F);
blob = widened;
modules/dnn/src/layers/quantlizelinear_layer.cpp
+117 -7
View File
@@ -6,6 +6,7 @@
#include "../precomp.hpp"
#include "layers_common.hpp"
#include "../net_impl.hpp"
#include "../onnx/onnx_dtype_convert.hpp"
#if defined(__x86_64__) || defined(_M_X64)
#include <immintrin.h>
@@ -151,9 +152,53 @@ static void quantizeLinear(const _InpTp* inp_, const _ScaleTp* scale_,
});
}
// Rounding is the fp8-grid snap itself, done by f32ToFp8 - no separate int round step.
template <typename _InpTp, typename _ScaleTp>
static void quantizeLinearToFp8Native(const _InpTp* inp, const _ScaleTp* scale, const uchar* zp,
uchar* out, const onnx_dtype::Fp8Fmt& fmt, bool saturate,
int64_t nslices, int sz_a, int64_t slice_size)
{
parallel_for_(Range(0, (int)nslices), [&](const Range& r) {
for (int slice = r.start; slice < r.end; slice++) {
size_t base = (size_t)slice * sz_a * slice_size;
for (int a = 0; a < sz_a; a++) {
float sc = (float)scale[a];
float zpv = zp ? onnx_dtype::fp8ToF32(zp[a], fmt) : 0.f;
for (int64_t j = 0; j < slice_size; j++) {
size_t idx = base + (size_t)a * slice_size + j;
out[idx] = onnx_dtype::f32ToFp8((float)inp[idx] / sc + zpv, fmt, saturate);
}
}
}
});
}
// E5M2/E5M2FNUZ have no native depth; store the grid-snapped value in CV_16F, or
// in CV_32F when setInput has widened the graph's zero_point and output.
template <typename _InpTp, typename _ScaleTp, typename _OutTp>
static void quantizeLinearToFp8Wide(const _InpTp* inp, const _ScaleTp* scale, const _OutTp* zp,
_OutTp* out, const onnx_dtype::Fp8Fmt& fmt, bool saturate,
int64_t nslices, int sz_a, int64_t slice_size)
{
parallel_for_(Range(0, (int)nslices), [&](const Range& r) {
for (int slice = r.start; slice < r.end; slice++) {
size_t base = (size_t)slice * sz_a * slice_size;
for (int a = 0; a < sz_a; a++) {
float sc = (float)scale[a];
float zpv = zp ? (float)zp[a] : 0.f;
for (int64_t j = 0; j < slice_size; j++) {
size_t idx = base + (size_t)a * slice_size + j;
uint8_t code = onnx_dtype::f32ToFp8((float)inp[idx] / sc + zpv, fmt, saturate);
out[idx] = _OutTp(onnx_dtype::fp8ToF32(code, fmt));
}
}
}
});
}
// Dequantize INT8/UINT8 to FP32/FP16; out must be preallocated
static void quantizeLinear(const Mat& inp, const Mat& scale_, const Mat& zp,
int axis, int block_size, Mat& out)
int axis, int block_size, int outputOnnxDtype, bool saturate, Mat& out)
{
Mat scale = scale_;
CV_Assert(inp.isContinuous());
@@ -172,7 +217,20 @@ static void quantizeLinear(const Mat& inp, const Mat& scale_, const Mat& zp,
CV_Assert(inptype == CV_32F || inptype == CV_16F);
CV_Assert(sctype == CV_32F || sctype == CV_16F);
CV_Assert(outtype == CV_8U || outtype == CV_8S);
// E5M2/E5M2FNUZ have no native depth, so they travel in CV_16F, or in CV_32F
// once setInput has widened a graph tensor. Both are ambiguous on their own -
// FLOAT4E2M1 and FLOAT8E8M0 map onto the same depths - so the raw ONNX dtype
// has to confirm this really is an FP8 request.
const bool wideFp8Out = (outtype == CV_16F || outtype == CV_32F) &&
onnx_dtype::isFp8(outputOnnxDtype);
// CV_16F/CV_32F are let through here so an unsupported request (FLOAT4E2M1,
// FLOAT8E8M0, plain fp16) reaches the descriptive error below instead of
// tripping a bare assertion.
CV_Assert(outtype == CV_8U || outtype == CV_8S ||
outtype == CV_8F_E4M3FN || outtype == CV_8F_E4M3FNUZ ||
outtype == CV_16F || outtype == CV_32F);
if (outtype == CV_8F_E4M3FN || outtype == CV_8F_E4M3FNUZ || wideFp8Out)
CV_Assert(block_size == 0); // block-wise FP8 quantization not yet supported
if (!zp.empty()) {
CV_Assert(zp.isContinuous());
@@ -289,13 +347,65 @@ static void quantizeLinear(const Mat& inp, const Mat& scale_, const Mat& zp,
reinterpret_cast<const int8_t*>(zp.data),
reinterpret_cast<int8_t*>(out.data),
nslices, sz_a, slice_size, block_size);
else if (outtype == CV_8F_E4M3FN || outtype == CV_8F_E4M3FNUZ) {
const onnx_dtype::Fp8Fmt fmt = onnx_dtype::fp8FmtFor(outtype == CV_8F_E4M3FN ? 17 : 18);
const uchar* zpdata = zp.empty() ? nullptr : reinterpret_cast<const uchar*>(zp.data);
uchar* d = reinterpret_cast<uchar*>(out.data);
if (inptype == CV_32F && sctype == CV_32F)
quantizeLinearToFp8Native(reinterpret_cast<const float*>(inp.data), reinterpret_cast<const float*>(scale.data),
zpdata, d, fmt, saturate, nslices, sz_a, slice_size);
else if (inptype == CV_32F && sctype == CV_16F)
quantizeLinearToFp8Native(reinterpret_cast<const float*>(inp.data), reinterpret_cast<const hfloat*>(scale.data),
zpdata, d, fmt, saturate, nslices, sz_a, slice_size);
else if (inptype == CV_16F && sctype == CV_32F)
quantizeLinearToFp8Native(reinterpret_cast<const hfloat*>(inp.data), reinterpret_cast<const float*>(scale.data),
zpdata, d, fmt, saturate, nslices, sz_a, slice_size);
else
quantizeLinearToFp8Native(reinterpret_cast<const hfloat*>(inp.data), reinterpret_cast<const hfloat*>(scale.data),
zpdata, d, fmt, saturate, nslices, sz_a, slice_size);
}
else if (wideFp8Out && outtype == CV_16F) {
const onnx_dtype::Fp8Fmt fmt = onnx_dtype::fp8FmtFor(outputOnnxDtype);
const hfloat* zpdata = zp.empty() ? nullptr : reinterpret_cast<const hfloat*>(zp.data);
hfloat* d = reinterpret_cast<hfloat*>(out.data);
if (inptype == CV_32F && sctype == CV_32F)
quantizeLinearToFp8Wide(reinterpret_cast<const float*>(inp.data), reinterpret_cast<const float*>(scale.data),
zpdata, d, fmt, saturate, nslices, sz_a, slice_size);
else if (inptype == CV_32F && sctype == CV_16F)
quantizeLinearToFp8Wide(reinterpret_cast<const float*>(inp.data), reinterpret_cast<const hfloat*>(scale.data),
zpdata, d, fmt, saturate, nslices, sz_a, slice_size);
else if (inptype == CV_16F && sctype == CV_32F)
quantizeLinearToFp8Wide(reinterpret_cast<const hfloat*>(inp.data), reinterpret_cast<const float*>(scale.data),
zpdata, d, fmt, saturate, nslices, sz_a, slice_size);
else
quantizeLinearToFp8Wide(reinterpret_cast<const hfloat*>(inp.data), reinterpret_cast<const hfloat*>(scale.data),
zpdata, d, fmt, saturate, nslices, sz_a, slice_size);
}
else if (wideFp8Out) {
const onnx_dtype::Fp8Fmt fmt = onnx_dtype::fp8FmtFor(outputOnnxDtype);
const float* zpdata = zp.empty() ? nullptr : reinterpret_cast<const float*>(zp.data);
float* d = reinterpret_cast<float*>(out.data);
if (inptype == CV_32F && sctype == CV_32F)
quantizeLinearToFp8Wide(reinterpret_cast<const float*>(inp.data), reinterpret_cast<const float*>(scale.data),
zpdata, d, fmt, saturate, nslices, sz_a, slice_size);
else if (inptype == CV_32F && sctype == CV_16F)
quantizeLinearToFp8Wide(reinterpret_cast<const float*>(inp.data), reinterpret_cast<const hfloat*>(scale.data),
zpdata, d, fmt, saturate, nslices, sz_a, slice_size);
else if (inptype == CV_16F && sctype == CV_32F)
quantizeLinearToFp8Wide(reinterpret_cast<const hfloat*>(inp.data), reinterpret_cast<const float*>(scale.data),
zpdata, d, fmt, saturate, nslices, sz_a, slice_size);
else
quantizeLinearToFp8Wide(reinterpret_cast<const hfloat*>(inp.data), reinterpret_cast<const hfloat*>(scale.data),
zpdata, d, fmt, saturate, nslices, sz_a, slice_size);
}
else {
CV_Error_(Error::StsNotImplemented,
("the following combination of types is not supported in "
"QuantizeLinear: inp=%s, scale=%s, out=%s",
"QuantizeLinear: inp=%s, scale=%s, out=%s, onnx_output_dtype=%d",
typeToString(inptype).c_str(),
typeToString(sctype).c_str(),
typeToString(outtype).c_str()));
typeToString(outtype).c_str(),
outputOnnxDtype));
}
}
@@ -310,8 +420,8 @@ public:
block_size = params.get<int>("block_size", 0);
saturate = params.get<bool>("saturate", true);
output_dtype = params.get<int>("output_dtype", -1);
output_onnx_dtype = params.get<int>("output_onnx_dtype", -1);
CV_Assert(block_size >= 0);
CV_Assert(saturate);
}
virtual bool supportBackend(int backendId) CV_OVERRIDE
@@ -383,13 +493,13 @@ public:
std::vector<Mat>& outs = outputs_arr.getMatVecRef();
outs.resize(1);
outs[0].fit(inpshape, outtype);
quantizeLinear(inp, scale, zeropoint, axis, block_size, outs[0]);
quantizeLinear(inp, scale, zeropoint, axis, block_size, output_onnx_dtype, saturate, outs[0]);
} else if (kind == _InputArray::STD_VECTOR_UMAT) {
std::vector<UMat>& outs = outputs_arr.getUMatVecRef();
outs.resize(1);
outs[0].fit(inpshape, outtype);
Mat temp(inpshape, outtype);
quantizeLinear(inp, scale, zeropoint, axis, block_size, temp);
quantizeLinear(inp, scale, zeropoint, axis, block_size, output_onnx_dtype, saturate, temp);
temp.copyTo(outs[0]);
} else {
CV_Error(Error::StsNotImplemented, "");
modules/dnn/src/onnx/onnx_graph_simplifier.cpp
+28 -7
View File
@@ -2023,6 +2023,22 @@ Mat getMatFromTensor(const opencv_onnx::TensorProto& tensor_proto, bool uint8ToI
"DNN/ONNX: tensor payload is smaller than its declared shape");
};
// The sub-byte and FP8 dtypes below are stored either in raw_data or, one byte's
// worth per entry, in int32_data. Narrow int32_data to bytes so both spellings
// reach the decode loops as the same buffer.
std::vector<uchar> narrowed;
const uchar* payload = reinterpret_cast<const uchar*>(rawdata);
size_t payload_bytes = raw_data_size;
if (!tensor_proto.int32_data().empty() && onnx_dtype::isExotic(datatype))
{
const auto& i32 = tensor_proto.int32_data();
narrowed.resize(i32.size());
for (int i = 0; i < i32.size(); i++)
narrowed[i] = static_cast<uchar>(i32[i] & 0xFF);
payload = narrowed.data();
payload_bytes = narrowed.size();
}
if (datatype == opencv_onnx::TensorProto_DataType_FLOAT) {
if (!tensor_proto.float_data().empty()) {
checkPayloadSize(tensor_proto.float_data().size());
@@ -2236,50 +2252,55 @@ Mat getMatFromTensor(const opencv_onnx::TensorProto& tensor_proto, bool uint8ToI
datatype == opencv_onnx::TensorProto_DataType_FLOAT8E4M3FNUZ)
{
// E4M3FN/E4M3FNUZ have a native depth: keep the raw FP8 bytes.
checkPayloadSize(raw_data_size);
checkPayloadSize(payload_bytes);
blob.create((int)sizes.size(), sizes.data(),
CV_MAKETYPE(onnx_dtype::fp8NativeDepth(datatype), 1));
memcpy(blob.data, rawdata, (size_t)blob.total() * blob.elemSize());
memcpy(blob.data, payload, (size_t)blob.total() * blob.elemSize());
}
else if (datatype == opencv_onnx::TensorProto_DataType_FLOAT8E5M2 ||
datatype == opencv_onnx::TensorProto_DataType_FLOAT8E5M2FNUZ)
{
// E5M2/E5M2FNUZ have no native depth: decode losslessly into CV_16F.
const onnx_dtype::Fp8Fmt fmt = onnx_dtype::fp8FmtFor(datatype);
checkPayloadSize(payload_bytes);
blob.create((int)sizes.size(), sizes.data(), CV_16FC1);
const uchar* src = (const uchar*)rawdata;
const uchar* src = payload;
hfloat* dst = blob.ptr<hfloat>();
for (size_t i = 0, total = blob.total(); i < total; i++)
dst[i] = hfloat(onnx_dtype::fp8ToF32(src[i], fmt));
}
else if (datatype == onnx_dtype::ONNX_FLOAT8E8M0)
{
checkPayloadSize(payload_bytes);
blob.create((int)sizes.size(), sizes.data(), CV_32FC1);
const uchar* src = (const uchar*)rawdata;
const uchar* src = payload;
float* dst = blob.ptr<float>();
for (size_t i = 0, total = blob.total(); i < total; i++)
dst[i] = onnx_dtype::e8m0ToF32(src[i]);
}
else if (datatype == opencv_onnx::TensorProto_DataType_FLOAT4E2M1)
{
checkPayloadSize(payload_bytes * 2); // two 4-bit elements per byte
blob.create((int)sizes.size(), sizes.data(), CV_16FC1);
const uchar* src = (const uchar*)rawdata;
const uchar* src = payload;
hfloat* dst = blob.ptr<hfloat>();
for (size_t i = 0, total = blob.total(); i < total; i++)
dst[i] = hfloat(onnx_dtype::fp4ToF32(onnx_dtype::unpackNibble(src, i)));
}
else if (datatype == opencv_onnx::TensorProto_DataType_INT4)
{
checkPayloadSize(payload_bytes * 2); // two 4-bit elements per byte
blob.create((int)sizes.size(), sizes.data(), CV_8SC1);
const uchar* src = (const uchar*)rawdata;
const uchar* src = payload;
schar* dst = blob.ptr<schar>();
for (size_t i = 0, total = blob.total(); i < total; i++)
dst[i] = onnx_dtype::int4SignExtend(onnx_dtype::unpackNibble(src, i));
}
else if (datatype == opencv_onnx::TensorProto_DataType_UINT4)
{
checkPayloadSize(payload_bytes * 2); // two 4-bit elements per byte
blob.create((int)sizes.size(), sizes.data(), CV_8UC1);
const uchar* src = (const uchar*)rawdata;
const uchar* src = payload;
uchar* dst = blob.ptr<uchar>();
for (size_t i = 0, total = blob.total(); i < total; i++)
dst[i] = onnx_dtype::unpackNibble(src, i);
modules/dnn/src/onnx/onnx_importer2.cpp
+31 -5
View File
@@ -109,6 +109,7 @@ protected:
void parseNode(const opencv_onnx::NodeProto& node_proto);
bool parseValueInfo(const opencv_onnx::ValueInfoProto& valueInfoProto, ArgData& data);
int findGraphTensorOnnxType(const std::string& name) const;
void rememberProducedOnnxType(const opencv_onnx::NodeProto& node_proto, int onnx_type);
Mat parseTensor(const opencv_onnx::TensorProto& tensorProto);
void rememberMissingOp(const std::string& opname);
@@ -160,6 +161,9 @@ protected:
std::unordered_map<std::string, const opencv_onnx::NodeProto*> const_producers;
std::vector<Ptr<LayerInfo> > curr_prog;
std::vector<Arg> node_inputs, node_outputs;
// Raw ONNX dtype per produced tensor, keyed by remapped name: exports often omit
// value_info for intermediates, and FP8 formats sharing one depth need it.
std::unordered_map<std::string, int> produced_onnx_type;
std::string framework_name;
std::set<std::string> missing_ops;
@@ -310,7 +314,7 @@ protected:
const opencv_onnx::NodeProto& node_proto, int axis = -1);
void addQuantize(const std::string& name, const Arg& data,
const std::vector<Arg>& scale_zp, const std::vector<Arg>& outputs,
const opencv_onnx::NodeProto& node_proto);
const opencv_onnx::NodeProto& node_proto, int output_onnx_dtype = -1);
std::map<std::string, int> onnx_opset_map; // map from OperatorSetIdProto
void parseOperatorSet();
@@ -1715,9 +1719,17 @@ void ONNXImporter2::parseShape(LayerParams& layerParams, const opencv_onnx::Node
void ONNXImporter2::parseCast2(LayerParams& layerParams, const opencv_onnx::NodeProto& node_proto)
{
layerParams.type = "Cast2";
rememberProducedOnnxType(node_proto, layerParams.get<int>("to", -1));
addLayer(layerParams, node_proto);
}
// First output only - the FP8-producing ops (Cast/CastLike/QuantizeLinear) have just one.
void ONNXImporter2::rememberProducedOnnxType(const opencv_onnx::NodeProto& node_proto, int onnx_type)
{
if (onnx_type > 0 && node_proto.output_size() > 0 && !node_proto.output(0).empty())
produced_onnx_type[remap(node_proto.output(0))] = onnx_type;
}
// Returns a graph tensor's ONNX data_type by name, or -1 if unknown.
int ONNXImporter2::findGraphTensorOnnxType(const std::string& name) const
{
@@ -1736,7 +1748,8 @@ int ONNXImporter2::findGraphTensorOnnxType(const std::string& name) const
for (int i = 0; i < g.initializer_size(); i++)
if (g.initializer(i).name() == name)
return g.initializer(i).data_type();
return -1;
auto it = produced_onnx_type.find(remap(name));
return it != produced_onnx_type.end() ? it->second : -1;
}
void ONNXImporter2::parseCastLike(LayerParams& layerParams, const opencv_onnx::NodeProto& node_proto)
@@ -1749,6 +1762,7 @@ void ONNXImporter2::parseCastLike(LayerParams& layerParams, const opencv_onnx::N
if (elemType > 0)
layerParams.set("to", elemType);
}
rememberProducedOnnxType(node_proto, layerParams.get<int>("to", -1));
addLayer(layerParams, node_proto);
}
@@ -2442,11 +2456,19 @@ void ONNXImporter2::parseDequantizeLinear(LayerParams& layerParams, const opencv
void ONNXImporter2::parseQuantizeLinear(LayerParams& layerParams, const opencv_onnx::NodeProto& node_proto)
{
// Raw dtype disambiguates E5M2 vs E5M2FNUZ, which both collapse to CV_16F.
int dt = layerParams.get<int>("output_dtype", -1);
if (dt < 0 && node_proto.input_size() >= 3)
dt = findGraphTensorOnnxType(node_proto.input(2));
if (dt >= 0)
layerParams.set<int>("output_onnx_dtype", dt);
int attrDt = layerParams.get<int>("output_dtype", -1);
if (attrDt >= 0)
{
layerParams.set<int>("output_dtype", dataType2cv((opencv_onnx::TensorProto_DataType)dt));
layerParams.set<int>("output_dtype", dataType2cv((opencv_onnx::TensorProto_DataType)attrDt));
}
rememberProducedOnnxType(node_proto, dt);
addLayer(layerParams, node_proto);
}
@@ -2500,11 +2522,14 @@ Arg ONNXImporter2::addDequantize(const std::string& name, const std::vector<Arg>
void ONNXImporter2::addQuantize(const std::string& name, const Arg& data,
const std::vector<Arg>& scale_zp,
const std::vector<Arg>& outputs,
const opencv_onnx::NodeProto& node_proto)
const opencv_onnx::NodeProto& node_proto,
int output_onnx_dtype)
{
LayerParams lp;
lp.name = name;
lp.type = "QuantizeLinear";
if (output_onnx_dtype >= 0)
lp.set<int>("output_onnx_dtype", output_onnx_dtype);
node_inputs = {data};
node_inputs.insert(node_inputs.end(), scale_zp.begin(), scale_zp.end());
node_outputs = outputs;
@@ -2593,7 +2618,8 @@ void ONNXImporter2::parseQMatMul(LayerParams& layerParams, const opencv_onnx::No
node_outputs = {mm_out};
addLayer(mmLp, node_proto);
addQuantize(bn + "/quant_y", mm_out, {inp[6], inp[7]}, out, node_proto);
addQuantize(bn + "/quant_y", mm_out, {inp[6], inp[7]}, out, node_proto,
findGraphTensorOnnxType(node_proto.input(7)));
}
void ONNXImporter2::parseQGemm(LayerParams& layerParams, const opencv_onnx::NodeProto& node_proto)
modules/dnn/test/test_onnx_conformance_layer_filter__openvino.inl.hpp
+12
View File
@@ -804,6 +804,14 @@ CASE(test_dequantizelinear_axis)
SKIP;
CASE(test_dequantizelinear_blocked)
SKIP;
CASE(test_dequantizelinear_e4m3fn)
SKIP;
CASE(test_dequantizelinear_e4m3fn_float16)
SKIP;
CASE(test_dequantizelinear_e4m3fn_zero_point)
SKIP;
CASE(test_dequantizelinear_e5m2)
SKIP;
CASE(test_det_2d)
SKIP;
CASE(test_det_nd)
@@ -2107,6 +2115,10 @@ CASE(test_quantizelinear_axis)
SKIP;
CASE(test_quantizelinear_blocked)
SKIP;
CASE(test_quantizelinear_e4m3fn)
SKIP;
CASE(test_quantizelinear_e5m2)
SKIP;
CASE(test_range_float_type_positive_delta)
SKIP;
CASE(test_range_float_type_positive_delta_expanded)
modules/dnn/test/test_onnx_conformance_layer_parser_denylist.inl.hpp
-6
View File
@@ -87,10 +87,6 @@
"test_convinteger_without_padding", //Issues::Layer::Can't create layer "onnx_node_output_0!y" of type "ConvInteger" in function 'getLayerInstance'
"test_deform_conv_with_mask_bias",
"test_deform_conv_with_multiple_offset_groups",
"test_dequantizelinear_e4m3fn",
"test_dequantizelinear_e4m3fn_float16",
"test_dequantizelinear_e4m3fn_zero_point",
"test_dequantizelinear_e5m2",
"test_dequantizelinear_float4e2m1",
"test_dequantizelinear_int16",
"test_dequantizelinear_int4",
@@ -139,8 +135,6 @@
"test_qlinearmatmul_3D_uint8_float32",
"test_quantizelinear_blocked_asymmetric",
"test_quantizelinear_blocked_symmetric",
"test_quantizelinear_e4m3fn",
"test_quantizelinear_e5m2",
"test_quantizelinear_float4e2m1",
"test_quantizelinear_int16",
"test_quantizelinear_int4",