samples: add OpenVX graph pipelining CPU+GPU sample - #7
Open
simonCatBot wants to merge 18 commits into
Open
Conversation
simonCatBot
force-pushed
the
feature/pipelining-sample
branch
from
August 6, 2026 17:38
8b51ead to
ab8c164
Compare
|
Tested on Radeon RX 9070 XT / Ryzen 9 9950X 16-Core Processor, using TheRock Nightly 10.1.0 2026-08-05:
Both backends pass the checksum-equality test and show speed-up when pipelining is enabled. One note is that CPU outperformed HIP backend, @kiritigowda is that reasonable? |
simonCatBot
force-pushed
the
feature/pipelining-sample
branch
3 times, most recently
from
August 6, 2026 22:06
4b11c3c to
00322f5
Compare
Add samples/c_samples/pipelining, a self-contained demonstration of the vx_khr_pipelining extension on a mixed CPU+GPU workload. The same vision graph (RGB -> ColorConvert -> ChannelExtract(Y) -> Box3x3 -> U8) is run synchronously with vxProcessGraph (--pipeline 0) and asynchronously with QUEUE_AUTO enqueue/dequeue (--pipeline 1). Both paths report identical per-frame checksums so correctness can be verified, while the pipelined path prints a higher fps to make the performance benefit visible. Also update samples/c_samples/README.md with build/run instructions.
simonCatBot
force-pushed
the
feature/pipelining-sample
branch
from
August 6, 2026 23:30
00322f5 to
7f75313
Compare
… to fix non-deterministic HIP pipe checksums The heavy preset previously chained two GPU Box3x3 nodes through a shared virtual intermediate. In QUEUE_AUTO pipelined mode on the HIP backend this produced slightly different aggregate checksums on every run, while sync and the CPU backend were stable. Replace the heavy graph with a single Gaussian3x3 node. It keeps the preset compute-heavy enough to show pipelining speed-up, and the sync and pipelined paths now produce identical, deterministic checksums on both HIP and CPU backends. The README is updated to describe the new heavy graph and to remove the caveat about sync/pipe checksum differences.
… preset change - Update pipelining --help to describe the heavy preset as Gaussian3x3 instead of the old Box3x3 -> Box3x3 chain. - Remove the misleading compare-mode message that claimed heavy sync/pipe checksum differences were expected. - Fix pipelining_hybrid header comment: the CPU stage is Threshold (U8) -> Box3x3, not a U1 conversion. - README: mention both Box3x3 and Gaussian3x3 as the heavier convolutions.
These GDF files demonstrate the proposed GDF syntax for vx_khr_pipelining support in runvx. They are not executable today because runvx does not yet implement QUEUE_AUTO, graph parameters, or enqueue/dequeue primitives. Files: - utilities/runvx/gdf/pipelining/README.md - utilities/runvx/gdf/pipelining/pipelining_box3x3.gdf - utilities/runvx/gdf/pipelining/pipelining_gaussian3x3.gdf - utilities/runvx/gdf/pipelining/pipelining_hybrid.gdf See ROCm#1734 for the enhancement request.
…EUE_MANUAL over-scheduling in graph pipelining (ROCm#1729) * Fix cross-target hand-off, queued OpenCL args, and QUEUE_MANUAL over-scheduling Three defects kept a pipelined graph from being usable once its work was split across the CPU and the GPU. A CPU node only recorded that it had written a buffer when that buffer already had device memory. A graph that hands its output to another graph usually writes an image no GPU node has bound yet, so the record was skipped, and the device memory reserved later by the consuming graph was read without ever being uploaded. The CPU node now records the write regardless, and a device buffer reserved for a queued reference starts out needing an upload. On OpenCL, buffer kernel arguments are bound when the graph is verified. That holds for an object that stays put, and the code already re-bound delay slots, but a queued graph parameter names a different reference on every execution. Every execution after the first therefore wrote the reference chosen at verification time. Queued references are now re-bound the same way delay slots are. In QUEUE_MANUAL, vxScheduleGraph runs every complete set of enqueued references it finds rather than a single one, so a request could carry out work that later requests were made for; those requests then found empty queues and reported an error for work that had in fact been done. The extra executions are credited and spent by the requests that follow, which keeps the error for the case it was written for - an application that scheduled the graph without enqueueing a reference everywhere. Two tests cover the parts that were silently wrong: a two-graph GPU-after- CPU chain whose result is only exact if the hand-off uploads, and a QUEUE_MANUAL run that checks each enqueued set runs once, with correct data and no failure logged. Co-authored-by: Cursor <cursoragent@cursor.com> * Make the QUEUE_MANUAL execution accounting explicit Review pointed out that crediting the executions a request performs beyond the one it asked for leaves credits behind that nothing is waiting for, so a later request with empty queues reports success instead of the missing references. A request cannot be limited to a single execution instead. The conformance suite's GraphPipeline.testManualSchedule enqueues one set per buffer, schedules the graph once, and requires every one of those sets to have run by the time it dequeues; one execution per request fails 7 of the 109 pipelining tests. With the number of executions decoupled from the number of requests that way, an application that enqueues N sets and schedules once and then N-1 times, and one that enqueues and schedules N times where a request ran several sets, make exactly the same calls in the same order. Nothing in the queues separates them, so what separates a request whose work an earlier one already did from a request the application enqueued nothing for is whether an execution is there to claim. The counter now says that directly: every request claims one execution, the executions no request has claimed are carried, and reconfiguring the graph starts the count again. Test 27 now covers both directions, a request per enqueued set and one request for every enqueued set, and checks that a request with nothing enqueued for it is still reported as an error. Co-authored-by: Cursor <cursoragent@cursor.com> --------- Co-authored-by: Cursor <cursoragent@cursor.com>
* HIP: add RGB/RGBX Remap nearest and bilinear kernels Add GPU Remap support for packed RGB and RGBX images in the HIP backend. Previously Remap only had a GPU implementation for VX_DF_IMAGE_U8. This change adds: - VX_KERNEL_AMD_REMAP_U24_U24_NEAREST - VX_KERNEL_AMD_REMAP_U32_U32_NEAREST Implementation details: - Added enum IDs, kernel list entries, and API declarations. - Implemented per-channel nearest and bilinear interpolation HIP kernels for 3-byte RGB and 4-byte RGBX in hipvx/geometric_kernels.cpp. - Wired validate and hip_execute in ago_kernel_api.cpp for the existing Remap_U24_U24_Bilinear and Remap_U32_U32_Bilinear kernels plus the new nearest variants. - Updated agoDramaDivideRemapNode to dispatch the new kernels for RGB/RGBX with undefined borders and nearest or bilinear interpolation. - Added GDF test cases under tests/amd_openvx_gdfs/geometric/. Tested on gfx1151 (Radeon 8060S) with runvx GPU affinity. * HIP Remap U8: use round-half-up packing to match CPU/OpenCV The CPU SSE path computes bilinear interpolation with fixed-point weights and rounds with `+32 >> 6` (round-half-up). The HIP U8 remap kernel used `__builtin_amdgcn_cvt_pk_u8_f32`, which rounds to nearest-even, so .5 ties (e.g. 62.5 -> 62) could differ from CPU by one. Add a `hip_pack_half_up()` helper that converts each float with `(uint)(x + 0.5f)` and clamps to [0,255], then wire it into the U8 bilinear remap kernels. RGB/RGBX bilinear kernels already used per-channel `(uchar)(x + 0.5f)` and are unaffected. Validation on 16x16 probes and the OpenVX 1280x720 demo mirror map shows CPU/HIP match exactly for U8, RGB, and RGBX. * HIP Remap RGB: fix output store to write all 24 bytes The RGB bilinear kernel packed 8 pixels into two uint3 registers (24 bytes) but then wrote only 20 bytes: out[0].x/out[0].y (8 bytes), out[1].x/out[1].y (8 bytes), and out[1].z (4 bytes). This dropped out[0].z (bytes 8-11), which contains the blue channel of pixel 2 and all of pixel 3, causing systematic corruption visible on the demo map. Write all six uints out[0].x/y/z and out[1].x/y/z to the destination so the 24 RGB bytes are preserved exactly. After the fix the HIP GPU output matches the CPU reference on U8, RGB, and RGBX bilinear remap. * OpenVX CPU: add scalar, SSE, and AVX2 RGB/RGBX Remap fallbacks Add CPU implementations for the new RGB/RGBX Remap kernels so the PR provides both CPU and HIP coverage for U8/RGB/RGBX remap: - HafCpu_Remap_U24_U24_Bilinear / Nearest - HafCpu_Remap_U32_U32_Bilinear / Nearest Includes SSE and AVX2 fast paths plus a scalar tail, source-coordinate clamping to keep the 2x2 neighborhood in-bounds, and wiring in ago_kernel_api.cpp so CPU affinity dispatches correctly. Validation: all Remap_* GDFs pass on both CPU and GPU affinity for both the CPU-only and HIP builds; CPU and HIP outputs match exactly on the 1280x720 demo map. * Remap RGB/RGBX: add constant-border CPU and HIP kernels Extend the PR so RGB and RGBX remap support VX_BORDER_MODE_CONSTANT, not just VX_BORDER_MODE_UNDEFINED. This brings them in line with the existing U8 remap kernels and lets them match OpenCV-style BORDER_CONSTANT behavior. Changes: - Added HIP kernels: Hip_Remap_RGB*_Bilinear_Constant / Nearest_Constant, plus constant-border bilinear sampling helpers. - Added CPU scalar fallbacks: HafCpu_Remap_U24/U32 *_Constant. - Registered new kernel IDs, list entries, and API declarations. - Wired CPU and HIP execute paths and drama divide dispatch. - Added GDF tests: Remap_RGB*_Bilinear_Constant and *_Nearest_Constant. Validation: - All 12 remap GDFs pass on CPU-only and HIP builds (CPU/GPU affinity). - CPU and HIP outputs match on the 1280x966 dewarp demo for U8 and RGB constant-border bilinear. * Remap RGB/RGBX: fix OOB stores, misaligned loads, and clamp; address review - hip_pack_half_up: clamp in the float domain before the uint cast, so a negative round-off value pins to 0 instead of underflowing to 255 - RGB/RGBX Remap kernels stored a full 8-pixel block unconditionally, writing past the row end and overflowing the row stride when dstWidth is not a multiple of 8; store partial tail blocks pixel-wise - RGBX bilinear read source pixels via 16-byte uint4 loads off 4-byte-aligned addresses (misaligned/UB); sample byte-wise through hip_bilinear_sample_RGBX - RGB bilinear dropped two unused uint2 loads that were also misaligned - GDF tests: RGB (RGB2) uniform-image fills use 24-bit values per repo convention, not 32-bit - CHANGELOG: add the RGB/RGBX Remap GPU support entry Co-Authored-By: Claude <noreply@anthropic.com> * Remap RGB/RGBX CPU: fix boundary over-reads in SSE/AVX bilinear paths The SIMD bilinear paths clamped the sampled coordinates but not the width of the loads, so near the allocation edge they read past the buffer: - RGB 4-pixel block packed each 3-byte pixel with a 16-byte loadu; at the bottom-right pixel this read up to 15 bytes past the image. Pack exactly three bytes via AgoLoadRGB24. - RGB 2-pixel tail read pixels through vx_uint32*, taking a 4th byte off the final 3-byte pixel; same three-byte packing applied. - RGB and RGBX 2-pixel tails decoded the remap through AgoRemapComputeOffsets_SSE, which loads four entries (and pMapY_X+1 a further four); with fewer than four entries left in a row this over-read the remap table. Decode one entry at a time via AgoRemapComputeOffset1. Behaviour is unchanged for in-bounds pixels; only the out-of-allocation reads are removed. Verified with a clean HIP build of libopenvx. Co-Authored-By: Claude <noreply@anthropic.com> --------- Co-authored-by: Kiriti Gowda <kiriti.gowda@gmail.com> Co-authored-by: Kiriti Gowda <kiritigowda@gmail.com> Co-authored-by: Claude <noreply@anthropic.com>
* CI - Use rocm 10.1 * CI - update comments
…Cm#1750) The software emulation of amd_bytealign (used on platforms that do not support cl_amd_media_ops, e.g. non-AMD OpenCL devices) masked the byte offset argument with & 31 instead of & 3. The hardware amd_bytealign uses only the low 2 bits of the third argument to select a byte offset (0-3) into the concatenated {src0:src1} pair; masking with & 31 allows shifts of up to 31 bytes, reading the wrong bytes whenever the offset is >= 4. This produced incorrect results in kernels that call amd_bytealign with a full (non-masked) byte index -- e.g. the ORB Gaussian scale kernel used by vxGaussianPyramidNode -- when running on a non-AMD device. The native hardware path is unaffected because the emulation is only injected when cl_amd_media_ops is unavailable. amd_bitalign (bit funnel shift, valid range 0-31 bits) correctly keeps its & 31 mask.
…m#1752) Bumps [rocm-docs-core](https://github.com/ROCm/rocm-docs-core) from 1.39.0 to 1.40.0. - [Release notes](https://github.com/ROCm/rocm-docs-core/releases) - [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md) - [Commits](ROCm/rocm-docs-core@v1.39.0...v1.40.0) --- updated-dependencies: - dependency-name: rocm-docs-core dependency-version: 1.40.0 dependency-type: direct:production update-type: version-update:semver-minor ... Signed-off-by: dependabot[bot] <support@github.com> Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
The OpenCL codegen for the ORB Gaussian scale kernel (used by vxGaussianPyramidNode with VX_SCALE_PYRAMID_ORB) performed the horizontal filter in place: each work-item read neighbouring image pixels from the local buffer (lbuf_ptr[ilx & ~3]) and, in the same pass with no barrier, wrote its packed 16-bit horizontal sums back to the same local buffer (((uint2*)lbuf_ptr)[lx]). Because the read footprint of one work-item overlaps the write footprint of its neighbours (e.g. lx=2 reads local bytes [8..19] while lx=1 writes [8..15]), a work-item could read sums another work-item had already written instead of the original pixels. On a lock-step SIMD device the reads happen to complete before the writes, so the result was correct; on a scalar device where work-items are not executed in lock-step (e.g. a CPU OpenCL driver such as PoCL) the writes raced ahead and corrupted the result -- the right-hand output columns were wrong. Fix: compute both the main-row and the +16-row horizontal sums into private registers (isum, isum2), then barrier, then store. All reads of the shared local buffer now complete before any work-item writes to it. This also removes the previous reuse of lbuf_ptr across the two stores. No behavioural change on AMD hardware (verified: GaussianPyramid ORB still 12/12 on gfx1151). Fixes the 8 GaussianPyramid ORB VX_BORDER_UNDEFINED OpenVX-CTS cases that failed on PoCL; full vision conformance now passes 5824/5824 on a PoCL CPU device.
…ROCm#1747) Follow-up to ROCm#1723. Six of seven memcpy(x, desc, s - desc) sites in the GDF data-descriptor parser were bounds-checked; this closes the gaps in the same class that were left unfixed: - ago_util.cpp agoGetDataFromDescription (threshold:): add the missing `(size_t)(s - desc) >= sizeof(thresh_type)` guard before the memcpy, matching the six sibling sites. - ago_interface.cpp agoUpdateLine: was expanding $VAR / $! values into the caller's fixed char line[2048] via unbounded strcpy; add a line_size capacity parameter and bound both expansions (callers pass sizeof(line)), mirroring the agoUpdateN hardening. - ago_interface.cpp def-macro: strncpy(macro_name, arg[1], sizeof) left the 256-byte buffer unterminated; copy sizeof-1 and NUL-terminate. - ago_util.cpp agoGetDataName: char tmp[512] strcpy of a name written up to MAX_MODULE_NAME_SIZE (1024); size tmp to MAX_MODULE_NAME_SIZE and bound the copy. Verified: full build + ctest 100% (32/32) against ROCm 7.2, gfx1100. Co-authored-by: Kiriti Gowda <kiritigowda@gmail.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Adds a new C++ sample under
samples/c_samples/pipeliningthat demonstrates the OpenVXvx_khr_pipeliningextension on a mixed CPU+GPU workload.The sample runs the same vision graph in two modes:
--pipeline 0— synchronousvxProcessGraphloop (baseline).--pipeline 1—QUEUE_AUTOpipelined enqueue/dequeue with multiple input/output buffers.Graph:
RGB -> ColorConvert -> ChannelExtract(Y) -> Box3x3 -> U8Both paths compute an identical per-frame checksum aggregate, and the pipelined path reports a higher fps so the benefit is immediately visible.
What was added
samples/c_samples/pipelining/pipelining.cpp— the sample application.samples/c_samples/pipelining/CMakeLists.txt— build file matching the existingcannysample.samples/c_samples/README.mdwith build/run instructions.Verification performed
Built and tested locally on a Radeon gfx1151 GPU and on the CPU-only backend:
Both backends pass the checksum-equality test and show measurable speed-up when pipelining is enabled.
Checklist
--pipeline 0and--pipeline 1produce the same aggregate checksum.--pipeline 1reports higher fps than--pipeline 0.