Add openvino_hggd module: HGGD OpenVINO iGPU enablement by deepaks2 · Pull Request #1069 · openvinotoolkit/openvino_contrib

deepaks2 · 2026-06-10T05:16:48Z

This pull request adds a new module, openvino_hggd, to the contrib repository, providing OpenVINO support for the HGGD (Hybrid Grasp Detection and Generation) model, a state-of-the-art Efficient Heatmap-Guided 6-Dof Grasp Detection in Cluttered Scenes. The changes include documentation updates, usage guides, code for HGGD (Hybrid Grasp Detection and Generation) with OpenVINO custom extensions, and supporting files for building and running the module.

Major additions and updates:

New HGGD (Hybrid Grasp Detection and Generation) module integration

Added the openvino_hggd module
Implemented OpenVino extensions for pointcloud, written optimized kernels for farthest point sampling, knn and ball query.

Documentation and Usage Guides
Added EXPORT_AND_INFERENCE_GUIDE.md with step-by-step instructions for exporting HGGD models to OpenVINO, building required OpenVINO extensions, and running inference/evaluation.
Supporting files

Added a .gitignore to the module directory to exclude build artifacts, logs, and virtual environments.

- Intel iGPU inference for HGGD (Hybrid Grasp Detection and Generation) via OpenVINO with custom C++/OpenCL point cloud extensions - Includes: export_models.py, infer.py, run.sh, setup.sh, ov_gpu_extensions/ - Users must copy customgraspnetAPI, dataset, models from upstream https://github.com/THU-VCLab/HGGD (see README) - Add 15 missing third-party entries to third-party-programs.txt (HGGD, graspnetAPI, dexnet, scipy, scikit-image, matplotlib, pandas, tensorboardX, torchsummary, transforms3d, trimesh, autolab_core, cvxopt, grasp-nms, Pillow)

allnes

Code review of the new openvino_hggd module (Python/ML + C++/OpenCL custom extension for OpenVINO GPU) plus the third-party-programs.txt additions. Read the affected files in full; didn't build (needs an Intel GPU, a torch+xpu conda env, the built .so, and the copied upstream GraspNet dataset).

Verdict: changes needed, with a few items for maintainer/legal review. The blockers are on the licensing side, not the code:

grasp-nms is listed as MIT but the pinned package declares no license (see inline).
dexnet is a UC Berkeley non-commercial license, incompatible with Apache-2.0 (see inline).
cvxopt is GPLv3 (see inline).

Two items without a line to attach to:

The module isn't listed in the root README.md where every other module is — please add an entry. The PR description also mentions an EXPORT_AND_INFERENCE_GUIDE.md, but there's no such file in the diff; either add it or drop the claim.
The feature commit has no Signed-off-by. If this repo enforces DCO, rebase with git commit -s so the check passes.

No tests are wired into CI (no job, no labeler.yml/CODEOWNERS entry), and export_models.py prints max_diff(PT vs OV) without asserting a threshold — worth adding a numeric gate. Detailed findings inline.

allnes · 2026-07-21T18:14:25Z

+
+-------------------------------------------------------------
+
+dexnet (Berkeley AUTOLAB)


dexnet ships under the UC Berkeley research/non-commercial license — it grants use for "educational, research, and not-for-profit purposes" only and points to Berkeley OTL for commercial licensing. That is not compatible with Apache-2.0, which grants commercial use and redistribution.

Two things:

It's not in setup.sh's install list and no committed code imports it, so it's unclear why it's listed at all — if the module doesn't depend on dexnet, drop the entry.

If the upstream pipeline copied from HGGD pulls it in at runtime, this needs legal review before the module lands: a non-commercial dependency in an Apache-2.0 module is a real conflict.

allnes · 2026-07-21T18:14:25Z

+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#define WORKGROUP_SIZE 256
+#define MAX_N 8192


min_dist lives in __local memory sized MAX_N=8192, and the loop clamps to n_eff = min(N, MAX_N). If a padded cloud exceeds 8192 points (fps rounds N up to the next power of two ≥1024, and all_points_num=25600 → N_pad up to 32768), FPS silently samples only the first 8192 points with no error — a wrong result that looks successful. The CPU FPSSingle::evaluate has no such cap, so CPU and GPU diverge for large N.

Either assert/return an error when N > MAX_N, or keep min_dist in global memory. What's the max N that actually reaches FPS in the pipeline?

fixed and tested

allnes · 2026-07-21T18:14:25Z

+
+#define WORKGROUP_SIZE 256
+
+__kernel void fps_single_v2_kernel(


This kernel (and fps_with_lengths.cl, knn_points_single_v2.cl) isn't referenced by pointcloud_ops_gpu_v2.xml or any code — the XML uses the _tiled and _optimized variants. Dead files add maintenance surface, and fps_single_v2/fps_with_lengths read the sample index back from an FP16 output, which is exactly the precision trap the optimized variants avoid. Drop them unless they're intentionally kept as references.

fixed and tested

- Remove hard 8192-point cap from fps_single_optimized.cl and fps_with_lengths_optimized.cl; recompute distances per iteration so large padded clouds (e.g., 32768 points) are handled correctly - Add PyTorch3D BSD-3 license/copyright notices to the affected OpenCL kernels and to pointcloud_ops.cpp - Sync pointcloud_ops_gpu_v2.xml for the updated kernels

Zero out invalid idx slots instead of clamping them to 0 Clamps were returning features[0] for invalid positions; masked assignment now matches MaskedGather::evaluate behavior

Remove -march=native to avoid ISA lock-in on the build machine Drop -ffast-math to preserve IEEE FP semantics needed for distance comparisons

Align export_models.py docstring with setup.sh and README.md so copy-paste setup instructions work

Replace lingering hggd_xpu references with hggd_intel for consistency

Signed-off-by: Deepak Soma Reddy <deepak.s@intel.com>

FP16 silently rounds integer indices above 2048, but HGGD point clouds reach 25000+ points and KNN/BallQuery/FPS indices are packed as floats. Switch shim default to f32; keep f16 opt-in only for small N.

Add NODE_VALIDATION_CHECK so partial_sort does not walk past dist_idx.end() when K > N2, which is UB and would read out of bounds in the output loop.

allnes · 2026-08-04T19:05:57Z

+    transforms3d==0.4.2 \
+    trimesh==4.11.4 \
+    autolab_core==1.1.1 \
+    cvxopt==1.3.3 \


cvxopt is GPL-3.0, a strong copyleft license, and setup.sh pins it as a runtime dependency of the hggd_intel environment. Shipping a module whose documented runtime pulls in GPL-3.0 needs maintainer/legal review for an Apache-2.0 project. Please confirm whether cvxopt is actually imported at runtime for this pipeline or only comes in transitively via graspnetAPI/grasp-nms; if it's not a direct requirement, drop the explicit pin, and if it is required, flag the GPL interaction for review before merge.

allnes · 2026-08-04T19:05:57Z

+
+-------------------------------------------------------------
+
+dexnet (Berkeley AUTOLAB)


This dexnet entry is a UC Berkeley non-commercial license (use limited to educational/research/not-for-profit; commercial use requires a signed agreement). Nothing in this module imports dexnet and setup.sh doesn't install it (only autolab_core), so either the entry is spurious and should be removed, or dexnet is a real transitive dependency, in which case a non-commercial license attached to a component of an Apache-2.0 distribution needs to be escalated rather than silently listed. Please clarify which, and don't ship a non-commercial dependency as a hard requirement.

allnes · 2026-08-04T19:05:57Z

@@ -0,0 +1,91 @@
+/*
+ * Copyright (C) 2018-2026 Intel Corporation


This kernel is one of the four actually installed and referenced by pointcloud_ops_gpu_v2.xml, but unlike knn_points_tiled.cl, fps_single_optimized.cl and fps_with_lengths_optimized.cl it carries only the Intel/Apache header with no PyTorch3D BSD-3 attribution, even though it implements the same ball_query algorithm ported from PyTorch3D and BallQuery in pointcloud_ops.cpp does carry that attribution. Add the same BSD-3 attribution block here for consistency.

allnes · 2026-08-04T19:05:57Z

+            model = self._get_model(key, xml)
+            result = model({"points": points_padded, "lengths": lengths})[0]
+        else:
+            # Use original FPSSingle kernel (pad N similarly)


When lengths is None, fps() zero-pads points up to a power-of-two bucket and then dispatches FPSSingle, which has no notion of valid length. The padded (0,0,0) rows become real candidates in the farthest-point search: for a cloud whose points sit away from the origin, an all-zero pad row is frequently the farthest point and gets selected, and its index (>= N) is only masked afterwards by np.clip(idx, 0, N-1), silently collapsing it to a valid-but-wrong index. Either always route padded inputs through FPSWithLengths passing the true N as length, or have FPSSingle ignore rows beyond N.

allnes · 2026-08-04T19:05:57Z

+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#define TILE_SIZE 256
+#define MAX_K 64


The GPU kernels store neighbors in fixed private arrays best_dists[MAX_K]/best_idx[MAX_K] with MAX_K=64, but nothing enforces K <= 64. validate_and_infer_types() only checks k <= N2 for KNN and nothing at all for BallQuery's k. The default k=32 is fine, but the op accepts arbitrary k and any K > 64 writes past the private arrays, which is out-of-bounds on GPU. Add a K <= 64 NODE_VALIDATION_CHECK in validate_and_infer_types for KNNPoints(Single) and BallQuery(Single), or size the arrays from K.

allnes · 2026-08-04T19:05:57Z

 * [**ollama_OpenVINO**](./modules/ollama_openvino)
 * [**openvino-langchain**](./modules/openvino-langchain): LangChain.js integrations for OpenVINO™
 * [**OpenVino BEVFusion**](./modules/openvino_bevfusion): Check the [INSTRUCTIONS](./modules/openvino_bevfusion/EXPORT_AND_INFERENCE_GUIDE.md) for detailed usage and build instructions.
+* [**OpenVino HGGD**](./modules/openvino_hggd): Check the [INSTRUCTIONS](./modules/openvino_hggd/README.md) for detailed usage and build instructions.


HGGD is added to the "Additional build instructions" section only, but not to the main module overview list at the top of the file where every other module (including openvino_bevfusion) is listed with a one-line description. Add the one-line overview entry so the module appears in the canonical list too.

allnes · 2026-08-04T19:05:57Z

@@ -0,0 +1,134 @@
+/*


This kernel, along with fps_with_lengths.cl and knn_points_single_v2.cl, is neither referenced by pointcloud_ops_gpu_v2.xml nor installed by CMakeLists, so it's dead code. fps_single_v2.cl documents a MAX_N=8192 truncation that the shipped kernels were specifically rewritten to remove, so keeping it around is misleading. Remove these three files (or wire them in if they're intended). knn_points_single_v2.cl additionally lacks the PyTorch3D attribution.

allnes · 2026-08-04T19:05:57Z

+# Transform Operations (pytorch3d API compatible)
+# ═══════════════════════════════════════════════════════════════════════════════
+
+def euler_angles_to_matrix(euler: torch.Tensor, convention: str = "XYZ") -> torch.Tensor:


euler_angles_to_matrix and matrix_to_quaternion reproduce the pytorch3d.transforms API and the standard Shepperd quaternion-from-matrix branches. The math is standard, but since these are drop-in replacements for pytorch3d and the rest of the module attributes PyTorch3D wherever code was ported, a one-line "API-compatible with pytorch3d.transforms" note here would keep provenance consistent. Separately, the init.py docstring references module names that don't exist (ov_shim_v3, pointcloud_ops_native).

allnes · 2026-08-04T19:05:57Z

+        if key in self._models:
+            return self._models[key]
+
+        model = self.core.read_model(xml.encode('utf-8'))


Each op builds an IR XML by string-formatting shape and attribute values and reads it via read_model on a weightless bytes blob. Is this intended over building an ov.Model programmatically from opset ops? The string path is brittle (radius/K interpolated into XML, one compiled model cached per unique (B, N, K)) and bypasses the ONNX frontend that ov_extension.cpp registers. Does inference actually exercise the ov::frontend::OpExtension paths, or only the direct OpExtension via read_model?

allnes · 2026-08-04T19:05:57Z

+OPENVINO_CREATE_EXTENSIONS(
+    std::vector<ov::Extension::Ptr>({
+        // KNN Points - K-nearest neighbors (multi-output)
+        std::make_shared<ov::OpExtension<HGGDExtension::KNNPoints>>(),


Six multi-output ops (KNNPoints, BallQuery, FPS, MaskedGather, GatherMaxPool, PointGather) are registered here and fully implemented in pointcloud_ops.cpp, but no .py or .xml uses them; only the *Single and FPSWithLengths variants are wired up. Are the multi-output ops part of the intended public surface or leftovers from development? If unused, dropping them shrinks the extension and the maintenance/attribution surface.

deepaks2 requested a review from a team as a code owner June 10, 2026 05:16

github-actions Bot added the category: build OpenVINO cmake script / infra label Jun 10, 2026

deepaks2 and others added 4 commits June 10, 2026 10:54

Merge branch 'master' into hggd_ov

50d3278

Merge branch 'master' into hggd_ov

727b3ef

Merge branch 'master' into hggd_ov

8f24be7

Merge branch 'master' into hggd_ov

b616135

allnes reviewed Jul 21, 2026

View reviewed changes

allnes self-assigned this Jul 21, 2026

deepaks2 added 10 commits July 29, 2026 17:52

Merge branch 'master' into hggd_ov

8b00fba

Fix masked_gather -1 index handling in GPU shim

389568e

Zero out invalid idx slots instead of clamping them to 0 Clamps were returning features[0] for invalid positions; masked assignment now matches MaskedGather::evaluate behavior

Make compiler flags portable

c46855f

Remove -march=native to avoid ISA lock-in on the build machine Drop -ffast-math to preserve IEEE FP semantics needed for distance comparisons

Normalize conda env name to hggd_intel

a739643

Align export_models.py docstring with setup.sh and README.md so copy-paste setup instructions work

Normalize conda env name in setup.sh comments

bb5ddcd

Replace lingering hggd_xpu references with hggd_intel for consistency

Adding hggd to README

68750d8

Signed-off-by: Deepak Soma Reddy <deepak.s@intel.com>

Force FP32 for index-carrying GPU point-cloud ops

fa257ee

FP16 silently rounds integer indices above 2048, but HGGD point clouds reach 25000+ points and KNN/BallQuery/FPS indices are packed as floats. Switch shim default to f32; keep f16 opt-in only for small N.

Guard KNN K <= N2 in validate_and_infer_types

f15e408

Add NODE_VALIDATION_CHECK so partial_sort does not walk past dist_idx.end() when K > N2, which is UB and would read out of bounds in the output loop.

Fix install file list to real kernel/config names

05c7ff4

deepaks2 requested a review from allnes July 30, 2026 11:41

allnes requested changes Aug 4, 2026

View reviewed changes


		-------------------------------------------------------------

		dexnet (Berkeley AUTOLAB)


		#define WORKGROUP_SIZE 256

		__kernel void fps_single_v2_kernel(

		@@ -0,0 +1,91 @@
		/*
		* Copyright (C) 2018-2026 Intel Corporation

Uh oh!

Conversation

deepaks2 commented Jun 10, 2026

Uh oh!

allnes left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants