mawei/onepilot
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

rm pyopencl package (#37058)

Browse Source 
rm pyopencl
This commit is contained in:
Adeeb Shihadeh
2026-02-01 15:42:42 -08:00
committed by GitHub
parent 422de59898
commit 948d42b3e5
6 changed files with 35 additions and 174 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
.gitignore vendored
View File

@@ -95,3 +95,7 @@ Pipfile
# Ignore all local history of files
.history
.ionide
.claude/
PLAN.md
TASK.md

1
pyproject.toml
View File

@@ -112,7 +112,6 @@ dev = [
"opencv-python-headless",
"parameterized >=0.8, <0.9",
"pyautogui",
"pyopencl",
"pytools>=2025.1.6; platform_machine != 'aarch64'",
"pywinctl",
"pyprof2calltree",

54
tools/sim/lib/camerad.py
View File

@@ -1,14 +1,39 @@
import numpy as np
import os
import pyopencl as cl
import pyopencl.array as cl_array
from msgq.visionipc import VisionIpcServer, VisionStreamType
from cereal import messaging
from openpilot.common.basedir import BASEDIR
from openpilot.tools.sim.lib.common import W, H
def rgb_to_nv12(rgb):
"""Convert RGB image to NV12 (YUV420) format using BT.601 coefficients."""
h, w = rgb.shape[:2]
r = rgb[:, :, 0].astype(np.int32)
g = rgb[:, :, 1].astype(np.int32)
b = rgb[:, :, 2].astype(np.int32)
# Y plane - BT.601 coefficients (matches original OpenCL kernel)
y = (((b * 13 + g * 65 + r * 33) + 64) >> 7) + 16
y = np.clip(y, 0, 255).astype(np.uint8)
# Subsample RGB for UV (2x2 box filter)
r_sub = (r[0::2, 0::2] + r[0::2, 1::2] + r[1::2, 0::2] + r[1::2, 1::2] + 2) >> 2
g_sub = (g[0::2, 0::2] + g[0::2, 1::2] + g[1::2, 0::2] + g[1::2, 1::2] + 2) >> 2
b_sub = (b[0::2, 0::2] + b[0::2, 1::2] + b[1::2, 0::2] + b[1::2, 1::2] + 2) >> 2
# U and V planes
u = np.clip((b_sub * 56 - g_sub * 37 - r_sub * 19 + 0x8080) >> 8, 0, 255).astype(np.uint8)
v = np.clip((r_sub * 56 - g_sub * 47 - b_sub * 9 + 0x8080) >> 8, 0, 255).astype(np.uint8)
# Interleave UV for NV12 format
uv = np.empty((h // 2, w), dtype=np.uint8)
uv[:, 0::2] = u
uv[:, 1::2] = v
return np.concatenate([y.ravel(), uv.ravel()]).tobytes()
class Camerad:
"""Simulates the camerad daemon"""
def __init__(self, dual_camera):
@@ -24,18 +49,6 @@ class Camerad:
self.vipc_server.start_listener()
# set up for pyopencl rgb to yuv conversion
self.ctx = cl.create_some_context()
self.queue = cl.CommandQueue(self.ctx)
cl_arg = f" -DHEIGHT={H} -DWIDTH={W} -DRGB_STRIDE={W * 3} -DUV_WIDTH={W // 2} -DUV_HEIGHT={H // 2} -DRGB_SIZE={W * H} -DCL_DEBUG "
kernel_fn = os.path.join(BASEDIR, "tools/sim/rgb_to_nv12.cl")
with open(kernel_fn) as f:
prg = cl.Program(self.ctx, f.read()).build(cl_arg)
self.krnl = prg.rgb_to_nv12
self.Wdiv4 = W // 4 if (W % 4 == 0) else (W + (4 - W % 4)) // 4
self.Hdiv4 = H // 4 if (H % 4 == 0) else (H + (4 - H % 4)) // 4
def cam_send_yuv_road(self, yuv):
self._send_yuv(yuv, self.frame_road_id, 'roadCameraState', VisionStreamType.VISION_STREAM_ROAD)
self.frame_road_id += 1
@@ -44,16 +57,11 @@ class Camerad:
self._send_yuv(yuv, self.frame_wide_id, 'wideRoadCameraState', VisionStreamType.VISION_STREAM_WIDE_ROAD)
self.frame_wide_id += 1
# Returns: yuv bytes
def rgb_to_yuv(self, rgb):
"""Convert RGB to NV12 YUV format."""
assert rgb.shape == (H, W, 3), f"{rgb.shape}"
assert rgb.dtype == np.uint8
rgb_cl = cl_array.to_device(self.queue, rgb)
yuv_cl = cl_array.empty_like(rgb_cl)
self.krnl(self.queue, (self.Wdiv4, self.Hdiv4), None, rgb_cl.data, yuv_cl.data).wait()
yuv = np.resize(yuv_cl.get(), rgb.size // 2)
return yuv.data.tobytes()
return rgb_to_nv12(rgb)
def _send_yuv(self, yuv, frame_id, pub_type, yuv_type):
eof = int(frame_id * 0.05 * 1e9)

119
tools/sim/rgb_to_nv12.cl
View File

@@ -1,119 +0,0 @@
#define RGB_TO_Y(r, g, b) ((((mul24(b, 13) + mul24(g, 65) + mul24(r, 33)) + 64) >> 7) + 16)
#define RGB_TO_U(r, g, b) ((mul24(b, 56) - mul24(g, 37) - mul24(r, 19) + 0x8080) >> 8)
#define RGB_TO_V(r, g, b) ((mul24(r, 56) - mul24(g, 47) - mul24(b, 9) + 0x8080) >> 8)
#define AVERAGE(x, y, z, w) ((convert_ushort(x) + convert_ushort(y) + convert_ushort(z) + convert_ushort(w) + 1) >> 1)
inline void convert_2_ys(__global uchar * out_yuv, int yi, const uchar8 rgbs1) {
uchar2 yy = (uchar2)(
RGB_TO_Y(rgbs1.s2, rgbs1.s1, rgbs1.s0),
RGB_TO_Y(rgbs1.s5, rgbs1.s4, rgbs1.s3)
);
#ifdef CL_DEBUG
if(yi >= RGB_SIZE)
printf("Y vector2 overflow, %d > %d\n", yi, RGB_SIZE);
#endif
vstore2(yy, 0, out_yuv + yi);
}
inline void convert_4_ys(__global uchar * out_yuv, int yi, const uchar8 rgbs1, const uchar8 rgbs3) {
const uchar4 yy = (uchar4)(
RGB_TO_Y(rgbs1.s2, rgbs1.s1, rgbs1.s0),
RGB_TO_Y(rgbs1.s5, rgbs1.s4, rgbs1.s3),
RGB_TO_Y(rgbs3.s0, rgbs1.s7, rgbs1.s6),
RGB_TO_Y(rgbs3.s3, rgbs3.s2, rgbs3.s1)
);
#ifdef CL_DEBUG
if(yi > RGB_SIZE - 4)
printf("Y vector4 overflow, %d > %d\n", yi, RGB_SIZE - 4);
#endif
vstore4(yy, 0, out_yuv + yi);
}
inline void convert_uv(__global uchar * out_yuv, int uvi,
const uchar8 rgbs1, const uchar8 rgbs2) {
// U & V: average of 2x2 pixels square
const short ab = AVERAGE(rgbs1.s0, rgbs1.s3, rgbs2.s0, rgbs2.s3);
const short ag = AVERAGE(rgbs1.s1, rgbs1.s4, rgbs2.s1, rgbs2.s4);
const short ar = AVERAGE(rgbs1.s2, rgbs1.s5, rgbs2.s2, rgbs2.s5);
#ifdef CL_DEBUG
if(uvi >= RGB_SIZE + RGB_SIZE / 2)
printf("UV overflow, %d >= %d\n", uvi, RGB_SIZE + RGB_SIZE / 2);
#endif
out_yuv[uvi] = RGB_TO_U(ar, ag, ab);
out_yuv[uvi+1] = RGB_TO_V(ar, ag, ab);
}
inline void convert_2_uvs(__global uchar * out_yuv, int uvi,
const uchar8 rgbs1, const uchar8 rgbs2, const uchar8 rgbs3, const uchar8 rgbs4) {
// U & V: average of 2x2 pixels square
const short ab1 = AVERAGE(rgbs1.s0, rgbs1.s3, rgbs2.s0, rgbs2.s3);
const short ag1 = AVERAGE(rgbs1.s1, rgbs1.s4, rgbs2.s1, rgbs2.s4);
const short ar1 = AVERAGE(rgbs1.s2, rgbs1.s5, rgbs2.s2, rgbs2.s5);
const short ab2 = AVERAGE(rgbs1.s6, rgbs3.s1, rgbs2.s6, rgbs4.s1);
const short ag2 = AVERAGE(rgbs1.s7, rgbs3.s2, rgbs2.s7, rgbs4.s2);
const short ar2 = AVERAGE(rgbs3.s0, rgbs3.s3, rgbs4.s0, rgbs4.s3);
uchar4 uv = (uchar4)(
RGB_TO_U(ar1, ag1, ab1),
RGB_TO_V(ar1, ag1, ab1),
RGB_TO_U(ar2, ag2, ab2),
RGB_TO_V(ar2, ag2, ab2)
);
#ifdef CL_DEBUG1
if(uvi > RGB_SIZE + RGB_SIZE / 2 - 4)
printf("UV2 overflow, %d >= %d\n", uvi, RGB_SIZE + RGB_SIZE / 2 - 2);
#endif
vstore4(uv, 0, out_yuv + uvi);
}
__kernel void rgb_to_nv12(__global uchar const * const rgb,
__global uchar * out_yuv)
{
const int dx = get_global_id(0);
const int dy = get_global_id(1);
const int col = mul24(dx, 4); // Current column in rgb image
const int row = mul24(dy, 4); // Current row in rgb image
const int bgri_start = mad24(row, RGB_STRIDE, mul24(col, 3)); // Start offset of rgb data being converted
const int yi_start = mad24(row, WIDTH, col); // Start offset in the target yuv buffer
int uvi = mad24(row / 2, WIDTH, RGB_SIZE + col);
int num_col = min(WIDTH - col, 4);
int num_row = min(HEIGHT - row, 4);
if(num_row == 4) {
const uchar8 rgbs0_0 = vload8(0, rgb + bgri_start);
const uchar8 rgbs0_1 = vload8(0, rgb + bgri_start + 8);
const uchar8 rgbs1_0 = vload8(0, rgb + bgri_start + RGB_STRIDE);
const uchar8 rgbs1_1 = vload8(0, rgb + bgri_start + RGB_STRIDE + 8);
const uchar8 rgbs2_0 = vload8(0, rgb + bgri_start + RGB_STRIDE * 2);
const uchar8 rgbs2_1 = vload8(0, rgb + bgri_start + RGB_STRIDE * 2 + 8);
const uchar8 rgbs3_0 = vload8(0, rgb + bgri_start + RGB_STRIDE * 3);
const uchar8 rgbs3_1 = vload8(0, rgb + bgri_start + RGB_STRIDE * 3 + 8);
if(num_col == 4) {
convert_4_ys(out_yuv, yi_start, rgbs0_0, rgbs0_1);
convert_4_ys(out_yuv, yi_start + WIDTH, rgbs1_0, rgbs1_1);
convert_4_ys(out_yuv, yi_start + WIDTH * 2, rgbs2_0, rgbs2_1);
convert_4_ys(out_yuv, yi_start + WIDTH * 3, rgbs3_0, rgbs3_1);
convert_2_uvs(out_yuv, uvi, rgbs0_0, rgbs1_0, rgbs0_1, rgbs1_1);
convert_2_uvs(out_yuv, uvi + WIDTH, rgbs2_0, rgbs3_0, rgbs2_1, rgbs3_1);
} else if(num_col == 2) {
convert_2_ys(out_yuv, yi_start, rgbs0_0);
convert_2_ys(out_yuv, yi_start + WIDTH, rgbs1_0);
convert_2_ys(out_yuv, yi_start + WIDTH * 2, rgbs2_0);
convert_2_ys(out_yuv, yi_start + WIDTH * 3, rgbs3_0);
convert_uv(out_yuv, uvi, rgbs0_0, rgbs1_0);
convert_uv(out_yuv, uvi + WIDTH, rgbs2_0, rgbs3_0);
}
} else {
const uchar8 rgbs0_0 = vload8(0, rgb + bgri_start);
const uchar8 rgbs0_1 = vload8(0, rgb + bgri_start + 8);
const uchar8 rgbs1_0 = vload8(0, rgb + bgri_start + RGB_STRIDE);
const uchar8 rgbs1_1 = vload8(0, rgb + bgri_start + RGB_STRIDE + 8);
if(num_col == 4) {
convert_4_ys(out_yuv, yi_start, rgbs0_0, rgbs0_1);
convert_4_ys(out_yuv, yi_start + WIDTH, rgbs1_0, rgbs1_1);
convert_2_uvs(out_yuv, uvi, rgbs0_0, rgbs1_0, rgbs0_1, rgbs1_1);
} else if(num_col == 2) {
convert_2_ys(out_yuv, yi_start, rgbs0_0);
convert_2_ys(out_yuv, yi_start + WIDTH, rgbs1_0);
convert_uv(out_yuv, uvi, rgbs0_0, rgbs1_0);
}
}
}

1
tools/sim/tests/test_metadrive_bridge.py
View File

@@ -8,7 +8,6 @@ from openpilot.tools.sim.bridge.metadrive.metadrive_bridge import MetaDriveBridg
from openpilot.tools.sim.tests.test_sim_bridge import TestSimBridgeBase
@pytest.mark.slow
@pytest.mark.filterwarnings("ignore::pyopencl.CompilerWarning") # Unimportant warning of non-empty compile log
class TestMetaDriveBridge(TestSimBridgeBase):
@pytest.fixture(autouse=True)
def setup_create_bridge(self, test_duration):

30
uv.lock generated
View File

@@ -1336,7 +1336,6 @@ dev = [
{ name = "opencv-python-headless" },
{ name = "parameterized" },
{ name = "pyautogui" },
{ name = "pyopencl" },
{ name = "pyprof2calltree" },
{ name = "pytools", marker = "platform_machine != 'aarch64'" },
{ name = "pywinctl" },
@@ -1409,7 +1408,6 @@ requires-dist = [
{ name = "pycapnp", specifier = "==2.1.0" },
{ name = "pycryptodome" },
{ name = "pyjwt" },
{ name = "pyopencl", marker = "extra == 'dev'" },
{ name = "pyopenssl", specifier = "<24.3.0" },
{ name = "pyprof2calltree", marker = "extra == 'dev'" },
{ name = "pyserial" },
@@ -4247,34 +4245,6 @@ wheels = [
{ url = "https://files.pythonhosted.org/packages/db/67/64920c8d201a7fc27962f467c636c4e763b43845baba2e091a50a97a5d52/pyobjc_framework_webkit-12.1-cp312-cp312-macosx_10_13_universal2.whl", hash = "sha256:af2c7197447638b92aafbe4847c063b6dd5e1ed83b44d3ce7e71e4c9b042ab5a", size = 50084, upload-time = "2025-11-14T10:07:05.868Z" },
]
[[package]]
name = "pyopencl"
version = "2026.1.2"
source = { registry = "https://pypi.org/simple" }
dependencies = [
{ name = "numpy" },
{ name = "platformdirs" },
{ name = "pytools" },
{ name = "typing-extensions" },
]
sdist = { url = "https://files.pythonhosted.org/packages/d8/81/fd8a2a695916a82e861bcf17b5b8fd9f81e12c9e5931f9ba536678d7b43a/pyopencl-2026.1.2.tar.gz", hash = "sha256:4397dd0b4cbb8b55f3e09bf87114a2465574506b363890b805b860c348b61970", size = 445132, upload-time = "2026-01-16T22:52:24.765Z" }
wheels = [
{ url = "https://files.pythonhosted.org/packages/e3/88/abf34e31d572c59203774a66cd81c1e3b3d60b911241483675151149c6f1/pyopencl-2026.1.2-cp311-cp311-macosx_10_14_x86_64.whl", hash = "sha256:8052e8b402b3ed33ee0807d87d4734f66f67dbafbfb3f5a8b81e478e4d417372", size = 437029, upload-time = "2026-01-16T22:51:30.953Z" },
{ url = "https://files.pythonhosted.org/packages/5c/3d/2dd2d8bbf05a190681582b40fc1ee55b210d00ccebcbb416c62b9f9c81a1/pyopencl-2026.1.2-cp311-cp311-macosx_11_0_arm64.whl", hash = "sha256:d5e03681c3fe22d5185b16a727d96783e3787e0b65e7a29e4afe01ae0cb4e802", size = 429031, upload-time = "2026-01-16T22:51:32.674Z" },
{ url = "https://files.pythonhosted.org/packages/41/16/e554b3bd20be2e858cfb6683ee6549aeebbe5f769e5b95f561f79340ab20/pyopencl-2026.1.2-cp311-cp311-manylinux_2_26_aarch64.manylinux_2_28_aarch64.whl", hash = "sha256:1c8c209d517d1421b17d20b80589a2c39e09ea33350f0367314e1caeed3bc741", size = 689596, upload-time = "2026-01-16T22:51:33.913Z" },
{ url = "https://files.pythonhosted.org/packages/22/a8/1df41cf6c7b25b3bfda14aa0183c6a90eaf849528ba27753eaa25fb26e20/pyopencl-2026.1.2-cp311-cp311-manylinux_2_27_x86_64.manylinux_2_28_x86_64.whl", hash = "sha256:e64e2e34bcfad426bd24b71fdb6b02aa5cb02475147742fe07ef93e81866fc7e", size = 736427, upload-time = "2026-01-16T22:51:36.595Z" },
{ url = "https://files.pythonhosted.org/packages/fc/3d/177b6a675691f7b6f708faef33f981e72fbc4bfed2b1dfa94dc70d0e8a25/pyopencl-2026.1.2-cp311-cp311-musllinux_1_2_aarch64.whl", hash = "sha256:65b151c56b936481d6b6050c2b9bc520840e1402be78c282ba5c01921c25477d", size = 1163888, upload-time = "2026-01-16T22:51:37.973Z" },
{ url = "https://files.pythonhosted.org/packages/e9/fa/5905571d9fa48827c0427a3e664c0213dd045940d581b3b739d83df9c0f6/pyopencl-2026.1.2-cp311-cp311-musllinux_1_2_x86_64.whl", hash = "sha256:cc40003446037f391ca0970694efb0627e2870fabb20ee21be75bc445a39d8f4", size = 1228235, upload-time = "2026-01-16T22:51:39.786Z" },
{ url = "https://files.pythonhosted.org/packages/1e/3d/538c675d078b91680d8d82962110d0c9fd42e1584763d515d6e2e82d8c57/pyopencl-2026.1.2-cp311-cp311-win_amd64.whl", hash = "sha256:b6a8e109ade7db60e8b1beb48df8f080941d0cd77fb2c225ad509c80cdef603e", size = 474753, upload-time = "2026-01-16T22:51:41.771Z" },
{ url = "https://files.pythonhosted.org/packages/cd/34/1497070e44d1689ddbd01d24a2265910e84ebc53457a489b9d2b6e1ac675/pyopencl-2026.1.2-cp312-cp312-macosx_10_14_x86_64.whl", hash = "sha256:7d88e59901bfe1f9296fd89acd9968f008dc7cfee7995f8cd09c3f1a77119aa6", size = 438145, upload-time = "2026-01-16T22:51:43.658Z" },
{ url = "https://files.pythonhosted.org/packages/5b/a3/71d6af8741b52d3bef443518c1ccfda003adcfa9cc1d0df83dac7005d08c/pyopencl-2026.1.2-cp312-cp312-macosx_11_0_arm64.whl", hash = "sha256:3f96a3bff8a09d2fa924e7c33dafac6ea3ef7ec70e746d6d8e17ce2d959a6836", size = 428820, upload-time = "2026-01-16T22:51:45.326Z" },
{ url = "https://files.pythonhosted.org/packages/db/ea/c8dbabeceac9cad3dbb368e08e0aa208cc6c6251c5134cc25eb15da03639/pyopencl-2026.1.2-cp312-cp312-manylinux_2_26_aarch64.manylinux_2_28_aarch64.whl", hash = "sha256:4d4e8e8215ec4fdee4b235b61977cdb1c4f041b487bdcf357be799f45b423d61", size = 685478, upload-time = "2026-01-16T22:51:46.545Z" },
{ url = "https://files.pythonhosted.org/packages/64/c7/5854ef7471dfee195bcef6348a107525ca4d1b73c15240e6444d490f9920/pyopencl-2026.1.2-cp312-cp312-manylinux_2_27_x86_64.manylinux_2_28_x86_64.whl", hash = "sha256:d0052a8ccbd282d8ab196705e31f4c3ab344113ea5d5c3ddaeede00cdcab068b", size = 734017, upload-time = "2026-01-16T22:51:48.277Z" },
{ url = "https://files.pythonhosted.org/packages/3d/79/42d4eec282ed299b38d8136d05545113ec8771a1bd6b10bb4ba83ae1236c/pyopencl-2026.1.2-cp312-cp312-musllinux_1_2_aarch64.whl", hash = "sha256:e43da12a376e9283407c2820b24cceeaa129b042ac710947cf8e07b13e294689", size = 1159871, upload-time = "2026-01-16T22:51:49.569Z" },
{ url = "https://files.pythonhosted.org/packages/a0/9a/fdc5d3bed0440d6206109e051008aa0a54ca131d64314bbd42177b8f0763/pyopencl-2026.1.2-cp312-cp312-musllinux_1_2_x86_64.whl", hash = "sha256:1b14b2cf11dec9e0b75cbd14223d1b3c93950fc3e2f7a306b54fa1b17a2cae0f", size = 1225288, upload-time = "2026-01-16T22:51:51.125Z" },
{ url = "https://files.pythonhosted.org/packages/2d/e3/358c19180e0dab5c7dd1fcacc569e6a7ab02a7fddcb9c954f393ceddb2fa/pyopencl-2026.1.2-cp312-cp312-win_amd64.whl", hash = "sha256:d02d7ecabc8d34590dccffe12346689adc5a1ceb07df5acc4ea6c4db8aa28277", size = 474876, upload-time = "2026-01-16T22:51:52.912Z" },
]
[[package]]
name = "pyopenssl"
version = "24.2.1"