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 <noreply@anthropic.com>
This commit is contained in:
@@ -8,4 +8,4 @@
|
||||
|
||||
# from .ms_flash_deform_attn_func import FlashMSDeformAttnFunction
|
||||
from .flash_deform_attn_func import FlashDeformAttnFunction
|
||||
from .dcnv4_func import DCNv4Function
|
||||
from .dcnv4_func import DCNv4Function, dcnv4_int8_forward
|
||||
@@ -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
|
||||
|
||||
@@ -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)
|
||||
|
||||
|
||||
159
DCNv4_op/scripts/test_dcnv4_int8.py
Normal file
159
DCNv4_op/scripts/test_dcnv4_int8.py
Normal file
@@ -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()
|
||||
188
DCNv4_op/src/cuda/dcnv4_int8_cuda.cu
Normal file
188
DCNv4_op/src/cuda/dcnv4_int8_cuda.cu
Normal file
@@ -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 <ATen/cuda/CUDAContext.h>
|
||||
#include <c10/cuda/CUDAException.h>
|
||||
#include <cuda.h>
|
||||
#include <cuda_fp16.h>
|
||||
#include <cuda_runtime.h>
|
||||
#include <torch/extension.h>
|
||||
|
||||
#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<const char4 *>( \
|
||||
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<char4 *>(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<int8_t>(),
|
||||
reinterpret_cast<const __half *>(p_offset.data_ptr<at::Half>()),
|
||||
output.data_ptr<int8_t>(), 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;
|
||||
}
|
||||
23
DCNv4_op/src/cuda/dcnv4_int8_cuda.h
Normal file
23
DCNv4_op/src/cuda/dcnv4_int8_cuda.h
Normal file
@@ -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 <torch/extension.h>
|
||||
|
||||
// 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);
|
||||
@@ -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<at::Tensor>
|
||||
dcnv4_backward(
|
||||
const at::Tensor &value,
|
||||
|
||||
@@ -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)");
|
||||
}
|
||||
|
||||
87
README.md
87
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% ускорения без переобучения
|
||||
|
||||
### Модели и бенчмарки
|
||||
|
||||
|
||||
Reference in New Issue
Block a user