warp-compile-time-optimizer

Par nvidia · skills

À utiliser lorsque le problème concerne les temps de compilation ou de démarrage dans du code qui utilise Warp : une demande d'amélioration, d'optimisation ou de réduction des temps de compilation ; une application lente au démarrage ou qui se bloque au premier `wp.launch` ; des secondes de compilation avant que le vrai travail commence ; des modules JIT qui recompilent à chaque exécution ou à chaque job CI. S'applique uniquement lorsque le code optimisé utilise des kernels Warp. Ne concerne pas les performances du kernel en régime établi, la mémoire, la correction, la compilation de Warp lui-même depuis les sources, ni les temps de build nvcc/C++.

npx skills add https://github.com/nvidia/skills --skill warp-compile-time-optimizer

Temps de compilation à froid Warp

Commencer par une commande exécutable

La sonde exécute la cible dans un sous-processus, n'écrit que dans des répertoires temporaires, et ne nécessite aucun réseau, serveur d'outils externe, ou checkout Warp.

Script Objectif Arguments
scripts/warp_compile_probe.py Mesurer la compilation et les lancements isolés en froid/chaud. measure [OPTIONS] -- COMMAND...; utilisez --help.

Utilisez run_script("scripts/warp_compile_probe.py", args=[...]) quand c'est supporté; sinon utilisez la commande Python ci-dessous. La cible doit s'exécuter jusqu'à son terme.

Modèle de compilation

Warp compile les modules, non les kernels individuels. L'identité d'un module est :

(live kernel & function set) x (module options) x (CUDA block_dim) x (generic instances)

Chaque identité nécessite une génération de code et une compilation native pour le module complet.

Le coût du démarrage à froid est approximativement :

nombre d'identités de module distinctes que vous touchez  x  taille de chaque module

Réduisez-le de deux façons :

  1. Arrêtez l'agitation d'identité. Un module qui change après le chargement se compile à nouveau.
  2. Arrêtez la duplication de source. Un module utilisé à trois dimensions de bloc se compile chaque kernel trois fois.

Supprimer un kernel d'un module qui se construit toujours n'économise qu'une partie d'une compilation. Supprimer une identité de module inutile économise la compilation complète.

Quand aucune des deux ne s'applique, chevachez les builds CUDA indépendants (CS-13). Cela change quand le travail se fait, pas combien est compilé, alors jugez-le sur le temps écoulé.

Quand les options d'un module sont fixées

Les options ont deux délais :

  1. Création du module (import Python). Le module copie enable_backward, max_unroll, lineinfo, deterministic, deterministic_max_records, et compile_time_trace depuis warp.config. Définir un global plus tard est silencieusement ignoré par ce module. default_grid_stride est l'exception.
  2. Premier chargement. Avant la compilation, changez le module existant avec wp.set_module_options() ou wp.get_module(name).options. Changer une option après le chargement crée une nouvelle identité et reconstruit le module (CS-3).
Délai manqué Symptôme Coût
wp.config.* défini après import hash inchangé, option silencieusement absente l'avantage entier, invisiblement
options de module définies après chargement un second hash, module compilé deux fois une compilation supplémentaire, visible dans la trace

Après avoir changé une option, confirmez que le hash a bougé pour chaque module cible. Un hash inchangé signifie que l'option n'est jamais arrivée.

Préserver le comportement

Changez comment Warp compile le code, pas la charge de travail.

Ne supprimez pas ou ne fusionnez pas les kernels pour affirmer un gain. Les étapes apparemment redondantes peuvent préserver la propriété, l'aliasing, les sorties conservées, les limites numériques, ou le comportement de l'API. Réparez la duplication au niveau du module.

Préservez chaque lancement et son ordre, ses dimensions, dtypes, appareils, dimensions de bloc, gradients, modes numériques, comportement dynamique/plugin, et signatures d'API publiques. Gardez les noms des kernels en déplaçant les définitions vers la portée du module car les logs, les artefacts de cache, et les outils externes les exposent.

Instructions

1. Trouvez le coût réel et confirmez qu'il s'agit de compilation

Demandez quelle commande l'utilisateur attend réellement, puis mesurez-la à froid :

python scripts/warp_compile_probe.py measure --samples 3 \
    --json baseline.json -- <the user's command>

La sonde donne à chaque exemple ses propres répertoires WARP_CACHE_PATH, WARP_CACHE_ROOT, et CUDA_CACHE_PATH, active les timeurs de module, et enregistre les lancements. Elle crée ces répertoires sous le répertoire temporaire système et les supprime elle-même, donc isoler un exemple ne nécessite jamais d'écrire un cache dans le projet ou de supprimer quoi que ce soit pour re-mesurer. Isolez un exemple fait à la main de la même façon. N'effacez jamais un cache actif avec wp.clear_kernel_cache() ou wp.clear_lto_cache(); l'effacement n'est pas isolé et peut perturber d'autres processus.

Lisez la sortie de la sonde avant la source. Si la compilation est une petite partie du temps mural, rapportez le goulot d'étranglement réel et arrêtez. Pour les bibliothèques et les tests, utilisez la plus petite commande qui compile les modules de la charge de travail.

Les modules qui se sont chacun compilés une fois, sans hashes répétés, variantes de dimension de bloc, ou LTO, n'ont pas d'agitation structurelle. Cela exclut les builds redondants, pas les builds surdimensionnées; vérifiez toujours la réutilisation du cache (CS-2), la codegen rétroactive (CS-10), le déroulement (CS-11), l'en-tête précompilé (CS-12), et le chevauchement quand plusieurs modules CUDA restent (CS-13).

Chaque exemple ré-exécute également la commande contre le cache qu'il vient de remplir. Si le travail du module chaud n'est pas près de zéro, diagnostiquez la réutilisation du cache (CS-2) avant de changer la structure du module.

2. Posez des questions sur les compromis d'exécution si nécessaire

Appliquez les correctifs d'ordonnancement, le regroupement du cycle de vie, et les options hoistées sans demander. Posez une question avant de changer fast_math, max_unroll, ou une implémentation MathDx/tile :

Certains de ces réglages réduisent le temps de compilation mais peuvent rendre les kernels compilés plus lents ou changer la numériquement. Optimisez-vous une boucle d'édition-exécution rapide (où les kernels plus lents sont généralement acceptables), ou un démarrage en production (où ils généralement ne le sont pas)?

Si l'utilisateur n'est pas disponible :

  • Laissez les options qui changent la numériquement et les échanges d'implémentation tranquilles.
  • Avant de supprimer une capacité telle que la codegen rétroactive, testez si elle est utilisée; rapportez ce qui est supprimé, l'avantage mesuré, et comment revenir.
  • Définissez les options globales wp.config.* aux points d'entrée de l'application, pas dans le code de la bibliothèque.

Enregistrez les options déclinées et leurs avantages mesurés dans le grand livre de l'étape 6.

Correspondez la portée du changement à la portée de la preuve

Un profil supporte un changement à l'application mesurée, pas à chaque consommateur d'une bibliothèque partagée. Les recherches de dépôt manquent également les appelants hors arbre et futurs. Par exemple, une application orientée avant uniquement ne justifie pas de désactiver les gradients à l'intérieur d'une bibliothèque de solveur qu'une autre application différencie.

Limitez l'option au processus mesuré, avant d'importer la bibliothèque :

import warp as wp

wp.config.enable_backward = False   # must precede the library import

import the_library

Cela atteint également chaque module que l'application charge. CS-10 couvre le piège de l'ordre d'import silencieux. Si seul un changement de bibliothèque fonctionne, envoyez sa mesure aux responsables et laissez-les décider du contrat.

3. Diagnostiquez à partir de la mesure, pas à partir de la lecture de la source

La sonde imprime chaque identité de module compilée avec son nom, hash, appareil, et dimension de bloc, puis nomme quels modules se sont construits plus d'une fois. Correspondez à ce que vous voyez :

Ce que la sonde affiche Ce que cela signifie Où regarder
Un nom de module, plusieurs hashes Agitation d'identité : son ensemble de kernel, options, ou instances génériques ont changé après son premier chargement CS-1, CS-3, CS-6
Un nom de module, plusieurs valeurs block_dim Le module entier est recompilé par dimension de bloc (CUDA) CS-5
Beaucoup de modules à un kernel dans une fonctionnalité Coût fixe par module répété CS-4
Un module nommé par hash par kernel module="unique" utilisé sur les kernels stables CS-9
Grand écart entre le temps du module et le temps de compilation native, plus des artefacts .lto Configuration MathDx/LTO CS-7
(compiled) sur une exécution qui aurait dû être chaude Le cache n'est pas réutilisé CS-2
Les modules se chargent, puis "Failed to find module" Race de première utilisation CPU JIT concurrente CS-8
Source générée volumineuse, aucun problème de reconstruction Budget de déroulement CS-11
Code adjoint dans un module que rien ne différencie Codegen rétroactive CS-10
Compile lentement dans l'ensemble, ou quelques petits modules sur CUDA sous la boîte à outils 13 L'en-tête précompilé est désactivé, ou ne s'amortit pas CS-12
Plusieurs modules indépendants, chacun construit une fois, overlap_factor proche de 1,0 Les builds s'exécutent un à la fois; le chargement parallèle est désactivé par défaut CS-13
Une option que vous avez définie n'a rien changé, et le hash de ce module est inchangé Il a été attribué après la création du module, donc il n'est jamais arrivé "Quand les options d'un module sont fixées"
Aucune ligne ci-dessus ne s'active Rien n'est construit de façon redondante; le coût est la taille des constructions elles-mêmes Étape 6

references/mechanisms.md a une section par mécanisme : comment le confirmer, le correctif, ses limites, et son mode défaillance. Lisez uniquement les sections sélectionnées par la mesure.

4. Choisissez les limites du module délibérément

Regroupez les kernels dans un seul module uniquement quand ils partagent :

  • cycle de vie : ils sont définis, chargés et invalidés ensemble;
  • ensemble d'options : ils ont besoin des mêmes paramètres fast_math, enable_backward, max_unroll, et MathDx;
  • dimension de bloc stable sur CUDA.

Les kernels avec le même cycle de vie mais des dimensions de bloc stables différentes ne devraient pas partager un module car chacun se compilerait deux fois. Séparez également les kernels avec des cycles de vie indépendants.

Les kernels dont la dimension de bloc varie au moment de l'exécution (choisie à partir de la taille d'entrée, par exemple) n'ont pas de mappage stable, alors gardez-les dans leur propre module plutôt que de traîner un module partagé entier dans une variante supplémentaire.

Préférez le changement le moins invasif qui supprime une construction. Les correctifs d'ordonnancement et les options hoistées sont moins chers et plus sûrs que de ré-architecturer la disposition du module; regroupez uniquement quand le coût fixe par module ou la duplication de dimension de bloc domine.

wp.set_module_options() cible son module Python appelant, pas les kernels avec un module="pkg.name" explicite. Soit utilisez un vrai module Python soit mettez à jour le module nommé avant qu'il se charge :

wp.set_module_options({"enable_backward": False})  # at module scope

wp.get_module("pkg.name").options.update({"enable_backward": False})

Ne passez pas wp.get_module() à wp.set_module_options(module=...), ou utilisez @wp.kernel(module_options={...}) sans module="unique". enable_backward=False par kernel a une exception tile-module couverte par CS-10. Confirmez le hash du module après chaque changement d'option.

5. Vérifiez

python scripts/warp_compile_probe.py measure --samples 3 \
    --json candidate.json -- <the same command>
python scripts/warp_compile_probe.py compare baseline.json candidate.json

compare rejette la topologie de lancement changée et traite un résultat à l'intérieur de max(1% de la ligne de base, 2 x ligne de base MAD) comme inconclusive.

Pour BUILDS OVERLAPPED, jugez les changements d'ordonnancement sur la compilation écoulée plutôt que sur les timeurs de module sommés. La passe chaude requise fournit cette horloge. Voir references/measurement.md.

Vérifiez ensuite ce que la sonde ne peut pas voir :

  • Diff la sortie numérique et exécutez les tests ou les points d'entrée du projet.
  • Après avoir changé enable_backward ou les limites, vérifiez un chemin de gradient.
  • Après avoir changé fast_math, max_unroll, MathDx, ou une implémentation, faites un benchmark en état stable.
  • Exercez les kernels dynamiques non couverts, dtypes, profils, et modes.

6. Rapportez les résultats

Lisez "Reporting results" dans references/measurement.md. Rapportez :

Décrivez chaque optimisation en langage clair : nommez le comportement, la preuve, et l'effet. Par exemple, écrivez "options du module déplacées avant le premier chargement pour éviter une reconstruction redondante," pas "appliqué CS-3." Traitez les étiquettes CS-* comme des aides à la navigation interne, pas comme des explications visibles par l'utilisateur.

Option Mesuré Pourquoi pas pris Pour le prendre
enable_backward=False sur pkg.solver −38% froid une bande enregistrée en direct parcourt ces kernels définir au point d'entrée, puis vérifier à nouveau les adjoints
max_unroll=4 −2%, dans le bruit change le code généré sans gain mesuré
  • médianes avant/après et nombres d'exemples;
  • la réduction et le coût résiduel, dimensionnés par rapport à la plainte initiale;
  • mécanismes corrigés, écartés, mesurés et déclinés, ou non atteints;
  • compromis et comportement non vérifiés;
  • le grand livre ci-dessus pour les options déclinées ou incomplètement vérifiées;
  • la valeur mesurée de relâcher toute contrainte qui a bloqué un correctif.

Énoncez uniquement ce que la preuve supporte. "Pas d'agitation structurelle" ne signifie pas "optimal" ou "irréductible." Mesurez les leviers déclinés quand c'est pratique; étiquetez toute estimation non testée. Avant de rapporter qu'aucun correctif n'est disponible, vérifiez CS-13. Pour un split de module possible, mesurez d'abord un module à un kernel avec les mêmes options pour établir le coût fixe répété.

Dépannage

Exécutez la commande cible directement avant de déboguer la sonde. Voir references/measurement.md pour les problèmes de cache/bruit et references/mechanisms.md pour les défaillances spécifiques au mécanisme.

Limitations

Deux règles supplantent tout gain :

  • Isolez à la fois les caches Warp et CUDA pour chaque exemple à froid.
  • Gardez max_workers <= 1 quand une charge peut cibler CPU, y compris device=None et les listes d'appareils mixtes. Les premiers chargements CPU concurrents peuvent perdre des kernels; les tentatives n'en font pas pour autant que ce soit sûr. Le chargement CUDA uniquement n'est pas affecté.

Les mesures sont spécifiques à l'environnement : les temps froids bougent avec CPU, GPU, driver, toolchain, et version Warp. Les mécanismes se transfèrent; les chiffres ne le font pas. Lisez les inconnues connues dans references/mechanisms.md avant de faire de larges affirmations.

Fichiers de référence

  • references/mechanisms.md : les treize mécanismes de temps de compilation, chacun avec son signal de confirmation, correctif, limites d'applicabilité, et mode de défaillance. Lisez les sections vers lesquelles votre mesure pointe.
  • references/measurement.md : protocole de mesure, ce que chaque métrique fait et ne signifie pas, conseils de rapport, exemples de log, et mesure manuelle.

Skills similaires