orisys-tech/orisys-sdk
3
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

tune for dq6a to reach 30 fps for 2 sensors

Browse Source
This commit is contained in:
radxa-d6a
2026-04-02 16:44:02 +08:00
parent 51c259e523
commit 80e951d6cd
15 changed files with 474 additions and 70 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/orisys/HHD.py
+35 -17
View File
@@ -1626,7 +1626,8 @@ class PoissonSolver(object):
def _get_opencl_batch_mul_kernel(self, queue):
"""Batched complex multiply: multiplies each batch plane by the same Green's spectrum.
Uses multiple elements per work item and explicit local size for better throughput."""
Uses multiple elements per work item, __constant for broadcast-read Green's spectrum,
and fma for better throughput/precision."""
if self._opencl_batch_mul_kernel is not None:
return self._opencl_batch_mul_kernel
@@ -1637,13 +1638,14 @@ class PoissonSolver(object):
const int n
) {
const int b = get_global_id(0);
const int base = b * n;
const int i0 = get_global_id(1) * 4;
for (int i = i0; i < i0 + 4 && i < n; i++) {
const int i_end = min(i0 + 4, n);
for (int i = i0; i < i_end; i++) {
float2 gv = g[i];
int idx = b * n + i;
float2 a = fft_batch[idx];
fft_batch[idx] = (float2)(a.x * gv.x - a.y * gv.y,
a.x * gv.y + a.y * gv.x);
float2 a = fft_batch[base + i];
fft_batch[base + i] = (float2)(fma(a.x, gv.x, -(a.y * gv.y)),
fma(a.x, gv.y, a.y * gv.x));
}
}
"""
@@ -1752,10 +1754,13 @@ class PoissonSolver(object):
'src_plane': src_plane,
'dst_plane': dst_plane,
'f_batch_dev': cla.zeros(queue, (2,) + full_shape, dtype=numpy.float32),
'host_pad': numpy.zeros((2,) + full_shape, dtype=numpy.float32),
'host_stage': numpy.empty((2,) + fshape, dtype=numpy.float32),
'f_stage_dev': cla.empty(queue, (2,) + fshape, dtype=numpy.float32),
'f_fft_batch_dev': cla.empty(queue, (2,) + fft_shape, dtype=numpy.complex64),
'conv_batch_dev': cla.empty(queue, (2,) + full_shape, dtype=numpy.float32),
'p_batch_dev': cla.empty(queue, (2,) + fshape, dtype=numpy.float32, allocator=alloc_pinned),
'p1_host': numpy.empty(fshape, dtype=numpy.float32),
'p2_host': numpy.empty(fshape, dtype=numpy.float32),
}
timing['cache_prepare_ms'] = (timeit.default_timer() - t0) * 1000.0
@@ -1763,15 +1768,26 @@ class PoissonSolver(object):
fshape_c = c['fshape']
full_shape_c = c['full_shape']
# H2D: write into pre-zeroed host pad buffer, single enqueue_copy (no map/unmap/scatter)
# H2D: compact transfer of unpadded data, then scatter to padded buffer on device
t0 = timeit.default_timer()
host_pad = c['host_pad']
H, W = fshape_c
f1_f32 = numpy.ascontiguousarray(f1, dtype=numpy.float32)
f2_f32 = numpy.ascontiguousarray(f2, dtype=numpy.float32)
host_pad[0, :H, :W] = f1_f32
host_pad[1, :H, :W] = f2_f32
cl.enqueue_copy(queue, c['f_batch_dev'].data, host_pad, is_blocking=False)
f1_f32 = f1 if (f1.flags['C_CONTIGUOUS'] and f1.dtype == numpy.float32) else numpy.ascontiguousarray(f1, dtype=numpy.float32)
f2_f32 = f2 if (f2.flags['C_CONTIGUOUS'] and f2.dtype == numpy.float32) else numpy.ascontiguousarray(f2, dtype=numpy.float32)
host_stage = c['host_stage']
numpy.copyto(host_stage[0], f1_f32)
numpy.copyto(host_stage[1], f2_f32)
f_stage_dev = c['f_stage_dev']
cl.enqueue_copy(queue, f_stage_dev.data, host_stage, is_blocking=False)
c['f_batch_dev'].fill(0, queue=queue)
scatter_batch = self._get_opencl_scatter_batch_kernel(queue)
scatter_batch(
queue, (W, H, 2), None,
f_stage_dev.data, c['f_batch_dev'].data,
numpy.int32(W), numpy.int32(H),
numpy.int32(full_shape_c[1]),
numpy.int32(H * W),
numpy.int32(c['dst_plane']),
)
if profile:
queue.finish()
timing['h2d_set_ms'] = (timeit.default_timer() - t0) * 1000.0
@@ -1827,11 +1843,13 @@ class PoissonSolver(object):
queue.finish()
timing['crop_scale_ms'] = (timeit.default_timer() - t0) * 1000.0
# D2H: map pinned buffer (ALLOC_HOST_PTR) for read, copy out p1/p2, unmap
# D2H: map pinned buffer, copy into pre-allocated host arrays (zero allocation)
t0 = timeit.default_timer()
p1 = c['p1_host']
p2 = c['p2_host']
mapped_p = c['p_batch_dev'].map_to_host(queue=queue, flags=cl.map_flags.READ, is_blocking=True)
p1 = numpy.asarray(mapped_p[0], dtype=numpy.float32).copy()
p2 = numpy.asarray(mapped_p[1], dtype=numpy.float32).copy()
numpy.copyto(p1, mapped_p[0])
numpy.copyto(p2, mapped_p[1])
mapped_p.base.release(queue=queue)
timing['d2h_get_ms'] = (timeit.default_timer() - t0) * 1000.0
src/orisys/Stitcher.py
+79 -13
View File
@@ -406,7 +406,7 @@ class Stitcher:
wx[-m:] = ramp[::-1]
return wy[:, None] * wx[None, :]
def _blend_tile_inplace(self, blend, wsum, cov, src, weight, y0, x0):
def _blend_tile_inplace(self, blend, wsum, cov, src, weight, y0, x0, scratch=None):
"""Copy non-overlap tiles directly; blend only when this tile touches existing content."""
h, w = src.shape[:2]
H, W = blend.shape[:2]
@@ -426,6 +426,19 @@ class Stitcher:
if not np.any(cov_patch):
blend_patch[:] = src_patch
wsum_patch[:] = w_patch
elif scratch is not None:
s_src, s_alpha = scratch
sf = s_src[:sh, :sw]
sa = s_alpha[:sh, :sw]
sf[:] = src_patch
np.add(wsum_patch, w_patch, out=wsum_patch)
np.maximum(wsum_patch, 1e-6, out=sa)
np.divide(w_patch, sa, out=sa)
np.subtract(sf, blend_patch, out=sf)
sf[:, :, 0] *= sa
sf[:, :, 1] *= sa
sf[:, :, 2] *= sa
np.add(blend_patch, sf, out=blend_patch)
else:
total = wsum_patch + w_patch
alpha = w_patch / np.maximum(total, 1e-6)
@@ -436,13 +449,15 @@ class Stitcher:
cov_patch[:] = 1
def _mean_gray_bgr(self, bgr):
return float(cv2.cvtColor(bgr, cv2.COLOR_BGR2GRAY).mean())
ch = cv2.mean(bgr)
return (ch[0] + ch[1] + ch[2]) / 3.0
def _match_bgr_brightness(self, tile_bgr, ref_mean_gray):
delta = int(round(ref_mean_gray - self._mean_gray_bgr(tile_bgr)))
if delta == 0:
return tile_bgr
return np.clip(tile_bgr.astype(np.int16) + delta, 0, 255).astype(np.uint8)
lut = np.clip(np.arange(256, dtype=np.int16) + delta, 0, 255).astype(np.uint8)
return cv2.LUT(tile_bgr, lut)
def _largest_fully_covered_rect(self, cov):
if cov.size == 0:
@@ -519,13 +534,22 @@ class Stitcher:
do_profile = verbose and frame_idx % print_every == 0
t_profile0 = cv2.getTickCount() / cv2.getTickFrequency() if do_profile else 0.0
t_ref_bgr = self._tile_to_bgr(t_ref_bgr)
if t_ref_bgr.ndim == 2:
t_ref_bgr = cv2.cvtColor(t_ref_bgr, cv2.COLOR_GRAY2BGR)
h0c, w0c = t_ref_bgr.shape[:2]
if match_brightness:
ref_mean_gray = self._mean_gray_bgr(t_ref_bgr)
tiles_bgr_c = [self._match_bgr_brightness(self._tile_to_bgr(tb), ref_mean_gray) for tb in tiles_bgr]
tiles_bgr_c = [
self._match_bgr_brightness(
tb if tb.ndim == 3 else cv2.cvtColor(tb, cv2.COLOR_GRAY2BGR),
ref_mean_gray,
)
for tb in tiles_bgr
]
elif any(tb.ndim == 2 for tb in tiles_bgr):
tiles_bgr_c = [tb if tb.ndim == 3 else cv2.cvtColor(tb, cv2.COLOR_GRAY2BGR) for tb in tiles_bgr]
else:
tiles_bgr_c = [self._tile_to_bgr(tb) for tb in tiles_bgr]
tiles_bgr_c = tiles_bgr
t_after_brightness = cv2.getTickCount() / cv2.getTickFrequency() if do_profile else 0.0
geom_key = self._geom_cache_key(t_ref_bgr, tiles_bgr_c, pad)
@@ -563,6 +587,20 @@ class Stitcher:
tx, ty = self._translation_from_H_to_ref(H_use)
pastes.append((int(np.floor(ty - y_min_f)), int(np.floor(tx - x_min_f))))
_dr_y0 = max(0, ref_y0)
_dr_x0 = max(0, ref_x0)
_dr_y1 = min(out_h, ref_y0 + h0c)
_dr_x1 = min(out_w, ref_x0 + w0c)
for _tb, _pa in zip(tiles_bgr_c, pastes):
if _pa is None:
continue
_py, _px = _pa
_th, _tw = _tb.shape[:2]
_dr_y0 = min(_dr_y0, max(0, _py))
_dr_x0 = min(_dr_x0, max(0, _px))
_dr_y1 = max(_dr_y1, min(out_h, _py + _th))
_dr_x1 = max(_dr_x1, min(out_w, _px + _tw))
canvas_cache["geom_key"] = geom_key
canvas_cache["out_w"] = out_w
canvas_cache["out_h"] = out_h
@@ -572,6 +610,7 @@ class Stitcher:
canvas_cache["x_min_f"] = x_min_f
canvas_cache["y_min_f"] = y_min_f
canvas_cache["crop_rect"] = None
canvas_cache["dirty_rect"] = (_dr_y0, _dr_y1, _dr_x0, _dr_x1)
canvas_cache["blend_buf"] = np.zeros((out_h, out_w, 3), dtype=np.float32)
canvas_cache["wsum_buf"] = np.zeros((out_h, out_w), dtype=np.float32)
canvas_cache["cov_buf"] = np.zeros((out_h, out_w), dtype=np.uint8)
@@ -588,12 +627,32 @@ class Stitcher:
wsum = canvas_cache["wsum_buf"]
cov = canvas_cache["cov_buf"]
weight_maps = canvas_cache["weight_maps"]
blend_buf.fill(0)
wsum.fill(0)
cov.fill(0)
if geom_cache_hit:
_dr = canvas_cache.get("dirty_rect")
if _dr is not None:
_dy0, _dy1, _dx0, _dx1 = _dr
blend_buf[_dy0:_dy1, _dx0:_dx1] = 0
wsum[_dy0:_dy1, _dx0:_dx1] = 0
cov[_dy0:_dy1, _dx0:_dx1] = 0
else:
blend_buf.fill(0)
wsum.fill(0)
cov.fill(0)
t_after_reset = cv2.getTickCount() / cv2.getTickFrequency() if do_profile else 0.0
blend_margin = 16 if blend else 0
max_th = max(h0c, *(tb.shape[0] for tb in tiles_bgr_c)) if tiles_bgr_c else h0c
max_tw = max(w0c, *(tb.shape[1] for tb in tiles_bgr_c)) if tiles_bgr_c else w0c
_ss = canvas_cache.get("scratch_src")
_sa = canvas_cache.get("scratch_alpha")
if _ss is None or _ss.shape[0] < max_th or _ss.shape[1] < max_tw:
_ss = np.empty((max_th, max_tw, 3), dtype=np.float32)
canvas_cache["scratch_src"] = _ss
if _sa is None or _sa.shape[0] < max_th or _sa.shape[1] < max_tw:
_sa = np.empty((max_th, max_tw), dtype=np.float32)
canvas_cache["scratch_alpha"] = _sa
_scratch = (_ss, _sa)
blended_tiles = 0
for tile_bgr, H_use, paste in zip(tiles_bgr_c, H_use_list, pastes):
if H_use is None or paste is None:
@@ -603,13 +662,13 @@ class Stitcher:
wkey = (th, tw, blend_margin)
if wkey not in weight_maps:
weight_maps[wkey] = self._tile_weight_map(th, tw, blend_margin)
self._blend_tile_inplace(blend_buf, wsum, cov, tile_bgr, weight_maps[wkey], y0, x0)
self._blend_tile_inplace(blend_buf, wsum, cov, tile_bgr, weight_maps[wkey], y0, x0, _scratch)
blended_tiles += 1
ref_wkey = (h0c, w0c, blend_margin)
if ref_wkey not in weight_maps:
weight_maps[ref_wkey] = self._tile_weight_map(h0c, w0c, blend_margin)
self._blend_tile_inplace(blend_buf, wsum, cov, t_ref_bgr, weight_maps[ref_wkey], ref_y0, ref_x0)
self._blend_tile_inplace(blend_buf, wsum, cov, t_ref_bgr, weight_maps[ref_wkey], ref_y0, ref_x0, _scratch)
t_after_blend = cv2.getTickCount() / cv2.getTickFrequency() if do_profile else 0.0
crop_rect = canvas_cache.get("crop_rect")
@@ -618,9 +677,16 @@ class Stitcher:
canvas_cache["crop_rect"] = crop_rect
if crop_rect is not None:
y0, y1, x0, x1 = crop_rect
stitched = np.clip(blend_buf[y0:y1, x0:x1], 0, 255).astype(np.uint8)
roi = blend_buf[y0:y1, x0:x1]
else:
stitched = np.clip(blend_buf, 0, 255).astype(np.uint8)
roi = blend_buf
rh, rw = roi.shape[:2]
stitch_buf = canvas_cache.get("stitch_buf")
if stitch_buf is None or stitch_buf.shape[0] != rh or stitch_buf.shape[1] != rw:
stitch_buf = np.empty((rh, rw, 3), dtype=np.uint8)
canvas_cache["stitch_buf"] = stitch_buf
cv2.convertScaleAbs(roi, dst=stitch_buf)
stitched = stitch_buf
t_after_finalize = cv2.getTickCount() / cv2.getTickFrequency() if do_profile else 0.0
if do_profile:
src/orisys/sensor.py
+37 -21
View File
@@ -60,7 +60,8 @@ class Sensor:
# =========================================================================
def __init__(self, vid_src, cuda=True, isstitch=True, verbose=True, debug=False,
config_name="prototype", cal_path=None, override=False, box=None, mirror=True, backend='auto'):
config_name="prototype", cal_path=None, override=False, box=None, mirror=True,
backend='auto'):
"""
Initialize Sensor with video source and configuration.
@@ -78,6 +79,8 @@ class Sensor:
"cupy": Use CuPy for GPU acceleration
"opencl": Use OpenCL for GPU acceleration
"cpu": Use CPU for computation
hhd_downsample: Downsample factor for flow field before HHD (default: 1, no downsampling).
Factor of 2 halves each dimension, giving ~4x speedup in Poisson solve.
"""
self.CUPY_AVAILABLE = check_cupy_available()
@@ -147,7 +150,12 @@ class Sensor:
self.__tracker = Tracker(adaptive=True, cuda=cuda, verbose=self.verbose, config_tracker=self.__config_tracker)
self.__contact3d = Contact3D(config_contact=self.__config_contact, verbose=self.verbose)
self.__hhd = None
self.poisson_backend = backend
# cfg_ds = self.__config_tracker.get("hhd_downsample", None)
# self._hhd_downsample = max(1, int(cfg_ds if cfg_ds is not None else hhd_downsample))
self._hhd_downsample = self.__config_tracker.get("hhd_downsample", None)
self._hhd_grid_shape = None
cfg_backend = self.__config_tracker.get("backend", None)
self.poisson_backend = cfg_backend if cfg_backend is not None else backend
self._use_pyav = False
self._av_container = None
self._av_stream = None
@@ -287,28 +295,36 @@ class Sensor:
flow = self.flow
if flow is None:
raise Exception("Warning: Compute flow first")
# Initialize HHD solver
if self.__hhd is None:
self.__hhd = nHHD(cuda=self.cuda, grid=flow.shape[:2], spacings=(18, 18), poisson_backend=self.poisson_backend, debug=self.debug)
# Force computation
# When OpenCL backend is selected, use host force path to ensure
# Poisson solve goes through PoissonSolver.solve_opencl().
_, forces = self.__hhd.get_force(flow)
# if self.poisson_backend == 'opencl':
# _, forces = self.__hhd.get_force_host(flow)
# elif self.cuda and self.CUPY_AVAILABLE:
# _, forces = self.__hhd.get_force_gpu_v2(flow)
# else:
# _, forces = self.__hhd.get_force_host(flow)
orig_h, orig_w = flow.shape[:2]
k = self._hhd_downsample
if k > 1:
ds_w = max(1, orig_w // k)
ds_h = max(1, orig_h // k)
flow_ds = cv2.resize(flow, (ds_w, ds_h), interpolation=cv2.INTER_AREA)
else:
flow_ds = flow
ds_shape = flow_ds.shape[:2]
if self.__hhd is None or self._hhd_grid_shape != ds_shape:
spacings = (18 * k, 18 * k)
self.__hhd = nHHD(cuda=self.cuda, grid=ds_shape, spacings=spacings, poisson_backend=self.poisson_backend, debug=self.debug)
self._hhd_grid_shape = ds_shape
_, forces = self.__hhd.get_force(flow_ds)
self.fnormal, self.fshearx, self.fsheary, self.ftor = forces
# Get components
self.vnormal = self.__hhd.get_normal()
self.vshear = self.__hhd.get_shear()
# Update state
vnormal_ds = self.__hhd.get_normal()
vshear_ds = self.__hhd.get_shear()
if k > 1:
self.vnormal = cv2.resize(vnormal_ds, (orig_w, orig_h), interpolation=cv2.INTER_LINEAR)
self.vshear = cv2.resize(vshear_ds, (orig_w, orig_h), interpolation=cv2.INTER_LINEAR)
else:
self.vnormal = vnormal_ds
self.vshear = vshear_ds
self.compute_deformation_valid = True
return True
src/orisys/tracker.py
+2 -2
View File
@@ -63,8 +63,8 @@ class Tracker:
# self.__tracking_inst.setUseSpatialPropagation(False)
# self.__tracking_inst.setUseMeanNormalization(False)
# self.__tracking_inst.setGradientDescentIterations(16) #16
self.__tracking_inst.setPatchSize(8) # 8
self.__tracking_inst.setPatchStride(1) # 3
# self.__tracking_inst.setPatchSize(12) # 8
# self.__tracking_inst.setPatchStride(5) # 3
# self.__tracking_inst.setVariationalRefinementIterations(5)#5
# else: