diff --git a/.gitignore b/.gitignore index e4992a3d05..e2a30fb70a 100644 --- a/.gitignore +++ b/.gitignore @@ -95,3 +95,7 @@ Pipfile # Ignore all local history of files .history .ionide + +.claude/ +PLAN.md +TASK.md diff --git a/pyproject.toml b/pyproject.toml index 57fd0b8355..2239770ac9 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -112,7 +112,6 @@ dev = [ "opencv-python-headless", "parameterized >=0.8, <0.9", "pyautogui", - "pyopencl", "pytools>=2025.1.6; platform_machine != 'aarch64'", "pywinctl", "pyprof2calltree", diff --git a/tools/sim/lib/camerad.py b/tools/sim/lib/camerad.py index be4e1a610c..7634b8524d 100644 --- a/tools/sim/lib/camerad.py +++ b/tools/sim/lib/camerad.py @@ -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) diff --git a/tools/sim/rgb_to_nv12.cl b/tools/sim/rgb_to_nv12.cl deleted file mode 100644 index 54816d5d7d..0000000000 --- a/tools/sim/rgb_to_nv12.cl +++ /dev/null @@ -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); - } - } -} diff --git a/tools/sim/tests/test_metadrive_bridge.py b/tools/sim/tests/test_metadrive_bridge.py index 04ce5d584f..9be640d736 100644 --- a/tools/sim/tests/test_metadrive_bridge.py +++ b/tools/sim/tests/test_metadrive_bridge.py @@ -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): diff --git a/uv.lock b/uv.lock index e488d1d78b..da387a3906 100644 --- a/uv.lock +++ b/uv.lock @@ -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"