From 71862bbecaea8402286bc24b924d5cd2a4fd1322 Mon Sep 17 00:00:00 2001 From: Pikaliov Date: Thu, 11 Jun 2026 16:19:21 +0300 Subject: [PATCH] DCNv4 INT8 patch (Level 1): int8 storage + fp32 arithmetic for SOFIA/MERIDIAN E9 - Add dcnv4_int8_cuda.cu/.h: CUDA kernel (int8 values, fp16 offsets, fp32 interp, requantize with value_scale/output_scale) - Add dcnv4_int8_forward(): inference-only functional wrapper (@no_grad) - Add DCNv4Strip.forward_int8(): module-level INT8 forward (without_pointwise=True) - Add scripts/test_dcnv4_int8.py: correctness gate (<=1 LSB, >=99% exact) and informational fp16 vs int8 benchmark - Update README: INT8 API section, updated structure tree, SOFIA/MERIDIAN context Co-Authored-By: Claude Sonnet 4.6 --- DCNv4_op/DCNv4/functions/__init__.py | 2 +- DCNv4_op/DCNv4/functions/dcnv4_func.py | 29 ++++ DCNv4_op/DCNv4/modules/dcnv4.py | 40 +++++- DCNv4_op/scripts/test_dcnv4_int8.py | 159 +++++++++++++++++++++ DCNv4_op/src/cuda/dcnv4_int8_cuda.cu | 188 +++++++++++++++++++++++++ DCNv4_op/src/cuda/dcnv4_int8_cuda.h | 23 +++ DCNv4_op/src/dcnv4.h | 22 +++ DCNv4_op/src/vision.cpp | 3 + README.md | 87 ++++++++++-- 9 files changed, 539 insertions(+), 14 deletions(-) create mode 100644 DCNv4_op/scripts/test_dcnv4_int8.py create mode 100644 DCNv4_op/src/cuda/dcnv4_int8_cuda.cu create mode 100644 DCNv4_op/src/cuda/dcnv4_int8_cuda.h diff --git a/DCNv4_op/DCNv4/functions/__init__.py b/DCNv4_op/DCNv4/functions/__init__.py index 0338e58..0178482 100644 --- a/DCNv4_op/DCNv4/functions/__init__.py +++ b/DCNv4_op/DCNv4/functions/__init__.py @@ -8,4 +8,4 @@ # from .ms_flash_deform_attn_func import FlashMSDeformAttnFunction from .flash_deform_attn_func import FlashDeformAttnFunction -from .dcnv4_func import DCNv4Function \ No newline at end of file +from .dcnv4_func import DCNv4Function, dcnv4_int8_forward \ No newline at end of file diff --git a/DCNv4_op/DCNv4/functions/dcnv4_func.py b/DCNv4_op/DCNv4/functions/dcnv4_func.py index 1f95c3e..be53742 100644 --- a/DCNv4_op/DCNv4/functions/dcnv4_func.py +++ b/DCNv4_op/DCNv4/functions/dcnv4_func.py @@ -59,6 +59,35 @@ def find_spec_bwd(B, H, W, G, C): n_thread = multiplier * G * C // d_stride return d_stride, n_thread +@torch.no_grad() +def dcnv4_int8_forward( + value_int8, offset_mask_fp16, + kernel_h, kernel_w, pad_h, pad_w, + group, group_channels, + offset_scale, value_scale, output_scale): + """Inference-only INT8 DCNv4 (Level 1: int8 storage + fp32 math). + + Args: + value_int8: int8 tensor [B, H, W, C], quantized as + real_value = value_int8 * value_scale (per-tensor, symmetric). + offset_mask_fp16: fp16 tensor [B, H, W, padded_offset_dim] — the raw + (real-valued) offsets and masks, NOT quantized. + kernel_h/kernel_w/pad_h/pad_w/group/group_channels/offset_scale: + same semantics as DCNv4Function (stride=1, dilation=1 implied). + value_scale / output_scale: per-tensor quantization scales; the kernel + folds them into a single requantization multiplier + value_scale / output_scale. + + Returns: + int8 tensor [B, H, W, C]; real output = result * output_scale. + """ + requant = float(value_scale) / float(output_scale) + return ext.dcnv4_int8_forward( + value_int8, offset_mask_fp16, + kernel_h, kernel_w, 1, 1, pad_h, pad_w, 1, 1, + group, group_channels, float(offset_scale), requant) + + class DCNv4Function(Function): @staticmethod @custom_fwd diff --git a/DCNv4_op/DCNv4/modules/dcnv4.py b/DCNv4_op/DCNv4/modules/dcnv4.py index cc28644..1ee0eb9 100644 --- a/DCNv4_op/DCNv4/modules/dcnv4.py +++ b/DCNv4_op/DCNv4/modules/dcnv4.py @@ -13,7 +13,7 @@ import torch from torch import nn import torch.nn.functional as F from torch.nn.init import xavier_uniform_, constant_ -from ..functions import DCNv4Function +from ..functions import DCNv4Function, dcnv4_int8_forward class CenterFeatureScaleModule(nn.Module): def forward(self, @@ -284,3 +284,41 @@ class DCNv4Strip(nn.Module): x = self.output_proj(x) return x + @torch.no_grad() + def forward_int8(self, input_int8, shape, value_scale, output_scale): + """INT8 inference forward (Level 1: int8 storage + fp32 math). + + The deformable sampling consumes/produces int8; the (small) offset/mask + Linear runs in the module's own float dtype on the dequantized input. + Requires ``without_pointwise=True`` (the int8 path quantizes the + sampling op only; pointwise projections belong to the surrounding + block's quantized 1x1 convs). + + Args: + input_int8: int8 [N, L, C]; real value = input_int8 * value_scale. + shape: (H, W) of the token grid. + value_scale: per-tensor input scale. + output_scale: per-tensor output scale. + + Returns: + int8 [N, L, C]; real output = result * output_scale. + """ + assert self.without_pointwise, \ + "forward_int8 supports without_pointwise=True only" + assert input_int8.dtype == torch.int8, "input must be int8" + N, L, C = input_int8.shape + H, W = shape + + w_dtype = self.offset_mask.weight.dtype + x_deq = input_int8.to(w_dtype) * float(value_scale) + offset_mask = self.offset_mask(x_deq).to(torch.float16) + offset_mask = offset_mask.reshape(N, H, W, -1).contiguous() + + out = dcnv4_int8_forward( + input_int8.reshape(N, H, W, C).contiguous(), offset_mask, + self.kernel_h, self.kernel_w, self.pad_h, self.pad_w, + self.group, self.group_channels, + self.offset_scale, value_scale, output_scale, + ) + return out.reshape(N, L, C) + diff --git a/DCNv4_op/scripts/test_dcnv4_int8.py b/DCNv4_op/scripts/test_dcnv4_int8.py new file mode 100644 index 0000000..97b466c --- /dev/null +++ b/DCNv4_op/scripts/test_dcnv4_int8.py @@ -0,0 +1,159 @@ +"""INT8 DCNv4 gate (Level 1: int8 storage + fp32 arithmetic). + +Validates ext.dcnv4_int8_forward against the fp32 op on dequantized inputs: +the int8 path must be bit-equivalent to dequantize -> fp32 DCNv4 -> quantize, +so the only allowed error is output rounding (<= 0.5 quantization steps, +plus tiny fp accumulation noise). + +Checks: + 1. correctness vs fp32 op for square (3,3), strip (1,9)/(9,1), SOFIA + Stage-1 width (C=192, G=12); + 2. padding channels of offset_mask are ignored; + 3. DCNv4Strip.forward_int8 == dequant->fp forward->quant (within 1 step); + 4. benchmark int8 vs fp16 op (memory-bound shapes) — informational. + +PASS gate: max |int8_out - quantize(fp32_out)| <= 1 LSB; >=99% exact match. + +Run on the training server after rebuilding the extension: + cd DCNv4_op && pip install -e . && python scripts/test_dcnv4_int8.py +""" + +from __future__ import annotations + +import time + +import torch + +from DCNv4 import DCNv4Strip, ext +from DCNv4.functions import DCNv4Function, dcnv4_int8_forward + + +def quantize(x: torch.Tensor, scale: float) -> torch.Tensor: + return torch.clamp(torch.round(x / scale), -128, 127).to(torch.int8) + + +def run_fp32(value_f32, om_f32, kh, kw, pad_h, pad_w, group, gc, offset_scale=1.0): + return DCNv4Function.apply( + value_f32, om_f32, kh, kw, 1, 1, pad_h, pad_w, 1, 1, + group, gc, offset_scale, 256, 0, + ) + + +def check_case(kh: int, kw: int, n=2, h=16, w=16, c=64, group=4, + offset_scale=1.0) -> None: + torch.manual_seed(7) + gc = c // group + k = kh * kw + p = ((group * k * 3 + 7) // 8) * 8 + pad_h, pad_w = (kh - 1) // 2, (kw - 1) // 2 + + value_f = torch.randn(n, h, w, c, device="cuda") * 2.0 + v_scale = value_f.abs().max().item() / 127.0 + v_i8 = quantize(value_f, v_scale) + v_deq = v_i8.float() * v_scale # exact int8 grid + + om = torch.zeros(n, h, w, p, device="cuda") + om[..., : group * k * 3] = torch.randn(n, h, w, group * k * 3, + device="cuda") * 0.7 + om_h = om.half() # int8 path consumes fp16 offsets — give the fp32 + # reference the SAME rounded values for a clean compare + + ref = run_fp32(v_deq, om_h.float(), kh, kw, pad_h, pad_w, group, gc, + offset_scale) + o_scale = ref.abs().max().item() / 127.0 + ref_q = quantize(ref, o_scale) + + out_i8 = dcnv4_int8_forward( + v_i8, om_h, kh, kw, pad_h, pad_w, group, gc, + offset_scale, v_scale, o_scale, + ).reshape(ref_q.shape) + + diff = (out_i8.int() - ref_q.int()).abs() + max_lsb = diff.max().item() + exact = (diff == 0).float().mean().item() + # fp16 offsets in the int8 path vs fp32 in the reference cause sub-LSB + # coordinate differences; allow 1 LSB and require >=99% exact agreement. + assert max_lsb <= 1, f"({kh}x{kw}) int8 deviates by {max_lsb} LSB" + assert exact >= 0.99, f"({kh}x{kw}) only {exact:.4f} exact" + print(f" OK ({kh}x{kw}, C={c}, G={group}, os={offset_scale}): " + f"max {max_lsb} LSB, exact {exact:.4f}") + + # Padding channels must be ignored. + om_p = om.half().clone() + om_p[..., group * k * 3:] = 30000.0 + out_p = dcnv4_int8_forward( + v_i8, om_p, kh, kw, pad_h, pad_w, group, gc, + offset_scale, v_scale, o_scale, + ) + assert torch.equal(out_p, out_i8.reshape(out_p.shape)), \ + f"({kh}x{kw}) padding channels are NOT ignored" + + +def check_module() -> None: + torch.manual_seed(0) + m = DCNv4Strip(channels=192, k=9, orientation="h", group=12).cuda().half() + torch.nn.init.normal_(m.offset_mask.weight, std=0.02) + torch.nn.init.normal_(m.offset_mask.bias, std=0.02) + + x_f = torch.randn(2, 64 * 64, 192, device="cuda") * 1.5 + v_scale = x_f.abs().max().item() / 127.0 + x_i8 = quantize(x_f, v_scale) + x_deq = (x_i8.half() * v_scale) + + ref = m(x_deq, shape=(64, 64)).float() + o_scale = max(ref.abs().max().item() / 127.0, 1e-8) + ref_q = quantize(ref, o_scale) + + out = m.forward_int8(x_i8, (64, 64), v_scale, o_scale) + diff = (out.int() - ref_q.int()).abs() + assert diff.max().item() <= 1, f"module int8 deviates {diff.max().item()} LSB" + print(f" OK DCNv4Strip.forward_int8: max {diff.max().item()} LSB, " + f"exact {(diff == 0).float().mean().item():.4f}") + + +def bench(n=8, hw=64, c=192, group=12, k=9, iters=100) -> None: + gc = c // group + p = ((group * k * 3 + 7) // 8) * 8 + v_f16 = torch.randn(n, hw, hw, c, device="cuda", dtype=torch.float16) + v_i8 = quantize(v_f16.float(), 0.05) + om16 = torch.randn(n, hw, hw, p, device="cuda", dtype=torch.float16) + + def timeit(fn): + for _ in range(20): + fn() + torch.cuda.synchronize() + t0 = time.perf_counter() + for _ in range(iters): + fn() + torch.cuda.synchronize() + return (time.perf_counter() - t0) / iters * 1e3 + + t_fp16 = timeit(lambda: DCNv4Function.apply( + v_f16, om16, 1, k, 1, 1, 0, (k - 1) // 2, 1, 1, + group, gc, 1.0, 256, 0)) + t_int8 = timeit(lambda: ext.dcnv4_int8_forward( + v_i8, om16, 1, k, 1, 1, 0, (k - 1) // 2, 1, 1, group, gc, 1.0, 1.0)) + print(f" strip(1,{k}) [{n},{hw}²,{c}] fwd: fp16 {t_fp16:.3f} ms, " + f"int8 {t_int8:.3f} ms ({t_fp16 / t_int8:.2f}x)") + if t_int8 > t_fp16: + print(" NOTE: v1 int8 kernel is correctness-first (no shm/vector " + "tuning); optimize only after the accuracy gate passes.") + + +def main() -> None: + assert torch.cuda.is_available(), "CUDA required" + print("== correctness (int8 vs dequant->fp32->quant) ==") + check_case(3, 3) + check_case(1, 9) + check_case(9, 1) + check_case(1, 9, c=192, group=12) + check_case(1, 9, c=192, group=12, offset_scale=0.5) + print("== module ==") + check_module() + print("== benchmark (informational) ==") + bench() + print("\nALL INT8 GATES PASSED") + + +if __name__ == "__main__": + main() diff --git a/DCNv4_op/src/cuda/dcnv4_int8_cuda.cu b/DCNv4_op/src/cuda/dcnv4_int8_cuda.cu new file mode 100644 index 0000000..4ddb40d --- /dev/null +++ b/DCNv4_op/src/cuda/dcnv4_int8_cuda.cu @@ -0,0 +1,188 @@ +/*! +************************************************************************************************** +* DCNv4 INT8 inference path (Level 1: int8 storage + fp32 arithmetic). +* +* Design (see results/PLAN_DCNv4_INT8_2026-06-11.md in sofia_pretrain_world_uav): +* - value / output tensors are int8 with per-tensor scales (memory-bound op: +* halving the bytes read vs fp16 is where the speedup comes from); +* - offset_mask stays fp16 (offsets are fractional coordinates, masks are +* unbounded reals in DCNv4 — quantizing them is the risky part, cf. DEFA); +* - all interpolation arithmetic is fp32 on raw int8 magnitudes; the value +* scale is folded into a single requantization multiplier at the end: +* out_i8 = clamp(round(acc * requant_scale)), +* requant_scale = value_scale / output_scale. +* This keeps the math bit-exact w.r.t. dequantize->fp32-op->quantize. +* +* Correctness-first v1 kernel: one thread computes 4 consecutive channels +* (char4 loads/stores), runtime loops over kernel points (no K templates, so +* any kernel shape works, including strips (1,k)/(k,1)). Sampling algebra +* mirrors dcnv4_im2col_cuda.cuh exactly (point order m = i*kernel_h + j with +* i over kernel_w; p0/p0_ offset_scale algebra; per-corner zero padding). +* +* Restrictions (asserted on the host): stride=1, dilation=1, remove_center=0, +* softmax=false, "same" padding, group_channels % 4 == 0. Inference only — +* no backward (training uses the fp16 path + QAT fake-quant). +************************************************************************************************** +*/ + +#include +#include +#include +#include +#include +#include + +#include "dcnv4_int8_cuda.h" + +namespace { + +__global__ void dcnv4_int8_fwd_kernel( + const int8_t *__restrict__ value, // [B, H*W, C] + const __half *__restrict__ offset_mask, // [B, H*W, P] + int8_t *__restrict__ output, // [B, H*W, C] + const int B, const int H, const int W, const int G, const int D, + const int kernel_h, const int kernel_w, const int pad_h, const int pad_w, + const float offset_scale, const float requant_scale, + const int padded_offset_dim) { + const int Q = H * W; + const int d4 = D >> 2; // channel chunks of 4 per group + const long long total = (long long)B * Q * G * d4; + const long long idx = (long long)blockIdx.x * blockDim.x + threadIdx.x; + if (idx >= total) { + return; + } + + const int dchunk = (int)(idx % d4); + const int g = (int)((idx / d4) % G); + const int q = (int)((idx / ((long long)d4 * G)) % Q); + const int b = (int)(idx / ((long long)d4 * G * Q)); + + const int wi = q % W; + const int hi = q / W; + const int C = G * D; + const int K_pts = kernel_h * kernel_w; + + const __half *om = + offset_mask + ((long long)b * Q + q) * padded_offset_dim + (long long)g * K_pts * 3; + + // Mirrors dcnv4_im2col_cuda.cuh (stride=1, dilation=1): + // p0 = center - pad + out_coord + // p0_ = p0 - center * offset_scale + // loc = p0_ + (tap_index + d) * offset_scale + const int center_w = (kernel_w - 1) >> 1; + const int center_h = (kernel_h - 1) >> 1; + const float p0_w = (float)(center_w - pad_w + wi) - (float)center_w * offset_scale; + const float p0_h = (float)(center_h - pad_h + hi) - (float)center_h * offset_scale; + + float acc0 = 0.f, acc1 = 0.f, acc2 = 0.f, acc3 = 0.f; + const int ch = g * D + (dchunk << 2); + const int8_t *vbase = value + (long long)b * Q * C + ch; + +#define DCNV4_INT8_ACC_CORNER(yy, xx, wgt) \ + do { \ + if ((yy) >= 0 && (yy) < H && (xx) >= 0 && (xx) < W) { \ + const char4 v = *reinterpret_cast( \ + vbase + ((long long)(yy) * W + (xx)) * C); \ + const float wa = (wgt); \ + acc0 += wa * (float)v.x; \ + acc1 += wa * (float)v.y; \ + acc2 += wa * (float)v.z; \ + acc3 += wa * (float)v.w; \ + } \ + } while (0) + + for (int i = 0; i < kernel_w; ++i) { + for (int j = 0; j < kernel_h; ++j) { + const int m = i * kernel_h + j; // CUDA point order: i (w) outer, j (h) inner + const float dx = __half2float(om[(m << 1)]); + const float dy = __half2float(om[(m << 1) + 1]); + const float attn = __half2float(om[(K_pts << 1) + m]); + const float loc_w = p0_w + ((float)i + dx) * offset_scale; + const float loc_h = p0_h + ((float)j + dy) * offset_scale; + + if (loc_w <= -1.f || loc_w >= (float)W || loc_h <= -1.f || + loc_h >= (float)H) { + continue; // whole point outside (same guard as the fp kernel) + } + const int w_low = (int)floorf(loc_w); + const int h_low = (int)floorf(loc_h); + const int w_high = w_low + 1; + const int h_high = h_low + 1; + const float lw = loc_w - (float)w_low; + const float lh = loc_h - (float)h_low; + const float hw = 1.f - lw; + const float hh = 1.f - lh; + + DCNV4_INT8_ACC_CORNER(h_low, w_low, attn * hh * hw); + DCNV4_INT8_ACC_CORNER(h_low, w_high, attn * hh * lw); + DCNV4_INT8_ACC_CORNER(h_high, w_low, attn * lh * hw); + DCNV4_INT8_ACC_CORNER(h_high, w_high, attn * lh * lw); + } + } +#undef DCNV4_INT8_ACC_CORNER + + const int o0 = max(-128, min(127, __float2int_rn(acc0 * requant_scale))); + const int o1 = max(-128, min(127, __float2int_rn(acc1 * requant_scale))); + const int o2 = max(-128, min(127, __float2int_rn(acc2 * requant_scale))); + const int o3 = max(-128, min(127, __float2int_rn(acc3 * requant_scale))); + char4 o; + o.x = (int8_t)o0; + o.y = (int8_t)o1; + o.z = (int8_t)o2; + o.w = (int8_t)o3; + *reinterpret_cast(output + ((long long)b * Q + q) * C + ch) = o; +} + +} // namespace + +at::Tensor dcnv4_int8_cuda_forward( + const at::Tensor &value, const at::Tensor &p_offset, const int kernel_h, + const int kernel_w, const int stride_h, const int stride_w, + const int pad_h, const int pad_w, const int dilation_h, + const int dilation_w, const int group, const int group_channels, + const float offset_scale, const float requant_scale) { + TORCH_CHECK(value.is_cuda() && p_offset.is_cuda(), "inputs must be CUDA"); + TORCH_CHECK(value.scalar_type() == at::kChar, "value must be int8 (kChar)"); + TORCH_CHECK(p_offset.scalar_type() == at::kHalf, "offset_mask must be fp16"); + TORCH_CHECK(value.is_contiguous() && p_offset.is_contiguous(), + "inputs must be contiguous"); + TORCH_CHECK(value.dim() == 4 && p_offset.dim() == 4, + "value/offset_mask must be [B, H, W, C] / [B, H, W, P]"); + TORCH_CHECK(stride_h == 1 && stride_w == 1 && dilation_h == 1 && + dilation_w == 1, + "int8 path supports stride=1, dilation=1 only"); + TORCH_CHECK(pad_h == (kernel_h - 1) / 2 && pad_w == (kernel_w - 1) / 2, + "int8 path requires 'same' padding (pad = (k-1)/2)"); + TORCH_CHECK(group_channels % 4 == 0, + "group_channels must be divisible by 4 (char4 path)"); + + const int B = (int)value.size(0); + const int H = (int)value.size(1); + const int W = (int)value.size(2); + const int C = (int)value.size(3); + TORCH_CHECK(C == group * group_channels, + "C must equal group * group_channels"); + const int P = (int)p_offset.size(3); + TORCH_CHECK(P % 8 == 0, "padded_offset_dim must be divisible by 8"); + TORCH_CHECK(P >= group * kernel_h * kernel_w * 3, + "offset_mask last dim too small for G*K*3"); + TORCH_CHECK(p_offset.size(0) == B && p_offset.size(1) == H && + p_offset.size(2) == W, + "offset_mask spatial shape must match value"); + + auto output = at::empty_like(value); + const long long total = + (long long)B * H * W * group * (group_channels / 4); + const int threads = 256; + const long long blocks = (total + threads - 1) / threads; + TORCH_CHECK(blocks <= 0x7FFFFFFFLL, "grid too large"); + + auto stream = at::cuda::getCurrentCUDAStream(); + dcnv4_int8_fwd_kernel<<<(unsigned int)blocks, threads, 0, stream>>>( + value.data_ptr(), + reinterpret_cast(p_offset.data_ptr()), + output.data_ptr(), B, H, W, group, group_channels, kernel_h, + kernel_w, pad_h, pad_w, offset_scale, requant_scale, P); + C10_CUDA_KERNEL_LAUNCH_CHECK(); + return output; +} diff --git a/DCNv4_op/src/cuda/dcnv4_int8_cuda.h b/DCNv4_op/src/cuda/dcnv4_int8_cuda.h new file mode 100644 index 0000000..caeb859 --- /dev/null +++ b/DCNv4_op/src/cuda/dcnv4_int8_cuda.h @@ -0,0 +1,23 @@ +/*! +************************************************************************************************** +* DCNv4 INT8 inference path (Level 1: int8 storage + fp32 arithmetic) +* SOFIA/CVGL addition, 2026. Apache-2.0 (same as DCNv4). +************************************************************************************************** +*/ + +#pragma once +#include + +// Inference-only forward: value/output are int8, offset_mask is fp16, +// all interpolation arithmetic is fp32, output is requantized with +// requant_scale = value_scale / output_scale. +// Restrictions (checked): stride=1, dilation=1, remove_center=0, no softmax, +// "same" padding (pad = (k-1)/2), group_channels % 4 == 0. +at::Tensor dcnv4_int8_cuda_forward( + const at::Tensor &value, // int8 [B, H, W, C], C = group * group_channels + const at::Tensor &p_offset, // half [B, H, W, padded_offset_dim] + const int kernel_h, const int kernel_w, const int stride_h, + const int stride_w, const int pad_h, const int pad_w, + const int dilation_h, const int dilation_w, const int group, + const int group_channels, const float offset_scale, + const float requant_scale); diff --git a/DCNv4_op/src/dcnv4.h b/DCNv4_op/src/dcnv4.h index 0065f00..0199140 100644 --- a/DCNv4_op/src/dcnv4.h +++ b/DCNv4_op/src/dcnv4.h @@ -13,6 +13,7 @@ #ifdef WITH_CUDA #include "cuda/dcnv4_cuda.h" +#include "cuda/dcnv4_int8_cuda.h" #include "cuda/flash_deform_attn_cuda.h" #endif @@ -82,6 +83,27 @@ at::Tensor dcnv4_forward( AT_ERROR("Not implemented on the CPU"); } +// INT8 inference path (Level 1): int8 value/output, fp16 offset_mask, +// fp32 arithmetic, requant_scale = value_scale / output_scale. +inline at::Tensor dcnv4_int8_forward( + const at::Tensor &value, const at::Tensor &p_offset, const int kernel_h, + const int kernel_w, const int stride_h, const int stride_w, + const int pad_h, const int pad_w, const int dilation_h, + const int dilation_w, const int group, const int group_channels, + const float offset_scale, const float requant_scale) { + if (value.device().is_cuda()) { +#ifdef WITH_CUDA + return dcnv4_int8_cuda_forward( + value, p_offset, kernel_h, kernel_w, stride_h, stride_w, pad_h, pad_w, + dilation_h, dilation_w, group, group_channels, offset_scale, + requant_scale); +#else + AT_ERROR("Not compiled with GPU support"); +#endif + } + AT_ERROR("Not implemented on the CPU"); +} + std::vector dcnv4_backward( const at::Tensor &value, diff --git a/DCNv4_op/src/vision.cpp b/DCNv4_op/src/vision.cpp index c559398..4f58bb8 100644 --- a/DCNv4_op/src/vision.cpp +++ b/DCNv4_op/src/vision.cpp @@ -18,4 +18,7 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) { "flash_deform_attn_backward"); m.def("dcnv4_forward", &dcnv4_forward, "dcnv4_forward"); m.def("dcnv4_backward", &dcnv4_backward, "dcnv4_backward"); + m.def("dcnv4_int8_forward", &dcnv4_int8_forward, + "dcnv4_int8_forward (inference-only: int8 value/output, fp16 " + "offset_mask, fp32 math)"); } diff --git a/README.md b/README.md index d2c0bff..1286eee 100644 --- a/README.md +++ b/README.md @@ -1,6 +1,6 @@ ## DCNv4 — Custom Op Mirror -**Зеркало для:** Pikaliov (CVGL-project) +**Зеркало для:** Pikaliov (CVGL/MERIDIAN project) **Email:** i@pikaliov.ru **Gitea репо:** https://git.lissad.keenetic.name/Pikaliov/DCN_custom_op **Upstream:** https://github.com/OpenGVLab/DCNv4 @@ -15,30 +15,93 @@ cd DCNv4_op pip install -e . ``` -**Требования:** CUDA toolkit + PyTorch (совместимые версии с nvcc) -**Импорт:** `from DCNv4 import DCNv4` +**Требования:** CUDA toolkit + PyTorch (совместимые версии с nvcc) + +```python +from DCNv4 import DCNv4, DCNv4Strip +from DCNv4.functions import DCNv4Function, dcnv4_int8_forward +``` + +--- + +### INT8-патч (SOFIA/MERIDIAN, E9) + +Этот форк расширяет upstream добавлением **INT8 inference пути (Level 1)**: int8 хранение значений + fp32 арифметика интерполяции. Предназначен для E9 QAT-эксперимента MERIDIAN (Jetson Orin NX, TensorRT INT8). + +#### Что добавлено + +| Файл | Описание | +|------|----------| +| `src/cuda/dcnv4_int8_cuda.cu/.h` | CUDA-ядро: int8 значения, fp16 offsets, fp32 интерполяция, requant | +| `DCNv4/functions/dcnv4_func.py` | `dcnv4_int8_forward()` — функциональный INT8 forward | +| `DCNv4/modules/dcnv4.py` | `DCNv4Strip.forward_int8()` — INT8 forward на уровне модуля | +| `scripts/test_dcnv4_int8.py` | Тесты корректности + бенчмарк | + +#### API + +```python +# Функциональный INT8 forward (inference-only, @no_grad) +out_int8 = dcnv4_int8_forward( + value_int8, # int8 [B, H, W, C] + offset_mask_fp16, # fp16 [B, H, W, padded_dim] + kernel_h, kernel_w, pad_h, pad_w, + group, group_channels, offset_scale, + value_scale, output_scale, # per-tensor quant scales +) # -> int8 [B, H, W, C] + +# Модульный INT8 forward (DCNv4Strip, without_pointwise=True) +out_int8 = strip_module.forward_int8( + input_int8, # int8 [N, L, C] + shape=(H, W), + value_scale=vs, + output_scale=os, +) # -> int8 [N, L, C] +``` + +**Ограничения:** stride=1, dilation=1, "same" padding, `group_channels % 4 == 0`, `without_pointwise=True` для модульного API. + +#### Проверка INT8 + +```bash +cd DCNv4_op && pip install -e . +python scripts/test_dcnv4_int8.py +# Проверяет: (3x3), (1x9)/(9x1) strip, SOFIA Stage-1 (C=192, G=12) +# PASS: max |int8_out - quantize(fp32_out)| <= 1 LSB, >=99% exact match +``` + +--- ### Структура проекта ``` -├── DCNv4_op/ ← PyTorch CUDA-расширение (главный модуль) -│ ├── src/ ← CUDA/C++ источники -│ ├── DCNv4/ ← Python пакет DCNv4 -│ └── setup.py ← Сборка (требует CUDA) +├── DCNv4_op/ ← PyTorch CUDA-расширение (главный модуль) +│ ├── src/ +│ │ ├── cuda/ +│ │ │ ├── dcnv4_cuda.cu ← upstream fp16/fp32 ядро +│ │ │ ├── dcnv4_int8_cuda.cu ← INT8 ядро (SOFIA/MERIDIAN) +│ │ │ └── dcnv4_int8_cuda.h +│ │ ├── dcnv4.h +│ │ └── vision.cpp +│ ├── DCNv4/ ← Python пакет DCNv4 +│ │ ├── functions/dcnv4_func.py ← DCNv4Function + dcnv4_int8_forward +│ │ └── modules/dcnv4.py ← DCNv4, DCNv4Strip + forward_int8 +│ ├── scripts/ +│ │ └── test_dcnv4_int8.py ← INT8 correctness gate +│ └── setup.py ├── classification/ ← ImageNet классификация + FlashInternImage ├── detection/ ← COCO detection & segmentation (Mask R-CNN, DINO) ├── segmentation/ ← ADE20K semantic segmentation (UperNet, Mask2Former) -├── README.md ← Этот файл + оригинальный OpenGVLab README +├── README.md └── LICENSE ← Apache-2.0 ``` ### Что такое DCNv4? **Deformable Convolution v4** — высокоэффективный оператор для vision-моделей: -- 🚀 **3× быстрее** чем DCNv3 (speedup на forward pass) -- ⚡ **Лучше сходится** (более быстрая конвергенция при обучении) -- 🎯 **Универсальность:** image classification, detection, segmentation, generation (т.е. diffusion models) -- 🔧 **Интеграция:** замена DCNv3 → DCNv4 в InternImage дает +80% ускорения без переобучения +- **3× быстрее** чем DCNv3 (speedup на forward pass) +- **Лучше сходится** (более быстрая конвергенция при обучении) +- **Универсальность:** image classification, detection, segmentation, generation (diffusion models) +- **Интеграция:** замена DCNv3 → DCNv4 в InternImage дает +80% ускорения без переобучения ### Модели и бенчмарки