Metal-cpp in One Weekend

Version 0.1, 2026-05-02

Overview

这本书用 metal-cpp 写第一个 GPU 程序。我们不会一开始就画三角形。三角形看起来简单，但它会同时引入窗口、drawable、render pass、render pipeline、顶点 shader、片元 shader 和坐标系。对第一次接触 Metal 的读者来说，这些概念会一次性挤在一起。

我们从图像处理开始。图像处理的输入和输出都很具体：一块像素数据进去，另一块像素数据出来。这样你可以先理解 Metal 的核心模型：Device、Command Queue、Command Buffer、Buffer、compute shader 和 Compute Pipeline State。

这本书的风格是一步一步改程序。每章会给出关键代码和必要改动，并在代码后解释这一步新增了什么；完整最终源码放在 src/MetalCppInOneWeekend/ 里作为对照。你应该尽量亲手输入这些代码；敲错之后再修正，是学习图形和 GPU 编程很有效的一部分。

本书使用 C++17。metal-cpp 的实现宏必须只在一个 .cpp 文件中定义一次：

main.cpp

#define NS_PRIVATE_IMPLEMENTATION#define CA_PRIVATE_IMPLEMENTATION#define MTL_PRIVATE_IMPLEMENTATION

第一册不创建窗口，所以所有示例都可以保持为普通 C++。后两册也继续保持离屏输出：第二册用 render pipeline 生成静态渲染图，第三册用 compute workload 做性能实验。

本教程用 CMake 管理参考代码。你不需要先理解复杂的工程结构；第一册只需要三个文件：CMakeLists.txt、main.cpp、Shaders.metal。正文会先按概念逐步写代码，书末给出最终参考代码位置和运行结果。

CMakeLists.txt

add_executable(MetalCppInOneWeekend main.cpp)

后面如果新增 .metal shader，CMake 还要多做一步：调用 Apple 的 metal 和 metallib 工具生成 default.metallib，并把路径传给 C++。

Output an Image

The PPM Image Format

任何图形程序都需要先看到结果。最简单的方式不是窗口，而是写一个图片文件。

PPM 是一个非常朴素的图片格式。它可以用纯文本表示，也可以用二进制表示。我们先用二进制 P6，文件头长这样：

gradient.ppm header

P6256 160255

第一行 P6 表示二进制 RGB。第二行是宽和高。第三行是每个颜色通道的最大值。后面紧跟 width * height 个 RGB 像素。

下面是第一个完整程序。它不使用 Metal，只是确认我们能生成一张图片。

main.cpp

#include <cstdint>#include <filesystem>#include <fstream>#include <iostream>int main(){    constexpr int width = 256;    constexpr int height = 160;    std::filesystem::create_directories("build");    std::ofstream out("build/gradient.ppm", std::ios::binary);    out << "P6\n" << width << " " << height << "\n255\n";    for (int y = 0; y < height; ++y)    {        for (int x = 0; x < width; ++x)        {            const uint8_t rgb[3] = {                static_cast<uint8_t>(x),                static_cast<uint8_t>(y * 255 / (height - 1)),                static_cast<uint8_t>(180),            };            out.write(reinterpret_cast<const char*>(rgb), sizeof(rgb));        }    }    std::cout << "Wrote build/gradient.ppm\n";    return 0;}

运行后打开 build/gradient.ppm，你应该看到一张水平红色渐变、垂直绿色渐变、带固定蓝色分量的图片。

这个阶段还没有 Metal shader，所以 CMake target 只需要 main.cpp：

CMakeLists.txt

add_executable(MetalCppInOneWeekend main.cpp)

这段代码有几个关键点：

像素按行写出。
每一行从左到右写。
行从上到下写。
每个像素有三个字节：red、green、blue。
文件里的颜色是整数 0..255，但 shader 里我们常常会先用 0.0..1.0 思考颜色。

到目前为止，我们还没有用 GPU。下一步才开始把计算交给 Metal。

A Metal Device

Creating the Device

Device 是 Metal 里最重要的对象之一。你可以把它理解成当前 GPU 的入口。几乎所有 GPU 资源都从 Device 创建出来，包括 Buffer、Texture 和 pipeline。

最小的 Metal 程序只需要创建一个 Device：

main.cpp

#define NS_PRIVATE_IMPLEMENTATION#define CA_PRIVATE_IMPLEMENTATION#define MTL_PRIVATE_IMPLEMENTATION#include <Foundation/Foundation.hpp>#include <QuartzCore/QuartzCore.hpp>#include <Metal/Metal.hpp>#include <iostream>int main(){    NS::AutoreleasePool* pool = NS::AutoreleasePool::alloc()->init();    MTL::Device* device = MTL::CreateSystemDefaultDevice();    if (!device)    {        std::cerr << "Metal is not available on this Mac.\n";        pool->release();        return 1;    }    std::cout << "Device: " << device->name()->utf8String() << "\n";    device->release();    pool->release();    return 0;}

这个章节开始依赖 metal-cpp 头文件和 Apple frameworks，所以 CMake 需要增加 include path 和 framework 链接：

CMakeLists.txt

target_include_directories(MetalCppInOneWeekend PRIVATE "${METAL_CPP_ROOT}")target_link_libraries(MetalCppInOneWeekend PRIVATE    "-framework Foundation"    "-framework QuartzCore"    "-framework Metal")

和上一章相比，这一步新增了三件事：

定义 metal-cpp 的三个实现宏。
创建 NS::AutoreleasePool。
调用 MTL::CreateSystemDefaultDevice() 获取默认 GPU。

metal-cpp 是对 Objective-C Metal API 的 C++ 包装，所以你仍然会看到 alloc()、release() 这类 Cocoa 风格的生命周期操作。这里先手动 release()，后面每次新增资源时都保持同一条规则：谁创建，谁在程序结束前释放。

Adding a Command Queue

只有 Device 还不能提交工作。CPU 需要通过 Command Queue 创建 Command Buffer，再把命令交给 GPU。

把下面这几行加在 Device 创建成功之后：

main.cpp

MTL::CommandQueue* queue = device->newCommandQueue();if (!queue){    std::cerr << "Could not create a command queue.\n";    device->release();    pool->release();    return 1;}std::cout << "Command Queue: " << queue << "\n";

程序结束前释放它：

main.cpp

queue->release();device->release();pool->release();

Device, Command Queue, and Command Buffer relationship

现在我们已经有了 CPU 向 GPU 提交工作的入口，但还没有任何 GPU 能执行的代码。下一章会加入第一个 compute shader。

A Buffer Round Trip

Round-Trip Goal

先让 GPU 改一小段数字。CPU 准备数组：

CPU input values

1 2 3 4

GPU 把每个数字乘以 2，CPU 再读回：

Expected CPU readback

2 4 6 8

这个例子很小，但它包含完整的 Metal compute 路径：

创建 Buffer
加载 .metal shader
创建 Compute Pipeline State
创建 Command Buffer
编码 compute 命令
提交并等待完成
从 shared buffer 读回结果

The Shader

新建 Shaders.metal：

Shaders.metal

#include <metal_stdlib>using namespace metal;kernel void double_values(device uint* values [[buffer(0)]],                          uint id [[thread_position_in_grid]]){    values[id] *= 2;}

加入 shader 后，CMake 需要声明 METALLIB_PATH，并添加从 Shaders.metal 生成 default.metallib 的命令：

CMakeLists.txt绿色 = 新增行

set(METALLIB "${CMAKE_BINARY_DIR}/default.metallib")target_compile_definitions(MetalCppInOneWeekend PRIVATE    METALLIB_PATH="${METALLIB}")add_custom_command(    OUTPUT "${METALLIB}"    COMMAND "${CMAKE_COMMAND}" -E make_directory "${CMAKE_BINARY_DIR}/ModuleCache"    COMMAND xcrun -sdk macosx metal            "-fmodules-cache-path=${CMAKE_BINARY_DIR}/ModuleCache"            -c Shaders.metal -o "${CMAKE_BINARY_DIR}/Shaders.air"    COMMAND xcrun -sdk macosx metallib "${CMAKE_BINARY_DIR}/Shaders.air" -o "${METALLIB}"    DEPENDS Shaders.metal)

这会在构建目录里执行两步：先用 xcrun -sdk macosx metal 把 .metal 编译成 .air，再用 xcrun -sdk macosx metallib 生成 default.metallib。

kernel 表示这是 compute shader 入口。values [[buffer(0)]] 表示第 0 号 buffer。thread_position_in_grid 给出当前 thread 的索引。我们会 dispatch 4 个 thread，所以 id 会是 0, 1, 2, 3。

Loading the Shader

主程序需要从 .metallib 里加载函数，再创建 pipeline：

main.cpp

NS::String* makeString(const char* value){    return NS::String::string(value, NS::UTF8StringEncoding);}MTL::ComputePipelineState* makePipeline(MTL::Device* device, const char* functionName){    NS::Error* error = nullptr;    MTL::Library* library = device->newLibrary(makeString(METALLIB_PATH), &error);    if (!library)    {        std::cerr << "Could not load metallib.\n";        return nullptr;    }    MTL::Function* function = library->newFunction(makeString(functionName));    MTL::ComputePipelineState* pipeline = device->newComputePipelineState(function, &error);    function->release();    library->release();    return pipeline;}

这一步新增的概念是 Compute Pipeline State。Metal 不会在每次 dispatch 时临时解释 shader。你先把 shader 函数编译成 pipeline state，后面编码命令时绑定这个 state。

Encoding Work

下面是完整的计算和提交部分：

main.cpp

std::array<uint32_t, 4> input = {1, 2, 3, 4};MTL::Buffer* buffer = device->newBuffer(    input.data(),    input.size() * sizeof(uint32_t),    MTL::ResourceStorageModeShared);MTL::CommandBuffer* commandBuffer = queue->commandBuffer();MTL::ComputeCommandEncoder* encoder = commandBuffer->computeCommandEncoder();encoder->setComputePipelineState(pipeline);encoder->setBuffer(buffer, 0, 0);encoder->dispatchThreads(    MTL::Size::Make(input.size(), 1, 1),    MTL::Size::Make(input.size(), 1, 1));encoder->endEncoding();commandBuffer->commit();commandBuffer->waitUntilCompleted();auto* output = static_cast<uint32_t*>(buffer->contents());std::cout << output[0] << " "          << output[1] << " "          << output[2] << " "          << output[3] << "\n";

这里最容易混淆的是 setBuffer(buffer, 0, 0)。第一个 0 是 buffer offset，第二个 0 是 shader 里的 [[buffer(0)]]。换句话说，我们把 C++ 里的 buffer 绑定到 shader 的第 0 个 buffer 参数上。

dispatchThreads 里的第一个 size 是总 thread 数，第二个 size 是每个 threadgroup 的 thread 数。这个例子只有 4 个元素，所以我们直接用一个 threadgroup 包住 4 个 thread。

First Compute Image

Moving From Numbers to Pixels

上一章用 GPU 改了 4 个整数。现在用同样的思路让 GPU 写一张图片。

我们仍然先用 Buffer，因为它比 Texture 更容易看清内存布局。每个像素用四个字节存储：

RGBA buffer layout

red, green, blue, alpha

用 C++ 表示就是 uchar4 或四个 uint8_t。

The Gradient Kernel

新 shader 根据当前 thread 的二维坐标生成颜色：

Shaders.metal

#include <metal_stdlib>using namespace metal;kernel void make_gradient(device uchar4* pixels [[buffer(0)]],                          constant uint2& size [[buffer(1)]],                          uint2 id [[thread_position_in_grid]]){    if (id.x >= size.x || id.y >= size.y)    {        return;    }    const uint index = id.y * size.x + id.x;    pixels[index] = uchar4(uchar(id.x * 255 / max(size.x - 1, 1u)),                           uchar(id.y * 255 / max(size.y - 1, 1u)),                           uchar(180),                           uchar(255));}

和 double_values 相比，这里修改了三件事：

thread id 从 uint 变成 uint2，因为图片有 x 和 y。
多了一个 size buffer，用来告诉 shader 图片宽高。
一维数组下标用 id.y * size.x + id.x 算出来。

Dispatching a 2D Grid

C++ 侧也要从一维 dispatch 改成二维 dispatch：

main.cpp

constexpr uint32_t width = 256;constexpr uint32_t height = 160;const UInt2 size{width, height};MTL::Buffer* pixels = device->newBuffer(    width * height * 4,    MTL::ResourceStorageModeShared);MTL::Buffer* sizeBuffer = device->newBuffer(    &size,    sizeof(size),    MTL::ResourceStorageModeShared);encoder->setComputePipelineState(pipeline);encoder->setBuffer(pixels, 0, 0);encoder->setBuffer(sizeBuffer, 0, 1);encoder->dispatchThreads(    MTL::Size::Make(width, height, 1),    MTL::Size::Make(16, 16, 1));

这次总 grid 是 width * height 个 thread。每个 threadgroup 是 16 * 16 个 thread。shader 里的边界判断很重要，因为图片宽高不一定刚好被 threadgroup 大小整除。

Compute grid, threadgroups, and image pixels

Writing the Result

GPU 写的是 RGBA，但 PPM 只需要 RGB，所以保存时每 4 个字节写前 3 个：

main.cpp

void writePPM(const char* path, const uint8_t* rgba, uint32_t width, uint32_t height){    std::ofstream out(path, std::ios::binary);    out << "P6\n" << width << " " << height << "\n255\n";    for (uint32_t i = 0; i < width * height; ++i)    {        out.write(reinterpret_cast<const char*>(rgba + i * 4), 3);    }}

现在我们已经让 GPU 写出第一张图片。

Grayscale

Two Kernels, One Buffer

接下来做第一个真正的图像处理：灰度化。

我们先用一个 kernel 生成彩色测试图，再用第二个 kernel 把同一块 buffer 改成灰度图。这样可以看到同一个 GPU buffer 在两个 command 之间逐步变化。

彩色图 kernel：

Shaders.metal

kernel void make_color_test(device uchar4* pixels [[buffer(0)]],                            constant uint2& size [[buffer(1)]],                            uint2 id [[thread_position_in_grid]]){    if (id.x >= size.x || id.y >= size.y)    {        return;    }    const uint index = id.y * size.x + id.x;    const float2 uv = float2(id) / float2(max(size.x - 1, 1u), max(size.y - 1, 1u));    pixels[index] = uchar4(uchar(uv.x * 255.0),                           uchar(uv.y * 255.0),                           uchar((1.0 - uv.x) * 220.0),                           uchar(255));}

灰度化 kernel：

Shaders.metal

kernel void grayscale(device uchar4* pixels [[buffer(0)]],                      constant uint2& size [[buffer(1)]],                      uint2 id [[thread_position_in_grid]]){    if (id.x >= size.x || id.y >= size.y)    {        return;    }    const uint index = id.y * size.x + id.x;    const uchar4 src = pixels[index];    const float gray = 0.299 * float(src.r)                     + 0.587 * float(src.g)                     + 0.114 * float(src.b);    const uchar value = uchar(gray);    pixels[index] = uchar4(value, value, value, 255);}

这一步只修改 shader，不需要改变 buffer 的格式。灰度公式来自人眼对不同颜色亮度的敏感度：绿色贡献最大，蓝色贡献最小。

Reusing Dispatch Code

两个 kernel 的 dispatch 过程完全一样，所以可以把编码命令抽成函数：

main.cpp

void dispatch(MTL::CommandQueue* queue,              MTL::ComputePipelineState* pipeline,              MTL::Buffer* pixels,              MTL::Buffer* sizeBuffer,              uint32_t width,              uint32_t height){    MTL::CommandBuffer* commandBuffer = queue->commandBuffer();    MTL::ComputeCommandEncoder* encoder = commandBuffer->computeCommandEncoder();    encoder->setComputePipelineState(pipeline);    encoder->setBuffer(pixels, 0, 0);    encoder->setBuffer(sizeBuffer, 0, 1);    encoder->dispatchThreads(        MTL::Size::Make(width, height, 1),        MTL::Size::Make(16, 16, 1));    encoder->endEncoding();    commandBuffer->commit();    commandBuffer->waitUntilCompleted();}

然后主程序依次调用：

main.cpp

dispatch(queue, makeColor, pixels, sizeBuffer, width, height);writePPM("build/color.ppm", static_cast<const uint8_t*>(pixels->contents()), width, height);dispatch(queue, makeGray, pixels, sizeBuffer, width, height);writePPM("build/grayscale.ppm", static_cast<const uint8_t*>(pixels->contents()), width, height);

这就是最小图像处理 pipeline：生成输入、执行处理、写出输出。

Color test image generated by the Metal compute shader — 运行彩色 kernel 后得到的 `color.ppm`，红色从左到右增加，绿色从上到下增加，蓝色从左到右减弱。

Grayscale image generated by the Metal compute shader — 运行灰度 kernel 后得到的 `grayscale.ppm`。它复用同一块 GPU buffer，把彩色图改写成灰度图。

Gaussian Blur

Sampling Neighbors

灰度化只看当前像素。模糊不一样，它必须读取周围邻居。最简单的模糊是 box blur，但视觉上更常见的是给中心更高权重的 3x3 Gaussian kernel：

为了避免一边读一边写把输入覆盖掉，模糊这一步不能继续原地改写同一个 buffer。我们保留灰度图在 pixels 里，再新建一块 blurred buffer 作为输出。

The Blur Kernel

在 Shaders.metal 里加入第三个 kernel：

Shaders.metal

kernel void gaussian_blur(const device uchar4* sourcePixels [[buffer(0)]],                          device uchar4* destinationPixels [[buffer(1)]],                          constant uint2& size [[buffer(2)]],                          uint2 id [[thread_position_in_grid]]){    if (id.x >= size.x || id.y >= size.y)    {        return;    }    const int weights[3][3] = {        {1, 2, 1},        {2, 4, 2},        {1, 2, 1},    };    float3 accum = float3(0.0);    float totalWeight = 0.0;    for (int offsetY = -1; offsetY <= 1; ++offsetY)    {        for (int offsetX = -1; offsetX <= 1; ++offsetX)        {            const uint sampleX = uint(clamp(int(id.x) + offsetX, 0, int(size.x) - 1));            const uint sampleY = uint(clamp(int(id.y) + offsetY, 0, int(size.y) - 1));            const uint sampleIndex = sampleY * size.x + sampleX;            const float weight = float(weights[offsetY + 1][offsetX + 1]);            const uchar4 sample = sourcePixels[sampleIndex];            accum += float3(sample.r, sample.g, sample.b) * weight;            totalWeight += weight;        }    }    const float3 color = accum / totalWeight;    const uint index = id.y * size.x + id.x;    destinationPixels[index] = uchar4(uchar(color.r), uchar(color.g), uchar(color.b), uchar(255));}

这里最重要的变化是从“当前像素自己决定颜色”变成“当前像素聚合邻居颜色”。边界像素不能越界读取，所以我们用 clamp 把采样位置限制在图片范围内。

Dispatching Into a Second Buffer

C++ 侧也要增加第二块输出 buffer，并给 blur kernel 绑定三个参数：

main.cpp绿色 = 新增行

MTL::Buffer* pixels = device->newBuffer(width * height * 4, MTL::ResourceStorageModeShared);MTL::Buffer* blurred = device->newBuffer(width * height * 4, MTL::ResourceStorageModeShared);MTL::Buffer* sizeBuffer = device->newBuffer(&size, sizeof(size), MTL::ResourceStorageModeShared);MTL::ComputePipelineState* makeBlur = makePipeline(device, library, "gaussian_blur");dispatch(queue, makeColor, pixels, sizeBuffer, width, height);writePPM("build/color.ppm", static_cast<const uint8_t*>(pixels->contents()), width, height);dispatch(queue, makeGray, pixels, sizeBuffer, width, height);writePPM("build/grayscale.ppm", static_cast<const uint8_t*>(pixels->contents()), width, height);dispatchBlur(queue, makeBlur, pixels, blurred, sizeBuffer, width, height);writePPM("build/blur.ppm", static_cast<const uint8_t*>(blurred->contents()), width, height);

dispatchBlur 和前面的 dispatch 很像，只是多绑定了一块输出 buffer。这个改动值得注意，因为它是从单输入单输出计算，过渡到“读一个资源、写另一个资源”的第一个例子。

运行完成后，你应该多得到一张 build/blur.ppm。它会把灰度图里相邻像素之间的对比压平，让边缘变软。

Blurred grayscale image generated by the Metal compute shader — `blur.ppm` 的视觉结果。和灰度图相比，边缘过渡更平滑，说明 kernel 已经开始读取邻域而不只是当前像素。

Reference Code

到这里，你已经在正文里写过完整程序的关键部分。最终参考实现放在 src/MetalCppInOneWeekend/，用于核对文件组织和细节，不需要在书页里重复阅读一整份源码。

project layout

src/MetalCppInOneWeekend/  main.cpp  Shaders.metalCMakeLists.txtthird_party/metal-cpp/

根目录 CMakeLists.txt 负责三件事：编译 main.cpp、链接 Metal 需要的 frameworks、把 Shaders.metal 编译成 default.metallib。

CMakeLists.txt

add_executable(MetalCppInOneWeekend    src/MetalCppInOneWeekend/main.cpp)target_include_directories(MetalCppInOneWeekend PRIVATE "${METAL_CPP_ROOT}")target_link_libraries(MetalCppInOneWeekend PRIVATE ${METAL_FRAMEWORKS})target_compile_definitions(MetalCppInOneWeekend PRIVATE    METALLIB_PATH="${BOOK1_METALLIB}")

最终代码包含这些文件：

src/MetalCppInOneWeekend/main.cpp：创建 Device、Command Queue、Buffer 和 Compute Pipeline，运行彩色、灰度和模糊三个 compute pass，然后写出 PPM 图片。
src/MetalCppInOneWeekend/Shaders.metal：提供 make_color_test、grayscale 和 gaussian_blur 三个 GPU kernel。
CMakeLists.txt：定义目标、framework 链接、shader 编译和运行时输出路径。

构建和运行：

terminal

cmake -S . -B buildcmake --build build./build/MetalCppInOneWeekend/MetalCppInOneWeekend

预期输出：

build/color.ppm：彩色测试图。
build/grayscale.ppm：灰度结果。
build/blur.ppm：3x3 Gaussian blur 结果。