{"id":1293,"date":"2026-08-21T14:31:18","date_gmt":"2026-08-21T14:31:18","guid":{"rendered":"https:\/\/matforge.org\/?p=1293","raw":"https:\/\/matforge.org\/?p=1293"},"modified":"2026-08-21T14:31:18","modified_gmt":"2026-08-21T14:31:18","slug":"gpu-kernel-programming-custom-physics-simulation","status":"publish","type":"post","link":"https:\/\/matforge.org\/fr\/gpu-kernel-programming-custom-physics-simulation\/","title":{"rendered":"Programmation du noyau GPU pour la simulation physique personnalis\u00e9e\u00a0: \u00e9criture de noyaux personnalis\u00e9s hautes performances en Python","raw":"Programmation du noyau GPU pour la simulation physique personnalis\u00e9e\u00a0: \u00e9criture de noyaux personnalis\u00e9s hautes performances en Python"},"content":{"rendered":"<span class=\"span-reading-time rt-reading-time\" style=\"display: block;\"><span class=\"rt-label rt-prefix\">Reading Time: <\/span> <span class=\"rt-time\"> 14<\/span> <span class=\"rt-label rt-postfix\">minutes<\/span><\/span><article>\n<p>L&rsquo;acc\u00e9l\u00e9ration des GPU en informatique scientifique commence souvent par des biblioth\u00e8ques de haut niveau. Un chercheur remplace NumPy par CUPY, utilise une FFT acc\u00e9l\u00e9r\u00e9e ou appelle une routine d&rsquo;alg\u00e8bre lin\u00e9aire activ\u00e9e par GPU. Cela peut produire une am\u00e9lioration substantielle sans n\u00e9cessiter de connaissances d\u00e9taill\u00e9es sur le mat\u00e9riel GPU.<\/p>\n<p>La programmation du noyau GPU personnalis\u00e9 va d&rsquo;un niveau plus profond\u00e9ment. Au lieu d&rsquo;appeler une op\u00e9ration pr\u00e9d\u00e9finie, le d\u00e9veloppeur \u00e9crit la fonction que chaque thread GPU ex\u00e9cute. Cela permet de contr\u00f4ler l&rsquo;indexation des threads, la disposition des donn\u00e9es, les transferts de m\u00e9moire, la synchronisation, la m\u00e9moire partag\u00e9e et le nombre d&rsquo;op\u00e9rations effectu\u00e9es lors de chaque lancement du noyau.<\/p>\n<p>Ce contr\u00f4le est utile lorsqu&rsquo;un mod\u00e8le physique contient des mod\u00e8les d&rsquo;acc\u00e8s qui ne peuvent pas \u00eatre exprim\u00e9s efficacement par le biais d&rsquo;op\u00e9rations de tableau ordinaires. Les interactions de particules, les pochoirs de diff\u00e9rences finies, les r\u00e8gles de collision, les mod\u00e8les de fluides bas\u00e9s sur le r\u00e9seau, les mises \u00e0 jour des mat\u00e9riaux et les requ\u00eates de g\u00e9om\u00e9trie sont des exemples courants.<\/p>\n<p>Un noyau personnalis\u00e9 n&rsquo;est pas automatiquement plus rapide qu&rsquo;une biblioth\u00e8que optimis\u00e9e. Cela devient utile lorsque l&rsquo;algorithme expose suffisamment de travail parall\u00e8le et lorsque le d\u00e9veloppeur peut r\u00e9duire le trafic m\u00e9moire, combiner les op\u00e9rations ou adapter le calcul \u00e0 l&rsquo;architecture GPU.<\/p>\n<h2>Qu&rsquo;est-ce qu&rsquo;un noyau GPU&nbsp;?<\/h2>\n<p>Un noyau GPU est une fonction ex\u00e9cut\u00e9e par de nombreux threads l\u00e9gers. Chaque thread g\u00e8re normalement une particule, une cellule de grille, un \u00e9l\u00e9ment, une face, un pixel ou une entr\u00e9e dans un tableau.<\/p>\n<p>Le programme h\u00f4te lance le noyau avec une grille de blocs de thread :<\/p>\n<pre><code>kernel[blocks, threads_per_block](arguments)<\/code><\/pre>\n<p>Chaque thread d\u00e9termine les donn\u00e9es qu&rsquo;il doit traiter \u00e0 partir de ses indices de blocs et de threads. Un index unidimensionnel typique est :<\/p>\n<pre><code>index =\n    block_index * block_size\n    + thread_index<\/code><\/pre>\n<p>Les threads sont organis\u00e9s en groupes appel\u00e9s blocs. Les threads \u00e0 l&rsquo;int\u00e9rieur d&rsquo;un seul bloc peuvent synchroniser et \u00e9changer des donn\u00e9es via la m\u00e9moire partag\u00e9e. Les blocs s\u00e9par\u00e9s doivent g\u00e9n\u00e9ralement s&rsquo;ex\u00e9cuter de mani\u00e8re ind\u00e9pendante.<\/p>\n<h2>Lorsque les noyaux personnalis\u00e9s valent la peine d&rsquo;\u00eatre \u00e9crits<\/h2>\n<p>Les biblioth\u00e8ques de GPU de haut niveau devraient g\u00e9n\u00e9ralement \u00eatre la premi\u00e8re option. Ils fournissent d\u00e9j\u00e0 une multiplication matricielle optimis\u00e9e, des FFT, des r\u00e9ductions, des op\u00e9rations rares, une g\u00e9n\u00e9ration de nombres al\u00e9atoires et de nombreuses fonctions par \u00e9l\u00e9ment.<\/p>\n<p>Un noyau personnalis\u00e9 devient utile lorsque :<\/p>\n<ul>\n<li>La simulation utilise un mod\u00e8le d&rsquo;acc\u00e8s \u00e0 la m\u00e9moire non standard.<\/li>\n<li>Plusieurs petites op\u00e9rations peuvent \u00eatre fusionn\u00e9es en un seul passage sur la m\u00e9moire.<\/li>\n<li>Les threads voisins ont \u00e0 plusieurs reprises besoin des m\u00eames donn\u00e9es locales.<\/li>\n<li>Une r\u00e8gle sp\u00e9cialis\u00e9e de particule, de collision ou de pochoir domine le temps d&rsquo;ex\u00e9cution.<\/li>\n<li>Les baies interm\u00e9diaires consomment trop de m\u00e9moire de p\u00e9riph\u00e9rique.<\/li>\n<li>La simulation n\u00e9cessite une op\u00e9ration diff\u00e9renciable personnalis\u00e9e.<\/li>\n<li>Une biblioth\u00e8que existante ne peut pas exprimer efficacement la limite ou la logique mat\u00e9rielle requise.<\/li>\n<\/ul>\n<p>Avant d&rsquo;\u00e9crire un noyau, profilez le programme. L&rsquo;optimisation d&rsquo;une fonction visuellement complexe qui ne repr\u00e9sente qu&rsquo;une petite partie du temps d&rsquo;ex\u00e9cution ne produira pas une acc\u00e9l\u00e9ration globale significative.<\/p>\n<h2>Mod\u00e8les de programmation GPU<\/h2>\n<p>Les logiciels scientifiques peuvent acc\u00e9der aux GPU via plusieurs niveaux d&rsquo;abstraction.<\/p>\n<table class=\"custom-table\">\n<tbody>\n<tr>\n<th>Approche<\/th>\n<th>Exemples<\/th>\n<th>force principale<\/th>\n<th>Principal compromis<\/th>\n<\/tr>\n<tr>\n<td>D\u00e9chargement bas\u00e9 sur les directives<\/td>\n<td>OpenMP, OpenACC<\/td>\n<td>Acc\u00e9l\u00e9ration incr\u00e9mentale du code C, C++ ou Fortran existant<\/td>\n<td>Moins de contr\u00f4le direct sur les noyaux g\u00e9n\u00e9r\u00e9s<\/td>\n<\/tr>\n<tr>\n<td>Programmation du noyau orient\u00e9e fournisseur<\/td>\n<td>cud\u00e9, hanche<\/td>\n<td>Contr\u00f4le d\u00e9taill\u00e9 de l&rsquo;ex\u00e9cution du GPU<\/td>\n<td>Plus de co\u00fbts de mise en \u0153uvre et de maintenance<\/td>\n<\/tr>\n<tr>\n<td>C++ portable de performances<\/td>\n<td>SYCL, Kokkos, Alpaka<\/td>\n<td>Un mod\u00e8le source pour plusieurs backends mat\u00e9riels<\/td>\n<td>La portabilit\u00e9 ne garantit pas une performance \u00e9gale partout<\/td>\n<\/tr>\n<tr>\n<td>Tableaux et noyaux de GPU Python<\/td>\n<td>CUPY, NUMBA-CUDA, NVIDIA WARP<\/td>\n<td>D\u00e9veloppement rapide avec acc\u00e8s au code GPU compil\u00e9<\/td>\n<td>Restrictions sp\u00e9cifiques au framework et comportement de compilation<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n<p><a href=\"https:\/\/enccs.github.io\/gpu-programming\/5-intro-to-gpu-prog-models\/\" rel=\"nofollow\" target=\"_blank\">ENCCS Introduction \u00e0 la programmation GPU Les mod\u00e8les<\/a> fournissent une comparaison plus large des approches CUDA, HIP, OpenCL, SYCL, Kokkos, OpenMP et associ\u00e9es.<\/p>\n<h2>Cuda<\/h2>\n<p>CUDA fournit un acc\u00e8s direct \u00e0 la programmation NVIDIA GPU via C++, Fortran et plusieurs liaisons de langage. Il expose les grilles, les blocs, les d\u00e9formations, la m\u00e9moire partag\u00e9e, les flux, les \u00e9v\u00e9nements, les biblioth\u00e8ques de p\u00e9riph\u00e9riques et les outils d&rsquo;optimisation sp\u00e9cifiques au mat\u00e9riel.<\/p>\n<p>CUDA est souvent s\u00e9lectionn\u00e9 lorsque des performances maximales sp\u00e9cifiques \u00e0 NVIDIA et des outils matures sont plus importants que la portabilit\u00e9 mat\u00e9rielle. Les principaux co\u00fbts sont la gestion de la m\u00e9moire manuelle, un mod\u00e8le de programmation de niveau inf\u00e9rieur et la d\u00e9pendance \u00e0 la plate-forme NVIDIA.<\/p>\n<h2>Hanche et ROCM<\/h2>\n<p>HIP fournit un mod\u00e8le de programmation C++ de type CUDA dans l&rsquo;\u00e9cosyst\u00e8me ROCM d&rsquo;AMD. Une grande quantit\u00e9 de code de style CUDA peut \u00eatre adapt\u00e9e \u00e0 la hanche, mais la portabilit\u00e9 de la source ne signifie pas qu&rsquo;un binaire s&rsquo;ex\u00e9cute inchang\u00e9 sur chaque GPU.<\/p>\n<p>Une compilation, des tests et un r\u00e9glage des performances s\u00e9par\u00e9s peuvent toujours \u00eatre n\u00e9cessaires pour chaque architecture cible. Les d\u00e9tails sont disponibles dans la <a href=\"https:\/\/rocm.docs.amd.com\/projects\/HIP\/en\/latest\/what_is_hip.html\" rel=\"nofollow\" target=\"_blank\">documentation officielle de la hanche<\/a>.<\/p>\n<p>HIP est une option pratique lorsque les GPU AMD doivent \u00eatre pris en charge ou lorsqu&rsquo;un projet souhaite r\u00e9duire la d\u00e9pendance \u00e0 un seul fournisseur d&rsquo;acc\u00e9l\u00e9rateurs tout en conservant un style de programmation de type CUDA.<\/p>\n<h2>Cupy Cernels personnalis\u00e9s<\/h2>\n<p>Cupy est surtout connu pour ses tableaux compatibles avec NumPy et ses fonctions num\u00e9riques acc\u00e9l\u00e9r\u00e9es. Il prend \u00e9galement en charge la programmation GPU personnalis\u00e9e.<\/p>\n<p>Les d\u00e9veloppeurs peuvent utiliser :<\/p>\n<ul>\n<li><code>ElementwiseKernel<\/code> pour les expressions personnalis\u00e9es par \u00e9l\u00e9ment<\/li>\n<li><code>ReductionKernel<\/code> pour les r\u00e9ductions personnalis\u00e9es<\/li>\n<li><code>RawKernel<\/code> pour les noyaux \u00e9crits en CUDA C++<\/li>\n<li><code>cupyx.jit.rawkernel<\/code> pour les d\u00e9finitions du noyau JIT de style Python<\/li>\n<\/ul>\n<p>Un simple noyau brut Cupy peut \u00eatre d\u00e9fini comme suit :<\/p>\n<pre><code class=\"language-python\">import cupy as cp\n\nscale_kernel = cp.RawKernel(\n    r'''\n    extern \"C\" __global__\n    void scale_values(\n        const float* input,\n        float* output,\n        const float factor,\n        const int size\n    ) {\n        int index =\n            blockDim.x * blockIdx.x\n            + threadIdx.x;\n\n        if (index &lt; size) {\n            output[index] =\n                factor * input[index];\n        }\n    }\n    ''',\n    \"scale_values\"\n)\n\nsize = 1_000_000\nthreads_per_block = 256\nblocks = (\n    size + threads_per_block - 1\n) \/\/ threads_per_block\n\ninput_values = cp.arange(\n    size,\n    dtype=cp.float32\n)\n\noutput_values = cp.empty_like(\n    input_values\n)\n\nscale_kernel(\n    (blocks,),\n    (threads_per_block,),\n    (\n        input_values,\n        output_values,\n        cp.float32(2.0),\n        cp.int32(size)\n    )\n)<\/code><\/pre>\n<p>CUPY est utile lorsque la plupart de l&rsquo;application utilise d\u00e9j\u00e0 des baies de GPU et que seules les op\u00e9rations s\u00e9lectionn\u00e9es ont besoin de noyaux personnalis\u00e9s.<\/p>\n<h2>Numba-cuda<\/h2>\n<p>numba-cuda compile un sous-ensemble restreint de python dans les noyaux GPU. Son mod\u00e8le de programmation suit de pr\u00e8s CUDA C. Les d\u00e9veloppeurs d\u00e9finissent une fonction avec un d\u00e9corateur de noyau, calculent un index de fil et lancent la fonction avec une configuration de grille et de bloc.<\/p>\n<p>La documentation active du projet est disponible \u00e0 <a href=\"https:\/\/nvidia.github.io\/numba-cuda\/\" rel=\"nofollow\" target=\"_blank\">numba-cuda<\/a>. \u00c9tant donn\u00e9 que son statut de d\u00e9veloppement peut changer, les \u00e9quipes cr\u00e9ant des logiciels de longue dur\u00e9e doivent revoir les directives de maintenance et de migration actuelles avant de s&rsquo;engager dans le cadre.<\/p>\n<h2>Un noyau de particules Numba<\/h2>\n<p>L&rsquo;exemple \u00e9ducatif suivant calcule l&rsquo;acc\u00e9l\u00e9ration gravitationnelle avec des interactions directes toutes paires :<\/p>\n<pre><code class=\"language-python\">import math\nimport numpy as np\nfrom numba import cuda\n\n@cuda.jit\ndef gravity_kernel(\n    positions,\n    masses,\n    accelerations,\n    gravitational_constant,\n    softening_squared\n):\n    particle = cuda.grid(1)\n    particle_count = positions.shape[0]\n\n    if particle &gt;= particle_count:\n        return\n\n    px = positions[particle, 0]\n    py = positions[particle, 1]\n    pz = positions[particle, 2]\n\n    ax = 0.0\n    ay = 0.0\n    az = 0.0\n\n    for other in range(particle_count):\n        if other == particle:\n            continue\n\n        dx = positions[other, 0] - px\n        dy = positions[other, 1] - py\n        dz = positions[other, 2] - pz\n\n        distance_squared = (\n            dx * dx\n            + dy * dy\n            + dz * dz\n            + softening_squared\n        )\n\n        inverse_distance = (\n            1.0\n            \/ math.sqrt(distance_squared)\n        )\n\n        inverse_distance_cubed = (\n            inverse_distance\n            * inverse_distance\n            * inverse_distance\n        )\n\n        scale = (\n            gravitational_constant\n            * masses[other]\n            * inverse_distance_cubed\n        )\n\n        ax += scale * dx\n        ay += scale * dy\n        az += scale * dz\n\n    accelerations[particle, 0] = ax\n    accelerations[particle, 1] = ay\n    accelerations[particle, 2] = az\n\n\nparticle_count = 10_000\nthreads_per_block = 256\n\nblocks = (\n    particle_count\n    + threads_per_block\n    - 1\n) \/\/ threads_per_block\n\nhost_positions = np.random.random(\n    (particle_count, 3)\n).astype(np.float32)\n\nhost_masses = np.ones(\n    particle_count,\n    dtype=np.float32\n)\n\ndevice_positions = cuda.to_device(\n    host_positions\n)\n\ndevice_masses = cuda.to_device(\n    host_masses\n)\n\ndevice_accelerations = cuda.device_array(\n    (particle_count, 3),\n    dtype=np.float32\n)\n\ngravity_kernel[\n    blocks,\n    threads_per_block\n](\n    device_positions,\n    device_masses,\n    device_accelerations,\n    np.float32(1.0),\n    np.float32(1e-4)\n)\n\ncuda.synchronize()\n\naccelerations = (\n    device_accelerations.copy_to_host()\n)<\/code><\/pre>\n<p>Ce noyau a une complexit\u00e9 de calcul de <code>O(N\u00b2)<\/code>. Il est adapt\u00e9 pour expliquer la cartographie des threads, mais ce n&rsquo;est pas un algorithme de production efficace pour les tr\u00e8s gros nombres de particules.<\/p>\n<p>Les grands syst\u00e8mes de gravitation ou de particules n\u00e9cessitent souvent une m\u00e9thode d&rsquo;arbre, une m\u00e9thode multip\u00f4le rapide, des bacs spatiaux ou des listes de voisins. Un GPU ne peut pas supprimer le probl\u00e8me de mise \u00e0 l&rsquo;\u00e9chelle d&rsquo;un algorithme inefficace.<\/p>\n<h2>NVIDIA WARP<\/h2>\n<p>NVIDIA WARP permet aux d\u00e9veloppeurs de d\u00e9finir des noyaux fortement typ\u00e9s avec la syntaxe Python. WARP compile ces fonctions en CPU ou CUDA et fournit des primitives sp\u00e9cialis\u00e9es pour la g\u00e9om\u00e9trie, la simulation, les op\u00e9rations rares, les \u00e9l\u00e9ments finis, l&rsquo;optimisation et la diff\u00e9renciation automatique.<\/p>\n<p>La <a href=\"https:\/\/developer.nvidia.com\/warp-python\" rel=\"nofollow\" target=\"_blank\">page produit NVIDIA WARP<\/a> et <a href=\"https:\/\/github.com\/nvidia\/warp\" rel=\"nofollow\" target=\"_blank\">R\u00e9f\u00e9rentiel GitHub de Warp<\/a> Inclut des exemples de particules, de fluides, de maillages, d&rsquo;optimisation, de simulation diff\u00e9renciable et Programmation GPU bas\u00e9e sur des tuiles.<\/p>\n<h2>Le m\u00eame noyau de particules dans WARP<\/h2>\n<pre><code class=\"language-python\">import numpy as np\nimport warp as wp\n\nwp.init()\n\n@wp.kernel\ndef gravity_kernel(\n    positions: wp.array(dtype=wp.vec3),\n    masses: wp.array(dtype=wp.float32),\n    accelerations: wp.array(dtype=wp.vec3),\n    gravitational_constant: wp.float32,\n    softening_squared: wp.float32\n):\n    particle = wp.tid()\n    particle_count = positions.shape[0]\n\n    position = positions[particle]\n    acceleration = wp.vec3(0.0, 0.0, 0.0)\n\n    for other in range(particle_count):\n        if other != particle:\n            displacement = (\n                positions[other]\n                - position\n            )\n\n            distance_squared = (\n                wp.dot(\n                    displacement,\n                    displacement\n                )\n                + softening_squared\n            )\n\n            inverse_distance = (\n                1.0\n                \/ wp.sqrt(distance_squared)\n            )\n\n            inverse_distance_cubed = (\n                inverse_distance\n                * inverse_distance\n                * inverse_distance\n            )\n\n            acceleration += (\n                gravitational_constant\n                * masses[other]\n                * inverse_distance_cubed\n                * displacement\n            )\n\n    accelerations[particle] = acceleration\n\n\nparticle_count = 10_000\n\nhost_positions = np.random.random(\n    (particle_count, 3)\n).astype(np.float32)\n\nhost_masses = np.ones(\n    particle_count,\n    dtype=np.float32\n)\n\npositions = wp.array(\n    host_positions,\n    dtype=wp.vec3,\n    device=\"cuda\"\n)\n\nmasses = wp.array(\n    host_masses,\n    dtype=wp.float32,\n    device=\"cuda\"\n)\n\naccelerations = wp.zeros(\n    particle_count,\n    dtype=wp.vec3,\n    device=\"cuda\"\n)\n\nwp.launch(\n    kernel=gravity_kernel,\n    dim=particle_count,\n    inputs=[\n        positions,\n        masses,\n        accelerations,\n        wp.float32(1.0),\n        wp.float32(1e-4)\n    ],\n    device=\"cuda\"\n)\n\nwp.synchronize()<\/code><\/pre>\n<p>WARP d\u00e9termine en interne une configuration de lancement appropri\u00e9e pour un lancement de noyau ordinaire. Les utilisateurs avanc\u00e9s peuvent toujours travailler avec des dimensions de blocs, des graphiques de commandes, une ex\u00e9cution en mosa\u00efque et des op\u00e9rations de p\u00e9riph\u00e9rique sp\u00e9cialis\u00e9es si n\u00e9cessaire.<\/p>\n<h2>Boucles de bordure de grille<\/h2>\n<p>Un noyau n&rsquo;a pas besoin d&rsquo;un thread assign\u00e9 de fa\u00e7on permanente pour chaque \u00e9l\u00e9ment. Une boucle Grid-Stride permet \u00e0 chaque thread de traiter plusieurs entr\u00e9es&nbsp;:<\/p>\n<pre><code class=\"language-python\">from numba import cuda\n\n@cuda.jit\ndef scale_with_stride(\n    input_values,\n    output_values,\n    factor\n):\n    index = cuda.grid(1)\n    stride = cuda.gridsize(1)\n\n    for position in range(\n        index,\n        input_values.size,\n        stride\n    ):\n        output_values[position] = (\n            factor\n            * input_values[position]\n        )<\/code><\/pre>\n<p>Ce mod\u00e8le s\u00e9pare le nombre de threads lanc\u00e9s de la taille totale des donn\u00e9es. Il est utile lors du traitement de tr\u00e8s grands tableaux ou de la r\u00e9utilisation d&rsquo;une configuration de lancement fixe.<\/p>\n<p>Une introduction pratique \u00e0 ce mod\u00e8le est disponible dans <a href=\"https:\/\/thedatafrog.com\/en\/articles\/cuda-kernel-python\/\" rel=\"nofollow\" target=\"_blank\">le Numba et le CUDA de la grenouille de donn\u00e9es Tutoriel<\/a>.<\/p>\n<h2>Coalition de m\u00e9moire<\/h2>\n<p>Les noyaux GPU sont souvent limit\u00e9s par la bande passante de m\u00e9moire plut\u00f4t que par le d\u00e9bit arithm\u00e9tique. Les threads \u00e0 l&rsquo;int\u00e9rieur d&rsquo;un distorsion doivent id\u00e9alement acc\u00e9der aux adresses m\u00e9moire \u00e0 proximit\u00e9 afin que le mat\u00e9riel puisse combiner leurs demandes.<\/p>\n<p>Consid\u00e9rez les donn\u00e9es de particules stock\u00e9es comme&nbsp;:<\/p>\n<pre><code>particle_0: x, y, z, mass\nparticle_1: x, y, z, mass\nparticle_2: x, y, z, mass<\/code><\/pre>\n<p>Cette disposition de tableau de structures peut \u00eatre pratique pour le code orient\u00e9 objet. Une mise en page de la structure des tableaux stocke des tableaux contigus distincts&nbsp;:<\/p>\n<pre><code>x_positions[]\ny_positions[]\nz_positions[]\nmasses[]<\/code><\/pre>\n<p>La deuxi\u00e8me mise en page peut fournir une meilleure coalescence lorsque chaque thread lit le m\u00eame champ pour une particule diff\u00e9rente. La meilleure mise en page d\u00e9pend toujours des champs accessibles ensemble.<\/p>\n<p>Les types de vecteurs ne corrigent pas automatiquement un mauvais mod\u00e8le d&rsquo;acc\u00e8s. Les d\u00e9veloppeurs doivent inspecter les adresses r\u00e9elles demand\u00e9es par les threads voisins.<\/p>\n<h2>M\u00e9moire partag\u00e9e<\/h2>\n<p>La m\u00e9moire partag\u00e9e est une petite zone de m\u00e9moire \u00e0 faible latence accessible par les threads dans le m\u00eame bloc. Il peut r\u00e9duire les lectures r\u00e9p\u00e9t\u00e9es de la m\u00e9moire globale.<\/p>\n<p>Un pochoir unidimensionnel peut charger un bloc de valeurs et son halo dans la m\u00e9moire partag\u00e9e :<\/p>\n<pre><code class=\"language-python\">import numpy as np\nfrom numba import cuda, float32\n\nBLOCK_SIZE = 256\n\n@cuda.jit\ndef three_point_stencil(\n    input_values,\n    output_values\n):\n    shared = cuda.shared.array(\n        shape=BLOCK_SIZE + 2,\n        dtype=float32\n    )\n\n    local_index = cuda.threadIdx.x\n    global_index = cuda.grid(1)\n    size = input_values.size\n\n    center = local_index + 1\n\n    if global_index &lt; size:\n        shared[center] = (\n            input_values[global_index]\n        )\n    else:\n        shared[center] = 0.0\n\n    if local_index == 0:\n        left_index = global_index - 1\n\n        shared[0] = (\n            input_values[left_index]\n            if left_index &gt;= 0\n            else 0.0\n        )\n\n    if local_index == BLOCK_SIZE - 1:\n        right_index = global_index + 1\n\n        shared[BLOCK_SIZE + 1] = (\n            input_values[right_index]\n            if right_index &lt; size\n            else 0.0\n        )\n\n    cuda.syncthreads()\n\n    if global_index &lt; size:\n        output_values[global_index] = (\n            shared[center - 1]\n            + shared[center]\n            + shared[center + 1]\n        ) \/ 3.0<\/code><\/pre>\n<p>L&rsquo;appel \u00e0 <code>cuda.syncthreads()<\/code> garantit que chaque thread finit de charger ses valeurs avant le d\u00e9but du calcul du pochoir.<\/p>\n<p>La m\u00e9moire partag\u00e9e ne doit pas \u00eatre utilis\u00e9e automatiquement. Une allocation de m\u00e9moire partag\u00e9e excessive peut r\u00e9duire l&rsquo;occupation, augmenter les co\u00fbts de synchronisation et ralentir le noyau. Le profilage est requis.<\/p>\n<h2>Fusion du noyau<\/h2>\n<p>Des op\u00e9rations de tableaux distinctes cr\u00e9ent souvent plusieurs lancements de noyau et des tableaux interm\u00e9diaires&nbsp;:<\/p>\n<pre><code>velocity += dt * acceleration\nposition += dt * velocity\nenergy = compute_energy(position, velocity)<\/code><\/pre>\n<p>Un noyau fusionn\u00e9 peut calculer les trois mises \u00e0 jour alors que les valeurs n\u00e9cessaires restent dans les registres. Cela r\u00e9duit les frais g\u00e9n\u00e9raux de lancement et le trafic de m\u00e9moire globale.<\/p>\n<p>La fusion est plus avantageuse lorsque les op\u00e9rations sont simples et li\u00e9es \u00e0 la m\u00e9moire. Fusionner trop de travail peut augmenter l&rsquo;utilisation des registres, r\u00e9duire l&rsquo;occupation et rendre le noyau difficile \u00e0 entretenir.<\/p>\n<h2>Choisir la taille du bloc<\/h2>\n<p>Des valeurs telles que 128 ou 256 threads par bloc sont des points de d\u00e9part raisonnables, et non des optima universels.<\/p>\n<p>La meilleure taille de bloc d\u00e9pend de :<\/p>\n<ul>\n<li>Registres utilis\u00e9s par fil<\/li>\n<li>M\u00e9moire partag\u00e9e utilis\u00e9e par bloc<\/li>\n<li>Divergence de branche<\/li>\n<li>m\u00e9lange d&rsquo;instructions<\/li>\n<li>Comportement d&rsquo;acc\u00e8s \u00e0 la m\u00e9moire<\/li>\n<li>L&rsquo;architecture GPU cible<\/li>\n<\/ul>\n<p>Une r\u00e8gle telle que le lancement d&rsquo;au moins deux fois plus de blocs que les multiprocesseurs de streaming peut \u00eatre une exp\u00e9rience initiale utile, mais elle ne garantit pas des performances maximales. Les calculatrices d&rsquo;occupation et les outils de profilage doivent guider la configuration finale.<\/p>\n<h2>Physique diff\u00e9rentiable<\/h2>\n<p>La simulation diff\u00e9rentiable calcule comment une sortie change par rapport aux entr\u00e9es telles que les propri\u00e9t\u00e9s des mat\u00e9riaux, les forces, la g\u00e9om\u00e9trie ou les conditions initiales.<\/p>\n<p>WARP peut enregistrer les op\u00e9rations du noyau pris en charge et ex\u00e9cuter la diff\u00e9renciation automatique en mode inverse. Cela peut \u00eatre utilis\u00e9 pour les probl\u00e8mes inverses, l&rsquo;optimisation de la conception, l&rsquo;estimation des param\u00e8tres et l&rsquo;int\u00e9gration avec les flux de travail d&rsquo;apprentissage automatique.<\/p>\n<p>Un motif simplifi\u00e9 utilise une bande :<\/p>\n<pre><code class=\"language-python\">with wp.Tape() as tape:\n    wp.launch(\n        kernel=simulation_kernel,\n        dim=element_count,\n        inputs=[state, parameters],\n        outputs=[result],\n        device=\"cuda\"\n    )\n\n    wp.launch(\n        kernel=loss_kernel,\n        dim=element_count,\n        inputs=[result, target],\n        outputs=[loss],\n        device=\"cuda\"\n    )\n\ntape.backward(loss)<\/code><\/pre>\n<p>La diff\u00e9renciation automatique n&rsquo;est pas garantie pour chaque noyau. Les \u00e9crasements sur place, les atomes non d\u00e9terministes, le code natif externe, la logique discontinue et les op\u00e9rations non pris en charge peuvent n\u00e9cessiter une reformulation ou des gradients personnalis\u00e9s.<\/p>\n<p>Cette fonctionnalit\u00e9 est directement connect\u00e9e \u00e0 des applications telles que l&rsquo;optimisation des associations et les <a href=\"https:\/\/matforge.org\/physics-informed-neural-networks-pinns-for-scientific-simulations\/\" rel=\"nofollow\" target=\"_blank\">r\u00e9seaux de neurones \u00e0 base de physique<\/a>.<\/p>\n<h2>Lorsqu&rsquo;un GPU n&rsquo;est peut-\u00eatre pas plus rapide<\/h2>\n<p>Il n&rsquo;y a pas de seuil universel de taille de probl\u00e8me auquel un GPU devient plus rapide qu&rsquo;un processeur. Le crossover d\u00e9pend du mat\u00e9riel, de la pr\u00e9cision, du mouvement des donn\u00e9es, de la structure des algorithmes, de la qualit\u00e9 du compilateur et de la fr\u00e9quence \u00e0 laquelle les m\u00eames donn\u00e9es de r\u00e9sidents sont r\u00e9utilis\u00e9es.<\/p>\n<table class=\"custom-table\">\n<tbody>\n<tr>\n<th>La charge de travail<\/th>\n<th>aptitude au GPU<\/th>\n<th>consid\u00e9ration principale<\/th>\n<\/tr>\n<tr>\n<td>Mise \u00e0 jour des grandes particules<\/td>\n<td>Souvent fort<\/td>\n<td>De nombreuses op\u00e9rations ind\u00e9pendantes similaires<\/td>\n<\/tr>\n<tr>\n<td>Grand pochoir r\u00e9gulier<\/td>\n<td>Souvent fort<\/td>\n<td>Acc\u00e8s \u00e0 la m\u00e9moire parall\u00e8le pr\u00e9visible<\/td>\n<\/tr>\n<tr>\n<td>Alg\u00e8bre lin\u00e9aire dense<\/td>\n<td>Fort lors de l&rsquo;utilisation de biblioth\u00e8ques optimis\u00e9es<\/td>\n<td>Intensit\u00e9 arithm\u00e9tique \u00e9lev\u00e9e<\/td>\n<\/tr>\n<tr>\n<td>Petite simulation<\/td>\n<td>d\u00e9pendant du mat\u00e9riel<\/td>\n<td>Les frais g\u00e9n\u00e9raux de lancement et de transfert peuvent dominer<\/td>\n<\/tr>\n<tr>\n<td>Travers\u00e9e de graphes irr\u00e9guli\u00e8res<\/td>\n<td>Mixte<\/td>\n<td>Divergence et acc\u00e8s m\u00e9moire impr\u00e9visible<\/td>\n<\/tr>\n<tr>\n<td>Flux de travail li\u00e9 aux E\/S<\/td>\n<td>g\u00e9n\u00e9ralement limit\u00e9<\/td>\n<td>Le GPU ne peut pas supprimer un goulot d&rsquo;\u00e9tranglement de stockage ou de r\u00e9seau<\/td>\n<\/tr>\n<tr>\n<td>Algorithme fortement s\u00e9quentiel<\/td>\n<td>g\u00e9n\u00e9ralement faible<\/td>\n<td>Travail ind\u00e9pendant insuffisant<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n<p>Le flux de travail correct consiste \u00e0 mesurer la version du processeur, la version de la biblioth\u00e8que GPU et la version du noyau personnalis\u00e9 avec des donn\u00e9es repr\u00e9sentatives. Le <a href=\"https:\/\/matforge.org\/scientific-computing-performance-profiling-identifying-bottlenecks\/\" rel=\"nofollow\" target=\"_blank\">Guide de profilage de performances de calcul scientifique<\/a> explique comment identifier le v\u00e9ritable goulot d&rsquo;\u00e9tranglement avant de l&rsquo;optimiser.<\/p>\n<h2>\u00c9viter les frais g\u00e9n\u00e9raux de transfert<\/h2>\n<p>Les transferts r\u00e9p\u00e9t\u00e9s entre l&rsquo;h\u00f4te et la m\u00e9moire de l&rsquo;appareil peuvent \u00e9liminer les avantages du calcul du GPU.<\/p>\n<p>Une boucle inefficace peut suivre ce mod\u00e8le :<\/p>\n<ol>\n<li>Copiez les donn\u00e9es dans le GPU.<\/li>\n<li>Ex\u00e9cutez un petit noyau.<\/li>\n<li>Copiez le r\u00e9sultat sur la CPU.<\/li>\n<li>Modifiez-le sur la CPU.<\/li>\n<li>Copiez-le dans le GPU.<\/li>\n<\/ol>\n<p>Une meilleure conception maintient l&rsquo;\u00e9tat de simulation sur l&rsquo;appareil pendant de nombreuses \u00e9tapes de temps et ne transf\u00e8re que la sortie requise pour la visualisation, le point de contr\u00f4le ou l&rsquo;analyse.<\/p>\n<p>Les flux asynchrones, la m\u00e9moire h\u00f4te \u00e9pingl\u00e9e, la communication qui se chevauche et les fonctionnalit\u00e9s de m\u00e9moire unifi\u00e9e peuvent aider, mais elles ne doivent \u00eatre introduites qu&rsquo;apr\u00e8s la mesure des co\u00fbts de transfert ordinaires.<\/p>\n<h2>L&rsquo;\u00e9tude de cas XLB<\/h2>\n<p>Le projet XLB d&rsquo;Autodesk Research fournit un exemple utile de simulation GPU native Python. XLB est une biblioth\u00e8que de r\u00e9seau open source Boltzmann avec plusieurs backends informatiques, dont NVIDIA WARP.<\/p>\n<p>Dans les configurations de r\u00e9f\u00e9rence signal\u00e9es par Autodesk Research et NVIDIA, le backend Warp a atteint des performances proches de la mise en \u0153uvre comparative C++\/OpenCL FluidX3D pour un cas de cavit\u00e9 sp\u00e9cifiquement pilot\u00e9 par LID. Une comparaison distincte a rapport\u00e9 une acc\u00e9l\u00e9ration approximative de huit fois par rapport au backend JAX de XLB sur le mat\u00e9riel et les configurations s\u00e9lectionn\u00e9s.<\/p>\n<p>L&rsquo;\u00e9quipe a \u00e9galement d\u00e9montr\u00e9 une approche erron\u00e9e sur un cluster GH200 \u00e0 huit n\u0153uds avec un domaine d&rsquo;environ 50&nbsp;milliards de cellules de r\u00e9seau. Ces r\u00e9sultats s&rsquo;appliquent au solveur, au probl\u00e8me, au mat\u00e9riel et aux choix d&rsquo;impl\u00e9mentation signal\u00e9s. Ils ne doivent pas \u00eatre trait\u00e9s comme des garanties d&rsquo;acc\u00e9l\u00e9ration g\u00e9n\u00e9rales pour les noyaux Python.<\/p>\n<p>L&rsquo;\u00e9tude de cas compl\u00e8te est disponible dans <a href=\"https:\/\/developer.nvidia.com\/blog\/autodesk-research-brings-warp-speed-to-computational-fluid-dynamics-on-nvidia-gh200\/\" rel=\"nofollow\" target=\"_blank\">La recherche d&rsquo;Autodesk apporte une vitesse de distorsion \u00e0 la dynamique des fluides de calcul sur NVIDIA GH200<\/a>. Le code de projet actuel est disponible dans le <a href=\"https:\/\/github.com\/Autodesk\/XLB\" rel=\"nofollow\" target=\"_blank\">Autodesk XLB Repository<\/a>.<\/p>\n<h2>Comment comparer correctement un noyau<\/h2>\n<p>Les op\u00e9rations GPU sont g\u00e9n\u00e9ralement asynchrones. Mesure uniquement la dur\u00e9e de l&rsquo;appel de fonctions Python peut indiquer le temps n\u00e9cessaire pour mettre le noyau en file plut\u00f4t que le temps n\u00e9cessaire pour l&rsquo;ex\u00e9cuter.<\/p>\n<p>Une r\u00e9f\u00e9rence de base devrait :<\/p>\n<ol>\n<li>Ex\u00e9cutez le noyau plusieurs fois pour d\u00e9clencher la compilation et l&rsquo;\u00e9chauffement.<\/li>\n<li>Synchronisez avant de d\u00e9marrer la minuterie.<\/li>\n<li>Ex\u00e9cutez plusieurs it\u00e9rations mesur\u00e9es.<\/li>\n<li>Synchronisez avant d&rsquo;arr\u00eater la minuterie.<\/li>\n<li>Signaler la moyenne et la variation des r\u00e9p\u00e9titions.<\/li>\n<li>S\u00e9parez le temps de transfert de donn\u00e9es du temps d&rsquo;ex\u00e9cution du noyau.<\/li>\n<li>V\u00e9rifiez que les versions du CPU et du GPU produisent des r\u00e9sultats \u00e9quivalents.<\/li>\n<\/ol>\n<pre><code class=\"language-python\">import time\nfrom numba import cuda\n\n# Warm-up and JIT compilation\nkernel[blocks, threads](*arguments)\ncuda.synchronize()\n\nstart = time.perf_counter()\n\nfor _ in range(100):\n    kernel[blocks, threads](*arguments)\n\ncuda.synchronize()\n\nelapsed = time.perf_counter() - start\naverage = elapsed \/ 100\n\nprint(\"Average kernel time:\", average)<\/code><\/pre>\n<p>La comparaison doit utiliser une taille de simulation r\u00e9aliste et inclure le flux de travail complet lorsque les co\u00fbts de transfert ou de pr\u00e9traitement sont importants.<\/p>\n<p>La m\u00e9thodologie de pr\u00e9cision du travail peut \u00e9galement comparer le temps d&rsquo;ex\u00e9cution \u00e0 l&rsquo;erreur num\u00e9rique. La <a href=\"https:\/\/docs.sciml.ai\/DiffEqDevDocs\/stable\/alg_dev\/benchmarks\/\" rel=\"nofollow\" target=\"_blank\">Documentation de r\u00e9f\u00e9rence SCIML<\/a> fournit des exemples de ce style d&rsquo;\u00e9valuation.<\/p>\n<h2>Un flux de travail de d\u00e9veloppement pratique<\/h2>\n<ol>\n<li>Impl\u00e9mentez et v\u00e9rifiez une version claire de la r\u00e9f\u00e9rence CPU.<\/li>\n<li>Profilez l&rsquo;application pour localiser l&rsquo;op\u00e9ration dominante.<\/li>\n<li>Essayez une biblioth\u00e8que GPU optimis\u00e9e avant d&rsquo;\u00e9crire un noyau.<\/li>\n<li>Conservez les donn\u00e9es fr\u00e9quemment r\u00e9utilis\u00e9es sur l&rsquo;appareil.<\/li>\n<li>\u00c9crivez le noyau correct le plus simple.<\/li>\n<li>Validez-le par rapport au r\u00e9sultat du CPU.<\/li>\n<li>Mesurez les transferts de m\u00e9moire et l&rsquo;ex\u00e9cution s\u00e9par\u00e9ment.<\/li>\n<li>Inspectez la coalescence, l&rsquo;occupation, la divergence et l&rsquo;utilisation des registres.<\/li>\n<li>Testez la m\u00e9moire partag\u00e9e ou la fusion uniquement lorsque le profilage les prend en charge.<\/li>\n<li>Benchmark plusieurs tailles de probl\u00e8mes et architectures GPU.<\/li>\n<\/ol>\n<h2>Choisir un cadre<\/h2>\n<table class=\"custom-table\">\n<tbody>\n<tr>\n<th>Exigence<\/th>\n<th>Point de d\u00e9part possible<\/th>\n<\/tr>\n<tr>\n<td>Flux de travail GPU de style Numpy existant<\/td>\n<td>CUPY avec des op\u00e9rations int\u00e9gr\u00e9es ou des noyaux personnalis\u00e9s<\/td>\n<\/tr>\n<tr>\n<td>Petit nombre de noyaux python de style CUDA<\/td>\n<td>NUMBA-CUDA, apr\u00e8s avoir examin\u00e9 son \u00e9tat de support actuel<\/td>\n<\/tr>\n<tr>\n<td>Simulation, g\u00e9om\u00e9trie et noyaux diff\u00e9rentiables<\/td>\n<td>NVIDIA WARP<\/td>\n<\/tr>\n<tr>\n<td>Contr\u00f4le sp\u00e9cifique \u00e0 NVIDIA maximum<\/td>\n<td>CUDA C++<\/td>\n<\/tr>\n<tr>\n<td>Cible GPU AMD avec une source de type CUDA<\/td>\n<td>Hanche et ROCM<\/td>\n<\/tr>\n<tr>\n<td>C++ portable sur plusieurs backends<\/td>\n<td>Sycl, Kokkos ou une autre couche de portabilit\u00e9 des performances<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n<h2>Erreurs courantes de programmation du noyau<\/h2>\n<ul>\n<li>D\u00e9placement des donn\u00e9es entre le CPU et le GPU \u00e0 chaque pas de temps<\/li>\n<li>Ignorer les threads hors limites<\/li>\n<li>Utiliser un algorithme inefficace et s&rsquo;attendre \u00e0 ce que le mat\u00e9riel corrige sa mise \u00e0 l&rsquo;\u00e9chelle<\/li>\n<li>Acc\u00e9der \u00e0 la m\u00e9moire avec une mise en page non fusionn\u00e9e<\/li>\n<li>Ajouter de la m\u00e9moire partag\u00e9e sans mesurer si cela aide<\/li>\n<li>Lancement de nombreux minuscules noyaux au lieu d&rsquo;envisager la fusion<\/li>\n<li>Utilisation de registres excessifs ou de m\u00e9moire partag\u00e9e par bloc<\/li>\n<li>Analyse comparative du code asynchrone sans synchronisation<\/li>\n<li>Comparaison des sorties sans v\u00e9rification de la pr\u00e9cision num\u00e9rique<\/li>\n<li>Utilisation de revendications de performances fixes sur diff\u00e9rents GPU et charges de travail<\/li>\n<li>En supposant que la source portable offre des performances portables<\/li>\n<li>Application de la diff\u00e9renciation automatique aux op\u00e9rations sur place non prises en charge<\/li>\n<\/ul>\n<h2>Guides connexes<\/h2>\n<ul>\n<li><a href=\"https:\/\/matforge.org\/gpu-accelerated-scientific-computing-cupy-numba-cudf-compared\/\" rel=\"nofollow\" target=\"_blank\">scientifique acc\u00e9l\u00e9r\u00e9 par GPU Informatique&nbsp;: CUPY, NUMBA et CUDF compar\u00e9s<\/a> &#8211; Comparez les outils Python de haut niveau pour les charges de travail GPU.<\/li>\n<li><a href=\"https:\/\/matforge.org\/gpu-acceleration-for-fipy-simulations-cupy-and-numba-integration-guide\/\" rel=\"nofollow\" target=\"_blank\">Acc\u00e9l\u00e9ration GPU pour les simulations FIPY&nbsp;: cupy et Numba Integration Guide<\/a> \u2014 Explorez les composants GPU possibles autour d&rsquo;un flux de travail Fipy.<\/li>\n<li><a href=\"https:\/\/matforge.org\/physics-informed-neural-networks-pinns-for-scientific-simulations\/\" rel=\"nofollow\" target=\"_blank\">R\u00e9seaux de neurones \u00e0 base de physique pour les simulations scientifiques<\/a> \u2014 Passez en revue la relation entre les mod\u00e8les physiques et l&rsquo;apprentissage diff\u00e9rentiable.<\/li>\n<li><a href=\"https:\/\/matforge.org\/when-to-use-fem-fvm-fdm\/\" rel=\"nofollow\" target=\"_blank\">Quand utiliser FEM, FVM et FDM<\/a> Discr\u00e9tisation avant d&rsquo;optimiser sa mise en \u0153uvre.<\/li>\n<li><a href=\"https:\/\/matforge.org\/scientific-computing-performance-profiling-identifying-bottlenecks\/\" rel=\"nofollow\" target=\"_blank\">Profilage des performances&nbsp;: identifier les goulots d&rsquo;\u00e9tranglement dans le code scientifique<\/a>&nbsp;: localiser les limitations du processeur, de la m\u00e9moire, du transfert et des E\/S.<\/li>\n<li><a href=\"https:\/\/matforge.org\/hpc-python-workflows-from-laptop-to-supercomputer\/\" rel=\"nofollow\" target=\"_blank\">Flows Python HPC&nbsp;: depuis un ordinateur portable au superordinateur<\/a> \u2014 Planifiez la progression du d\u00e9veloppement local vers les grands syst\u00e8mes.<\/li>\n<\/ul>\n<h2>Conclusion<\/h2>\n<p>Les noyaux GPU personnalis\u00e9s permettent aux d\u00e9veloppeurs scientifiques de traduire directement des particules, des cellules de la grille, des r\u00e8gles de mat\u00e9riaux et d&rsquo;autres op\u00e9rations physiques en fonctions massivement parall\u00e8les. Ils permettent de contr\u00f4ler l&rsquo;indexation, l&rsquo;acc\u00e8s \u00e0 la m\u00e9moire, la synchronisation, la fusion du noyau et l&rsquo;optimisation sp\u00e9cifique \u00e0 l&rsquo;appareil.<\/p>\n<p>CUPY propose \u00e0 la fois des op\u00e9rations de baies de haut niveau et des interfaces de noyau personnalis\u00e9es. Numba-CUDA fournit un mod\u00e8le de type CUDA en Python, tandis que NVIDIA WARP ajoute des primitives de simulation, des op\u00e9rations de tuiles et une diff\u00e9renciation automatique. CUDA et HIP offrent un contr\u00f4le C++ de niveau inf\u00e9rieur, et les cadres portables de performances prennent en charge des objectifs mat\u00e9riels plus larges.<\/p>\n<p>Les gains de performances les plus importants proviennent rarement de la modification de la syntaxe seule. Ils proviennent du choix d&rsquo;un algorithme parall\u00e8le, de la conservation des donn\u00e9es sur l&rsquo;appareil, de la r\u00e9duction du trafic m\u00e9moire, de l&rsquo;utilisation d&rsquo;une disposition de donn\u00e9es appropri\u00e9e et de l&rsquo;\u00e9limination des op\u00e9rations interm\u00e9diaires inutiles.<\/p>\n<p>Un noyau personnalis\u00e9 ne doit \u00eatre d\u00e9velopp\u00e9 qu&rsquo;apr\u00e8s le profilage identifiant un v\u00e9ritable goulot d&rsquo;\u00e9tranglement. Il doit \u00eatre valid\u00e9 par rapport \u00e0 une solution de r\u00e9f\u00e9rence et \u00e9talonn\u00e9 avec la synchronisation, la taille repr\u00e9sentative des probl\u00e8mes et la comptabilisation compl\u00e8te des co\u00fbts de transfert de m\u00e9moire.<\/p>\n<p>Lorsque ces conditions sont remplies, les outils de noyau bas\u00e9s sur Python peuvent prendre en charge la simulation physique s\u00e9rieuse sans obliger les chercheurs \u00e0 d\u00e9placer chaque partie de l&rsquo;application en C++ de bas niveau.<\/p>\n<\/article>\n","protected":false,"raw":"<article>\n<p>L'acc\u00e9l\u00e9ration des GPU en informatique scientifique commence souvent par des biblioth\u00e8ques de haut niveau. Un chercheur remplace NumPy par CUPY, utilise une FFT acc\u00e9l\u00e9r\u00e9e ou appelle une routine d'alg\u00e8bre lin\u00e9aire activ\u00e9e par GPU. Cela peut produire une am\u00e9lioration substantielle sans n\u00e9cessiter de connaissances d\u00e9taill\u00e9es sur le mat\u00e9riel GPU.<\/p>\n<p>La programmation du noyau GPU personnalis\u00e9 va d'un niveau plus profond\u00e9ment. Au lieu d'appeler une op\u00e9ration pr\u00e9d\u00e9finie, le d\u00e9veloppeur \u00e9crit la fonction que chaque thread GPU ex\u00e9cute. Cela permet de contr\u00f4ler l'indexation des threads, la disposition des donn\u00e9es, les transferts de m\u00e9moire, la synchronisation, la m\u00e9moire partag\u00e9e et le nombre d'op\u00e9rations effectu\u00e9es lors de chaque lancement du noyau.<\/p>\n<p>Ce contr\u00f4le est utile lorsqu'un mod\u00e8le physique contient des mod\u00e8les d'acc\u00e8s qui ne peuvent pas \u00eatre exprim\u00e9s efficacement par le biais d'op\u00e9rations de tableau ordinaires. Les interactions de particules, les pochoirs de diff\u00e9rences finies, les r\u00e8gles de collision, les mod\u00e8les de fluides bas\u00e9s sur le r\u00e9seau, les mises \u00e0 jour des mat\u00e9riaux et les requ\u00eates de g\u00e9om\u00e9trie sont des exemples courants.<\/p>\n<p>Un noyau personnalis\u00e9 n'est pas automatiquement plus rapide qu'une biblioth\u00e8que optimis\u00e9e. Cela devient utile lorsque l'algorithme expose suffisamment de travail parall\u00e8le et lorsque le d\u00e9veloppeur peut r\u00e9duire le trafic m\u00e9moire, combiner les op\u00e9rations ou adapter le calcul \u00e0 l'architecture GPU.<\/p>\n<h2>Qu'est-ce qu'un noyau GPU&nbsp;?<\/h2>\n<p>Un noyau GPU est une fonction ex\u00e9cut\u00e9e par de nombreux threads l\u00e9gers. Chaque thread g\u00e8re normalement une particule, une cellule de grille, un \u00e9l\u00e9ment, une face, un pixel ou une entr\u00e9e dans un tableau.<\/p>\n<p>Le programme h\u00f4te lance le noyau avec une grille de blocs de thread :<\/p>\n<pre><code>kernel[blocks, threads_per_block](arguments)<\/code><\/pre>\n<p>Chaque thread d\u00e9termine les donn\u00e9es qu'il doit traiter \u00e0 partir de ses indices de blocs et de threads. Un index unidimensionnel typique est :<\/p>\n<pre><code>index =\n    block_index * block_size\n    + thread_index<\/code><\/pre>\n<p>Les threads sont organis\u00e9s en groupes appel\u00e9s blocs. Les threads \u00e0 l'int\u00e9rieur d'un seul bloc peuvent synchroniser et \u00e9changer des donn\u00e9es via la m\u00e9moire partag\u00e9e. Les blocs s\u00e9par\u00e9s doivent g\u00e9n\u00e9ralement s'ex\u00e9cuter de mani\u00e8re ind\u00e9pendante.<\/p>\n<h2>Lorsque les noyaux personnalis\u00e9s valent la peine d'\u00eatre \u00e9crits<\/h2>\n<p>Les biblioth\u00e8ques de GPU de haut niveau devraient g\u00e9n\u00e9ralement \u00eatre la premi\u00e8re option. Ils fournissent d\u00e9j\u00e0 une multiplication matricielle optimis\u00e9e, des FFT, des r\u00e9ductions, des op\u00e9rations rares, une g\u00e9n\u00e9ration de nombres al\u00e9atoires et de nombreuses fonctions par \u00e9l\u00e9ment.<\/p>\n<p>Un noyau personnalis\u00e9 devient utile lorsque :<\/p>\n<ul>\n<li>La simulation utilise un mod\u00e8le d'acc\u00e8s \u00e0 la m\u00e9moire non standard.<\/li>\n<li>Plusieurs petites op\u00e9rations peuvent \u00eatre fusionn\u00e9es en un seul passage sur la m\u00e9moire.<\/li>\n<li>Les threads voisins ont \u00e0 plusieurs reprises besoin des m\u00eames donn\u00e9es locales.<\/li>\n<li>Une r\u00e8gle sp\u00e9cialis\u00e9e de particule, de collision ou de pochoir domine le temps d'ex\u00e9cution.<\/li>\n<li>Les baies interm\u00e9diaires consomment trop de m\u00e9moire de p\u00e9riph\u00e9rique.<\/li>\n<li>La simulation n\u00e9cessite une op\u00e9ration diff\u00e9renciable personnalis\u00e9e.<\/li>\n<li>Une biblioth\u00e8que existante ne peut pas exprimer efficacement la limite ou la logique mat\u00e9rielle requise.<\/li>\n<\/ul>\n<p>Avant d'\u00e9crire un noyau, profilez le programme. L'optimisation d'une fonction visuellement complexe qui ne repr\u00e9sente qu'une petite partie du temps d'ex\u00e9cution ne produira pas une acc\u00e9l\u00e9ration globale significative.<\/p>\n<h2>Mod\u00e8les de programmation GPU<\/h2>\n<p>Les logiciels scientifiques peuvent acc\u00e9der aux GPU via plusieurs niveaux d'abstraction.<\/p>\n<table class=\"custom-table\">\n<tbody><tr>\n<th>Approche<\/th>\n<th>Exemples<\/th>\n<th>force principale<\/th>\n<th>Principal compromis<\/th>\n<\/tr>\n<tr>\n<td>D\u00e9chargement bas\u00e9 sur les directives<\/td>\n<td>OpenMP, OpenACC<\/td>\n<td>Acc\u00e9l\u00e9ration incr\u00e9mentale du code C, C++ ou Fortran existant<\/td>\n<td>Moins de contr\u00f4le direct sur les noyaux g\u00e9n\u00e9r\u00e9s<\/td>\n<\/tr>\n<tr>\n<td>Programmation du noyau orient\u00e9e fournisseur<\/td>\n<td>cud\u00e9, hanche<\/td>\n<td>Contr\u00f4le d\u00e9taill\u00e9 de l'ex\u00e9cution du GPU<\/td>\n<td>Plus de co\u00fbts de mise en \u0153uvre et de maintenance<\/td>\n<\/tr>\n<tr>\n<td>C++ portable de performances<\/td>\n<td>SYCL, Kokkos, Alpaka<\/td>\n<td>Un mod\u00e8le source pour plusieurs backends mat\u00e9riels<\/td>\n<td>La portabilit\u00e9 ne garantit pas une performance \u00e9gale partout<\/td>\n<\/tr>\n<tr>\n<td>Tableaux et noyaux de GPU Python<\/td>\n<td>CUPY, NUMBA-CUDA, NVIDIA WARP<\/td>\n<td>D\u00e9veloppement rapide avec acc\u00e8s au code GPU compil\u00e9<\/td>\n<td>Restrictions sp\u00e9cifiques au framework et comportement de compilation<\/td>\n<\/tr>\n<\/tbody><\/table>\n<p><a href=\"https:\/\/enccs.github.io\/gpu-programming\/5-intro-to-gpu-prog-models\/\" rel=\"nofollow\" target=\"_blank\">ENCCS Introduction \u00e0 la programmation GPU Les mod\u00e8les<\/a> fournissent une comparaison plus large des approches CUDA, HIP, OpenCL, SYCL, Kokkos, OpenMP et associ\u00e9es.<\/p>\n<h2>Cuda<\/h2>\n<p>CUDA fournit un acc\u00e8s direct \u00e0 la programmation NVIDIA GPU via C++, Fortran et plusieurs liaisons de langage. Il expose les grilles, les blocs, les d\u00e9formations, la m\u00e9moire partag\u00e9e, les flux, les \u00e9v\u00e9nements, les biblioth\u00e8ques de p\u00e9riph\u00e9riques et les outils d'optimisation sp\u00e9cifiques au mat\u00e9riel.<\/p>\n<p>CUDA est souvent s\u00e9lectionn\u00e9 lorsque des performances maximales sp\u00e9cifiques \u00e0 NVIDIA et des outils matures sont plus importants que la portabilit\u00e9 mat\u00e9rielle. Les principaux co\u00fbts sont la gestion de la m\u00e9moire manuelle, un mod\u00e8le de programmation de niveau inf\u00e9rieur et la d\u00e9pendance \u00e0 la plate-forme NVIDIA.<\/p>\n<h2>Hanche et ROCM<\/h2>\n<p>HIP fournit un mod\u00e8le de programmation C++ de type CUDA dans l'\u00e9cosyst\u00e8me ROCM d'AMD. Une grande quantit\u00e9 de code de style CUDA peut \u00eatre adapt\u00e9e \u00e0 la hanche, mais la portabilit\u00e9 de la source ne signifie pas qu'un binaire s'ex\u00e9cute inchang\u00e9 sur chaque GPU.<\/p>\n<p>Une compilation, des tests et un r\u00e9glage des performances s\u00e9par\u00e9s peuvent toujours \u00eatre n\u00e9cessaires pour chaque architecture cible. Les d\u00e9tails sont disponibles dans la <a href=\"https:\/\/rocm.docs.amd.com\/projects\/HIP\/en\/latest\/what_is_hip.html\" rel=\"nofollow\" target=\"_blank\">documentation officielle de la hanche<\/a>.<\/p>\n<p>HIP est une option pratique lorsque les GPU AMD doivent \u00eatre pris en charge ou lorsqu'un projet souhaite r\u00e9duire la d\u00e9pendance \u00e0 un seul fournisseur d'acc\u00e9l\u00e9rateurs tout en conservant un style de programmation de type CUDA.<\/p>\n<h2>Cupy Cernels personnalis\u00e9s<\/h2>\n<p>Cupy est surtout connu pour ses tableaux compatibles avec NumPy et ses fonctions num\u00e9riques acc\u00e9l\u00e9r\u00e9es. Il prend \u00e9galement en charge la programmation GPU personnalis\u00e9e.<\/p>\n<p>Les d\u00e9veloppeurs peuvent utiliser :<\/p>\n<ul>\n<li><code>ElementwiseKernel<\/code> pour les expressions personnalis\u00e9es par \u00e9l\u00e9ment<\/li>\n<li><code>ReductionKernel<\/code> pour les r\u00e9ductions personnalis\u00e9es<\/li>\n<li><code>RawKernel<\/code> pour les noyaux \u00e9crits en CUDA C++<\/li>\n<li><code>cupyx.jit.rawkernel<\/code> pour les d\u00e9finitions du noyau JIT de style Python<\/li>\n<\/ul>\n<p>Un simple noyau brut Cupy peut \u00eatre d\u00e9fini comme suit :<\/p>\n<pre><code class=\"language-python\">import cupy as cp\n\nscale_kernel = cp.RawKernel(\n    r'''\n    extern \"C\" __global__\n    void scale_values(\n        const float* input,\n        float* output,\n        const float factor,\n        const int size\n    ) {\n        int index =\n            blockDim.x * blockIdx.x\n            + threadIdx.x;\n\n        if (index &lt; size) {\n            output[index] =\n                factor * input[index];\n        }\n    }\n    ''',\n    \"scale_values\"\n)\n\nsize = 1_000_000\nthreads_per_block = 256\nblocks = (\n    size + threads_per_block - 1\n) \/\/ threads_per_block\n\ninput_values = cp.arange(\n    size,\n    dtype=cp.float32\n)\n\noutput_values = cp.empty_like(\n    input_values\n)\n\nscale_kernel(\n    (blocks,),\n    (threads_per_block,),\n    (\n        input_values,\n        output_values,\n        cp.float32(2.0),\n        cp.int32(size)\n    )\n)<\/code><\/pre>\n<p>CUPY est utile lorsque la plupart de l'application utilise d\u00e9j\u00e0 des baies de GPU et que seules les op\u00e9rations s\u00e9lectionn\u00e9es ont besoin de noyaux personnalis\u00e9s.<\/p>\n<h2>Numba-cuda<\/h2>\n<p>numba-cuda compile un sous-ensemble restreint de python dans les noyaux GPU. Son mod\u00e8le de programmation suit de pr\u00e8s CUDA C. Les d\u00e9veloppeurs d\u00e9finissent une fonction avec un d\u00e9corateur de noyau, calculent un index de fil et lancent la fonction avec une configuration de grille et de bloc.<\/p>\n<p>La documentation active du projet est disponible \u00e0 <a href=\"https:\/\/nvidia.github.io\/numba-cuda\/\" rel=\"nofollow\" target=\"_blank\">numba-cuda<\/a>. \u00c9tant donn\u00e9 que son statut de d\u00e9veloppement peut changer, les \u00e9quipes cr\u00e9ant des logiciels de longue dur\u00e9e doivent revoir les directives de maintenance et de migration actuelles avant de s'engager dans le cadre.<\/p>\n<h2>Un noyau de particules Numba<\/h2>\n<p>L'exemple \u00e9ducatif suivant calcule l'acc\u00e9l\u00e9ration gravitationnelle avec des interactions directes toutes paires :<\/p>\n<pre><code class=\"language-python\">import math\nimport numpy as np\nfrom numba import cuda\n\n@cuda.jit\ndef gravity_kernel(\n    positions,\n    masses,\n    accelerations,\n    gravitational_constant,\n    softening_squared\n):\n    particle = cuda.grid(1)\n    particle_count = positions.shape[0]\n\n    if particle &gt;= particle_count:\n        return\n\n    px = positions[particle, 0]\n    py = positions[particle, 1]\n    pz = positions[particle, 2]\n\n    ax = 0.0\n    ay = 0.0\n    az = 0.0\n\n    for other in range(particle_count):\n        if other == particle:\n            continue\n\n        dx = positions[other, 0] - px\n        dy = positions[other, 1] - py\n        dz = positions[other, 2] - pz\n\n        distance_squared = (\n            dx * dx\n            + dy * dy\n            + dz * dz\n            + softening_squared\n        )\n\n        inverse_distance = (\n            1.0\n            \/ math.sqrt(distance_squared)\n        )\n\n        inverse_distance_cubed = (\n            inverse_distance\n            * inverse_distance\n            * inverse_distance\n        )\n\n        scale = (\n            gravitational_constant\n            * masses[other]\n            * inverse_distance_cubed\n        )\n\n        ax += scale * dx\n        ay += scale * dy\n        az += scale * dz\n\n    accelerations[particle, 0] = ax\n    accelerations[particle, 1] = ay\n    accelerations[particle, 2] = az\n\n\nparticle_count = 10_000\nthreads_per_block = 256\n\nblocks = (\n    particle_count\n    + threads_per_block\n    - 1\n) \/\/ threads_per_block\n\nhost_positions = np.random.random(\n    (particle_count, 3)\n).astype(np.float32)\n\nhost_masses = np.ones(\n    particle_count,\n    dtype=np.float32\n)\n\ndevice_positions = cuda.to_device(\n    host_positions\n)\n\ndevice_masses = cuda.to_device(\n    host_masses\n)\n\ndevice_accelerations = cuda.device_array(\n    (particle_count, 3),\n    dtype=np.float32\n)\n\ngravity_kernel[\n    blocks,\n    threads_per_block\n](\n    device_positions,\n    device_masses,\n    device_accelerations,\n    np.float32(1.0),\n    np.float32(1e-4)\n)\n\ncuda.synchronize()\n\naccelerations = (\n    device_accelerations.copy_to_host()\n)<\/code><\/pre>\n<p>Ce noyau a une complexit\u00e9 de calcul de <code>O(N\u00b2)<\/code>. Il est adapt\u00e9 pour expliquer la cartographie des threads, mais ce n'est pas un algorithme de production efficace pour les tr\u00e8s gros nombres de particules.<\/p>\n<p>Les grands syst\u00e8mes de gravitation ou de particules n\u00e9cessitent souvent une m\u00e9thode d'arbre, une m\u00e9thode multip\u00f4le rapide, des bacs spatiaux ou des listes de voisins. Un GPU ne peut pas supprimer le probl\u00e8me de mise \u00e0 l'\u00e9chelle d'un algorithme inefficace.<\/p>\n<h2>NVIDIA WARP<\/h2>\n<p>NVIDIA WARP permet aux d\u00e9veloppeurs de d\u00e9finir des noyaux fortement typ\u00e9s avec la syntaxe Python. WARP compile ces fonctions en CPU ou CUDA et fournit des primitives sp\u00e9cialis\u00e9es pour la g\u00e9om\u00e9trie, la simulation, les op\u00e9rations rares, les \u00e9l\u00e9ments finis, l'optimisation et la diff\u00e9renciation automatique.<\/p>\n<p>La <a href=\"https:\/\/developer.nvidia.com\/warp-python\" rel=\"nofollow\" target=\"_blank\">page produit NVIDIA WARP<\/a> et <a href=\"https:\/\/github.com\/nvidia\/warp\" rel=\"nofollow\" target=\"_blank\">R\u00e9f\u00e9rentiel GitHub de Warp<\/a> Inclut des exemples de particules, de fluides, de maillages, d'optimisation, de simulation diff\u00e9renciable et Programmation GPU bas\u00e9e sur des tuiles.<\/p>\n<h2>Le m\u00eame noyau de particules dans WARP<\/h2>\n<pre><code class=\"language-python\">import numpy as np\nimport warp as wp\n\nwp.init()\n\n@wp.kernel\ndef gravity_kernel(\n    positions: wp.array(dtype=wp.vec3),\n    masses: wp.array(dtype=wp.float32),\n    accelerations: wp.array(dtype=wp.vec3),\n    gravitational_constant: wp.float32,\n    softening_squared: wp.float32\n):\n    particle = wp.tid()\n    particle_count = positions.shape[0]\n\n    position = positions[particle]\n    acceleration = wp.vec3(0.0, 0.0, 0.0)\n\n    for other in range(particle_count):\n        if other != particle:\n            displacement = (\n                positions[other]\n                - position\n            )\n\n            distance_squared = (\n                wp.dot(\n                    displacement,\n                    displacement\n                )\n                + softening_squared\n            )\n\n            inverse_distance = (\n                1.0\n                \/ wp.sqrt(distance_squared)\n            )\n\n            inverse_distance_cubed = (\n                inverse_distance\n                * inverse_distance\n                * inverse_distance\n            )\n\n            acceleration += (\n                gravitational_constant\n                * masses[other]\n                * inverse_distance_cubed\n                * displacement\n            )\n\n    accelerations[particle] = acceleration\n\n\nparticle_count = 10_000\n\nhost_positions = np.random.random(\n    (particle_count, 3)\n).astype(np.float32)\n\nhost_masses = np.ones(\n    particle_count,\n    dtype=np.float32\n)\n\npositions = wp.array(\n    host_positions,\n    dtype=wp.vec3,\n    device=\"cuda\"\n)\n\nmasses = wp.array(\n    host_masses,\n    dtype=wp.float32,\n    device=\"cuda\"\n)\n\naccelerations = wp.zeros(\n    particle_count,\n    dtype=wp.vec3,\n    device=\"cuda\"\n)\n\nwp.launch(\n    kernel=gravity_kernel,\n    dim=particle_count,\n    inputs=[\n        positions,\n        masses,\n        accelerations,\n        wp.float32(1.0),\n        wp.float32(1e-4)\n    ],\n    device=\"cuda\"\n)\n\nwp.synchronize()<\/code><\/pre>\n<p>WARP d\u00e9termine en interne une configuration de lancement appropri\u00e9e pour un lancement de noyau ordinaire. Les utilisateurs avanc\u00e9s peuvent toujours travailler avec des dimensions de blocs, des graphiques de commandes, une ex\u00e9cution en mosa\u00efque et des op\u00e9rations de p\u00e9riph\u00e9rique sp\u00e9cialis\u00e9es si n\u00e9cessaire.<\/p>\n<h2>Boucles de bordure de grille<\/h2>\n<p>Un noyau n'a pas besoin d'un thread assign\u00e9 de fa\u00e7on permanente pour chaque \u00e9l\u00e9ment. Une boucle Grid-Stride permet \u00e0 chaque thread de traiter plusieurs entr\u00e9es&nbsp;:<\/p>\n<pre><code class=\"language-python\">from numba import cuda\n\n@cuda.jit\ndef scale_with_stride(\n    input_values,\n    output_values,\n    factor\n):\n    index = cuda.grid(1)\n    stride = cuda.gridsize(1)\n\n    for position in range(\n        index,\n        input_values.size,\n        stride\n    ):\n        output_values[position] = (\n            factor\n            * input_values[position]\n        )<\/code><\/pre>\n<p>Ce mod\u00e8le s\u00e9pare le nombre de threads lanc\u00e9s de la taille totale des donn\u00e9es. Il est utile lors du traitement de tr\u00e8s grands tableaux ou de la r\u00e9utilisation d'une configuration de lancement fixe.<\/p>\n<p>Une introduction pratique \u00e0 ce mod\u00e8le est disponible dans <a href=\"https:\/\/thedatafrog.com\/en\/articles\/cuda-kernel-python\/\" rel=\"nofollow\" target=\"_blank\">le Numba et le CUDA de la grenouille de donn\u00e9es Tutoriel<\/a>.<\/p>\n<h2>Coalition de m\u00e9moire<\/h2>\n<p>Les noyaux GPU sont souvent limit\u00e9s par la bande passante de m\u00e9moire plut\u00f4t que par le d\u00e9bit arithm\u00e9tique. Les threads \u00e0 l'int\u00e9rieur d'un distorsion doivent id\u00e9alement acc\u00e9der aux adresses m\u00e9moire \u00e0 proximit\u00e9 afin que le mat\u00e9riel puisse combiner leurs demandes.<\/p>\n<p>Consid\u00e9rez les donn\u00e9es de particules stock\u00e9es comme&nbsp;:<\/p>\n<pre><code>particle_0: x, y, z, mass\nparticle_1: x, y, z, mass\nparticle_2: x, y, z, mass<\/code><\/pre>\n<p>Cette disposition de tableau de structures peut \u00eatre pratique pour le code orient\u00e9 objet. Une mise en page de la structure des tableaux stocke des tableaux contigus distincts&nbsp;:<\/p>\n<pre><code>x_positions[]\ny_positions[]\nz_positions[]\nmasses[]<\/code><\/pre>\n<p>La deuxi\u00e8me mise en page peut fournir une meilleure coalescence lorsque chaque thread lit le m\u00eame champ pour une particule diff\u00e9rente. La meilleure mise en page d\u00e9pend toujours des champs accessibles ensemble.<\/p>\n<p>Les types de vecteurs ne corrigent pas automatiquement un mauvais mod\u00e8le d'acc\u00e8s. Les d\u00e9veloppeurs doivent inspecter les adresses r\u00e9elles demand\u00e9es par les threads voisins.<\/p>\n<h2>M\u00e9moire partag\u00e9e<\/h2>\n<p>La m\u00e9moire partag\u00e9e est une petite zone de m\u00e9moire \u00e0 faible latence accessible par les threads dans le m\u00eame bloc. Il peut r\u00e9duire les lectures r\u00e9p\u00e9t\u00e9es de la m\u00e9moire globale.<\/p>\n<p>Un pochoir unidimensionnel peut charger un bloc de valeurs et son halo dans la m\u00e9moire partag\u00e9e :<\/p>\n<pre><code class=\"language-python\">import numpy as np\nfrom numba import cuda, float32\n\nBLOCK_SIZE = 256\n\n@cuda.jit\ndef three_point_stencil(\n    input_values,\n    output_values\n):\n    shared = cuda.shared.array(\n        shape=BLOCK_SIZE + 2,\n        dtype=float32\n    )\n\n    local_index = cuda.threadIdx.x\n    global_index = cuda.grid(1)\n    size = input_values.size\n\n    center = local_index + 1\n\n    if global_index &lt; size:\n        shared[center] = (\n            input_values[global_index]\n        )\n    else:\n        shared[center] = 0.0\n\n    if local_index == 0:\n        left_index = global_index - 1\n\n        shared[0] = (\n            input_values[left_index]\n            if left_index &gt;= 0\n            else 0.0\n        )\n\n    if local_index == BLOCK_SIZE - 1:\n        right_index = global_index + 1\n\n        shared[BLOCK_SIZE + 1] = (\n            input_values[right_index]\n            if right_index &lt; size\n            else 0.0\n        )\n\n    cuda.syncthreads()\n\n    if global_index &lt; size:\n        output_values[global_index] = (\n            shared[center - 1]\n            + shared[center]\n            + shared[center + 1]\n        ) \/ 3.0<\/code><\/pre>\n<p>L'appel \u00e0 <code>cuda.syncthreads()<\/code> garantit que chaque thread finit de charger ses valeurs avant le d\u00e9but du calcul du pochoir.<\/p>\n<p>La m\u00e9moire partag\u00e9e ne doit pas \u00eatre utilis\u00e9e automatiquement. Une allocation de m\u00e9moire partag\u00e9e excessive peut r\u00e9duire l'occupation, augmenter les co\u00fbts de synchronisation et ralentir le noyau. Le profilage est requis.<\/p>\n<h2>Fusion du noyau<\/h2>\n<p>Des op\u00e9rations de tableaux distinctes cr\u00e9ent souvent plusieurs lancements de noyau et des tableaux interm\u00e9diaires&nbsp;:<\/p>\n<pre><code>velocity += dt * acceleration\nposition += dt * velocity\nenergy = compute_energy(position, velocity)<\/code><\/pre>\n<p>Un noyau fusionn\u00e9 peut calculer les trois mises \u00e0 jour alors que les valeurs n\u00e9cessaires restent dans les registres. Cela r\u00e9duit les frais g\u00e9n\u00e9raux de lancement et le trafic de m\u00e9moire globale.<\/p>\n<p>La fusion est plus avantageuse lorsque les op\u00e9rations sont simples et li\u00e9es \u00e0 la m\u00e9moire. Fusionner trop de travail peut augmenter l'utilisation des registres, r\u00e9duire l'occupation et rendre le noyau difficile \u00e0 entretenir.<\/p>\n<h2>Choisir la taille du bloc<\/h2>\n<p>Des valeurs telles que 128 ou 256 threads par bloc sont des points de d\u00e9part raisonnables, et non des optima universels.<\/p>\n<p>La meilleure taille de bloc d\u00e9pend de :<\/p>\n<ul>\n<li>Registres utilis\u00e9s par fil<\/li>\n<li>M\u00e9moire partag\u00e9e utilis\u00e9e par bloc<\/li>\n<li>Divergence de branche<\/li>\n<li>m\u00e9lange d'instructions<\/li>\n<li>Comportement d'acc\u00e8s \u00e0 la m\u00e9moire<\/li>\n<li>L'architecture GPU cible<\/li>\n<\/ul>\n<p>Une r\u00e8gle telle que le lancement d'au moins deux fois plus de blocs que les multiprocesseurs de streaming peut \u00eatre une exp\u00e9rience initiale utile, mais elle ne garantit pas des performances maximales. Les calculatrices d'occupation et les outils de profilage doivent guider la configuration finale.<\/p>\n<h2>Physique diff\u00e9rentiable<\/h2>\n<p>La simulation diff\u00e9rentiable calcule comment une sortie change par rapport aux entr\u00e9es telles que les propri\u00e9t\u00e9s des mat\u00e9riaux, les forces, la g\u00e9om\u00e9trie ou les conditions initiales.<\/p>\n<p>WARP peut enregistrer les op\u00e9rations du noyau pris en charge et ex\u00e9cuter la diff\u00e9renciation automatique en mode inverse. Cela peut \u00eatre utilis\u00e9 pour les probl\u00e8mes inverses, l'optimisation de la conception, l'estimation des param\u00e8tres et l'int\u00e9gration avec les flux de travail d'apprentissage automatique.<\/p>\n<p>Un motif simplifi\u00e9 utilise une bande :<\/p>\n<pre><code class=\"language-python\">with wp.Tape() as tape:\n    wp.launch(\n        kernel=simulation_kernel,\n        dim=element_count,\n        inputs=[state, parameters],\n        outputs=[result],\n        device=\"cuda\"\n    )\n\n    wp.launch(\n        kernel=loss_kernel,\n        dim=element_count,\n        inputs=[result, target],\n        outputs=[loss],\n        device=\"cuda\"\n    )\n\ntape.backward(loss)<\/code><\/pre>\n<p>La diff\u00e9renciation automatique n'est pas garantie pour chaque noyau. Les \u00e9crasements sur place, les atomes non d\u00e9terministes, le code natif externe, la logique discontinue et les op\u00e9rations non pris en charge peuvent n\u00e9cessiter une reformulation ou des gradients personnalis\u00e9s.<\/p>\n<p>Cette fonctionnalit\u00e9 est directement connect\u00e9e \u00e0 des applications telles que l'optimisation des associations et les <a href=\"https:\/\/matforge.org\/physics-informed-neural-networks-pinns-for-scientific-simulations\/\" rel=\"nofollow\" target=\"_blank\">r\u00e9seaux de neurones \u00e0 base de physique<\/a>.<\/p>\n<h2>Lorsqu'un GPU n'est peut-\u00eatre pas plus rapide<\/h2>\n<p>Il n'y a pas de seuil universel de taille de probl\u00e8me auquel un GPU devient plus rapide qu'un processeur. Le crossover d\u00e9pend du mat\u00e9riel, de la pr\u00e9cision, du mouvement des donn\u00e9es, de la structure des algorithmes, de la qualit\u00e9 du compilateur et de la fr\u00e9quence \u00e0 laquelle les m\u00eames donn\u00e9es de r\u00e9sidents sont r\u00e9utilis\u00e9es.<\/p>\n<table class=\"custom-table\">\n<tbody><tr>\n<th>La charge de travail<\/th>\n<th>aptitude au GPU<\/th>\n<th>consid\u00e9ration principale<\/th>\n<\/tr>\n<tr>\n<td>Mise \u00e0 jour des grandes particules<\/td>\n<td>Souvent fort<\/td>\n<td>De nombreuses op\u00e9rations ind\u00e9pendantes similaires<\/td>\n<\/tr>\n<tr>\n<td>Grand pochoir r\u00e9gulier<\/td>\n<td>Souvent fort<\/td>\n<td>Acc\u00e8s \u00e0 la m\u00e9moire parall\u00e8le pr\u00e9visible<\/td>\n<\/tr>\n<tr>\n<td>Alg\u00e8bre lin\u00e9aire dense<\/td>\n<td>Fort lors de l'utilisation de biblioth\u00e8ques optimis\u00e9es<\/td>\n<td>Intensit\u00e9 arithm\u00e9tique \u00e9lev\u00e9e<\/td>\n<\/tr>\n<tr>\n<td>Petite simulation<\/td>\n<td>d\u00e9pendant du mat\u00e9riel<\/td>\n<td>Les frais g\u00e9n\u00e9raux de lancement et de transfert peuvent dominer<\/td>\n<\/tr>\n<tr>\n<td>Travers\u00e9e de graphes irr\u00e9guli\u00e8res<\/td>\n<td>Mixte<\/td>\n<td>Divergence et acc\u00e8s m\u00e9moire impr\u00e9visible<\/td>\n<\/tr>\n<tr>\n<td>Flux de travail li\u00e9 aux E\/S<\/td>\n<td>g\u00e9n\u00e9ralement limit\u00e9<\/td>\n<td>Le GPU ne peut pas supprimer un goulot d'\u00e9tranglement de stockage ou de r\u00e9seau<\/td>\n<\/tr>\n<tr>\n<td>Algorithme fortement s\u00e9quentiel<\/td>\n<td>g\u00e9n\u00e9ralement faible<\/td>\n<td>Travail ind\u00e9pendant insuffisant<\/td>\n<\/tr>\n<\/tbody><\/table>\n<p>Le flux de travail correct consiste \u00e0 mesurer la version du processeur, la version de la biblioth\u00e8que GPU et la version du noyau personnalis\u00e9 avec des donn\u00e9es repr\u00e9sentatives. Le <a href=\"https:\/\/matforge.org\/scientific-computing-performance-profiling-identifying-bottlenecks\/\" rel=\"nofollow\" target=\"_blank\">Guide de profilage de performances de calcul scientifique<\/a> explique comment identifier le v\u00e9ritable goulot d'\u00e9tranglement avant de l'optimiser.<\/p>\n<h2>\u00c9viter les frais g\u00e9n\u00e9raux de transfert<\/h2>\n<p>Les transferts r\u00e9p\u00e9t\u00e9s entre l'h\u00f4te et la m\u00e9moire de l'appareil peuvent \u00e9liminer les avantages du calcul du GPU.<\/p>\n<p>Une boucle inefficace peut suivre ce mod\u00e8le :<\/p>\n<ol>\n<li>Copiez les donn\u00e9es dans le GPU.<\/li>\n<li>Ex\u00e9cutez un petit noyau.<\/li>\n<li>Copiez le r\u00e9sultat sur la CPU.<\/li>\n<li>Modifiez-le sur la CPU.<\/li>\n<li>Copiez-le dans le GPU.<\/li>\n<\/ol>\n<p>Une meilleure conception maintient l'\u00e9tat de simulation sur l'appareil pendant de nombreuses \u00e9tapes de temps et ne transf\u00e8re que la sortie requise pour la visualisation, le point de contr\u00f4le ou l'analyse.<\/p>\n<p>Les flux asynchrones, la m\u00e9moire h\u00f4te \u00e9pingl\u00e9e, la communication qui se chevauche et les fonctionnalit\u00e9s de m\u00e9moire unifi\u00e9e peuvent aider, mais elles ne doivent \u00eatre introduites qu'apr\u00e8s la mesure des co\u00fbts de transfert ordinaires.<\/p>\n<h2>L'\u00e9tude de cas XLB<\/h2>\n<p>Le projet XLB d'Autodesk Research fournit un exemple utile de simulation GPU native Python. XLB est une biblioth\u00e8que de r\u00e9seau open source Boltzmann avec plusieurs backends informatiques, dont NVIDIA WARP.<\/p>\n<p>Dans les configurations de r\u00e9f\u00e9rence signal\u00e9es par Autodesk Research et NVIDIA, le backend Warp a atteint des performances proches de la mise en \u0153uvre comparative C++\/OpenCL FluidX3D pour un cas de cavit\u00e9 sp\u00e9cifiquement pilot\u00e9 par LID. Une comparaison distincte a rapport\u00e9 une acc\u00e9l\u00e9ration approximative de huit fois par rapport au backend JAX de XLB sur le mat\u00e9riel et les configurations s\u00e9lectionn\u00e9s.<\/p>\n<p>L'\u00e9quipe a \u00e9galement d\u00e9montr\u00e9 une approche erron\u00e9e sur un cluster GH200 \u00e0 huit n\u0153uds avec un domaine d'environ 50&nbsp;milliards de cellules de r\u00e9seau. Ces r\u00e9sultats s'appliquent au solveur, au probl\u00e8me, au mat\u00e9riel et aux choix d'impl\u00e9mentation signal\u00e9s. Ils ne doivent pas \u00eatre trait\u00e9s comme des garanties d'acc\u00e9l\u00e9ration g\u00e9n\u00e9rales pour les noyaux Python.<\/p>\n<p>L'\u00e9tude de cas compl\u00e8te est disponible dans <a href=\"https:\/\/developer.nvidia.com\/blog\/autodesk-research-brings-warp-speed-to-computational-fluid-dynamics-on-nvidia-gh200\/\" rel=\"nofollow\" target=\"_blank\">La recherche d'Autodesk apporte une vitesse de distorsion \u00e0 la dynamique des fluides de calcul sur NVIDIA GH200<\/a>. Le code de projet actuel est disponible dans le <a href=\"https:\/\/github.com\/Autodesk\/XLB\" rel=\"nofollow\" target=\"_blank\">Autodesk XLB Repository<\/a>.<\/p>\n<h2>Comment comparer correctement un noyau<\/h2>\n<p>Les op\u00e9rations GPU sont g\u00e9n\u00e9ralement asynchrones. Mesure uniquement la dur\u00e9e de l'appel de fonctions Python peut indiquer le temps n\u00e9cessaire pour mettre le noyau en file plut\u00f4t que le temps n\u00e9cessaire pour l'ex\u00e9cuter.<\/p>\n<p>Une r\u00e9f\u00e9rence de base devrait :<\/p>\n<ol>\n<li>Ex\u00e9cutez le noyau plusieurs fois pour d\u00e9clencher la compilation et l'\u00e9chauffement.<\/li>\n<li>Synchronisez avant de d\u00e9marrer la minuterie.<\/li>\n<li>Ex\u00e9cutez plusieurs it\u00e9rations mesur\u00e9es.<\/li>\n<li>Synchronisez avant d'arr\u00eater la minuterie.<\/li>\n<li>Signaler la moyenne et la variation des r\u00e9p\u00e9titions.<\/li>\n<li>S\u00e9parez le temps de transfert de donn\u00e9es du temps d'ex\u00e9cution du noyau.<\/li>\n<li>V\u00e9rifiez que les versions du CPU et du GPU produisent des r\u00e9sultats \u00e9quivalents.<\/li>\n<\/ol>\n<pre><code class=\"language-python\">import time\nfrom numba import cuda\n\n# Warm-up and JIT compilation\nkernel[blocks, threads](*arguments)\ncuda.synchronize()\n\nstart = time.perf_counter()\n\nfor _ in range(100):\n    kernel[blocks, threads](*arguments)\n\ncuda.synchronize()\n\nelapsed = time.perf_counter() - start\naverage = elapsed \/ 100\n\nprint(\"Average kernel time:\", average)<\/code><\/pre>\n<p>La comparaison doit utiliser une taille de simulation r\u00e9aliste et inclure le flux de travail complet lorsque les co\u00fbts de transfert ou de pr\u00e9traitement sont importants.<\/p>\n<p>La m\u00e9thodologie de pr\u00e9cision du travail peut \u00e9galement comparer le temps d'ex\u00e9cution \u00e0 l'erreur num\u00e9rique. La <a href=\"https:\/\/docs.sciml.ai\/DiffEqDevDocs\/stable\/alg_dev\/benchmarks\/\" rel=\"nofollow\" target=\"_blank\">Documentation de r\u00e9f\u00e9rence SCIML<\/a> fournit des exemples de ce style d'\u00e9valuation.<\/p>\n<h2>Un flux de travail de d\u00e9veloppement pratique<\/h2>\n<ol>\n<li>Impl\u00e9mentez et v\u00e9rifiez une version claire de la r\u00e9f\u00e9rence CPU.<\/li>\n<li>Profilez l'application pour localiser l'op\u00e9ration dominante.<\/li>\n<li>Essayez une biblioth\u00e8que GPU optimis\u00e9e avant d'\u00e9crire un noyau.<\/li>\n<li>Conservez les donn\u00e9es fr\u00e9quemment r\u00e9utilis\u00e9es sur l'appareil.<\/li>\n<li>\u00c9crivez le noyau correct le plus simple.<\/li>\n<li>Validez-le par rapport au r\u00e9sultat du CPU.<\/li>\n<li>Mesurez les transferts de m\u00e9moire et l'ex\u00e9cution s\u00e9par\u00e9ment.<\/li>\n<li>Inspectez la coalescence, l'occupation, la divergence et l'utilisation des registres.<\/li>\n<li>Testez la m\u00e9moire partag\u00e9e ou la fusion uniquement lorsque le profilage les prend en charge.<\/li>\n<li>Benchmark plusieurs tailles de probl\u00e8mes et architectures GPU.<\/li>\n<\/ol>\n<h2>Choisir un cadre<\/h2>\n<table class=\"custom-table\">\n<tbody><tr>\n<th>Exigence<\/th>\n<th>Point de d\u00e9part possible<\/th>\n<\/tr>\n<tr>\n<td>Flux de travail GPU de style Numpy existant<\/td>\n<td>CUPY avec des op\u00e9rations int\u00e9gr\u00e9es ou des noyaux personnalis\u00e9s<\/td>\n<\/tr>\n<tr>\n<td>Petit nombre de noyaux python de style CUDA<\/td>\n<td>NUMBA-CUDA, apr\u00e8s avoir examin\u00e9 son \u00e9tat de support actuel<\/td>\n<\/tr>\n<tr>\n<td>Simulation, g\u00e9om\u00e9trie et noyaux diff\u00e9rentiables<\/td>\n<td>NVIDIA WARP<\/td>\n<\/tr>\n<tr>\n<td>Contr\u00f4le sp\u00e9cifique \u00e0 NVIDIA maximum<\/td>\n<td>CUDA C++<\/td>\n<\/tr>\n<tr>\n<td>Cible GPU AMD avec une source de type CUDA<\/td>\n<td>Hanche et ROCM<\/td>\n<\/tr>\n<tr>\n<td>C++ portable sur plusieurs backends<\/td>\n<td>Sycl, Kokkos ou une autre couche de portabilit\u00e9 des performances<\/td>\n<\/tr>\n<\/tbody><\/table>\n<h2>Erreurs courantes de programmation du noyau<\/h2>\n<ul>\n<li>D\u00e9placement des donn\u00e9es entre le CPU et le GPU \u00e0 chaque pas de temps<\/li>\n<li>Ignorer les threads hors limites<\/li>\n<li>Utiliser un algorithme inefficace et s'attendre \u00e0 ce que le mat\u00e9riel corrige sa mise \u00e0 l'\u00e9chelle<\/li>\n<li>Acc\u00e9der \u00e0 la m\u00e9moire avec une mise en page non fusionn\u00e9e<\/li>\n<li>Ajouter de la m\u00e9moire partag\u00e9e sans mesurer si cela aide<\/li>\n<li>Lancement de nombreux minuscules noyaux au lieu d'envisager la fusion<\/li>\n<li>Utilisation de registres excessifs ou de m\u00e9moire partag\u00e9e par bloc<\/li>\n<li>Analyse comparative du code asynchrone sans synchronisation<\/li>\n<li>Comparaison des sorties sans v\u00e9rification de la pr\u00e9cision num\u00e9rique<\/li>\n<li>Utilisation de revendications de performances fixes sur diff\u00e9rents GPU et charges de travail<\/li>\n<li>En supposant que la source portable offre des performances portables<\/li>\n<li>Application de la diff\u00e9renciation automatique aux op\u00e9rations sur place non prises en charge<\/li>\n<\/ul>\n<h2>Guides connexes<\/h2>\n<ul>\n<li><a href=\"https:\/\/matforge.org\/gpu-accelerated-scientific-computing-cupy-numba-cudf-compared\/\" rel=\"nofollow\" target=\"_blank\">scientifique acc\u00e9l\u00e9r\u00e9 par GPU Informatique&nbsp;: CUPY, NUMBA et CUDF compar\u00e9s<\/a> - Comparez les outils Python de haut niveau pour les charges de travail GPU.<\/li>\n<li><a href=\"https:\/\/matforge.org\/gpu-acceleration-for-fipy-simulations-cupy-and-numba-integration-guide\/\" rel=\"nofollow\" target=\"_blank\">Acc\u00e9l\u00e9ration GPU pour les simulations FIPY&nbsp;: cupy et Numba Integration Guide<\/a> \u2014 Explorez les composants GPU possibles autour d'un flux de travail Fipy.<\/li>\n<li><a href=\"https:\/\/matforge.org\/physics-informed-neural-networks-pinns-for-scientific-simulations\/\" rel=\"nofollow\" target=\"_blank\">R\u00e9seaux de neurones \u00e0 base de physique pour les simulations scientifiques<\/a> \u2014 Passez en revue la relation entre les mod\u00e8les physiques et l'apprentissage diff\u00e9rentiable.<\/li>\n<li><a href=\"https:\/\/matforge.org\/when-to-use-fem-fvm-fdm\/\" rel=\"nofollow\" target=\"_blank\">Quand utiliser FEM, FVM et FDM<\/a> Discr\u00e9tisation avant d'optimiser sa mise en \u0153uvre.<\/li>\n<li><a href=\"https:\/\/matforge.org\/scientific-computing-performance-profiling-identifying-bottlenecks\/\" rel=\"nofollow\" target=\"_blank\">Profilage des performances&nbsp;: identifier les goulots d'\u00e9tranglement dans le code scientifique<\/a>&nbsp;: localiser les limitations du processeur, de la m\u00e9moire, du transfert et des E\/S.<\/li>\n<li><a href=\"https:\/\/matforge.org\/hpc-python-workflows-from-laptop-to-supercomputer\/\" rel=\"nofollow\" target=\"_blank\">Flows Python HPC&nbsp;: depuis un ordinateur portable au superordinateur<\/a> \u2014 Planifiez la progression du d\u00e9veloppement local vers les grands syst\u00e8mes.<\/li>\n<\/ul>\n<h2>Conclusion<\/h2>\n<p>Les noyaux GPU personnalis\u00e9s permettent aux d\u00e9veloppeurs scientifiques de traduire directement des particules, des cellules de la grille, des r\u00e8gles de mat\u00e9riaux et d'autres op\u00e9rations physiques en fonctions massivement parall\u00e8les. Ils permettent de contr\u00f4ler l'indexation, l'acc\u00e8s \u00e0 la m\u00e9moire, la synchronisation, la fusion du noyau et l'optimisation sp\u00e9cifique \u00e0 l'appareil.<\/p>\n<p>CUPY propose \u00e0 la fois des op\u00e9rations de baies de haut niveau et des interfaces de noyau personnalis\u00e9es. Numba-CUDA fournit un mod\u00e8le de type CUDA en Python, tandis que NVIDIA WARP ajoute des primitives de simulation, des op\u00e9rations de tuiles et une diff\u00e9renciation automatique. CUDA et HIP offrent un contr\u00f4le C++ de niveau inf\u00e9rieur, et les cadres portables de performances prennent en charge des objectifs mat\u00e9riels plus larges.<\/p>\n<p>Les gains de performances les plus importants proviennent rarement de la modification de la syntaxe seule. Ils proviennent du choix d'un algorithme parall\u00e8le, de la conservation des donn\u00e9es sur l'appareil, de la r\u00e9duction du trafic m\u00e9moire, de l'utilisation d'une disposition de donn\u00e9es appropri\u00e9e et de l'\u00e9limination des op\u00e9rations interm\u00e9diaires inutiles.<\/p>\n<p>Un noyau personnalis\u00e9 ne doit \u00eatre d\u00e9velopp\u00e9 qu'apr\u00e8s le profilage identifiant un v\u00e9ritable goulot d'\u00e9tranglement. Il doit \u00eatre valid\u00e9 par rapport \u00e0 une solution de r\u00e9f\u00e9rence et \u00e9talonn\u00e9 avec la synchronisation, la taille repr\u00e9sentative des probl\u00e8mes et la comptabilisation compl\u00e8te des co\u00fbts de transfert de m\u00e9moire.<\/p>\n<p>Lorsque ces conditions sont remplies, les outils de noyau bas\u00e9s sur Python peuvent prendre en charge la simulation physique s\u00e9rieuse sans obliger les chercheurs \u00e0 d\u00e9placer chaque partie de l'application en C++ de bas niveau.<\/p>\n<\/article>\n"},"excerpt":{"rendered":"<p><span class=\"span-reading-time rt-reading-time\" style=\"display: block;\"><span class=\"rt-label rt-prefix\">Reading Time: <\/span> <span class=\"rt-time\"> 14<\/span> <span class=\"rt-label rt-postfix\">minutes<\/span><\/span>Apprenez \u00e0 \u00e9crire des noyaux GPU personnalis\u00e9s pour la simulation physique avec NVIDIA WARP, CUDA Python et HIP\/ROCM. Couvre les mod\u00e8les de performances, la physique diff\u00e9rentiable et les d\u00e9ploiements de production.<\/p>\n","protected":false,"raw":"Apprenez \u00e0 \u00e9crire des noyaux GPU personnalis\u00e9s pour la simulation physique avec NVIDIA WARP, CUDA Python et HIP\/ROCM. Couvre les mod\u00e8les de performances, la physique diff\u00e9rentiable et les d\u00e9ploiements de production."},"author":2,"featured_media":0,"comment_status":"open","ping_status":"open","sticky":false,"template":"","format":"standard","meta":{"_locale":"fr_FR","_original_post":"https:\/\/matforge.org\/?p=482","iawp_total_views":0,"footnotes":""},"categories":[1],"tags":[],"class_list":["post-1293","post","type-post","status-publish","format-standard","hentry","category-simulation-modeling-projects","fr-FR"],"yoast_head":"<!-- This site is optimized with the Yoast SEO plugin v28.3 - https:\/\/yoast.com\/product\/yoast-seo-wordpress\/ -->\n<title>Noyaux de GPU personnalis\u00e9s pour les simulations physiques<\/title>\n<meta name=\"description\" content=\"Apprenez \u00e0 \u00e9crire et \u00e0 optimiser les noyaux GPU pour des simulations physiques \u00e0 l&#039;aide de Numba, Warp, Cupy, CUDA, HIP et m\u00e9moire.\" \/>\n<meta name=\"robots\" content=\"index, follow, max-snippet:-1, max-image-preview:large, max-video-preview:-1\" \/>\n<link rel=\"canonical\" href=\"https:\/\/matforge.org\/fr\/gpu-kernel-programming-custom-physics-simulation\/\" \/>\n<meta property=\"og:locale\" content=\"fr_FR\" \/>\n<meta property=\"og:type\" content=\"article\" \/>\n<meta property=\"og:title\" content=\"Noyaux de GPU personnalis\u00e9s pour les simulations physiques\" \/>\n<meta property=\"og:description\" content=\"Apprenez \u00e0 \u00e9crire et \u00e0 optimiser les noyaux GPU pour des simulations physiques \u00e0 l&#039;aide de Numba, Warp, Cupy, CUDA, HIP et m\u00e9moire.\" \/>\n<meta property=\"og:url\" content=\"https:\/\/matforge.org\/fr\/gpu-kernel-programming-custom-physics-simulation\/\" \/>\n<meta property=\"og:site_name\" content=\"matforge.org\" \/>\n<meta property=\"article:published_time\" content=\"2026-08-21T14:31:18+00:00\" \/>\n<meta name=\"author\" content=\"Elena Markovska\" \/>\n<meta name=\"twitter:card\" content=\"summary_large_image\" \/>\n<meta name=\"twitter:label1\" content=\"\u00c9crit par\" \/>\n\t<meta name=\"twitter:data1\" content=\"Elena Markovska\" \/>\n\t<meta name=\"twitter:label2\" content=\"Dur\u00e9e de lecture estim\u00e9e\" \/>\n\t<meta name=\"twitter:data2\" content=\"23 minutes\" \/>\n<script type=\"application\/ld+json\" class=\"yoast-schema-graph\">{\"@context\":\"https:\\\/\\\/schema.org\",\"@graph\":[{\"@type\":\"Article\",\"@id\":\"https:\\\/\\\/matforge.org\\\/fr\\\/gpu-kernel-programming-custom-physics-simulation\\\/#article\",\"isPartOf\":{\"@id\":\"https:\\\/\\\/matforge.org\\\/fr\\\/gpu-kernel-programming-custom-physics-simulation\\\/\"},\"author\":{\"name\":\"Elena Markovska\",\"@id\":\"https:\\\/\\\/matforge.org\\\/#\\\/schema\\\/person\\\/980162bb5de46742daece973661d93da\"},\"headline\":\"Programmation du noyau GPU pour la simulation physique personnalis\u00e9e\u00a0: \u00e9criture de noyaux personnalis\u00e9s hautes performances en Python\",\"datePublished\":\"2026-08-21T14:31:18+00:00\",\"mainEntityOfPage\":{\"@id\":\"https:\\\/\\\/matforge.org\\\/fr\\\/gpu-kernel-programming-custom-physics-simulation\\\/\"},\"wordCount\":3868,\"commentCount\":0,\"articleSection\":[\"Simulation &amp; Projets de mod\u00e9lisation\"],\"inLanguage\":\"fr-FR\",\"potentialAction\":[{\"@type\":\"CommentAction\",\"name\":\"Comment\",\"target\":[\"https:\\\/\\\/matforge.org\\\/fr\\\/gpu-kernel-programming-custom-physics-simulation\\\/#respond\"]}]},{\"@type\":\"WebPage\",\"@id\":\"https:\\\/\\\/matforge.org\\\/fr\\\/gpu-kernel-programming-custom-physics-simulation\\\/\",\"url\":\"https:\\\/\\\/matforge.org\\\/fr\\\/gpu-kernel-programming-custom-physics-simulation\\\/\",\"name\":\"Noyaux de GPU personnalis\u00e9s pour les simulations physiques\",\"isPartOf\":{\"@id\":\"https:\\\/\\\/matforge.org\\\/#website\"},\"datePublished\":\"2026-08-21T14:31:18+00:00\",\"author\":{\"@id\":\"https:\\\/\\\/matforge.org\\\/#\\\/schema\\\/person\\\/980162bb5de46742daece973661d93da\"},\"description\":\"Apprenez \u00e0 \u00e9crire et \u00e0 optimiser les noyaux GPU pour des simulations physiques \u00e0 l'aide de Numba, Warp, Cupy, CUDA, HIP et m\u00e9moire.\",\"breadcrumb\":{\"@id\":\"https:\\\/\\\/matforge.org\\\/fr\\\/gpu-kernel-programming-custom-physics-simulation\\\/#breadcrumb\"},\"inLanguage\":\"fr-FR\",\"potentialAction\":[{\"@type\":\"ReadAction\",\"target\":[\"https:\\\/\\\/matforge.org\\\/fr\\\/gpu-kernel-programming-custom-physics-simulation\\\/\"]}]},{\"@type\":\"BreadcrumbList\",\"@id\":\"https:\\\/\\\/matforge.org\\\/fr\\\/gpu-kernel-programming-custom-physics-simulation\\\/#breadcrumb\",\"itemListElement\":[{\"@type\":\"ListItem\",\"position\":1,\"name\":\"Home\",\"item\":\"https:\\\/\\\/matforge.org\\\/\"},{\"@type\":\"ListItem\",\"position\":2,\"name\":\"Programmation du noyau GPU pour la simulation physique personnalis\u00e9e\u00a0: \u00e9criture de noyaux personnalis\u00e9s hautes performances en Python\"}]},{\"@type\":\"WebSite\",\"@id\":\"https:\\\/\\\/matforge.org\\\/#website\",\"url\":\"https:\\\/\\\/matforge.org\\\/\",\"name\":\"matforge.org\",\"description\":\"\",\"potentialAction\":[{\"@type\":\"SearchAction\",\"target\":{\"@type\":\"EntryPoint\",\"urlTemplate\":\"https:\\\/\\\/matforge.org\\\/?s={search_term_string}\"},\"query-input\":{\"@type\":\"PropertyValueSpecification\",\"valueRequired\":true,\"valueName\":\"search_term_string\"}}],\"inLanguage\":\"fr-FR\"},{\"@type\":\"Person\",\"@id\":\"https:\\\/\\\/matforge.org\\\/#\\\/schema\\\/person\\\/980162bb5de46742daece973661d93da\",\"name\":\"Elena Markovska\",\"image\":{\"@type\":\"ImageObject\",\"inLanguage\":\"fr-FR\",\"@id\":\"https:\\\/\\\/secure.gravatar.com\\\/avatar\\\/2de4e35b6581d7d8a839335156c6b5834cfaa0aef537a1c837e882dc57eea1e7?s=96&d=mm&r=g\",\"url\":\"https:\\\/\\\/secure.gravatar.com\\\/avatar\\\/2de4e35b6581d7d8a839335156c6b5834cfaa0aef537a1c837e882dc57eea1e7?s=96&d=mm&r=g\",\"contentUrl\":\"https:\\\/\\\/secure.gravatar.com\\\/avatar\\\/2de4e35b6581d7d8a839335156c6b5834cfaa0aef537a1c837e882dc57eea1e7?s=96&d=mm&r=g\",\"caption\":\"Elena Markovska\"},\"sameAs\":[\"http:\\\/\\\/matforge.org\"],\"url\":\"https:\\\/\\\/matforge.org\\\/author\\\/elena-markovska\\\/\"}]}<\/script>\n<!-- \/ Yoast SEO plugin. -->","yoast_head_json":{"title":"Noyaux de GPU personnalis\u00e9s pour les simulations physiques","description":"Apprenez \u00e0 \u00e9crire et \u00e0 optimiser les noyaux GPU pour des simulations physiques \u00e0 l'aide de Numba, Warp, Cupy, CUDA, HIP et m\u00e9moire.","robots":{"index":"index","follow":"follow","max-snippet":"max-snippet:-1","max-image-preview":"max-image-preview:large","max-video-preview":"max-video-preview:-1"},"canonical":"https:\/\/matforge.org\/fr\/gpu-kernel-programming-custom-physics-simulation\/","og_locale":"fr_FR","og_type":"article","og_title":"Noyaux de GPU personnalis\u00e9s pour les simulations physiques","og_description":"Apprenez \u00e0 \u00e9crire et \u00e0 optimiser les noyaux GPU pour des simulations physiques \u00e0 l'aide de Numba, Warp, Cupy, CUDA, HIP et m\u00e9moire.","og_url":"https:\/\/matforge.org\/fr\/gpu-kernel-programming-custom-physics-simulation\/","og_site_name":"matforge.org","article_published_time":"2026-08-21T14:31:18+00:00","author":"Elena Markovska","twitter_card":"summary_large_image","twitter_misc":{"\u00c9crit par":"Elena Markovska","Dur\u00e9e de lecture estim\u00e9e":"23 minutes"},"schema":{"@context":"https:\/\/schema.org","@graph":[{"@type":"Article","@id":"https:\/\/matforge.org\/fr\/gpu-kernel-programming-custom-physics-simulation\/#article","isPartOf":{"@id":"https:\/\/matforge.org\/fr\/gpu-kernel-programming-custom-physics-simulation\/"},"author":{"name":"Elena Markovska","@id":"https:\/\/matforge.org\/#\/schema\/person\/980162bb5de46742daece973661d93da"},"headline":"Programmation du noyau GPU pour la simulation physique personnalis\u00e9e\u00a0: \u00e9criture de noyaux personnalis\u00e9s hautes performances en Python","datePublished":"2026-08-21T14:31:18+00:00","mainEntityOfPage":{"@id":"https:\/\/matforge.org\/fr\/gpu-kernel-programming-custom-physics-simulation\/"},"wordCount":3868,"commentCount":0,"articleSection":["Simulation &amp; Projets de mod\u00e9lisation"],"inLanguage":"fr-FR","potentialAction":[{"@type":"CommentAction","name":"Comment","target":["https:\/\/matforge.org\/fr\/gpu-kernel-programming-custom-physics-simulation\/#respond"]}]},{"@type":"WebPage","@id":"https:\/\/matforge.org\/fr\/gpu-kernel-programming-custom-physics-simulation\/","url":"https:\/\/matforge.org\/fr\/gpu-kernel-programming-custom-physics-simulation\/","name":"Noyaux de GPU personnalis\u00e9s pour les simulations physiques","isPartOf":{"@id":"https:\/\/matforge.org\/#website"},"datePublished":"2026-08-21T14:31:18+00:00","author":{"@id":"https:\/\/matforge.org\/#\/schema\/person\/980162bb5de46742daece973661d93da"},"description":"Apprenez \u00e0 \u00e9crire et \u00e0 optimiser les noyaux GPU pour des simulations physiques \u00e0 l'aide de Numba, Warp, Cupy, CUDA, HIP et m\u00e9moire.","breadcrumb":{"@id":"https:\/\/matforge.org\/fr\/gpu-kernel-programming-custom-physics-simulation\/#breadcrumb"},"inLanguage":"fr-FR","potentialAction":[{"@type":"ReadAction","target":["https:\/\/matforge.org\/fr\/gpu-kernel-programming-custom-physics-simulation\/"]}]},{"@type":"BreadcrumbList","@id":"https:\/\/matforge.org\/fr\/gpu-kernel-programming-custom-physics-simulation\/#breadcrumb","itemListElement":[{"@type":"ListItem","position":1,"name":"Home","item":"https:\/\/matforge.org\/"},{"@type":"ListItem","position":2,"name":"Programmation du noyau GPU pour la simulation physique personnalis\u00e9e\u00a0: \u00e9criture de noyaux personnalis\u00e9s hautes performances en Python"}]},{"@type":"WebSite","@id":"https:\/\/matforge.org\/#website","url":"https:\/\/matforge.org\/","name":"matforge.org","description":"","potentialAction":[{"@type":"SearchAction","target":{"@type":"EntryPoint","urlTemplate":"https:\/\/matforge.org\/?s={search_term_string}"},"query-input":{"@type":"PropertyValueSpecification","valueRequired":true,"valueName":"search_term_string"}}],"inLanguage":"fr-FR"},{"@type":"Person","@id":"https:\/\/matforge.org\/#\/schema\/person\/980162bb5de46742daece973661d93da","name":"Elena Markovska","image":{"@type":"ImageObject","inLanguage":"fr-FR","@id":"https:\/\/secure.gravatar.com\/avatar\/2de4e35b6581d7d8a839335156c6b5834cfaa0aef537a1c837e882dc57eea1e7?s=96&d=mm&r=g","url":"https:\/\/secure.gravatar.com\/avatar\/2de4e35b6581d7d8a839335156c6b5834cfaa0aef537a1c837e882dc57eea1e7?s=96&d=mm&r=g","contentUrl":"https:\/\/secure.gravatar.com\/avatar\/2de4e35b6581d7d8a839335156c6b5834cfaa0aef537a1c837e882dc57eea1e7?s=96&d=mm&r=g","caption":"Elena Markovska"},"sameAs":["http:\/\/matforge.org"],"url":"https:\/\/matforge.org\/author\/elena-markovska\/"}]}},"_links":{"self":[{"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/posts\/1293","targetHints":{"allow":["GET"]}}],"collection":[{"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/posts"}],"about":[{"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/types\/post"}],"author":[{"embeddable":true,"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/users\/2"}],"replies":[{"embeddable":true,"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/comments?post=1293"}],"version-history":[{"count":1,"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/posts\/1293\/revisions"}],"predecessor-version":[{"id":1397,"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/posts\/1293\/revisions\/1397"}],"wp:attachment":[{"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/media?parent=1293"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/categories?post=1293"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/tags?post=1293"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}