GPU 基础

const std = @import("std");
const builtin = @import("builtin");

pub fn main() !void {
    // Query the compile-time target information to inspect the environment
    // 查询编译时目标信息以检查环境
    // this binary is being compiled for (host or cross-compilation target)
    // 此二进制文件正在为其编译（主机或交叉编译目标）
    const target = builtin.target;

    // Display basic target information: CPU architecture, OS, and object format
    // 显示基本目标信息：CPU架构、操作系统和对象格式
    std.debug.print("host architecture: {s}\n", .{@tagName(target.cpu.arch)});
    std.debug.print("host operating system: {s}\n", .{@tagName(target.os.tag)});
    std.debug.print("default object format: {s}\n", .{@tagName(target.ofmt)});

    // Check if we're compiling for a GPU backend by examining the target CPU architecture.
    // 通过检查目标 CPU 架构，检查我们是否正在为 GPU 后端编译。
    // GPU architectures include AMD GCN, NVIDIA PTX variants, and SPIR-V targets.
    // GPU 架构包括 AMD GCN、NVIDIA PTX 变体和 SPIR-V 目标。
    const is_gpu_backend = switch (target.cpu.arch) {
        .amdgcn, .nvptx, .nvptx64, .spirv32, .spirv64 => true,
        else => false,
    };
    std.debug.print("compiling as GPU backend: {}\n", .{is_gpu_backend});

    // Import address space types for querying GPU-specific memory capabilities
    // 导入地址空间类型以查询 GPU 特定的内存功能
    const AddressSpace = std.builtin.AddressSpace;
    const Context = AddressSpace.Context;

    // Query whether the target supports GPU-specific address spaces:
    // 查询目标是否支持 GPU 特定的地址空间：
    // - shared: memory shared within a workgroup/threadblock
    // - shared: 工作组/线程块内共享的内存
    // - constant: read-only memory optimized for uniform access across threads
    // - constant: 为跨线程统一访问优化的只读内存
    const shared_ok = target.cpu.supportsAddressSpace(AddressSpace.shared, null);
    const constant_ok = target.cpu.supportsAddressSpace(AddressSpace.constant, Context.constant);

    std.debug.print("supports shared address space: {}\n", .{shared_ok});
    std.debug.print("supports constant address space: {}\n", .{constant_ok});

    // Construct a custom target query for SPIR-V 64-bit targeting Vulkan
    // 构造一个针对 Vulkan 的 SPIR-V 64 位自定义目标查询
    const gpa = std.heap.page_allocator;
    const query = std.Target.Query{
        .cpu_arch = .spirv64,
        .os_tag = .vulkan,
        .abi = .none,
    };

    // Convert the target query to a triple string (e.g., "spirv64-vulkan")
    // 将目标查询转换为三元组字符串（例如，“spirv64-vulkan”）
    const triple = try query.zigTriple(gpa);
    defer gpa.free(triple);
    std.debug.print("example SPIR-V triple: {s}\n", .{triple});
}

// ! GPU内核：坐标捕获
//!
//! 此模块演示一个最小的SPIR-V计算内核，它将GPU调度坐标捕获到存储缓冲区中。
//! 它展示了如何使用Zig的GPU特定内置函数和地址空间注释来编写编译为SPIR-V的内核。

const builtin = @import("builtin");

/// 表示单个调用的GPU调度坐标
///
/// 使用`extern`布局来保证内存布局与主机端期望匹配，
/// 确保内核的输出可以被读取缓冲区的CPU代码安全解释。
const Coordinates = extern struct {
    // 工作组ID（此调用所属的组）
    group: u32,
    // 工作组大小（此维度中每组的调用数）
    group_size: u32,
    // 工作组内的本地调用ID（0到group_size-1）
    local: u32,
    // 所有调用的全局线性ID（group * group_size + local）
    linear: u32,
};

/// 捕获调度坐标的GPU内核入口点
///
/// 此函数必须被导出，以便SPIR-V编译器生成入口点。
/// `callconv(.kernel)`调用约定告诉Zig发出GPU特定函数属性，
/// 并根据计算着色器ABI处理参数传递。
///
/// 参数：
///   - out：指向将写入坐标的存储缓冲区的指针。
///          `.storage_buffer`地址空间注释确保适当的
///          设备可见GPU内存的内存访问模式。
pub export fn captureCoordinates(out: *addrspace(.storage_buffer) Coordinates) callconv(.kernel) void {
    // 查询X维度中的工作组ID（第一维度）
    // @workGroupId是GPU特定的内置函数，返回当前工作组的坐标
    const group = @workGroupId(0);

    // 查询工作组大小（此维度中每组的调用数）
    // 这在调度时由主机设置并在此处查询以保持完整性
    const group_size = @workGroupSize(0);

    // 查询此工作组内的本地调用ID（0到group_size-1）
    // @workItemId是每个工作组的线程索引
    const local = @workItemId(0);

    // 计算所有调用的全局线性索引
    // 此公式将2D坐标（group, local）转换为扁平1D索引
    const linear = group * group_size + local;

    // 将所有捕获的坐标写入输出缓冲区
    // GPU将确保此写入在同步后对主机可见
    out.* = .{
        .group = group,
        .group_size = group_size,
        .local = local,
        .linear = linear,
    };
}

// 编译时验证，确保此模块仅针对SPIR-V目标编译
// 这可以防止意外编译到GPU内置函数不可用的CPU架构
comptime {
    switch (builtin.target.cpu.arch) {
        // 接受32位和64位SPIR-V架构
        .spirv32, .spirv64 => {},
        // 使用有用的错误消息拒绝所有其他架构
        else => @compileError("captureCoordinates must be compiled with a SPIR-V target, e.g. -target spirv32-vulkan-none"),
    }
}

// ! GPU 调度规划工具
//!
// ! 该模块演示了如何计算 GPU 计算着色器的工作组调度参数。
// ! 它展示了总工作项、工作组大小和由此产生的调度
// ! 配置之间的关系，包括处理未填满完整工作组的“尾部”元素。

const std = @import("std");

// / 表示用于并行执行的完整调度配置
// / 包含启动计算内核或并行任务所需的所有参数
const DispatchPlan = struct {
    // / 每个工作组的大小（每组的线程/调用数）
    workgroup_size: u32,
    // / 覆盖所有项目所需的工作组数量
    group_count: u32,
    // / 包括填充在内的总调用次数（始终是 workgroup_size 的倍数）
    padded_invocations: u32,
    // / 最后一个工作组中填充/未使用的调用次数
    tail: u32,
};

// / 为给定的问题大小和工作组配置计算最佳调度参数
///
// / 计算处理所有项目所需的工作组数量，考虑到
// / 最后一个工作组可能部分填充的事实。这对于 GPU 计算着色器至关重要，
// / 其中工作必须以工作组大小的倍数进行调度。
fn computeDispatch(total_items: u32, workgroup_size: u32) DispatchPlan {
    // 确保工作组大小有效（GPU 工作组不能为空）
    std.debug.assert(workgroup_size > 0);

    // 计算所需工作组的数量，向上取整以确保所有项目都被覆盖
    const groups = std.math.divCeil(u32, total_items, workgroup_size) catch unreachable;

    // 计算包括填充在内的总调用次数（GPU 总是启动完整的工作组）
    const padded = groups * workgroup_size;

    return .{
        .workgroup_size = workgroup_size,
        .group_count = groups,
        .padded_invocations = padded,
        // 尾部表示必须通过边界检查处理的浪费的调用
        .tail = padded - total_items,
    };
}

// / 使用相同的调度逻辑模拟 CPU 侧并行执行规划
///
// / 演示了工作组调度公式同样适用于 CPU 线程批处理，
// / 确保 GPU 和 CPU 回退实现之间的一致行为。
fn simulateCpuFallback(total_items: u32, lanes: u32) DispatchPlan {
    // 重用 GPU 公式，使主机侧分块与设备调度匹配。
    return computeDispatch(total_items, lanes);
}

pub fn main() !void {
    // 定义一个示例问题：处理 1000 个项目
    const problem_size: u32 = 1000;

    // 典型的 GPU 工作组大小（通常为 32、64 或 256，取决于硬件）
    const workgroup_size: u32 = 64;

    // 计算 GPU 调度配置
    const plan = computeDispatch(problem_size, workgroup_size);
    std.debug.print(
        "gpu dispatch: {d} groups × {d} lanes => {d} invocations (tail {d})\n",
        .{ plan.group_count, plan.workgroup_size, plan.padded_invocations, plan.tail },
    );

    // 模拟 CPU 回退，并行通道更少
    const fallback_threads: u32 = 16;
    const cpu = simulateCpuFallback(problem_size, fallback_threads);
    std.debug.print(
        "cpu chunks: {d} batches × {d} lanes => {d} logical tasks (tail {d})\n",
        .{ cpu.group_count, cpu.workgroup_size, cpu.padded_invocations, cpu.tail },
    );
}