{"id":906,"date":"2026-07-30T12:23:15","date_gmt":"2026-07-30T12:23:15","guid":{"rendered":"https:\/\/matforge.org\/?p=906","raw":"https:\/\/matforge.org\/?p=906"},"modified":"2026-07-30T12:23:15","modified_gmt":"2026-07-30T12:23:15","slug":"gpu-kernel-programming-custom-physics-simulation","status":"publish","type":"post","link":"https:\/\/matforge.org\/de\/gpu-kernel-programming-custom-physics-simulation\/","title":{"rendered":"GPU-Kernel-Programmierung f\u00fcr benutzerdefinierte Physiksimulation: Schreiben von leistungsstarken benutzerdefinierten Kerneln in Python","raw":"GPU-Kernel-Programmierung f\u00fcr benutzerdefinierte Physiksimulation: Schreiben von leistungsstarken benutzerdefinierten Kerneln in 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\"> 11<\/span> <span class=\"rt-label rt-postfix\">minutes<\/span><\/span><article>\n<p>Die GPU-Beschleunigung im wissenschaftlichen Rechnen beginnt h\u00e4ufig mit Bibliotheken auf hoher Ebene. Ein Forscher ersetzt Numpy durch Cupy, verwendet eine beschleunigte FFT oder ruft eine GPU-f\u00e4hige lineare Algebra-Routine auf. Dies kann zu einer wesentlichen Verbesserung f\u00fchren, ohne dass detaillierte Kenntnisse der GPU-Hardware erforderlich sind.<\/p>\n<p>Die benutzerdefinierte GPU-Kernel-Programmierung geht eine Ebene tiefer. Anstatt eine vordefinierte Operation aufzurufen, schreibt der Entwickler die Funktion, die jeder GPU-Thread ausf\u00fchrt. Dies bietet Kontrolle \u00fcber Thread-Indizierung, Datenlayout, Speicher\u00fcbertragungen, Synchronisation, gemeinsam genutzten Speicher und die Anzahl der w\u00e4hrend jedes Kernels gestarteten Operationen.<\/p>\n<p>Diese Steuerung ist n\u00fctzlich, wenn ein physikalisches Modell Zugriffsmuster enth\u00e4lt, die nicht effizient durch gew\u00f6hnliche Array-Operationen ausgedr\u00fcckt werden k\u00f6nnen. Partikelwechselwirkungen, Finite-Differenz-Schablonen, Kollisionsregeln, gitterbasierte Fluidmodelle, Materialaktualisierungen und Geometrieabfragen sind g\u00e4ngige Beispiele.<\/p>\n<p>Ein benutzerdefinierter Kernel ist nicht automatisch schneller als eine optimierte Bibliothek. Es wird wertvoll, wenn der Algorithmus gen\u00fcgend parallele Arbeit aufzeigt und der Entwickler den Speicherverkehr reduzieren, Operationen kombinieren oder die Berechnung an die GPU-Architektur anpassen kann.<\/p>\n<h2>Was ist ein GPU-Kernel?<\/h2>\n<p>Ein GPU-Kernel ist eine Funktion, die von vielen leichten Threads ausgef\u00fchrt wird. Jeder Thread behandelt normalerweise ein Teilchen, eine Gitterzelle, ein Element, eine Fl\u00e4che, ein Pixel oder einen Eintrag in einem Array.<\/p>\n<p>Das Host-Programm startet den Kernel mit einem Raster von Thread-Bl\u00f6cken:<\/p>\n<pre><code>kernel[blocks, threads_per_block](arguments)<\/code><\/pre>\n<p>Jeder Thread bestimmt, welche Daten er aus seinen Block- und Thread-Indizes verarbeiten soll. Ein typischer eindimensionaler Index ist:<\/p>\n<pre><code>index =\n    block_index * block_size\n    + thread_index<\/code><\/pre>\n<p>Threads sind in Gruppen unterteilt, die als Bl\u00f6cke bezeichnet werden. Threads in einem Block k\u00f6nnen Daten \u00fcber gemeinsam genutzten Speicher synchronisieren und austauschen. Es wird allgemein erwartet, dass separate Bl\u00f6cke unabh\u00e4ngig ausgef\u00fchrt werden.<\/p>\n<h2>Wenn benutzerdefinierte Kernel es wert sind, geschrieben zu werden<\/h2>\n<p>GPU-Bibliotheken auf hoher Ebene sollten in der Regel die erste Option sein. Sie bieten bereits optimierte Matrixmultiplikation, FFTs, Reduktionen, sp\u00e4rliche Operationen, Zufallszahlengenerierung und viele elementweise Funktionen.<\/p>\n<p>Ein benutzerdefinierter Kernel wird n\u00fctzlich, wenn:<\/p>\n<ul>\n<li>Die Simulation verwendet ein nicht standardm\u00e4\u00dfiges Speicherzugriffsmuster.<\/li>\n<li>Mehrere kleine Operationen k\u00f6nnen zu einem Durchgang \u00fcber Speicher verschmolzen werden.<\/li>\n<li>Benachbarte Threads ben\u00f6tigen wiederholt dieselben lokalen Daten.<\/li>\n<li>Eine spezielle Partikel-, Kollisions- oder Schablonenregel dominiert die Laufzeit.<\/li>\n<li>Zwischen Arrays verbrauchen zu viel Ger\u00e4tespeicher.<\/li>\n<li>Die Simulation ben\u00f6tigt eine benutzerdefinierte differenzierbare Operation.<\/li>\n<li>Eine vorhandene Bibliothek kann die erforderliche Grenz- oder Materiallogik nicht effizient ausdr\u00fccken.<\/li>\n<\/ul>\n<p>Profilieren Sie das Programm, bevor Sie einen Kernel schreiben. Die Optimierung einer optisch komplexen Funktion, die nur einen kleinen Teil der Laufzeit ausmacht, f\u00fchrt nicht zu einer aussagekr\u00e4ftigen Gesamtbeschleunigung.<\/p>\n<h2>GPU-Programmiermodelle<\/h2>\n<p>Wissenschaftliche Software kann \u00fcber mehrere Abstraktionsebenen auf GPUs zugreifen.<\/p>\n<table class=\"custom-table\">\n<tbody>\n<tr>\n<th>Sich n\u00e4hern<\/th>\n<th>Beispiele<\/th>\n<th>Hauptst\u00e4rke<\/th>\n<th>Hauptkompromiss<\/th>\n<\/tr>\n<tr>\n<td>Direktives Offloading<\/td>\n<td>OpenMP, OpenACC<\/td>\n<td>Inkrementelle Beschleunigung des vorhandenen C, C++ oder FORTRAN-Codes<\/td>\n<td>Weniger direkte Kontrolle \u00fcber generierte Kernel<\/td>\n<\/tr>\n<tr>\n<td>Herstellerorientierte Kernel-Programmierung<\/td>\n<td>Cuda, H\u00fcfte<\/td>\n<td>Detaillierte Kontrolle \u00fcber die GPU-Ausf\u00fchrung<\/td>\n<td>H\u00f6here Implementierungs- und Wartungskosten<\/td>\n<\/tr>\n<tr>\n<td>Leistungstragbares C ++<\/td>\n<td>Sycl, Kokkos, Alpaka<\/td>\n<td>Ein Quellmodell f\u00fcr mehrere Hardware-Backends<\/td>\n<td>Portabilit\u00e4t garantiert nicht \u00fcberall gleiche Leistung<\/td>\n<\/tr>\n<tr>\n<td>Python-GPU-Arrays und -Kernel<\/td>\n<td>Cupy, Numba-Cuda, Nvidia Warp<\/td>\n<td>Schnelle Entwicklung mit Zugriff auf kompilierten GPU-Code<\/td>\n<td>Framework-spezifische Einschr\u00e4nkungen und Zusammenstellungsverhalten<\/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 Einf\u00fchrung in die GPU-Programmierung Modelle <\/a> bieten einen umfassenderen Vergleich von CUDA, HIP, OpenCL, Sycl, Kokkos, OpenMP und verwandten Ans\u00e4tzen.<\/p>\n<h2>Cuda<\/h2>\n<p>CUDA bietet direkten Zugriff auf die NVIDIA-GPU-Programmierung \u00fcber C++, FORTRAN und mehrere Sprachbindungen. Es stellt Gitter, Bl\u00f6cke, Warps, gemeinsam genutzten Speicher, Streams, Ereignisse, Ger\u00e4tebibliotheken und hardwarespezifische Optimierungstools bereit.<\/p>\n<p>CUDA wird h\u00e4ufig ausgew\u00e4hlt, wenn maximale NVIDIA-spezifische Leistung und ausgereifte Tools wichtiger sind als die Portabilit\u00e4t der Hardware. Die Hauptkosten sind das manuelle Speichermanagement, ein Programmiermodell auf niedrigerer Ebene und die Abh\u00e4ngigkeit von der NVIDIA-Plattform.<\/p>\n<h2>H\u00fcfte und ROCM<\/h2>\n<p>HIP bietet ein CUDA-\u00e4hnliches C ++ &#8211; Programmiermodell im ROCM-\u00d6kosystem von AMD. Eine gro\u00dfe Menge an CUDA-Code kann an HIP angepasst werden, aber die Portabilit\u00e4t der Quellen bedeutet nicht, dass eine Bin\u00e4rdatei auf jeder GPU unver\u00e4ndert ausgef\u00fchrt wird.<\/p>\n<p>F\u00fcr jede Zielarchitektur ist m\u00f6glicherweise weiterhin eine separate Kompilierung, Pr\u00fcfung und Leistungsoptimierung erforderlich. Details finden Sie in der <a href=\"https:\/\/rocm.docs.amd.com\/projects\/HIP\/en\/latest\/what_is_hip.html\" rel=\"nofollow\" target=\"_blank\">Official Hip Documentation <\/a>.<\/p>\n<p>HIP ist eine praktische Option, wenn AMD-GPUs unterst\u00fctzt werden m\u00fcssen oder wenn ein Projekt die Abh\u00e4ngigkeit von einem einzelnen Beschleunigerhersteller verringern m\u00f6chte, w\u00e4hrend ein CUDA-\u00e4hnlicher Programmierstil beibehalten wird.<\/p>\n<h2>Cupy benutzerdefinierte Kernel<\/h2>\n<p>Cupy ist am bekanntesten f\u00fcr seine numpy-kompatiblen Arrays und beschleunigten numerischen Funktionen. Es unterst\u00fctzt auch die benutzerdefinierte GPU-Programmierung.<\/p>\n<p>Entwickler k\u00f6nnen verwenden:<\/p>\n<ul>\n<li><code>ElementwiseKernel<\/code> f\u00fcr benutzerdefinierte elementweise Ausdr\u00fccke<\/li>\n<li><code>ReductionKernel<\/code> f\u00fcr benutzerdefinierte Reduzierungen<\/li>\n<li><code>RawKernel<\/code> f\u00fcr in CUDA C++ geschriebene Kernel<\/li>\n<li><code>cupyx.jit.rawkernel<\/code> f\u00fcr JIT-Kernel-Definitionen im Python-Stil<\/li>\n<\/ul>\n<p>Ein einfacher Cupy-Raw-Kernel kann wie folgt definiert werden:<\/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 ist n\u00fctzlich, wenn die meisten Anwendungen bereits GPU-Arrays verwenden und nur ausgew\u00e4hlte Operationen benutzerdefinierte Kernel ben\u00f6tigen.<\/p>\n<h2>Numba-Cuda<\/h2>\n<p>Numba-Cuda kompiliert eine eingeschr\u00e4nkte Teilmenge von Python in GPU-Kernel. Das Programmiermodell folgt CUDA C. Die Entwickler definieren eine Funktion mit einem Kerneldekorator, berechnen einen Thread-Index und starten die Funktion mit einer Raster- und Blockkonfiguration.<\/p>\n<p>Die aktive Projektdokumentation ist unter <a href=\"https:\/\/nvidia.github.io\/numba-cuda\/\" rel=\"nofollow\" target=\"_blank\">Numba-cuda<\/a> verf\u00fcgbar. Da sich der Entwicklungsstatus \u00e4ndern kann, sollten Teams, die langlebige Software erstellen, die aktuellen Wartungs- und Migrationsrichtlinien \u00fcberpr\u00fcfen, bevor sie sich f\u00fcr das Framework engagieren.<\/p>\n<h2>ein Numba-Teilchenkern<\/h2>\n<p>Das folgende Lehrbeispiel berechnet die Gravitationsbeschleunigung mit direkten All-Pairs-Interaktionen:<\/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>Dieser Kernel hat rechnerische Komplexit\u00e4t von <code>O(N\u00b2)<\/code> . Es ist zur Erkl\u00e4rung der Thread-Mapping geeignet, aber es ist kein effizienter Produktionsalgorithmus f\u00fcr sehr gro\u00dfe Partikelzahlen.<\/p>\n<p>Gro\u00dfe Gravitations- oder Partikelsysteme erfordern h\u00e4ufig eine Baummethode, eine schnelle Multipolmethode, r\u00e4umliche Beh\u00e4lter oder Nachbarlisten. Eine GPU kann das Skalierungsproblem eines ineffizienten Algorithmus nicht beseitigen.<\/p>\n<h2>Nvidia Warp<\/h2>\n<p>NVIDIA Warp erm\u00f6glicht Entwicklern, stark typisierte Kernel mit Python-Syntax zu definieren. Warp kompiliert diese Funktionen in CPU- oder CUDA-Code und bietet spezielle Primitive f\u00fcr Geometrie, Simulation, sp\u00e4rliche Operationen, endliche Elemente, Optimierung und automatische Differenzierung.<\/p>\n<p>Die aktuelle <a href=\"https:\/\/developer.nvidia.com\/warp-python\" rel=\"nofollow\" target=\"_blank\">Nvidia Warp-Produktseite<\/a> und <a href=\"https:\/\/github.com\/nvidia\/warp\" rel=\"nofollow\" target=\"_blank\">Warp-GitHub-Repository<\/a> Enth\u00e4lt Beispiele f\u00fcr Partikel, Fl\u00fcssigkeiten, Netze, Optimierungen, differenzierbare Simulationen und Tile-based GPU-Programmierung.<\/p>\n<h2>der gleiche Partikelkern in 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 bestimmt intern eine geeignete Startkonfiguration f\u00fcr einen gew\u00f6hnlichen Kernel-Start. Fortgeschrittene Benutzer k\u00f6nnen bei Bedarf weiterhin mit Blockdimensionen, Befehlsdiagrammen, gekachelten Ausf\u00fchrungen und speziellen Ger\u00e4teoperationen arbeiten.<\/p>\n<h2>Grid-Stride-Schleifen<\/h2>\n<p>Ein Kernel ben\u00f6tigt nicht f\u00fcr jedes Element einen fest zugewiesenen Thread. Eine Grid-Stride-Schleife erm\u00f6glicht es jedem Thread, mehrere Eintr\u00e4ge zu verarbeiten:<\/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>Dieses Muster trennt die Anzahl der gestarteten Threads von der Gesamtdatengr\u00f6\u00dfe. Es ist n\u00fctzlich, wenn sehr gro\u00dfe Arrays verarbeitet oder eine feste Startkonfiguration wiederverwendet werden.<\/p>\n<p>Eine praktische Einf\u00fchrung in dieses Muster finden Sie in <a href=\"https:\/\/thedatafrog.com\/en\/articles\/cuda-kernel-python\/\" rel=\"nofollow\" target=\"_blank\"> Das Numba- und CUDA-Tutorial des Datenfroschs <\/a>.<\/p>\n<h2>Ged\u00e4chtnis verschmelzen<\/h2>\n<p>GPU-Kernel sind h\u00e4ufig eher durch Speicherbandbreite als durch den arithmetischen Durchsatz begrenzt. Threads in einem Warp sollten idealerweise auf die nahe gelegenen Speicheradressen zugreifen, damit die Hardware ihre Anfragen kombinieren kann.<\/p>\n<p>Betrachten Sie Partikeldaten als gespeichert:<\/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>Dieses Array-of-Struktur-Layout kann f\u00fcr objektorientierten Code zweckm\u00e4\u00dfig sein. Ein Layout mit Struktur-of-Arrays speichert separate zusammenh\u00e4ngende Arrays:<\/p>\n<pre><code>x_positions[]\ny_positions[]\nz_positions[]\nmasses[]<\/code><\/pre>\n<p>Das zweite Layout kann eine bessere Koaleszenz liefern, wenn jeder Thread das gleiche Feld f\u00fcr ein anderes Teilchen liest. Das beste Layout h\u00e4ngt immer noch davon ab, auf welche Felder zusammen zugegriffen wird.<\/p>\n<p>Vektortypen korrigieren nicht automatisch ein schlechtes Zugriffsmuster. Entwickler sollten die tats\u00e4chlichen Adressen \u00fcberpr\u00fcfen, die von benachbarten Threads angefordert werden.<\/p>\n<h2>Geteilter Speicher<\/h2>\n<p>Shared Memory ist ein kleiner Speicherbereich mit geringer Latenz, auf den Threads im selben Block zugegriffen werden kann. Es kann wiederholte Lesevorg\u00e4nge aus dem globalen Speicher reduzieren.<\/p>\n<p>Eine eindimensionale Schablone kann einen Werteblock und seinen Halo in den gemeinsamen Speicher laden:<\/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>Der Aufruf von <code>cuda.syncthreads()<\/code> stellt sicher, dass jeder Thread seine Werte l\u00e4dt, bevor die Schablonenberechnung beginnt.<\/p>\n<p>Shared Memory sollte nicht automatisch verwendet werden. \u00dcberm\u00e4\u00dfige Zuweisung von Shared-Memory kann die Belegung verringern, die Synchronisierungskosten erh\u00f6hen und den Kernel langsamer machen. Profilierung erforderlich.<\/p>\n<h2>Kernfusion<\/h2>\n<p>Separate Array-Operationen erzeugen h\u00e4ufig mehrere Kernel-Starts und Zwischenarrays:<\/p>\n<pre><code>velocity += dt * acceleration\nposition += dt * velocity\nenergy = compute_energy(position, velocity)<\/code><\/pre>\n<p>Ein fusionierter Kernel kann alle drei Aktualisierungen berechnen, w\u00e4hrend die notwendigen Werte in Registern verbleiben. Dies reduziert den Start-Overhead und den globalen Speicherverkehr.<\/p>\n<p>Fusion ist am vorteilhaftesten, wenn die Operationen einfach und speichergebunden sind. Zu viel Arbeit kann die Nutzung des Registers erh\u00f6hen, die Belegung verringern und die Wartung des Kernels erschweren.<\/p>\n<h2>Blockgr\u00f6\u00dfe w\u00e4hlen<\/h2>\n<p>Werte wie 128 oder 256 Threads pro Block sind vern\u00fcnftige Ausgangspunkte, nicht universelle Optima.<\/p>\n<p>Die beste Blockgr\u00f6\u00dfe h\u00e4ngt ab von:<\/p>\n<ul>\n<li>Register pro Thread verwendet<\/li>\n<li>Shared Memory pro Block verwendet<\/li>\n<li>Zweigdivergenz<\/li>\n<li>Anleitungsmix<\/li>\n<li>Speicherzugriffsverhalten<\/li>\n<li>Die Ziel-GPU-Architektur<\/li>\n<\/ul>\n<p>Eine Regel wie das Starten von mindestens doppelt so vielen Bl\u00f6cken wie Streaming-Multiprozessoren kann ein n\u00fctzliches anf\u00e4ngliches Experiment sein, garantiert jedoch keine maximale Leistung. Belegungsrechner und Profilierungswerkzeuge sollten die endg\u00fcltige Konfiguration leiten.<\/p>\n<h2>differenzierbare Physik<\/h2>\n<p>Die differenzierbare Simulation berechnet, wie sich eine Ausgabe in Bezug auf Eingaben wie Materialeigenschaften, Kr\u00e4fte, Geometrie oder Anfangsbedingungen \u00e4ndert.<\/p>\n<p>Warp kann unterst\u00fctzte Kerneloperationen aufzeichnen und die automatische Differenzierung des Reverse-Modus ausf\u00fchren. Dies kann f\u00fcr inverse Probleme, Designoptimierung, Parametersch\u00e4tzung und Integration in maschinelle Lern-Workflows verwendet werden.<\/p>\n<p>Ein vereinfachtes Muster verwendet ein Band:<\/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>Eine automatische Differenzierung ist nicht f\u00fcr jeden Kernel garantiert. In-Place-Overwrites, nichtdeterministische Atome, externer nativer Code, diskontinuierliche Logik und nicht unterst\u00fctzte Operationen erfordern m\u00f6glicherweise eine Neuformulierung oder benutzerdefinierte Gradienten.<\/p>\n<p>Diese F\u00e4higkeit verbindet sich direkt mit Anwendungen wie Adjoint Optimization und <a href=\"https:\/\/matforge.org\/physics-informed-neural-networks-pinns-for-scientific-simulations\/\" rel=\"nofollow\" target=\"_blank\"> Physikinformierte neuronale Netze <\/a>.<\/p>\n<h2>Wenn eine GPU m\u00f6glicherweise nicht schneller ist<\/h2>\n<p>Es gibt keinen universellen Schwellenwert f\u00fcr die Problemgr\u00f6\u00dfe, bei dem eine GPU schneller wird als eine CPU. Der Crossover h\u00e4ngt von Hardware, Pr\u00e4zision, Datenbewegung, Algorithmenstruktur, Compilerqualit\u00e4t und der H\u00e4ufigkeit der Wiederverwendung derselben ger\u00e4teresidenten Daten ab.<\/p>\n<table class=\"custom-table\">\n<tbody>\n<tr>\n<th>Arbeitsanfall<\/th>\n<th>GPU-Eignung<\/th>\n<th>Haupt\u00fcberlegung<\/th>\n<\/tr>\n<tr>\n<td>Gro\u00dfe Partikelaktualisierung<\/td>\n<td>Oft stark<\/td>\n<td>viele \u00e4hnliche unabh\u00e4ngige Operationen<\/td>\n<\/tr>\n<tr>\n<td>Gro\u00dfe regul\u00e4re Schablone<\/td>\n<td>Oft stark<\/td>\n<td>Vorhersehbarer paralleler Speicherzugriff<\/td>\n<\/tr>\n<tr>\n<td>Dichte lineare Algebra<\/td>\n<td>Stark bei der Verwendung optimierter Bibliotheken<\/td>\n<td>Hohe arithmetische Intensit\u00e4t<\/td>\n<\/tr>\n<tr>\n<td>Kleine Simulation<\/td>\n<td>Hardwareabh\u00e4ngig<\/td>\n<td>Start und Transfer-Overhead k\u00f6nnen dominieren<\/td>\n<\/tr>\n<tr>\n<td>Unregelm\u00e4\u00dfige Graphendurchquerung<\/td>\n<td>Gemischt<\/td>\n<td>Divergenz und unvorhersehbarer Speicherzugriff<\/td>\n<\/tr>\n<tr>\n<td>I\/O-gebundener Workflow<\/td>\n<td>in der Regel begrenzt<\/td>\n<td>Die GPU kann einen Speicher- oder Netzwerkengpass nicht entfernen<\/td>\n<\/tr>\n<tr>\n<td>stark sequentieller Algorithmus<\/td>\n<td>normalerweise schwach<\/td>\n<td>Unzureichende unabh\u00e4ngige Arbeit<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n<p>Der richtige Workflow besteht darin, die CPU-Version, die GPU-Bibliotheksversion und die benutzerdefinierten Kernel-Version mit repr\u00e4sentativen Daten zu messen. <a href=\"https:\/\/matforge.org\/scientific-computing-performance-profiling-identifying-bottlenecks\/\" rel=\"nofollow\" target=\"_blank\">Forscherhandbuch f\u00fcr wissenschaftliche Rechenleistung<\/a> Erkl\u00e4rt, wie Sie den tats\u00e4chlichen Engpass identifizieren, bevor Sie ihn optimieren.<\/p>\n<h2>Vermeidung von Transfer-Overhead<\/h2>\n<p>Wiederholte \u00dcbertragungen zwischen Host- und Ger\u00e4tespeicher k\u00f6nnen den Nutzen der GPU-Berechnung beseitigen.<\/p>\n<p>Eine ineffiziente Schleife kann diesem Muster folgen:<\/p>\n<ol>\n<li>Daten auf die GPU kopieren.<\/li>\n<li>F\u00fchren Sie einen kleinen Kernel aus.<\/li>\n<li>Kopieren Sie das Ergebnis in die CPU.<\/li>\n<li>\u00c4ndern Sie es auf der CPU.<\/li>\n<li>Kopieren Sie es wieder auf die GPU.<\/li>\n<\/ol>\n<p>Ein besseres Design beh\u00e4lt den Simulationszustand \u00fcber viele Zeitschritte auf dem Ger\u00e4t und \u00fcbertr\u00e4gt nur die Ausgabe, die f\u00fcr Visualisierung, Checkpointing oder Analyse erforderlich ist.<\/p>\n<p>Asynchrone Streams, festgesteckter Host-Speicher, \u00fcberlappende Kommunikation und Unified-Memory-Funktionen k\u00f6nnen helfen, sollten jedoch erst eingef\u00fchrt werden, nachdem die normalen \u00dcbertragungskosten gemessen wurden.<\/p>\n<h2>Die XLB-Fallstudie<\/h2>\n<p>Das XLB-Projekt von Autodesk Research bietet ein n\u00fctzliches Beispiel f\u00fcr die Python-native GPU-Simulation. XLB ist eine Open-Source-Lattice-Boltzmann-Bibliothek mit mehreren rechnerischen Backends, darunter Nvidia Warp.<\/p>\n<p>In von Autodesk Research und NVIDIA gemeldeten Benchmark-Konfigurationen erreichte das Warp-Backend eine Leistung in der N\u00e4he der verglichenen C ++ \/ OpenCL FluidX3D-Implementierung f\u00fcr einen bestimmten Lid-gesteuerten Hohlraumfall. Ein separater Vergleich ergab eine ungef\u00e4hre achtfache Beschleunigung \u00fcber das JAX-Backend von XLB bei ausgew\u00e4hlten Hardware und Konfigurationen.<\/p>\n<p>Das Team zeigte auch einen Out-of-Core-Ansatz f\u00fcr einen GH200-Cluster mit acht Knoten mit einer Dom\u00e4ne von ungef\u00e4hr 50 Milliarden Gitterzellen. Diese Ergebnisse gelten f\u00fcr den gemeldeten Solver, Problem, Hardware und Implementierungsoptionen. Sie sollten nicht als allgemeine Beschleunigungsgarantien f\u00fcr Python-Kernel behandelt werden.<\/p>\n<p>Die vollst\u00e4ndige Fallstudie ist in <a href=\"https:\/\/developer.nvidia.com\/blog\/autodesk-research-brings-warp-speed-to-computational-fluid-dynamics-on-nvidia-gh200\/\" rel=\"nofollow\" target=\"_blank\"> verf\u00fcgbar. Autodesk-Forschung bringt die Warp-Geschwindigkeit in die rechnerische Fluiddynamik auf NVIDIA GH200 <\/a>. Der aktuelle Projektcode ist im <a href=\"https:\/\/github.com\/Autodesk\/XLB\" rel=\"nofollow\" target=\"_blank\"> Autodesk XLB-Repository <\/a> verf\u00fcgbar.<\/p>\n<h2>Wie man einen Kernel richtig bewertet<\/h2>\n<p>GPU-Operationen sind normalerweise asynchron. Das Messen nur der Python-Funktionsdauer kann die Zeit angeben, die erforderlich ist, um den Kernel in die Warteschlange zu stellen, und nicht die Zeit, die f\u00fcr die Ausf\u00fchrung erforderlich ist.<\/p>\n<p>Ein grundlegender Benchmark sollte:<\/p>\n<ol>\n<li>F\u00fchren Sie den Kernel mehrmals aus, um die Kompilierung und das Aufw\u00e4rmen auszul\u00f6sen.<\/li>\n<li>Synchronisieren, bevor Sie den Timer starten.<\/li>\n<li>F\u00fchren Sie mehrere gemessene Iterationen aus.<\/li>\n<li>Synchronisieren, bevor der Timer gestoppt wird.<\/li>\n<li>Berichten Sie \u00fcber Durchschnitt und Variation \u00fcber Wiederholungen.<\/li>\n<li>Trennen Sie die Daten\u00fcbertragungszeit von der Kernel-Ausf\u00fchrungszeit.<\/li>\n<li>Stellen Sie sicher, dass CPU- und GPU-Versionen gleichwertige Ergebnisse liefern.<\/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>Der Vergleich sollte eine realistische Simulationsgr\u00f6\u00dfe verwenden und den vollst\u00e4ndigen Workflow einschlie\u00dfen, wenn die Kosten f\u00fcr die \u00dcbertragung oder Vorverarbeitung von Bedeutung sind.<\/p>\n<p>Die Arbeitspr\u00e4zisionsmethode kann auch die Laufzeit mit numerischen Fehlern vergleichen. Die <a href=\"https:\/\/docs.sciml.ai\/DiffEqDevDocs\/stable\/alg_dev\/benchmarks\/\" rel=\"nofollow\" target=\"_blank\"> SCIML-Benchmark-Dokumentation <\/a> enth\u00e4lt Beispiele f\u00fcr diesen Bewertungsstil.<\/p>\n<h2>Ein praktischer Entwicklungsworkflow<\/h2>\n<ol>\n<li>Implementieren und \u00fcberpr\u00fcfen Sie eine klare CPU-Referenzversion.<\/li>\n<li>Profil Die Anwendung, um die dominante Operation zu lokalisieren.<\/li>\n<li>Versuchen Sie es mit einer optimierten GPU-Bibliothek, bevor Sie einen Kernel schreiben.<\/li>\n<li>Bewahren Sie h\u00e4ufig wiederverwendete Daten auf dem Ger\u00e4t auf.<\/li>\n<li>Schreiben Sie den einfachsten korrekten Kernel.<\/li>\n<li>\u00dcberpr\u00fcfen Sie es gegen das CPU-Ergebnis.<\/li>\n<li>Messen Sie Speicher\u00fcbertragungen und Ausf\u00fchrung separat.<\/li>\n<li>\u00dcberpr\u00fcfen Sie die Koaleszenz, Belegung, Divergenz und Registernutzung.<\/li>\n<li>Testen Sie den gemeinsamen Speicher oder die Fusion nur, wenn die Profilerstellung sie unterst\u00fctzt.<\/li>\n<li>Vergleichen Sie mehrere Problemgr\u00f6\u00dfen und GPU-Architekturen.<\/li>\n<\/ol>\n<h2>Auswahl eines Frameworks<\/h2>\n<table class=\"custom-table\">\n<tbody>\n<tr>\n<th>Anforderung<\/th>\n<th>M\u00f6glicher Ausgangspunkt<\/th>\n<\/tr>\n<tr>\n<td>Vorhandener GPU-Workflow im Numpy-Stil<\/td>\n<td>Cupy mit integrierten Operationen oder benutzerdefinierten Kernel<\/td>\n<\/tr>\n<tr>\n<td>Kleine Anzahl von Python-Kernel im Cuda-Stil<\/td>\n<td>NUMBA-CUDA, nachdem er seinen aktuellen Support-Status \u00fcberpr\u00fcft hat<\/td>\n<\/tr>\n<tr>\n<td>Simulation, Geometrie und differenzierbare Kernel<\/td>\n<td>Nvidia Warp<\/td>\n<\/tr>\n<tr>\n<td>Maximale NVIDIA-spezifische Steuerung<\/td>\n<td>CUDA C ++<\/td>\n<\/tr>\n<tr>\n<td>AMD-GPU-Ziel mit CUDA-\u00e4hnlicher Quelle<\/td>\n<td>H\u00fcfte und ROCM<\/td>\n<\/tr>\n<tr>\n<td>Portable C++ \u00fcber mehrere Backends<\/td>\n<td>SYCL, KOKKOS oder eine andere Performance-Portabilit\u00e4tsschicht<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n<h2>H\u00e4ufige Fehler beim Programmieren von Kernel<\/h2>\n<ul>\n<li>Verschieben von Daten zwischen CPU und GPU innerhalb jedes Zeitschritts<\/li>\n<li>Ignorieren von Out-of-Bounds-Threads<\/li>\n<li>Verwenden eines ineffizienten Algorithmus und Erwarten von Hardware, um seine Skalierung zu reparieren<\/li>\n<li>Zugriff auf Speicher mit einem nicht koaleszierten Layout<\/li>\n<li>Hinzuf\u00fcgen eines gemeinsamen Speichers, ohne zu messen, ob es hilft<\/li>\n<li>Viele winzige Kernel starten, anstatt Fusion in Betracht zu ziehen<\/li>\n<li>Verwenden \u00fcberm\u00e4\u00dfiger Register oder gemeinsam genutzter Speicher pro Block<\/li>\n<li>Benchmarking asynchroner Code ohne Synchronisation<\/li>\n<li>Vergleichen von Ausgaben ohne numerische Genauigkeit zu \u00fcberpr\u00fcfen<\/li>\n<li>Verwenden von festen Leistungsanspr\u00fcchen \u00fcber verschiedene GPUs und Workloads<\/li>\n<li>Angenommen, die tragbare Quelle bietet tragbare Leistung<\/li>\n<li>Anwenden der automatischen Differenzierung auf nicht unterst\u00fctzte In-Place-Operationen<\/li>\n<\/ul>\n<h2>Verwandte Anleitungen<\/h2>\n<ul>\n<li><a href=\"https:\/\/matforge.org\/gpu-accelerated-scientific-computing-cupy-numba-cudf-compared\/\" rel=\"nofollow\" target=\"_blank\">GPU-beschleunigt Scientific Computing: Cupy, Numba und CUDF im Vergleich <\/a> &#8211; Vergleichen Sie hochrangige Python-Tools f\u00fcr GPU-Workloads.<\/li>\n<li><a href=\"https:\/\/matforge.org\/gpu-acceleration-for-fipy-simulations-cupy-and-numba-integration-guide\/\" rel=\"nofollow\" target=\"_blank\"> GPU-Beschleunigung f\u00fcr FIPY Simulationen: Cupy- und NUMBA-Integrationshandbuch <\/a> &#8211; Erkunden Sie m\u00f6gliche GPU-Komponenten rund um einen FIPY-Workflow.<\/li>\n<li><a href=\"https:\/\/matforge.org\/physics-informed-neural-networks-pinns-for-scientific-simulations\/\" rel=\"nofollow\" target=\"_blank\"> Physikinformierte neuronale Netzwerke f\u00fcr wissenschaftliche Simulationen <\/a> &#8211; \u00dcberpr\u00fcfen Sie die Beziehung zwischen physikalischen Modellen und differenzierbarem Lernen.<\/li>\n<li><a href=\"https:\/\/matforge.org\/when-to-use-fem-fvm-fdm\/\" rel=\"nofollow\" target=\"_blank\">Wann verwendet werden FEM, FVM und FDM<\/a> \u2014 W\u00e4hlen Sie eine r\u00e4umliche Diskretisierung vor der Optimierung der Implementierung.<\/li>\n<li><a href=\"https:\/\/matforge.org\/scientific-computing-performance-profiling-identifying-bottlenecks\/\" rel=\"nofollow\" target=\"_blank\">Performance Profiling: Identifizierung von Engp\u00e4ssen in wissenschaftlichen Code <\/a> &#8211; Suchen Sie die Einschr\u00e4nkungen der CPU, des Speichers, der \u00dcbertragung und der E \/ A.<\/li>\n<li><a href=\"https:\/\/matforge.org\/hpc-python-workflows-from-laptop-to-supercomputer\/\" rel=\"nofollow\" target=\"_blank\">HPC-Python-Workflows: vom Laptop nach Supercomputer <\/a> &#8211; Planen Sie den Fortschritt von der lokalen Entwicklung zu gro\u00dfen Systemen.<\/li>\n<\/ul>\n<h2>Schlussfolgerung<\/h2>\n<p>Benutzerdefinierte GPU-Kernel erm\u00f6glichen es wissenschaftlichen Entwicklern, Partikel, Gitterzellen, Materialregeln und andere physikalische Operationen direkt in massiv parallele Funktionen zu \u00fcbersetzen. Sie bieten Kontrolle \u00fcber Indexierung, Speicherzugriff, Synchronisation, Kernelfusion und ger\u00e4tespezifische Optimierung.<\/p>\n<p>Cupy bietet sowohl hochrangige Array-Operationen als auch benutzerdefinierte Kernel-Schnittstellen. NUMBA-CUDA bietet ein CUDA-\u00e4hnliches Modell in Python, w\u00e4hrend Nvidia Warp Simulationsprimitive, Kacheloperationen und automatische Differenzierung hinzuf\u00fcgt. CUDA und HIP bieten C++-Steuerung auf niedrigerer Ebene, und leistungstragbare Frameworks unterst\u00fctzen breitere Hardware-Ziele.<\/p>\n<p>Die gr\u00f6\u00dften Leistungssteigerungen ergeben sich selten allein aus einer sich \u00e4ndernden Syntax. Sie stammen aus der Auswahl eines parallelen Algorithmus, der Speicherung von Daten auf dem Ger\u00e4t, der Reduzierung des Speicherverkehrs, der Verwendung eines geeigneten Datenlayouts und der Eliminierung unn\u00f6tiger Zwischenoperationen.<\/p>\n<p>Ein benutzerdefinierter Kernel sollte erst entwickelt werden, nachdem die Profilerstellung einen echten Engpass identifiziert hat. Es sollte anhand einer Referenzl\u00f6sung validiert und mit Synchronisation, repr\u00e4sentativen Problemgr\u00f6\u00dfen und vollst\u00e4ndiger Abrechnung der Speicher\u00fcbertragungskosten verglichen werden.<\/p>\n<p>Wenn diese Bedingungen erf\u00fcllt sind, k\u00f6nnen Python-basierte Kernel-Tools ernsthafte Physiksimulationen unterst\u00fctzen, ohne die Forscher zu zwingen, jeden Teil der Anwendung in C ++ auf niedriger Ebene zu verschieben.<\/p>\n<\/article>\n","protected":false,"raw":"<article>\n<p>Die GPU-Beschleunigung im wissenschaftlichen Rechnen beginnt h\u00e4ufig mit Bibliotheken auf hoher Ebene. Ein Forscher ersetzt Numpy durch Cupy, verwendet eine beschleunigte FFT oder ruft eine GPU-f\u00e4hige lineare Algebra-Routine auf. Dies kann zu einer wesentlichen Verbesserung f\u00fchren, ohne dass detaillierte Kenntnisse der GPU-Hardware erforderlich sind.<\/p>\n<p>Die benutzerdefinierte GPU-Kernel-Programmierung geht eine Ebene tiefer. Anstatt eine vordefinierte Operation aufzurufen, schreibt der Entwickler die Funktion, die jeder GPU-Thread ausf\u00fchrt. Dies bietet Kontrolle \u00fcber Thread-Indizierung, Datenlayout, Speicher\u00fcbertragungen, Synchronisation, gemeinsam genutzten Speicher und die Anzahl der w\u00e4hrend jedes Kernels gestarteten Operationen.<\/p>\n<p>Diese Steuerung ist n\u00fctzlich, wenn ein physikalisches Modell Zugriffsmuster enth\u00e4lt, die nicht effizient durch gew\u00f6hnliche Array-Operationen ausgedr\u00fcckt werden k\u00f6nnen. Partikelwechselwirkungen, Finite-Differenz-Schablonen, Kollisionsregeln, gitterbasierte Fluidmodelle, Materialaktualisierungen und Geometrieabfragen sind g\u00e4ngige Beispiele.<\/p>\n<p>Ein benutzerdefinierter Kernel ist nicht automatisch schneller als eine optimierte Bibliothek. Es wird wertvoll, wenn der Algorithmus gen\u00fcgend parallele Arbeit aufzeigt und der Entwickler den Speicherverkehr reduzieren, Operationen kombinieren oder die Berechnung an die GPU-Architektur anpassen kann.<\/p>\n<h2>Was ist ein GPU-Kernel?<\/h2>\n<p>Ein GPU-Kernel ist eine Funktion, die von vielen leichten Threads ausgef\u00fchrt wird. Jeder Thread behandelt normalerweise ein Teilchen, eine Gitterzelle, ein Element, eine Fl\u00e4che, ein Pixel oder einen Eintrag in einem Array.<\/p>\n<p>Das Host-Programm startet den Kernel mit einem Raster von Thread-Bl\u00f6cken:<\/p>\n<pre><code>kernel[blocks, threads_per_block](arguments)<\/code><\/pre>\n<p>Jeder Thread bestimmt, welche Daten er aus seinen Block- und Thread-Indizes verarbeiten soll. Ein typischer eindimensionaler Index ist:<\/p>\n<pre><code>index =\n    block_index * block_size\n    + thread_index<\/code><\/pre>\n<p>Threads sind in Gruppen unterteilt, die als Bl\u00f6cke bezeichnet werden. Threads in einem Block k\u00f6nnen Daten \u00fcber gemeinsam genutzten Speicher synchronisieren und austauschen. Es wird allgemein erwartet, dass separate Bl\u00f6cke unabh\u00e4ngig ausgef\u00fchrt werden.<\/p>\n<h2>Wenn benutzerdefinierte Kernel es wert sind, geschrieben zu werden<\/h2>\n<p>GPU-Bibliotheken auf hoher Ebene sollten in der Regel die erste Option sein. Sie bieten bereits optimierte Matrixmultiplikation, FFTs, Reduktionen, sp\u00e4rliche Operationen, Zufallszahlengenerierung und viele elementweise Funktionen.<\/p>\n<p>Ein benutzerdefinierter Kernel wird n\u00fctzlich, wenn:<\/p>\n<ul>\n<li>Die Simulation verwendet ein nicht standardm\u00e4\u00dfiges Speicherzugriffsmuster.<\/li>\n<li>Mehrere kleine Operationen k\u00f6nnen zu einem Durchgang \u00fcber Speicher verschmolzen werden.<\/li>\n<li>Benachbarte Threads ben\u00f6tigen wiederholt dieselben lokalen Daten.<\/li>\n<li>Eine spezielle Partikel-, Kollisions- oder Schablonenregel dominiert die Laufzeit.<\/li>\n<li>Zwischen Arrays verbrauchen zu viel Ger\u00e4tespeicher.<\/li>\n<li>Die Simulation ben\u00f6tigt eine benutzerdefinierte differenzierbare Operation.<\/li>\n<li>Eine vorhandene Bibliothek kann die erforderliche Grenz- oder Materiallogik nicht effizient ausdr\u00fccken.<\/li>\n<\/ul>\n<p>Profilieren Sie das Programm, bevor Sie einen Kernel schreiben. Die Optimierung einer optisch komplexen Funktion, die nur einen kleinen Teil der Laufzeit ausmacht, f\u00fchrt nicht zu einer aussagekr\u00e4ftigen Gesamtbeschleunigung.<\/p>\n<h2>GPU-Programmiermodelle<\/h2>\n<p>Wissenschaftliche Software kann \u00fcber mehrere Abstraktionsebenen auf GPUs zugreifen.<\/p>\n<table class=\"custom-table\">\n<tbody><tr>\n<th>Sich n\u00e4hern<\/th>\n<th>Beispiele<\/th>\n<th>Hauptst\u00e4rke<\/th>\n<th>Hauptkompromiss<\/th>\n<\/tr>\n<tr>\n<td>Direktives Offloading<\/td>\n<td>OpenMP, OpenACC<\/td>\n<td>Inkrementelle Beschleunigung des vorhandenen C, C++ oder FORTRAN-Codes<\/td>\n<td>Weniger direkte Kontrolle \u00fcber generierte Kernel<\/td>\n<\/tr>\n<tr>\n<td>Herstellerorientierte Kernel-Programmierung<\/td>\n<td>Cuda, H\u00fcfte<\/td>\n<td>Detaillierte Kontrolle \u00fcber die GPU-Ausf\u00fchrung<\/td>\n<td>H\u00f6here Implementierungs- und Wartungskosten<\/td>\n<\/tr>\n<tr>\n<td>Leistungstragbares C ++<\/td>\n<td>Sycl, Kokkos, Alpaka<\/td>\n<td>Ein Quellmodell f\u00fcr mehrere Hardware-Backends<\/td>\n<td>Portabilit\u00e4t garantiert nicht \u00fcberall gleiche Leistung<\/td>\n<\/tr>\n<tr>\n<td>Python-GPU-Arrays und -Kernel<\/td>\n<td>Cupy, Numba-Cuda, Nvidia Warp<\/td>\n<td>Schnelle Entwicklung mit Zugriff auf kompilierten GPU-Code<\/td>\n<td>Framework-spezifische Einschr\u00e4nkungen und Zusammenstellungsverhalten<\/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 Einf\u00fchrung in die GPU-Programmierung Modelle <\/a> bieten einen umfassenderen Vergleich von CUDA, HIP, OpenCL, Sycl, Kokkos, OpenMP und verwandten Ans\u00e4tzen.<\/p>\n<h2>Cuda<\/h2>\n<p>CUDA bietet direkten Zugriff auf die NVIDIA-GPU-Programmierung \u00fcber C++, FORTRAN und mehrere Sprachbindungen. Es stellt Gitter, Bl\u00f6cke, Warps, gemeinsam genutzten Speicher, Streams, Ereignisse, Ger\u00e4tebibliotheken und hardwarespezifische Optimierungstools bereit.<\/p>\n<p>CUDA wird h\u00e4ufig ausgew\u00e4hlt, wenn maximale NVIDIA-spezifische Leistung und ausgereifte Tools wichtiger sind als die Portabilit\u00e4t der Hardware. Die Hauptkosten sind das manuelle Speichermanagement, ein Programmiermodell auf niedrigerer Ebene und die Abh\u00e4ngigkeit von der NVIDIA-Plattform.<\/p>\n<h2>H\u00fcfte und ROCM<\/h2>\n<p>HIP bietet ein CUDA-\u00e4hnliches C ++ - Programmiermodell im ROCM-\u00d6kosystem von AMD. Eine gro\u00dfe Menge an CUDA-Code kann an HIP angepasst werden, aber die Portabilit\u00e4t der Quellen bedeutet nicht, dass eine Bin\u00e4rdatei auf jeder GPU unver\u00e4ndert ausgef\u00fchrt wird.<\/p>\n<p>F\u00fcr jede Zielarchitektur ist m\u00f6glicherweise weiterhin eine separate Kompilierung, Pr\u00fcfung und Leistungsoptimierung erforderlich. Details finden Sie in der <a href=\"https:\/\/rocm.docs.amd.com\/projects\/HIP\/en\/latest\/what_is_hip.html\" rel=\"nofollow\" target=\"_blank\">Official Hip Documentation <\/a>.<\/p>\n<p>HIP ist eine praktische Option, wenn AMD-GPUs unterst\u00fctzt werden m\u00fcssen oder wenn ein Projekt die Abh\u00e4ngigkeit von einem einzelnen Beschleunigerhersteller verringern m\u00f6chte, w\u00e4hrend ein CUDA-\u00e4hnlicher Programmierstil beibehalten wird.<\/p>\n<h2>Cupy benutzerdefinierte Kernel<\/h2>\n<p>Cupy ist am bekanntesten f\u00fcr seine numpy-kompatiblen Arrays und beschleunigten numerischen Funktionen. Es unterst\u00fctzt auch die benutzerdefinierte GPU-Programmierung.<\/p>\n<p>Entwickler k\u00f6nnen verwenden:<\/p>\n<ul>\n<li><code>ElementwiseKernel<\/code> f\u00fcr benutzerdefinierte elementweise Ausdr\u00fccke<\/li>\n<li><code>ReductionKernel<\/code> f\u00fcr benutzerdefinierte Reduzierungen<\/li>\n<li><code>RawKernel<\/code> f\u00fcr in CUDA C++ geschriebene Kernel<\/li>\n<li><code>cupyx.jit.rawkernel<\/code> f\u00fcr JIT-Kernel-Definitionen im Python-Stil<\/li>\n<\/ul>\n<p>Ein einfacher Cupy-Raw-Kernel kann wie folgt definiert werden:<\/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 ist n\u00fctzlich, wenn die meisten Anwendungen bereits GPU-Arrays verwenden und nur ausgew\u00e4hlte Operationen benutzerdefinierte Kernel ben\u00f6tigen.<\/p>\n<h2>Numba-Cuda<\/h2>\n<p>Numba-Cuda kompiliert eine eingeschr\u00e4nkte Teilmenge von Python in GPU-Kernel. Das Programmiermodell folgt CUDA C. Die Entwickler definieren eine Funktion mit einem Kerneldekorator, berechnen einen Thread-Index und starten die Funktion mit einer Raster- und Blockkonfiguration.<\/p>\n<p>Die aktive Projektdokumentation ist unter <a href=\"https:\/\/nvidia.github.io\/numba-cuda\/\" rel=\"nofollow\" target=\"_blank\">Numba-cuda<\/a> verf\u00fcgbar. Da sich der Entwicklungsstatus \u00e4ndern kann, sollten Teams, die langlebige Software erstellen, die aktuellen Wartungs- und Migrationsrichtlinien \u00fcberpr\u00fcfen, bevor sie sich f\u00fcr das Framework engagieren.<\/p>\n<h2>ein Numba-Teilchenkern<\/h2>\n<p>Das folgende Lehrbeispiel berechnet die Gravitationsbeschleunigung mit direkten All-Pairs-Interaktionen:<\/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>Dieser Kernel hat rechnerische Komplexit\u00e4t von <code>O(N\u00b2)<\/code> . Es ist zur Erkl\u00e4rung der Thread-Mapping geeignet, aber es ist kein effizienter Produktionsalgorithmus f\u00fcr sehr gro\u00dfe Partikelzahlen.<\/p>\n<p>Gro\u00dfe Gravitations- oder Partikelsysteme erfordern h\u00e4ufig eine Baummethode, eine schnelle Multipolmethode, r\u00e4umliche Beh\u00e4lter oder Nachbarlisten. Eine GPU kann das Skalierungsproblem eines ineffizienten Algorithmus nicht beseitigen.<\/p>\n<h2>Nvidia Warp<\/h2>\n<p>NVIDIA Warp erm\u00f6glicht Entwicklern, stark typisierte Kernel mit Python-Syntax zu definieren. Warp kompiliert diese Funktionen in CPU- oder CUDA-Code und bietet spezielle Primitive f\u00fcr Geometrie, Simulation, sp\u00e4rliche Operationen, endliche Elemente, Optimierung und automatische Differenzierung.<\/p>\n<p>Die aktuelle <a href=\"https:\/\/developer.nvidia.com\/warp-python\" rel=\"nofollow\" target=\"_blank\">Nvidia Warp-Produktseite<\/a> und <a href=\"https:\/\/github.com\/nvidia\/warp\" rel=\"nofollow\" target=\"_blank\">Warp-GitHub-Repository<\/a> Enth\u00e4lt Beispiele f\u00fcr Partikel, Fl\u00fcssigkeiten, Netze, Optimierungen, differenzierbare Simulationen und Tile-based GPU-Programmierung.<\/p>\n<h2>der gleiche Partikelkern in 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 bestimmt intern eine geeignete Startkonfiguration f\u00fcr einen gew\u00f6hnlichen Kernel-Start. Fortgeschrittene Benutzer k\u00f6nnen bei Bedarf weiterhin mit Blockdimensionen, Befehlsdiagrammen, gekachelten Ausf\u00fchrungen und speziellen Ger\u00e4teoperationen arbeiten.<\/p>\n<h2>Grid-Stride-Schleifen<\/h2>\n<p>Ein Kernel ben\u00f6tigt nicht f\u00fcr jedes Element einen fest zugewiesenen Thread. Eine Grid-Stride-Schleife erm\u00f6glicht es jedem Thread, mehrere Eintr\u00e4ge zu verarbeiten:<\/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>Dieses Muster trennt die Anzahl der gestarteten Threads von der Gesamtdatengr\u00f6\u00dfe. Es ist n\u00fctzlich, wenn sehr gro\u00dfe Arrays verarbeitet oder eine feste Startkonfiguration wiederverwendet werden.<\/p>\n<p>Eine praktische Einf\u00fchrung in dieses Muster finden Sie in <a href=\"https:\/\/thedatafrog.com\/en\/articles\/cuda-kernel-python\/\" rel=\"nofollow\" target=\"_blank\"> Das Numba- und CUDA-Tutorial des Datenfroschs <\/a>.<\/p>\n<h2>Ged\u00e4chtnis verschmelzen<\/h2>\n<p>GPU-Kernel sind h\u00e4ufig eher durch Speicherbandbreite als durch den arithmetischen Durchsatz begrenzt. Threads in einem Warp sollten idealerweise auf die nahe gelegenen Speicheradressen zugreifen, damit die Hardware ihre Anfragen kombinieren kann.<\/p>\n<p>Betrachten Sie Partikeldaten als gespeichert:<\/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>Dieses Array-of-Struktur-Layout kann f\u00fcr objektorientierten Code zweckm\u00e4\u00dfig sein. Ein Layout mit Struktur-of-Arrays speichert separate zusammenh\u00e4ngende Arrays:<\/p>\n<pre><code>x_positions[]\ny_positions[]\nz_positions[]\nmasses[]<\/code><\/pre>\n<p>Das zweite Layout kann eine bessere Koaleszenz liefern, wenn jeder Thread das gleiche Feld f\u00fcr ein anderes Teilchen liest. Das beste Layout h\u00e4ngt immer noch davon ab, auf welche Felder zusammen zugegriffen wird.<\/p>\n<p>Vektortypen korrigieren nicht automatisch ein schlechtes Zugriffsmuster. Entwickler sollten die tats\u00e4chlichen Adressen \u00fcberpr\u00fcfen, die von benachbarten Threads angefordert werden.<\/p>\n<h2>Geteilter Speicher<\/h2>\n<p>Shared Memory ist ein kleiner Speicherbereich mit geringer Latenz, auf den Threads im selben Block zugegriffen werden kann. Es kann wiederholte Lesevorg\u00e4nge aus dem globalen Speicher reduzieren.<\/p>\n<p>Eine eindimensionale Schablone kann einen Werteblock und seinen Halo in den gemeinsamen Speicher laden:<\/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>Der Aufruf von <code>cuda.syncthreads()<\/code> stellt sicher, dass jeder Thread seine Werte l\u00e4dt, bevor die Schablonenberechnung beginnt.<\/p>\n<p>Shared Memory sollte nicht automatisch verwendet werden. \u00dcberm\u00e4\u00dfige Zuweisung von Shared-Memory kann die Belegung verringern, die Synchronisierungskosten erh\u00f6hen und den Kernel langsamer machen. Profilierung erforderlich.<\/p>\n<h2>Kernfusion<\/h2>\n<p>Separate Array-Operationen erzeugen h\u00e4ufig mehrere Kernel-Starts und Zwischenarrays:<\/p>\n<pre><code>velocity += dt * acceleration\nposition += dt * velocity\nenergy = compute_energy(position, velocity)<\/code><\/pre>\n<p>Ein fusionierter Kernel kann alle drei Aktualisierungen berechnen, w\u00e4hrend die notwendigen Werte in Registern verbleiben. Dies reduziert den Start-Overhead und den globalen Speicherverkehr.<\/p>\n<p>Fusion ist am vorteilhaftesten, wenn die Operationen einfach und speichergebunden sind. Zu viel Arbeit kann die Nutzung des Registers erh\u00f6hen, die Belegung verringern und die Wartung des Kernels erschweren.<\/p>\n<h2>Blockgr\u00f6\u00dfe w\u00e4hlen<\/h2>\n<p>Werte wie 128 oder 256 Threads pro Block sind vern\u00fcnftige Ausgangspunkte, nicht universelle Optima.<\/p>\n<p>Die beste Blockgr\u00f6\u00dfe h\u00e4ngt ab von:<\/p>\n<ul>\n<li>Register pro Thread verwendet<\/li>\n<li>Shared Memory pro Block verwendet<\/li>\n<li>Zweigdivergenz<\/li>\n<li>Anleitungsmix<\/li>\n<li>Speicherzugriffsverhalten<\/li>\n<li>Die Ziel-GPU-Architektur<\/li>\n<\/ul>\n<p>Eine Regel wie das Starten von mindestens doppelt so vielen Bl\u00f6cken wie Streaming-Multiprozessoren kann ein n\u00fctzliches anf\u00e4ngliches Experiment sein, garantiert jedoch keine maximale Leistung. Belegungsrechner und Profilierungswerkzeuge sollten die endg\u00fcltige Konfiguration leiten.<\/p>\n<h2>differenzierbare Physik<\/h2>\n<p>Die differenzierbare Simulation berechnet, wie sich eine Ausgabe in Bezug auf Eingaben wie Materialeigenschaften, Kr\u00e4fte, Geometrie oder Anfangsbedingungen \u00e4ndert.<\/p>\n<p>Warp kann unterst\u00fctzte Kerneloperationen aufzeichnen und die automatische Differenzierung des Reverse-Modus ausf\u00fchren. Dies kann f\u00fcr inverse Probleme, Designoptimierung, Parametersch\u00e4tzung und Integration in maschinelle Lern-Workflows verwendet werden.<\/p>\n<p>Ein vereinfachtes Muster verwendet ein Band:<\/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>Eine automatische Differenzierung ist nicht f\u00fcr jeden Kernel garantiert. In-Place-Overwrites, nichtdeterministische Atome, externer nativer Code, diskontinuierliche Logik und nicht unterst\u00fctzte Operationen erfordern m\u00f6glicherweise eine Neuformulierung oder benutzerdefinierte Gradienten.<\/p>\n<p>Diese F\u00e4higkeit verbindet sich direkt mit Anwendungen wie Adjoint Optimization und <a href=\"https:\/\/matforge.org\/physics-informed-neural-networks-pinns-for-scientific-simulations\/\" rel=\"nofollow\" target=\"_blank\"> Physikinformierte neuronale Netze <\/a>.<\/p>\n<h2>Wenn eine GPU m\u00f6glicherweise nicht schneller ist<\/h2>\n<p>Es gibt keinen universellen Schwellenwert f\u00fcr die Problemgr\u00f6\u00dfe, bei dem eine GPU schneller wird als eine CPU. Der Crossover h\u00e4ngt von Hardware, Pr\u00e4zision, Datenbewegung, Algorithmenstruktur, Compilerqualit\u00e4t und der H\u00e4ufigkeit der Wiederverwendung derselben ger\u00e4teresidenten Daten ab.<\/p>\n<table class=\"custom-table\">\n<tbody><tr>\n<th>Arbeitsanfall<\/th>\n<th>GPU-Eignung<\/th>\n<th>Haupt\u00fcberlegung<\/th>\n<\/tr>\n<tr>\n<td>Gro\u00dfe Partikelaktualisierung<\/td>\n<td>Oft stark<\/td>\n<td>viele \u00e4hnliche unabh\u00e4ngige Operationen<\/td>\n<\/tr>\n<tr>\n<td>Gro\u00dfe regul\u00e4re Schablone<\/td>\n<td>Oft stark<\/td>\n<td>Vorhersehbarer paralleler Speicherzugriff<\/td>\n<\/tr>\n<tr>\n<td>Dichte lineare Algebra<\/td>\n<td>Stark bei der Verwendung optimierter Bibliotheken<\/td>\n<td>Hohe arithmetische Intensit\u00e4t<\/td>\n<\/tr>\n<tr>\n<td>Kleine Simulation<\/td>\n<td>Hardwareabh\u00e4ngig<\/td>\n<td>Start und Transfer-Overhead k\u00f6nnen dominieren<\/td>\n<\/tr>\n<tr>\n<td>Unregelm\u00e4\u00dfige Graphendurchquerung<\/td>\n<td>Gemischt<\/td>\n<td>Divergenz und unvorhersehbarer Speicherzugriff<\/td>\n<\/tr>\n<tr>\n<td>I\/O-gebundener Workflow<\/td>\n<td>in der Regel begrenzt<\/td>\n<td>Die GPU kann einen Speicher- oder Netzwerkengpass nicht entfernen<\/td>\n<\/tr>\n<tr>\n<td>stark sequentieller Algorithmus<\/td>\n<td>normalerweise schwach<\/td>\n<td>Unzureichende unabh\u00e4ngige Arbeit<\/td>\n<\/tr>\n<\/tbody><\/table>\n<p>Der richtige Workflow besteht darin, die CPU-Version, die GPU-Bibliotheksversion und die benutzerdefinierten Kernel-Version mit repr\u00e4sentativen Daten zu messen. <a href=\"https:\/\/matforge.org\/scientific-computing-performance-profiling-identifying-bottlenecks\/\" rel=\"nofollow\" target=\"_blank\">Forscherhandbuch f\u00fcr wissenschaftliche Rechenleistung<\/a> Erkl\u00e4rt, wie Sie den tats\u00e4chlichen Engpass identifizieren, bevor Sie ihn optimieren.<\/p>\n<h2>Vermeidung von Transfer-Overhead<\/h2>\n<p>Wiederholte \u00dcbertragungen zwischen Host- und Ger\u00e4tespeicher k\u00f6nnen den Nutzen der GPU-Berechnung beseitigen.<\/p>\n<p>Eine ineffiziente Schleife kann diesem Muster folgen:<\/p>\n<ol>\n<li>Daten auf die GPU kopieren.<\/li>\n<li>F\u00fchren Sie einen kleinen Kernel aus.<\/li>\n<li>Kopieren Sie das Ergebnis in die CPU.<\/li>\n<li>\u00c4ndern Sie es auf der CPU.<\/li>\n<li>Kopieren Sie es wieder auf die GPU.<\/li>\n<\/ol>\n<p>Ein besseres Design beh\u00e4lt den Simulationszustand \u00fcber viele Zeitschritte auf dem Ger\u00e4t und \u00fcbertr\u00e4gt nur die Ausgabe, die f\u00fcr Visualisierung, Checkpointing oder Analyse erforderlich ist.<\/p>\n<p>Asynchrone Streams, festgesteckter Host-Speicher, \u00fcberlappende Kommunikation und Unified-Memory-Funktionen k\u00f6nnen helfen, sollten jedoch erst eingef\u00fchrt werden, nachdem die normalen \u00dcbertragungskosten gemessen wurden.<\/p>\n<h2>Die XLB-Fallstudie<\/h2>\n<p>Das XLB-Projekt von Autodesk Research bietet ein n\u00fctzliches Beispiel f\u00fcr die Python-native GPU-Simulation. XLB ist eine Open-Source-Lattice-Boltzmann-Bibliothek mit mehreren rechnerischen Backends, darunter Nvidia Warp.<\/p>\n<p>In von Autodesk Research und NVIDIA gemeldeten Benchmark-Konfigurationen erreichte das Warp-Backend eine Leistung in der N\u00e4he der verglichenen C ++ \/ OpenCL FluidX3D-Implementierung f\u00fcr einen bestimmten Lid-gesteuerten Hohlraumfall. Ein separater Vergleich ergab eine ungef\u00e4hre achtfache Beschleunigung \u00fcber das JAX-Backend von XLB bei ausgew\u00e4hlten Hardware und Konfigurationen.<\/p>\n<p>Das Team zeigte auch einen Out-of-Core-Ansatz f\u00fcr einen GH200-Cluster mit acht Knoten mit einer Dom\u00e4ne von ungef\u00e4hr 50 Milliarden Gitterzellen. Diese Ergebnisse gelten f\u00fcr den gemeldeten Solver, Problem, Hardware und Implementierungsoptionen. Sie sollten nicht als allgemeine Beschleunigungsgarantien f\u00fcr Python-Kernel behandelt werden.<\/p>\n<p>Die vollst\u00e4ndige Fallstudie ist in <a href=\"https:\/\/developer.nvidia.com\/blog\/autodesk-research-brings-warp-speed-to-computational-fluid-dynamics-on-nvidia-gh200\/\" rel=\"nofollow\" target=\"_blank\"> verf\u00fcgbar. Autodesk-Forschung bringt die Warp-Geschwindigkeit in die rechnerische Fluiddynamik auf NVIDIA GH200 <\/a>. Der aktuelle Projektcode ist im <a href=\"https:\/\/github.com\/Autodesk\/XLB\" rel=\"nofollow\" target=\"_blank\"> Autodesk XLB-Repository <\/a> verf\u00fcgbar.<\/p>\n<h2>Wie man einen Kernel richtig bewertet<\/h2>\n<p>GPU-Operationen sind normalerweise asynchron. Das Messen nur der Python-Funktionsdauer kann die Zeit angeben, die erforderlich ist, um den Kernel in die Warteschlange zu stellen, und nicht die Zeit, die f\u00fcr die Ausf\u00fchrung erforderlich ist.<\/p>\n<p>Ein grundlegender Benchmark sollte:<\/p>\n<ol>\n<li>F\u00fchren Sie den Kernel mehrmals aus, um die Kompilierung und das Aufw\u00e4rmen auszul\u00f6sen.<\/li>\n<li>Synchronisieren, bevor Sie den Timer starten.<\/li>\n<li>F\u00fchren Sie mehrere gemessene Iterationen aus.<\/li>\n<li>Synchronisieren, bevor der Timer gestoppt wird.<\/li>\n<li>Berichten Sie \u00fcber Durchschnitt und Variation \u00fcber Wiederholungen.<\/li>\n<li>Trennen Sie die Daten\u00fcbertragungszeit von der Kernel-Ausf\u00fchrungszeit.<\/li>\n<li>Stellen Sie sicher, dass CPU- und GPU-Versionen gleichwertige Ergebnisse liefern.<\/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>Der Vergleich sollte eine realistische Simulationsgr\u00f6\u00dfe verwenden und den vollst\u00e4ndigen Workflow einschlie\u00dfen, wenn die Kosten f\u00fcr die \u00dcbertragung oder Vorverarbeitung von Bedeutung sind.<\/p>\n<p>Die Arbeitspr\u00e4zisionsmethode kann auch die Laufzeit mit numerischen Fehlern vergleichen. Die <a href=\"https:\/\/docs.sciml.ai\/DiffEqDevDocs\/stable\/alg_dev\/benchmarks\/\" rel=\"nofollow\" target=\"_blank\"> SCIML-Benchmark-Dokumentation <\/a> enth\u00e4lt Beispiele f\u00fcr diesen Bewertungsstil.<\/p>\n<h2>Ein praktischer Entwicklungsworkflow<\/h2>\n<ol>\n<li>Implementieren und \u00fcberpr\u00fcfen Sie eine klare CPU-Referenzversion.<\/li>\n<li>Profil Die Anwendung, um die dominante Operation zu lokalisieren.<\/li>\n<li>Versuchen Sie es mit einer optimierten GPU-Bibliothek, bevor Sie einen Kernel schreiben.<\/li>\n<li>Bewahren Sie h\u00e4ufig wiederverwendete Daten auf dem Ger\u00e4t auf.<\/li>\n<li>Schreiben Sie den einfachsten korrekten Kernel.<\/li>\n<li>\u00dcberpr\u00fcfen Sie es gegen das CPU-Ergebnis.<\/li>\n<li>Messen Sie Speicher\u00fcbertragungen und Ausf\u00fchrung separat.<\/li>\n<li>\u00dcberpr\u00fcfen Sie die Koaleszenz, Belegung, Divergenz und Registernutzung.<\/li>\n<li>Testen Sie den gemeinsamen Speicher oder die Fusion nur, wenn die Profilerstellung sie unterst\u00fctzt.<\/li>\n<li>Vergleichen Sie mehrere Problemgr\u00f6\u00dfen und GPU-Architekturen.<\/li>\n<\/ol>\n<h2>Auswahl eines Frameworks<\/h2>\n<table class=\"custom-table\">\n<tbody><tr>\n<th>Anforderung<\/th>\n<th>M\u00f6glicher Ausgangspunkt<\/th>\n<\/tr>\n<tr>\n<td>Vorhandener GPU-Workflow im Numpy-Stil<\/td>\n<td>Cupy mit integrierten Operationen oder benutzerdefinierten Kernel<\/td>\n<\/tr>\n<tr>\n<td>Kleine Anzahl von Python-Kernel im Cuda-Stil<\/td>\n<td>NUMBA-CUDA, nachdem er seinen aktuellen Support-Status \u00fcberpr\u00fcft hat<\/td>\n<\/tr>\n<tr>\n<td>Simulation, Geometrie und differenzierbare Kernel<\/td>\n<td>Nvidia Warp<\/td>\n<\/tr>\n<tr>\n<td>Maximale NVIDIA-spezifische Steuerung<\/td>\n<td>CUDA C ++<\/td>\n<\/tr>\n<tr>\n<td>AMD-GPU-Ziel mit CUDA-\u00e4hnlicher Quelle<\/td>\n<td>H\u00fcfte und ROCM<\/td>\n<\/tr>\n<tr>\n<td>Portable C++ \u00fcber mehrere Backends<\/td>\n<td>SYCL, KOKKOS oder eine andere Performance-Portabilit\u00e4tsschicht<\/td>\n<\/tr>\n<\/tbody><\/table>\n<h2>H\u00e4ufige Fehler beim Programmieren von Kernel<\/h2>\n<ul>\n<li>Verschieben von Daten zwischen CPU und GPU innerhalb jedes Zeitschritts<\/li>\n<li>Ignorieren von Out-of-Bounds-Threads<\/li>\n<li>Verwenden eines ineffizienten Algorithmus und Erwarten von Hardware, um seine Skalierung zu reparieren<\/li>\n<li>Zugriff auf Speicher mit einem nicht koaleszierten Layout<\/li>\n<li>Hinzuf\u00fcgen eines gemeinsamen Speichers, ohne zu messen, ob es hilft<\/li>\n<li>Viele winzige Kernel starten, anstatt Fusion in Betracht zu ziehen<\/li>\n<li>Verwenden \u00fcberm\u00e4\u00dfiger Register oder gemeinsam genutzter Speicher pro Block<\/li>\n<li>Benchmarking asynchroner Code ohne Synchronisation<\/li>\n<li>Vergleichen von Ausgaben ohne numerische Genauigkeit zu \u00fcberpr\u00fcfen<\/li>\n<li>Verwenden von festen Leistungsanspr\u00fcchen \u00fcber verschiedene GPUs und Workloads<\/li>\n<li>Angenommen, die tragbare Quelle bietet tragbare Leistung<\/li>\n<li>Anwenden der automatischen Differenzierung auf nicht unterst\u00fctzte In-Place-Operationen<\/li>\n<\/ul>\n<h2>Verwandte Anleitungen<\/h2>\n<ul>\n<li><a href=\"https:\/\/matforge.org\/gpu-accelerated-scientific-computing-cupy-numba-cudf-compared\/\" rel=\"nofollow\" target=\"_blank\">GPU-beschleunigt Scientific Computing: Cupy, Numba und CUDF im Vergleich <\/a> - Vergleichen Sie hochrangige Python-Tools f\u00fcr GPU-Workloads.<\/li>\n<li><a href=\"https:\/\/matforge.org\/gpu-acceleration-for-fipy-simulations-cupy-and-numba-integration-guide\/\" rel=\"nofollow\" target=\"_blank\"> GPU-Beschleunigung f\u00fcr FIPY Simulationen: Cupy- und NUMBA-Integrationshandbuch <\/a> - Erkunden Sie m\u00f6gliche GPU-Komponenten rund um einen FIPY-Workflow.<\/li>\n<li><a href=\"https:\/\/matforge.org\/physics-informed-neural-networks-pinns-for-scientific-simulations\/\" rel=\"nofollow\" target=\"_blank\"> Physikinformierte neuronale Netzwerke f\u00fcr wissenschaftliche Simulationen <\/a> - \u00dcberpr\u00fcfen Sie die Beziehung zwischen physikalischen Modellen und differenzierbarem Lernen.<\/li>\n<li><a href=\"https:\/\/matforge.org\/when-to-use-fem-fvm-fdm\/\" rel=\"nofollow\" target=\"_blank\">Wann verwendet werden FEM, FVM und FDM<\/a> \u2014 W\u00e4hlen Sie eine r\u00e4umliche Diskretisierung vor der Optimierung der Implementierung.<\/li>\n<li><a href=\"https:\/\/matforge.org\/scientific-computing-performance-profiling-identifying-bottlenecks\/\" rel=\"nofollow\" target=\"_blank\">Performance Profiling: Identifizierung von Engp\u00e4ssen in wissenschaftlichen Code <\/a> - Suchen Sie die Einschr\u00e4nkungen der CPU, des Speichers, der \u00dcbertragung und der E \/ A.<\/li>\n<li><a href=\"https:\/\/matforge.org\/hpc-python-workflows-from-laptop-to-supercomputer\/\" rel=\"nofollow\" target=\"_blank\">HPC-Python-Workflows: vom Laptop nach Supercomputer <\/a> - Planen Sie den Fortschritt von der lokalen Entwicklung zu gro\u00dfen Systemen.<\/li>\n<\/ul>\n<h2>Schlussfolgerung<\/h2>\n<p>Benutzerdefinierte GPU-Kernel erm\u00f6glichen es wissenschaftlichen Entwicklern, Partikel, Gitterzellen, Materialregeln und andere physikalische Operationen direkt in massiv parallele Funktionen zu \u00fcbersetzen. Sie bieten Kontrolle \u00fcber Indexierung, Speicherzugriff, Synchronisation, Kernelfusion und ger\u00e4tespezifische Optimierung.<\/p>\n<p>Cupy bietet sowohl hochrangige Array-Operationen als auch benutzerdefinierte Kernel-Schnittstellen. NUMBA-CUDA bietet ein CUDA-\u00e4hnliches Modell in Python, w\u00e4hrend Nvidia Warp Simulationsprimitive, Kacheloperationen und automatische Differenzierung hinzuf\u00fcgt. CUDA und HIP bieten C++-Steuerung auf niedrigerer Ebene, und leistungstragbare Frameworks unterst\u00fctzen breitere Hardware-Ziele.<\/p>\n<p>Die gr\u00f6\u00dften Leistungssteigerungen ergeben sich selten allein aus einer sich \u00e4ndernden Syntax. Sie stammen aus der Auswahl eines parallelen Algorithmus, der Speicherung von Daten auf dem Ger\u00e4t, der Reduzierung des Speicherverkehrs, der Verwendung eines geeigneten Datenlayouts und der Eliminierung unn\u00f6tiger Zwischenoperationen.<\/p>\n<p>Ein benutzerdefinierter Kernel sollte erst entwickelt werden, nachdem die Profilerstellung einen echten Engpass identifiziert hat. Es sollte anhand einer Referenzl\u00f6sung validiert und mit Synchronisation, repr\u00e4sentativen Problemgr\u00f6\u00dfen und vollst\u00e4ndiger Abrechnung der Speicher\u00fcbertragungskosten verglichen werden.<\/p>\n<p>Wenn diese Bedingungen erf\u00fcllt sind, k\u00f6nnen Python-basierte Kernel-Tools ernsthafte Physiksimulationen unterst\u00fctzen, ohne die Forscher zu zwingen, jeden Teil der Anwendung in C ++ auf niedriger Ebene zu verschieben.<\/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\"> 11<\/span> <span class=\"rt-label rt-postfix\">minutes<\/span><\/span>Erfahren Sie, wie Sie benutzerdefinierte GPU-Kernel f\u00fcr die Physiksimulation mit NVIDIA Warp, CUDA Python und HIP \/ ROCM schreiben. Deckt Leistungsmuster, differenzierbare Physik und Produktionseins\u00e4tze ab.<\/p>\n","protected":false,"raw":"Erfahren Sie, wie Sie benutzerdefinierte GPU-Kernel f\u00fcr die Physiksimulation mit NVIDIA Warp, CUDA Python und HIP \/ ROCM schreiben. Deckt Leistungsmuster, differenzierbare Physik und Produktionseins\u00e4tze ab."},"author":2,"featured_media":0,"comment_status":"open","ping_status":"open","sticky":false,"template":"","format":"standard","meta":{"_locale":"de_DE","_original_post":"https:\/\/matforge.org\/?p=482","iawp_total_views":0,"footnotes":""},"categories":[1],"tags":[],"class_list":["post-906","post","type-post","status-publish","format-standard","hentry","category-simulation-modeling-projects","de-DE"],"yoast_head":"<!-- This site is optimized with the Yoast SEO plugin v28.1 - https:\/\/yoast.com\/product\/yoast-seo-wordpress\/ -->\n<title>Benutzerdefinierte GPU-Kernel f\u00fcr Physiksimulationen<\/title>\n<meta name=\"description\" content=\"Erfahren Sie, wie Sie GPU-Kernel f\u00fcr Physiksimulationen schreiben und optimieren, indem Sie Numba, Warp, Cupy, CUDA, HIP und Memory-aware Design verwenden.\" \/>\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\/de\/gpu-kernel-programming-custom-physics-simulation\/\" \/>\n<meta property=\"og:locale\" content=\"de_DE\" \/>\n<meta property=\"og:type\" content=\"article\" \/>\n<meta property=\"og:title\" content=\"Benutzerdefinierte GPU-Kernel f\u00fcr Physiksimulationen\" \/>\n<meta property=\"og:description\" content=\"Erfahren Sie, wie Sie GPU-Kernel f\u00fcr Physiksimulationen schreiben und optimieren, indem Sie Numba, Warp, Cupy, CUDA, HIP und Memory-aware Design verwenden.\" \/>\n<meta property=\"og:url\" content=\"https:\/\/matforge.org\/de\/gpu-kernel-programming-custom-physics-simulation\/\" \/>\n<meta property=\"og:site_name\" content=\"matforge.org\" \/>\n<meta property=\"article:published_time\" content=\"2026-07-30T12:23:15+00:00\" \/>\n<meta name=\"author\" content=\"Elena Markovska\" \/>\n<meta name=\"twitter:card\" content=\"summary_large_image\" \/>\n<meta name=\"twitter:label1\" content=\"Verfasst von\" \/>\n\t<meta name=\"twitter:data1\" content=\"Elena Markovska\" \/>\n\t<meta name=\"twitter:label2\" content=\"Gesch\u00e4tzte Lesezeit\" \/>\n\t<meta name=\"twitter:data2\" content=\"17\u00a0Minuten\" \/>\n<script type=\"application\/ld+json\" class=\"yoast-schema-graph\">{\"@context\":\"https:\\\/\\\/schema.org\",\"@graph\":[{\"@type\":\"Article\",\"@id\":\"https:\\\/\\\/matforge.org\\\/de\\\/gpu-kernel-programming-custom-physics-simulation\\\/#article\",\"isPartOf\":{\"@id\":\"https:\\\/\\\/matforge.org\\\/de\\\/gpu-kernel-programming-custom-physics-simulation\\\/\"},\"author\":{\"name\":\"Elena Markovska\",\"@id\":\"https:\\\/\\\/matforge.org\\\/#\\\/schema\\\/person\\\/980162bb5de46742daece973661d93da\"},\"headline\":\"GPU-Kernel-Programmierung f\u00fcr benutzerdefinierte Physiksimulation: Schreiben von leistungsstarken benutzerdefinierten Kerneln in Python\",\"datePublished\":\"2026-07-30T12:23:15+00:00\",\"mainEntityOfPage\":{\"@id\":\"https:\\\/\\\/matforge.org\\\/de\\\/gpu-kernel-programming-custom-physics-simulation\\\/\"},\"wordCount\":2694,\"commentCount\":0,\"articleSection\":[\"Simulation & amp; Modellierungsprojekte\"],\"inLanguage\":\"de\",\"potentialAction\":[{\"@type\":\"CommentAction\",\"name\":\"Comment\",\"target\":[\"https:\\\/\\\/matforge.org\\\/de\\\/gpu-kernel-programming-custom-physics-simulation\\\/#respond\"]}]},{\"@type\":\"WebPage\",\"@id\":\"https:\\\/\\\/matforge.org\\\/de\\\/gpu-kernel-programming-custom-physics-simulation\\\/\",\"url\":\"https:\\\/\\\/matforge.org\\\/de\\\/gpu-kernel-programming-custom-physics-simulation\\\/\",\"name\":\"Benutzerdefinierte GPU-Kernel f\u00fcr Physiksimulationen\",\"isPartOf\":{\"@id\":\"https:\\\/\\\/matforge.org\\\/#website\"},\"datePublished\":\"2026-07-30T12:23:15+00:00\",\"author\":{\"@id\":\"https:\\\/\\\/matforge.org\\\/#\\\/schema\\\/person\\\/980162bb5de46742daece973661d93da\"},\"description\":\"Erfahren Sie, wie Sie GPU-Kernel f\u00fcr Physiksimulationen schreiben und optimieren, indem Sie Numba, Warp, Cupy, CUDA, HIP und Memory-aware Design verwenden.\",\"breadcrumb\":{\"@id\":\"https:\\\/\\\/matforge.org\\\/de\\\/gpu-kernel-programming-custom-physics-simulation\\\/#breadcrumb\"},\"inLanguage\":\"de\",\"potentialAction\":[{\"@type\":\"ReadAction\",\"target\":[\"https:\\\/\\\/matforge.org\\\/de\\\/gpu-kernel-programming-custom-physics-simulation\\\/\"]}]},{\"@type\":\"BreadcrumbList\",\"@id\":\"https:\\\/\\\/matforge.org\\\/de\\\/gpu-kernel-programming-custom-physics-simulation\\\/#breadcrumb\",\"itemListElement\":[{\"@type\":\"ListItem\",\"position\":1,\"name\":\"Home\",\"item\":\"https:\\\/\\\/matforge.org\\\/de\\\/\"},{\"@type\":\"ListItem\",\"position\":2,\"name\":\"GPU-Kernel-Programmierung f\u00fcr benutzerdefinierte Physiksimulation: Schreiben von leistungsstarken benutzerdefinierten Kerneln in 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\":\"de\"},{\"@type\":\"Person\",\"@id\":\"https:\\\/\\\/matforge.org\\\/#\\\/schema\\\/person\\\/980162bb5de46742daece973661d93da\",\"name\":\"Elena Markovska\",\"image\":{\"@type\":\"ImageObject\",\"inLanguage\":\"de\",\"@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":"Benutzerdefinierte GPU-Kernel f\u00fcr Physiksimulationen","description":"Erfahren Sie, wie Sie GPU-Kernel f\u00fcr Physiksimulationen schreiben und optimieren, indem Sie Numba, Warp, Cupy, CUDA, HIP und Memory-aware Design verwenden.","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\/de\/gpu-kernel-programming-custom-physics-simulation\/","og_locale":"de_DE","og_type":"article","og_title":"Benutzerdefinierte GPU-Kernel f\u00fcr Physiksimulationen","og_description":"Erfahren Sie, wie Sie GPU-Kernel f\u00fcr Physiksimulationen schreiben und optimieren, indem Sie Numba, Warp, Cupy, CUDA, HIP und Memory-aware Design verwenden.","og_url":"https:\/\/matforge.org\/de\/gpu-kernel-programming-custom-physics-simulation\/","og_site_name":"matforge.org","article_published_time":"2026-07-30T12:23:15+00:00","author":"Elena Markovska","twitter_card":"summary_large_image","twitter_misc":{"Verfasst von":"Elena Markovska","Gesch\u00e4tzte Lesezeit":"17\u00a0Minuten"},"schema":{"@context":"https:\/\/schema.org","@graph":[{"@type":"Article","@id":"https:\/\/matforge.org\/de\/gpu-kernel-programming-custom-physics-simulation\/#article","isPartOf":{"@id":"https:\/\/matforge.org\/de\/gpu-kernel-programming-custom-physics-simulation\/"},"author":{"name":"Elena Markovska","@id":"https:\/\/matforge.org\/#\/schema\/person\/980162bb5de46742daece973661d93da"},"headline":"GPU-Kernel-Programmierung f\u00fcr benutzerdefinierte Physiksimulation: Schreiben von leistungsstarken benutzerdefinierten Kerneln in Python","datePublished":"2026-07-30T12:23:15+00:00","mainEntityOfPage":{"@id":"https:\/\/matforge.org\/de\/gpu-kernel-programming-custom-physics-simulation\/"},"wordCount":2694,"commentCount":0,"articleSection":["Simulation & amp; Modellierungsprojekte"],"inLanguage":"de","potentialAction":[{"@type":"CommentAction","name":"Comment","target":["https:\/\/matforge.org\/de\/gpu-kernel-programming-custom-physics-simulation\/#respond"]}]},{"@type":"WebPage","@id":"https:\/\/matforge.org\/de\/gpu-kernel-programming-custom-physics-simulation\/","url":"https:\/\/matforge.org\/de\/gpu-kernel-programming-custom-physics-simulation\/","name":"Benutzerdefinierte GPU-Kernel f\u00fcr Physiksimulationen","isPartOf":{"@id":"https:\/\/matforge.org\/#website"},"datePublished":"2026-07-30T12:23:15+00:00","author":{"@id":"https:\/\/matforge.org\/#\/schema\/person\/980162bb5de46742daece973661d93da"},"description":"Erfahren Sie, wie Sie GPU-Kernel f\u00fcr Physiksimulationen schreiben und optimieren, indem Sie Numba, Warp, Cupy, CUDA, HIP und Memory-aware Design verwenden.","breadcrumb":{"@id":"https:\/\/matforge.org\/de\/gpu-kernel-programming-custom-physics-simulation\/#breadcrumb"},"inLanguage":"de","potentialAction":[{"@type":"ReadAction","target":["https:\/\/matforge.org\/de\/gpu-kernel-programming-custom-physics-simulation\/"]}]},{"@type":"BreadcrumbList","@id":"https:\/\/matforge.org\/de\/gpu-kernel-programming-custom-physics-simulation\/#breadcrumb","itemListElement":[{"@type":"ListItem","position":1,"name":"Home","item":"https:\/\/matforge.org\/de\/"},{"@type":"ListItem","position":2,"name":"GPU-Kernel-Programmierung f\u00fcr benutzerdefinierte Physiksimulation: Schreiben von leistungsstarken benutzerdefinierten Kerneln in 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":"de"},{"@type":"Person","@id":"https:\/\/matforge.org\/#\/schema\/person\/980162bb5de46742daece973661d93da","name":"Elena Markovska","image":{"@type":"ImageObject","inLanguage":"de","@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\/906","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=906"}],"version-history":[{"count":1,"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/posts\/906\/revisions"}],"predecessor-version":[{"id":1002,"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/posts\/906\/revisions\/1002"}],"wp:attachment":[{"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/media?parent=906"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/categories?post=906"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/tags?post=906"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}