L'emplacement de ton curseur de souris dans tes scènes 3D
Le but de ce shader est de lire la texture de profondeur créée au préalable, et selon la nuance de blanc et de noir, retourner une position Z. Il n'affiche rien à l'écran, c'est un calcul/une fonction sur GPU : kernel.
Toutes les données sont créées et envoyées à notre classe principale. Tu les envois via la fonction update() pour remplir un MTL::Buffer unique à ce shader.
/// C++ — Cursor3D_shared.h — update:18/05/26
#include <simd/simd.h>
struct alignas(16) PickUniforms
{
simd::float4x4 invViewProjectionMatrix;
simd::float2 mousePosition;
simd::float2 windowSize;
};Le shader accède à la texture réglée au préalable sur MTL::TextureUsageShaderRead.
invViewProjectionMatrix car tu dois déconstituer dans le sens inverse du pipeline GPU : partir d'un pixel écran → retrouver la position world.
Technique indispensable pour diviser une surface plane (comme un sol carrelé ou un mur avec des fenêtres) de manière à respecter l'effet de profondeur et l'éloignement
Une depth texture est un buffer 2D alloué par Metal en même temps que le color attachment. Pour chaque fragment rendu, le GPU écrit la valeur z_ndc après la division de perspective, remappée dans [0.0, 1.0].
Le format que tu utilises : MTL::PixelFormatDepth32Float stocke un flottant par pixel. Le GPU compare automatiquement la profondeur des nouveaux fragments avec ce qui existe déjà dans la texture qui n'était que des zéros initialement au début de la frame (couleur noir), et n'écrit que si le fragment est plus proche.
/// MSL — Cursor3D.metal — update:18/05/26
#include <metal_stdlib>
#include "../includes/SharedGPU/Cursor3D_shared.h"
using namespace metal;
struct PickResult
{
float3 worldPosition;
float depth;
int isValid;
};
kernel void pickDepth(texture2d<float, access::read> depthTexture [[texture(0)]],
constant PickUniforms& uniforms [[buffer(0)]],
device PickResult& result [[buffer(1)]],
uint2 gid [[thread_position_in_grid]])
{
if (gid.x != 0 || gid.y != 0) // early-exit du dispatch
return; // seul le thread (0, 0) doit faire le travail
int2 pixelCoord = int2(uniforms.mousePosition);
if (pixelCoord.x < 0 || pixelCoord.x >= int(uniforms.windowSize.x) ||
pixelCoord.y < 0 || pixelCoord.y >= int(uniforms.windowSize.y))
{
result.isValid = 0;
return; // si le pointeur de la souris est hors écran, invalide
}
float depth = depthTexture.read(uint2(pixelCoord)).r;
if (depth >= 1.0)
{
result.isValid = 0;
result.depth = depth;
return;
}
// NDC: x [-1, 1], y [-1, 1]
float2 ndc;
ndc.x = (uniforms.mousePosition.x / uniforms.windowSize.x) * 2.0 - 1.0;
ndc.y = 1.0 - (uniforms.mousePosition.y / uniforms.windowSize.y) * 2.0;
// Clip space → World space
float4 clipPos = float4(ndc.x, ndc.y, depth, 1.0);
float4 worldPos = uniforms.invViewProjectionMatrix * clipPos;
worldPos /= worldPos.w; // perspective divide
result.worldPosition = worldPos.xyz;
result.depth = depth;
result.isValid = 1;
}Seul le premier canal, ici rouge est lu : .r car c'est une texture scalaire (un seul attribut).
Un pass de calcul peut exécuter une fonction de noyau (kernel) sur une grille 1D, 2D ou 3D. Chaque thread appartient à un groupe de threads qui fonctionne ensemble et partage un bloc de mémoire commun! On peut aussi les traiter indépendamment bien évidemment.
Dans le cas d'une grille 2D, gid sera un vecteur à 2 attributs.
Threads & ThreadGroups
thread_position_in_grid est un qualificatif d'attribut (attribute qualifier) remarquable par son double crochet [[]]. Il permet aux paramètres du noyau d'être liés aux ressources et aux variables intégrées → ici, la position du thread (fil) dans la grille de la fonction du noyau. Via read().
/// MSL — Cursor3D.hpp — update:18/05/26
#ifndef Cursor3D_hpp
#define Cursor3D_hpp
#include <simd/simd.h>
#include <Metal/Metal.hpp>
#include "../includes/SharedGPU/Cursor3D_shared.h"
struct PickResult
{
simd::float3 worldPosition;
float depth;
int isValid;
};
class MouseDepthPicker
{
public:
MouseDepthPicker(MTL::Device* device, MTL::Library* shaderLibrary,
MTL::PixelFormat pixelFormat, MTL::PixelFormat depthPixelFormat);
~MouseDepthPicker();
void updateUniforms(simd::float4x4 invViewProjectionMatrix, simd::float2 windowSize, simd::float2 mousePosition);
void pick(MTL::CommandBuffer* commandBuffer, MTL::Texture* depthTexture);
PickResult getResult();
private:
void buildPipeline(MTL::Device* device, MTL::Library* shaderLibrary,
MTL::PixelFormat pixelFormat, MTL::PixelFormat depthPixelFormat);
void buildBuffers(MTL::Device* device);
NS::SharedPtr<MTL::ComputePipelineState> m_mousePosComputePipelineState;
NS::SharedPtr<MTL::Buffer> m_uniformBuffer;
NS::SharedPtr<MTL::Buffer> m_resultBuffer;
PickUniforms m_uniforms;
};
#endif /* Cursor3D_hpp *//// C++ — Cursor3D.cpp — update:18/05/26
#include "Cursor3D.hpp"
MouseDepthPicker::MouseDepthPicker(MTL::Device* device, MTL::Library* shaderLibrary,
MTL::PixelFormat pixelFormat, MTL::PixelFormat depthPixelFormat)
{
m_uniforms = {}; // init à nullptr de la structure (sécurité)
buildPipeline(device, shaderLibrary, pixelFormat, depthPixelFormat);
buildBuffers(device);
}
MouseDepthPicker::~MouseDepthPicker() {}
void MouseDepthPicker::buildPipeline(MTL::Device* device, MTL::Library* shaderLibrary,
MTL::PixelFormat pixelFormat, MTL::PixelFormat depthPixelFormat)
{
auto function = NS::TransferPtr(shaderLibrary->newFunction(MTLSTR("pickDepth")));
NS::Error* error = nullptr;
m_mousePosComputePipelineState = NS::TransferPtr(device->newComputePipelineState(function.get(), &error));
assert(!error && "Pipeline creation failed");
}
void MouseDepthPicker::buildBuffers(MTL::Device* device)
{
m_uniformBuffer = NS::TransferPtr(device->newBuffer(sizeof(PickUniforms), MTL::ResourceStorageModeShared));
m_resultBuffer = NS::TransferPtr(device->newBuffer(sizeof(PickResult), MTL::ResourceStorageModeShared));
}
void MouseDepthPicker::updateUniforms(simd::float4x4 invViewProjectionMatrix, simd::float2 windowSize, simd::float2 mousePosition)
{
m_uniforms.windowSize = windowSize;
m_uniforms.mousePosition = mousePosition;
m_uniforms.invViewProjectionMatrix = invViewProjectionMatrix;
memcpy(m_uniformBuffer->contents(), &m_uniforms, sizeof(PickUniforms));
}
void MouseDepthPicker::pick(MTL::CommandBuffer* commandBuffer, MTL::Texture* depthTexture)
{
MTL::ComputeCommandEncoder* computeCommandEncoder = commandBuffer->computeCommandEncoder();
computeCommandEncoder->setComputePipelineState(m_mousePosComputePipelineState.get());
computeCommandEncoder->setTexture(depthTexture, 0);
computeCommandEncoder->setBuffer(m_uniformBuffer.get(), 0, 0);
computeCommandEncoder->setBuffer(m_resultBuffer.get(), 0, 1);
computeCommandEncoder->dispatchThreads(MTL::Size(1, 1, 1), MTL::Size(1, 1, 1));
computeCommandEncoder->endEncoding();
}
PickResult MouseDepthPicker::getResult()
{
PickResult result;
struct GPUResult
{
simd::float3 worldPosition;
float depth;
int isValid;
};
GPUResult* gpuResult = static_cast<GPUResult*>(m_resultBuffer->contents());
result.worldPosition = gpuResult->worldPosition;
result.depth = gpuResult->depth;
result.isValid = gpuResult->isValid != 0;
return result;
}
// MouseDepthPicker::~MouseDepthPicker() = default;
En C++ ce code ne dispatche qu'un Thread, dans le shader, il est alors redondant de vérifier s'il gid > 0.
Intégration à la classe principale
/// C++ — Renderer.hpp — update:18/05/26
#include "Cursor3D.hpp"
private:
MTL::PixelFormat m_depthPixelFormat;
MouseDepthPicker m_cursor3D;
simd::float3 m_cursorPosition3D;/// C++ — Renderer.cpp — update:18/05/26
m_text(m_device, m_shaderLibrary, m_pixelFormat, m_depthPixelFormat, resourcePath),
m_cursor3D(m_device, m_shaderLibrary, m_pixelFormat, m_depthPixelFormat), /// C++ — Renderer.cpp — update:18/05/26 — update()
m_camera.setPosition(position); // après m_globalUniforms = m_camera.updateUniforms();
m_cursor3D.updateUniforms(m_globalUniforms.cameraUniforms.invViewProjectionMatrix, m_windowSize, m_input.mousePosition);
// même nom de fn, ça peut arriver. pas les variables et constantes.
// pour m_windowSize → m_windowUniforms.m_windowSize est aussi valide,
// pour getInvViewProjectionMatrix() → ou m_camera.getInvViewProjectionMatrix()/// C++ — Renderer.cpp — update:18/05/26 — draw()
renderCommandEncoder->endEncoding();
}
m_cursor3D.pick(commandBuffer, m_depth.get());
commandBuffer->presentDrawable(view->currentDrawable());
commandBuffer->commit();
commandBuffer->waitUntilCompleted();
PickResult result = m_cursor3D.getResult();
m_cursorPosition3D.x = roundf(result.worldPosition.x);
m_cursorPosition3D.y = roundf(result.worldPosition.y);
m_cursorPosition3D.z = roundf(result.worldPosition.z);
// ne garde pas
if (frame % 240 == 0)
printf("mouse2D : x = %.10f, y = %.10f\n", m_input.mousePosition.x, m_input.mousePosition.y);
if (frame % 240 == 0)
printf("mouse3D : x = %.10f, y = %.10f, z = %.10f\n", m_cursorPosition3D.x, m_cursorPosition3D.y, m_cursorPosition3D.z);
autoreleasePool->release();
}/// Shell — données dans la console + entrées clavier
./LearnMSL_Makefile.app/Contents/MacOS/LearnMSL_Makefile