Triple buffering — éviter les race conditions

Si le CPU écrit dans un buffer pendant que le GPU le lit, on a une race condition. La solution : 3 copies du buffer d'uniforms, une par frame en vol.

static constexpr int MAX_FRAMES_IN_FLIGHT = 3;

MTL::Buffer* m_p_uniformBuffers[MAX_FRAMES_IN_FLIGHT];
dispatch_semaphore_t m_semaphore;
int m_frameIndex = 0;

// Init :
m_semaphore = dispatch_semaphore_create(MAX_FRAMES_IN_FLIGHT);
for (int i = 0; i < MAX_FRAMES_IN_FLIGHT; i++) {
    m_p_uniformBuffers[i] = m_p_device->newBuffer(
        sizeof(Uniforms), MTL::ResourceStorageModeShared);
}

// draw() :
dispatch_semaphore_wait(m_semaphore, DISPATCH_TIME_FOREVER);
m_frameIndex = (m_frameIndex + 1) % MAX_FRAMES_IN_FLIGHT;

Uniforms* u = (Uniforms*)m_p_uniformBuffers[m_frameIndex]->contents();
// ... remplir u ...

MTL::CommandBuffer* cmd = m_p_queue->commandBuffer();
cmd->addCompletedHandler(^(MTL::CommandBuffer*) {
    dispatch_semaphore_signal(m_semaphore); // libère le slot quand le GPU termine
});

Alignement mémoire — règle des 256 bytes

Les buffers d'uniforms doivent être alignés sur 256 bytes sur Apple Silicon. Si ta structure Uniforms fait 128 bytes, le stride entre deux frames doit quand même être 256.

size_t alignedUniformSize = (sizeof(Uniforms) + 255) & ~255;
// sizeof(Uniforms) = 128 → alignedUniformSize = 256

Pourquoi la même structure doit exister des deux côtés du bus — et comment le fichier _shared.h garantit le contrat

Le CPU écrit. Le GPU lit. Entre les deux, rien d'autre que de la mémoire brute. Ce chapitre explique pourquoi un même fichier d'en-tête est inclus par ton code C++ et par ton shader Metal, et pourquoi l'alignement mémoire est une question de survie.

Deux compilateurs, une seule mémoire

Ton projet compile en deux passes distinctes :

Passe 1 — metal compile tes shaders

shader.metal → metal → shader.metallib

Passe 2 — clang compile ton C++

→ binaire ARM → l'exécutable.

→ bundle app → dossier de fichiers d'application Apple.

Ces deux compilateurs ne se parlent pas. Quand le CPU remplit un buffer et que le GPU le lit, la seule chose qui garantit que les données correspondent octet pour octet, c'est toi — via le fichier partagé.

/// C++
/* Renderer_shared.h
 * Inclus par Renderer.cpp ET shader.metal
 * C'est le contrat de mémoire entre CPU et GPU */

#ifndef Renderer_shared_h
#define Renderer_shared_h

#include <simd/simd.h>

struct alignas(16) Sun
{
    simd::float4 sunDirection;
    simd::float4 sunColor;
};

struct alignas(16) CameraUniforms
{
    simd::float4x4 viewMatrix;
    simd::float4x4 projectionMatrix;
    simd::float4x4 viewProjectionMatrix;
    simd::float4x4 invViewProjectionMatrix;
    
    simd::float3   cameraPosition;
    float          _pad;
};

struct alignas(16) Uniforms
{   // total : 320 bytes
    Sun             sun;
    CameraUniforms  cameraUniforms;

    simd::float3    mouseState;

    float           frameTime = 0.f;
};

Comment l'inclusion fonctionne des deux côtés

Côté C++, l'include est classique :

// Renderer.hpp
#include "Renderer_shared.h" // struct Uniforms disponible en C++

class Renderer
{
private:
    Uniforms      m_uniforms;
    MTL::Buffer*  m_uniformBuffer;
};

Côté Metal, le compilateur accepte aussi les headers C standard et simd :

// shader.metal
#include <metal_stdlib>

#include "Renderer_shared.h" // même struct, même layout

using namespace metal;

struct VertexOut
{
    float4 position [[position]];
    float4 color;
};

vertex VertexOut vertexShader(VertexIn            in        [[stage_in]],
                              constant Uniforms&  uniforms  [[buffer(1)]])
{
    VertexOut out;
    out.position = uniforms.cameraUniforms.projectionMatrix * uniforms.cameraUniforms.viewMatrix *
                   uniforms.cameraUniforms.modelMatrix * float4(in.position, 1.0);
    return out;
}

Pourquoi l'alignement peut tout casser silencieusement

Le danger principal n'est pas un crash — c'est que le GPU lit des données décalées sans le signaler :

// ⚠️  DANGER — float3 a un comportement différent CPU vs GPU

// côté C++ (clang) :
struct Bad
{
    float   value;      // offset 0,  4 bytes
    float3  position;   // offset 4,  12 bytes  ← clang aligne sur 4
};                      // total : 16 bytes côté CPU

// côté Metal (GPU Apple Silicon) :
struct Bad
{
    float   value;      // offset 0,  4 bytes
    float3  position;   // offset 16, 12 bytes  ← Metal aligne float3 sur 16 !
};                      // total : 28 bytes côté GPU

// → GPU lit position à l'offset 16, CPU l'a écrit à l'offset 4 → garbage

Pourquoi les structs contiennent autant de données

Trois raisons qui s'accumulent :

Un seul aller-retour CPU → GPU

Chaque appel setVertexBuffer a un coût d'encodage. Mieux vaut un buffer dense qu'une dizaine de petits appels par frame.

Les matrices sont incompressibles

Une float4x4 = 64 bytes. Une scène typique aux Matrices, parfois lightSpaceMatrix pour les shadow maps — soit 256 bytes minimum rien que pour les matrices caméra basique.

Les règles d'alignement gonflent la struct

Le padding implicite entre membres peut ajouter 12 à 48 bytes selon l'ordre des champs. On préfère le rendre explicite avec _pad[] plutôt que de subir une surprise.

// Dans Renderer.cpp — assertion de sécurité supplémentaire
static_assert(sizeof(Uniforms) % 16 == 0, "Uniforms doit être alignée sur 16 bytes");

// Dans Renderer_shared.h
static_assert(sizeof(Uniforms) % 272);

On retient : dans le sens des aiguilles d'une montre, en commençant par le haut gauche

Face 0 Top-Left > Top-Right > Back-Left

Face 1 Top-Right > Back-Right > Back-Left

S'il n'y a qu'un triangle, ce serait le numéro 0*.

Dans une vue orthographique 3D

TL > TF > BF > BL

TL > TT > TR > TF

TF > TR > BR > BF

Viewport coordinate system on Metal

Shader Cores & SIMD Groups

Le GPU est organisé en Shader Cores. Chaque core exécute des groupes de threads appelés SIMD groups (équivalent des "warps" NVIDIA ou "wavefronts" AMD). Sur Apple Silicon, un SIMD group = 32 threads qui exécutent exactement la même instruction au même cycle.

/// .metal
// Dans un compute shader Metal, tu déclares la taille de ton threadgroup :
// [[threads_per_threadgroup]] = combien de threads dans un groupe
// Un SIMD group = 32 threads → multiple de 32 = optimal

kernel void mon_compute(uint tid [[thread_position_in_grid]])
{
    // Ce code tourne sur des milliers de threads en parallèle
    // Chaque thread a son propre 'tid' (thread_id)
}

Threadgroups & grille de dispatch

Quand tu lances un compute shader, tu dispatches une grille de threads organisée en threadgroups. Chaque threadgroup partage une mémoire locale ultra-rapide (threadgroup memory) inaccessible aux autres groupes.

/// C++
// Côté CPU (metal-cpp) — on dispatche une grille 2D :
MTL::Size gridSize        = MTL::Size::Make(textureWidth, textureHeight, 1);
MTL::Size threadgroupSize = MTL::Size::Make(16, 16, 1); // 256 threads/groupe

computeCommandEncoder->dispatchThreads(gridSize, threadgroupSize);
// → le GPU lance (textureWidth/16) × (textureHeight/16) threadgroups