{"id":539,"date":"2026-07-22T08:18:02","date_gmt":"2026-07-22T08:18:02","guid":{"rendered":"https:\/\/matforge.org\/?p=539","raw":"https:\/\/matforge.org\/?p=539"},"modified":"2026-07-22T08:18:02","modified_gmt":"2026-07-22T08:18:02","slug":"gpu-kernel-programming-custom-physics-simulation","status":"publish","type":"post","link":"https:\/\/matforge.org\/es\/gpu-kernel-programming-custom-physics-simulation\/","title":{"rendered":"Programaci\u00f3n de kernel de GPU para simulaci\u00f3n de f\u00edsica personalizada: escribir kernels personalizados de alto rendimiento en Python","raw":"Programaci\u00f3n de kernel de GPU para simulaci\u00f3n de f\u00edsica personalizada: escribir kernels personalizados de alto rendimiento 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\"> 13<\/span> <span class=\"rt-label rt-postfix\">minutes<\/span><\/span><article>\n<p>La aceleraci\u00f3n de GPU en la computaci\u00f3n cient\u00edfica a menudo comienza con bibliotecas de alto nivel. Un investigador reemplaza a Numpy con Cupy, usa una FFT acelerada o llama a una rutina de \u00e1lgebra lineal habilitada para GPU. Esto puede producir una mejora sustancial sin requerir un conocimiento detallado del hardware de GPU.<\/p>\n<p>La programaci\u00f3n personalizada del kernel de GPU va un nivel m\u00e1s profundo. En lugar de llamar a una operaci\u00f3n predefinida, el desarrollador escribe la funci\u00f3n que ejecuta cada subproceso de GPU. Esto proporciona control sobre la indexaci\u00f3n de subprocesos, el dise\u00f1o de datos, las transferencias de memoria, la sincronizaci\u00f3n, la memoria compartida y el n\u00famero de operaciones realizadas durante el lanzamiento de cada kernel.<\/p>\n<p>Este control es \u00fatil cuando un modelo f\u00edsico contiene patrones de acceso que no se pueden expresar de manera eficiente a trav\u00e9s de operaciones de arreglos ordinarios. Las interacciones de part\u00edculas, las plantillas de diferencias finitas, las reglas de colisi\u00f3n, los modelos de fluidos basados en la red, las actualizaciones de materiales y las consultas de geometr\u00eda son ejemplos comunes.<\/p>\n<p>Un kernel personalizado no es autom\u00e1ticamente m\u00e1s r\u00e1pido que una biblioteca optimizada. Se vuelve valioso cuando el algoritmo expone suficiente trabajo paralelo y cuando el desarrollador puede reducir el tr\u00e1fico de memoria, combinar operaciones o adaptar el c\u00e1lculo a la arquitectura de la GPU.<\/p>\n<h2>\u00bfQu\u00e9 es un kernel de GPU?<\/h2>\n<p>Un kernel de GPU es una funci\u00f3n ejecutada por muchos subprocesos ligeros. Cada subproceso normalmente maneja una part\u00edcula, celda de cuadr\u00edcula, elemento, cara, p\u00edxel o entrada en una matriz.<\/p>\n<p>El programa host lanza el kernel con una cuadr\u00edcula de bloques de subprocesos:<\/p>\n<pre><code>kernel[blocks, threads_per_block](arguments)<\/code><\/pre>\n<p>Cada subproceso determina qu\u00e9 datos debe procesar desde su bloque e \u00edndices de subprocesos. Un \u00edndice unidimensional t\u00edpico es:<\/p>\n<pre><code>index =\n    block_index * block_size\n    + thread_index<\/code><\/pre>\n<p>Los hilos se organizan en grupos llamados bloques. Los hilos dentro de un bloque pueden sincronizar e intercambiar datos a trav\u00e9s de la memoria compartida. Por lo general, se espera que los bloques separados se ejecuten de forma independiente.<\/p>\n<h2>Cuando vale la pena escribir los kernels personalizados<\/h2>\n<p>Las bibliotecas de GPU de alto nivel generalmente deber\u00edan ser la primera opci\u00f3n. Ya proporcionan multiplicaci\u00f3n de matrices optimizada, FFT, reducciones, operaciones escasas, generaci\u00f3n de n\u00fameros aleatorios y muchas funciones de elementos.<\/p>\n<p>Un kernel personalizado se vuelve \u00fatil cuando:<\/p>\n<ul>\n<li>La simulaci\u00f3n utiliza un patr\u00f3n de acceso a la memoria no est\u00e1ndar.<\/li>\n<li>Se pueden fusionar varias operaciones peque\u00f1as en una sola pasada de memoria.<\/li>\n<li>Los subprocesos vecinos necesitan repetidamente los mismos datos locales.<\/li>\n<li>Una regla especializada de part\u00edculas, colisi\u00f3n o plantilla domina el tiempo de ejecuci\u00f3n.<\/li>\n<li>Las matrices intermedias consumen demasiada memoria de dispositivo.<\/li>\n<li>La simulaci\u00f3n necesita una operaci\u00f3n diferenciable personalizada.<\/li>\n<li>Una biblioteca existente no puede expresar la l\u00f3gica de los l\u00edmites o material de manera eficiente.<\/li>\n<\/ul>\n<p>Antes de escribir un kernel, perfile el programa. La optimizaci\u00f3n de una funci\u00f3n visualmente compleja que representa solo una peque\u00f1a parte del tiempo de ejecuci\u00f3n no producir\u00e1 una aceleraci\u00f3n general significativa.<\/p>\n<h2>Modelos de programaci\u00f3n de GPU<\/h2>\n<p>El software cient\u00edfico puede acceder a las GPU a trav\u00e9s de varios niveles de abstracci\u00f3n.<\/p>\n<table class=\"custom-table\">\n<tbody>\n<tr>\n<th>Enfoque<\/th>\n<th>ejemplos<\/th>\n<th>Fuerza principal<\/th>\n<th>compensaci\u00f3n principal<\/th>\n<\/tr>\n<tr>\n<td>Descarga basada en directivas<\/td>\n<td>OpenMP, OpenACC<\/td>\n<td>Aceleraci\u00f3n incremental de c\u00f3digo C, C++ o Fortran existente<\/td>\n<td>Menos control directo sobre los n\u00facleos generados<\/td>\n<\/tr>\n<tr>\n<td>Programaci\u00f3n de kernel orientado a proveedores<\/td>\n<td>cuda, cadera<\/td>\n<td>Control detallado sobre la ejecuci\u00f3n de GPU<\/td>\n<td>Mayor coste de implementaci\u00f3n y mantenimiento<\/td>\n<\/tr>\n<tr>\n<td>C++ port\u00e1til de rendimiento<\/td>\n<td>Sycl, Kokkos, Alpaka<\/td>\n<td>Un modelo de fuente para varios backends de hardware<\/td>\n<td>La portabilidad no garantiza el mismo rendimiento en todas partes<\/td>\n<\/tr>\n<tr>\n<td>Arreglos y kernels de GPU de Python<\/td>\n<td>Cupy, Numba-cuda, urdimbre de Nvidia<\/td>\n<td>Desarrollo r\u00e1pido con acceso a c\u00f3digo GPU compilado<\/td>\n<td>Restricciones espec\u00edficas del marco y comportamiento de compilaci\u00f3n<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n<p>La <a href=\"https:\/\/enccs.github.io\/gpu-programming\/5-intro-to-gpu-prog-models\/\" rel=\"nofollow\" target=\"_blank\">ENCCS Introducci\u00f3n a los modelos de programaci\u00f3n de GPU<\/a> proporciona una comparaci\u00f3n m\u00e1s amplia de enfoques CUDA, HIP, OpenCL, SYCL, Kokkos, OpenMP y relacionados.<\/p>\n<h2>cuda<\/h2>\n<p>CUDA proporciona acceso directo a la programaci\u00f3n de GPU de NVIDIA a trav\u00e9s de C++, Fortran y varios enlaces de idioma. Expone cuadr\u00edculas, bloques, warps, memoria compartida, flujos, eventos, bibliotecas de dispositivos y herramientas de optimizaci\u00f3n espec\u00edficas de hardware.<\/p>\n<p>A menudo se selecciona CUDA cuando el m\u00e1ximo rendimiento espec\u00edfico de NVIDIA y las herramientas maduras son m\u00e1s importantes que la portabilidad del hardware. Los principales costos son la gesti\u00f3n de memoria manual, un modelo de programaci\u00f3n de nivel inferior y la dependencia de la plataforma NVIDIA.<\/p>\n<h2>Cadera y Rocm<\/h2>\n<p>HIP proporciona un modelo de programaci\u00f3n C++ similar a CUDA dentro del ecosistema ROCM de AMD. Una gran cantidad de c\u00f3digo de estilo CUDA se puede adaptar a HIP, pero la portabilidad de la fuente no significa que un binario se ejecute sin cambios en cada GPU.<\/p>\n<p>Es posible que a\u00fan se requiera una compilaci\u00f3n, pruebas y ajustes de rendimiento separados para cada arquitectura de destino. Los detalles est\u00e1n disponibles en la <a href=\"https:\/\/rocm.docs.amd.com\/projects\/HIP\/en\/latest\/what_is_hip.html\" rel=\"nofollow\" target=\"_blank\">Documentaci\u00f3n oficial de HIP<\/a>.<\/p>\n<p>HIP es una opci\u00f3n pr\u00e1ctica cuando se debe admitir las GPU AMD o cuando un proyecto quiere reducir la dependencia de un solo proveedor de aceleraci\u00f3n mientras se conserva un estilo de programaci\u00f3n similar a CUDA.<\/p>\n<h2>Granos personalizados de Cupy<\/h2>\n<p>Cupy es mejor conocido por sus arreglos compatibles con Numpy y sus funciones num\u00e9ricas aceleradas. Tambi\u00e9n es compatible con la programaci\u00f3n de GPU personalizada.<\/p>\n<p>Los desarrolladores pueden usar:<\/p>\n<ul>\n<li><code>ElementwiseKernel<\/code> para expresiones personalizadas de elementos<\/li>\n<li><code>ReductionKernel<\/code> para reducciones personalizadas<\/li>\n<li><code>RawKernel<\/code> para n\u00facleos escritos en CUDA C++<\/li>\n<li><code>cupyx.jit.rawkernel<\/code> para las definiciones de kernel JIT de estilo Python<\/li>\n<\/ul>\n<p>Un kernel crudo simple se puede definir de la siguiente manera:<\/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 es \u00fatil cuando la mayor\u00eda de la aplicaci\u00f3n ya usa matrices de GPU y solo las operaciones seleccionadas necesitan kernels personalizados.<\/p>\n<h2>numba-cuda<\/h2>\n<p>NUMBA-CUDA compila un subconjunto restringido de Python en kernels de GPU. Su modelo de programaci\u00f3n sigue de cerca a CUDA C. Los desarrolladores definen una funci\u00f3n con un decorador de kernel, calculan un \u00edndice de subprocesos e inician la funci\u00f3n con una configuraci\u00f3n de cuadr\u00edcula y bloque.<\/p>\n<p>La documentaci\u00f3n activa del proyecto est\u00e1 disponible en <a href=\"https:\/\/nvidia.github.io\/numba-cuda\/\" rel=\"nofollow\" target=\"_blank\">numba-cuda<\/a>. Debido a que su estado de desarrollo puede cambiar, los equipos que crean software de larga duraci\u00f3n deben revisar la gu\u00eda actual de mantenimiento y migraci\u00f3n antes de comprometerse con el marco.<\/p>\n<h2>un n\u00facleo de part\u00edculas numba<\/h2>\n<p>El siguiente ejemplo educativo calcula la aceleraci\u00f3n gravitacional con las interacciones directas de todos los pares:<\/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>Este kernel tiene complejidad computacional de <code>O(N\u00b2)<\/code>. Es adecuado para explicar el mapeo de subprocesos, pero no es un algoritmo de producci\u00f3n eficiente para recuentos de part\u00edculas muy grandes.<\/p>\n<p>Los grandes sistemas gravitatorios o de part\u00edculas a menudo requieren un m\u00e9todo de \u00e1rbol, un m\u00e9todo multipolar r\u00e1pido, contenedores espaciales o listas de vecinos. Una GPU no puede eliminar el problema de escala de un algoritmo ineficiente.<\/p>\n<h2>Deformaci\u00f3n de Nvidia<\/h2>\n<p>NVIDIA WARP permite a los desarrolladores definir kernels fuertemente tipeados con sintaxis de Python. Warp compila esas funciones en CPU o c\u00f3digo CUDA y proporciona primitivas especializadas para geometr\u00eda, simulaci\u00f3n, operaciones escasas, elementos finitos, optimizaci\u00f3n y diferenciaci\u00f3n autom\u00e1tica.<\/p>\n<p>La <a href=\"https:\/\/developer.nvidia.com\/warp-python\" rel=\"nofollow\" target=\"_blank\">nvidia warp product page<\/a> y <a href=\"https:\/\/github.com\/nvidia\/warp\" rel=\"nofollow\" target=\"_blank\">Warp GitHub Repository<\/a> Incluye ejemplos de part\u00edculas, fluidos, mallas, optimizaci\u00f3n, simulaci\u00f3n diferenciable y Programaci\u00f3n de GPU basada en azulejos.<\/p>\n<h2>El mismo grano de part\u00edculas en urdimbre<\/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 determina internamente una configuraci\u00f3n de lanzamiento adecuada para un lanzamiento de kernel ordinario. Los usuarios avanzados a\u00fan pueden trabajar con dimensiones de bloques, gr\u00e1ficos de comandos, ejecuci\u00f3n en mosaico y operaciones de dispositivos especializados cuando sea necesario.<\/p>\n<h2>Bucles de zancada<\/h2>\n<p>Un kernel no necesita un hilo asignado permanentemente para cada elemento. Un bucle de tracci\u00f3n de cuadr\u00edcula permite que cada subproceso procese varias entradas:<\/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>Este patr\u00f3n separa el n\u00famero de subprocesos lanzados del tama\u00f1o total de los datos. Es \u00fatil cuando se procesan arreglos muy grandes o se reutiliza una configuraci\u00f3n de lanzamiento fijo.<\/p>\n<p>Una introducci\u00f3n pr\u00e1ctica a este patr\u00f3n est\u00e1 disponible en <a href=\"https:\/\/thedatafrog.com\/en\/articles\/cuda-kernel-python\/\" rel=\"nofollow\" target=\"_blank\">la NUMBA y CUDA de Data Frog Tutorial<\/a>.<\/p>\n<h2>Coalescente de memoria<\/h2>\n<p>Los kernels de GPU a menudo est\u00e1n limitados por el ancho de banda de la memoria en lugar del rendimiento aritm\u00e9tico. Los subprocesos dentro de un warp deber\u00edan acceder idealmente a las direcciones de memoria cercanas para que el hardware pueda combinar sus solicitudes.<\/p>\n<p>Considere los datos de part\u00edculas almacenados como:<\/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>Este dise\u00f1o de matriz de estructuras puede ser conveniente para el c\u00f3digo orientado a objetos. Un dise\u00f1o de estructura de arreglos almacena matrices contiguas separadas:<\/p>\n<pre><code>x_positions[]\ny_positions[]\nz_positions[]\nmasses[]<\/code><\/pre>\n<p>El segundo dise\u00f1o puede proporcionar un mejor coalescencia cuando cada subproceso lee el mismo campo para una part\u00edcula diferente. El mejor dise\u00f1o a\u00fan depende de qu\u00e9 campos se accedan juntos.<\/p>\n<p>Los tipos de vectores no corrigen autom\u00e1ticamente un patr\u00f3n de acceso deficiente. Los desarrolladores deben inspeccionar las direcciones reales solicitadas por los hilos vecinos.<\/p>\n<h2>memoria compartida<\/h2>\n<p>La memoria compartida es una peque\u00f1a \u00e1rea de memoria de baja latencia accesible por subprocesos en el mismo bloque. Puede reducir las lecturas repetidas de la memoria global.<\/p>\n<p>Una plantilla unidimensional puede cargar un bloque de valores y su halo en la memoria compartida:<\/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>La llamada a <code>cuda.syncthreads()<\/code> garantiza que cada subproceso termine de cargar sus valores antes de que comience el c\u00e1lculo de la plantilla.<\/p>\n<p>La memoria compartida no debe usarse autom\u00e1ticamente. La asignaci\u00f3n excesiva de memoria compartida puede reducir la ocupaci\u00f3n, agregar el costo de sincronizaci\u00f3n y hacer que el kernel sea m\u00e1s lento. Se requiere perfilado.<\/p>\n<h2>Fusi\u00f3n de grano<\/h2>\n<p>Las operaciones de matriz separadas a menudo crean varios lanzamientos de kernel y matrices intermedios:<\/p>\n<pre><code>velocity += dt * acceleration\nposition += dt * velocity\nenergy = compute_energy(position, velocity)<\/code><\/pre>\n<p>Un kernel fusionado puede calcular las tres actualizaciones mientras los valores necesarios permanecen en los registros. Esto reduce la sobrecarga de lanzamiento y el tr\u00e1fico de memoria global.<\/p>\n<p>La fusi\u00f3n es m\u00e1s beneficiosa cuando las operaciones son simples y est\u00e1n vinculadas a la memoria. Fusionar demasiado trabajo puede aumentar el uso de registros, reducir la ocupaci\u00f3n y hacer que el kernel sea dif\u00edcil de mantener.<\/p>\n<h2>Elegir el tama\u00f1o del bloque<\/h2>\n<p>Los valores como 128 o 256 hilos por bloque son puntos de partida razonables, no como Optima universal.<\/p>\n<p>El mejor tama\u00f1o de bloque depende de:<\/p>\n<ul>\n<li>Registros utilizados por hilo<\/li>\n<li>Memoria compartida utilizada por bloque<\/li>\n<li>Divergencia de ramas<\/li>\n<li>Mezcla de instrucciones<\/li>\n<li>Comportamiento de acceso a la memoria<\/li>\n<li>La arquitectura de GPU objetivo<\/li>\n<\/ul>\n<p>Una regla como lanzar al menos el doble de bloques que los multiprocesadores de transmisi\u00f3n puede ser un experimento inicial \u00fatil, pero no garantiza el m\u00e1ximo rendimiento. Las calculadoras de ocupaci\u00f3n y las herramientas de perfilado deben guiar la configuraci\u00f3n final.<\/p>\n<h2>F\u00edsica diferenciable<\/h2>\n<p>La simulaci\u00f3n diferenciable calcula c\u00f3mo cambia una salida con respecto a las entradas como las propiedades del material, las fuerzas, la geometr\u00eda o las condiciones iniciales.<\/p>\n<p>Warp puede registrar las operaciones del kernel compatible y ejecutar la diferenciaci\u00f3n autom\u00e1tica en modo inverso. Esto se puede utilizar para problemas inversos, optimizaci\u00f3n del dise\u00f1o, estimaci\u00f3n de par\u00e1metros e integraci\u00f3n con flujos de trabajo de aprendizaje autom\u00e1tico.<\/p>\n<p>Un patr\u00f3n simplificado utiliza una cinta:<\/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>No se garantiza la diferenciaci\u00f3n autom\u00e1tica para cada kernel. Las sobreescrituras en el lugar, los at\u00f3micos no deterministas, el c\u00f3digo nativo externo, la l\u00f3gica discontinua y las operaciones no admitidas pueden requerir reformulaci\u00f3n o gradientes personalizados.<\/p>\n<p>Esta capacidad se conecta directamente con aplicaciones como la optimizaci\u00f3n adjunta y <a href=\"https:\/\/matforge.org\/physics-informed-neural-networks-pinns-for-scientific-simulations\/\" rel=\"nofollow\" target=\"_blank\">Redes neuronales informadas por la f\u00edsica<\/a>.<\/p>\n<h2>Cuando una GPU puede no ser m\u00e1s r\u00e1pida<\/h2>\n<p>No existe un umbral de tama\u00f1o de problema universal en el que una GPU se vuelva m\u00e1s r\u00e1pida que una CPU. El cruce depende del hardware, la precisi\u00f3n, el movimiento de datos, la estructura del algoritmo, la calidad del compilador y la frecuencia con la que se reutilizan los mismos datos residentes en el dispositivo.<\/p>\n<table class=\"custom-table\">\n<tbody>\n<tr>\n<th>carga de trabajo<\/th>\n<th>Adecuaci\u00f3n de GPU<\/th>\n<th>Consideraci\u00f3n principal<\/th>\n<\/tr>\n<tr>\n<td>Actualizaci\u00f3n de part\u00edculas grandes<\/td>\n<td>A menudo fuerte<\/td>\n<td>muchas operaciones independientes similares<\/td>\n<\/tr>\n<tr>\n<td>Plantilla grande regular<\/td>\n<td>A menudo fuerte<\/td>\n<td>Acceso predecible a la memoria en paralelo<\/td>\n<\/tr>\n<tr>\n<td>\u00c1lgebra lineal densa<\/td>\n<td>Fuerte al usar bibliotecas optimizadas<\/td>\n<td>Alta intensidad aritm\u00e9tica<\/td>\n<\/tr>\n<tr>\n<td>peque\u00f1a simulaci\u00f3n<\/td>\n<td>dependiente<\/td>\n<td>Lanzamiento y transferencia Los gastos generales pueden dominar<\/td>\n<\/tr>\n<tr>\n<td>Trasversal de gr\u00e1ficos irregulares<\/td>\n<td>Mezclado<\/td>\n<td>Divergencia y acceso a la memoria impredecible<\/td>\n<\/tr>\n<tr>\n<td>Flujo de trabajo enlazado a E\/S<\/td>\n<td>por lo general limitado<\/td>\n<td>La GPU no puede eliminar un cuello de botella de almacenamiento o de red<\/td>\n<\/tr>\n<tr>\n<td>Algoritmo fuertemente secuencial<\/td>\n<td>por lo general d\u00e9bil<\/td>\n<td>trabajo independiente insuficiente<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n<p>El flujo de trabajo correcto es medir la versi\u00f3n de la CPU, la versi\u00f3n de la biblioteca GPU y la versi\u00f3n del kernel personalizado con datos representativos. La <a href=\"https:\/\/matforge.org\/scientific-computing-performance-profiling-identifying-bottlenecks\/\" rel=\"nofollow\" target=\"_blank\">Gu\u00eda de perfiles de rendimiento de computaci\u00f3n cient\u00edfica<\/a> explica c\u00f3mo identificar el cuello de botella real antes de optimizarlo.<\/p>\n<h2>Evitar la transferencia de gastos generales<\/h2>\n<p>Las transferencias repetidas entre la memoria del dispositivo y del dispositivo pueden eliminar el beneficio del c\u00e1lculo de la GPU.<\/p>\n<p>Un bucle ineficiente puede seguir este patr\u00f3n:<\/p>\n<ol>\n<li>Copiar datos a la GPU.<\/li>\n<li>Ejecute un kernel peque\u00f1o.<\/li>\n<li>Copie el resultado a la CPU.<\/li>\n<li>modificarlo en la CPU.<\/li>\n<li>c\u00f3pielo de nuevo a la GPU.<\/li>\n<\/ol>\n<p>Un mejor dise\u00f1o mantiene el estado de simulaci\u00f3n en el dispositivo durante muchos pasos y transfiere solo la salida requerida para la visualizaci\u00f3n, los puntos de control o el an\u00e1lisis.<\/p>\n<p>Los flujos asincr\u00f3nicos, la memoria del host anclada, la comunicaci\u00f3n superpuesta y las funciones de memoria unificada pueden ayudar, pero solo deben introducirse despu\u00e9s de que se hayan medido los costos de transferencia ordinarios.<\/p>\n<h2>El estudio de caso XLB<\/h2>\n<p>El proyecto XLB de Autodesk Research proporciona un ejemplo \u00fatil de simulaci\u00f3n de GPU nativa de Python. XLB es una biblioteca Boltzmann de c\u00f3digo abierto con varios backends computacionales, incluido Nvidia Warp.<\/p>\n<p>En las configuraciones de referencia informadas por Autodesk Research y NVIDIA, el backend warp alcanz\u00f3 un rendimiento cercano a la implementaci\u00f3n comparada de C++\/OpenCl FluidX3D para un caso espec\u00edfico de cavidad controlada por TAPA. Una comparaci\u00f3n separada inform\u00f3 una aceleraci\u00f3n de ocho veces aproximada sobre el backend JAX de XLB en el hardware y las configuraciones seleccionados.<\/p>\n<p>El equipo tambi\u00e9n demostr\u00f3 un enfoque fuera de n\u00facleo en un grupo GH200 de ocho nodos con un dominio de aproximadamente 50 mil millones de celdas de red. Estos resultados se aplican a las opciones de soluci\u00f3n, problema, hardware y la implementaci\u00f3n reportada. No deben tratarse como garant\u00edas generales de aceleraci\u00f3n para los n\u00facleos de Python.<\/p>\n<p>El estudio de caso completo est\u00e1 disponible en <a href=\"https:\/\/developer.nvidia.com\/blog\/autodesk-research-brings-warp-speed-to-computational-fluid-dynamics-on-nvidia-gh200\/\" rel=\"nofollow\" target=\"_blank\">Autodesk Research lleva la velocidad de la deformaci\u00f3n a la din\u00e1mica de fluidos computacional en NVIDIA GH200<\/a>. El c\u00f3digo de proyecto actual est\u00e1 disponible en el repositorio <a href=\"https:\/\/github.com\/Autodesk\/XLB\" rel=\"nofollow\" target=\"_blank\">Autodesk XLB<\/a>.<\/p>\n<h2>C\u00f3mo comparar correctamente un kernel<\/h2>\n<p>Las operaciones de GPU suelen ser as\u00edncronas. Medir solo la duraci\u00f3n de la llamada a la funci\u00f3n de Python puede informar el tiempo necesario para poner en cola el kernel en lugar del tiempo necesario para ejecutarlo.<\/p>\n<p>Un punto de referencia b\u00e1sico debe:<\/p>\n<ol>\n<li>Ejecute el kernel varias veces para activar la compilaci\u00f3n y el calentamiento.<\/li>\n<li>Sincronice antes de iniciar el temporizador.<\/li>\n<li>Ejecute varias iteraciones medidas.<\/li>\n<li>Sincronice antes de detener el temporizador.<\/li>\n<li>Reporte el promedio y la variaci\u00f3n entre las repeticiones.<\/li>\n<li>Separe el tiempo de transferencia de datos desde el tiempo de ejecuci\u00f3n del kernel.<\/li>\n<li>Verifique que las versiones de CPU y GPU produzcan resultados equivalentes.<\/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 comparaci\u00f3n debe usar un tama\u00f1o de simulaci\u00f3n realista e incluir el flujo de trabajo completo cuando los costos de transferencia o preprocesamiento importan.<\/p>\n<p>La metodolog\u00eda de precisi\u00f3n de trabajo tambi\u00e9n puede comparar el tiempo de ejecuci\u00f3n con el error num\u00e9rico. La <a href=\"https:\/\/docs.sciml.ai\/DiffEqDevDocs\/stable\/alg_dev\/benchmarks\/\" rel=\"nofollow\" target=\"_blank\">Documentaci\u00f3n de referencia de SCML<\/a> proporciona ejemplos de este estilo de evaluaci\u00f3n.<\/p>\n<h2>un flujo de trabajo de desarrollo pr\u00e1ctico<\/h2>\n<ol>\n<li>Implemente y verifique una versi\u00f3n de referencia de CPU clara.<\/li>\n<li>Perfil de la aplicaci\u00f3n para localizar la operaci\u00f3n dominante.<\/li>\n<li>Pruebe una biblioteca de GPU optimizada antes de escribir un kernel.<\/li>\n<li>Mantenga los datos reutilizados con frecuencia en el dispositivo.<\/li>\n<li>Escriba el kernel correcto m\u00e1s simple.<\/li>\n<li>Validarlo contra el resultado de la CPU.<\/li>\n<li>Mida las transferencias de memoria y la ejecuci\u00f3n por separado.<\/li>\n<li>Inspeccione el uso de coalescencia, ocupaci\u00f3n, divergencia y registro.<\/li>\n<li>Pruebe la memoria compartida o la fusi\u00f3n solo cuando la creaci\u00f3n de perfiles los admita.<\/li>\n<li>Compare varios tama\u00f1os de problemas y arquitecturas de GPU.<\/li>\n<\/ol>\n<h2>Elegir un marco<\/h2>\n<table class=\"custom-table\">\n<tbody>\n<tr>\n<th>Requisito<\/th>\n<th>Posible punto de partida<\/th>\n<\/tr>\n<tr>\n<td>Flujo de trabajo de GPU de estilo numpy existente<\/td>\n<td>CUPY con operaciones integradas o kernels personalizados<\/td>\n<\/tr>\n<tr>\n<td>Peque\u00f1a cantidad de kernels de Python de estilo CUDA<\/td>\n<td>Numba-CUDA, despu\u00e9s de revisar su estado de soporte actual<\/td>\n<\/tr>\n<tr>\n<td>Simulaci\u00f3n, geometr\u00eda y kernels diferenciables<\/td>\n<td>Deformaci\u00f3n de Nvidia<\/td>\n<\/tr>\n<tr>\n<td>Control m\u00e1ximo espec\u00edfico de NVIDIA<\/td>\n<td>cud\u00e1 C++<\/td>\n<\/tr>\n<tr>\n<td>Objetivo de GPU AMD con fuente CUDA<\/td>\n<td>Cadera y Rocm<\/td>\n<\/tr>\n<tr>\n<td>C++ port\u00e1til en varios backends<\/td>\n<td>SYCL, Kokkos u otra capa de rendimiento-portabilidad<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n<h2>Errores comunes de programaci\u00f3n del kernel<\/h2>\n<ul>\n<li>Mover datos entre la CPU y la GPU dentro de cada paso del tiempo<\/li>\n<li>Ignorar hilos fuera de los l\u00edmites<\/li>\n<li>Usando un algoritmo ineficiente y esperando hardware para arreglar su escala<\/li>\n<li>Acceder a la memoria con un dise\u00f1o no coalescido<\/li>\n<li>Agregar memoria compartida sin medir si ayuda<\/li>\n<li>lanzando muchos granos peque\u00f1os en lugar de considerar la fusi\u00f3n<\/li>\n<li>Uso de registros excesivos o memoria compartida por bloque<\/li>\n<li>Benchmarking c\u00f3digo as\u00edncrono sin sincronizaci\u00f3n<\/li>\n<li>Comparaci\u00f3n de salidas sin comprobar la precisi\u00f3n num\u00e9rica<\/li>\n<li>Uso de reclamos de rendimiento fijo en diferentes GPU y cargas de trabajo<\/li>\n<li>Suponiendo que la fuente port\u00e1til proporcione un rendimiento port\u00e1til<\/li>\n<li>Aplicaci\u00f3n de diferenciaci\u00f3n autom\u00e1tica a operaciones in situ no soportadas<\/li>\n<\/ul>\n<h2>Gu\u00edas relacionadas<\/h2>\n<ul>\n<li><a href=\"https:\/\/matforge.org\/gpu-accelerated-scientific-computing-cupy-numba-cudf-compared\/\" rel=\"nofollow\" target=\"_blank\">acelerado de GPU Computaci\u00f3n cient\u00edfica: comparado con CUPY, NUMBA y CUDF<\/a>: compare las herramientas de Python de alto nivel para las cargas de trabajo de GPU.<\/li>\n<li><a href=\"https:\/\/matforge.org\/gpu-acceleration-for-fipy-simulations-cupy-and-numba-integration-guide\/\" rel=\"nofollow\" target=\"_blank\">aceleraci\u00f3n de la GPU Para simulaciones FIPY: Gu\u00eda de integraci\u00f3n de CUPY y Numba<\/a>: explore los posibles componentes de la GPU en torno a un flujo de trabajo de Fipy.<\/li>\n<li><a href=\"https:\/\/matforge.org\/physics-informed-neural-networks-pinns-for-scientific-simulations\/\" rel=\"nofollow\" target=\"_blank\">informado de f\u00edsica Redes neuronales para simulaciones cient\u00edficas<\/a> \u2014 Revisar la relaci\u00f3n entre los modelos f\u00edsicos y el aprendizaje diferenciable.<\/li>\n<li><a href=\"https:\/\/matforge.org\/when-to-use-fem-fvm-fdm\/\" rel=\"nofollow\" target=\"_blank\">cu\u00e1ndo utilizar FEM, FVM y FDM<\/a> \u2014 Seleccione un espacio Discretizaci\u00f3n antes de optimizar su implementaci\u00f3n.<\/li>\n<li><a href=\"https:\/\/matforge.org\/scientific-computing-performance-profiling-identifying-bottlenecks\/\" rel=\"nofollow\" target=\"_blank\">Perfil de rendimiento: Identificaci\u00f3n de cuellos de botella en el c\u00f3digo cient\u00edfico<\/a>: busque las limitaciones de CPU, memoria, transferencia y E\/S.<\/li>\n<li><a href=\"https:\/\/matforge.org\/hpc-python-workflows-from-laptop-to-supercomputer\/\" rel=\"nofollow\" target=\"_blank\">Flujos de trabajo de HPC Python: desde la computadora port\u00e1til a la supercomputadora<\/a>: planifique la progresi\u00f3n del desarrollo local a los grandes sistemas.<\/li>\n<\/ul>\n<h2>Conclusi\u00f3n<\/h2>\n<p>Los kernels de GPU personalizados permiten a los desarrolladores cient\u00edficos traducir part\u00edculas, celdas de cuadr\u00edcula, reglas de materiales y otras operaciones f\u00edsicas directamente en funciones masivamente paralelas. Proporcionan control sobre la indexaci\u00f3n, el acceso a la memoria, la sincronizaci\u00f3n, la fusi\u00f3n del kernel y la optimizaci\u00f3n espec\u00edfica del dispositivo.<\/p>\n<p>CUPY ofrece operaciones de matriz de alto nivel e interfaces de kernel personalizadas. Numba-Cuda proporciona un modelo similar a CUDA en Python, mientras que Nvidia Warp agrega primitivas de simulaci\u00f3n, operaciones de mosaico y diferenciaci\u00f3n autom\u00e1tica. CUDA y HIP proporcionan control C++ de nivel inferior, y los marcos port\u00e1tiles de rendimiento admiten objetivos de hardware m\u00e1s amplios.<\/p>\n<p>Las mayores ganancias de rendimiento rara vez provienen de cambiar la sintaxis por s\u00ed sola. Provienen de elegir un algoritmo paralelo, mantener datos en el dispositivo, reducir el tr\u00e1fico de memoria, utilizar un dise\u00f1o de datos adecuado y eliminar operaciones intermedias innecesarias.<\/p>\n<p>Un kernel personalizado debe desarrollarse solo despu\u00e9s de que la creaci\u00f3n de perfiles identifique un cuello de botella real. Debe validarse contra una soluci\u00f3n de referencia y compararse con la sincronizaci\u00f3n, los tama\u00f1os de problemas representativos y la contabilidad completa de los costos de transferencia de memoria.<\/p>\n<p>Cuando se cumplen esas condiciones, las herramientas de kernel basadas en Python pueden respaldar una simulaci\u00f3n f\u00edsica seria sin obligar a los investigadores a mover cada parte de la aplicaci\u00f3n a C++ de bajo nivel.<\/p>\n<\/article>\n","protected":false,"raw":"<article>\n<p>La aceleraci\u00f3n de GPU en la computaci\u00f3n cient\u00edfica a menudo comienza con bibliotecas de alto nivel. Un investigador reemplaza a Numpy con Cupy, usa una FFT acelerada o llama a una rutina de \u00e1lgebra lineal habilitada para GPU. Esto puede producir una mejora sustancial sin requerir un conocimiento detallado del hardware de GPU.<\/p>\n<p>La programaci\u00f3n personalizada del kernel de GPU va un nivel m\u00e1s profundo. En lugar de llamar a una operaci\u00f3n predefinida, el desarrollador escribe la funci\u00f3n que ejecuta cada subproceso de GPU. Esto proporciona control sobre la indexaci\u00f3n de subprocesos, el dise\u00f1o de datos, las transferencias de memoria, la sincronizaci\u00f3n, la memoria compartida y el n\u00famero de operaciones realizadas durante el lanzamiento de cada kernel.<\/p>\n<p>Este control es \u00fatil cuando un modelo f\u00edsico contiene patrones de acceso que no se pueden expresar de manera eficiente a trav\u00e9s de operaciones de arreglos ordinarios. Las interacciones de part\u00edculas, las plantillas de diferencias finitas, las reglas de colisi\u00f3n, los modelos de fluidos basados en la red, las actualizaciones de materiales y las consultas de geometr\u00eda son ejemplos comunes.<\/p>\n<p>Un kernel personalizado no es autom\u00e1ticamente m\u00e1s r\u00e1pido que una biblioteca optimizada. Se vuelve valioso cuando el algoritmo expone suficiente trabajo paralelo y cuando el desarrollador puede reducir el tr\u00e1fico de memoria, combinar operaciones o adaptar el c\u00e1lculo a la arquitectura de la GPU.<\/p>\n<h2>\u00bfQu\u00e9 es un kernel de GPU?<\/h2>\n<p>Un kernel de GPU es una funci\u00f3n ejecutada por muchos subprocesos ligeros. Cada subproceso normalmente maneja una part\u00edcula, celda de cuadr\u00edcula, elemento, cara, p\u00edxel o entrada en una matriz.<\/p>\n<p>El programa host lanza el kernel con una cuadr\u00edcula de bloques de subprocesos:<\/p>\n<pre><code>kernel[blocks, threads_per_block](arguments)<\/code><\/pre>\n<p>Cada subproceso determina qu\u00e9 datos debe procesar desde su bloque e \u00edndices de subprocesos. Un \u00edndice unidimensional t\u00edpico es:<\/p>\n<pre><code>index =\n    block_index * block_size\n    + thread_index<\/code><\/pre>\n<p>Los hilos se organizan en grupos llamados bloques. Los hilos dentro de un bloque pueden sincronizar e intercambiar datos a trav\u00e9s de la memoria compartida. Por lo general, se espera que los bloques separados se ejecuten de forma independiente.<\/p>\n<h2>Cuando vale la pena escribir los kernels personalizados<\/h2>\n<p>Las bibliotecas de GPU de alto nivel generalmente deber\u00edan ser la primera opci\u00f3n. Ya proporcionan multiplicaci\u00f3n de matrices optimizada, FFT, reducciones, operaciones escasas, generaci\u00f3n de n\u00fameros aleatorios y muchas funciones de elementos.<\/p>\n<p>Un kernel personalizado se vuelve \u00fatil cuando:<\/p>\n<ul>\n<li>La simulaci\u00f3n utiliza un patr\u00f3n de acceso a la memoria no est\u00e1ndar.<\/li>\n<li>Se pueden fusionar varias operaciones peque\u00f1as en una sola pasada de memoria.<\/li>\n<li>Los subprocesos vecinos necesitan repetidamente los mismos datos locales.<\/li>\n<li>Una regla especializada de part\u00edculas, colisi\u00f3n o plantilla domina el tiempo de ejecuci\u00f3n.<\/li>\n<li>Las matrices intermedias consumen demasiada memoria de dispositivo.<\/li>\n<li>La simulaci\u00f3n necesita una operaci\u00f3n diferenciable personalizada.<\/li>\n<li>Una biblioteca existente no puede expresar la l\u00f3gica de los l\u00edmites o material de manera eficiente.<\/li>\n<\/ul>\n<p>Antes de escribir un kernel, perfile el programa. La optimizaci\u00f3n de una funci\u00f3n visualmente compleja que representa solo una peque\u00f1a parte del tiempo de ejecuci\u00f3n no producir\u00e1 una aceleraci\u00f3n general significativa.<\/p>\n<h2>Modelos de programaci\u00f3n de GPU<\/h2>\n<p>El software cient\u00edfico puede acceder a las GPU a trav\u00e9s de varios niveles de abstracci\u00f3n.<\/p>\n<table class=\"custom-table\">\n<tbody><tr>\n<th>Enfoque<\/th>\n<th>ejemplos<\/th>\n<th>Fuerza principal<\/th>\n<th>compensaci\u00f3n principal<\/th>\n<\/tr>\n<tr>\n<td>Descarga basada en directivas<\/td>\n<td>OpenMP, OpenACC<\/td>\n<td>Aceleraci\u00f3n incremental de c\u00f3digo C, C++ o Fortran existente<\/td>\n<td>Menos control directo sobre los n\u00facleos generados<\/td>\n<\/tr>\n<tr>\n<td>Programaci\u00f3n de kernel orientado a proveedores<\/td>\n<td>cuda, cadera<\/td>\n<td>Control detallado sobre la ejecuci\u00f3n de GPU<\/td>\n<td>Mayor coste de implementaci\u00f3n y mantenimiento<\/td>\n<\/tr>\n<tr>\n<td>C++ port\u00e1til de rendimiento<\/td>\n<td>Sycl, Kokkos, Alpaka<\/td>\n<td>Un modelo de fuente para varios backends de hardware<\/td>\n<td>La portabilidad no garantiza el mismo rendimiento en todas partes<\/td>\n<\/tr>\n<tr>\n<td>Arreglos y kernels de GPU de Python<\/td>\n<td>Cupy, Numba-cuda, urdimbre de Nvidia<\/td>\n<td>Desarrollo r\u00e1pido con acceso a c\u00f3digo GPU compilado<\/td>\n<td>Restricciones espec\u00edficas del marco y comportamiento de compilaci\u00f3n<\/td>\n<\/tr>\n<\/tbody><\/table>\n<p>La <a href=\"https:\/\/enccs.github.io\/gpu-programming\/5-intro-to-gpu-prog-models\/\" rel=\"nofollow\" target=\"_blank\">ENCCS Introducci\u00f3n a los modelos de programaci\u00f3n de GPU<\/a> proporciona una comparaci\u00f3n m\u00e1s amplia de enfoques CUDA, HIP, OpenCL, SYCL, Kokkos, OpenMP y relacionados.<\/p>\n<h2>cuda<\/h2>\n<p>CUDA proporciona acceso directo a la programaci\u00f3n de GPU de NVIDIA a trav\u00e9s de C++, Fortran y varios enlaces de idioma. Expone cuadr\u00edculas, bloques, warps, memoria compartida, flujos, eventos, bibliotecas de dispositivos y herramientas de optimizaci\u00f3n espec\u00edficas de hardware.<\/p>\n<p>A menudo se selecciona CUDA cuando el m\u00e1ximo rendimiento espec\u00edfico de NVIDIA y las herramientas maduras son m\u00e1s importantes que la portabilidad del hardware. Los principales costos son la gesti\u00f3n de memoria manual, un modelo de programaci\u00f3n de nivel inferior y la dependencia de la plataforma NVIDIA.<\/p>\n<h2>Cadera y Rocm<\/h2>\n<p>HIP proporciona un modelo de programaci\u00f3n C++ similar a CUDA dentro del ecosistema ROCM de AMD. Una gran cantidad de c\u00f3digo de estilo CUDA se puede adaptar a HIP, pero la portabilidad de la fuente no significa que un binario se ejecute sin cambios en cada GPU.<\/p>\n<p>Es posible que a\u00fan se requiera una compilaci\u00f3n, pruebas y ajustes de rendimiento separados para cada arquitectura de destino. Los detalles est\u00e1n disponibles en la <a href=\"https:\/\/rocm.docs.amd.com\/projects\/HIP\/en\/latest\/what_is_hip.html\" rel=\"nofollow\" target=\"_blank\">Documentaci\u00f3n oficial de HIP<\/a>.<\/p>\n<p>HIP es una opci\u00f3n pr\u00e1ctica cuando se debe admitir las GPU AMD o cuando un proyecto quiere reducir la dependencia de un solo proveedor de aceleraci\u00f3n mientras se conserva un estilo de programaci\u00f3n similar a CUDA.<\/p>\n<h2>Granos personalizados de Cupy<\/h2>\n<p>Cupy es mejor conocido por sus arreglos compatibles con Numpy y sus funciones num\u00e9ricas aceleradas. Tambi\u00e9n es compatible con la programaci\u00f3n de GPU personalizada.<\/p>\n<p>Los desarrolladores pueden usar:<\/p>\n<ul>\n<li><code>ElementwiseKernel<\/code> para expresiones personalizadas de elementos<\/li>\n<li><code>ReductionKernel<\/code> para reducciones personalizadas<\/li>\n<li><code>RawKernel<\/code> para n\u00facleos escritos en CUDA C++<\/li>\n<li><code>cupyx.jit.rawkernel<\/code> para las definiciones de kernel JIT de estilo Python<\/li>\n<\/ul>\n<p>Un kernel crudo simple se puede definir de la siguiente manera:<\/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 es \u00fatil cuando la mayor\u00eda de la aplicaci\u00f3n ya usa matrices de GPU y solo las operaciones seleccionadas necesitan kernels personalizados.<\/p>\n<h2>numba-cuda<\/h2>\n<p>NUMBA-CUDA compila un subconjunto restringido de Python en kernels de GPU. Su modelo de programaci\u00f3n sigue de cerca a CUDA C. Los desarrolladores definen una funci\u00f3n con un decorador de kernel, calculan un \u00edndice de subprocesos e inician la funci\u00f3n con una configuraci\u00f3n de cuadr\u00edcula y bloque.<\/p>\n<p>La documentaci\u00f3n activa del proyecto est\u00e1 disponible en <a href=\"https:\/\/nvidia.github.io\/numba-cuda\/\" rel=\"nofollow\" target=\"_blank\">numba-cuda<\/a>. Debido a que su estado de desarrollo puede cambiar, los equipos que crean software de larga duraci\u00f3n deben revisar la gu\u00eda actual de mantenimiento y migraci\u00f3n antes de comprometerse con el marco.<\/p>\n<h2>un n\u00facleo de part\u00edculas numba<\/h2>\n<p>El siguiente ejemplo educativo calcula la aceleraci\u00f3n gravitacional con las interacciones directas de todos los pares:<\/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>Este kernel tiene complejidad computacional de <code>O(N\u00b2)<\/code>. Es adecuado para explicar el mapeo de subprocesos, pero no es un algoritmo de producci\u00f3n eficiente para recuentos de part\u00edculas muy grandes.<\/p>\n<p>Los grandes sistemas gravitatorios o de part\u00edculas a menudo requieren un m\u00e9todo de \u00e1rbol, un m\u00e9todo multipolar r\u00e1pido, contenedores espaciales o listas de vecinos. Una GPU no puede eliminar el problema de escala de un algoritmo ineficiente.<\/p>\n<h2>Deformaci\u00f3n de Nvidia<\/h2>\n<p>NVIDIA WARP permite a los desarrolladores definir kernels fuertemente tipeados con sintaxis de Python. Warp compila esas funciones en CPU o c\u00f3digo CUDA y proporciona primitivas especializadas para geometr\u00eda, simulaci\u00f3n, operaciones escasas, elementos finitos, optimizaci\u00f3n y diferenciaci\u00f3n autom\u00e1tica.<\/p>\n<p>La <a href=\"https:\/\/developer.nvidia.com\/warp-python\" rel=\"nofollow\" target=\"_blank\">nvidia warp product page<\/a> y <a href=\"https:\/\/github.com\/nvidia\/warp\" rel=\"nofollow\" target=\"_blank\">Warp GitHub Repository<\/a> Incluye ejemplos de part\u00edculas, fluidos, mallas, optimizaci\u00f3n, simulaci\u00f3n diferenciable y Programaci\u00f3n de GPU basada en azulejos.<\/p>\n<h2>El mismo grano de part\u00edculas en urdimbre<\/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 determina internamente una configuraci\u00f3n de lanzamiento adecuada para un lanzamiento de kernel ordinario. Los usuarios avanzados a\u00fan pueden trabajar con dimensiones de bloques, gr\u00e1ficos de comandos, ejecuci\u00f3n en mosaico y operaciones de dispositivos especializados cuando sea necesario.<\/p>\n<h2>Bucles de zancada<\/h2>\n<p>Un kernel no necesita un hilo asignado permanentemente para cada elemento. Un bucle de tracci\u00f3n de cuadr\u00edcula permite que cada subproceso procese varias entradas:<\/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>Este patr\u00f3n separa el n\u00famero de subprocesos lanzados del tama\u00f1o total de los datos. Es \u00fatil cuando se procesan arreglos muy grandes o se reutiliza una configuraci\u00f3n de lanzamiento fijo.<\/p>\n<p>Una introducci\u00f3n pr\u00e1ctica a este patr\u00f3n est\u00e1 disponible en <a href=\"https:\/\/thedatafrog.com\/en\/articles\/cuda-kernel-python\/\" rel=\"nofollow\" target=\"_blank\">la NUMBA y CUDA de Data Frog Tutorial<\/a>.<\/p>\n<h2>Coalescente de memoria<\/h2>\n<p>Los kernels de GPU a menudo est\u00e1n limitados por el ancho de banda de la memoria en lugar del rendimiento aritm\u00e9tico. Los subprocesos dentro de un warp deber\u00edan acceder idealmente a las direcciones de memoria cercanas para que el hardware pueda combinar sus solicitudes.<\/p>\n<p>Considere los datos de part\u00edculas almacenados como:<\/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>Este dise\u00f1o de matriz de estructuras puede ser conveniente para el c\u00f3digo orientado a objetos. Un dise\u00f1o de estructura de arreglos almacena matrices contiguas separadas:<\/p>\n<pre><code>x_positions[]\ny_positions[]\nz_positions[]\nmasses[]<\/code><\/pre>\n<p>El segundo dise\u00f1o puede proporcionar un mejor coalescencia cuando cada subproceso lee el mismo campo para una part\u00edcula diferente. El mejor dise\u00f1o a\u00fan depende de qu\u00e9 campos se accedan juntos.<\/p>\n<p>Los tipos de vectores no corrigen autom\u00e1ticamente un patr\u00f3n de acceso deficiente. Los desarrolladores deben inspeccionar las direcciones reales solicitadas por los hilos vecinos.<\/p>\n<h2>memoria compartida<\/h2>\n<p>La memoria compartida es una peque\u00f1a \u00e1rea de memoria de baja latencia accesible por subprocesos en el mismo bloque. Puede reducir las lecturas repetidas de la memoria global.<\/p>\n<p>Una plantilla unidimensional puede cargar un bloque de valores y su halo en la memoria compartida:<\/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>La llamada a <code>cuda.syncthreads()<\/code> garantiza que cada subproceso termine de cargar sus valores antes de que comience el c\u00e1lculo de la plantilla.<\/p>\n<p>La memoria compartida no debe usarse autom\u00e1ticamente. La asignaci\u00f3n excesiva de memoria compartida puede reducir la ocupaci\u00f3n, agregar el costo de sincronizaci\u00f3n y hacer que el kernel sea m\u00e1s lento. Se requiere perfilado.<\/p>\n<h2>Fusi\u00f3n de grano<\/h2>\n<p>Las operaciones de matriz separadas a menudo crean varios lanzamientos de kernel y matrices intermedios:<\/p>\n<pre><code>velocity += dt * acceleration\nposition += dt * velocity\nenergy = compute_energy(position, velocity)<\/code><\/pre>\n<p>Un kernel fusionado puede calcular las tres actualizaciones mientras los valores necesarios permanecen en los registros. Esto reduce la sobrecarga de lanzamiento y el tr\u00e1fico de memoria global.<\/p>\n<p>La fusi\u00f3n es m\u00e1s beneficiosa cuando las operaciones son simples y est\u00e1n vinculadas a la memoria. Fusionar demasiado trabajo puede aumentar el uso de registros, reducir la ocupaci\u00f3n y hacer que el kernel sea dif\u00edcil de mantener.<\/p>\n<h2>Elegir el tama\u00f1o del bloque<\/h2>\n<p>Los valores como 128 o 256 hilos por bloque son puntos de partida razonables, no como Optima universal.<\/p>\n<p>El mejor tama\u00f1o de bloque depende de:<\/p>\n<ul>\n<li>Registros utilizados por hilo<\/li>\n<li>Memoria compartida utilizada por bloque<\/li>\n<li>Divergencia de ramas<\/li>\n<li>Mezcla de instrucciones<\/li>\n<li>Comportamiento de acceso a la memoria<\/li>\n<li>La arquitectura de GPU objetivo<\/li>\n<\/ul>\n<p>Una regla como lanzar al menos el doble de bloques que los multiprocesadores de transmisi\u00f3n puede ser un experimento inicial \u00fatil, pero no garantiza el m\u00e1ximo rendimiento. Las calculadoras de ocupaci\u00f3n y las herramientas de perfilado deben guiar la configuraci\u00f3n final.<\/p>\n<h2>F\u00edsica diferenciable<\/h2>\n<p>La simulaci\u00f3n diferenciable calcula c\u00f3mo cambia una salida con respecto a las entradas como las propiedades del material, las fuerzas, la geometr\u00eda o las condiciones iniciales.<\/p>\n<p>Warp puede registrar las operaciones del kernel compatible y ejecutar la diferenciaci\u00f3n autom\u00e1tica en modo inverso. Esto se puede utilizar para problemas inversos, optimizaci\u00f3n del dise\u00f1o, estimaci\u00f3n de par\u00e1metros e integraci\u00f3n con flujos de trabajo de aprendizaje autom\u00e1tico.<\/p>\n<p>Un patr\u00f3n simplificado utiliza una cinta:<\/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>No se garantiza la diferenciaci\u00f3n autom\u00e1tica para cada kernel. Las sobreescrituras en el lugar, los at\u00f3micos no deterministas, el c\u00f3digo nativo externo, la l\u00f3gica discontinua y las operaciones no admitidas pueden requerir reformulaci\u00f3n o gradientes personalizados.<\/p>\n<p>Esta capacidad se conecta directamente con aplicaciones como la optimizaci\u00f3n adjunta y <a href=\"https:\/\/matforge.org\/physics-informed-neural-networks-pinns-for-scientific-simulations\/\" rel=\"nofollow\" target=\"_blank\">Redes neuronales informadas por la f\u00edsica<\/a>.<\/p>\n<h2>Cuando una GPU puede no ser m\u00e1s r\u00e1pida<\/h2>\n<p>No existe un umbral de tama\u00f1o de problema universal en el que una GPU se vuelva m\u00e1s r\u00e1pida que una CPU. El cruce depende del hardware, la precisi\u00f3n, el movimiento de datos, la estructura del algoritmo, la calidad del compilador y la frecuencia con la que se reutilizan los mismos datos residentes en el dispositivo.<\/p>\n<table class=\"custom-table\">\n<tbody><tr>\n<th>carga de trabajo<\/th>\n<th>Adecuaci\u00f3n de GPU<\/th>\n<th>Consideraci\u00f3n principal<\/th>\n<\/tr>\n<tr>\n<td>Actualizaci\u00f3n de part\u00edculas grandes<\/td>\n<td>A menudo fuerte<\/td>\n<td>muchas operaciones independientes similares<\/td>\n<\/tr>\n<tr>\n<td>Plantilla grande regular<\/td>\n<td>A menudo fuerte<\/td>\n<td>Acceso predecible a la memoria en paralelo<\/td>\n<\/tr>\n<tr>\n<td>\u00c1lgebra lineal densa<\/td>\n<td>Fuerte al usar bibliotecas optimizadas<\/td>\n<td>Alta intensidad aritm\u00e9tica<\/td>\n<\/tr>\n<tr>\n<td>peque\u00f1a simulaci\u00f3n<\/td>\n<td>dependiente<\/td>\n<td>Lanzamiento y transferencia Los gastos generales pueden dominar<\/td>\n<\/tr>\n<tr>\n<td>Trasversal de gr\u00e1ficos irregulares<\/td>\n<td>Mezclado<\/td>\n<td>Divergencia y acceso a la memoria impredecible<\/td>\n<\/tr>\n<tr>\n<td>Flujo de trabajo enlazado a E\/S<\/td>\n<td>por lo general limitado<\/td>\n<td>La GPU no puede eliminar un cuello de botella de almacenamiento o de red<\/td>\n<\/tr>\n<tr>\n<td>Algoritmo fuertemente secuencial<\/td>\n<td>por lo general d\u00e9bil<\/td>\n<td>trabajo independiente insuficiente<\/td>\n<\/tr>\n<\/tbody><\/table>\n<p>El flujo de trabajo correcto es medir la versi\u00f3n de la CPU, la versi\u00f3n de la biblioteca GPU y la versi\u00f3n del kernel personalizado con datos representativos. La <a href=\"https:\/\/matforge.org\/scientific-computing-performance-profiling-identifying-bottlenecks\/\" rel=\"nofollow\" target=\"_blank\">Gu\u00eda de perfiles de rendimiento de computaci\u00f3n cient\u00edfica<\/a> explica c\u00f3mo identificar el cuello de botella real antes de optimizarlo.<\/p>\n<h2>Evitar la transferencia de gastos generales<\/h2>\n<p>Las transferencias repetidas entre la memoria del dispositivo y del dispositivo pueden eliminar el beneficio del c\u00e1lculo de la GPU.<\/p>\n<p>Un bucle ineficiente puede seguir este patr\u00f3n:<\/p>\n<ol>\n<li>Copiar datos a la GPU.<\/li>\n<li>Ejecute un kernel peque\u00f1o.<\/li>\n<li>Copie el resultado a la CPU.<\/li>\n<li>modificarlo en la CPU.<\/li>\n<li>c\u00f3pielo de nuevo a la GPU.<\/li>\n<\/ol>\n<p>Un mejor dise\u00f1o mantiene el estado de simulaci\u00f3n en el dispositivo durante muchos pasos y transfiere solo la salida requerida para la visualizaci\u00f3n, los puntos de control o el an\u00e1lisis.<\/p>\n<p>Los flujos asincr\u00f3nicos, la memoria del host anclada, la comunicaci\u00f3n superpuesta y las funciones de memoria unificada pueden ayudar, pero solo deben introducirse despu\u00e9s de que se hayan medido los costos de transferencia ordinarios.<\/p>\n<h2>El estudio de caso XLB<\/h2>\n<p>El proyecto XLB de Autodesk Research proporciona un ejemplo \u00fatil de simulaci\u00f3n de GPU nativa de Python. XLB es una biblioteca Boltzmann de c\u00f3digo abierto con varios backends computacionales, incluido Nvidia Warp.<\/p>\n<p>En las configuraciones de referencia informadas por Autodesk Research y NVIDIA, el backend warp alcanz\u00f3 un rendimiento cercano a la implementaci\u00f3n comparada de C++\/OpenCl FluidX3D para un caso espec\u00edfico de cavidad controlada por TAPA. Una comparaci\u00f3n separada inform\u00f3 una aceleraci\u00f3n de ocho veces aproximada sobre el backend JAX de XLB en el hardware y las configuraciones seleccionados.<\/p>\n<p>El equipo tambi\u00e9n demostr\u00f3 un enfoque fuera de n\u00facleo en un grupo GH200 de ocho nodos con un dominio de aproximadamente 50 mil millones de celdas de red. Estos resultados se aplican a las opciones de soluci\u00f3n, problema, hardware y la implementaci\u00f3n reportada. No deben tratarse como garant\u00edas generales de aceleraci\u00f3n para los n\u00facleos de Python.<\/p>\n<p>El estudio de caso completo est\u00e1 disponible en <a href=\"https:\/\/developer.nvidia.com\/blog\/autodesk-research-brings-warp-speed-to-computational-fluid-dynamics-on-nvidia-gh200\/\" rel=\"nofollow\" target=\"_blank\">Autodesk Research lleva la velocidad de la deformaci\u00f3n a la din\u00e1mica de fluidos computacional en NVIDIA GH200<\/a>. El c\u00f3digo de proyecto actual est\u00e1 disponible en el repositorio <a href=\"https:\/\/github.com\/Autodesk\/XLB\" rel=\"nofollow\" target=\"_blank\">Autodesk XLB<\/a>.<\/p>\n<h2>C\u00f3mo comparar correctamente un kernel<\/h2>\n<p>Las operaciones de GPU suelen ser as\u00edncronas. Medir solo la duraci\u00f3n de la llamada a la funci\u00f3n de Python puede informar el tiempo necesario para poner en cola el kernel en lugar del tiempo necesario para ejecutarlo.<\/p>\n<p>Un punto de referencia b\u00e1sico debe:<\/p>\n<ol>\n<li>Ejecute el kernel varias veces para activar la compilaci\u00f3n y el calentamiento.<\/li>\n<li>Sincronice antes de iniciar el temporizador.<\/li>\n<li>Ejecute varias iteraciones medidas.<\/li>\n<li>Sincronice antes de detener el temporizador.<\/li>\n<li>Reporte el promedio y la variaci\u00f3n entre las repeticiones.<\/li>\n<li>Separe el tiempo de transferencia de datos desde el tiempo de ejecuci\u00f3n del kernel.<\/li>\n<li>Verifique que las versiones de CPU y GPU produzcan resultados equivalentes.<\/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 comparaci\u00f3n debe usar un tama\u00f1o de simulaci\u00f3n realista e incluir el flujo de trabajo completo cuando los costos de transferencia o preprocesamiento importan.<\/p>\n<p>La metodolog\u00eda de precisi\u00f3n de trabajo tambi\u00e9n puede comparar el tiempo de ejecuci\u00f3n con el error num\u00e9rico. La <a href=\"https:\/\/docs.sciml.ai\/DiffEqDevDocs\/stable\/alg_dev\/benchmarks\/\" rel=\"nofollow\" target=\"_blank\">Documentaci\u00f3n de referencia de SCML<\/a> proporciona ejemplos de este estilo de evaluaci\u00f3n.<\/p>\n<h2>un flujo de trabajo de desarrollo pr\u00e1ctico<\/h2>\n<ol>\n<li>Implemente y verifique una versi\u00f3n de referencia de CPU clara.<\/li>\n<li>Perfil de la aplicaci\u00f3n para localizar la operaci\u00f3n dominante.<\/li>\n<li>Pruebe una biblioteca de GPU optimizada antes de escribir un kernel.<\/li>\n<li>Mantenga los datos reutilizados con frecuencia en el dispositivo.<\/li>\n<li>Escriba el kernel correcto m\u00e1s simple.<\/li>\n<li>Validarlo contra el resultado de la CPU.<\/li>\n<li>Mida las transferencias de memoria y la ejecuci\u00f3n por separado.<\/li>\n<li>Inspeccione el uso de coalescencia, ocupaci\u00f3n, divergencia y registro.<\/li>\n<li>Pruebe la memoria compartida o la fusi\u00f3n solo cuando la creaci\u00f3n de perfiles los admita.<\/li>\n<li>Compare varios tama\u00f1os de problemas y arquitecturas de GPU.<\/li>\n<\/ol>\n<h2>Elegir un marco<\/h2>\n<table class=\"custom-table\">\n<tbody><tr>\n<th>Requisito<\/th>\n<th>Posible punto de partida<\/th>\n<\/tr>\n<tr>\n<td>Flujo de trabajo de GPU de estilo numpy existente<\/td>\n<td>CUPY con operaciones integradas o kernels personalizados<\/td>\n<\/tr>\n<tr>\n<td>Peque\u00f1a cantidad de kernels de Python de estilo CUDA<\/td>\n<td>Numba-CUDA, despu\u00e9s de revisar su estado de soporte actual<\/td>\n<\/tr>\n<tr>\n<td>Simulaci\u00f3n, geometr\u00eda y kernels diferenciables<\/td>\n<td>Deformaci\u00f3n de Nvidia<\/td>\n<\/tr>\n<tr>\n<td>Control m\u00e1ximo espec\u00edfico de NVIDIA<\/td>\n<td>cud\u00e1 C++<\/td>\n<\/tr>\n<tr>\n<td>Objetivo de GPU AMD con fuente CUDA<\/td>\n<td>Cadera y Rocm<\/td>\n<\/tr>\n<tr>\n<td>C++ port\u00e1til en varios backends<\/td>\n<td>SYCL, Kokkos u otra capa de rendimiento-portabilidad<\/td>\n<\/tr>\n<\/tbody><\/table>\n<h2>Errores comunes de programaci\u00f3n del kernel<\/h2>\n<ul>\n<li>Mover datos entre la CPU y la GPU dentro de cada paso del tiempo<\/li>\n<li>Ignorar hilos fuera de los l\u00edmites<\/li>\n<li>Usando un algoritmo ineficiente y esperando hardware para arreglar su escala<\/li>\n<li>Acceder a la memoria con un dise\u00f1o no coalescido<\/li>\n<li>Agregar memoria compartida sin medir si ayuda<\/li>\n<li>lanzando muchos granos peque\u00f1os en lugar de considerar la fusi\u00f3n<\/li>\n<li>Uso de registros excesivos o memoria compartida por bloque<\/li>\n<li>Benchmarking c\u00f3digo as\u00edncrono sin sincronizaci\u00f3n<\/li>\n<li>Comparaci\u00f3n de salidas sin comprobar la precisi\u00f3n num\u00e9rica<\/li>\n<li>Uso de reclamos de rendimiento fijo en diferentes GPU y cargas de trabajo<\/li>\n<li>Suponiendo que la fuente port\u00e1til proporcione un rendimiento port\u00e1til<\/li>\n<li>Aplicaci\u00f3n de diferenciaci\u00f3n autom\u00e1tica a operaciones in situ no soportadas<\/li>\n<\/ul>\n<h2>Gu\u00edas relacionadas<\/h2>\n<ul>\n<li><a href=\"https:\/\/matforge.org\/gpu-accelerated-scientific-computing-cupy-numba-cudf-compared\/\" rel=\"nofollow\" target=\"_blank\">acelerado de GPU Computaci\u00f3n cient\u00edfica: comparado con CUPY, NUMBA y CUDF<\/a>: compare las herramientas de Python de alto nivel para las cargas de trabajo de GPU.<\/li>\n<li><a href=\"https:\/\/matforge.org\/gpu-acceleration-for-fipy-simulations-cupy-and-numba-integration-guide\/\" rel=\"nofollow\" target=\"_blank\">aceleraci\u00f3n de la GPU Para simulaciones FIPY: Gu\u00eda de integraci\u00f3n de CUPY y Numba<\/a>: explore los posibles componentes de la GPU en torno a un flujo de trabajo de Fipy.<\/li>\n<li><a href=\"https:\/\/matforge.org\/physics-informed-neural-networks-pinns-for-scientific-simulations\/\" rel=\"nofollow\" target=\"_blank\">informado de f\u00edsica Redes neuronales para simulaciones cient\u00edficas<\/a> \u2014 Revisar la relaci\u00f3n entre los modelos f\u00edsicos y el aprendizaje diferenciable.<\/li>\n<li><a href=\"https:\/\/matforge.org\/when-to-use-fem-fvm-fdm\/\" rel=\"nofollow\" target=\"_blank\">cu\u00e1ndo utilizar FEM, FVM y FDM<\/a> \u2014 Seleccione un espacio Discretizaci\u00f3n antes de optimizar su implementaci\u00f3n.<\/li>\n<li><a href=\"https:\/\/matforge.org\/scientific-computing-performance-profiling-identifying-bottlenecks\/\" rel=\"nofollow\" target=\"_blank\">Perfil de rendimiento: Identificaci\u00f3n de cuellos de botella en el c\u00f3digo cient\u00edfico<\/a>: busque las limitaciones de CPU, memoria, transferencia y E\/S.<\/li>\n<li><a href=\"https:\/\/matforge.org\/hpc-python-workflows-from-laptop-to-supercomputer\/\" rel=\"nofollow\" target=\"_blank\">Flujos de trabajo de HPC Python: desde la computadora port\u00e1til a la supercomputadora<\/a>: planifique la progresi\u00f3n del desarrollo local a los grandes sistemas.<\/li>\n<\/ul>\n<h2>Conclusi\u00f3n<\/h2>\n<p>Los kernels de GPU personalizados permiten a los desarrolladores cient\u00edficos traducir part\u00edculas, celdas de cuadr\u00edcula, reglas de materiales y otras operaciones f\u00edsicas directamente en funciones masivamente paralelas. Proporcionan control sobre la indexaci\u00f3n, el acceso a la memoria, la sincronizaci\u00f3n, la fusi\u00f3n del kernel y la optimizaci\u00f3n espec\u00edfica del dispositivo.<\/p>\n<p>CUPY ofrece operaciones de matriz de alto nivel e interfaces de kernel personalizadas. Numba-Cuda proporciona un modelo similar a CUDA en Python, mientras que Nvidia Warp agrega primitivas de simulaci\u00f3n, operaciones de mosaico y diferenciaci\u00f3n autom\u00e1tica. CUDA y HIP proporcionan control C++ de nivel inferior, y los marcos port\u00e1tiles de rendimiento admiten objetivos de hardware m\u00e1s amplios.<\/p>\n<p>Las mayores ganancias de rendimiento rara vez provienen de cambiar la sintaxis por s\u00ed sola. Provienen de elegir un algoritmo paralelo, mantener datos en el dispositivo, reducir el tr\u00e1fico de memoria, utilizar un dise\u00f1o de datos adecuado y eliminar operaciones intermedias innecesarias.<\/p>\n<p>Un kernel personalizado debe desarrollarse solo despu\u00e9s de que la creaci\u00f3n de perfiles identifique un cuello de botella real. Debe validarse contra una soluci\u00f3n de referencia y compararse con la sincronizaci\u00f3n, los tama\u00f1os de problemas representativos y la contabilidad completa de los costos de transferencia de memoria.<\/p>\n<p>Cuando se cumplen esas condiciones, las herramientas de kernel basadas en Python pueden respaldar una simulaci\u00f3n f\u00edsica seria sin obligar a los investigadores a mover cada parte de la aplicaci\u00f3n a C++ de bajo nivel.<\/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\"> 13<\/span> <span class=\"rt-label rt-postfix\">minutes<\/span><\/span>Aprenda a escribir kernels de GPU personalizados para simulaci\u00f3n de f\u00edsica con Nvidia Warp, CUDA Python y HIP\/ROCM. Cubre patrones de rendimiento, f\u00edsica diferenciable y despliegues de producci\u00f3n.<\/p>\n","protected":false,"raw":"Aprenda a escribir kernels de GPU personalizados para simulaci\u00f3n de f\u00edsica con Nvidia Warp, CUDA Python y HIP\/ROCM. Cubre patrones de rendimiento, f\u00edsica diferenciable y despliegues de producci\u00f3n."},"author":2,"featured_media":0,"comment_status":"open","ping_status":"open","sticky":false,"template":"","format":"standard","meta":{"_locale":"es_ES","_original_post":"https:\/\/matforge.org\/?p=482","iawp_total_views":0,"footnotes":""},"categories":[1],"tags":[],"class_list":["post-539","post","type-post","status-publish","format-standard","hentry","category-simulation-modeling-projects","es-ES"],"yoast_head":"<!-- This site is optimized with the Yoast SEO plugin v28.1 - https:\/\/yoast.com\/product\/yoast-seo-wordpress\/ -->\n<title>Kernels de GPU personalizados para simulaciones de f\u00edsica<\/title>\n<meta name=\"description\" content=\"Aprenda c\u00f3mo escribir y optimizar los n\u00facleos de GPU para simulaciones de f\u00edsica utilizando el dise\u00f1o numba, warp, cupy, CUDA, HIP y memoria.\" \/>\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\/es\/gpu-kernel-programming-custom-physics-simulation\/\" \/>\n<meta property=\"og:locale\" content=\"es_ES\" \/>\n<meta property=\"og:type\" content=\"article\" \/>\n<meta property=\"og:title\" content=\"Kernels de GPU personalizados para simulaciones de f\u00edsica\" \/>\n<meta property=\"og:description\" content=\"Aprenda c\u00f3mo escribir y optimizar los n\u00facleos de GPU para simulaciones de f\u00edsica utilizando el dise\u00f1o numba, warp, cupy, CUDA, HIP y memoria.\" \/>\n<meta property=\"og:url\" content=\"https:\/\/matforge.org\/es\/gpu-kernel-programming-custom-physics-simulation\/\" \/>\n<meta property=\"og:site_name\" content=\"matforge.org\" \/>\n<meta property=\"article:published_time\" content=\"2026-07-22T08:18:02+00:00\" \/>\n<meta name=\"author\" content=\"Elena Markovska\" \/>\n<meta name=\"twitter:card\" content=\"summary_large_image\" \/>\n<meta name=\"twitter:label1\" content=\"Escrito por\" \/>\n\t<meta name=\"twitter:data1\" content=\"Elena Markovska\" \/>\n\t<meta name=\"twitter:label2\" content=\"Tiempo de lectura\" \/>\n\t<meta name=\"twitter:data2\" content=\"21 minutos\" \/>\n<script type=\"application\/ld+json\" class=\"yoast-schema-graph\">{\"@context\":\"https:\\\/\\\/schema.org\",\"@graph\":[{\"@type\":\"Article\",\"@id\":\"https:\\\/\\\/matforge.org\\\/es\\\/gpu-kernel-programming-custom-physics-simulation\\\/#article\",\"isPartOf\":{\"@id\":\"https:\\\/\\\/matforge.org\\\/es\\\/gpu-kernel-programming-custom-physics-simulation\\\/\"},\"author\":{\"name\":\"Elena Markovska\",\"@id\":\"https:\\\/\\\/matforge.org\\\/#\\\/schema\\\/person\\\/980162bb5de46742daece973661d93da\"},\"headline\":\"Programaci\u00f3n de kernel de GPU para simulaci\u00f3n de f\u00edsica personalizada: escribir kernels personalizados de alto rendimiento en Python\",\"datePublished\":\"2026-07-22T08:18:02+00:00\",\"mainEntityOfPage\":{\"@id\":\"https:\\\/\\\/matforge.org\\\/es\\\/gpu-kernel-programming-custom-physics-simulation\\\/\"},\"wordCount\":3525,\"commentCount\":0,\"articleSection\":[\"Simulaci\u00f3n &amp; Proyectos de modelado\"],\"inLanguage\":\"es\",\"potentialAction\":[{\"@type\":\"CommentAction\",\"name\":\"Comment\",\"target\":[\"https:\\\/\\\/matforge.org\\\/es\\\/gpu-kernel-programming-custom-physics-simulation\\\/#respond\"]}]},{\"@type\":\"WebPage\",\"@id\":\"https:\\\/\\\/matforge.org\\\/es\\\/gpu-kernel-programming-custom-physics-simulation\\\/\",\"url\":\"https:\\\/\\\/matforge.org\\\/es\\\/gpu-kernel-programming-custom-physics-simulation\\\/\",\"name\":\"Kernels de GPU personalizados para simulaciones de f\u00edsica\",\"isPartOf\":{\"@id\":\"https:\\\/\\\/matforge.org\\\/#website\"},\"datePublished\":\"2026-07-22T08:18:02+00:00\",\"author\":{\"@id\":\"https:\\\/\\\/matforge.org\\\/#\\\/schema\\\/person\\\/980162bb5de46742daece973661d93da\"},\"description\":\"Aprenda c\u00f3mo escribir y optimizar los n\u00facleos de GPU para simulaciones de f\u00edsica utilizando el dise\u00f1o numba, warp, cupy, CUDA, HIP y memoria.\",\"breadcrumb\":{\"@id\":\"https:\\\/\\\/matforge.org\\\/es\\\/gpu-kernel-programming-custom-physics-simulation\\\/#breadcrumb\"},\"inLanguage\":\"es\",\"potentialAction\":[{\"@type\":\"ReadAction\",\"target\":[\"https:\\\/\\\/matforge.org\\\/es\\\/gpu-kernel-programming-custom-physics-simulation\\\/\"]}]},{\"@type\":\"BreadcrumbList\",\"@id\":\"https:\\\/\\\/matforge.org\\\/es\\\/gpu-kernel-programming-custom-physics-simulation\\\/#breadcrumb\",\"itemListElement\":[{\"@type\":\"ListItem\",\"position\":1,\"name\":\"Home\",\"item\":\"https:\\\/\\\/matforge.org\\\/es\\\/\"},{\"@type\":\"ListItem\",\"position\":2,\"name\":\"Programaci\u00f3n de kernel de GPU para simulaci\u00f3n de f\u00edsica personalizada: escribir kernels personalizados de alto rendimiento 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\":\"es\"},{\"@type\":\"Person\",\"@id\":\"https:\\\/\\\/matforge.org\\\/#\\\/schema\\\/person\\\/980162bb5de46742daece973661d93da\",\"name\":\"Elena Markovska\",\"image\":{\"@type\":\"ImageObject\",\"inLanguage\":\"es\",\"@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":"Kernels de GPU personalizados para simulaciones de f\u00edsica","description":"Aprenda c\u00f3mo escribir y optimizar los n\u00facleos de GPU para simulaciones de f\u00edsica utilizando el dise\u00f1o numba, warp, cupy, CUDA, HIP y memoria.","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\/es\/gpu-kernel-programming-custom-physics-simulation\/","og_locale":"es_ES","og_type":"article","og_title":"Kernels de GPU personalizados para simulaciones de f\u00edsica","og_description":"Aprenda c\u00f3mo escribir y optimizar los n\u00facleos de GPU para simulaciones de f\u00edsica utilizando el dise\u00f1o numba, warp, cupy, CUDA, HIP y memoria.","og_url":"https:\/\/matforge.org\/es\/gpu-kernel-programming-custom-physics-simulation\/","og_site_name":"matforge.org","article_published_time":"2026-07-22T08:18:02+00:00","author":"Elena Markovska","twitter_card":"summary_large_image","twitter_misc":{"Escrito por":"Elena Markovska","Tiempo de lectura":"21 minutos"},"schema":{"@context":"https:\/\/schema.org","@graph":[{"@type":"Article","@id":"https:\/\/matforge.org\/es\/gpu-kernel-programming-custom-physics-simulation\/#article","isPartOf":{"@id":"https:\/\/matforge.org\/es\/gpu-kernel-programming-custom-physics-simulation\/"},"author":{"name":"Elena Markovska","@id":"https:\/\/matforge.org\/#\/schema\/person\/980162bb5de46742daece973661d93da"},"headline":"Programaci\u00f3n de kernel de GPU para simulaci\u00f3n de f\u00edsica personalizada: escribir kernels personalizados de alto rendimiento en Python","datePublished":"2026-07-22T08:18:02+00:00","mainEntityOfPage":{"@id":"https:\/\/matforge.org\/es\/gpu-kernel-programming-custom-physics-simulation\/"},"wordCount":3525,"commentCount":0,"articleSection":["Simulaci\u00f3n &amp; Proyectos de modelado"],"inLanguage":"es","potentialAction":[{"@type":"CommentAction","name":"Comment","target":["https:\/\/matforge.org\/es\/gpu-kernel-programming-custom-physics-simulation\/#respond"]}]},{"@type":"WebPage","@id":"https:\/\/matforge.org\/es\/gpu-kernel-programming-custom-physics-simulation\/","url":"https:\/\/matforge.org\/es\/gpu-kernel-programming-custom-physics-simulation\/","name":"Kernels de GPU personalizados para simulaciones de f\u00edsica","isPartOf":{"@id":"https:\/\/matforge.org\/#website"},"datePublished":"2026-07-22T08:18:02+00:00","author":{"@id":"https:\/\/matforge.org\/#\/schema\/person\/980162bb5de46742daece973661d93da"},"description":"Aprenda c\u00f3mo escribir y optimizar los n\u00facleos de GPU para simulaciones de f\u00edsica utilizando el dise\u00f1o numba, warp, cupy, CUDA, HIP y memoria.","breadcrumb":{"@id":"https:\/\/matforge.org\/es\/gpu-kernel-programming-custom-physics-simulation\/#breadcrumb"},"inLanguage":"es","potentialAction":[{"@type":"ReadAction","target":["https:\/\/matforge.org\/es\/gpu-kernel-programming-custom-physics-simulation\/"]}]},{"@type":"BreadcrumbList","@id":"https:\/\/matforge.org\/es\/gpu-kernel-programming-custom-physics-simulation\/#breadcrumb","itemListElement":[{"@type":"ListItem","position":1,"name":"Home","item":"https:\/\/matforge.org\/es\/"},{"@type":"ListItem","position":2,"name":"Programaci\u00f3n de kernel de GPU para simulaci\u00f3n de f\u00edsica personalizada: escribir kernels personalizados de alto rendimiento 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":"es"},{"@type":"Person","@id":"https:\/\/matforge.org\/#\/schema\/person\/980162bb5de46742daece973661d93da","name":"Elena Markovska","image":{"@type":"ImageObject","inLanguage":"es","@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\/539","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=539"}],"version-history":[{"count":1,"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/posts\/539\/revisions"}],"predecessor-version":[{"id":750,"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/posts\/539\/revisions\/750"}],"wp:attachment":[{"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/media?parent=539"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/categories?post=539"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/matforge.org\/wp-json\/wp\/v2\/tags?post=539"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}