feat(reconstruction): add resource-aware matcher runtime and Vulkan ORB backend

AGENTS.md

@@ -63,6 +63,24 @@ disponible. Ne jamais annoncer une vérification non exécutée.
 - Éviter le code-golf et garder des fonctions confortables à relire dans Neovim.
 - Commenter les invariants non évidents plutôt que paraphraser le code.
 
+## Lisibilité du code
+
+Ces règles s'appliquent aussi bien au code du dépôt qu'au code temporaire
+créé pour les benchmarks, diagnostics, expérimentations ou probes matériels.
+
+- Viser environ 100 colonnes par ligne.
+- Limite absolue : 120 colonnes.
+- Ne pas compresser plusieurs instructions logiques sur une même ligne.
+- Aucun code-golf, même pour un prototype.
+- Les fichiers temporaires C, C++, GLSL et les scripts doivent rester
+  lisibles et auditables.
+- Découper les grosses fonctions en responsabilités claires lorsque cela
+  améliore la compréhension.
+- Les benchmarks et prototypes doivent pouvoir être relus et débogués
+  facilement.
+- Un reformatage ne doit jamais modifier le comportement.
+- Ces règles s'appliquent également aux fichiers créés sous `/tmp`.
+
 ## Long run OpenCode
 
 - `.opencode/work/current_ticket.md` est la mémoire durable du ticket et doit

README.md

@@ -100,6 +100,7 @@ Acquisitions
 - [Registry des types de tâches](docs/architecture/task_kind_registry.md)
 - [File de tâches](docs/architecture/task_queue.md)
 - [Resource Governor](docs/architecture/resource_governor.md)
+- [Pipeline sensible aux ressources](docs/architecture/resource_aware_pipeline.md)
 - [Intégration Scheduler ↔ Governor](docs/architecture/scheduler_resource_integration.md)
 - [Pipeline de reconstruction](docs/architecture/reconstruction_pipeline.md)
 - [Persistance](docs/architecture/persistence.md)
@@ -110,6 +111,7 @@ Acquisitions
 - [Candidate Pair](docs/architecture/candidate_pair.md)
 - [Match Result](docs/architecture/match_result.md)
 - [Matcher](docs/architecture/matcher.md)
+- [Backend Vulkan ORB](docs/architecture/vulkan_matcher.md)
 - [Viewer](docs/architecture/viewer.md)
 - [Revue des fondations](docs/architecture/foundation_review.md)
 
@@ -124,6 +126,7 @@ Acquisitions
 - [Build](docs/development/build.md)
 - [Tests](docs/development/testing.md)
 - [Concurrence](docs/development/concurrency.md)
+- [Profil de performance de la machine cible](docs/performance/target_hardware.md)
 - [OpenCode long run](docs/development/opencode_long_run.md)
 
 ### Roadmap
@@ -149,8 +152,10 @@ Pour les changements sensibles à la mémoire ou à la concurrence, ajouter ASan
 
 Lardon3D est en développement actif. La persistance des tâches, le catalogue,
 le Feature Store multipasse, le Visual Index ORB, Candidate Pair Generator
-et Matcher v1 sont implémentés. DAG générique, vérification géométrique,
+et Matcher v1 sont implémentés. Le runtime Feature + Matcher emploie des tâches
-SfM et viewer restent des tickets séparés planifiés.
+durables, de petits lots, le Resource Governor interactif et un hot path Vulkan
+ORB exact avec fallback CPU. DAG générique,
+vérification géométrique, SfM et viewer restent des tickets séparés planifiés.
 
 ## Licence
 

docs/architecture/matcher.md

@@ -76,13 +76,30 @@ matérialisée.
 Le Match File complet est sérialisé dans un buffer heap borné à 98336 octets et
 écrit par un unique `write_exact`, puis synchronisé une fois. Les mesures locales
 restent dans `.opencode/work/current_ticket.md`, pas dans ce contrat canonique.
-À 8192 features, le coût dominant mesuré reste l'évaluation exacte des
+À 8192 features, le coût CPU dominant reste l'évaluation exacte des distances
-distances dans `cv::BFMatcher::knnMatch`; v1 ne remplace pas OpenCV ni BFMatcher.
+dans `cv::BFMatcher::knnMatch`. ORB peut remplacer ce seul hot path par Vulkan ;
+SIFT et RootSIFT restent intégralement sur BFMatcher CPU.
 
-Le Matcher n'est pas encore un task kind autonome. Lorsqu'il est orchestré par
+`matcher.run` v1 orchestre le Matcher sans connaître son backend interne. La
-une tâche, celle-ci doit utiliser l'unique Resource Governor existant avec une
+tâche persiste uniquement la configuration, l'identité des Feature Sets à
-estimation couvrant ce working set; aucune seconde logique de budget n'est
+sélectionner et un curseur Candidate Pair. Une paire est atomique et publiée
-introduite ici.
+immédiatement. Les lots 1/2/4/8 sont séparés par checkpoint et
+`task_sequence_break()`. La tâche utilise l'unique Resource Governor avec une
+estimation couvrant ce working set ; aucune seconde logique de budget n'est
+introduite. Sa ligne durable `matcher_tasks` appartient au schéma Project DB
+v11 ; le Match Result reste le contrat publié en v10.
+
+### Backend Vulkan ORB
+
+La frontière évaluée remplace uniquement KNN Hamming par un compute top-2 : un
+thread GPU par feature A parcourt B, conserve deux indices/distances et applique
+le tie-break du plus petit index. Elle ne matérialise jamais A×B. Lowe,
+canonicalisation et persistance restent communs. Sur Radeon 780M, la parité
+top-2 avec OpenCV est exacte et le gain warm est supérieur à 90 % à 4096/8192.
+Le backend de production conserve exactement cette frontière. Sa parité entière
+permet au CPU et à Vulkan de partager l'identité persistante. La sélection est
+une politique runtime ; le CPU reste le fallback portable. Le contrat détaillé
+est décrit dans [vulkan_matcher.md](vulkan_matcher.md).
 
 ## Déterminisme et fingerprint
 

docs/architecture/project_database.md

@@ -1,7 +1,8 @@
 # Base de données projet Lardon3D
 
-> Version courante : **v10**. La migration transactionnelle v9→v10 ajoute
+> Version courante : **v11**. La migration transactionnelle v10→v11 ajoute
-> la table `match_results` pour le Match Result Model. La migration v8→v9
+> `matcher_tasks` pour la tâche Matcher durable. La version v10 publiée ajoute
+> uniquement `match_results` pour le Match Result Model. La migration v8→v9
 > ajoute la table `candidate_pair_generate_tasks` pour la tâche durable
 > Candidate Pair. La migration v7→v8 ajoute la table `candidate_pairs` pour
 > le sous-système Candidate Pair.
@@ -306,10 +307,37 @@ CREATE TABLE candidate_pair_generate_tasks(
 - `lardon3d_project_db_record_candidate_pair_generate_task()` — UPSERT checkpoint
 - `lardon3d_project_db_load_candidate_pair_generate_task()` — SELECT par task_id
 
-## Schéma v10 implémenté
+## Schéma v10 publié
 
-La migration v9→v10 ajoute la table `match_results` pour le Match Result
+La migration v9→v10 ajoute uniquement `match_results` pour le Match Result
-Model :
+Model. Son schéma et ses invariants restent inchangés.
+
+## Schéma v11 implémenté
+
+La migration v10→v11 ajoute `matcher_tasks`. Cette table conserve uniquement
+la configuration immutable et le curseur durable :
+
+```sql
+CREATE TABLE matcher_tasks(
+    task_id INTEGER PRIMARY KEY REFERENCES tasks(task_id) ON DELETE CASCADE,
+    after_candidate_pair_id INTEGER NOT NULL
+        CHECK(after_candidate_pair_id>=0),
+    feature_extractor_kind TEXT NOT NULL,
+    feature_extractor_version INTEGER NOT NULL
+        CHECK(feature_extractor_version>0),
+    feature_parameter_fingerprint BLOB NOT NULL
+        CHECK(length(feature_parameter_fingerprint)=32),
+    matcher_kind INTEGER NOT NULL CHECK(matcher_kind BETWEEN 0 AND 2),
+    ratio_threshold REAL NOT NULL
+        CHECK(ratio_threshold>0.0 AND ratio_threshold<1.0)
+);
+```
+
+Le curseur est le dernier `candidate_pair_id` checkpointé. Il n'implique ni
+continuité des IDs ni liste persistée de Candidate Pairs. La configuration ne
+peut pas changer lors d'un UPSERT ; seul le curseur avance.
+
+Le Match Result ci-dessous reste le contrat publié de v10 :
 
 ```sql
 CREATE TABLE match_results(
@@ -365,18 +393,22 @@ CREATE INDEX match_results_feature_set_b_idx
 - `lardon3d_project_db_load_match_result()` — SELECT par ID
 - `lardon3d_project_db_find_match_result()` — SELECT par (candidate_pair_id, feature_set_id_a, feature_set_id_b, matcher_kind, matcher_version, parameter_fingerprint)
 - `lardon3d_project_db_list_match_results()` — SELECT paginé ORDER BY id
+- `lardon3d_project_db_record_matcher_task()` — UPSERT configuration/curseur
+- `lardon3d_project_db_load_matcher_task()` — SELECT par task_id
 
 ## Ouverture et migrations
 
-Une DB vide reçoit directement le schéma v7 dans une transaction
+Une DB vide reçoit la chaîne de schémas jusqu'à v11 dans une transaction
 `BEGIN IMMEDIATE`. Une DB v1 reçoit transactionnellement les colonnes nullable
 `task_kind` et `task_kind_version`, puis les migrations v2→v3. Les anciennes lignes restent
 `NULL/NULL`, sans type inventé et sans perte des projets, tâches, checkpoints ou
 artefacts. Une interruption ou erreur provoque un rollback complet. Les DB v1,
-v2, v3, v4, v5 et v6 sont migrées séquentiellement vers v7. Une version future est refusée et une DB contenant
+v2, v3, v4, v5, v6, v7, v8, v9 et v10 sont migrées séquentiellement vers v11.
+Une v10 publiée est validée comme telle avant que v10→v11 crée
+`matcher_tasks` ; son absence n'est donc pas une corruption. Une version future est refusée et une DB contenant
 des tables sans métadonnée de version est considérée corrompue. La fonction
 interne de migration applique uniquement la chaîne séquentielle connue jusqu'à
-v7 ; une valeur hors de 1..7 est refusée.
+v11 ; une valeur hors de 1..11 est refusée.
 
 Migration v1→v2 exacte, exécutée entre `BEGIN IMMEDIATE` et `COMMIT` :
 
@@ -513,9 +545,9 @@ ouvert.
 
 ## Statut
 
-**IMPLEMENTED** — SQLite système, schéma v10 et migrations v1→v2→v3→v4→v5→v6→v7→v8→v9→v10, identité
+**IMPLEMENTED** — SQLite système, schéma v11 et migrations séquentielles
-projet, transactions tâche+checkpoint, pagination de reprise et artefacts
+v1→v2→v3→v4→v5→v6→v7→v8→v9→v10→v11, identité projet, transactions
-génériques.
+tâche+checkpoint, pagination de reprise et artefacts génériques.
 
 **IMPLEMENTED** — ouverture/fermeture avec le projet, identité INI/DB cohérente,
 publication de checkpoints par le projet et inventaire de reprise validé.

docs/architecture/reconstruction_pipeline.md

@@ -56,7 +56,7 @@ content-addressed. Les états Feature/Matching/Reconstruction restent planifiés
 **Extension v1A :** ORB reste la passe coarse et l'unique entrée du Visual
 Index ORB-LSH. SIFT/RootSIFT F32×128 sont des passes précises indépendantes,
 suivies d'une consolidation spatiale intra-image qui ne mélange jamais les
-descriptors. Candidate Pair Generator et matching restent planifiés.
+descriptors.
 
 ---
 
@@ -79,24 +79,27 @@ updates incrémentales et query top-K bornée. Le générateur de paires reste p
 |--------|-------------|
 | **Sources de paires candidates** | (1) Visual index : paires visuellement proches. (2) Proximité temporelle. (3) Scan set commun. (4) Géométrie approximative (si GPS/IMU disponible). |
 | **Matching coûteux limité** | Le nombre de paires soumises au matching géométrique (étape F) doit être borné. Le candidate generator filtre et classe pour ne garder que les paires les plus prometteuses. |
-| **Persistance** | Les paires candidates sont persistées dans la table `candidate_pairs` (Project DB v10). Ordre canonique : `image_id_a < image_id_b`. Self-pairs interdits. Unicité garantie. |
+| **Persistance** | Les paires candidates sont persistées dans la table `candidate_pairs` (Project DB v8). Ordre canonique : `image_id_a < image_id_b`. Self-pairs interdits. Unicité garantie. |
 | **Déterminisme** | Pour mêmes entrées et configuration, le générateur produit les mêmes paires dans le même ordre. |
 | **Idempotence** | L'exécution répétée ne crée pas de doublons. |
 
-**Statut :** IMPLEMENTED v1 — génération single-source depuis Visual Index,
+**Statut :** IMPLEMENTED v1 — génération single-source et batch depuis Visual
-persistance, canonicalisation et idempotence. La génération batch projet
+Index, persistance, canonicalisation, idempotence et tâche durable.
-et l'intégration tâche durable restent planifiées.
 
 ---
 
-### F. Matching et vérification géométrique
+### F. Matching de descripteurs
 
 | Aspect | Description |
 |--------|-------------|
-| **Distinction des étapes** | (1) *Feature matching* : appariement brut des descripteurs entre deux images. (2) *Geometric verification* : estimation de la transformation rigide (RANSAC ou équivalent) et validation de la compatibilité épipolaire. |
+| **Matcher v1** | ORB/Hamming exact CPU ou Vulkan, SIFT/RootSIFT L2 CPU, KNN k=2 et Lowe, sans vérification géométrique. |
-| **Validation ou rejet** | Une paire validée produit une *edge* dans le graphe de visibilité. Une paire rejetée est marquée comme telle pour éviter les retraitements inutiles. |
+| **Persistance** | Match Result NO_MATCH/MATCHED et Match File content-addressed validé. |
+| **Orchestration** | `matcher.run` traite les Candidate Pairs par pages et lots durables de 1/2/4/8. |
 
-**Statut :** PLANNED — aucune implémentation existante.
+**Statut :** IMPLEMENTED v1 — Matcher, Match Store, reprise idempotente et Task
+durable. La vérification géométrique reste l'étape suivante, non commencée.
+Le Match Result appartient à Project DB v10 et la tâche durable `matcher.run`
+à Project DB v11.
 
 ---
 

docs/architecture/resource_aware_pipeline.md

@@ -0,0 +1,86 @@
+# Pipeline Feature + Matcher sensible aux ressources
+
+## Contrat portable
+
+Une unité lourde ne démarre qu'avec une réservation active. Elle termine son
+petit travail courant sans être tuée sur une mesure instantanée, publie le
+résultat atomiquement, checkpoint, libère ses buffers, puis repasse par le
+Governor avant la séquence suivante. Les files restent bornées et le swap n'est
+jamais ajouté au budget de travail.
+
+Le mode normal est interactif : il réserve de la RAM et des threads logiques au
+desktop. Les signaux d'admission combinent `MemAvailable`, charge CPU, PSI CPU,
+PSI mémoire, PSI I/O et deltas `pswpin`/`pswpout`. Un seuil dépassé empêche une
+nouvelle admission ; il ne rompt pas une réservation saine déjà active.
+
+Le Governor maintient trois zones. GREEN emploie le lot adapté normal. La soft
+floor RAM, un PSI au seuil ou un premier intervalle avec swap actif produit
+YELLOW et interdit toute croissance. Deux observations de pression
+consécutives, ou `MemAvailable` sous la hard floor, produisent RED et suspendent
+toute admission. Le premier snapshot swap établit seulement la baseline.
+
+La récupération possède deux phases distinctes : trois observations saines
+font `RED → YELLOW`, puis trois nouvelles observations saines font
+`YELLOW → GREEN`. Après RED, le plafond de lot reste 1. En GREEN, trois
+observations saines sont nécessaires à chaque palier `1 → 2 → 4 → 8`. Une
+nouvelle pression réinitialise cette progression. Cette mémoire est
+process-local, bornée et protégée par le mutex du Governor.
+
+## Feature Extraction
+
+ORB est déjà une tâche durable par image : source validée, extraction,
+publication Feature Store, métadonnées DB, checkpoint terminal et libération du
+buffer. Le batch vaut donc une image et la granularité de reprise est une image.
+Le worker unique et la file bornée fournissent la backpressure actuelle.
+
+OpenCV est configuré une seule fois avant le démarrage des workers. La tâche
+réserve le nombre réel de threads OpenCV au lieu d'annoncer artificiellement un
+thread pendant qu'une primitive interne en utilise davantage.
+
+## Matcher
+
+`matcher.run` v1 est une tâche durable. Son unité atomique est une Candidate
+Pair et son lot vaut 1, 2, 4 ou 8 paires. La tâche page la DB par
+`candidate_pair_id`, sans supposer des IDs continus, et ne conserve jamais la
+liste entière. Chaque paire publie immédiatement son Match Result avant que le
+curseur ne soit avancé en mémoire.
+
+Project DB v10 porte le Match Result publié. La migration transactionnelle
+v10→v11 ajoute uniquement `matcher_tasks`, qui porte la configuration et ce
+curseur durable.
+
+Après chaque lot, la tâche persiste le curseur, checkpoint, puis appelle
+`lardon3d_task_sequence_break()`. Pause et annulation sont vérifiées avant
+chaque paire et entre les lots. Un crash après publication mais avant le
+checkpoint revoit la paire : le Matcher réutilise alors le Match Result et ne
+recalcule pas les descripteurs.
+
+## GPU et files
+
+La Radeon 780M est UMA : toute mémoire GPU compte aussi comme pression RAM. Un
+backend GPU emploie un unique job actif, des dispatchs courts, puis publie avant
+de continuer. Vulkan 1.4.354 énumère la 780M RADV et une file compute dédiée. Le
+backend ORB top-2 de production possède un contexte lazy réutilisable, 640 Kio
+de buffers bornés et un fallback CPU exact. Le CPU reste le fallback portable
+si Vulkan est absent, incompatible ou désactivé pour la session.
+
+## Profil interactif 8845HS mesuré
+
+- budget CPU Lardon3D : 12 threads logiques, 4 réservés au desktop ;
+- réserve `MemAvailable` : un quart de la RAM, environ 3,8 Gio ;
+- hard floor `MemAvailable` : un huitième, environ 1,9 Gio ;
+- Feature workers : 1 ; batch : 1 image ;
+- Matcher workers : 1 ; lots adaptatifs 1, 2, 4 ou 8 Candidate Pairs ;
+- profondeur de la Task Queue : 64 tâches légères, un seul callback actif ;
+- PSI CPU avg10 : nouvelle admission suspendue à 20 % ;
+- PSI mémoire avg10 : nouvelle admission suspendue à 1 % ;
+- PSI I/O avg10 : seuil existant 80 %.
+
+Le benchmark Matcher 8192 mesure environ 70 ms ORB et 135 ms SIFT à 12 threads,
+contre 68 ms et 127 ms à 16 threads : le profil interactif abandonne environ
+3–7 % de latence isolée pour réserver quatre threads logiques au desktop.
+
+## Limites
+
+Le profil maximal explicite et les pools multi-workers restent hors périmètre.
+SIFT/RootSIFT et Feature Extraction Vulkan restent hors de ce contrat.

docs/architecture/resource_governor.md

@@ -9,7 +9,26 @@ OpenCV peut employer son parallélisme interne ; aucun état global n'est modifi
 
 ## Responsabilité
 
-Le Resource Governor est l'unique propriétaire des budgets (RAM, GPU, CPU, IO). Il arbitre les ressources disponibles et calcule les lots adaptatifs pour chaque tâche.
+Le Resource Governor est l'unique propriétaire des budgets (RAM, GPU, CPU,
+IO). Il arbitre les ressources disponibles et calcule les lots adaptatifs pour
+chaque tâche.
+
+Le profil interactif par défaut conserve un quart de la RAM détectée et un
+quart des threads logiques pour le système hôte. Sur 16 Gio/16 threads, cela
+donne environ 3,8 Gio et 4 threads de headroom. Une nouvelle admission attend
+également lorsque PSI CPU `some avg10` atteint 20 %, ou PSI mémoire 1 %. Ces
+signaux n'interrompent jamais le petit job déjà réservé.
+
+La soft floor vaut un quart et la hard floor un huitième de la RAM détectée. La
+soft floor place le Governor au minimum en YELLOW ; la hard floor le place
+immédiatement en RED. Le premier delta swap entre deux snapshots produit
+YELLOW. Un second delta consécutif produit RED. Le premier snapshot ne constitue
+qu'une baseline et n'est jamais interprété comme une activité récente.
+
+La récupération interdit `RED → GREEN` : trois observations saines produisent
+RED vers YELLOW, puis trois autres YELLOW vers GREEN. Le plafond reste 1 pendant
+ces phases. Une fois GREEN, chaque groupe de trois observations saines double
+le plafond : 1, 2, 4, 8, puis les paliers supérieurs utiles aux autres kinds.
 
 ## API principale
 
@@ -34,6 +53,7 @@ Le Resource Governor est l'unique propriétaire des budgets (RAM, GPU, CPU, IO).
 - `lardon3d_resource_governor_record_batch()` - Enregistrer les métriques d'un lot
 - `lardon3d_resource_governor_generation()` - Obtenir la génération actuelle
 - `lardon3d_resource_governor_wait_for_change()` - Attendre un changement
+- `lardon3d_resource_governor_pressure()` - Lire GREEN, YELLOW ou RED
 
 ## Invariants
 
@@ -94,11 +114,15 @@ Le Resource Governor est l'unique propriétaire des budgets (RAM, GPU, CPU, IO).
   réelle du lot ; `peak_memory_bytes == 0` signifie « mesure inconnue ».
   Chaque séquence interroge le Visual Index pour jusqu'à 64 Feature Sets et
   persiste les paires candidates avec idempotence.
-- Le Matcher v1 possède un working set contrôlé inférieur à environ 10 Mio au
+- `matcher.run` réserve douze threads CPU, un slot IO et un working set
+  contrôlé inférieur à environ 10 Mio au
   maximum SIFT/RootSIFT (8 Mio de descripteurs contigus, KNN `k=2`, sorties et
-  fichier bornés), hors scratch interne OpenCV. Il n'est pas encore exposé comme
+  fichier bornés), hors scratch interne OpenCV. Ses lots sont bornés à 1, 2, 4
-  task kind autonome; sa future orchestration devra réserver CPU+IO et cette
+  ou 8 Candidate Pairs et chaque paire libère ses buffers avant la suivante.
-  estimation via ce Governor, sans budget parallèle.
+  Quand le profil détecte un GPU et que le runtime possède le backend ORB, la
+  réservation ajoute un slot GPU et 640 Kio. Sur UMA ces 640 Kio sont aussi
+  débités du budget RAM. Sans GPU/backend, l'estimation reste CPU-only afin que
+  le fallback portable ne soit jamais refusé artificiellement.
 
 ## Limites actuelles
 

docs/architecture/task_kind_registry.md

@@ -64,4 +64,8 @@ depuis `image_id` et ses paramètres bornés.
 + scanset_filter + exclude_same_asset` depuis `candidate_pair_generate_tasks`
 et reconstruit un contexte boundé.
 
-**PLANNED** — kinds de matching et reconstruction.
+**IMPLEMENTED** — `matcher.run`, version 1, recharge la configuration Matcher,
+l'identité Feature Set et le curseur `after_candidate_pair_id`. Il traite une
+Candidate Pair atomique à la fois dans des lots bornés à huit, checkpoint le
+curseur et repasse par le Governor entre les lots. La table durable
+`matcher_tasks` est introduite par Project DB v11, après le Match Result v10.

docs/architecture/task_system.md

@@ -105,6 +105,12 @@ candidates avec idempotence, checkpoint après chaque lot et repasse par le
 Governor via `sequence_break`. La reprise est idempotente avec le curseur
 `after_feature_set_id` rechargé depuis la DB.
 
+**IMPLEMENTED** — `matcher.run` traite une Candidate Pair atomique à la fois,
+par lots adaptatifs de 1, 2, 4 ou 8. Il persiste le curseur
+`after_candidate_pair_id`, checkpoint après publication de chaque lot et
+effectue une rupture de séquence avant le suivant. Une paire repassée après un
+crash est réutilisée par son Match Result.
+
 Le chemin de production de l'import ne possède plus de thread ni de drapeau
 d'annulation privés. Son wrapper TUI ne fait qu'enqueue/cancel/observer la
 tâche générique. Chaque callback traite un lot borné, checkpoint hors mutex de

docs/architecture/vulkan_matcher.md

@@ -0,0 +1,68 @@
+# Backend Vulkan ORB v1
+
+## Frontière de correction
+
+Le backend Vulkan remplace uniquement la recherche exacte des deux plus proches
+voisins ORB/Hamming. La lecture Feature Store, le filtre Lowe, l'ordre canonique,
+le Match File, le SHA-256, le Match Result et la reprise restent communs au
+backend CPU OpenCV.
+
+Pour chaque feature A, le shader parcourt B et conserve seulement les deux
+couples `(distance, feature_index_b)` minimaux. La distance est une somme exacte
+de huit `bitCount` sur les 32 octets ORB. L'ordre total est distance croissante,
+puis index B croissant. Aucune matrice A×B n'est matérialisée. CPU et Vulkan
+doivent donc produire le même top-2, la même décision Lowe et les mêmes octets
+persistés. Le backend ne participe pas à l'identité scientifique.
+
+## Contexte et sélection
+
+Le runtime possède un contexte opaque, initialement sans device. La première
+paire éligible initialise Vulkan une fois ; les paires suivantes réutilisent le
+device, la file, le pipeline, le command buffer et trois buffers
+bornés. Un mutex impose un dispatch à la fois. Une famille compute sans graphics
+est préférée, avec fallback vers toute famille compute compatible.
+
+Le sélecteur utilise le travail `feature_count_a × feature_count_b`. Sous le
+seuil mesuré de `768 × 768` comparaisons, OpenCV reste utilisé afin d'éviter le
+coût fixe du dispatch. Le seuil est une politique d'exécution et ne modifie ni
+fingerprint ni Match Result. L'initialisation est lazy : une application qui ne
+matche aucun grand couple ORB ne paie aucun cold start.
+
+## Mémoire et pannes
+
+Les buffers maximaux contiennent 256 Kio pour A, 256 Kio pour B et 128 Kio pour
+8192 sorties top-2, soit 640 Kio de payload, hors petits objets du driver. Une
+mémoire host-visible, cohérente et cached est préférée sur UMA, car elle réduit
+nettement le coût CPU de copie/readback observé sur RADV. Le backend sait
+appliquer flush/invalidate lorsque le type retenu n'est pas cohérent.
+
+Une absence de loader/device/queue ou un échec d'initialisation est mémorisé et
+utilise le CPU sans nouvelle tentative par paire. Une panne de soumission ou un
+device lost désactive Vulkan pour la session ; la paire courante est reprise sur
+CPU avant toute publication. Un Match Result n'est jamais créé depuis une
+sortie GPU partielle.
+
+Le job est synchrone au niveau du Matcher. Il soumet sur l'unique queue détenue
+par le contexte puis attend cette queue, sans `vkDeviceWaitIdle` sur le chemin
+normal. Le Resource Governor admet la tâche avant le callback ; en RED aucun
+nouveau batch ne démarre, tandis qu'une paire déjà soumise finit et se publie.
+Le worker unique et le mutex interdisent plusieurs dispatchs concurrents en v1.
+
+## Shader et déterminisme
+
+Le shader parcourt les indices B dans l'ordre croissant. Il ne remplace le best
+ou le second que pour une distance strictement inférieure : la première égalité,
+donc le plus petit index, est conservée. Ce contrat a été comparé exactement à
+`BFMatcher(NORM_HAMMING).knnMatch(k=2)` pour 0, 1, 2, 16, 64, 256, 1024, 4096 et
+8192 descriptors, ainsi que pour les égalités et descriptors identiques. Le
+Matcher complet produit les mêmes Match File et SHA-256 sur CPU et Vulkan.
+
+## Build portable
+
+Meson active Vulkan seulement si le loader de développement et `glslc` sont
+disponibles. Le GLSL versionné est compilé en SPIR-V puis incorporé dans un
+header généré ; le SPIR-V est contrôlé par `spirv-val` pendant la validation
+lorsque l'outil est présent. Sans cette chaîne,
+le même code compile avec un stub indisponible et tous les Matchers restent CPU.
+L'option Meson `-Dvulkan_orb=disabled` force ce build CPU-only ; `auto` est le
+défaut portable et `enabled` exige explicitement le loader et `glslc`.

docs/performance/target_hardware.md

@@ -0,0 +1,78 @@
+# Profil de performance de la machine cible
+
+Ce document décrit une politique de performance mesurée. Il ne modifie aucun
+contrat de correction du Matcher, du Match Store ou du Match Result.
+
+## Cible principale actuelle
+
+- AMD Ryzen 7 8845HS, Zen 4, 8 cœurs et 16 threads SMT ;
+- Radeon 780M à mémoire système partagée ;
+- 16 Gio de RAM et zram d'environ 6 Gio ;
+- Arch Linux, Clang 22, OpenCV 5.0.0.
+
+Le build OpenCV observé emploie TBB 2023.1, expose 16 threads par défaut et les
+chemins SIMD jusqu'à AVX512-SKX. Il a été compilé avec OpenCL, mais
+`cv::ocl::haveOpenCL()` retourne faux. Vulkan énumère en revanche
+`AMD Radeon 780M Graphics (RADV PHOENIX)`, API 1.4.354, avec une file compute
+dédiée et de la mémoire UMA host-visible/cohérente.
+
+Le runtime actuel de Lardon3D possède un worker. Le profil interactif conserve
+12 threads OpenCV process-wide, quatre threads logiques pour le desktop et un
+seul Matcher actif. Lorsque
+les pools multi-workers seront introduits, la cible de départ recommandée est
+deux Matchers avec huit threads OpenCV chacun pour une charge mixte. Les mesures
+montrent toutefois que quatre Matchers à quatre threads favorisent SIFT et les
+cas 4096, tandis que deux à huit favorisent ORB 8192. Le Governor devra donc
+choisir à partir de la classe de charge, pas d'une constante universelle. Le
+nombre de threads OpenCV devra être réglé une fois au démarrage :
+`cv::setNumThreads()` est une configuration globale et ne doit jamais être
+modifiée concurremment par des workers.
+
+Quatre Matchers avec un ou deux threads chacun dégradent fortement les grands
+cas ORB. Quatre fois quatre threads augmente le working set et la variance, mais
+peut améliorer SIFT soutenu. Le Governor devra compter les threads OpenCV dans
+le budget CPU afin d'éviter `workers × threads` supérieur aux 16 threads
+matériels.
+
+Le working set directement contrôlé d'un Matcher SIFT/RootSIFT reste inférieur
+à environ 10 Mio, hors scratch TBB/OpenCV. Même plusieurs Matchers restent loin
+de la pression mémoire sur 16 Gio ; les mesures n'ont produit aucun swap-in ni
+swap-out. La limite pratique observée est le CPU, pas la RAM.
+
+## Fallback portable
+
+Le build portable conserve les réglages Meson génériques et ne force ni
+`-march=native`, ni OpenCL, ni un nombre de threads spécifique au 8845HS. Sur une
+autre machine, laisser OpenCV choisir son backend et limiter l'orchestration à
+un Matcher reste le fallback sûr. Une future configuration multi-worker devra
+être dérivée du profil matériel par le Resource Governor, sans second scheduler.
+
+## Méthode et portée
+
+`benchmark-matcher` utilise des Feature Sets synthétiques persistés, un warm-up
+et sept répétitions dont il rapporte la médiane. Les mesures absolues sont
+locales et sensibles au boost et à la température ; la décision repose surtout
+sur le scaling et le débit soutenu. La campagne a été exécutée avec le governor
+Linux `powersave`; après charge soutenue, 61–72 °C ont été observés et aucun
+swap-in/swap-out. Le coût dominant reste l'évaluation exacte des distances dans
+`cv::BFMatcher::knnMatch`.
+
+Le backend Vulkan ORB de production est borné, utilise une invocation par query
+et ne matérialise aucune matrice A×B. Sur la 780M, le workgroup 32 est le meilleur
+des quatre candidats 32/64/128/256 à 4096 et 8192. Le chemin complet warm mesure
+environ 0,18 ms à 256, 0,45 ms à 1024, 1,6 ms à 4096 et 4,0 ms à 8192, contre
+environ 0,10, 0,96, 14,7 et 59,6 ms pour BFMatcher CPU lors de la campagne
+production. L'initialisation lazy mesurée vaut environ 129–136 ms. Le seuil
+`feature_count_a × feature_count_b >= 768²` évite le GPU pour les petits travaux.
+
+La mémoire permanente directement contrôlée vaut 640 Kio de payload. Les tests
+de parité couvrent exactement le top-2 jusqu'à 8192, le Match File complet et le
+fallback CPU. Ces nombres décrivent la machine mesurée et ne sont pas un contrat
+portable de latence.
+
+Un run soutenu de 5000 dispatchs mixtes 1024/4096/8192/4096 a traité environ
+1232 paires/s en 4,06 s. Le processus de benchmark complet a culminé à environ
+202 Mio RSS ; les compteurs `pswpin` et `pswpout` sont restés à zéro. Après le
+run, PSI CPU `some avg10` valait 0,14 %, PSI mémoire et I/O 0 %, avec 76 °C CPU
+et 64 °C au bord GPU. Ces mesures sont des observations ponctuelles, pas des
+seuils du Resource Governor.

include/lardon3d/app_state.h

@@ -10,6 +10,7 @@ typedef struct Lardon3DImageView Lardon3DImageView;
 typedef struct Lardon3DTaskQueue Lardon3DTaskQueue;
 typedef struct Lardon3DResourceGovernor Lardon3DResourceGovernor;
 typedef struct Lardon3DProjectDb Lardon3DProjectDb;
+typedef struct Lardon3DOrbVulkanBackend Lardon3DOrbVulkanBackend;
 
 typedef enum {
     LARDON3D_SCREEN_HOME = 0,
@@ -35,6 +36,7 @@ typedef struct {
     Lardon3DHardwareProfile hardware_profile;
     Lardon3DResourceGovernor *resource_governor;
     Lardon3DProjectDb *project_db;
+    Lardon3DOrbVulkanBackend *orb_vulkan_backend;
     size_t recovery_inspected;
     size_t recovery_resumed;
     size_t recovery_skipped;

include/lardon3d/feature_extractor.h

@@ -68,6 +68,10 @@ typedef enum {
   LARDON3D_FEATURE_EXTRACT_ERROR
 } Lardon3DFeatureExtractResult;
 
+/* Configuration OpenCV process-wide. À appeler avant le démarrage des workers. */
+bool lardon3d_feature_opencv_configure_threads(unsigned int threads);
+unsigned int lardon3d_feature_opencv_thread_count(void);
+
 bool lardon3d_feature_extractor_parameters_valid(
     const Lardon3DFeatureExtractorParameters *parameters);
 void lardon3d_feature_extractor_parameter_fingerprint(

include/lardon3d/matcher.h

@@ -3,6 +3,7 @@
 
 #include <lardon3d/feature_store.h>
 #include <lardon3d/match_file.h>
+#include <lardon3d/orb_vulkan_backend.h>
 #include <lardon3d/project_db.h>
 
 #ifdef __cplusplus
@@ -40,6 +41,8 @@ typedef struct {
     uint64_t publication_ns;
     uint64_t database_ns;
     uint64_t total_ns;
+    bool used_vulkan;
+    bool vulkan_fallback;
 } Lardon3DMatcherStats;
 
 typedef enum {
@@ -65,6 +68,15 @@ Lardon3DMatcherResult lardon3d_matcher_run(
     const char *match_file_path,
     Lardon3DMatcherStats *stats);
 
+Lardon3DMatcherResult lardon3d_matcher_run_with_backend(
+    const char *project_path,
+    const Lardon3DProjectDbFeatureSet *feature_set_a,
+    const Lardon3DProjectDbFeatureSet *feature_set_b,
+    const Lardon3DMatcherParams *params,
+    const char *match_file_path,
+    Lardon3DOrbVulkanBackend *backend,
+    Lardon3DMatcherStats *stats);
+
 Lardon3DMatcherResult lardon3d_matcher_match_and_publish(
     const char *project_path,
     Lardon3DProjectDb *database,
@@ -84,6 +96,17 @@ Lardon3DMatcherResult lardon3d_matcher_match_and_publish_profiled(
     Lardon3DProjectDbMatchResult *result,
     Lardon3DMatcherStats *stats);
 
+Lardon3DMatcherResult lardon3d_matcher_match_and_publish_with_backend(
+    const char *project_path,
+    Lardon3DProjectDb *database,
+    const Lardon3DProjectDbCandidatePair *pair,
+    const Lardon3DProjectDbFeatureSet *feature_set_a,
+    const Lardon3DProjectDbFeatureSet *feature_set_b,
+    const Lardon3DMatcherParams *params,
+    Lardon3DOrbVulkanBackend *backend,
+    Lardon3DProjectDbMatchResult *result,
+    Lardon3DMatcherStats *stats);
+
 #ifdef __cplusplus
 }
 #endif

include/lardon3d/matcher_task.h

@@ -0,0 +1,38 @@
+#ifndef LARDON3D_MATCHER_TASK_H
+#define LARDON3D_MATCHER_TASK_H
+
+#include <stdbool.h>
+#include <stdint.h>
+
+#include <lardon3d/app_state.h>
+#include <lardon3d/matcher.h>
+#include <lardon3d/task_kind_registry.h>
+
+#define LARDON3D_MATCHER_TASK_KIND "matcher.run"
+
+enum {
+  LARDON3D_MATCHER_TASK_KIND_VERSION = 1,
+  LARDON3D_MATCHER_TASK_MINIMUM_BATCH = 1,
+  LARDON3D_MATCHER_TASK_MAXIMUM_BATCH = 8,
+};
+
+typedef struct {
+  char feature_extractor_kind[LARDON3D_PROJECT_DB_KIND_CAPACITY];
+  uint32_t feature_extractor_version;
+  unsigned char feature_parameter_fingerprint[LARDON3D_PROJECT_DB_SHA256_SIZE];
+  Lardon3DMatcherParams matcher;
+} Lardon3DMatcherTaskConfiguration;
+
+Lardon3DTask *lardon3d_project_create_matcher_task(
+    Lardon3DAppState *state,
+    const Lardon3DMatcherTaskConfiguration *configuration, uint64_t *task_id);
+
+bool lardon3d_project_enqueue_matcher_task(
+    Lardon3DAppState *state,
+    const Lardon3DMatcherTaskConfiguration *configuration, uint64_t *task_id);
+
+bool lardon3d_matcher_task_reconstruct(
+    const Lardon3DTaskDurableSnapshot *snapshot, void *context,
+    Lardon3DTaskKindBinding *binding);
+
+#endif

include/lardon3d/orb_vulkan_backend.h

@@ -0,0 +1,63 @@
+#ifndef LARDON3D_ORB_VULKAN_BACKEND_H
+#define LARDON3D_ORB_VULKAN_BACKEND_H
+
+#include <stdbool.h>
+#include <stddef.h>
+#include <stdint.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+typedef struct Lardon3DOrbVulkanBackend Lardon3DOrbVulkanBackend;
+
+enum {
+  LARDON3D_ORB_VULKAN_PERMANENT_BUFFER_BYTES = 640 * 1024,
+};
+
+typedef struct {
+  uint32_t neighbor_count;
+  uint32_t best_index;
+  uint32_t best_distance;
+  uint32_t second_index;
+  uint32_t second_distance;
+} Lardon3DOrbTop2;
+
+typedef enum {
+  LARDON3D_ORB_VULKAN_OK = 0,
+  LARDON3D_ORB_VULKAN_UNAVAILABLE,
+  LARDON3D_ORB_VULKAN_FAILED,
+  LARDON3D_ORB_VULKAN_INVALID_ARGUMENT
+} Lardon3DOrbVulkanResult;
+
+typedef struct {
+  bool available;
+  bool initialized;
+  bool dedicated_compute_queue;
+  char device_name[256];
+  uint32_t workgroup_size;
+  uint64_t permanent_payload_bytes;
+  uint64_t initialization_ns;
+  uint64_t dispatch_ns;
+  uint64_t gpu_ns;
+} Lardon3DOrbVulkanInfo;
+
+Lardon3DOrbVulkanBackend *lardon3d_orb_vulkan_backend_create(void);
+void lardon3d_orb_vulkan_backend_destroy(Lardon3DOrbVulkanBackend *backend);
+
+bool lardon3d_orb_vulkan_should_use(uint32_t feature_count_a,
+                                    uint32_t feature_count_b);
+
+Lardon3DOrbVulkanResult lardon3d_orb_vulkan_top2(
+    Lardon3DOrbVulkanBackend *backend, const unsigned char *descriptors_a,
+    uint32_t feature_count_a, const unsigned char *descriptors_b,
+    uint32_t feature_count_b, Lardon3DOrbTop2 *output, size_t output_capacity);
+
+bool lardon3d_orb_vulkan_backend_info(Lardon3DOrbVulkanBackend *backend,
+                                      Lardon3DOrbVulkanInfo *info);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif

include/lardon3d/project.h

@@ -70,6 +70,9 @@ Lardon3DProjectTaskCheckpointResult lardon3d_project_checkpoint_visual_index_upd
 Lardon3DProjectTaskCheckpointResult lardon3d_project_checkpoint_candidate_pair_generate_task(
     Lardon3DAppState *state, const Lardon3DTask *task,
     const Lardon3DProjectDbCandidatePairGenerateTask *parameters);
+Lardon3DProjectTaskCheckpointResult lardon3d_project_checkpoint_matcher_task(
+    Lardon3DAppState *state, const Lardon3DTask *task,
+    const Lardon3DProjectDbMatcherTask *parameters);
 Lardon3DProjectDbResult lardon3d_project_list_recoverable(Lardon3DAppState *state,
                                                           const Lardon3DTaskKindRegistry *registry,
                                                           uint64_t after_task_id,

include/lardon3d/project_db.h

@@ -9,7 +9,7 @@
 #include <lardon3d/task.h>
 
 enum {
-  LARDON3D_PROJECT_DB_SCHEMA_VERSION = 10,
+  LARDON3D_PROJECT_DB_SCHEMA_VERSION = 11,
   LARDON3D_PROJECT_DB_ID_CAPACITY = 65,
   LARDON3D_PROJECT_DB_KIND_CAPACITY = 65,
   LARDON3D_PROJECT_DB_PATH_CAPACITY = 4096,
@@ -131,6 +131,16 @@ typedef struct {
   int64_t created_at;
 } Lardon3DProjectDbCandidatePair;
 
+typedef struct {
+  uint64_t task_id;
+  uint64_t after_candidate_pair_id;
+  char feature_extractor_kind[LARDON3D_PROJECT_DB_KIND_CAPACITY];
+  uint32_t feature_extractor_version;
+  unsigned char feature_parameter_fingerprint[LARDON3D_PROJECT_DB_SHA256_SIZE];
+  int matcher_kind;
+  float ratio_threshold;
+} Lardon3DProjectDbMatcherTask;
+
 typedef struct {
   uint64_t match_result_id;
   uint64_t candidate_pair_id;
@@ -463,6 +473,14 @@ Lardon3DProjectDbResult lardon3d_project_db_record_candidate_pair_generate_task(
 Lardon3DProjectDbResult lardon3d_project_db_load_candidate_pair_generate_task(
     Lardon3DProjectDb *database, uint64_t task_id,
     Lardon3DProjectDbCandidatePairGenerateTask *parameters);
+Lardon3DProjectDbResult lardon3d_project_db_record_matcher_task(
+    Lardon3DProjectDb *database, const Lardon3DTaskDurableSnapshot *snapshot,
+    const char *task_kind, uint32_t task_kind_version,
+    const Lardon3DProjectDbCheckpoint *checkpoint,
+    const Lardon3DProjectDbMatcherTask *parameters, int64_t updated_at);
+Lardon3DProjectDbResult lardon3d_project_db_load_matcher_task(
+    Lardon3DProjectDb *database, uint64_t task_id,
+    Lardon3DProjectDbMatcherTask *parameters);
 
 Lardon3DProjectDbResult lardon3d_project_db_create_match_result(
     Lardon3DProjectDb *database, uint64_t candidate_pair_id, uint64_t feature_set_id_a,

include/lardon3d/resource_governor.h

@@ -18,14 +18,23 @@ typedef struct Lardon3DResourceReservation Lardon3DResourceReservation;
 
 typedef struct {
     uint64_t system_memory_reserve_bytes;
+    uint64_t emergency_memory_floor_bytes;
     uint64_t gpu_memory_reserve_bytes;
     unsigned int system_cpu_reserve;
     double maximum_cpu_load_ratio;
+    double maximum_cpu_pressure_avg10;
+    double maximum_memory_pressure_avg10;
     double maximum_io_pressure_avg10;
     unsigned int gpu_slot_capacity;
     unsigned int io_slot_capacity;
 } Lardon3DResourcePolicy;
 
+typedef enum {
+    LARDON3D_RESOURCE_PRESSURE_GREEN = 0,
+    LARDON3D_RESOURCE_PRESSURE_YELLOW,
+    LARDON3D_RESOURCE_PRESSURE_RED
+} Lardon3DResourcePressure;
+
 typedef enum {
     LARDON3D_RESOURCE_TASK_GENERAL = 0,
     LARDON3D_RESOURCE_TASK_IMPORT,
@@ -187,6 +196,9 @@ bool lardon3d_resource_governor_wait_for_change(
 const char *lardon3d_resource_decision_name(
     Lardon3DResourceDecisionKind kind
 );
+Lardon3DResourcePressure lardon3d_resource_governor_pressure(
+    Lardon3DResourceGovernor *governor
+);
 /* Enregistre les métriques d'un lot terminé pour l'adaptation dynamique
  * de la taille des lots futurs. batch_size est le nombre d'éléments dont le
  * traitement a été validé dans ce lot. peak_memory_bytes == 0 signifie que

include/lardon3d/resource_snapshot.h

@@ -18,8 +18,15 @@ typedef struct {
     double cpu_load_1m;
     double cpu_load_5m;
     double cpu_load_15m;
+    bool cpu_pressure_known;
+    double cpu_pressure_avg10;
+    bool memory_pressure_known;
+    double memory_pressure_avg10;
     bool io_pressure_known;
     double io_pressure_avg10;
+    bool swap_activity_known;
+    uint64_t swap_pages_in;
+    uint64_t swap_pages_out;
 } Lardon3DResourceSnapshot;
 
 bool lardon3d_resource_snapshot_capture(

include/lardon3d/task_kind_registry.h

@@ -8,6 +8,7 @@
 
 typedef struct Lardon3DProjectDb Lardon3DProjectDb;
 typedef struct Lardon3DResourceGovernor Lardon3DResourceGovernor;
+typedef struct Lardon3DOrbVulkanBackend Lardon3DOrbVulkanBackend;
 
 enum {
   LARDON3D_TASK_KIND_REGISTRY_MAX = 64,
@@ -39,6 +40,7 @@ typedef struct {
   const char *project_path;
   Lardon3DProjectDb *project_db;
   Lardon3DResourceGovernor *resource_governor;
+  Lardon3DOrbVulkanBackend *orb_vulkan_backend;
 } Lardon3DTaskReconstructionContext;
 
 typedef enum {

meson.build

@@ -37,6 +37,39 @@ opencv_benchmark = dependency(
     modules: ['opencv_core', 'opencv_imgcodecs', 'opencv_features2d', 'opencv_imgproc'],
 )
 
+vulkan_orb = get_option('vulkan_orb')
+vulkan = dependency('vulkan', required: vulkan_orb)
+glslc = find_program('glslc', required: vulkan_orb)
+python = find_program('python3', required: true)
+matcher_vulkan_enabled = vulkan_orb.allowed() and vulkan.found() and glslc.found()
+matcher_vulkan_configuration = configuration_data()
+matcher_vulkan_configuration.set10('LARDON3D_HAVE_VULKAN', matcher_vulkan_enabled)
+matcher_vulkan_config = configure_file(
+    output: 'matcher_vulkan_config.h',
+    configuration: matcher_vulkan_configuration,
+)
+matcher_backend_sources = [
+    'src/orb_vulkan_backend.cpp',
+    matcher_vulkan_config,
+]
+matcher_backend_dependencies = []
+if matcher_vulkan_enabled
+    orb_top2_spv = custom_target(
+        'orb-top2-spv',
+        input: 'shaders/orb_top2.comp',
+        output: 'orb_top2.spv',
+        command: [glslc, '-fshader-stage=compute', '@INPUT@', '-o', '@OUTPUT@'],
+    )
+    orb_top2_header = custom_target(
+        'orb-top2-header',
+        input: orb_top2_spv,
+        output: 'orb_top2_spv.h',
+        command: [python, files('tools/embed_spirv.py'), '@INPUT@', '@OUTPUT@'],
+    )
+    matcher_backend_sources += [orb_top2_header]
+    matcher_backend_dependencies += [vulkan]
+endif
+
 opencv_test_environment = environment()
 if get_option('b_sanitize').contains('thread')
     opencv_test_environment.set(
@@ -65,6 +98,7 @@ executable(
         'src/visual_index_task.c',
         'src/candidate_pair_gen.c',
         'src/candidate_pair_task.c',
+        'src/matcher_task.c',
         'src/image_view.c',
         'src/project.c',
         'src/project_db.c',
@@ -78,9 +112,10 @@ executable(
         'src/hardware_profile.c',
         'src/match_file.c',
         'src/matcher.cpp',
-    ],
+    ] + matcher_backend_sources,
     include_directories: include_directories('include'),
-    dependencies: [ncursesw, threads, sqlite3, openssl, opencv],
+    dependencies: [ncursesw, threads, sqlite3, openssl, opencv]
+        + matcher_backend_dependencies,
 )
 
 import_test = executable(
@@ -114,6 +149,7 @@ import_task_test = executable(
         'src/task_checkpoint.c',
         'src/task_kind_registry.c',
         'src/task_kinds.c',
+        'src/matcher_task.c', 'src/matcher.cpp', 'src/match_file.c',
         'src/feature_task.c', 'src/sift_task.c', 'src/precision_features.c',
         'src/feature_store.c',
         'src/visual_index.c',
@@ -126,14 +162,14 @@ import_task_test = executable(
         'src/image_catalog.c',
         'src/image_catalog_persistent.c',
         'src/image_view.c',
-    ],
+    ] + matcher_backend_sources,
     c_args: [
         '-DLARDON3D_IMPORT_TASK_TESTING',
         '-DLARDON3D_PROJECT_DB_TESTING',
         '-DLARDON3D_CHECKPOINT_TESTING',
     ],
     include_directories: include_directories('include'),
-    dependencies: [threads, sqlite3, openssl, opencv],
+    dependencies: [threads, sqlite3, openssl, opencv] + matcher_backend_dependencies,
 )
 
 test('import-task', import_task_test, timeout: 30)
@@ -252,7 +288,8 @@ candidate_pair_task_test = executable(
         'tests/test_candidate_pair_task.c', 'src/app_state.c',
         'src/project.c', 'src/project_db.c', 'src/task.c',
         'src/task_checkpoint.c', 'src/task_kind_registry.c',
-        'src/task_kinds.c', 'src/task_queue.c', 'src/import.c',
+        'src/task_kinds.c', 'src/matcher_task.c', 'src/matcher.cpp',
+        'src/match_file.c', 'src/task_queue.c', 'src/import.c',
         'src/import_task.c', 'src/image_catalog.c',
         'src/image_catalog_persistent.c', 'src/image_view.c',
         'src/feature_task.c', 'src/sift_task.c',
@@ -261,7 +298,7 @@ candidate_pair_task_test = executable(
         'src/candidate_pair_gen.c', 'src/candidate_pair_task.c',
         'src/feature_extractor_opencv.cpp',
         'src/resource_governor.c', 'src/resource_snapshot.c',
-    ],
+    ] + matcher_backend_sources,
     c_args: [
         '-DLARDON3D_PROJECT_DB_TESTING',
         '-DLARDON3D_CANDIDATE_PAIR_TASK_TESTING',
@@ -269,17 +306,43 @@ candidate_pair_task_test = executable(
         '-DLARDON3D_FEATURE_TASK_TESTING',
     ],
     include_directories: include_directories('include'),
-    dependencies: [threads, sqlite3, openssl, opencv],
+    dependencies: [threads, sqlite3, openssl, opencv] + matcher_backend_dependencies,
 )
 test('candidate-pair-task', candidate_pair_task_test, timeout: 60,
      env: opencv_test_environment)
 
+matcher_task_test = executable(
+    'test-matcher-task',
+    sources: [
+        'tests/test_matcher_task.c', 'src/app_state.c', 'src/project.c',
+        'src/project_db.c', 'src/task.c', 'src/task_checkpoint.c',
+        'src/task_kind_registry.c', 'src/task_kinds.c', 'src/task_queue.c',
+        'src/matcher_task.c', 'src/matcher.cpp', 'src/match_file.c',
+        'src/feature_store.c', 'src/resource_governor.c',
+        'src/resource_snapshot.c', 'src/hardware_profile.c',
+        'src/import.c', 'src/import_task.c', 'src/image_catalog.c',
+        'src/image_catalog_persistent.c', 'src/image_view.c',
+        'src/feature_task.c', 'src/sift_task.c', 'src/precision_features.c',
+        'src/feature_extractor_opencv.cpp', 'src/visual_index.c',
+        'src/visual_index_task.c', 'src/candidate_pair_gen.c',
+        'src/candidate_pair_task.c',
+    ] + matcher_backend_sources,
+    c_args: [
+        '-DLARDON3D_PROJECT_DB_TESTING',
+        '-DLARDON3D_MATCHER_TASK_TESTING',
+    ],
+    include_directories: include_directories('include'),
+    dependencies: [threads, sqlite3, openssl, opencv] + matcher_backend_dependencies,
+)
+test('matcher-task', matcher_task_test, timeout: 60, env: opencv_test_environment)
+
 feature_task_test = executable(
     'test-feature-task',
     sources: [
         'tests/test_feature_task.c', 'src/app_state.c', 'src/project.c',
         'src/project_db.c', 'src/task.c', 'src/task_checkpoint.c',
         'src/task_kind_registry.c', 'src/task_kinds.c', 'src/task_queue.c',
+        'src/matcher_task.c', 'src/matcher.cpp', 'src/match_file.c',
         'src/import.c', 'src/import_task.c', 'src/image_catalog.c',
         'src/image_catalog_persistent.c', 'src/image_view.c',
         'src/feature_task.c', 'src/sift_task.c', 'src/precision_features.c',
@@ -288,13 +351,13 @@ feature_task_test = executable(
         'src/candidate_pair_gen.c', 'src/candidate_pair_task.c',
         'src/feature_extractor_opencv.cpp', 'src/resource_governor.c',
         'src/resource_snapshot.c',
-    ],
+    ] + matcher_backend_sources,
     c_args: [
         '-DLARDON3D_PROJECT_DB_TESTING', '-DLARDON3D_FEATURE_TASK_TESTING',
         '-DLARDON3D_VISUAL_INDEX_TASK_TESTING',
     ],
     include_directories: include_directories('include'),
-    dependencies: [threads, sqlite3, openssl, opencv],
+    dependencies: [threads, sqlite3, openssl, opencv] + matcher_backend_dependencies,
 )
 test('feature-task', feature_task_test, timeout: 60, env: opencv_test_environment)
 
@@ -312,7 +375,8 @@ precision_consolidation_test = executable(
         'tests/test_precision_consolidation.c', 'src/app_state.c',
         'src/project.c', 'src/project_db.c', 'src/task.c',
         'src/task_checkpoint.c', 'src/task_kind_registry.c',
-        'src/task_kinds.c', 'src/task_queue.c', 'src/import.c',
+        'src/task_kinds.c', 'src/matcher_task.c', 'src/matcher.cpp',
+        'src/match_file.c', 'src/task_queue.c', 'src/import.c',
         'src/import_task.c', 'src/image_catalog.c',
         'src/image_catalog_persistent.c', 'src/image_view.c',
         'src/feature_task.c', 'src/sift_task.c',
@@ -321,14 +385,14 @@ precision_consolidation_test = executable(
         'src/candidate_pair_gen.c', 'src/candidate_pair_task.c',
         'src/feature_extractor_opencv.cpp',
         'src/resource_governor.c', 'src/resource_snapshot.c',
-    ],
+    ] + matcher_backend_sources,
     c_args: [
         '-DLARDON3D_PROJECT_DB_TESTING',
         '-DLARDON3D_FEATURE_TASK_TESTING',
         '-DLARDON3D_VISUAL_INDEX_TASK_TESTING',
     ],
     include_directories: include_directories('include'),
-    dependencies: [threads, sqlite3, openssl, opencv],
+    dependencies: [threads, sqlite3, openssl, opencv] + matcher_backend_dependencies,
 )
 test('precision-consolidation', precision_consolidation_test,
      timeout: 120, env: opencv_test_environment)
@@ -419,6 +483,7 @@ project_test = executable(
         'src/task_checkpoint.c',
         'src/task_kind_registry.c',
         'src/task_kinds.c',
+        'src/matcher_task.c', 'src/matcher.cpp', 'src/match_file.c',
         'src/feature_task.c', 'src/sift_task.c', 'src/precision_features.c',
         'src/feature_store.c',
         'src/visual_index.c',
@@ -433,13 +498,13 @@ project_test = executable(
         'src/image_view.c',
         'src/resource_governor.c',
         'src/resource_snapshot.c',
-    ],
+    ] + matcher_backend_sources,
     c_args: [
         '-DLARDON3D_PROJECT_DB_TESTING',
         '-DLARDON3D_CHECKPOINT_TESTING',
     ],
     include_directories: include_directories('include'),
-    dependencies: [threads, sqlite3, openssl, opencv],
+    dependencies: [threads, sqlite3, openssl, opencv] + matcher_backend_dependencies,
 )
 
 test('project', project_test, timeout: 30)
@@ -551,14 +616,38 @@ matcher_test = executable(
         'src/resource_snapshot.c',
         'src/image_catalog_persistent.c',
         'src/app_state.c',
-    ],
+    ] + matcher_backend_sources,
     c_args: ['-DLARDON3D_PROJECT_DB_TESTING'],
     include_directories: include_directories('include'),
-    dependencies: [threads, sqlite3, openssl, opencv],
+    dependencies: [threads, sqlite3, openssl, opencv] + matcher_backend_dependencies,
 )
 
 test('matcher', matcher_test, timeout: 60)
 
+if matcher_vulkan_enabled
+    orb_vulkan_backend_test = executable(
+        'test-orb-vulkan-backend',
+        sources: [
+            'tests/test_orb_vulkan_backend.cpp',
+        ] + matcher_backend_sources,
+        cpp_args: ['-DLARDON3D_ORB_VULKAN_TESTING'],
+        include_directories: include_directories('include'),
+        dependencies: [threads, opencv] + matcher_backend_dependencies,
+    )
+
+    test('orb-vulkan-backend', orb_vulkan_backend_test, timeout: 120,
+         env: opencv_test_environment)
+
+    executable(
+        'benchmark-orb-vulkan',
+        sources: [
+            'tests/benchmark_orb_vulkan.cpp',
+        ] + matcher_backend_sources,
+        include_directories: include_directories('include'),
+        dependencies: [opencv] + matcher_backend_dependencies,
+    )
+endif
+
 matcher_e2e_test = executable(
     'test-matcher-e2e',
     sources: [
@@ -573,9 +662,10 @@ matcher_e2e_test = executable(
         'src/resource_snapshot.c',
         'src/image_catalog_persistent.c',
         'src/app_state.c',
-    ],
+    ] + matcher_backend_sources,
+    c_args: matcher_vulkan_enabled ? ['-DLARDON3D_MATCHER_E2E_VULKAN'] : [],
     include_directories: include_directories('include'),
-    dependencies: [threads, sqlite3, openssl, opencv],
+    dependencies: [threads, sqlite3, openssl, opencv] + matcher_backend_dependencies,
 )
 
 test('matcher-e2e', matcher_e2e_test, timeout: 60)
@@ -594,7 +684,7 @@ executable(
         'src/resource_snapshot.c',
         'src/image_catalog_persistent.c',
         'src/app_state.c',
-    ],
+    ] + matcher_backend_sources,
     include_directories: include_directories('include'),
-    dependencies: [threads, sqlite3, openssl, opencv],
+    dependencies: [threads, sqlite3, openssl, opencv] + matcher_backend_dependencies,
 )

meson_options.txt

@@ -0,0 +1,6 @@
+option(
+    'vulkan_orb',
+    type: 'feature',
+    value: 'auto',
+    description: 'Build the optional Vulkan ORB matcher backend',
+)

shaders/orb_top2.comp

@@ -0,0 +1,63 @@
+#version 450
+
+layout(local_size_x_id = 0) in;
+
+layout(set = 0, binding = 0, std430) readonly buffer DescriptorsA {
+  uint descriptors_a[];
+};
+
+layout(set = 0, binding = 1, std430) readonly buffer DescriptorsB {
+  uint descriptors_b[];
+};
+
+layout(set = 0, binding = 2, std430) writeonly buffer Top2Output {
+  uvec4 top2[];
+};
+
+layout(push_constant) uniform Counts {
+  uint count_a;
+  uint count_b;
+} counts;
+
+void main() {
+  uint query_index = gl_GlobalInvocationID.x;
+  if (query_index >= counts.count_a) {
+    return;
+  }
+
+  uint best_index = 0xffffffffu;
+  uint best_distance = 0xffffffffu;
+  uint second_index = 0xffffffffu;
+  uint second_distance = 0xffffffffu;
+
+  for (uint train_index = 0; train_index < counts.count_b; ++train_index) {
+    uint distance = 0;
+    for (uint word = 0; word < 8; ++word) {
+      uint a = descriptors_a[query_index * 8 + word];
+      uint b = descriptors_b[train_index * 8 + word];
+      distance += bitCount(a ^ b);
+    }
+
+    // train_index increases monotonically, so keeping the first equal distance
+    // implements the required smallest-index tie-break without another comparison.
+    if (distance < best_distance) {
+      second_index = best_index;
+      second_distance = best_distance;
+      best_index = train_index;
+      best_distance = distance;
+      continue;
+    }
+
+    if (distance < second_distance) {
+      second_index = train_index;
+      second_distance = distance;
+    }
+  }
+
+  top2[query_index] = uvec4(
+      best_index,
+      best_distance,
+      second_index,
+      second_distance
+  );
+}

src/app.c

@@ -3,8 +3,10 @@
 
 #include <lardon3d/app.h>
 #include <lardon3d/app_state.h>
+#include <lardon3d/feature_extractor.h>
 #include <lardon3d/image_catalog.h>
 #include <lardon3d/image_view.h>
+#include <lardon3d/orb_vulkan_backend.h>
 #include <lardon3d/project.h>
 #include <lardon3d/resource_governor.h>
 #include <lardon3d/task_queue.h>
@@ -29,6 +31,10 @@ lardon3d_app_run(void)
         || !lardon3d_resource_policy_default(
             &state.hardware_profile,
             &resource_policy
+        )
+        || !lardon3d_feature_opencv_configure_threads(
+            state.hardware_profile.logical_cpu_count
+                - resource_policy.system_cpu_reserve
         )) {
         return EXIT_FAILURE;
     }
@@ -39,14 +45,21 @@ lardon3d_app_run(void)
     if (!state.resource_governor) {
         return EXIT_FAILURE;
     }
+    state.orb_vulkan_backend = lardon3d_orb_vulkan_backend_create();
+    if (!state.orb_vulkan_backend) {
+        lardon3d_resource_governor_destroy(state.resource_governor);
+        return EXIT_FAILURE;
+    }
     state.task_queue = lardon3d_task_queue_create(state.resource_governor, 64);
     if (!state.task_queue) {
+        lardon3d_orb_vulkan_backend_destroy(state.orb_vulkan_backend);
         lardon3d_resource_governor_destroy(state.resource_governor);
         return EXIT_FAILURE;
     }
 
     if (!lardon3d_tui_init()) {
         lardon3d_task_queue_destroy(state.task_queue);
+        lardon3d_orb_vulkan_backend_destroy(state.orb_vulkan_backend);
         lardon3d_resource_governor_destroy(state.resource_governor);
         return EXIT_FAILURE;
     }
@@ -58,6 +71,7 @@ lardon3d_app_run(void)
     if (state.project_loaded) {
         lardon3d_project_close(&state);
     }
+    lardon3d_orb_vulkan_backend_destroy(state.orb_vulkan_backend);
     lardon3d_resource_governor_destroy(state.resource_governor);
 
     return success ? EXIT_SUCCESS : EXIT_FAILURE;

src/candidate_pair_task.c

@@ -268,8 +268,11 @@ Lardon3DTask *lardon3d_project_create_candidate_pair_generate_task(
         .exclude_same_asset = query_options->exclude_same_asset,
     };
     Lardon3DTaskReconstructionContext runtime = {
-        state->project_path, state->project_db,
+        .project_path = state->project_path,
-        state->resource_governor};
+        .project_db = state->project_db,
+        .resource_governor = state->resource_governor,
+        .orb_vulkan_backend = state->orb_vulkan_backend,
+    };
     Lardon3DCandidatePairTaskContext *context =
         make_context(&runtime, &parameters);
     if (!context) return NULL;

src/feature_extractor_opencv.cpp

@@ -16,6 +16,19 @@ extern "C" {
 #include <lardon3d/feature_extractor.h>
 }
 
+extern "C" bool lardon3d_feature_opencv_configure_threads(unsigned int threads) {
+  if (threads == 0 || threads > static_cast<unsigned int>(std::numeric_limits<int>::max())) {
+    return false;
+  }
+  cv::setNumThreads(static_cast<int>(threads));
+  return cv::getNumThreads() == static_cast<int>(threads);
+}
+
+extern "C" unsigned int lardon3d_feature_opencv_thread_count(void) {
+  int threads = cv::getNumThreads();
+  return threads > 0 ? static_cast<unsigned int>(threads) : 1U;
+}
+
 extern "C" bool
 lardon3d_feature_extractor_parameters_valid(const Lardon3DFeatureExtractorParameters *parameters) {
   return parameters && parameters->max_features > 0 &&

src/feature_task.c

@@ -279,8 +279,12 @@ lardon3d_project_create_feature_extract_task(Lardon3DAppState *state, uint64_t i
   snprintf(durable.extractor_kind, sizeof(durable.extractor_kind), "%s",
            LARDON3D_FEATURE_EXTRACTOR_KIND);
   lardon3d_feature_extractor_parameter_fingerprint(parameters, durable.parameter_fingerprint);
-  Lardon3DTaskReconstructionContext runtime = {state->project_path, state->project_db,
+  Lardon3DTaskReconstructionContext runtime = {
-                                               state->resource_governor};
+      .project_path = state->project_path,
+      .project_db = state->project_db,
+      .resource_governor = state->resource_governor,
+      .orb_vulkan_backend = state->orb_vulkan_backend,
+  };
   Lardon3DFeatureTaskContext *context = make_context(&runtime, &durable);
   if (!context) {
     return NULL;
@@ -289,7 +293,8 @@ lardon3d_project_create_feature_extract_task(Lardon3DAppState *state, uint64_t i
                                        .memory_bytes_per_item = 512ULL * 1024 * 1024,
                                        .minimum_batch_size = 1,
                                        .maximum_batch_size = 1,
-                                       .desired_cpu_threads = 1,
+                                       .desired_cpu_threads =
+                                           lardon3d_feature_opencv_thread_count(),
                                        .desired_io_slots = 1,
                                        .task_class = LARDON3D_RESOURCE_TASK_CPU};
   Lardon3DTask *task = lardon3d_task_create_typed("Extraction de features", &estimate,

src/matcher.cpp

@@ -210,6 +210,26 @@ static bool accept_knn_match(const std::vector<cv::DMatch> &knn, float threshold
     return true;
 }
 
+static bool accept_orb_top2(const Lardon3DOrbTop2 &top2, uint32_t query_index,
+                            float threshold, MatchEntry *entry) {
+    if (!entry || top2.neighbor_count == 0 || top2.neighbor_count > 2 ||
+        !std::isfinite(threshold) || threshold <= 0.0F || threshold >= 1.0F) {
+        return false;
+    }
+    if (top2.neighbor_count == 2 &&
+        (top2.second_distance == 0 ||
+         !(static_cast<float>(top2.best_distance) <
+           threshold * static_cast<float>(top2.second_distance)))) {
+        return false;
+    }
+    *entry = {
+        static_cast<int>(query_index),
+        static_cast<int>(top2.best_index),
+        static_cast<float>(top2.best_distance),
+    };
+    return true;
+}
+
 static bool operator<(const MatchEntry &a, const MatchEntry &b) {
     if (a.query_idx != b.query_idx) {
         return a.query_idx < b.query_idx;
@@ -272,6 +292,7 @@ static Lardon3DMatcherResult matcher_run_impl(
     const Lardon3DProjectDbFeatureSet *feature_set_b,
     const Lardon3DMatcherParams *params,
     const char *match_file_path,
+    Lardon3DOrbVulkanBackend *backend,
     Lardon3DMatcherStats *stats) {
     if (stats) {
         memset(stats, 0, sizeof(*stats));
@@ -347,6 +368,34 @@ static Lardon3DMatcherResult matcher_run_impl(
             result = LARDON3D_MATCHER_IO_ERROR;
         } else {
             stats->descriptor_read_ns = elapsed_ns(phase_start);
+            bool completed = false;
+            if (backend && lardon3d_orb_vulkan_should_use(
+                               feature_set_a->feature_count,
+                               feature_set_b->feature_count)) {
+                std::vector<Lardon3DOrbTop2> top2(feature_set_a->feature_count);
+                phase_start = std::chrono::steady_clock::now();
+                Lardon3DOrbVulkanResult vulkan_result = lardon3d_orb_vulkan_top2(
+                    backend, desc_a.data(), feature_set_a->feature_count,
+                    desc_b.data(), feature_set_b->feature_count, top2.data(), top2.size());
+                stats->knn_ns = elapsed_ns(phase_start);
+                if (vulkan_result == LARDON3D_ORB_VULKAN_OK) {
+                    stats->used_vulkan = true;
+                    stats->knn_query_count = feature_set_a->feature_count;
+                    phase_start = std::chrono::steady_clock::now();
+                    for (uint32_t index = 0; index < feature_set_a->feature_count; ++index) {
+                        MatchEntry entry;
+                        if (accept_orb_top2(top2[index], index, params->ratio_threshold,
+                                            &entry)) {
+                            filtered_matches.push_back(entry);
+                        }
+                    }
+                    stats->filter_ns = elapsed_ns(phase_start);
+                    completed = true;
+                } else {
+                    stats->vulkan_fallback = true;
+                }
+            }
+            if (!completed) {
                 cv::Mat mat_a((int)feature_set_a->feature_count, 32, CV_8UC1, desc_a.data());
                 cv::Mat mat_b((int)feature_set_b->feature_count, 32, CV_8UC1, desc_b.data());
                 cv::BFMatcher matcher(cv::NORM_HAMMING, false);
@@ -356,14 +405,15 @@ static Lardon3DMatcherResult matcher_run_impl(
                 stats->knn_ns = elapsed_ns(phase_start);
                 stats->knn_query_count = (uint32_t)matches.size();
                 phase_start = std::chrono::steady_clock::now();
-            for (size_t i = 0; i < matches.size(); ++i) {
+                for (const auto &knn : matches) {
-                const auto &knn = matches[i];
                     MatchEntry entry;
-                if (accept_knn_match(knn, params->ratio_threshold, &entry))
+                    if (accept_knn_match(knn, params->ratio_threshold, &entry)) {
                         filtered_matches.push_back(entry);
                     }
+                }
                 stats->filter_ns = elapsed_ns(phase_start);
             }
+        }
     } else {
         std::vector<float> desc_a, desc_b;
         phase_start = std::chrono::steady_clock::now();
@@ -454,9 +504,21 @@ extern "C" Lardon3DMatcherResult lardon3d_matcher_run(
     const Lardon3DMatcherParams *params,
     const char *match_file_path,
     Lardon3DMatcherStats *stats) {
+    return lardon3d_matcher_run_with_backend(
+        project_path, feature_set_a, feature_set_b, params, match_file_path, nullptr, stats);
+}
+
+extern "C" Lardon3DMatcherResult lardon3d_matcher_run_with_backend(
+    const char *project_path,
+    const Lardon3DProjectDbFeatureSet *feature_set_a,
+    const Lardon3DProjectDbFeatureSet *feature_set_b,
+    const Lardon3DMatcherParams *params,
+    const char *match_file_path,
+    Lardon3DOrbVulkanBackend *backend,
+    Lardon3DMatcherStats *stats) {
     try {
         return matcher_run_impl(project_path, feature_set_a, feature_set_b, params,
-                                match_file_path, stats);
+                                match_file_path, backend, stats);
     } catch (const cv::Exception &) {
         return LARDON3D_MATCHER_FAILED;
     } catch (const std::bad_alloc &) {
@@ -473,6 +535,21 @@ extern "C" Lardon3DMatcherResult lardon3d_matcher_match_and_publish_profiled(
     const Lardon3DMatcherParams *params,
     Lardon3DProjectDbMatchResult *result,
     Lardon3DMatcherStats *profile) {
+    return lardon3d_matcher_match_and_publish_with_backend(
+        project_path, database, pair, feature_set_a, feature_set_b, params, nullptr,
+        result, profile);
+}
+
+extern "C" Lardon3DMatcherResult lardon3d_matcher_match_and_publish_with_backend(
+    const char *project_path,
+    Lardon3DProjectDb *database,
+    const Lardon3DProjectDbCandidatePair *pair,
+    const Lardon3DProjectDbFeatureSet *feature_set_a,
+    const Lardon3DProjectDbFeatureSet *feature_set_b,
+    const Lardon3DMatcherParams *params,
+    Lardon3DOrbVulkanBackend *backend,
+    Lardon3DProjectDbMatchResult *result,
+    Lardon3DMatcherStats *profile) {
     auto total_start = std::chrono::steady_clock::now();
     if (profile) memset(profile, 0, sizeof(*profile));
     if (result) {
@@ -548,8 +625,8 @@ extern "C" Lardon3DMatcherResult lardon3d_matcher_match_and_publish_profiled(
     }
 
     Lardon3DMatcherStats stats;
-    Lardon3DMatcherResult run_result = lardon3d_matcher_run(
+    Lardon3DMatcherResult run_result = lardon3d_matcher_run_with_backend(
-        project_path, feature_set_a, feature_set_b, params, tmp_path, &stats);
+        project_path, feature_set_a, feature_set_b, params, tmp_path, backend, &stats);
     if (run_result != LARDON3D_MATCHER_OK) {
         unlink(tmp_path);
         return run_result;
@@ -573,9 +650,9 @@ extern "C" Lardon3DMatcherResult lardon3d_matcher_match_and_publish_profiled(
                   feature_set_b->feature_set_id, matcher_kind, LARDON3D_MATCHER_VERSION, fp,
                   LARDON3D_MATCH_RESULT_STATUS_NO_MATCH, 0, NULL, NULL, 0, now, result);
         if (db_result == LARDON3D_PROJECT_DB_CONSTRAINT) {
-            return lardon3d_matcher_match_and_publish_profiled(
+            return lardon3d_matcher_match_and_publish_with_backend(
-                project_path, database, pair, feature_set_a, feature_set_b, params, result,
+                project_path, database, pair, feature_set_a, feature_set_b, params, backend,
-                profile);
+                result, profile);
         }
         if (profile) {
             profile->database_ns = elapsed_ns(database_start);
@@ -668,8 +745,9 @@ extern "C" Lardon3DMatcherResult lardon3d_matcher_match_and_publish_profiled(
               LARDON3D_MATCH_RESULT_STATUS_MATCHED, stats.match_count, file_hash, relative,
               (uint64_t)file_size, now, result);
     if (db_result == LARDON3D_PROJECT_DB_CONSTRAINT) {
-        return lardon3d_matcher_match_and_publish_profiled(
+        return lardon3d_matcher_match_and_publish_with_backend(
-            project_path, database, pair, feature_set_a, feature_set_b, params, result, profile);
+            project_path, database, pair, feature_set_a, feature_set_b, params, backend,
+            result, profile);
     }
     if (db_result != LARDON3D_PROJECT_DB_OK) {
         return LARDON3D_MATCHER_FAILED;

src/matcher_task.c

@@ -0,0 +1,375 @@
+#include <math.h>
+#include <stdbool.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <time.h>
+
+#include <lardon3d/matcher_task.h>
+#include <lardon3d/project.h>
+#include <lardon3d/task_queue.h>
+
+enum {
+  MATCHER_TASK_PAGE_CAPACITY = LARDON3D_MATCHER_TASK_MAXIMUM_BATCH + 1,
+  MATCHER_TASK_MEMORY_BYTES = 10 * 1024 * 1024,
+  MATCHER_TASK_CPU_THREADS = 12,
+};
+
+typedef struct {
+  char project_path[PATH_MAX];
+  Lardon3DProjectDb *database;
+  Lardon3DResourceGovernor *governor;
+  Lardon3DOrbVulkanBackend *orb_vulkan_backend;
+  Lardon3DProjectDbMatcherTask parameters;
+} Lardon3DMatcherTaskContext;
+
+static void destroy_context(void *userdata) { free(userdata); }
+
+static void runtime_state(const Lardon3DMatcherTaskContext *context,
+                          Lardon3DAppState *state) {
+  lardon3d_app_state_init(state);
+  state->project_loaded = true;
+  state->project_db = context->database;
+  state->resource_governor = context->governor;
+  state->orb_vulkan_backend = context->orb_vulkan_backend;
+  (void)snprintf(state->project_path, sizeof(state->project_path), "%s",
+                 context->project_path);
+}
+
+static void finished_callback(const Lardon3DTask *task, void *userdata) {
+#ifdef LARDON3D_MATCHER_TASK_TESTING
+  const char *skip = getenv("LARDON3D_TEST_MATCHER_SKIP_FINISHED_CHECKPOINT");
+  if (skip && strcmp(skip, "1") == 0) {
+    return;
+  }
+#endif
+  Lardon3DMatcherTaskContext *context = userdata;
+  Lardon3DAppState state;
+  runtime_state(context, &state);
+  (void)lardon3d_project_checkpoint_matcher_task(&state, task,
+                                                 &context->parameters);
+}
+
+static uint64_t elapsed_ns(struct timespec begin, struct timespec end) {
+  uint64_t seconds =
+      end.tv_sec >= begin.tv_sec ? (uint64_t)(end.tv_sec - begin.tv_sec) : 0;
+  long nanoseconds = end.tv_nsec - begin.tv_nsec;
+  if (nanoseconds < 0 && seconds > 0) {
+    --seconds;
+    nanoseconds += 1000000000L;
+  }
+  if (seconds > UINT64_MAX / 1000000000ULL) {
+    return UINT64_MAX;
+  }
+  return seconds * 1000000000ULL + (uint64_t)nanoseconds;
+}
+
+static bool load_feature_sets(Lardon3DMatcherTaskContext *context,
+                              const Lardon3DProjectDbCandidatePair *pair,
+                              Lardon3DProjectDbFeatureSet *feature_set_a,
+                              Lardon3DProjectDbFeatureSet *feature_set_b) {
+  return lardon3d_project_db_find_feature_set(
+             context->database, pair->image_id_a,
+             context->parameters.feature_extractor_kind,
+             context->parameters.feature_extractor_version,
+             context->parameters.feature_parameter_fingerprint,
+             feature_set_a) == LARDON3D_PROJECT_DB_OK &&
+         lardon3d_project_db_find_feature_set(
+             context->database, pair->image_id_b,
+             context->parameters.feature_extractor_kind,
+             context->parameters.feature_extractor_version,
+             context->parameters.feature_parameter_fingerprint,
+             feature_set_b) == LARDON3D_PROJECT_DB_OK;
+}
+
+static bool process_pair(Lardon3DTask *task,
+                         Lardon3DMatcherTaskContext *context,
+                         const Lardon3DProjectDbCandidatePair *pair) {
+  if (!lardon3d_task_checkpoint(task)) {
+    return false;
+  }
+  Lardon3DProjectDbFeatureSet feature_set_a;
+  Lardon3DProjectDbFeatureSet feature_set_b;
+  if (!load_feature_sets(context, pair, &feature_set_a, &feature_set_b)) {
+    return lardon3d_task_fail(task, "Feature Sets du Matcher introuvables.");
+  }
+  Lardon3DMatcherParams matcher = {
+      .kind = (Lardon3DMatcherKind)context->parameters.matcher_kind,
+      .ratio_threshold = context->parameters.ratio_threshold,
+  };
+  Lardon3DProjectDbMatchResult result;
+  if (lardon3d_matcher_match_and_publish_with_backend(
+          context->project_path, context->database, pair, &feature_set_a,
+          &feature_set_b, &matcher, context->orb_vulkan_backend, &result,
+          NULL) != LARDON3D_MATCHER_OK) {
+    return lardon3d_task_fail(task,
+                              "Matching de la Candidate Pair impossible.");
+  }
+#ifdef LARDON3D_MATCHER_TASK_TESTING
+  const char *pause = getenv("LARDON3D_TEST_MATCHER_PAUSE_AFTER_PUBLICATION");
+  if (pause && strcmp(pause, "1") == 0) {
+    (void)lardon3d_task_pause(task);
+    return lardon3d_task_checkpoint(task);
+  }
+#endif
+  return true;
+}
+
+static bool checkpoint_batch(Lardon3DTask *task,
+                             Lardon3DMatcherTaskContext *context,
+                             unsigned int progress, uint64_t processed) {
+  char message[LARDON3D_TASK_MESSAGE_CAPACITY];
+  (void)snprintf(message, sizeof(message), "Candidate Pairs traitées:%lu",
+                 (unsigned long)processed);
+  if (!lardon3d_task_set_progress(task, progress, message)) {
+    return false;
+  }
+  Lardon3DAppState state;
+  runtime_state(context, &state);
+  return lardon3d_project_checkpoint_matcher_task(&state, task,
+                                                  &context->parameters) ==
+         LARDON3D_PROJECT_TASK_CHECKPOINT_OK;
+}
+
+static bool run(Lardon3DTask *task, void *userdata) {
+  Lardon3DMatcherTaskContext *context = userdata;
+  uint64_t total_processed = 0;
+
+  for (;;) {
+    if (!lardon3d_task_checkpoint(task)) {
+      return false;
+    }
+    Lardon3DTaskExecutionContract contract;
+    if (!lardon3d_task_execution_contract(task, &contract) ||
+        contract.batch_size < LARDON3D_MATCHER_TASK_MINIMUM_BATCH ||
+        contract.batch_size > LARDON3D_MATCHER_TASK_MAXIMUM_BATCH) {
+      return lardon3d_task_fail(task, "Contrat de lot Matcher invalide.");
+    }
+
+    Lardon3DProjectDbCandidatePair page[MATCHER_TASK_PAGE_CAPACITY];
+    size_t count = 0;
+    size_t page_capacity = contract.batch_size + 1;
+    if (lardon3d_project_db_list_candidate_pairs(
+            context->database, context->parameters.after_candidate_pair_id,
+            page, page_capacity, &count) != LARDON3D_PROJECT_DB_OK) {
+      return lardon3d_task_fail(task, "Pagination Candidate Pair impossible.");
+    }
+    if (count == 0) {
+      return lardon3d_task_set_progress(task, 100, "Matching terminé.");
+    }
+
+    size_t batch_count =
+        count < contract.batch_size ? count : contract.batch_size;
+    struct timespec begin;
+    struct timespec end;
+    (void)clock_gettime(CLOCK_MONOTONIC, &begin);
+    for (size_t index = 0; index < batch_count; ++index) {
+      if (!process_pair(task, context, &page[index])) {
+        return false;
+      }
+      context->parameters.after_candidate_pair_id =
+          page[index].candidate_pair_id;
+      ++total_processed;
+    }
+    (void)clock_gettime(CLOCK_MONOTONIC, &end);
+    (void)lardon3d_resource_governor_record_batch(
+        context->governor, LARDON3D_RESOURCE_TASK_CPU, batch_count,
+        elapsed_ns(begin, end), 0);
+
+    bool exhausted = count <= contract.batch_size;
+    unsigned int progress = exhausted ? 100U : 99U;
+    if (!checkpoint_batch(task, context, progress, total_processed)) {
+      return lardon3d_task_fail(task, "Checkpoint Matcher impossible.");
+    }
+    if (exhausted) {
+      return lardon3d_task_set_progress(task, 100, "Matching terminé.");
+    }
+
+#ifdef LARDON3D_MATCHER_TASK_TESTING
+    const char *pause_after_batch =
+        getenv("LARDON3D_TEST_MATCHER_PAUSE_AFTER_BATCH");
+    if (pause_after_batch && strcmp(pause_after_batch, "1") == 0) {
+      (void)lardon3d_task_pause(task);
+      if (!lardon3d_task_checkpoint(task)) {
+        return false;
+      }
+    }
+#endif
+
+    Lardon3DResourceReservation *reservation = NULL;
+    if (!lardon3d_task_sequence_break(task, context->governor, &reservation,
+                                      &contract)) {
+      return false;
+    }
+  }
+}
+
+static bool
+valid_configuration(const Lardon3DMatcherTaskConfiguration *configuration) {
+  if (!configuration ||
+      !lardon3d_task_kind_is_valid(configuration->feature_extractor_kind) ||
+      configuration->feature_extractor_version == 0) {
+    return false;
+  }
+  bool kind_matches =
+      (configuration->matcher.kind == LARDON3D_MATCHER_ORB_BF &&
+       strcmp(configuration->feature_extractor_kind, "orb") == 0) ||
+      (configuration->matcher.kind == LARDON3D_MATCHER_SIFT_BF &&
+       strcmp(configuration->feature_extractor_kind, "sift") == 0) ||
+      (configuration->matcher.kind == LARDON3D_MATCHER_ROOTSIFT_BF &&
+       strcmp(configuration->feature_extractor_kind, "rootsift") == 0);
+  return kind_matches && isfinite(configuration->matcher.ratio_threshold) &&
+         configuration->matcher.ratio_threshold > 0.0F &&
+         configuration->matcher.ratio_threshold < 1.0F;
+}
+
+static Lardon3DMatcherTaskContext *
+make_context(const Lardon3DTaskReconstructionContext *runtime,
+             const Lardon3DProjectDbMatcherTask *parameters) {
+  if (!runtime || !runtime->project_path || !runtime->project_db ||
+      !runtime->resource_governor || !parameters) {
+    return NULL;
+  }
+  Lardon3DMatcherTaskContext *context = calloc(1, sizeof(*context));
+  if (!context) {
+    return NULL;
+  }
+  int written = snprintf(context->project_path, sizeof(context->project_path),
+                         "%s", runtime->project_path);
+  if (written <= 0 || (size_t)written >= sizeof(context->project_path)) {
+    free(context);
+    return NULL;
+  }
+  context->database = runtime->project_db;
+  context->governor = runtime->resource_governor;
+  context->orb_vulkan_backend = runtime->orb_vulkan_backend;
+  context->parameters = *parameters;
+  return context;
+}
+
+bool lardon3d_matcher_task_reconstruct(
+    const Lardon3DTaskDurableSnapshot *snapshot, void *userdata,
+    Lardon3DTaskKindBinding *binding) {
+  Lardon3DTaskReconstructionContext *runtime = userdata;
+  if (!snapshot || !runtime || !binding) {
+    return false;
+  }
+  Lardon3DProjectDbMatcherTask parameters;
+  if (lardon3d_project_db_load_matcher_task(runtime->project_db, snapshot->id,
+                                            &parameters) !=
+      LARDON3D_PROJECT_DB_OK) {
+    return false;
+  }
+  Lardon3DMatcherTaskConfiguration configuration = {
+      .feature_extractor_version = parameters.feature_extractor_version,
+      .matcher =
+          {
+              .kind = (Lardon3DMatcherKind)parameters.matcher_kind,
+              .ratio_threshold = parameters.ratio_threshold,
+          },
+  };
+  (void)snprintf(configuration.feature_extractor_kind,
+                 sizeof(configuration.feature_extractor_kind), "%s",
+                 parameters.feature_extractor_kind);
+  if (!valid_configuration(&configuration)) {
+    return false;
+  }
+  Lardon3DMatcherTaskContext *context = make_context(runtime, &parameters);
+  if (!context) {
+    return false;
+  }
+  *binding = (Lardon3DTaskKindBinding){
+      .callback = run,
+      .userdata = context,
+      .userdata_destroy = destroy_context,
+      .finished_callback = finished_callback,
+      .finished_userdata = context,
+  };
+  return true;
+}
+
+Lardon3DTask *lardon3d_project_create_matcher_task(
+    Lardon3DAppState *state,
+    const Lardon3DMatcherTaskConfiguration *configuration, uint64_t *task_id) {
+  if (task_id) {
+    *task_id = 0;
+  }
+  if (!state || !state->project_loaded || !state->project_db ||
+      !state->resource_governor || !task_id ||
+      !valid_configuration(configuration)) {
+    return NULL;
+  }
+  uint64_t id = 0;
+  if (lardon3d_project_db_allocate_task_id(state->project_db, &id) !=
+      LARDON3D_PROJECT_DB_OK) {
+    return NULL;
+  }
+  Lardon3DProjectDbMatcherTask parameters = {
+      .task_id = id,
+      .matcher_kind = (int)configuration->matcher.kind,
+      .ratio_threshold = configuration->matcher.ratio_threshold,
+      .feature_extractor_version = configuration->feature_extractor_version,
+  };
+  (void)snprintf(parameters.feature_extractor_kind,
+                 sizeof(parameters.feature_extractor_kind), "%s",
+                 configuration->feature_extractor_kind);
+  memcpy(parameters.feature_parameter_fingerprint,
+         configuration->feature_parameter_fingerprint,
+         sizeof(parameters.feature_parameter_fingerprint));
+  Lardon3DTaskReconstructionContext runtime = {
+      .project_path = state->project_path,
+      .project_db = state->project_db,
+      .resource_governor = state->resource_governor,
+      .orb_vulkan_backend = state->orb_vulkan_backend,
+  };
+  Lardon3DMatcherTaskContext *context = make_context(&runtime, &parameters);
+  if (!context) {
+    return NULL;
+  }
+  bool may_use_vulkan = configuration->matcher.kind == LARDON3D_MATCHER_ORB_BF &&
+                        state->orb_vulkan_backend &&
+                        state->hardware_profile.gpu_available;
+  const Lardon3DResourceEstimate estimate = {
+      .memory_fixed_bytes = MATCHER_TASK_MEMORY_BYTES,
+      .gpu_memory_fixed_bytes = may_use_vulkan
+                                    ? LARDON3D_ORB_VULKAN_PERMANENT_BUFFER_BYTES
+                                    : 0,
+      .minimum_batch_size = LARDON3D_MATCHER_TASK_MINIMUM_BATCH,
+      .maximum_batch_size = LARDON3D_MATCHER_TASK_MAXIMUM_BATCH,
+      .desired_cpu_threads = MATCHER_TASK_CPU_THREADS,
+      .desired_gpu_slots = may_use_vulkan ? 1U : 0U,
+      .desired_io_slots = 1,
+      .task_class = LARDON3D_RESOURCE_TASK_CPU,
+  };
+  Lardon3DTask *task = lardon3d_task_create_typed(
+      "Matching Candidate Pairs", &estimate, LARDON3D_MATCHER_TASK_KIND,
+      LARDON3D_MATCHER_TASK_KIND_VERSION, run, context, destroy_context);
+  if (!task || !lardon3d_task_assign_id(task, id) ||
+      !lardon3d_task_set_finished_callback(task, finished_callback, context) ||
+      lardon3d_project_checkpoint_matcher_task(state, task, &parameters) !=
+          LARDON3D_PROJECT_TASK_CHECKPOINT_OK) {
+    lardon3d_task_destroy(task);
+    return NULL;
+  }
+  *task_id = id;
+  return task;
+}
+
+bool lardon3d_project_enqueue_matcher_task(
+    Lardon3DAppState *state,
+    const Lardon3DMatcherTaskConfiguration *configuration, uint64_t *task_id) {
+  if (!state || !state->task_queue) {
+    return false;
+  }
+  Lardon3DTask *task =
+      lardon3d_project_create_matcher_task(state, configuration, task_id);
+  if (!task) {
+    return false;
+  }
+  if (!lardon3d_task_queue_add(state->task_queue, task, NULL)) {
+    lardon3d_task_destroy(task);
+    return false;
+  }
+  return true;
+}

src/orb_vulkan_backend.cpp

@@ -0,0 +1,819 @@
+#include <lardon3d/orb_vulkan_backend.h>
+
+#include <chrono>
+#include <cstdio>
+#include <cstring>
+#include <mutex>
+#include <new>
+
+#include <lardon3d/feature_extractor.h>
+
+#include "matcher_vulkan_config.h"
+
+#if LARDON3D_HAVE_VULKAN
+
+#include <algorithm>
+#include <cstdlib>
+#include <vector>
+
+#include <vulkan/vulkan.h>
+
+#include "orb_top2_spv.h"
+
+namespace {
+
+constexpr VkDeviceSize kDescriptorBufferBytes =
+    static_cast<VkDeviceSize>(LARDON3D_FEATURE_MAX_FEATURES) * 32;
+constexpr VkDeviceSize kOutputBufferBytes =
+    static_cast<VkDeviceSize>(LARDON3D_FEATURE_MAX_FEATURES) * 4 * sizeof(uint32_t);
+static_assert(kDescriptorBufferBytes * 2 + kOutputBufferBytes ==
+              LARDON3D_ORB_VULKAN_PERMANENT_BUFFER_BYTES);
+constexpr uint64_t kDefaultVulkanWorkThreshold = 768ULL * 768ULL;
+constexpr uint32_t kDefaultWorkgroupSize = 32;
+
+enum class BackendState {
+  kUninitialized,
+  kAvailable,
+  kUnavailable,
+  kFailed,
+};
+
+struct Buffer {
+  VkBuffer buffer = VK_NULL_HANDLE;
+  VkDeviceMemory memory = VK_NULL_HANDLE;
+  void *mapping = nullptr;
+  VkDeviceSize size = 0;
+  bool coherent = false;
+};
+
+struct RawTop2 {
+  uint32_t best_index;
+  uint32_t best_distance;
+  uint32_t second_index;
+  uint32_t second_distance;
+};
+
+static uint64_t elapsed_ns(std::chrono::steady_clock::time_point start) {
+  auto elapsed = std::chrono::steady_clock::now() - start;
+  return static_cast<uint64_t>(
+      std::chrono::duration_cast<std::chrono::nanoseconds>(elapsed).count());
+}
+
+static bool validation_requested() {
+  const char *value = std::getenv("LARDON3D_VULKAN_VALIDATION");
+  return value && std::strcmp(value, "1") == 0;
+}
+
+static uint32_t configured_workgroup_size() {
+  const char *value = std::getenv("LARDON3D_VULKAN_WORKGROUP_SIZE");
+  if (!value || value[0] == '\0') {
+    return kDefaultWorkgroupSize;
+  }
+  char *end = nullptr;
+  unsigned long parsed = std::strtoul(value, &end, 10);
+  if (!end || end[0] != '\0' ||
+      (parsed != 32 && parsed != 64 && parsed != 128 && parsed != 256)) {
+    return kDefaultWorkgroupSize;
+  }
+  return static_cast<uint32_t>(parsed);
+}
+
+static bool has_validation_layer() {
+  uint32_t count = 0;
+  if (vkEnumerateInstanceLayerProperties(&count, nullptr) != VK_SUCCESS) {
+    return false;
+  }
+  std::vector<VkLayerProperties> layers(count);
+  if (count > 0 &&
+      vkEnumerateInstanceLayerProperties(&count, layers.data()) != VK_SUCCESS) {
+    return false;
+  }
+  return std::any_of(layers.begin(), layers.end(), [](const auto &layer) {
+    return std::strcmp(layer.layerName, "VK_LAYER_KHRONOS_validation") == 0;
+  });
+}
+
+}  // namespace
+
+struct Lardon3DOrbVulkanBackend {
+  std::mutex mutex;
+  BackendState state = BackendState::kUninitialized;
+  VkInstance instance = VK_NULL_HANDLE;
+  VkPhysicalDevice physical_device = VK_NULL_HANDLE;
+  VkDevice device = VK_NULL_HANDLE;
+  VkQueue queue = VK_NULL_HANDLE;
+  uint32_t queue_family = UINT32_MAX;
+  bool dedicated_compute_queue = false;
+  VkPhysicalDeviceProperties properties{};
+  VkPhysicalDeviceMemoryProperties memory_properties{};
+  VkCommandPool command_pool = VK_NULL_HANDLE;
+  VkCommandBuffer command_buffer = VK_NULL_HANDLE;
+  VkDescriptorSetLayout descriptor_set_layout = VK_NULL_HANDLE;
+  VkPipelineLayout pipeline_layout = VK_NULL_HANDLE;
+  VkPipeline pipeline = VK_NULL_HANDLE;
+  VkDescriptorPool descriptor_pool = VK_NULL_HANDLE;
+  VkDescriptorSet descriptor_set = VK_NULL_HANDLE;
+  VkQueryPool query_pool = VK_NULL_HANDLE;
+  bool timestamps_available = false;
+  Buffer descriptors_a;
+  Buffer descriptors_b;
+  Buffer output;
+  uint32_t workgroup_size = kDefaultWorkgroupSize;
+  uint64_t initialization_ns = 0;
+  uint64_t last_dispatch_ns = 0;
+  uint64_t last_gpu_ns = 0;
+};
+
+namespace {
+
+static void destroy_buffer(Lardon3DOrbVulkanBackend *backend, Buffer *buffer) {
+  if (!backend || !buffer || backend->device == VK_NULL_HANDLE) {
+    return;
+  }
+  if (buffer->mapping) {
+    vkUnmapMemory(backend->device, buffer->memory);
+  }
+  if (buffer->buffer != VK_NULL_HANDLE) {
+    vkDestroyBuffer(backend->device, buffer->buffer, nullptr);
+  }
+  if (buffer->memory != VK_NULL_HANDLE) {
+    vkFreeMemory(backend->device, buffer->memory, nullptr);
+  }
+  *buffer = Buffer{};
+}
+
+static void destroy_vulkan(Lardon3DOrbVulkanBackend *backend) {
+  if (!backend) {
+    return;
+  }
+  if (backend->device != VK_NULL_HANDLE) {
+    (void)vkDeviceWaitIdle(backend->device);
+  }
+  destroy_buffer(backend, &backend->descriptors_a);
+  destroy_buffer(backend, &backend->descriptors_b);
+  destroy_buffer(backend, &backend->output);
+  if (backend->query_pool != VK_NULL_HANDLE) {
+    vkDestroyQueryPool(backend->device, backend->query_pool, nullptr);
+  }
+  if (backend->pipeline != VK_NULL_HANDLE) {
+    vkDestroyPipeline(backend->device, backend->pipeline, nullptr);
+  }
+  if (backend->pipeline_layout != VK_NULL_HANDLE) {
+    vkDestroyPipelineLayout(backend->device, backend->pipeline_layout, nullptr);
+  }
+  if (backend->descriptor_pool != VK_NULL_HANDLE) {
+    vkDestroyDescriptorPool(backend->device, backend->descriptor_pool, nullptr);
+  }
+  if (backend->descriptor_set_layout != VK_NULL_HANDLE) {
+    vkDestroyDescriptorSetLayout(backend->device, backend->descriptor_set_layout, nullptr);
+  }
+  if (backend->command_pool != VK_NULL_HANDLE) {
+    vkDestroyCommandPool(backend->device, backend->command_pool, nullptr);
+  }
+  if (backend->device != VK_NULL_HANDLE) {
+    vkDestroyDevice(backend->device, nullptr);
+  }
+  if (backend->instance != VK_NULL_HANDLE) {
+    vkDestroyInstance(backend->instance, nullptr);
+  }
+  backend->instance = VK_NULL_HANDLE;
+  backend->physical_device = VK_NULL_HANDLE;
+  backend->device = VK_NULL_HANDLE;
+  backend->queue = VK_NULL_HANDLE;
+  backend->command_pool = VK_NULL_HANDLE;
+  backend->command_buffer = VK_NULL_HANDLE;
+  backend->descriptor_set_layout = VK_NULL_HANDLE;
+  backend->pipeline_layout = VK_NULL_HANDLE;
+  backend->pipeline = VK_NULL_HANDLE;
+  backend->descriptor_pool = VK_NULL_HANDLE;
+  backend->descriptor_set = VK_NULL_HANDLE;
+  backend->query_pool = VK_NULL_HANDLE;
+  backend->timestamps_available = false;
+}
+
+static bool create_instance(Lardon3DOrbVulkanBackend *backend) {
+  VkApplicationInfo application{};
+  application.sType = VK_STRUCTURE_TYPE_APPLICATION_INFO;
+  application.pApplicationName = "Lardon3D ORB Matcher";
+  application.applicationVersion = VK_MAKE_API_VERSION(0, 1, 0, 0);
+  application.pEngineName = "Lardon3D";
+  application.engineVersion = VK_MAKE_API_VERSION(0, 1, 0, 0);
+  application.apiVersion = VK_API_VERSION_1_1;
+
+  const char *validation_layer = "VK_LAYER_KHRONOS_validation";
+  bool enable_validation = validation_requested() && has_validation_layer();
+  VkInstanceCreateInfo create_info{};
+  create_info.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO;
+  create_info.pApplicationInfo = &application;
+  create_info.enabledLayerCount = enable_validation ? 1U : 0U;
+  create_info.ppEnabledLayerNames = enable_validation ? &validation_layer : nullptr;
+  return vkCreateInstance(&create_info, nullptr, &backend->instance) == VK_SUCCESS;
+}
+
+static bool find_compute_queue(VkPhysicalDevice device, uint32_t *family,
+                               bool *dedicated) {
+  uint32_t count = 0;
+  vkGetPhysicalDeviceQueueFamilyProperties(device, &count, nullptr);
+  if (count == 0) {
+    return false;
+  }
+  std::vector<VkQueueFamilyProperties> families(count);
+  vkGetPhysicalDeviceQueueFamilyProperties(device, &count, families.data());
+  uint32_t fallback = UINT32_MAX;
+  for (uint32_t index = 0; index < count; ++index) {
+    VkQueueFlags flags = families[index].queueFlags;
+    if (families[index].queueCount == 0 || (flags & VK_QUEUE_COMPUTE_BIT) == 0) {
+      continue;
+    }
+    if ((flags & VK_QUEUE_GRAPHICS_BIT) == 0) {
+      *family = index;
+      *dedicated = true;
+      return true;
+    }
+    if (fallback == UINT32_MAX) {
+      fallback = index;
+    }
+  }
+  if (fallback == UINT32_MAX) {
+    return false;
+  }
+  *family = fallback;
+  *dedicated = false;
+  return true;
+}
+
+static int device_score(VkPhysicalDevice device, uint32_t *family,
+                        bool *dedicated) {
+  if (!find_compute_queue(device, family, dedicated)) {
+    return -1;
+  }
+  VkPhysicalDeviceProperties properties;
+  vkGetPhysicalDeviceProperties(device, &properties);
+  if (properties.limits.maxComputeWorkGroupInvocations < 32 ||
+      properties.limits.maxComputeWorkGroupSize[0] < 32 ||
+      properties.limits.maxStorageBufferRange < kDescriptorBufferBytes) {
+    return -1;
+  }
+  int score = *dedicated ? 100 : 0;
+  if (properties.deviceType == VK_PHYSICAL_DEVICE_TYPE_INTEGRATED_GPU) {
+    score += 30;
+  } else if (properties.deviceType == VK_PHYSICAL_DEVICE_TYPE_DISCRETE_GPU) {
+    score += 20;
+  }
+  const char *requested = std::getenv("LARDON3D_VULKAN_DEVICE");
+  if (requested && requested[0] != '\0' &&
+      std::strstr(properties.deviceName, requested)) {
+    score += 1000;
+  }
+  return score;
+}
+
+static bool select_device(Lardon3DOrbVulkanBackend *backend) {
+  uint32_t count = 0;
+  if (vkEnumeratePhysicalDevices(backend->instance, &count, nullptr) != VK_SUCCESS ||
+      count == 0) {
+    return false;
+  }
+  std::vector<VkPhysicalDevice> devices(count);
+  if (vkEnumeratePhysicalDevices(backend->instance, &count, devices.data()) != VK_SUCCESS) {
+    return false;
+  }
+  int best_score = -1;
+  for (VkPhysicalDevice device : devices) {
+    uint32_t family = UINT32_MAX;
+    bool dedicated = false;
+    int score = device_score(device, &family, &dedicated);
+    if (score > best_score) {
+      best_score = score;
+      backend->physical_device = device;
+      backend->queue_family = family;
+      backend->dedicated_compute_queue = dedicated;
+    }
+  }
+  if (backend->physical_device == VK_NULL_HANDLE) {
+    return false;
+  }
+  vkGetPhysicalDeviceProperties(backend->physical_device, &backend->properties);
+  vkGetPhysicalDeviceMemoryProperties(backend->physical_device,
+                                      &backend->memory_properties);
+  backend->workgroup_size = configured_workgroup_size();
+  return backend->workgroup_size <=
+             backend->properties.limits.maxComputeWorkGroupInvocations &&
+         backend->workgroup_size <=
+             backend->properties.limits.maxComputeWorkGroupSize[0];
+}
+
+static bool create_device_and_commands(Lardon3DOrbVulkanBackend *backend) {
+  float priority = 0.5F;
+  VkDeviceQueueCreateInfo queue_info{};
+  queue_info.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO;
+  queue_info.queueFamilyIndex = backend->queue_family;
+  queue_info.queueCount = 1;
+  queue_info.pQueuePriorities = &priority;
+  VkDeviceCreateInfo device_info{};
+  device_info.sType = VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO;
+  device_info.queueCreateInfoCount = 1;
+  device_info.pQueueCreateInfos = &queue_info;
+  if (vkCreateDevice(backend->physical_device, &device_info, nullptr,
+                     &backend->device) != VK_SUCCESS) {
+    return false;
+  }
+  vkGetDeviceQueue(backend->device, backend->queue_family, 0, &backend->queue);
+
+  VkCommandPoolCreateInfo pool_info{};
+  pool_info.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO;
+  pool_info.flags = VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT;
+  pool_info.queueFamilyIndex = backend->queue_family;
+  if (vkCreateCommandPool(backend->device, &pool_info, nullptr,
+                          &backend->command_pool) != VK_SUCCESS) {
+    return false;
+  }
+  VkCommandBufferAllocateInfo command_info{};
+  command_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO;
+  command_info.commandPool = backend->command_pool;
+  command_info.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY;
+  command_info.commandBufferCount = 1;
+  if (vkAllocateCommandBuffers(backend->device, &command_info,
+                               &backend->command_buffer) != VK_SUCCESS) {
+    return false;
+  }
+  return true;
+}
+
+static bool create_pipeline(Lardon3DOrbVulkanBackend *backend) {
+  VkDescriptorSetLayoutBinding bindings[3]{};
+  for (uint32_t index = 0; index < 3; ++index) {
+    bindings[index].binding = index;
+    bindings[index].descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;
+    bindings[index].descriptorCount = 1;
+    bindings[index].stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;
+  }
+  VkDescriptorSetLayoutCreateInfo descriptor_info{};
+  descriptor_info.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO;
+  descriptor_info.bindingCount = 3;
+  descriptor_info.pBindings = bindings;
+  if (vkCreateDescriptorSetLayout(backend->device, &descriptor_info, nullptr,
+                                  &backend->descriptor_set_layout) != VK_SUCCESS) {
+    return false;
+  }
+
+  VkPushConstantRange push_range{};
+  push_range.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;
+  push_range.size = 2 * sizeof(uint32_t);
+  VkPipelineLayoutCreateInfo layout_info{};
+  layout_info.sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO;
+  layout_info.setLayoutCount = 1;
+  layout_info.pSetLayouts = &backend->descriptor_set_layout;
+  layout_info.pushConstantRangeCount = 1;
+  layout_info.pPushConstantRanges = &push_range;
+  if (vkCreatePipelineLayout(backend->device, &layout_info, nullptr,
+                             &backend->pipeline_layout) != VK_SUCCESS) {
+    return false;
+  }
+
+  VkShaderModuleCreateInfo shader_info{};
+  shader_info.sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO;
+  shader_info.codeSize = lardon3d_orb_top2_spv_size;
+  shader_info.pCode = lardon3d_orb_top2_spv;
+  VkShaderModule shader = VK_NULL_HANDLE;
+  if (vkCreateShaderModule(backend->device, &shader_info, nullptr, &shader) != VK_SUCCESS) {
+    return false;
+  }
+  VkSpecializationMapEntry workgroup_entry{0, 0, sizeof(uint32_t)};
+  VkSpecializationInfo specialization{};
+  specialization.mapEntryCount = 1;
+  specialization.pMapEntries = &workgroup_entry;
+  specialization.dataSize = sizeof(backend->workgroup_size);
+  specialization.pData = &backend->workgroup_size;
+  VkPipelineShaderStageCreateInfo stage{};
+  stage.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO;
+  stage.stage = VK_SHADER_STAGE_COMPUTE_BIT;
+  stage.module = shader;
+  stage.pName = "main";
+  stage.pSpecializationInfo = &specialization;
+  VkComputePipelineCreateInfo pipeline_info{};
+  pipeline_info.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO;
+  pipeline_info.stage = stage;
+  pipeline_info.layout = backend->pipeline_layout;
+  VkResult result = vkCreateComputePipelines(backend->device, VK_NULL_HANDLE, 1,
+                                             &pipeline_info, nullptr,
+                                             &backend->pipeline);
+  vkDestroyShaderModule(backend->device, shader, nullptr);
+  return result == VK_SUCCESS;
+}
+
+static bool select_memory_type(Lardon3DOrbVulkanBackend *backend,
+                               uint32_t memory_type_bits, uint32_t *type_index,
+                               bool *coherent) {
+  int best_score = -1;
+  for (uint32_t index = 0; index < backend->memory_properties.memoryTypeCount; ++index) {
+    if ((memory_type_bits & (1U << index)) == 0) {
+      continue;
+    }
+    VkMemoryPropertyFlags flags =
+        backend->memory_properties.memoryTypes[index].propertyFlags;
+    if ((flags & VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT) == 0) {
+      continue;
+    }
+    int score = 0;
+    if ((flags & VK_MEMORY_PROPERTY_HOST_COHERENT_BIT) != 0) {
+      score += 2;
+    }
+    if ((flags & VK_MEMORY_PROPERTY_HOST_CACHED_BIT) != 0) {
+      score += 4;
+    }
+    if (score > best_score) {
+      best_score = score;
+      *type_index = index;
+      *coherent = (flags & VK_MEMORY_PROPERTY_HOST_COHERENT_BIT) != 0;
+    }
+  }
+  return best_score >= 0;
+}
+
+static bool create_buffer(Lardon3DOrbVulkanBackend *backend, VkDeviceSize size,
+                          Buffer *buffer) {
+  buffer->size = size;
+  VkBufferCreateInfo buffer_info{};
+  buffer_info.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO;
+  buffer_info.size = size;
+  buffer_info.usage = VK_BUFFER_USAGE_STORAGE_BUFFER_BIT;
+  buffer_info.sharingMode = VK_SHARING_MODE_EXCLUSIVE;
+  if (vkCreateBuffer(backend->device, &buffer_info, nullptr, &buffer->buffer) !=
+      VK_SUCCESS) {
+    return false;
+  }
+  VkMemoryRequirements requirements;
+  vkGetBufferMemoryRequirements(backend->device, buffer->buffer, &requirements);
+  uint32_t type_index = 0;
+  if (!select_memory_type(backend, requirements.memoryTypeBits, &type_index,
+                          &buffer->coherent)) {
+    return false;
+  }
+  VkMemoryAllocateInfo allocate_info{};
+  allocate_info.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO;
+  allocate_info.allocationSize = requirements.size;
+  allocate_info.memoryTypeIndex = type_index;
+  if (vkAllocateMemory(backend->device, &allocate_info, nullptr, &buffer->memory) !=
+          VK_SUCCESS ||
+      vkBindBufferMemory(backend->device, buffer->buffer, buffer->memory, 0) !=
+          VK_SUCCESS ||
+      vkMapMemory(backend->device, buffer->memory, 0, size, 0,
+                  &buffer->mapping) != VK_SUCCESS) {
+    return false;
+  }
+  return true;
+}
+
+static bool create_buffers_and_descriptors(Lardon3DOrbVulkanBackend *backend) {
+  VkDescriptorPoolSize pool_size{};
+  pool_size.type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;
+  pool_size.descriptorCount = 3;
+  VkDescriptorPoolCreateInfo pool_info{};
+  pool_info.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_POOL_CREATE_INFO;
+  pool_info.maxSets = 1;
+  pool_info.poolSizeCount = 1;
+  pool_info.pPoolSizes = &pool_size;
+  if (vkCreateDescriptorPool(backend->device, &pool_info, nullptr,
+                             &backend->descriptor_pool) != VK_SUCCESS) {
+    return false;
+  }
+  VkDescriptorSetAllocateInfo set_info{};
+  set_info.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_ALLOCATE_INFO;
+  set_info.descriptorPool = backend->descriptor_pool;
+  set_info.descriptorSetCount = 1;
+  set_info.pSetLayouts = &backend->descriptor_set_layout;
+  if (vkAllocateDescriptorSets(backend->device, &set_info,
+                               &backend->descriptor_set) != VK_SUCCESS ||
+      !create_buffer(backend, kDescriptorBufferBytes, &backend->descriptors_a) ||
+      !create_buffer(backend, kDescriptorBufferBytes, &backend->descriptors_b) ||
+      !create_buffer(backend, kOutputBufferBytes, &backend->output)) {
+    return false;
+  }
+
+  VkDescriptorBufferInfo buffer_info[3] = {
+      {backend->descriptors_a.buffer, 0, backend->descriptors_a.size},
+      {backend->descriptors_b.buffer, 0, backend->descriptors_b.size},
+      {backend->output.buffer, 0, backend->output.size},
+  };
+  VkWriteDescriptorSet writes[3]{};
+  for (uint32_t index = 0; index < 3; ++index) {
+    writes[index].sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET;
+    writes[index].dstSet = backend->descriptor_set;
+    writes[index].dstBinding = index;
+    writes[index].descriptorCount = 1;
+    writes[index].descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;
+    writes[index].pBufferInfo = &buffer_info[index];
+  }
+  vkUpdateDescriptorSets(backend->device, 3, writes, 0, nullptr);
+
+  uint32_t family_count = 0;
+  vkGetPhysicalDeviceQueueFamilyProperties(backend->physical_device, &family_count, nullptr);
+  std::vector<VkQueueFamilyProperties> families(family_count);
+  vkGetPhysicalDeviceQueueFamilyProperties(backend->physical_device, &family_count,
+                                           families.data());
+  backend->timestamps_available =
+      backend->queue_family < family_count &&
+      families[backend->queue_family].timestampValidBits > 0;
+  if (backend->timestamps_available) {
+    VkQueryPoolCreateInfo query_info{};
+    query_info.sType = VK_STRUCTURE_TYPE_QUERY_POOL_CREATE_INFO;
+    query_info.queryType = VK_QUERY_TYPE_TIMESTAMP;
+    query_info.queryCount = 2;
+    if (vkCreateQueryPool(backend->device, &query_info, nullptr,
+                          &backend->query_pool) != VK_SUCCESS) {
+      backend->timestamps_available = false;
+    }
+  }
+  return true;
+}
+
+static bool initialize_locked(Lardon3DOrbVulkanBackend *backend) {
+  if (backend->state == BackendState::kAvailable) {
+    return true;
+  }
+  if (backend->state != BackendState::kUninitialized) {
+    return false;
+  }
+  const char *disabled = std::getenv("LARDON3D_VULKAN_DISABLE");
+  if (disabled && std::strcmp(disabled, "1") == 0) {
+    backend->state = BackendState::kUnavailable;
+    return false;
+  }
+  auto start = std::chrono::steady_clock::now();
+  bool success = create_instance(backend) && select_device(backend) &&
+                 create_device_and_commands(backend) && create_pipeline(backend) &&
+                 create_buffers_and_descriptors(backend);
+  backend->initialization_ns = elapsed_ns(start);
+  if (!success) {
+    destroy_vulkan(backend);
+    backend->state = BackendState::kUnavailable;
+    return false;
+  }
+  backend->state = BackendState::kAvailable;
+  return true;
+}
+
+static Lardon3DOrbVulkanResult fail_session_locked(
+    Lardon3DOrbVulkanBackend *backend) {
+  destroy_vulkan(backend);
+  backend->state = BackendState::kFailed;
+  return LARDON3D_ORB_VULKAN_FAILED;
+}
+
+static bool synchronize_host_write(Lardon3DOrbVulkanBackend *backend,
+                                   const Buffer &buffer, VkDeviceSize size) {
+  if (buffer.coherent || size == 0) {
+    return true;
+  }
+  VkMappedMemoryRange range{};
+  range.sType = VK_STRUCTURE_TYPE_MAPPED_MEMORY_RANGE;
+  range.memory = buffer.memory;
+  range.offset = 0;
+  range.size = VK_WHOLE_SIZE;
+  return vkFlushMappedMemoryRanges(backend->device, 1, &range) == VK_SUCCESS;
+}
+
+static bool synchronize_host_read(Lardon3DOrbVulkanBackend *backend,
+                                  const Buffer &buffer) {
+  if (buffer.coherent) {
+    return true;
+  }
+  VkMappedMemoryRange range{};
+  range.sType = VK_STRUCTURE_TYPE_MAPPED_MEMORY_RANGE;
+  range.memory = buffer.memory;
+  range.offset = 0;
+  range.size = VK_WHOLE_SIZE;
+  return vkInvalidateMappedMemoryRanges(backend->device, 1, &range) == VK_SUCCESS;
+}
+
+static VkResult record_and_submit(Lardon3DOrbVulkanBackend *backend,
+                                  uint32_t count_a, uint32_t count_b) {
+  VkResult result = vkResetCommandBuffer(backend->command_buffer, 0);
+  if (result != VK_SUCCESS) {
+    return result;
+  }
+  VkCommandBufferBeginInfo begin_info{};
+  begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO;
+  begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT;
+  result = vkBeginCommandBuffer(backend->command_buffer, &begin_info);
+  if (result != VK_SUCCESS) {
+    return result;
+  }
+  if (backend->timestamps_available) {
+    vkCmdResetQueryPool(backend->command_buffer, backend->query_pool, 0, 2);
+    vkCmdWriteTimestamp(backend->command_buffer, VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT,
+                        backend->query_pool, 0);
+  }
+  vkCmdBindPipeline(backend->command_buffer, VK_PIPELINE_BIND_POINT_COMPUTE,
+                    backend->pipeline);
+  vkCmdBindDescriptorSets(backend->command_buffer, VK_PIPELINE_BIND_POINT_COMPUTE,
+                          backend->pipeline_layout, 0, 1,
+                          &backend->descriptor_set, 0, nullptr);
+  uint32_t counts[2] = {count_a, count_b};
+  vkCmdPushConstants(backend->command_buffer, backend->pipeline_layout,
+                     VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(counts), counts);
+  uint32_t groups = (count_a + backend->workgroup_size - 1) /
+                    backend->workgroup_size;
+  vkCmdDispatch(backend->command_buffer, groups, 1, 1);
+  VkMemoryBarrier barrier{};
+  barrier.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER;
+  barrier.srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT;
+  barrier.dstAccessMask = VK_ACCESS_HOST_READ_BIT;
+  vkCmdPipelineBarrier(backend->command_buffer, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT,
+                       VK_PIPELINE_STAGE_HOST_BIT, 0, 1, &barrier, 0, nullptr, 0,
+                       nullptr);
+  if (backend->timestamps_available) {
+    vkCmdWriteTimestamp(backend->command_buffer, VK_PIPELINE_STAGE_BOTTOM_OF_PIPE_BIT,
+                        backend->query_pool, 1);
+  }
+  result = vkEndCommandBuffer(backend->command_buffer);
+  if (result != VK_SUCCESS) {
+    return result;
+  }
+  VkSubmitInfo submit_info{};
+  submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO;
+  submit_info.commandBufferCount = 1;
+  submit_info.pCommandBuffers = &backend->command_buffer;
+  result = vkQueueSubmit(backend->queue, 1, &submit_info, VK_NULL_HANDLE);
+  if (result != VK_SUCCESS) {
+    return result;
+  }
+  return vkQueueWaitIdle(backend->queue);
+}
+
+static void read_gpu_time(Lardon3DOrbVulkanBackend *backend) {
+  backend->last_gpu_ns = 0;
+  if (!backend->timestamps_available) {
+    return;
+  }
+  uint64_t timestamps[2]{};
+  VkResult result = vkGetQueryPoolResults(
+      backend->device, backend->query_pool, 0, 2, sizeof(timestamps), timestamps,
+      sizeof(uint64_t), VK_QUERY_RESULT_64_BIT | VK_QUERY_RESULT_WAIT_BIT);
+  if (result == VK_SUCCESS && timestamps[1] >= timestamps[0]) {
+    double nanoseconds = static_cast<double>(timestamps[1] - timestamps[0]) *
+                         backend->properties.limits.timestampPeriod;
+    backend->last_gpu_ns = static_cast<uint64_t>(nanoseconds);
+  }
+}
+
+}  // namespace
+
+extern "C" Lardon3DOrbVulkanBackend *lardon3d_orb_vulkan_backend_create(void) {
+  return new (std::nothrow) Lardon3DOrbVulkanBackend();
+}
+
+extern "C" void lardon3d_orb_vulkan_backend_destroy(
+    Lardon3DOrbVulkanBackend *backend) {
+  if (!backend) {
+    return;
+  }
+  {
+    std::lock_guard<std::mutex> lock(backend->mutex);
+    destroy_vulkan(backend);
+  }
+  delete backend;
+}
+
+extern "C" bool lardon3d_orb_vulkan_should_use(uint32_t feature_count_a,
+                                                uint32_t feature_count_b) {
+  if (feature_count_a == 0 || feature_count_b == 0 ||
+      feature_count_a > LARDON3D_FEATURE_MAX_FEATURES ||
+      feature_count_b > LARDON3D_FEATURE_MAX_FEATURES) {
+    return false;
+  }
+  return static_cast<uint64_t>(feature_count_a) * feature_count_b >=
+         kDefaultVulkanWorkThreshold;
+}
+
+extern "C" Lardon3DOrbVulkanResult lardon3d_orb_vulkan_top2(
+    Lardon3DOrbVulkanBackend *backend, const unsigned char *descriptors_a,
+    uint32_t feature_count_a, const unsigned char *descriptors_b,
+    uint32_t feature_count_b, Lardon3DOrbTop2 *output, size_t output_capacity) {
+  if (!backend || feature_count_a > LARDON3D_FEATURE_MAX_FEATURES ||
+      feature_count_b > LARDON3D_FEATURE_MAX_FEATURES ||
+      (feature_count_a > 0 && (!descriptors_a || !output ||
+                              output_capacity < feature_count_a)) ||
+      (feature_count_b > 0 && !descriptors_b)) {
+    return LARDON3D_ORB_VULKAN_INVALID_ARGUMENT;
+  }
+  if (feature_count_a == 0) {
+    return LARDON3D_ORB_VULKAN_OK;
+  }
+  if (feature_count_b == 0) {
+    for (uint32_t index = 0; index < feature_count_a; ++index) {
+      output[index] = Lardon3DOrbTop2{};
+    }
+    return LARDON3D_ORB_VULKAN_OK;
+  }
+
+  std::lock_guard<std::mutex> lock(backend->mutex);
+  if (!initialize_locked(backend)) {
+    return LARDON3D_ORB_VULKAN_UNAVAILABLE;
+  }
+  VkDeviceSize bytes_a = static_cast<VkDeviceSize>(feature_count_a) * 32;
+  VkDeviceSize bytes_b = static_cast<VkDeviceSize>(feature_count_b) * 32;
+  std::memcpy(backend->descriptors_a.mapping, descriptors_a,
+              static_cast<size_t>(bytes_a));
+  std::memcpy(backend->descriptors_b.mapping, descriptors_b,
+              static_cast<size_t>(bytes_b));
+  if (!synchronize_host_write(backend, backend->descriptors_a, bytes_a) ||
+      !synchronize_host_write(backend, backend->descriptors_b, bytes_b)) {
+    return fail_session_locked(backend);
+  }
+
+  auto start = std::chrono::steady_clock::now();
+#ifdef LARDON3D_ORB_VULKAN_TESTING
+  const char *force_failure = std::getenv("LARDON3D_TEST_VULKAN_DEVICE_LOST");
+  if (force_failure && std::strcmp(force_failure, "1") == 0) {
+    return fail_session_locked(backend);
+  }
+#endif
+  VkResult dispatch_result = record_and_submit(backend, feature_count_a,
+                                               feature_count_b);
+  backend->last_dispatch_ns = elapsed_ns(start);
+  if (dispatch_result != VK_SUCCESS ||
+      !synchronize_host_read(backend, backend->output)) {
+    return fail_session_locked(backend);
+  }
+  read_gpu_time(backend);
+
+  const RawTop2 *raw = static_cast<const RawTop2 *>(backend->output.mapping);
+  uint32_t neighbors = std::min(feature_count_b, 2U);
+  for (uint32_t index = 0; index < feature_count_a; ++index) {
+    output[index].neighbor_count = neighbors;
+    output[index].best_index = raw[index].best_index;
+    output[index].best_distance = raw[index].best_distance;
+    output[index].second_index = neighbors == 2 ? raw[index].second_index : 0;
+    output[index].second_distance = neighbors == 2 ? raw[index].second_distance : 0;
+  }
+  return LARDON3D_ORB_VULKAN_OK;
+}
+
+extern "C" bool lardon3d_orb_vulkan_backend_info(
+    Lardon3DOrbVulkanBackend *backend, Lardon3DOrbVulkanInfo *info) {
+  if (!backend || !info) {
+    return false;
+  }
+  std::lock_guard<std::mutex> lock(backend->mutex);
+  std::memset(info, 0, sizeof(*info));
+  info->available = backend->state == BackendState::kAvailable;
+  info->initialized = backend->state != BackendState::kUninitialized;
+  info->dedicated_compute_queue = backend->dedicated_compute_queue;
+  info->workgroup_size = backend->workgroup_size;
+  info->permanent_payload_bytes = static_cast<uint64_t>(
+      kDescriptorBufferBytes * 2 + kOutputBufferBytes);
+  info->initialization_ns = backend->initialization_ns;
+  info->dispatch_ns = backend->last_dispatch_ns;
+  info->gpu_ns = backend->last_gpu_ns;
+  if (backend->physical_device != VK_NULL_HANDLE) {
+    std::snprintf(info->device_name, sizeof(info->device_name), "%s",
+                  backend->properties.deviceName);
+  }
+  return true;
+}
+
+#else
+
+struct Lardon3DOrbVulkanBackend {};
+
+extern "C" Lardon3DOrbVulkanBackend *lardon3d_orb_vulkan_backend_create(void) {
+  return new (std::nothrow) Lardon3DOrbVulkanBackend();
+}
+
+extern "C" void lardon3d_orb_vulkan_backend_destroy(
+    Lardon3DOrbVulkanBackend *backend) {
+  delete backend;
+}
+
+extern "C" bool lardon3d_orb_vulkan_should_use(uint32_t, uint32_t) {
+  return false;
+}
+
+extern "C" Lardon3DOrbVulkanResult lardon3d_orb_vulkan_top2(
+    Lardon3DOrbVulkanBackend *backend, const unsigned char *, uint32_t feature_count_a,
+    const unsigned char *, uint32_t feature_count_b, Lardon3DOrbTop2 *output,
+    size_t output_capacity) {
+  if (!backend || (feature_count_a > 0 && (!output || output_capacity < feature_count_a))) {
+    return LARDON3D_ORB_VULKAN_INVALID_ARGUMENT;
+  }
+  if (feature_count_a == 0 || feature_count_b == 0) {
+    for (uint32_t index = 0; index < feature_count_a; ++index) {
+      output[index] = Lardon3DOrbTop2{};
+    }
+    return LARDON3D_ORB_VULKAN_OK;
+  }
+  return LARDON3D_ORB_VULKAN_UNAVAILABLE;
+}
+
+extern "C" bool lardon3d_orb_vulkan_backend_info(
+    Lardon3DOrbVulkanBackend *backend, Lardon3DOrbVulkanInfo *info) {
+  if (!backend || !info) {
+    return false;
+  }
+  std::memset(info, 0, sizeof(*info));
+  return true;
+}
+
+#endif

src/project.c

@@ -672,7 +672,8 @@ checkpoint_task_internal(Lardon3DAppState *state, const Lardon3DTask *task,
                          const Lardon3DProjectDbFeatureExtractTask *feature_parameters,
                          const Lardon3DProjectDbSiftExtractTask *sift_parameters,
                          const Lardon3DProjectDbVisualIndexUpdateTask *visual_parameters,
-                         const Lardon3DProjectDbCandidatePairGenerateTask *candidate_parameters) {
+                         const Lardon3DProjectDbCandidatePairGenerateTask *candidate_parameters,
+                         const Lardon3DProjectDbMatcherTask *matcher_parameters) {
   if (!state || !state->project_loaded || !state->project_db) {
     return LARDON3D_PROJECT_TASK_CHECKPOINT_NO_PROJECT;
   }
@@ -731,6 +732,10 @@ checkpoint_task_internal(Lardon3DAppState *state, const Lardon3DTask *task,
           ? lardon3d_project_db_record_candidate_pair_generate_task(
                 state->project_db, &snapshot, task_kind, task_kind_version, &checkpoint,
                 candidate_parameters, now.tv_sec)
+      : matcher_parameters
+          ? lardon3d_project_db_record_matcher_task(
+                state->project_db, &snapshot, task_kind, task_kind_version, &checkpoint,
+                matcher_parameters, now.tv_sec)
           : lardon3d_project_db_record_task(state->project_db, &snapshot, task_kind,
                                             task_kind_version, &checkpoint, now.tv_sec);
   if (recorded == LARDON3D_PROJECT_DB_BUSY) {
@@ -746,7 +751,7 @@ checkpoint_task_internal(Lardon3DAppState *state, const Lardon3DTask *task,
 
 Lardon3DProjectTaskCheckpointResult lardon3d_project_checkpoint_task(Lardon3DAppState *state,
                                                                      const Lardon3DTask *task) {
-  return checkpoint_task_internal(state, task, NULL, 0, NULL, NULL, NULL, NULL);
+  return checkpoint_task_internal(state, task, NULL, 0, NULL, NULL, NULL, NULL, NULL);
 }
 
 Lardon3DProjectTaskCheckpointResult
@@ -755,7 +760,8 @@ lardon3d_project_checkpoint_image_import_task(Lardon3DAppState *state, const Lar
   if (!source_path || !source_path[0] || scanset_id == 0) {
     return LARDON3D_PROJECT_TASK_CHECKPOINT_INVALID_TASK;
   }
-  return checkpoint_task_internal(state, task, source_path, scanset_id, NULL, NULL, NULL, NULL);
+  return checkpoint_task_internal(
+      state, task, source_path, scanset_id, NULL, NULL, NULL, NULL, NULL);
 }
 
 Lardon3DProjectTaskCheckpointResult lardon3d_project_checkpoint_feature_extract_task(
@@ -764,14 +770,14 @@ Lardon3DProjectTaskCheckpointResult lardon3d_project_checkpoint_feature_extract_
   if (!parameters) {
     return LARDON3D_PROJECT_TASK_CHECKPOINT_INVALID_TASK;
   }
-  return checkpoint_task_internal(state, task, NULL, 0, parameters, NULL, NULL, NULL);
+  return checkpoint_task_internal(state, task, NULL, 0, parameters, NULL, NULL, NULL, NULL);
 }
 
 Lardon3DProjectTaskCheckpointResult lardon3d_project_checkpoint_sift_extract_task(
     Lardon3DAppState *state, const Lardon3DTask *task,
     const Lardon3DProjectDbSiftExtractTask *parameters) {
   if (!parameters) return LARDON3D_PROJECT_TASK_CHECKPOINT_INVALID_TASK;
-  return checkpoint_task_internal(state, task, NULL, 0, NULL, parameters, NULL, NULL);
+  return checkpoint_task_internal(state, task, NULL, 0, NULL, parameters, NULL, NULL, NULL);
 }
 
 Lardon3DProjectTaskCheckpointResult lardon3d_project_checkpoint_visual_index_update_task(
@@ -780,7 +786,7 @@ Lardon3DProjectTaskCheckpointResult lardon3d_project_checkpoint_visual_index_upd
   if (!parameters) {
     return LARDON3D_PROJECT_TASK_CHECKPOINT_INVALID_TASK;
   }
-  return checkpoint_task_internal(state, task, NULL, 0, NULL, NULL, parameters, NULL);
+  return checkpoint_task_internal(state, task, NULL, 0, NULL, NULL, parameters, NULL, NULL);
 }
 
 Lardon3DProjectTaskCheckpointResult lardon3d_project_checkpoint_candidate_pair_generate_task(
@@ -789,7 +795,16 @@ Lardon3DProjectTaskCheckpointResult lardon3d_project_checkpoint_candidate_pair_g
   if (!parameters) {
     return LARDON3D_PROJECT_TASK_CHECKPOINT_INVALID_TASK;
   }
-  return checkpoint_task_internal(state, task, NULL, 0, NULL, NULL, NULL, parameters);
+  return checkpoint_task_internal(state, task, NULL, 0, NULL, NULL, NULL, parameters, NULL);
+}
+
+Lardon3DProjectTaskCheckpointResult lardon3d_project_checkpoint_matcher_task(
+    Lardon3DAppState *state, const Lardon3DTask *task,
+    const Lardon3DProjectDbMatcherTask *parameters) {
+  if (!parameters) {
+    return LARDON3D_PROJECT_TASK_CHECKPOINT_INVALID_TASK;
+  }
+  return checkpoint_task_internal(state, task, NULL, 0, NULL, NULL, NULL, NULL, parameters);
 }
 
 static bool coherent_recovery(const Lardon3DProjectDbTask *database_task,
@@ -935,6 +950,7 @@ lardon3d_project_resume_recoverable_tasks(Lardon3DAppState *state,
           .project_path = state->project_path,
           .project_db = state->project_db,
           .resource_governor = state->resource_governor,
+          .orb_vulkan_backend = state->orb_vulkan_backend,
       };
       Lardon3DTask *task = NULL;
       Lardon3DTaskKindResult restored = lardon3d_task_kind_registry_restore(

src/project_db.c

@@ -226,6 +226,18 @@ static const char schema_match_result_v10[] =
     "CREATE INDEX match_results_feature_set_a_idx ON match_results(feature_set_id_a);"
     "CREATE INDEX match_results_feature_set_b_idx ON match_results(feature_set_id_b);";
 
+static const char schema_matcher_task_v11[] =
+    "CREATE TABLE matcher_tasks("
+    "task_id INTEGER PRIMARY KEY REFERENCES tasks(task_id) ON DELETE CASCADE,"
+    "after_candidate_pair_id INTEGER NOT NULL CHECK(after_candidate_pair_id>=0),"
+    "feature_extractor_kind TEXT NOT NULL CHECK(length(feature_extractor_kind)>0 AND "
+    "length(feature_extractor_kind)<65),"
+    "feature_extractor_version INTEGER NOT NULL CHECK(feature_extractor_version>0),"
+    "feature_parameter_fingerprint BLOB NOT NULL "
+    "CHECK(length(feature_parameter_fingerprint)=32),"
+    "matcher_kind INTEGER NOT NULL CHECK(matcher_kind BETWEEN 0 AND 2),"
+    "ratio_threshold REAL NOT NULL CHECK(ratio_threshold>0.0 AND ratio_threshold<1.0));";
+
 static void copy_error(char destination[LARDON3D_PROJECT_DB_ERROR_CAPACITY], const char *text) {
   if (destination) {
     (void)snprintf(destination, LARDON3D_PROJECT_DB_ERROR_CAPACITY, "%s", text ? text : "");
@@ -315,7 +327,7 @@ static Lardon3DProjectDbResult migrate(Lardon3DProjectDb *database, unsigned int
   }
   if (from_version != 0 && from_version != 1 && from_version != 2 && from_version != 3 &&
       from_version != 4 && from_version != 5 && from_version != 6 && from_version != 7 &&
-      from_version != 8 && from_version != 9) {
+      from_version != 8 && from_version != 9 && from_version != 10) {
     return LARDON3D_PROJECT_DB_CORRUPT;
   }
   Lardon3DProjectDbResult result = execute(database, "BEGIN IMMEDIATE", "begin migration");
@@ -541,6 +553,21 @@ static Lardon3DProjectDbResult migrate(Lardon3DProjectDb *database, unsigned int
                        "finish schema v10 migration");
     }
   }
+  if (result == LARDON3D_PROJECT_DB_OK && from_version < 11) {
+    result = execute(database, schema_matcher_task_v11, "migrate schema v10 to v11");
+#ifdef LARDON3D_PROJECT_DB_TESTING
+    const char *forced_failure = getenv("LARDON3D_TEST_PROJECT_DB_FAIL_MIGRATION_V11");
+    if (result == LARDON3D_PROJECT_DB_OK && forced_failure && strcmp(forced_failure, "1") == 0) {
+      result = execute(database, "INSERT INTO missing_test_table VALUES(1)",
+                       "forced migration v11 failure");
+    }
+#endif
+    if (result == LARDON3D_PROJECT_DB_OK) {
+      result = execute(database,
+                       "UPDATE metadata SET value=11 WHERE key='schema_version' AND value=10",
+                       "finish schema v11 migration");
+    }
+  }
   if (result == LARDON3D_PROJECT_DB_OK) {
     result = execute(database, "COMMIT", "commit migration");
   }
@@ -654,8 +681,11 @@ Lardon3DProjectDbResult lardon3d_project_db_open(const char *path, Lardon3DProje
                               "feature_support_groups",
                               "feature_support_members",
                               "candidate_pairs",
+                              "matcher_tasks",
                               "match_results"};
-    for (size_t index = 0; index < 17 && result == LARDON3D_PROJECT_DB_OK; ++index) {
+    for (size_t index = 0; index < sizeof(required) / sizeof(required[0]) &&
+                           result == LARDON3D_PROJECT_DB_OK;
+         ++index) {
       if (!table_exists(database->connection, required[index])) {
         copy_error(database->error, "Schéma v1 incomplet.");
         result = LARDON3D_PROJECT_DB_CORRUPT;
@@ -843,6 +873,7 @@ record_task_internal(Lardon3DProjectDb *database, const Lardon3DTaskDurableSnaps
                      const Lardon3DProjectDbSiftExtractTask *sift,
                      const Lardon3DProjectDbVisualIndexUpdateTask *visual,
                      const Lardon3DProjectDbCandidatePairGenerateTask *candidate_pair,
+                     const Lardon3DProjectDbMatcherTask *matcher,
                      int64_t updated_at) {
   bool typed = task_kind != NULL;
   if (!database || !valid_durable_task(snapshot, updated_at) ||
@@ -879,6 +910,13 @@ record_task_internal(Lardon3DProjectDb *database, const Lardon3DTaskDurableSnaps
         candidate_pair->top_k == 0 || candidate_pair->top_k > 256 ||
         candidate_pair->minimum_evidence_count > 1024 ||
         candidate_pair->scanset_filter < 0 || candidate_pair->scanset_filter > 2)) ||
+      (matcher &&
+       (!valid_task_id(matcher->task_id) || matcher->task_id != snapshot->id ||
+        matcher->after_candidate_pair_id > INT64_MAX ||
+        !lardon3d_task_kind_is_valid(matcher->feature_extractor_kind) ||
+        matcher->feature_extractor_version == 0 || matcher->matcher_kind < 0 ||
+        matcher->matcher_kind > 2 || !isfinite(matcher->ratio_threshold) ||
+        matcher->ratio_threshold <= 0.0F || matcher->ratio_threshold >= 1.0F)) ||
       (checkpoint && !valid_checkpoint(checkpoint))) {
     return LARDON3D_PROJECT_DB_INVALID_ARGUMENT;
   }
@@ -1046,6 +1084,35 @@ record_task_internal(Lardon3DProjectDb *database, const Lardon3DTaskDurableSnaps
       }
     }
   }
+  if (result == LARDON3D_PROJECT_DB_OK && matcher) {
+    result = prepare(
+        database,
+        "INSERT INTO matcher_tasks(task_id,after_candidate_pair_id,feature_extractor_kind,"
+        "feature_extractor_version,feature_parameter_fingerprint,matcher_kind,ratio_threshold) "
+        "VALUES(?1,?2,?3,?4,?5,?6,?7) ON CONFLICT(task_id) DO UPDATE SET "
+        "after_candidate_pair_id=excluded.after_candidate_pair_id WHERE "
+        "matcher_tasks.feature_extractor_kind=excluded.feature_extractor_kind AND "
+        "matcher_tasks.feature_extractor_version=excluded.feature_extractor_version AND "
+        "matcher_tasks.feature_parameter_fingerprint=excluded.feature_parameter_fingerprint AND "
+        "matcher_tasks.matcher_kind=excluded.matcher_kind AND "
+        "matcher_tasks.ratio_threshold=excluded.ratio_threshold",
+        &statement);
+    if (result == LARDON3D_PROJECT_DB_OK) {
+      sqlite3_bind_int64(statement, 1, (sqlite3_int64)matcher->task_id);
+      sqlite3_bind_int64(statement, 2, (sqlite3_int64)matcher->after_candidate_pair_id);
+      sqlite3_bind_text(statement, 3, matcher->feature_extractor_kind, -1, SQLITE_TRANSIENT);
+      sqlite3_bind_int64(statement, 4, matcher->feature_extractor_version);
+      sqlite3_bind_blob(statement, 5, matcher->feature_parameter_fingerprint, 32,
+                        SQLITE_TRANSIENT);
+      sqlite3_bind_int(statement, 6, matcher->matcher_kind);
+      sqlite3_bind_double(statement, 7, matcher->ratio_threshold);
+      result = step_done(database, statement, "upsert matcher task");
+      if (result == LARDON3D_PROJECT_DB_OK && sqlite3_changes(database->connection) != 1) {
+        copy_error(database->error, "Configuration Matcher de tâche immuable.");
+        result = LARDON3D_PROJECT_DB_CONSTRAINT;
+      }
+    }
+  }
   if (result == LARDON3D_PROJECT_DB_OK && sift) {
     result = prepare(
         database,
@@ -1115,7 +1182,7 @@ Lardon3DProjectDbResult lardon3d_project_db_record_task(
     Lardon3DProjectDb *database, const Lardon3DTaskDurableSnapshot *snapshot, const char *task_kind,
     uint32_t task_kind_version, const Lardon3DProjectDbCheckpoint *checkpoint, int64_t updated_at) {
   return record_task_internal(database, snapshot, task_kind, task_kind_version, checkpoint, NULL, 0,
-                              NULL, NULL, NULL, NULL, updated_at);
+                              NULL, NULL, NULL, NULL, NULL, updated_at);
 }
 
 Lardon3DProjectDbResult lardon3d_project_db_record_image_import_task(
@@ -1126,7 +1193,7 @@ Lardon3DProjectDbResult lardon3d_project_db_record_image_import_task(
     return LARDON3D_PROJECT_DB_INVALID_ARGUMENT;
   }
   return record_task_internal(database, snapshot, task_kind, task_kind_version, checkpoint,
-                              source_path, scanset_id, NULL, NULL, NULL, NULL, updated_at);
+                              source_path, scanset_id, NULL, NULL, NULL, NULL, NULL, updated_at);
 }
 
 Lardon3DProjectDbResult lardon3d_project_db_record_feature_extract_task(
@@ -1137,7 +1204,7 @@ Lardon3DProjectDbResult lardon3d_project_db_record_feature_extract_task(
     return LARDON3D_PROJECT_DB_INVALID_ARGUMENT;
   }
   return record_task_internal(database, snapshot, task_kind, task_kind_version, checkpoint, NULL, 0,
-                              parameters, NULL, NULL, NULL, updated_at);
+                              parameters, NULL, NULL, NULL, NULL, updated_at);
 }
 
 Lardon3DProjectDbResult lardon3d_project_db_record_sift_extract_task(
@@ -1147,7 +1214,7 @@ Lardon3DProjectDbResult lardon3d_project_db_record_sift_extract_task(
     const Lardon3DProjectDbSiftExtractTask *parameters, int64_t updated_at) {
   if (!parameters) return LARDON3D_PROJECT_DB_INVALID_ARGUMENT;
   return record_task_internal(database, snapshot, task_kind, task_kind_version, checkpoint, NULL, 0,
-                              NULL, parameters, NULL, NULL, updated_at);
+                              NULL, parameters, NULL, NULL, NULL, updated_at);
 }
 
 static bool read_task(sqlite3_stmt *statement, Lardon3DProjectDbTask *task) {
@@ -3129,7 +3196,7 @@ Lardon3DProjectDbResult lardon3d_project_db_record_visual_index_update_task(
     return LARDON3D_PROJECT_DB_INVALID_ARGUMENT;
   }
   return record_task_internal(db, snapshot, kind, version, checkpoint, NULL, 0, NULL,
-                              NULL, parameters, NULL, updated_at);
+                              NULL, parameters, NULL, NULL, updated_at);
 }
 
 Lardon3DProjectDbResult lardon3d_project_db_load_visual_index_update_task(
@@ -3176,7 +3243,7 @@ Lardon3DProjectDbResult lardon3d_project_db_record_candidate_pair_generate_task(
     return LARDON3D_PROJECT_DB_INVALID_ARGUMENT;
   }
   return record_task_internal(db, snapshot, kind, version, checkpoint, NULL, 0, NULL,
-                              NULL, NULL, parameters, updated_at);
+                              NULL, NULL, parameters, NULL, updated_at);
 }
 
 Lardon3DProjectDbResult lardon3d_project_db_load_candidate_pair_generate_task(
@@ -3223,6 +3290,64 @@ Lardon3DProjectDbResult lardon3d_project_db_load_candidate_pair_generate_task(
   return result;
 }
 
+Lardon3DProjectDbResult lardon3d_project_db_record_matcher_task(
+    Lardon3DProjectDb *db, const Lardon3DTaskDurableSnapshot *snapshot,
+    const char *kind, uint32_t version,
+    const Lardon3DProjectDbCheckpoint *checkpoint,
+    const Lardon3DProjectDbMatcherTask *parameters, int64_t updated_at) {
+  if (!snapshot || !parameters || parameters->task_id != snapshot->id) {
+    return LARDON3D_PROJECT_DB_INVALID_ARGUMENT;
+  }
+  return record_task_internal(db, snapshot, kind, version, checkpoint, NULL, 0, NULL,
+                              NULL, NULL, NULL, parameters, updated_at);
+}
+
+Lardon3DProjectDbResult lardon3d_project_db_load_matcher_task(
+    Lardon3DProjectDb *db, uint64_t task_id,
+    Lardon3DProjectDbMatcherTask *parameters) {
+  if (!db || !valid_task_id(task_id) || !parameters) {
+    return LARDON3D_PROJECT_DB_INVALID_ARGUMENT;
+  }
+  memset(parameters, 0, sizeof(*parameters));
+  (void)pthread_mutex_lock(&db->mutex);
+  sqlite3_stmt *statement = NULL;
+  Lardon3DProjectDbResult result = prepare(
+      db,
+      "SELECT after_candidate_pair_id,feature_extractor_kind,feature_extractor_version,"
+      "feature_parameter_fingerprint,matcher_kind,ratio_threshold FROM matcher_tasks "
+      "WHERE task_id=?1",
+      &statement);
+  if (result == LARDON3D_PROJECT_DB_OK) {
+    sqlite3_bind_int64(statement, 1, (sqlite3_int64)task_id);
+    int code = sqlite3_step(statement);
+    sqlite3_int64 after = sqlite3_column_int64(statement, 0);
+    sqlite3_int64 extractor_version = sqlite3_column_int64(statement, 2);
+    int matcher_kind = sqlite3_column_int(statement, 4);
+    double ratio = sqlite3_column_double(statement, 5);
+    if (code == SQLITE_DONE) {
+      result = LARDON3D_PROJECT_DB_NOT_FOUND;
+    } else if (code != SQLITE_ROW || after < 0 || extractor_version <= 0 ||
+               extractor_version > UINT32_MAX || matcher_kind < 0 || matcher_kind > 2 ||
+               !isfinite(ratio) || ratio <= 0.0 || ratio >= 1.0 ||
+               sqlite3_column_bytes(statement, 3) != LARDON3D_PROJECT_DB_SHA256_SIZE ||
+               !copy_column(statement, 1, parameters->feature_extractor_kind,
+                            sizeof(parameters->feature_extractor_kind))) {
+      result = LARDON3D_PROJECT_DB_CORRUPT;
+    } else {
+      parameters->task_id = task_id;
+      parameters->after_candidate_pair_id = (uint64_t)after;
+      parameters->feature_extractor_version = (uint32_t)extractor_version;
+      memcpy(parameters->feature_parameter_fingerprint,
+             sqlite3_column_blob(statement, 3), LARDON3D_PROJECT_DB_SHA256_SIZE);
+      parameters->matcher_kind = matcher_kind;
+      parameters->ratio_threshold = (float)ratio;
+    }
+    sqlite3_finalize(statement);
+  }
+  (void)pthread_mutex_unlock(&db->mutex);
+  return result;
+}
+
 static bool read_match_result(sqlite3_stmt *statement,
                               Lardon3DProjectDbMatchResult *result) {
   sqlite3_int64 id = sqlite3_column_int64(statement, 0);

src/resource_governor.c

@@ -38,6 +38,15 @@ struct Lardon3DResourceGovernor {
     size_t batch_metrics_count[LARDON3D_RESOURCE_TASK_MIXED + 1];
     size_t batch_metrics_head[LARDON3D_RESOURCE_TASK_MIXED + 1];
     size_t active_count;
+    bool swap_baseline_known;
+    uint64_t last_swap_pages_in;
+    uint64_t last_swap_pages_out;
+    unsigned int pressure_streak;
+    unsigned int recovery_streak;
+    unsigned int slow_start_streak;
+    size_t slow_start_limit;
+    bool slow_start_active;
+    Lardon3DResourcePressure pressure;
     Lardon3DResourceReservation *active;
     Lardon3DResourceReservation *released;
 };
@@ -60,9 +69,15 @@ valid_policy(
 {
     return valid_profile(profile) && policy
         && policy->system_memory_reserve_bytes < profile->memory_total_bytes
+        && policy->emergency_memory_floor_bytes
+            <= policy->system_memory_reserve_bytes
         && policy->system_cpu_reserve < profile->logical_cpu_count
         && policy->maximum_cpu_load_ratio > 0.0
         && policy->maximum_cpu_load_ratio <= 1.0
+        && policy->maximum_cpu_pressure_avg10 >= 0.0
+        && policy->maximum_cpu_pressure_avg10 <= 100.0
+        && policy->maximum_memory_pressure_avg10 >= 0.0
+        && policy->maximum_memory_pressure_avg10 <= 100.0
         && policy->maximum_io_pressure_avg10 >= 0.0
         && policy->maximum_io_pressure_avg10 <= 100.0
         && policy->io_slot_capacity > 0
@@ -79,6 +94,12 @@ valid_snapshot(
 )
 {
     return snapshot && snapshot->cpu_load_1m >= 0.0
+        && (!snapshot->cpu_pressure_known
+            || (snapshot->cpu_pressure_avg10 >= 0.0
+                && snapshot->cpu_pressure_avg10 <= 100.0))
+        && (!snapshot->memory_pressure_known
+            || (snapshot->memory_pressure_avg10 >= 0.0
+                && snapshot->memory_pressure_avg10 <= 100.0))
         && (!snapshot->io_pressure_known
             || (snapshot->io_pressure_avg10 >= 0.0
                 && snapshot->io_pressure_avg10 <= 100.0))
@@ -278,13 +299,20 @@ lardon3d_resource_policy_default(
     if (!valid_profile(profile) || !policy) {
         return false;
     }
+    unsigned int cpu_reserve = profile->logical_cpu_count / 4;
+    if (cpu_reserve == 0 && profile->logical_cpu_count > 2) {
+        cpu_reserve = 1;
+    }
     *policy = (Lardon3DResourcePolicy) {
-        .system_memory_reserve_bytes = profile->memory_total_bytes / 8,
+        .system_memory_reserve_bytes = profile->memory_total_bytes / 4,
+        .emergency_memory_floor_bytes = profile->memory_total_bytes / 8,
         .gpu_memory_reserve_bytes = profile->gpu_memory_known
             ? profile->gpu_memory_total_bytes / 8
             : 0,
-        .system_cpu_reserve = profile->logical_cpu_count > 2 ? 1 : 0,
+        .system_cpu_reserve = cpu_reserve,
         .maximum_cpu_load_ratio = 0.90,
+        .maximum_cpu_pressure_avg10 = 20.0,
+        .maximum_memory_pressure_avg10 = 1.0,
         .maximum_io_pressure_avg10 = 80.0,
         .gpu_slot_capacity = profile->gpu_available ? 1 : 0,
         .io_slot_capacity = 1,
@@ -332,6 +360,8 @@ lardon3d_resource_governor_create(
     governor->profile = *profile;
     governor->policy = *policy;
     governor->next_reservation_id = 1;
+    governor->slow_start_limit = SIZE_MAX;
+    governor->pressure = LARDON3D_RESOURCE_PRESSURE_GREEN;
     return governor;
 }
 
@@ -542,6 +572,93 @@ evaluate_locked(
         set_decision(decision, LARDON3D_RESOURCE_REJECT, 0, 0, 0, 0, "Estimation de ressources invalide.");
         return;
     }
+    bool swap_changed = false;
+    if (snapshot->swap_activity_known) {
+        if (governor->swap_baseline_known) {
+            swap_changed = snapshot->swap_pages_in > governor->last_swap_pages_in
+                || snapshot->swap_pages_out > governor->last_swap_pages_out;
+        }
+        governor->last_swap_pages_in = snapshot->swap_pages_in;
+        governor->last_swap_pages_out = snapshot->swap_pages_out;
+        governor->swap_baseline_known = true;
+    }
+    bool hard_memory_pressure = governor->policy.emergency_memory_floor_bytes > 0
+        && snapshot->memory_available_bytes
+            <= governor->policy.emergency_memory_floor_bytes;
+    bool soft_memory_pressure = governor->policy.system_memory_reserve_bytes > 0
+        && snapshot->memory_available_bytes
+            <= governor->policy.system_memory_reserve_bytes;
+    bool psi_pressure = (governor->policy.maximum_cpu_pressure_avg10 > 0.0
+            && snapshot->cpu_pressure_known
+            && snapshot->cpu_pressure_avg10
+                >= governor->policy.maximum_cpu_pressure_avg10)
+        || (governor->policy.maximum_memory_pressure_avg10 > 0.0
+            && snapshot->memory_pressure_known
+            && snapshot->memory_pressure_avg10
+                >= governor->policy.maximum_memory_pressure_avg10)
+        || (governor->policy.maximum_io_pressure_avg10 > 0.0
+            && snapshot->io_pressure_known
+            && snapshot->io_pressure_avg10
+                >= governor->policy.maximum_io_pressure_avg10);
+    bool pressure_signal = soft_memory_pressure || psi_pressure || swap_changed;
+
+    if (hard_memory_pressure) {
+        governor->pressure = LARDON3D_RESOURCE_PRESSURE_RED;
+        governor->pressure_streak = 0;
+        governor->recovery_streak = 0;
+        governor->slow_start_streak = 0;
+        governor->slow_start_limit = 1;
+        governor->slow_start_active = true;
+    } else if (pressure_signal) {
+        governor->recovery_streak = 0;
+        governor->slow_start_streak = 0;
+        ++governor->pressure_streak;
+        if (governor->pressure == LARDON3D_RESOURCE_PRESSURE_RED
+            || governor->pressure_streak >= 2) {
+            governor->pressure = LARDON3D_RESOURCE_PRESSURE_RED;
+            governor->slow_start_limit = 1;
+            governor->slow_start_active = true;
+        } else {
+            governor->pressure = LARDON3D_RESOURCE_PRESSURE_YELLOW;
+        }
+    } else {
+        governor->pressure_streak = 0;
+        if (governor->pressure == LARDON3D_RESOURCE_PRESSURE_RED) {
+            ++governor->recovery_streak;
+            if (governor->recovery_streak >= 3) {
+                governor->pressure = LARDON3D_RESOURCE_PRESSURE_YELLOW;
+                governor->recovery_streak = 0;
+            }
+        } else if (governor->pressure == LARDON3D_RESOURCE_PRESSURE_YELLOW) {
+            if (!governor->slow_start_active) {
+                governor->slow_start_active = true;
+                governor->slow_start_limit = 1;
+                governor->slow_start_streak = 0;
+            }
+            ++governor->recovery_streak;
+            if (governor->recovery_streak >= 3) {
+                governor->pressure = LARDON3D_RESOURCE_PRESSURE_GREEN;
+                governor->recovery_streak = 0;
+                governor->slow_start_streak = 0;
+            }
+        } else if (governor->slow_start_active) {
+            ++governor->slow_start_streak;
+            if (governor->slow_start_streak >= 3) {
+                governor->slow_start_streak = 0;
+                if (governor->slow_start_limit > SIZE_MAX / 2) {
+                    governor->slow_start_limit = SIZE_MAX;
+                    governor->slow_start_active = false;
+                } else {
+                    governor->slow_start_limit *= 2;
+                }
+            }
+        }
+    }
+    if (governor->pressure == LARDON3D_RESOURCE_PRESSURE_RED) {
+        set_decision(decision, LARDON3D_RESOURCE_WAIT, 0, 0, 0, 0,
+            "Pression mémoire ou swap persistante.");
+        return;
+    }
     if ((estimate->desired_gpu_slots > 0
             || estimate->gpu_memory_fixed_bytes > 0
             || estimate->gpu_memory_bytes_per_item > 0)
@@ -594,6 +711,22 @@ evaluate_locked(
         set_decision(decision, LARDON3D_RESOURCE_WAIT, 0, 0, 0, 0, "Charge CPU trop élevée.");
         return;
     }
+    if (governor->policy.maximum_cpu_pressure_avg10 > 0.0
+        && snapshot->cpu_pressure_known
+        && snapshot->cpu_pressure_avg10
+            >= governor->policy.maximum_cpu_pressure_avg10) {
+        set_decision(decision, LARDON3D_RESOURCE_WAIT, 0, 0, 0, 0,
+            "Pression CPU trop élevée.");
+        return;
+    }
+    if (governor->policy.maximum_memory_pressure_avg10 > 0.0
+        && snapshot->memory_pressure_known
+        && snapshot->memory_pressure_avg10
+            >= governor->policy.maximum_memory_pressure_avg10) {
+        set_decision(decision, LARDON3D_RESOURCE_WAIT, 0, 0, 0, 0,
+            "Pression mémoire trop élevée.");
+        return;
+    }
     if (estimate->desired_io_slots > 0 && snapshot->io_pressure_known
         && snapshot->io_pressure_avg10
             >= governor->policy.maximum_io_pressure_avg10) {
@@ -640,9 +773,27 @@ evaluate_locked(
     if (adapted_maximum < estimate->minimum_batch_size) {
         adapted_maximum = estimate->minimum_batch_size;
     }
+    if (governor->slow_start_active) {
+        adapted_maximum = minimum_size(adapted_maximum, governor->slow_start_limit);
+    }
+    if (governor->pressure == LARDON3D_RESOURCE_PRESSURE_YELLOW
+        && adapted_maximum > estimate->minimum_batch_size) {
+        adapted_maximum /= 2;
+        if (adapted_maximum < estimate->minimum_batch_size) {
+            adapted_maximum = estimate->minimum_batch_size;
+        }
+    }
     batch = minimum_size(batch, adapted_maximum);
     if (batch < estimate->minimum_batch_size) {
-        set_decision(decision, LARDON3D_RESOURCE_WAIT, 0, 0, 0, 0, "Ressources déjà réservées ou temporairement insuffisantes.");
+        set_decision(
+            decision,
+            LARDON3D_RESOURCE_WAIT,
+            0,
+            0,
+            0,
+            0,
+            "Ressources déjà réservées ou temporairement insuffisantes."
+        );
         return;
     }
     if (available.cpu_available == 0
@@ -1017,3 +1168,15 @@ lardon3d_resource_decision_name(Lardon3DResourceDecisionKind kind)
         return "Inconnue";
     }
 }
+
+Lardon3DResourcePressure
+lardon3d_resource_governor_pressure(Lardon3DResourceGovernor *governor)
+{
+    if (!governor) {
+        return LARDON3D_RESOURCE_PRESSURE_RED;
+    }
+    (void)pthread_mutex_lock(&governor->mutex);
+    Lardon3DResourcePressure pressure = governor->pressure;
+    (void)pthread_mutex_unlock(&governor->mutex);
+    return pressure;
+}

src/resource_snapshot.c

@@ -98,18 +98,38 @@ capture_load(Lardon3DResourceSnapshot *snapshot)
     ) == 3;
 }
 
-static void
+static bool
-capture_io_pressure(Lardon3DResourceSnapshot *snapshot)
+capture_pressure(const char *path, double *average)
 {
     char buffer[512];
-    if (!read_file("/proc/pressure/io", buffer, sizeof(buffer))) {
+    if (!read_file(path, buffer, sizeof(buffer))) {
-        return;
+        return false;
     }
-    double average;
+    return sscanf(buffer, "some avg10=%lf", average) == 1
-    if (sscanf(buffer, "some avg10=%lf", &average) == 1 && average >= 0.0) {
+        && *average >= 0.0 && *average <= 100.0;
-        snapshot->io_pressure_known = true;
+}
-        snapshot->io_pressure_avg10 = average;
+
+static bool
+vmstat_counter(const char *buffer, const char *key, uint64_t *value)
+{
+    const char *line = buffer;
+    size_t key_length = strlen(key);
+    while (*line) {
+        if (strncmp(line, key, key_length) == 0 && line[key_length] == ' ') {
+            errno = 0;
+            char *end;
+            unsigned long long parsed = strtoull(line + key_length + 1, &end, 10);
+            if (errno == 0 && end != line + key_length + 1) {
+                *value = (uint64_t)parsed;
+                return true;
             }
+            return false;
+        }
+        const char *newline = strchr(line, '\n');
+        if (!newline) break;
+        line = newline + 1;
+    }
+    return false;
 }
 
 static void
@@ -186,7 +206,17 @@ lardon3d_resource_snapshot_capture(
         set_error(error_message, error_message_size, "Instantané système impossible.");
         return false;
     }
-    capture_io_pressure(snapshot);
+    snapshot->cpu_pressure_known = capture_pressure(
+        "/proc/pressure/cpu", &snapshot->cpu_pressure_avg10);
+    snapshot->memory_pressure_known = capture_pressure(
+        "/proc/pressure/memory", &snapshot->memory_pressure_avg10);
+    snapshot->io_pressure_known = capture_pressure(
+        "/proc/pressure/io", &snapshot->io_pressure_avg10);
+    if (read_file("/proc/vmstat", buffer, sizeof(buffer))) {
+        snapshot->swap_activity_known = vmstat_counter(
+            buffer, "pswpin", &snapshot->swap_pages_in)
+            && vmstat_counter(buffer, "pswpout", &snapshot->swap_pages_out);
+    }
     capture_gpu_memory(profile, snapshot);
     return true;
 }

src/sift_task.c

@@ -234,8 +234,12 @@ Lardon3DTask *lardon3d_project_create_sift_extract_task(
   snprintf(durable.extractor_kind, sizeof(durable.extractor_kind), "%s",
            parameters->rootsift ? LARDON3D_ROOTSIFT_EXTRACTOR_KIND : LARDON3D_SIFT_EXTRACTOR_KIND);
   lardon3d_sift_extractor_parameter_fingerprint(parameters, durable.parameter_fingerprint);
-  Lardon3DTaskReconstructionContext runtime = {state->project_path, state->project_db,
+  Lardon3DTaskReconstructionContext runtime = {
-                                               state->resource_governor};
+      .project_path = state->project_path,
+      .project_db = state->project_db,
+      .resource_governor = state->resource_governor,
+      .orb_vulkan_backend = state->orb_vulkan_backend,
+  };
   SiftTaskContext *context = make_context(&runtime, &durable);
   if (!context) return NULL;
   Lardon3DResourceEstimate estimate = {.memory_fixed_bytes = 64ULL * 1024 * 1024,

src/task_kinds.c

@@ -1,6 +1,7 @@
 #include <lardon3d/candidate_pair_task.h>
 #include <lardon3d/feature_task.h>
 #include <lardon3d/import_task.h>
+#include <lardon3d/matcher_task.h>
 #include <lardon3d/sift_task.h>
 #include <lardon3d/visual_index_task.h>
 #include <lardon3d/task_kind_registry.h>
@@ -36,6 +37,11 @@ const Lardon3DTaskKindRegistry *lardon3d_task_kind_registry_production(void) {
           .kind = LARDON3D_CANDIDATE_PAIR_GENERATE_TASK_KIND,
           .kind_version = LARDON3D_CANDIDATE_PAIR_GENERATE_TASK_KIND_VERSION,
           .reconstruct = lardon3d_candidate_pair_generate_reconstruct,
+      },
+      {
+          .kind = LARDON3D_MATCHER_TASK_KIND,
+          .kind_version = LARDON3D_MATCHER_TASK_KIND_VERSION,
+          .reconstruct = lardon3d_matcher_task_reconstruct,
       }};
   static const Lardon3DTaskKindRegistry registry = {
       .descriptors = descriptors,

src/task_queue.c

@@ -40,7 +40,6 @@ terminal_state(Lardon3DTaskState state)
 static void
 unlink_pending(Lardon3DTaskQueue *queue, TaskNode *previous, TaskNode *node)
 {
-    bool was_full = queue->pending_count >= queue->capacity;
     if (previous) {
         previous->next_pending = node->next_pending;
     } else {
@@ -51,9 +50,10 @@ unlink_pending(Lardon3DTaskQueue *queue, TaskNode *previous, TaskNode *node)
     }
     node->next_pending = NULL;
     --queue->pending_count;
-    if (was_full) {
+    /* Chaque retrait libère une place. Plusieurs producteurs peuvent dormir
+     * pendant que le worker retire plusieurs tâches avant qu'ils reprennent
+     * le mutex ; chacun de ces retraits doit donc produire un réveil. */
     (void)pthread_cond_signal(&queue->not_full);
-    }
 }
 
 /* Parcourt la file d'attente et sélectionne la première tâche admissible.

src/visual_index_task.c

@@ -176,8 +176,12 @@ Lardon3DTask *lardon3d_project_create_visual_index_update_task(
       .visual_index_id = visual_index_id,
       .after_feature_set_id = 0,
   };
-  Lardon3DTaskReconstructionContext runtime = {state->project_path, state->project_db,
+  Lardon3DTaskReconstructionContext runtime = {
-                                               state->resource_governor};
+      .project_path = state->project_path,
+      .project_db = state->project_db,
+      .resource_governor = state->resource_governor,
+      .orb_vulkan_backend = state->orb_vulkan_backend,
+  };
   VisualIndexTaskContext *context = make_context(&runtime, &parameters);
   if (!context) {
     return NULL;

tests/benchmark_matcher.cpp

@@ -4,6 +4,7 @@
 #include <cstdlib>
 #include <cstring>
 #include <filesystem>
+#include <opencv2/core.hpp>
 #include <string>
 #include <vector>
 
@@ -109,7 +110,8 @@ static bool publish(Lardon3DAppState *state, uint64_t image_id, const char *kind
 }
 
 static bool benchmark_case(const char *kind, Lardon3DMatcherKind matcher_kind,
-                           Lardon3DFeatureDescriptorType type, uint32_t count) {
+                           Lardon3DFeatureDescriptorType type, uint32_t count,
+                           int repetitions) {
   char root_template[] = "/tmp/lardon3d-matcher-bench-XXXXXX";
   char *root = mkdtemp(root_template);
   if (!root) return false;
@@ -137,13 +139,14 @@ static bool benchmark_case(const char *kind, Lardon3DMatcherKind matcher_kind,
                  db, image_a.image_id, image_b.image_id, 1, &pair) == LARDON3D_PROJECT_DB_OK;
   std::vector<Sample> samples;
   Lardon3DMatcherParams params = {matcher_kind, lardon3d_matcher_default_ratio(matcher_kind)};
-  for (int repetition = 0; ok && repetition < 3; ++repetition) {
+  for (int repetition = 0; ok && repetition <= repetitions; ++repetition) {
     Lardon3DMatcherStats stats;
     params.ratio_threshold = lardon3d_matcher_default_ratio(matcher_kind) -
-                             (float)repetition * 0.01F;
+                             (float)repetition * 0.001F;
     Lardon3DProjectDbMatchResult result;
     ok = lardon3d_matcher_match_and_publish_profiled(
              root, db, &pair, &set_a, &set_b, &params, &result, &stats) == LARDON3D_MATCHER_OK;
+    if (repetition > 0)
       samples.push_back({stats.feature_open_ns, stats.descriptor_read_ns, stats.knn_ns,
                          stats.filter_ns, stats.canonicalize_ns, stats.serialize_ns,
                          stats.sha256_ns, stats.publication_ns, stats.database_ns,
@@ -184,17 +187,38 @@ static bool benchmark_case(const char *kind, Lardon3DMatcherKind matcher_kind,
   return ok;
 }
 
-int main(void) {
+int main(int argc, char **argv) {
+  int repetitions = argc > 3 ? atoi(argv[3]) : 7;
+  int threads = argc > 4 ? atoi(argv[4]) : cv::getNumThreads();
+  if (repetitions < 1 || threads < 1) return EXIT_FAILURE;
+  cv::setNumThreads(threads);
   printf("kind,count,open_ms,read_ms,knn_ms,filter_ms,canonicalize_ms,serialize_ms,"
          "sha_ms,publish_ms,db_ms,total_ms\n");
+  if (argc > 2) {
+    uint32_t count = (uint32_t)strtoul(argv[2], nullptr, 10);
+    if (strcmp(argv[1], "orb") == 0)
+      return benchmark_case("orb", LARDON3D_MATCHER_ORB_BF,
+                            LARDON3D_FEATURE_DESCRIPTOR_U8, count, repetitions)
+                 ? EXIT_SUCCESS
+                 : EXIT_FAILURE;
+    if (strcmp(argv[1], "sift") == 0 || strcmp(argv[1], "rootsift") == 0) {
+      bool rootsift = strcmp(argv[1], "rootsift") == 0;
+      return benchmark_case(argv[1], rootsift ? LARDON3D_MATCHER_ROOTSIFT_BF
+                                              : LARDON3D_MATCHER_SIFT_BF,
+                            LARDON3D_FEATURE_DESCRIPTOR_F32, count, repetitions)
+                 ? EXIT_SUCCESS
+                 : EXIT_FAILURE;
+    }
+    return EXIT_FAILURE;
+  }
   const uint32_t sizes[] = {64, 256, 1024, 4096, 8192};
   for (uint32_t count : sizes) {
     if (!benchmark_case("orb", LARDON3D_MATCHER_ORB_BF,
-                        LARDON3D_FEATURE_DESCRIPTOR_U8, count) ||
+                        LARDON3D_FEATURE_DESCRIPTOR_U8, count, repetitions) ||
         !benchmark_case("sift", LARDON3D_MATCHER_SIFT_BF,
-                        LARDON3D_FEATURE_DESCRIPTOR_F32, count) ||
+                        LARDON3D_FEATURE_DESCRIPTOR_F32, count, repetitions) ||
         !benchmark_case("rootsift", LARDON3D_MATCHER_ROOTSIFT_BF,
-                        LARDON3D_FEATURE_DESCRIPTOR_F32, count))
+                        LARDON3D_FEATURE_DESCRIPTOR_F32, count, repetitions))
       return EXIT_FAILURE;
   }
   return EXIT_SUCCESS;

tests/benchmark_orb_vulkan.cpp

@@ -0,0 +1,148 @@
+#include <lardon3d/orb_vulkan_backend.h>
+
+#include <opencv2/core.hpp>
+#include <opencv2/features2d.hpp>
+
+#include <algorithm>
+#include <chrono>
+#include <cstdint>
+#include <cstdio>
+#include <cstring>
+#include <iterator>
+#include <vector>
+
+namespace {
+
+constexpr uint32_t kDescriptorBytes = 32;
+
+static uint32_t next_random(uint32_t *state) {
+  *state = *state * 1664525U + 1013904223U;
+  return *state;
+}
+
+static std::vector<unsigned char> descriptors(uint32_t count, uint32_t seed) {
+  std::vector<unsigned char> result(static_cast<size_t>(count) * kDescriptorBytes);
+  for (unsigned char &value : result) {
+    value = static_cast<unsigned char>(next_random(&seed) >> 24);
+  }
+  return result;
+}
+
+template <typename Function>
+static double milliseconds(Function function) {
+  auto start = std::chrono::steady_clock::now();
+  function();
+  auto elapsed = std::chrono::steady_clock::now() - start;
+  return std::chrono::duration<double, std::milli>(elapsed).count();
+}
+
+static double median(std::vector<double> values) {
+  std::sort(values.begin(), values.end());
+  return values[values.size() / 2];
+}
+
+static double benchmark_cpu(const std::vector<unsigned char> &a,
+                            const std::vector<unsigned char> &b, uint32_t count) {
+  cv::Mat matrix_a(static_cast<int>(count), kDescriptorBytes, CV_8U,
+                   const_cast<unsigned char *>(a.data()));
+  cv::Mat matrix_b(static_cast<int>(count), kDescriptorBytes, CV_8U,
+                   const_cast<unsigned char *>(b.data()));
+  cv::BFMatcher matcher(cv::NORM_HAMMING, false);
+  std::vector<double> samples;
+  std::vector<std::vector<cv::DMatch>> output;
+  matcher.knnMatch(matrix_a, matrix_b, output, 2);
+  for (int repetition = 0; repetition < 7; ++repetition) {
+    samples.push_back(milliseconds([&] {
+      output.clear();
+      matcher.knnMatch(matrix_a, matrix_b, output, 2);
+    }));
+  }
+  return median(samples);
+}
+
+static bool benchmark_size(Lardon3DOrbVulkanBackend *backend, uint32_t count) {
+  std::vector<unsigned char> a = descriptors(count, 0x12340000U + count);
+  std::vector<unsigned char> b = descriptors(count, 0xabcd0000U + count);
+  std::vector<Lardon3DOrbTop2> output(count);
+  double cpu_ms = benchmark_cpu(a, b, count);
+  std::vector<double> samples;
+  for (int repetition = 0; repetition < 8; ++repetition) {
+    Lardon3DOrbVulkanResult result = LARDON3D_ORB_VULKAN_FAILED;
+    double elapsed = milliseconds([&] {
+      result = lardon3d_orb_vulkan_top2(backend, a.data(), count, b.data(), count,
+                                        output.data(), output.size());
+    });
+    if (result != LARDON3D_ORB_VULKAN_OK) {
+      return false;
+    }
+    if (repetition > 0) {
+      samples.push_back(elapsed);
+    }
+  }
+  Lardon3DOrbVulkanInfo info{};
+  if (!lardon3d_orb_vulkan_backend_info(backend, &info)) {
+    return false;
+  }
+  std::printf("%u cpu_ms=%.3f vulkan_ms=%.3f submit_ms=%.3f gpu_ms=%.3f "
+              "selector=%s\n",
+              count, cpu_ms, median(samples),
+              static_cast<double>(info.dispatch_ns) / 1.0e6,
+              static_cast<double>(info.gpu_ns) / 1.0e6,
+              lardon3d_orb_vulkan_should_use(count, count) ? "vulkan" : "cpu");
+  return true;
+}
+
+static bool benchmark_sustained(Lardon3DOrbVulkanBackend *backend) {
+  const uint32_t sizes[] = {1024, 4096, 8192, 4096};
+  std::vector<std::vector<unsigned char>> inputs_a;
+  std::vector<std::vector<unsigned char>> inputs_b;
+  std::vector<std::vector<Lardon3DOrbTop2>> outputs;
+  for (uint32_t size : sizes) {
+    inputs_a.push_back(descriptors(size, 0x98760000U + size));
+    inputs_b.push_back(descriptors(size, 0x67890000U + size));
+    outputs.emplace_back(size);
+  }
+  constexpr int repetitions = 5000;
+  double elapsed = milliseconds([&] {
+    for (int repetition = 0; repetition < repetitions; ++repetition) {
+      size_t slot = static_cast<size_t>(repetition) % std::size(sizes);
+      uint32_t size = sizes[slot];
+      Lardon3DOrbVulkanResult result = lardon3d_orb_vulkan_top2(
+          backend, inputs_a[slot].data(), size, inputs_b[slot].data(), size,
+          outputs[slot].data(), outputs[slot].size());
+      if (result != LARDON3D_ORB_VULKAN_OK) {
+        std::fprintf(stderr, "sustained dispatch failed at repetition %d\n", repetition);
+        std::abort();
+      }
+    }
+  });
+  std::printf("sustained jobs=%d total_ms=%.3f pairs_per_second=%.3f\n", repetitions,
+              elapsed, repetitions * 1000.0 / elapsed);
+  return true;
+}
+
+}  // namespace
+
+int main(int argc, char **argv) {
+  Lardon3DOrbVulkanBackend *backend = lardon3d_orb_vulkan_backend_create();
+  if (!backend) {
+    return 1;
+  }
+  bool ok = true;
+  const uint32_t sizes[] = {256, 512, 768, 1024, 4096, 8192};
+  for (uint32_t size : sizes) {
+    ok = benchmark_size(backend, size) && ok;
+  }
+  if (argc == 2 && std::strcmp(argv[1], "--sustained") == 0) {
+    ok = benchmark_sustained(backend) && ok;
+  }
+  Lardon3DOrbVulkanInfo info{};
+  if (lardon3d_orb_vulkan_backend_info(backend, &info)) {
+    std::printf("device=%s workgroup=%u cold_init_ms=%.3f payload=%llu\n",
+                info.device_name, info.workgroup_size,
+                static_cast<double>(info.initialization_ns) / 1.0e6,
+                static_cast<unsigned long long>(info.permanent_payload_bytes));
+  }
+  lardon3d_orb_vulkan_backend_destroy(backend);
+  return ok ? 0 : 1;
+}

tests/test_match_result.c

@@ -48,6 +48,7 @@ static bool create_v9_database(const char *path) {
   if (sqlite3_open(path, &connection) != SQLITE_OK) return false;
   static const char sql[] =
       "PRAGMA foreign_keys=OFF;BEGIN IMMEDIATE;"
+      "DROP TABLE matcher_tasks;"
       "DROP TABLE match_results;"
       "UPDATE metadata SET value=9 WHERE key='schema_version';COMMIT;PRAGMA foreign_keys=ON;";
   bool ok = sqlite3_exec(connection, sql, NULL, NULL, NULL) == SQLITE_OK;
@@ -92,7 +93,7 @@ static bool run_test(void) {
   char error[LARDON3D_PROJECT_DB_ERROR_CAPACITY];
   Lardon3DProjectDb *database = NULL;
   CHECK(lardon3d_project_db_open(database_path, &database, error) == LARDON3D_PROJECT_DB_OK);
-  CHECK(database && lardon3d_project_db_schema_version(database) == 10);
+  CHECK(database && lardon3d_project_db_schema_version(database) == 11);
 
   Lardon3DProjectDbScanSet scanset;
   CHECK(lardon3d_project_db_create_scanset(database, "Match-test", &scanset) ==
@@ -451,7 +452,7 @@ static bool run_test(void) {
   database = NULL;
 
   CHECK(lardon3d_project_db_open(database_path, &database, error) == LARDON3D_PROJECT_DB_OK);
-  CHECK(lardon3d_project_db_schema_version(database) == 10);
+  CHECK(lardon3d_project_db_schema_version(database) == 11);
 
   /* Verify persistence: load previously created results */
   CHECK(lardon3d_project_db_load_match_result(database, first_id, &loaded) ==
@@ -491,10 +492,10 @@ static bool run_test(void) {
   CHECK(create_v9_database(v9_path));
   CHECK(query_integer(v9_path, "SELECT value FROM metadata WHERE key='schema_version'", 9));
   CHECK(lardon3d_project_db_open(v9_path, &database, error) == LARDON3D_PROJECT_DB_OK);
-  CHECK(lardon3d_project_db_schema_version(database) == 10);
+  CHECK(lardon3d_project_db_schema_version(database) == 11);
   lardon3d_project_db_close(database);
   database = NULL;
-  CHECK(query_integer(v9_path, "SELECT value FROM metadata WHERE key='schema_version'", 10));
+  CHECK(query_integer(v9_path, "SELECT value FROM metadata WHERE key='schema_version'", 11));
   CHECK(query_integer(v9_path,
                       "SELECT count(*) FROM sqlite_master WHERE type='table' AND "
                       "name='match_results'", 1));
@@ -509,8 +510,11 @@ static bool run_test(void) {
   CHECK(query_integer(failed_v10_path,
                       "SELECT count(*) FROM sqlite_master WHERE type='table' AND "
                       "name='match_results'", 0));
+  CHECK(query_integer(failed_v10_path,
+                      "SELECT count(*) FROM sqlite_master WHERE type='table' AND "
+                      "name='matcher_tasks'", 0));
   CHECK(lardon3d_project_db_open(failed_v10_path, &database, error) == LARDON3D_PROJECT_DB_OK &&
-        lardon3d_project_db_schema_version(database) == 10);
+        lardon3d_project_db_schema_version(database) == 11);
   lardon3d_project_db_close(database);
   database = NULL;
 

tests/test_matcher_e2e.c

@@ -11,6 +11,8 @@
 #include <sys/stat.h>
 #include <unistd.h>
 
+#include <openssl/evp.h>
+
 #include <lardon3d/app_state.h>
 #include <lardon3d/feature_store.h>
 #include <lardon3d/matcher.h>
@@ -177,6 +179,128 @@ static bool read_result_asset(const Fixture *fixture, const Lardon3DProjectDbFea
   return read_result == LARDON3D_MATCH_FILE_OK;
 }
 
+static bool files_equal(const char *path_a, const char *path_b) {
+  int fd_a = open(path_a, O_RDONLY | O_CLOEXEC);
+  int fd_b = open(path_b, O_RDONLY | O_CLOEXEC);
+  if (fd_a < 0 || fd_b < 0) {
+    if (fd_a >= 0) (void)close(fd_a);
+    if (fd_b >= 0) (void)close(fd_b);
+    return false;
+  }
+  unsigned char a[4096];
+  unsigned char b[4096];
+  bool equal = true;
+  for (;;) {
+    ssize_t read_a = read(fd_a, a, sizeof(a));
+    ssize_t read_b = read(fd_b, b, sizeof(b));
+    if (read_a < 0 || read_b < 0 || read_a != read_b ||
+        (read_a > 0 && memcmp(a, b, (size_t)read_a) != 0)) {
+      equal = false;
+      break;
+    }
+    if (read_a == 0) break;
+  }
+  if (close(fd_a) != 0 || close(fd_b) != 0) equal = false;
+  return equal;
+}
+
+static bool file_sha256(const char *path, unsigned char output[32]) {
+  int fd = open(path, O_RDONLY | O_CLOEXEC);
+  EVP_MD_CTX *context = EVP_MD_CTX_new();
+  if (fd < 0 || !context) {
+    if (fd >= 0) (void)close(fd);
+    EVP_MD_CTX_free(context);
+    return false;
+  }
+  bool ok = EVP_DigestInit_ex(context, EVP_sha256(), NULL) == 1;
+  unsigned char buffer[4096];
+  while (ok) {
+    ssize_t count = read(fd, buffer, sizeof(buffer));
+    if (count < 0) {
+      ok = false;
+    } else if (count == 0) {
+      break;
+    } else {
+      ok = EVP_DigestUpdate(context, buffer, (size_t)count) == 1;
+    }
+  }
+  unsigned int length = 0;
+  ok = ok && EVP_DigestFinal_ex(context, output, &length) == 1 && length == 32;
+  EVP_MD_CTX_free(context);
+  return close(fd) == 0 && ok;
+}
+
+static uint32_t deterministic_random(uint32_t *state) {
+  *state = *state * 1664525U + 1013904223U;
+  return *state;
+}
+
+#ifdef LARDON3D_MATCHER_E2E_VULKAN
+static bool test_cpu_vulkan_match_file_parity(void) {
+  const uint32_t feature_count = 768;
+  const size_t bytes = (size_t)feature_count * 32;
+  Fixture fixture;
+  CHECK(fixture_open(&fixture));
+  unsigned char *descriptors = malloc(bytes);
+  CHECK(descriptors != NULL);
+  uint32_t random_state = 0x12345678U;
+  for (size_t index = 0; index < bytes; ++index) {
+    descriptors[index] = (unsigned char)(deterministic_random(&random_state) >> 24);
+  }
+  Lardon3DProjectDbFeatureSet set_a, set_b;
+  CHECK(publish_features(&fixture, fixture.image_a.image_id, "orb",
+                         LARDON3D_FEATURE_DESCRIPTOR_U8, descriptors, feature_count, 31,
+                         &set_a));
+  CHECK(publish_features(&fixture, fixture.image_b.image_id, "orb",
+                         LARDON3D_FEATURE_DESCRIPTOR_U8, descriptors, feature_count, 32,
+                         &set_b));
+  free(descriptors);
+
+  char cpu_path[PATH_MAX];
+  char vulkan_path[PATH_MAX];
+  char fallback_path[PATH_MAX];
+  CHECK(join_path(cpu_path, fixture.root, "cpu.match"));
+  CHECK(join_path(vulkan_path, fixture.root, "vulkan.match"));
+  CHECK(join_path(fallback_path, fixture.root, "fallback.match"));
+  Lardon3DMatcherParams params = {LARDON3D_MATCHER_ORB_BF, 0.75F};
+  Lardon3DMatcherStats cpu_stats;
+  Lardon3DMatcherStats vulkan_stats;
+  CHECK(lardon3d_matcher_run(fixture.root, &set_a, &set_b, &params, cpu_path,
+                             &cpu_stats) == LARDON3D_MATCHER_OK);
+  Lardon3DOrbVulkanBackend *backend = lardon3d_orb_vulkan_backend_create();
+  CHECK(backend != NULL);
+  CHECK(lardon3d_matcher_run_with_backend(fixture.root, &set_a, &set_b, &params,
+                                          vulkan_path, backend, &vulkan_stats) ==
+        LARDON3D_MATCHER_OK);
+  CHECK(vulkan_stats.used_vulkan && !vulkan_stats.vulkan_fallback);
+  CHECK(cpu_stats.match_count == vulkan_stats.match_count);
+  CHECK(files_equal(cpu_path, vulkan_path));
+  unsigned char cpu_sha[32];
+  unsigned char vulkan_sha[32];
+  CHECK(file_sha256(cpu_path, cpu_sha));
+  CHECK(file_sha256(vulkan_path, vulkan_sha));
+  CHECK(memcmp(cpu_sha, vulkan_sha, sizeof(cpu_sha)) == 0);
+  lardon3d_orb_vulkan_backend_destroy(backend);
+
+  CHECK(setenv("LARDON3D_VULKAN_DISABLE", "1", 1) == 0);
+  backend = lardon3d_orb_vulkan_backend_create();
+  CHECK(backend != NULL);
+  Lardon3DMatcherStats fallback_stats;
+  CHECK(lardon3d_matcher_run_with_backend(fixture.root, &set_a, &set_b, &params,
+                                          fallback_path, backend, &fallback_stats) ==
+        LARDON3D_MATCHER_OK);
+  CHECK(fallback_stats.vulkan_fallback && !fallback_stats.used_vulkan);
+  CHECK(files_equal(cpu_path, fallback_path));
+  lardon3d_orb_vulkan_backend_destroy(backend);
+  CHECK(unsetenv("LARDON3D_VULKAN_DISABLE") == 0);
+  CHECK(unlink(cpu_path) == 0);
+  CHECK(unlink(vulkan_path) == 0);
+  CHECK(unlink(fallback_path) == 0);
+  CHECK(fixture_close(&fixture));
+  return true;
+}
+#endif
+
 static bool run_kind_e2e(const char *kind, Lardon3DMatcherKind matcher_kind,
                          Lardon3DFeatureDescriptorType type) {
   Fixture fixture;
@@ -377,6 +501,9 @@ static bool run_tests(void) {
                      LARDON3D_FEATURE_DESCRIPTOR_F32));
   CHECK(test_ratio_single_neighbor_and_no_match());
   CHECK(test_kind_and_ownership_rejection());
+#ifdef LARDON3D_MATCHER_E2E_VULKAN
+  CHECK(test_cpu_vulkan_match_file_parity());
+#endif
   return true;
 }
 

tests/test_matcher_task.c

@@ -0,0 +1,482 @@
+#include <dirent.h>
+#include <errno.h>
+#include <limits.h>
+#include <sched.h>
+#include <sqlite3.h>
+#include <stdbool.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <sys/stat.h>
+#include <unistd.h>
+
+#include <lardon3d/feature_store.h>
+#include <lardon3d/matcher_task.h>
+#include <lardon3d/project.h>
+#include <lardon3d/task_queue.h>
+
+#define CHECK(condition)                                                       \
+  do {                                                                         \
+    if (!(condition)) {                                                        \
+      (void)fprintf(stderr, "Échec ligne %d : %s\n", __LINE__, #condition);    \
+      return false;                                                            \
+    }                                                                          \
+  } while (0)
+
+enum {
+  IMAGE_COUNT = 42,
+  PAIR_COUNT = IMAGE_COUNT - 1,
+  PERSISTED_PAIR_COUNT = PAIR_COUNT - 1,
+};
+
+typedef struct {
+  char root[PATH_MAX];
+  Lardon3DAppState state;
+  Lardon3DProjectDbScanSet scanset;
+  Lardon3DProjectDbImage images[IMAGE_COUNT];
+  Lardon3DProjectDbCandidatePair pairs[PAIR_COUNT];
+  unsigned char feature_fingerprint[32];
+} Fixture;
+
+static Lardon3DResourcePolicy interactive_policy(void) {
+  return (Lardon3DResourcePolicy){
+      .system_memory_reserve_bytes = 4ULL * 1024 * 1024 * 1024,
+      .emergency_memory_floor_bytes = 2ULL * 1024 * 1024 * 1024,
+      .system_cpu_reserve = 4,
+      .maximum_cpu_load_ratio = 1.0,
+      .maximum_cpu_pressure_avg10 = 100.0,
+      .maximum_memory_pressure_avg10 = 100.0,
+      .maximum_io_pressure_avg10 = 100.0,
+      .io_slot_capacity = 1,
+      .gpu_slot_capacity = 1,
+  };
+}
+
+static bool join_path(char output[PATH_MAX], const char *left,
+                      const char *right) {
+  int written = snprintf(output, PATH_MAX, "%s/%s", left, right);
+  return written > 0 && (size_t)written < PATH_MAX;
+}
+
+static bool query_integer(const char *path, const char *sql,
+                          sqlite3_int64 expected) {
+  sqlite3 *connection = NULL;
+  sqlite3_stmt *statement = NULL;
+  bool success = sqlite3_open_v2(path, &connection, SQLITE_OPEN_READONLY,
+                                 NULL) == SQLITE_OK &&
+                 sqlite3_prepare_v2(connection, sql, -1, &statement, NULL) ==
+                     SQLITE_OK &&
+                 sqlite3_step(statement) == SQLITE_ROW &&
+                 sqlite3_column_int64(statement, 0) == expected;
+  if (statement) {
+    (void)sqlite3_finalize(statement);
+  }
+  if (connection) {
+    (void)sqlite3_close(connection);
+  }
+  return success;
+}
+
+static bool downgrade_project_to_historical_v10(const char *database_path) {
+  sqlite3 *connection = NULL;
+  if (sqlite3_open_v2(database_path, &connection, SQLITE_OPEN_READWRITE,
+                      NULL) != SQLITE_OK) {
+    if (connection) {
+      (void)sqlite3_close(connection);
+    }
+    return false;
+  }
+  static const char sql[] =
+      "PRAGMA foreign_keys=OFF;BEGIN IMMEDIATE;"
+      "DROP TABLE matcher_tasks;"
+      "UPDATE metadata SET value=10 WHERE key='schema_version';"
+      "COMMIT;PRAGMA foreign_keys=ON;";
+  bool success = sqlite3_exec(connection, sql, NULL, NULL, NULL) == SQLITE_OK;
+  return sqlite3_close(connection) == SQLITE_OK && success;
+}
+
+static bool remove_tree(const char *path) {
+  struct stat information;
+  if (lstat(path, &information) != 0) {
+    return errno == ENOENT;
+  }
+  if (!S_ISDIR(information.st_mode)) {
+    return unlink(path) == 0;
+  }
+  DIR *directory = opendir(path);
+  if (!directory) {
+    return false;
+  }
+  bool success = true;
+  for (struct dirent *entry = readdir(directory); entry;
+       entry = readdir(directory)) {
+    if (strcmp(entry->d_name, ".") == 0 || strcmp(entry->d_name, "..") == 0) {
+      continue;
+    }
+    char child[PATH_MAX];
+    if (!join_path(child, path, entry->d_name) || !remove_tree(child)) {
+      success = false;
+    }
+  }
+  if (closedir(directory) != 0 || rmdir(path) != 0) {
+    success = false;
+  }
+  return success;
+}
+
+static bool create_runtime(Lardon3DAppState *state) {
+  state->hardware_profile = (Lardon3DHardwareProfile){
+      .logical_cpu_count = 16,
+      .page_size_bytes = 4096,
+      .memory_total_bytes = 16ULL * 1024 * 1024 * 1024,
+      .gpu_available = true,
+      .gpu_uses_shared_memory = true,
+      .cpu_architecture = "test",
+  };
+  Lardon3DResourcePolicy policy = interactive_policy();
+  state->resource_governor =
+      lardon3d_resource_governor_create(&state->hardware_profile, &policy);
+  state->orb_vulkan_backend = lardon3d_orb_vulkan_backend_create();
+  state->task_queue =
+      state->resource_governor && state->orb_vulkan_backend
+          ? lardon3d_task_queue_create(state->resource_governor, 16)
+          : NULL;
+  return state->task_queue != NULL;
+}
+
+static bool wait_state(Lardon3DTaskQueue *queue, uint64_t task_id,
+                       Lardon3DTaskState wanted,
+                       Lardon3DTaskSnapshot *snapshot) {
+  for (size_t attempt = 0; attempt < 2000000; ++attempt) {
+    if (lardon3d_task_queue_get(queue, task_id, snapshot) &&
+        snapshot->state == wanted) {
+      return true;
+    }
+    sched_yield();
+  }
+  return false;
+}
+
+static bool wait_durable_state(Lardon3DProjectDb *database, uint64_t task_id,
+                               Lardon3DTaskState wanted,
+                               Lardon3DProjectDbTask *durable_task) {
+  for (size_t attempt = 0; attempt < 2000000; ++attempt) {
+    if (lardon3d_project_db_load_task(database, task_id, durable_task) ==
+            LARDON3D_PROJECT_DB_OK &&
+        durable_task->saved_state == wanted) {
+      return true;
+    }
+    sched_yield();
+  }
+  return false;
+}
+
+static void image_asset_path(const unsigned char hash[32], char path[4096]) {
+  static const char digits[] = "0123456789abcdef";
+  char hex[65];
+  for (size_t index = 0; index < 32; ++index) {
+    hex[2 * index] = digits[hash[index] >> 4];
+    hex[2 * index + 1] = digits[hash[index] & 15U];
+  }
+  hex[64] = '\0';
+  (void)snprintf(path, 4096, "assets/images/%c%c/%s", hex[0], hex[1], hex);
+}
+
+static bool register_image(Fixture *fixture, unsigned char seed, size_t index) {
+  unsigned char hash[32];
+  memset(hash, seed, sizeof(hash));
+  char path[4096];
+  image_asset_path(hash, path);
+  Lardon3DProjectDbImageRegisterStatus status;
+  return lardon3d_project_db_register_image(
+             fixture->state.project_db, fixture->scanset.scanset_id, hash, path,
+             1, "fixture.bin", "/synthetic/fixture.bin", 0, seed, &status,
+             &fixture->images[index]) == LARDON3D_PROJECT_DB_OK;
+}
+
+static bool publish_features(Fixture *fixture, size_t image_index) {
+  unsigned char descriptor[32];
+  memset(descriptor, (int)image_index, sizeof(descriptor));
+  Lardon3DFeatureKeypoint keypoint = {
+      .size = 1.0F,
+  };
+  Lardon3DExtractedFeatures features = {
+      .image_width = 64,
+      .image_height = 64,
+      .feature_count = 1,
+      .keypoints = &keypoint,
+      .descriptors = descriptor,
+      .descriptor_bytes = sizeof(descriptor),
+  };
+  Lardon3DProjectDbFeatureSet feature_set;
+  return lardon3d_feature_store_publish_v2(
+             &fixture->state, fixture->images[image_index].image_id, 0, "orb",
+             1, fixture->feature_fingerprint, LARDON3D_FEATURE_DESCRIPTOR_U8,
+             32, 0, &features, &feature_set) == LARDON3D_FEATURE_STORE_OK;
+}
+
+static bool fixture_create(Fixture *fixture) {
+  memset(fixture, 0, sizeof(*fixture));
+  char root[] = "/tmp/lardon3d-matcher-task-XXXXXX";
+  char *created = mkdtemp(root);
+  if (!created ||
+      snprintf(fixture->root, sizeof(fixture->root), "%s", created) <= 0 ||
+      setenv("LARDON3D_PROJECTS_ROOT", fixture->root, 1) != 0) {
+    return false;
+  }
+  memset(fixture->feature_fingerprint, 0x5A,
+         sizeof(fixture->feature_fingerprint));
+  lardon3d_app_state_init(&fixture->state);
+  if (!create_runtime(&fixture->state) ||
+      !lardon3d_project_create(&fixture->state, "MatcherTask") ||
+      lardon3d_project_db_create_scanset(fixture->state.project_db,
+                                         "matcher-task", &fixture->scanset) !=
+          LARDON3D_PROJECT_DB_OK) {
+    return false;
+  }
+  for (size_t index = 0; index < IMAGE_COUNT; ++index) {
+    if (!register_image(fixture, (unsigned char)(index + 1), index) ||
+        !publish_features(fixture, index)) {
+      return false;
+    }
+  }
+  for (size_t index = 0; index < PAIR_COUNT; ++index) {
+    if (lardon3d_project_db_create_candidate_pair(
+            fixture->state.project_db, fixture->images[index].image_id,
+            fixture->images[index + 1].image_id, (int64_t)index,
+            &fixture->pairs[index]) != LARDON3D_PROJECT_DB_OK) {
+      return false;
+    }
+  }
+  char database_path[PATH_MAX];
+  if (!join_path(database_path, fixture->state.project_path, "project.db")) {
+    return false;
+  }
+  sqlite3 *connection = NULL;
+  if (sqlite3_open_v2(database_path, &connection, SQLITE_OPEN_READWRITE,
+                      NULL) != SQLITE_OK) {
+    if (connection) {
+      sqlite3_close(connection);
+    }
+    return false;
+  }
+  char sql[128];
+  int written =
+      snprintf(sql, sizeof(sql),
+               "DELETE FROM candidate_pairs WHERE candidate_pair_id=%lu",
+               (unsigned long)fixture->pairs[2].candidate_pair_id);
+  bool deleted = written > 0 && (size_t)written < sizeof(sql) &&
+                 sqlite3_exec(connection, sql, NULL, NULL, NULL) == SQLITE_OK &&
+                 sqlite3_changes(connection) == 1;
+  (void)sqlite3_close(connection);
+  if (!deleted) {
+    return false;
+  }
+  return true;
+}
+
+static void stop_runtime(Fixture *fixture) {
+  if (fixture->state.task_queue) {
+    lardon3d_task_queue_destroy(fixture->state.task_queue);
+    fixture->state.task_queue = NULL;
+  }
+  lardon3d_project_close(&fixture->state);
+  if (fixture->state.resource_governor) {
+    lardon3d_resource_governor_destroy(fixture->state.resource_governor);
+    fixture->state.resource_governor = NULL;
+  }
+  lardon3d_orb_vulkan_backend_destroy(fixture->state.orb_vulkan_backend);
+  fixture->state.orb_vulkan_backend = NULL;
+}
+
+static bool reopen_runtime(Fixture *fixture) {
+  lardon3d_app_state_init(&fixture->state);
+  return create_runtime(&fixture->state) &&
+         lardon3d_project_open(&fixture->state, "MatcherTask");
+}
+
+static Lardon3DMatcherTaskConfiguration configuration(const Fixture *fixture) {
+  Lardon3DMatcherTaskConfiguration result = {
+      .feature_extractor_version = 1,
+      .matcher =
+          {
+              .kind = LARDON3D_MATCHER_ORB_BF,
+              .ratio_threshold = 0.75F,
+          },
+  };
+  (void)snprintf(result.feature_extractor_kind,
+                 sizeof(result.feature_extractor_kind), "orb");
+  memcpy(result.feature_parameter_fingerprint, fixture->feature_fingerprint,
+         sizeof(result.feature_parameter_fingerprint));
+  return result;
+}
+
+static bool count_results(Fixture *fixture, size_t *count) {
+  Lardon3DProjectDbMatchResult results[16];
+  uint64_t cursor = 0;
+  *count = 0;
+  for (;;) {
+    size_t page_count = 0;
+    if (lardon3d_project_db_list_match_results(
+            fixture->state.project_db, cursor, results, 16, &page_count) !=
+        LARDON3D_PROJECT_DB_OK) {
+      return false;
+    }
+    *count += page_count;
+    if (page_count < 16) {
+      return true;
+    }
+    cursor = results[page_count - 1].match_result_id;
+  }
+}
+
+static bool run_test(void) {
+  Fixture fixture;
+  CHECK(fixture_create(&fixture));
+  char database_path[PATH_MAX];
+  CHECK(join_path(database_path, fixture.state.project_path, "project.db"));
+  stop_runtime(&fixture);
+  CHECK(downgrade_project_to_historical_v10(database_path));
+  CHECK(query_integer(database_path,
+                      "SELECT value FROM metadata WHERE key='schema_version'",
+                      10));
+  CHECK(query_integer(database_path,
+                      "SELECT count(*) FROM sqlite_master WHERE type='table' AND "
+                      "name='match_results'",
+                      1));
+  CHECK(query_integer(database_path,
+                      "SELECT count(*) FROM sqlite_master WHERE type='table' AND "
+                      "name='matcher_tasks'",
+                      0));
+  CHECK(reopen_runtime(&fixture));
+  CHECK(lardon3d_project_db_schema_version(fixture.state.project_db) == 11);
+  CHECK(query_integer(database_path,
+                      "SELECT count(*) FROM sqlite_master WHERE type='table' AND "
+                      "name='matcher_tasks'",
+                      1));
+  const Lardon3DTaskKindRegistry *registry =
+      lardon3d_task_kind_registry_production();
+  const Lardon3DTaskKindDescriptor *descriptor = NULL;
+  CHECK(lardon3d_task_kind_registry_lookup(registry, LARDON3D_MATCHER_TASK_KIND,
+                                           LARDON3D_MATCHER_TASK_KIND_VERSION,
+                                           &descriptor) ==
+            LARDON3D_TASK_KIND_OK &&
+        descriptor != NULL);
+
+  Lardon3DMatcherTaskConfiguration settings = configuration(&fixture);
+  Lardon3DMatcherTaskConfiguration invalid = settings;
+  invalid.matcher.kind = LARDON3D_MATCHER_SIFT_BF;
+  uint64_t invalid_task_id = 0;
+  CHECK(!lardon3d_project_create_matcher_task(&fixture.state, &invalid,
+                                              &invalid_task_id));
+  CHECK(invalid_task_id == 0);
+  CHECK(setenv("LARDON3D_TEST_MATCHER_PAUSE_AFTER_PUBLICATION", "1", 1) == 0);
+  CHECK(setenv("LARDON3D_TEST_MATCHER_SKIP_FINISHED_CHECKPOINT", "1", 1) == 0);
+  uint64_t task_id = 0;
+  CHECK(lardon3d_project_enqueue_matcher_task(&fixture.state, &settings,
+                                              &task_id));
+  Lardon3DTaskSnapshot snapshot;
+  CHECK(wait_state(fixture.state.task_queue, task_id, TASK_PAUSED, &snapshot));
+
+  size_t result_count = 0;
+  CHECK(count_results(&fixture, &result_count) && result_count == 1);
+  Lardon3DProjectDbMatcherTask saved;
+  CHECK(lardon3d_project_db_load_matcher_task(fixture.state.project_db, task_id,
+                                              &saved) ==
+        LARDON3D_PROJECT_DB_OK);
+  CHECK(saved.after_candidate_pair_id == 0);
+  CHECK(saved.matcher_kind == LARDON3D_MATCHER_ORB_BF);
+  CHECK(saved.ratio_threshold == 0.75F);
+
+  stop_runtime(&fixture);
+  CHECK(unsetenv("LARDON3D_TEST_MATCHER_PAUSE_AFTER_PUBLICATION") == 0);
+  CHECK(unsetenv("LARDON3D_TEST_MATCHER_SKIP_FINISHED_CHECKPOINT") == 0);
+  CHECK(reopen_runtime(&fixture));
+  Lardon3DProjectRecoverySummary recovery;
+  CHECK(lardon3d_project_last_recovery_summary(&fixture.state, &recovery));
+  CHECK(recovery.resumed == 1);
+  CHECK(
+      wait_state(fixture.state.task_queue, task_id, TASK_COMPLETED, &snapshot));
+  CHECK(snapshot.progress == 100);
+  Lardon3DProjectDbTask durable_task;
+  CHECK(wait_durable_state(fixture.state.project_db, task_id, TASK_COMPLETED,
+                           &durable_task));
+  CHECK(durable_task.sequence_count >= 1);
+  CHECK(count_results(&fixture, &result_count) &&
+        result_count == PERSISTED_PAIR_COUNT);
+  CHECK(lardon3d_project_db_load_matcher_task(fixture.state.project_db, task_id,
+                                              &saved) ==
+        LARDON3D_PROJECT_DB_OK);
+  CHECK(saved.after_candidate_pair_id ==
+        fixture.pairs[PAIR_COUNT - 1].candidate_pair_id);
+
+  CHECK(setenv("LARDON3D_TEST_MATCHER_PAUSE_AFTER_PUBLICATION", "1", 1) == 0);
+  uint64_t cancelled_id = 0;
+  CHECK(lardon3d_project_enqueue_matcher_task(&fixture.state, &settings,
+                                              &cancelled_id));
+  CHECK(wait_state(fixture.state.task_queue, cancelled_id, TASK_PAUSED,
+                   &snapshot));
+  CHECK(lardon3d_task_queue_cancel(fixture.state.task_queue, cancelled_id));
+  CHECK(wait_state(fixture.state.task_queue, cancelled_id, TASK_CANCELLED,
+                   &snapshot));
+  CHECK(unsetenv("LARDON3D_TEST_MATCHER_PAUSE_AFTER_PUBLICATION") == 0);
+
+  CHECK(setenv("LARDON3D_TEST_MATCHER_PAUSE_AFTER_PUBLICATION", "1", 1) == 0);
+  uint64_t resumed_id = 0;
+  CHECK(lardon3d_project_enqueue_matcher_task(&fixture.state, &settings,
+                                              &resumed_id));
+  CHECK(
+      wait_state(fixture.state.task_queue, resumed_id, TASK_PAUSED, &snapshot));
+  CHECK(unsetenv("LARDON3D_TEST_MATCHER_PAUSE_AFTER_PUBLICATION") == 0);
+  CHECK(lardon3d_task_queue_resume(fixture.state.task_queue, resumed_id));
+  CHECK(wait_state(fixture.state.task_queue, resumed_id, TASK_COMPLETED,
+                   &snapshot));
+
+  CHECK(setenv("LARDON3D_TEST_MATCHER_PAUSE_AFTER_BATCH", "1", 1) == 0);
+  uint64_t pressure_id = 0;
+  CHECK(lardon3d_project_enqueue_matcher_task(&fixture.state, &settings,
+                                              &pressure_id));
+  CHECK(wait_state(fixture.state.task_queue, pressure_id, TASK_PAUSED,
+                   &snapshot));
+  Lardon3DResourcePolicy pressure_policy = interactive_policy();
+  pressure_policy.system_memory_reserve_bytes =
+      fixture.state.hardware_profile.memory_total_bytes - 16ULL * 1024 * 1024;
+  pressure_policy.emergency_memory_floor_bytes =
+      pressure_policy.system_memory_reserve_bytes;
+  CHECK(lardon3d_resource_governor_set_policy(fixture.state.resource_governor,
+                                              &pressure_policy));
+  CHECK(unsetenv("LARDON3D_TEST_MATCHER_PAUSE_AFTER_BATCH") == 0);
+  CHECK(lardon3d_task_queue_resume(fixture.state.task_queue, pressure_id));
+  for (size_t attempt = 0; attempt < 2000000; ++attempt) {
+    if (lardon3d_resource_governor_pressure(fixture.state.resource_governor) ==
+        LARDON3D_RESOURCE_PRESSURE_RED) {
+      break;
+    }
+    sched_yield();
+  }
+  CHECK(lardon3d_resource_governor_pressure(fixture.state.resource_governor) ==
+        LARDON3D_RESOURCE_PRESSURE_RED);
+  Lardon3DResourcePolicy normal_policy = interactive_policy();
+  CHECK(lardon3d_resource_governor_set_policy(fixture.state.resource_governor,
+                                              &normal_policy));
+  CHECK(wait_state(fixture.state.task_queue, pressure_id, TASK_COMPLETED,
+                   &snapshot));
+
+  uint64_t estimate_task_id = 0;
+  Lardon3DTask *estimate_task = lardon3d_project_create_matcher_task(
+      &fixture.state, &settings, &estimate_task_id);
+  Lardon3DTaskDurableSnapshot estimate_snapshot;
+  CHECK(estimate_task != NULL && estimate_task_id != 0);
+  CHECK(lardon3d_task_durable_snapshot(estimate_task, &estimate_snapshot));
+  CHECK(estimate_snapshot.estimate.gpu_memory_fixed_bytes ==
+        LARDON3D_ORB_VULKAN_PERMANENT_BUFFER_BYTES);
+  CHECK(estimate_snapshot.estimate.desired_gpu_slots == 1);
+  lardon3d_task_destroy(estimate_task);
+
+  stop_runtime(&fixture);
+  CHECK(remove_tree(fixture.root));
+  return true;
+}
+
+int main(void) { return run_test() ? EXIT_SUCCESS : EXIT_FAILURE; }

tests/test_orb_vulkan_backend.cpp

@@ -0,0 +1,226 @@
+#include <lardon3d/orb_vulkan_backend.h>
+
+#include <opencv2/core.hpp>
+#include <opencv2/features2d.hpp>
+
+#include <algorithm>
+#include <cstdint>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <thread>
+#include <vector>
+
+namespace {
+
+constexpr uint32_t kDescriptorBytes = 32;
+
+static uint32_t next_random(uint32_t *state) {
+  *state = *state * 1664525U + 1013904223U;
+  return *state;
+}
+
+static std::vector<unsigned char> make_descriptors(uint32_t count, uint32_t seed) {
+  std::vector<unsigned char> descriptors(static_cast<size_t>(count) * kDescriptorBytes);
+  for (unsigned char &value : descriptors) {
+    value = static_cast<unsigned char>(next_random(&seed) >> 24);
+  }
+  return descriptors;
+}
+
+static std::vector<Lardon3DOrbTop2> opencv_top2(
+    const std::vector<unsigned char> &a, uint32_t count_a,
+    const std::vector<unsigned char> &b, uint32_t count_b) {
+  std::vector<Lardon3DOrbTop2> output(count_a);
+  if (count_a == 0 || count_b == 0) {
+    return output;
+  }
+  cv::Mat matrix_a(static_cast<int>(count_a), kDescriptorBytes, CV_8U,
+                   const_cast<unsigned char *>(a.data()));
+  cv::Mat matrix_b(static_cast<int>(count_b), kDescriptorBytes, CV_8U,
+                   const_cast<unsigned char *>(b.data()));
+  cv::BFMatcher matcher(cv::NORM_HAMMING, false);
+  std::vector<std::vector<cv::DMatch>> matches;
+  matcher.knnMatch(matrix_a, matrix_b, matches, 2);
+  for (uint32_t query = 0; query < count_a; ++query) {
+    const auto &knn = matches[query];
+    output[query].neighbor_count = static_cast<uint32_t>(knn.size());
+    if (!knn.empty()) {
+      output[query].best_index = static_cast<uint32_t>(knn[0].trainIdx);
+      output[query].best_distance = static_cast<uint32_t>(knn[0].distance);
+    }
+    if (knn.size() >= 2) {
+      output[query].second_index = static_cast<uint32_t>(knn[1].trainIdx);
+      output[query].second_distance = static_cast<uint32_t>(knn[1].distance);
+    }
+  }
+  return output;
+}
+
+static bool equal_top2(const Lardon3DOrbTop2 &a, const Lardon3DOrbTop2 &b) {
+  return a.neighbor_count == b.neighbor_count && a.best_index == b.best_index &&
+         a.best_distance == b.best_distance && a.second_index == b.second_index &&
+         a.second_distance == b.second_distance;
+}
+
+static bool check_case(Lardon3DOrbVulkanBackend *backend, uint32_t count_a,
+                       uint32_t count_b, uint32_t seed) {
+  std::vector<unsigned char> a = make_descriptors(count_a, seed);
+  std::vector<unsigned char> b = make_descriptors(count_b, seed ^ 0xa5a5a5a5U);
+  std::vector<Lardon3DOrbTop2> expected = opencv_top2(a, count_a, b, count_b);
+  std::vector<Lardon3DOrbTop2> actual(count_a);
+  Lardon3DOrbVulkanResult result = lardon3d_orb_vulkan_top2(
+      backend, a.data(), count_a, b.data(), count_b, actual.data(), actual.size());
+  if (result != LARDON3D_ORB_VULKAN_OK) {
+    std::fprintf(stderr, "Vulkan top-2 failed for %u x %u: %d\n", count_a,
+                 count_b, static_cast<int>(result));
+    return false;
+  }
+  for (uint32_t query = 0; query < count_a; ++query) {
+    if (!equal_top2(expected[query], actual[query])) {
+      std::fprintf(stderr,
+                   "top-2 mismatch at %u for %u x %u: "
+                   "CPU=(%u,%u,%u,%u,%u) Vulkan=(%u,%u,%u,%u,%u)\n",
+                   query, count_a, count_b, expected[query].neighbor_count,
+                   expected[query].best_index, expected[query].best_distance,
+                   expected[query].second_index, expected[query].second_distance,
+                   actual[query].neighbor_count, actual[query].best_index,
+                   actual[query].best_distance, actual[query].second_index,
+                   actual[query].second_distance);
+      return false;
+    }
+  }
+  return true;
+}
+
+static bool check_ties(Lardon3DOrbVulkanBackend *backend) {
+  constexpr uint32_t count_a = 4;
+  constexpr uint32_t count_b = 16;
+  std::vector<unsigned char> a(count_a * kDescriptorBytes, 0);
+  std::vector<unsigned char> b(count_b * kDescriptorBytes, 0);
+  std::memset(b.data() + 7 * kDescriptorBytes, 0xff, kDescriptorBytes);
+  std::memset(a.data() + 3 * kDescriptorBytes, 0xff, kDescriptorBytes);
+  std::vector<Lardon3DOrbTop2> expected = opencv_top2(a, count_a, b, count_b);
+  std::vector<Lardon3DOrbTop2> actual(count_a);
+  if (lardon3d_orb_vulkan_top2(backend, a.data(), count_a, b.data(), count_b,
+                               actual.data(), actual.size()) !=
+      LARDON3D_ORB_VULKAN_OK) {
+    return false;
+  }
+  return std::equal(expected.begin(), expected.end(), actual.begin(), equal_top2);
+}
+
+static bool check_serialized_threads(Lardon3DOrbVulkanBackend *backend) {
+  std::vector<unsigned char> a = make_descriptors(1024, 0x11111111U);
+  std::vector<unsigned char> b = make_descriptors(1024, 0x22222222U);
+  std::vector<Lardon3DOrbTop2> expected_a = opencv_top2(a, 1024, b, 1024);
+  std::vector<Lardon3DOrbTop2> expected_b = opencv_top2(b, 1024, a, 1024);
+  std::vector<Lardon3DOrbTop2> actual_a(1024);
+  std::vector<Lardon3DOrbTop2> actual_b(1024);
+  Lardon3DOrbVulkanResult result_a = LARDON3D_ORB_VULKAN_FAILED;
+  Lardon3DOrbVulkanResult result_b = LARDON3D_ORB_VULKAN_FAILED;
+  std::thread thread_a([&] {
+    result_a = lardon3d_orb_vulkan_top2(backend, a.data(), 1024, b.data(), 1024,
+                                        actual_a.data(), actual_a.size());
+  });
+  std::thread thread_b([&] {
+    result_b = lardon3d_orb_vulkan_top2(backend, b.data(), 1024, a.data(), 1024,
+                                        actual_b.data(), actual_b.size());
+  });
+  thread_a.join();
+  thread_b.join();
+  return result_a == LARDON3D_ORB_VULKAN_OK &&
+         result_b == LARDON3D_ORB_VULKAN_OK &&
+         std::equal(expected_a.begin(), expected_a.end(), actual_a.begin(), equal_top2) &&
+         std::equal(expected_b.begin(), expected_b.end(), actual_b.begin(), equal_top2);
+}
+
+static bool check_cached_unavailable() {
+  if (setenv("LARDON3D_VULKAN_DISABLE", "1", 1) != 0) {
+    return false;
+  }
+  Lardon3DOrbVulkanBackend *backend = lardon3d_orb_vulkan_backend_create();
+  std::vector<unsigned char> descriptors = make_descriptors(1024, 0x10101010U);
+  std::vector<Lardon3DOrbTop2> output(1024);
+  Lardon3DOrbVulkanResult first = lardon3d_orb_vulkan_top2(
+      backend, descriptors.data(), 1024, descriptors.data(), 1024, output.data(),
+      output.size());
+  unsetenv("LARDON3D_VULKAN_DISABLE");
+  Lardon3DOrbVulkanResult second = lardon3d_orb_vulkan_top2(
+      backend, descriptors.data(), 1024, descriptors.data(), 1024, output.data(),
+      output.size());
+  lardon3d_orb_vulkan_backend_destroy(backend);
+  return first == LARDON3D_ORB_VULKAN_UNAVAILABLE &&
+         second == LARDON3D_ORB_VULKAN_UNAVAILABLE;
+}
+
+static bool check_invalid_inputs(Lardon3DOrbVulkanBackend *backend) {
+  unsigned char descriptor[kDescriptorBytes]{};
+  Lardon3DOrbTop2 output{};
+  return lardon3d_orb_vulkan_top2(nullptr, descriptor, 1, descriptor, 1, &output, 1) ==
+             LARDON3D_ORB_VULKAN_INVALID_ARGUMENT &&
+         lardon3d_orb_vulkan_top2(backend, descriptor, 8193, descriptor, 1, &output,
+                                  1) == LARDON3D_ORB_VULKAN_INVALID_ARGUMENT &&
+         lardon3d_orb_vulkan_top2(backend, descriptor, 1, descriptor, 1, &output,
+                                  0) == LARDON3D_ORB_VULKAN_INVALID_ARGUMENT;
+}
+
+#ifdef LARDON3D_ORB_VULKAN_TESTING
+static bool check_cached_device_failure() {
+  Lardon3DOrbVulkanBackend *backend = lardon3d_orb_vulkan_backend_create();
+  std::vector<unsigned char> descriptors = make_descriptors(1024, 0x20202020U);
+  std::vector<Lardon3DOrbTop2> output(1024);
+  if (setenv("LARDON3D_TEST_VULKAN_DEVICE_LOST", "1", 1) != 0) {
+    return false;
+  }
+  Lardon3DOrbVulkanResult first = lardon3d_orb_vulkan_top2(
+      backend, descriptors.data(), 1024, descriptors.data(), 1024, output.data(),
+      output.size());
+  unsetenv("LARDON3D_TEST_VULKAN_DEVICE_LOST");
+  Lardon3DOrbVulkanResult second = lardon3d_orb_vulkan_top2(
+      backend, descriptors.data(), 1024, descriptors.data(), 1024, output.data(),
+      output.size());
+  lardon3d_orb_vulkan_backend_destroy(backend);
+  return first == LARDON3D_ORB_VULKAN_FAILED &&
+         second == LARDON3D_ORB_VULKAN_UNAVAILABLE;
+}
+#endif
+
+}  // namespace
+
+int main() {
+  Lardon3DOrbVulkanBackend *backend = lardon3d_orb_vulkan_backend_create();
+  if (!backend) {
+    return 1;
+  }
+  const uint32_t sizes[] = {0, 1, 2, 16, 64, 256, 1024, 4096, 8192};
+  bool ok = true;
+  for (uint32_t size : sizes) {
+    ok = check_case(backend, size, size, 0x12345678U + size) && ok;
+  }
+  ok = check_case(backend, 64, 1, 0x99112233U) && ok;
+  ok = check_case(backend, 64, 2, 0x88112233U) && ok;
+  ok = check_ties(backend) && ok;
+  ok = check_invalid_inputs(backend) && ok;
+  ok = check_serialized_threads(backend) && ok;
+  ok = !lardon3d_orb_vulkan_should_use(256, 256) && ok;
+  ok = !lardon3d_orb_vulkan_should_use(512, 512) && ok;
+  ok = lardon3d_orb_vulkan_should_use(768, 768) && ok;
+  ok = lardon3d_orb_vulkan_should_use(1024, 1024) && ok;
+
+  Lardon3DOrbVulkanInfo info{};
+  ok = lardon3d_orb_vulkan_backend_info(backend, &info) && info.available && ok;
+  if (info.available) {
+    std::printf("device=%s workgroup=%u payload=%llu init_ms=%.3f gpu_ms=%.3f\n",
+                info.device_name, info.workgroup_size,
+                static_cast<unsigned long long>(info.permanent_payload_bytes),
+                static_cast<double>(info.initialization_ns) / 1.0e6,
+                static_cast<double>(info.gpu_ns) / 1.0e6);
+  }
+  lardon3d_orb_vulkan_backend_destroy(backend);
+  ok = check_cached_unavailable() && ok;
+#ifdef LARDON3D_ORB_VULKAN_TESTING
+  ok = check_cached_device_failure() && ok;
+#endif
+  return ok ? 0 : 1;
+}

tests/test_precision_features.cpp

@@ -166,6 +166,10 @@ static cv::Mat affine_homography(const cv::Mat &affine) {
 }
 
 int main() {
+  unsigned int original_threads = lardon3d_feature_opencv_thread_count();
+  CHECK(lardon3d_feature_opencv_configure_threads(2));
+  CHECK(lardon3d_feature_opencv_thread_count() == 2);
+  CHECK(lardon3d_feature_opencv_configure_threads(original_threads));
   Lardon3DSiftExtractorParameters valid_parameters = lardon3d_sift_precision_classic_v1(false);
   CHECK(lardon3d_sift_extractor_parameters_valid(&valid_parameters));
   Lardon3DSiftExtractorParameters invalid_parameters = valid_parameters;

tests/test_project_db.c

@@ -89,7 +89,7 @@ static bool create_future_database(const char *path) {
   }
   bool ok = sqlite3_exec(connection,
                          "CREATE TABLE metadata(key TEXT PRIMARY KEY,value INTEGER NOT NULL);"
-                         "INSERT INTO metadata VALUES('schema_version',11);",
+                         "INSERT INTO metadata VALUES('schema_version',12);",
                          NULL, NULL, NULL) == SQLITE_OK;
   return sqlite3_close(connection) == SQLITE_OK && ok;
 }
@@ -103,6 +103,7 @@ static bool create_v7_database(const char *path) {
   if (sqlite3_open(path, &connection) != SQLITE_OK) return false;
   static const char sql[] =
       "PRAGMA foreign_keys=OFF;BEGIN IMMEDIATE;"
+      "DROP TABLE matcher_tasks;"
       "DROP TABLE match_results;"
       "DROP TABLE candidate_pair_generate_tasks;"
       "DROP TABLE candidate_pairs;"
@@ -120,6 +121,7 @@ static bool create_v6_database(const char *path) {
   if (sqlite3_open(path, &connection) != SQLITE_OK) return false;
   static const char sql[] =
       "PRAGMA foreign_keys=OFF;BEGIN IMMEDIATE;"
+      "DROP TABLE matcher_tasks;"
       "DROP TABLE match_results;"
       "DROP TABLE candidate_pair_generate_tasks;"
       "DROP TABLE candidate_pairs;"
@@ -134,6 +136,26 @@ static bool create_v6_database(const char *path) {
   return sqlite3_close(connection) == SQLITE_OK && ok;
 }
 
+static bool create_v10_database(const char *path) {
+  Lardon3DProjectDb *database = NULL;
+  char error[LARDON3D_PROJECT_DB_ERROR_CAPACITY];
+  if (lardon3d_project_db_open(path, &database, error) != LARDON3D_PROJECT_DB_OK) {
+    return false;
+  }
+  lardon3d_project_db_close(database);
+  sqlite3 *connection = NULL;
+  if (sqlite3_open(path, &connection) != SQLITE_OK) {
+    return false;
+  }
+  static const char sql[] =
+      "PRAGMA foreign_keys=OFF;BEGIN IMMEDIATE;"
+      "DROP TABLE matcher_tasks;"
+      "UPDATE metadata SET value=10 WHERE key='schema_version';"
+      "COMMIT;PRAGMA foreign_keys=ON;";
+  bool ok = sqlite3_exec(connection, sql, NULL, NULL, NULL) == SQLITE_OK;
+  return sqlite3_close(connection) == SQLITE_OK && ok;
+}
+
 static bool create_v5_database(const char *path) {
   if (!create_v6_database(path)) return false;
   sqlite3 *connection = NULL;
@@ -282,6 +304,7 @@ static bool run_test(void) {
   char failed_v3_migration_path[512], v3_path[512], failed_v4_path[512];
   char v4_path[512], failed_v5_path[512], failed_v6_path[512], failed_v7_path[512];
   char direct_v5_path[512], v8_path[512], failed_v8_path[512];
+  char v10_path[512], failed_v11_path[512];
   CHECK(snprintf(database_path, sizeof(database_path), "%s/project.db", directory) > 0);
   CHECK(snprintf(artifact_path, sizeof(artifact_path), "%s/artifact.bin", directory) > 0);
   CHECK(snprintf(future_path, sizeof(future_path), "%s/future.db", directory) > 0);
@@ -306,11 +329,13 @@ static bool run_test(void) {
   CHECK(snprintf(v8_path, sizeof(v8_path), "%s/v8.db", directory) > 0);
   CHECK(snprintf(failed_v8_path, sizeof(failed_v8_path), "%s/failed-v8-migration.db", directory) >
         0);
+  CHECK(snprintf(v10_path, sizeof(v10_path), "%s/v10.db", directory) > 0);
+  CHECK(snprintf(failed_v11_path, sizeof(failed_v11_path), "%s/failed-v11.db", directory) > 0);
 
   char error[LARDON3D_PROJECT_DB_ERROR_CAPACITY];
   Lardon3DProjectDb *database = NULL;
   CHECK(lardon3d_project_db_open(database_path, &database, error) == LARDON3D_PROJECT_DB_OK);
-  CHECK(database && lardon3d_project_db_schema_version(database) == 10);
+  CHECK(database && lardon3d_project_db_schema_version(database) == 11);
   bool legacy_pending = true;
   CHECK(lardon3d_project_db_legacy_catalog_pending(database, &legacy_pending) ==
             LARDON3D_PROJECT_DB_OK &&
@@ -577,7 +602,7 @@ static bool run_test(void) {
         LARDON3D_PROJECT_DB_INVALID_ARGUMENT);
   lardon3d_project_db_close(contexts[0].database);
   database = NULL;
-  CHECK(query_integer(database_path, "SELECT value FROM metadata WHERE key='schema_version'", 10));
+  CHECK(query_integer(database_path, "SELECT value FROM metadata WHERE key='schema_version'", 11));
   CHECK(query_integer(database_path, "SELECT count(*) FROM tasks WHERE task_id=1", 1));
   CHECK(lardon3d_project_db_open(database_path, &database, error) == LARDON3D_PROJECT_DB_OK);
   CHECK(lardon3d_project_db_load_task(database, 1, &task) == LARDON3D_PROJECT_DB_OK);
@@ -599,7 +624,7 @@ static bool run_test(void) {
 
   CHECK(create_v1_database(legacy_path));
   CHECK(lardon3d_project_db_open(legacy_path, &database, error) == LARDON3D_PROJECT_DB_OK);
-  CHECK(lardon3d_project_db_schema_version(database) == 10);
+  CHECK(lardon3d_project_db_schema_version(database) == 11);
   CHECK(lardon3d_project_db_get_project(database, &loaded_project) == LARDON3D_PROJECT_DB_OK &&
         strcmp(loaded_project.stable_id, "legacy-project") == 0);
   CHECK(lardon3d_project_db_load_task(database, 9, &task) == LARDON3D_PROJECT_DB_OK);
@@ -609,7 +634,7 @@ static bool run_test(void) {
         LARDON3D_PROJECT_DB_OK);
   lardon3d_project_db_close(database);
   database = NULL;
-  CHECK(query_integer(legacy_path, "SELECT value FROM metadata WHERE key='schema_version'", 10));
+  CHECK(query_integer(legacy_path, "SELECT value FROM metadata WHERE key='schema_version'", 11));
 
   CHECK(create_v1_database(failed_migration_path));
   CHECK(setenv("LARDON3D_TEST_PROJECT_DB_FAIL_MIGRATION_V2", "1", 1) == 0);
@@ -634,7 +659,7 @@ static bool run_test(void) {
         LARDON3D_PROJECT_DB_OK);
   lardon3d_project_db_close(database);
   database = NULL;
-  CHECK(query_integer(v2_path, "SELECT value FROM metadata WHERE key='schema_version'", 10));
+  CHECK(query_integer(v2_path, "SELECT value FROM metadata WHERE key='schema_version'", 11));
 
   CHECK(create_v2_database(failed_v3_migration_path));
   CHECK(setenv("LARDON3D_TEST_PROJECT_DB_FAIL_MIGRATION_V3", "1", 1) == 0);
@@ -662,7 +687,7 @@ static bool run_test(void) {
         LARDON3D_PROJECT_DB_OK);
   lardon3d_project_db_close(database);
   database = NULL;
-  CHECK(query_integer(v3_path, "SELECT value FROM metadata WHERE key='schema_version'", 10));
+  CHECK(query_integer(v3_path, "SELECT value FROM metadata WHERE key='schema_version'", 11));
 
   CHECK(create_v3_database(failed_v4_path));
   CHECK(setenv("LARDON3D_TEST_PROJECT_DB_FAIL_MIGRATION_V4", "1", 1) == 0);
@@ -679,7 +704,7 @@ static bool run_test(void) {
     fprintf(stderr, "Migration v4 (%d): %s\n", (int)v4_result, error);
   }
   CHECK(v4_result == LARDON3D_PROJECT_DB_OK);
-  CHECK(lardon3d_project_db_schema_version(database) == 10);
+  CHECK(lardon3d_project_db_schema_version(database) == 11);
   CHECK(lardon3d_project_db_load_task(database, 9, &task) == LARDON3D_PROJECT_DB_OK);
   CHECK(lardon3d_project_db_load_artifact(database, "legacy-artifact", &loaded_artifact) ==
         LARDON3D_PROJECT_DB_OK);
@@ -712,25 +737,30 @@ static bool run_test(void) {
   CHECK(query_integer(failed_v7_path,
                       "SELECT count(*) FROM sqlite_master WHERE type='table' AND "
                       "name='sift_extract_tasks'", 0));
-  CHECK(lardon3d_project_db_open(failed_v7_path, &database, error) == LARDON3D_PROJECT_DB_OK &&
+  Lardon3DProjectDbResult retry_v7 = lardon3d_project_db_open(
-        lardon3d_project_db_schema_version(database) == 10);
+      failed_v7_path, &database, error);
+  if (retry_v7 != LARDON3D_PROJECT_DB_OK) {
+    fprintf(stderr, "Nouvelle tentative migration v7 (%d): %s\n", (int)retry_v7, error);
+  }
+  CHECK(retry_v7 == LARDON3D_PROJECT_DB_OK &&
+        lardon3d_project_db_schema_version(database) == 11);
   lardon3d_project_db_close(database);
   database = NULL;
 
   CHECK(create_v5_database(direct_v5_path));
   CHECK(query_integer(direct_v5_path, "SELECT value FROM metadata WHERE key='schema_version'", 5));
   CHECK(lardon3d_project_db_open(direct_v5_path, &database, error) == LARDON3D_PROJECT_DB_OK &&
-        lardon3d_project_db_schema_version(database) == 10);
+        lardon3d_project_db_schema_version(database) == 11);
   lardon3d_project_db_close(database);
   database = NULL;
 
   CHECK(create_v7_database(v8_path));
   CHECK(query_integer(v8_path, "SELECT value FROM metadata WHERE key='schema_version'", 7));
   CHECK(lardon3d_project_db_open(v8_path, &database, error) == LARDON3D_PROJECT_DB_OK);
-  CHECK(lardon3d_project_db_schema_version(database) == 10);
+  CHECK(lardon3d_project_db_schema_version(database) == 11);
   lardon3d_project_db_close(database);
   database = NULL;
-  CHECK(query_integer(v8_path, "SELECT value FROM metadata WHERE key='schema_version'", 10));
+  CHECK(query_integer(v8_path, "SELECT value FROM metadata WHERE key='schema_version'", 11));
 
   CHECK(create_v7_database(failed_v8_path));
   CHECK(setenv("LARDON3D_TEST_PROJECT_DB_FAIL_MIGRATION_V8", "1", 1) == 0);
@@ -741,7 +771,48 @@ static bool run_test(void) {
                       "SELECT count(*) FROM sqlite_master WHERE type='table' AND "
                       "name='candidate_pairs'", 0));
   CHECK(lardon3d_project_db_open(failed_v8_path, &database, error) == LARDON3D_PROJECT_DB_OK &&
-        lardon3d_project_db_schema_version(database) == 10);
+        lardon3d_project_db_schema_version(database) == 11);
+  lardon3d_project_db_close(database);
+  database = NULL;
+
+  CHECK(create_v10_database(v10_path));
+  CHECK(query_integer(v10_path, "SELECT value FROM metadata WHERE key='schema_version'", 10));
+  CHECK(query_integer(v10_path,
+                      "SELECT count(*) FROM sqlite_master WHERE type='table' AND "
+                      "name='match_results'",
+                      1));
+  CHECK(query_integer(v10_path,
+                      "SELECT count(*) FROM sqlite_master WHERE type='table' AND "
+                      "name='matcher_tasks'",
+                      0));
+  CHECK(lardon3d_project_db_open(v10_path, &database, error) == LARDON3D_PROJECT_DB_OK &&
+        lardon3d_project_db_schema_version(database) == 11);
+  lardon3d_project_db_close(database);
+  database = NULL;
+  CHECK(query_integer(v10_path, "SELECT value FROM metadata WHERE key='schema_version'", 11));
+  CHECK(query_integer(v10_path,
+                      "SELECT count(*) FROM sqlite_master WHERE type='table' AND "
+                      "name='matcher_tasks'",
+                      1));
+
+  CHECK(create_v10_database(failed_v11_path));
+  CHECK(setenv("LARDON3D_TEST_PROJECT_DB_FAIL_MIGRATION_V11", "1", 1) == 0);
+  CHECK(lardon3d_project_db_open(failed_v11_path, &database, error) ==
+        LARDON3D_PROJECT_DB_IO_ERROR);
+  CHECK(unsetenv("LARDON3D_TEST_PROJECT_DB_FAIL_MIGRATION_V11") == 0);
+  CHECK(query_integer(failed_v11_path,
+                      "SELECT value FROM metadata WHERE key='schema_version'", 10));
+  CHECK(query_integer(failed_v11_path,
+                      "SELECT count(*) FROM sqlite_master WHERE type='table' AND "
+                      "name='matcher_tasks'",
+                      0));
+  CHECK(query_integer(failed_v11_path,
+                      "SELECT count(*) FROM sqlite_master WHERE type='table' AND "
+                      "name='match_results'",
+                      1));
+  CHECK(lardon3d_project_db_open(failed_v11_path, &database, error) ==
+        LARDON3D_PROJECT_DB_OK);
+  CHECK(lardon3d_project_db_schema_version(database) == 11);
   lardon3d_project_db_close(database);
   database = NULL;
 
@@ -762,6 +833,8 @@ static bool run_test(void) {
   CHECK(unlink(direct_v5_path) == 0);
   CHECK(unlink(v8_path) == 0);
   CHECK(unlink(failed_v8_path) == 0);
+  CHECK(unlink(v10_path) == 0);
+  CHECK(unlink(failed_v11_path) == 0);
   CHECK(rmdir(directory) == 0);
   return true;
 }

tests/test_resource_governor.c

@@ -84,13 +84,18 @@ run_test(void)
     };
     Lardon3DResourcePolicy policy;
     CHECK(lardon3d_resource_policy_default(&profile, &policy));
-    CHECK(policy.system_memory_reserve_bytes == GIBIBYTES(2));
+    CHECK(policy.system_memory_reserve_bytes == GIBIBYTES(4));
-    CHECK(policy.system_cpu_reserve == 1);
+    CHECK(policy.emergency_memory_floor_bytes == GIBIBYTES(2));
+    CHECK(policy.system_cpu_reserve == 4);
+    CHECK(policy.maximum_cpu_pressure_avg10 == 20.0);
+    CHECK(policy.maximum_memory_pressure_avg10 == 1.0);
     policy = (Lardon3DResourcePolicy) {
         .system_memory_reserve_bytes = GIBIBYTES(2),
         .gpu_memory_reserve_bytes = 0,
         .system_cpu_reserve = 1,
         .maximum_cpu_load_ratio = 0.90,
+        .maximum_cpu_pressure_avg10 = 20.0,
+        .maximum_memory_pressure_avg10 = 1.0,
         .maximum_io_pressure_avg10 = 80.0,
         .io_slot_capacity = 8,
     };
@@ -121,6 +126,54 @@ run_test(void)
     CHECK(decision.batch_size == 8);
     CHECK(decision.cpu_threads == 15);
     CHECK(decision.reason[0]);
+
+    snapshot.swap_activity_known = true;
+    snapshot.swap_pages_in = 100;
+    snapshot.swap_pages_out = 200;
+    CHECK(lardon3d_resource_governor_decide(governor, &snapshot, &request, &decision));
+    CHECK(lardon3d_resource_governor_pressure(governor) ==
+          LARDON3D_RESOURCE_PRESSURE_GREEN);
+    snapshot.swap_pages_out = 201;
+    CHECK(lardon3d_resource_governor_decide(governor, &snapshot, &request, &decision));
+    CHECK(decision.kind == LARDON3D_RESOURCE_REDUCE_BATCH);
+    CHECK(decision.batch_size == 4);
+    CHECK(lardon3d_resource_governor_pressure(governor) ==
+          LARDON3D_RESOURCE_PRESSURE_YELLOW);
+    snapshot.swap_pages_out = 202;
+    CHECK(lardon3d_resource_governor_decide(governor, &snapshot, &request, &decision));
+    CHECK(decision.kind == LARDON3D_RESOURCE_WAIT);
+    CHECK(lardon3d_resource_governor_pressure(governor) ==
+          LARDON3D_RESOURCE_PRESSURE_RED);
+    for (unsigned int observation = 0; observation < 2; ++observation) {
+        CHECK(lardon3d_resource_governor_decide(
+            governor, &snapshot, &request, &decision));
+        CHECK(decision.kind == LARDON3D_RESOURCE_WAIT);
+        CHECK(lardon3d_resource_governor_pressure(governor) ==
+              LARDON3D_RESOURCE_PRESSURE_RED);
+    }
+    CHECK(lardon3d_resource_governor_decide(governor, &snapshot, &request, &decision));
+    CHECK(decision.kind == LARDON3D_RESOURCE_WAIT);
+    CHECK(lardon3d_resource_governor_pressure(governor) ==
+          LARDON3D_RESOURCE_PRESSURE_YELLOW);
+    request.minimum_batch_size = 1;
+    CHECK(lardon3d_resource_governor_decide(governor, &snapshot, &request, &decision));
+    CHECK(decision.kind == LARDON3D_RESOURCE_REDUCE_BATCH);
+    CHECK(decision.batch_size == 1);
+    CHECK(lardon3d_resource_governor_decide(governor, &snapshot, &request, &decision));
+    CHECK(decision.batch_size == 1);
+    CHECK(lardon3d_resource_governor_decide(governor, &snapshot, &request, &decision));
+    CHECK(lardon3d_resource_governor_pressure(governor) ==
+          LARDON3D_RESOURCE_PRESSURE_GREEN);
+    CHECK(decision.batch_size == 1);
+    for (size_t expected = 2; expected <= 8; expected *= 2) {
+        for (unsigned int observation = 0; observation < 3; ++observation) {
+            CHECK(lardon3d_resource_governor_decide(
+                governor, &snapshot, &request, &decision));
+        }
+        CHECK(decision.batch_size == expected);
+    }
+    request.minimum_batch_size = 2;
+    snapshot.swap_activity_known = false;
     Lardon3DResourceSnapshot invalid_snapshot = snapshot;
     invalid_snapshot.cpu_load_1m = -1.0;
     CHECK(!lardon3d_resource_governor_decide(
@@ -143,12 +196,28 @@ run_test(void)
     CHECK(lardon3d_resource_governor_decide(governor, &snapshot, &request, &decision));
     CHECK(decision.kind == LARDON3D_RESOURCE_WAIT);
     snapshot.cpu_load_1m = 2.0;
+    snapshot.cpu_pressure_known = true;
+    snapshot.cpu_pressure_avg10 = 21.0;
+    CHECK(lardon3d_resource_governor_decide(governor, &snapshot, &request, &decision));
+    CHECK(decision.kind == LARDON3D_RESOURCE_WAIT);
+    snapshot.cpu_pressure_avg10 = 0.0;
+    snapshot.memory_pressure_known = true;
+    snapshot.memory_pressure_avg10 = 1.5;
+    CHECK(lardon3d_resource_governor_decide(governor, &snapshot, &request, &decision));
+    CHECK(decision.kind == LARDON3D_RESOURCE_WAIT);
+    snapshot.memory_pressure_avg10 = 0.0;
     snapshot.io_pressure_known = true;
     snapshot.io_pressure_avg10 = 90.0;
     request.io_intensive = true;
     CHECK(lardon3d_resource_governor_decide(governor, &snapshot, &request, &decision));
     CHECK(decision.kind == LARDON3D_RESOURCE_WAIT);
 
+    lardon3d_resource_governor_destroy(governor);
+    governor = lardon3d_resource_governor_create(&profile, &policy);
+    CHECK(governor);
+    snapshot.cpu_pressure_known = false;
+    snapshot.memory_pressure_known = false;
+    snapshot.io_pressure_known = false;
     request.io_intensive = false;
     request.memory_bytes_per_item = GIBIBYTES(8);
     CHECK(lardon3d_resource_governor_decide(governor, &snapshot, &request, &decision));
@@ -159,6 +228,54 @@ run_test(void)
     CHECK(decision.kind == LARDON3D_RESOURCE_REJECT);
     lardon3d_resource_governor_destroy(governor);
 
+    policy.emergency_memory_floor_bytes = GIBIBYTES(1);
+    governor = lardon3d_resource_governor_create(&profile, &policy);
+    CHECK(governor);
+    request = (Lardon3DResourceRequest) {
+        .minimum_batch_size = 1,
+        .preferred_batch_size = 8,
+        .requested_cpu_threads = 1,
+    };
+    snapshot = (Lardon3DResourceSnapshot) {
+        .memory_available_bytes = GIBIBYTES(2),
+        .cpu_load_1m = 0.0,
+        .swap_activity_known = true,
+        .swap_pages_in = 10,
+        .swap_pages_out = 10,
+    };
+    CHECK(lardon3d_resource_governor_decide(governor, &snapshot, &request, &decision));
+    CHECK(lardon3d_resource_governor_pressure(governor) ==
+          LARDON3D_RESOURCE_PRESSURE_YELLOW);
+    snapshot.memory_available_bytes = GIBIBYTES(1);
+    ++snapshot.swap_pages_out;
+    CHECK(lardon3d_resource_governor_decide(governor, &snapshot, &request, &decision));
+    CHECK(decision.kind == LARDON3D_RESOURCE_WAIT);
+    CHECK(lardon3d_resource_governor_pressure(governor) ==
+          LARDON3D_RESOURCE_PRESSURE_RED);
+    lardon3d_resource_governor_destroy(governor);
+
+    governor = lardon3d_resource_governor_create(&profile, &policy);
+    CHECK(governor);
+    snapshot.memory_available_bytes = GIBIBYTES(10);
+    snapshot.swap_activity_known = true;
+    snapshot.swap_pages_in = 50;
+    snapshot.swap_pages_out = 75;
+    CHECK(lardon3d_resource_governor_decide(governor, &snapshot, &request, &decision));
+    CHECK(lardon3d_resource_governor_pressure(governor) ==
+          LARDON3D_RESOURCE_PRESSURE_GREEN);
+    ++snapshot.swap_pages_in;
+    CHECK(lardon3d_resource_governor_decide(governor, &snapshot, &request, &decision));
+    CHECK(lardon3d_resource_governor_pressure(governor) ==
+          LARDON3D_RESOURCE_PRESSURE_YELLOW);
+    for (unsigned int observation = 0; observation < 3; ++observation) {
+        CHECK(lardon3d_resource_governor_decide(
+            governor, &snapshot, &request, &decision));
+    }
+    CHECK(lardon3d_resource_governor_pressure(governor) ==
+          LARDON3D_RESOURCE_PRESSURE_GREEN);
+    CHECK(decision.batch_size == 1);
+    lardon3d_resource_governor_destroy(governor);
+
     profile.gpu_available = true;
     profile.gpu_uses_shared_memory = true;
     policy.gpu_memory_reserve_bytes = 0;

tests/test_resource_snapshot.c

@@ -50,6 +50,14 @@ run_test(void)
         CHECK(snapshot.io_pressure_avg10 >= 0.0);
         CHECK(snapshot.io_pressure_avg10 <= 100.0);
     }
+    if (snapshot.cpu_pressure_known) {
+        CHECK(snapshot.cpu_pressure_avg10 >= 0.0);
+        CHECK(snapshot.cpu_pressure_avg10 <= 100.0);
+    }
+    if (snapshot.memory_pressure_known) {
+        CHECK(snapshot.memory_pressure_avg10 >= 0.0);
+        CHECK(snapshot.memory_pressure_avg10 <= 100.0);
+    }
     return true;
 }
 

tests/tsan-opencv.supp

@@ -4,3 +4,4 @@
 # all other libraries remain fully checked.
 race:libopencv_features.so
 race:libopencv_core.so
+race:libtbb.so

tools/embed_spirv.py

@@ -0,0 +1,45 @@
+#!/usr/bin/env python3
+
+import pathlib
+import struct
+import sys
+
+
+def main() -> int:
+    if len(sys.argv) != 3:
+        return 2
+
+    source = pathlib.Path(sys.argv[1]).read_bytes()
+    if len(source) == 0 or len(source) % 4 != 0:
+        return 1
+
+    words = struct.unpack(f"<{len(source) // 4}I", source)
+    output = pathlib.Path(sys.argv[2])
+    lines = [
+        "#ifndef LARDON3D_ORB_TOP2_SPV_H",
+        "#define LARDON3D_ORB_TOP2_SPV_H",
+        "",
+        "#include <stddef.h>",
+        "#include <stdint.h>",
+        "",
+        "static const uint32_t lardon3d_orb_top2_spv[] = {",
+    ]
+    for start in range(0, len(words), 6):
+        chunk = ", ".join(f"0x{word:08x}U" for word in words[start : start + 6])
+        lines.append(f"    {chunk},")
+    lines.extend(
+        [
+            "};",
+            "static const size_t lardon3d_orb_top2_spv_size =",
+            "    sizeof(lardon3d_orb_top2_spv);",
+            "",
+            "#endif",
+            "",
+        ]
+    )
+    output.write_text("\n".join(lines), encoding="utf-8")
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())