Tici camerad (#2048)
* fix camera config * typos * oops * more typo * lambless * forget to send * visualize hist * more typo * 0xC000 * simple * data loss prevention, clean up later * loggerd * back up code * backup * fixed memory leak * fix vsync * upload ecam * WB * 3stream * fix OMX crash on loggerd rotation * rewritten debayer kernel * update viewer * improved AE * no artifact lines/grids * standard trigger * cleanups * CCM * cleanups * slight tweak * upd push sock * build all these * update tele fl * update cereal * upd viewer * DualCameraState -> MultiCameraState * cameras_open * disable frame zmq push by default * more cleanup * no apks * fix submodule error * wat * clean up trash * remove junk * only build on qcom2 * no need to check these * update cereal * some more minor cleanup * bump panda * add todo * minor typo * Revert "minor typo" This reverts commit 9233a1df7cac214fae6827cdae3a10cb3bd060d9. * not care * use consistent hdr * some cleanup * Update selfdrive/camerad/main.cc Co-authored-by: Adeeb Shihadeh <adeebshihadeh@gmail.com> * more cleanups * remove irrelevant stuff * this too * cleanup * rerun ci Co-authored-by: ZwX1616 <zwx1616@gmail.com> Co-authored-by: Tici <robbe@comma.ai> Co-authored-by: Adeeb Shihadeh <adeebshihadeh@gmail.com>pull/49/head
robbederks
5 years ago
committed by
GitHub
parent
3014280d1a
commit
37d6472bfa
21 changed files with 1244 additions and 467 deletions
@ -0,0 +1,128 @@ |
||||
const __constant float3 color_correction[3] = { |
||||
// post wb CCM |
||||
(float3)(1.17898387, -0.19583185, -0.19881648), |
||||
(float3)(-0.03367879, 1.33692858, -0.02475203), |
||||
(float3)(-0.14530508, -0.14109673, 1.22356851), |
||||
}; |
||||
|
||||
float3 color_correct(float r, float g, float b) { |
||||
float3 ret = (0,0,0); |
||||
|
||||
ret += r * color_correction[0]; |
||||
ret += g * color_correction[1]; |
||||
ret += b * color_correction[2]; |
||||
ret = max(0.0, min(1.0, ret)); |
||||
return ret; |
||||
} |
||||
|
||||
uint int_from_10(const uchar * source, uint start, uint offset) { |
||||
// source: source |
||||
// start: starting address of 0 |
||||
// offset: 0 - 3 |
||||
uint major = (uint)source[start + offset] << 2; |
||||
uint minor = (source[start + 4] >> (2 * offset)) & 3; |
||||
return major + minor; |
||||
} |
||||
|
||||
float to_normal(uint x) { |
||||
float pv = (float)(x); |
||||
const float black_level = 0; |
||||
pv = max(0.0, pv - black_level); |
||||
pv /= (1024.0f - black_level); |
||||
pv = 1.6*pv / (1.0f + 1.6*pv); // reinhard |
||||
return pv; |
||||
} |
||||
|
||||
__kernel void debayer10(const __global uchar * in, |
||||
__global uchar * out, |
||||
__local float * cached |
||||
) |
||||
{ |
||||
const int x_global = get_global_id(0); |
||||
const int y_global = get_global_id(1); |
||||
|
||||
// const int globalOffset = ; |
||||
|
||||
const int localRowLen = 2 + get_local_size(0); // 2 padding |
||||
const int x_local = get_local_id(0); |
||||
const int y_local = get_local_id(1); |
||||
|
||||
const int localOffset = (y_local + 1) * localRowLen + x_local + 1; |
||||
|
||||
// cache local pixels first |
||||
// saves memory access and avoids repeated normalization |
||||
uint globalStart_10 = y_global * FRAME_STRIDE + (5 * (x_global / 4)); |
||||
uint offset_10 = x_global % 4; |
||||
uint raw_val = int_from_10(in, globalStart_10, offset_10); |
||||
cached[localOffset] = to_normal(raw_val); |
||||
|
||||
// edges |
||||
if (x_global < 1 || x_global > RGB_WIDTH - 2 || y_global < 1 || y_global > RGB_WIDTH - 2) { |
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
return; |
||||
} else { |
||||
int localColOffset = -1; |
||||
int globalColOffset = -1; |
||||
|
||||
// cache padding |
||||
if (x_local < 1) { |
||||
localColOffset = x_local; |
||||
globalColOffset = -1; |
||||
cached[(y_local + 1) * localRowLen + x_local] = to_normal(int_from_10(in, y_global * FRAME_STRIDE + (5 * ((x_global-1) / 4)), (offset_10 + 3) % 4)); |
||||
} else if (x_local >= get_local_size(0) - 1) { |
||||
localColOffset = x_local + 2; |
||||
globalColOffset = 1; |
||||
cached[localOffset + 1] = to_normal(int_from_10(in, y_global * FRAME_STRIDE + (5 * ((x_global+1) / 4)), (offset_10 + 1) % 4)); |
||||
} |
||||
|
||||
if (y_local < 1) { |
||||
cached[y_local * localRowLen + x_local + 1] = to_normal(int_from_10(in, globalStart_10 - FRAME_STRIDE, offset_10)); |
||||
if (localColOffset != -1) { |
||||
cached[y_local * localRowLen + localColOffset] = to_normal(int_from_10(in, (y_global-1) * FRAME_STRIDE + (5 * ((x_global+globalColOffset) / 4)), (offset_10+4+globalColOffset) % 4)); |
||||
} |
||||
} else if (y_local >= get_local_size(1) - 1) { |
||||
cached[(y_local + 2) * localRowLen + x_local + 1] = to_normal(int_from_10(in, globalStart_10 + FRAME_STRIDE, offset_10)); |
||||
if (localColOffset != -1) { |
||||
cached[(y_local + 2) * localRowLen + localColOffset] = to_normal(int_from_10(in, (y_global+1) * FRAME_STRIDE + (5 * ((x_global+globalColOffset) / 4)), (offset_10+4+globalColOffset) % 4)); |
||||
} |
||||
} |
||||
|
||||
// sync |
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
// perform debayer |
||||
float r; |
||||
float g; |
||||
float b; |
||||
|
||||
if (x_global % 2 == 0) { |
||||
if (y_global % 2 == 0) { // G1 |
||||
r = (cached[localOffset - 1] + cached[localOffset + 1]) / 2.0f; |
||||
g = (cached[localOffset] + cached[localOffset + localRowLen + 1]) / 2.0f; |
||||
b = (cached[localOffset - localRowLen] + cached[localOffset + localRowLen]) / 2.0f; |
||||
} else { // B |
||||
r = (cached[localOffset - localRowLen - 1] + cached[localOffset - localRowLen + 1] + cached[localOffset + localRowLen - 1] + cached[localOffset + localRowLen + 1]) / 4.0f; |
||||
g = (cached[localOffset - localRowLen] + cached[localOffset + localRowLen] + cached[localOffset - 1] + cached[localOffset + 1]) / 4.0f; |
||||
b = cached[localOffset]; |
||||
} |
||||
} else { |
||||
if (y_global % 2 == 0) { // R |
||||
r = cached[localOffset]; |
||||
g = (cached[localOffset - localRowLen] + cached[localOffset + localRowLen] + cached[localOffset - 1] + cached[localOffset + 1]) / 4.0f; |
||||
b = (cached[localOffset - localRowLen - 1] + cached[localOffset - localRowLen + 1] + cached[localOffset + localRowLen - 1] + cached[localOffset + localRowLen + 1]) / 4.0f; |
||||
} else { // G2 |
||||
r = (cached[localOffset - localRowLen] + cached[localOffset + localRowLen]) / 2.0f; |
||||
g = (cached[localOffset] + cached[localOffset - localRowLen - 1]) / 2.0f; |
||||
b = (cached[localOffset - 1] + cached[localOffset + 1]) / 2.0f; |
||||
} |
||||
} |
||||
|
||||
float3 rgb = color_correct(r, g, b); |
||||
// rgb = srgb_gamma(rgb); |
||||
|
||||
// BGR output |
||||
out[3 * x_global + 3 * y_global * RGB_WIDTH + 0] = (uchar)(255.0f * rgb.z); |
||||
out[3 * x_global + 3 * y_global * RGB_WIDTH + 1] = (uchar)(255.0f * rgb.y); |
||||
out[3 * x_global + 3 * y_global * RGB_WIDTH + 2] = (uchar)(255.0f * rgb.x); |
||||
} |
||||
} |
||||
@ -0,0 +1,47 @@ |
||||
# flake8: noqa |
||||
# pylint: disable=W |
||||
|
||||
#!/usr/bin/env python |
||||
import numpy as np |
||||
import cv2 |
||||
from time import time, sleep |
||||
|
||||
H, W = (604*2//6, 964*2//6) |
||||
# H, W = (604, 964) |
||||
|
||||
cam_bufs = np.zeros((3,H,W,3), dtype=np.uint8) |
||||
hist_bufs = np.zeros((3,H,200,3), dtype=np.uint8) |
||||
|
||||
if __name__ == '__main__': |
||||
import zmq |
||||
context = zmq.Context() |
||||
socket = context.socket(zmq.PULL) |
||||
socket.bind("tcp://192.168.3.4:7768") |
||||
while True: |
||||
try: |
||||
message = socket.recv() |
||||
except Exception as ex: |
||||
print(ex) |
||||
message = b"123" |
||||
|
||||
dat = np.frombuffer(message, dtype=np.uint8) |
||||
cam_id = (dat[0] + 1) % 3 |
||||
# import pdb; pdb.set_trace() |
||||
b = dat[::3].reshape(H, W) |
||||
g = dat[1::3].reshape(H, W) |
||||
r = dat[2::3].reshape(H, W) |
||||
cam_bufs[cam_id] = cv2.merge((r, g, b)) |
||||
cam_bufs[cam_id]= cv2.cvtColor(cam_bufs[cam_id], cv2.COLOR_RGB2BGR) |
||||
|
||||
hist = cv2.calcHist([cv2.cvtColor(cam_bufs[cam_id], cv2.COLOR_BGR2GRAY)],[0],None,[32],[0,256]) |
||||
hist = (H*hist/hist.max()).astype(np.uint8) |
||||
for i,bb in enumerate(hist): |
||||
hist_bufs[cam_id, H-bb[0]:,i*(200//32):(i+1)*(200//32), :] = (222,222,222) |
||||
|
||||
out = cam_bufs.reshape((3*H,W,3)) |
||||
hist_bufs_out = hist_bufs.reshape((3*H,200,3)) |
||||
out = np.hstack([out, hist_bufs_out]) |
||||
cv2.imshow('RGB', out) |
||||
cv2.waitKey(55) |
||||
#dat.tofile('/tmp/c3rgb.img') |
||||
#cv2.imwrite('/tmp/c3rgb.png', out) |
||||
@ -0,0 +1,34 @@ |
||||
# flake8: noqa |
||||
# pylint: disable=W |
||||
|
||||
#!/usr/bin/env python |
||||
import numpy as np |
||||
import cv2 |
||||
from time import time, sleep |
||||
|
||||
H, W = (256, 512) |
||||
|
||||
if __name__ == '__main__': |
||||
import zmq |
||||
context = zmq.Context() |
||||
socket = context.socket(zmq.PULL) |
||||
socket.bind("tcp://192.168.3.4:7769") |
||||
while True: |
||||
try: |
||||
message = socket.recv() |
||||
except Exception as ex: |
||||
print(ex) |
||||
message = b"123" |
||||
|
||||
dat = np.frombuffer(message, dtype=np.float32) |
||||
mc = (dat.reshape(H//2, W//2)).astype(np.uint8) |
||||
|
||||
hist = cv2.calcHist([mc],[0],None,[32],[0,256]) |
||||
hist = (H*hist/hist.max()).astype(np.uint8) |
||||
himg = np.zeros((H//2, W//2), dtype=np.uint8) |
||||
for i,b in enumerate(hist): |
||||
himg[H//2-b[0]:,i*(W//2//32):(i+1)*(W//2//32)] = 222 |
||||
|
||||
cv2.imshow('model fov', np.hstack([mc, himg])) |
||||
cv2.waitKey(20) |
||||
dat.tofile('/tmp/c3yuv.img') |
||||
Loading…
Reference in new issue