如何在Metal多实例多LOD间接命令缓冲区中过滤实例？

现有一段运行正常的Metal计算着色器，用于为不同网格/子网格计算间接命令缓冲区（ICB），每个网格对应多个实例：

struct ICBContainer {
    command_buffer icb [[id(0)]];
};

struct ICBMesh {
    // Mesh data
    constant float *vertices;
    constant float *uv;
    constant float *tangents;
    constant float *bitangents;
    // Submesh data
    constant uint *indices;
    constant float *materials;
};

kernel void model_encode(
    uint meshId [[thread_position_in_grid]],
    device ICBContainer *icbContainer [[buffer(BUF_ICB)]],
    constant ICBMesh *meshes [[buffer(BUF_MESHES)]],
    constant MTLDrawIndexedPrimitivesIndirectArguments *draw [[buffer(BUF_DRAW)]],
    constant InstanceData *instances [[buffer(BUF_INSTANCES)]],
    constant Uniforms &uniforms [[buffer(BUF_UNIFORMS)]])
{
    ICBMesh mesh = meshes[meshId];
    MTLDrawIndexedPrimitivesIndirectArguments drawArgs = draw[meshId];
    
    render_command cmd(icbContainer->icb, meshId);
    cmd.set_vertex_buffer(&uniforms, BUF_UNIFORMS);
    cmd.set_vertex_buffer(mesh.vertices, BUF_VERTEX);
    cmd.set_vertex_buffer(mesh.uv, BUF_UV);
    cmd.set_vertex_buffer(mesh.tangents, BUF_TANGENT);
    cmd.set_vertex_buffer(mesh.bitangents, BUF_BITANGENT);
    cmd.set_fragment_buffer(mesh.materials, BUF_MATERIAL);
    cmd.set_vertex_buffer(instances, BUF_INSTANCES);
    cmd.draw_indexed_primitives(
        primitive_type::triangle,
        drawArgs.indexCount,
        mesh.indices + drawArgs.indexStart,
        drawArgs.instanceCount,
        drawArgs.baseVertex,
        drawArgs.baseInstance);
}

现在需要为模型添加LOD（0、1、2）版本，使用同一实例列表，让计算着色器根据实例到相机的距离计算LOD，过滤绘制调用，仅渲染与当前网格LOD匹配的实例，无需CPU端逐帧计算LOD并提供不同缓冲区。

已为ICBMesh结构体添加lod字段，网格缓冲区包含所有LOD版本的网格：

struct ICBMesh {
    uint lod;
    // Mesh data
    constant float *vertices;
    constant float *uv;
    constant float *tangents;
    constant float *bitangents;
    // Submesh data
    constant uint *indices;
    constant float *materials;
};

但当前会对所有LOD版本调用model_encode()，导致每个实例的所有LOD都被渲染。希望在着色器中计算实例LOD，仅绘制与当前网格LOD匹配的实例，猜测需要生成过滤后的实例数组传入渲染命令，但不知道如何在提前未知数组大小的情况下创建该数组，同时询问此方法是否合理。

在GPU端过滤实例并生成对应绘制命令的思路是合理的，但Metal计算着色器无法动态分配内存，不能直接创建未知大小的数组。可以通过两步计算+间接绘制的方式实现：

步骤1：预计算每个实例的目标LOD并分类 #

新增一个计算着色器，遍历所有实例，根据实例到相机的距离计算目标LOD，用原子计数器跟踪每个LOD的实例数量，避免线程竞争，将实例数据写入预分配的缓冲区：

// 定义每个LOD的实例计数和存储缓冲区
struct LODInstanceData {
    device uint *instanceCounts; // 长度为3，对应LOD0、1、2的实例数
    device InstanceData *filteredInstances; // 预分配足够大的缓冲区，存放所有过滤后的实例
    device uint *writeOffsets; // 每个LOD的写入偏移量，初始为0
};

kernel void compute_instance_lod(
    uint instanceId [[thread_position_in_grid]],
    constant InstanceData *allInstances [[buffer(BUF_ALL_INSTANCES)]],
    constant Uniforms &uniforms [[buffer(BUF_UNIFORMS)]],
    device LODInstanceData *lodData [[buffer(BUF_LOD_DATA)]])
{
    InstanceData instance = allInstances[instanceId];
    
    // 计算实例到相机的距离（假设instance包含世界空间位置）
    float3 cameraPos = uniforms.cameraPosition;
    float3 instancePos = instance.worldPosition;
    float distance = length(cameraPos - instancePos);
    
    // 根据距离确定目标LOD（示例逻辑，可按需调整）
    uint targetLOD = 0;
    if (distance > 100.0) targetLOD = 1;
    if (distance > 200.0) targetLOD = 2;
    
    // 原子操作获取当前LOD的写入偏移，并递增计数
    uint offset = atomic_fetch_add_explicit(&lodData->writeOffsets[targetLOD], 1, memory_order_relaxed);
    // 将实例数据写入对应位置
    lodData->filteredInstances[offset] = instance;
    // 更新实例计数
    atomic_fetch_add_explicit(&lodData->instanceCounts[targetLOD], 1, memory_order_relaxed);
}

步骤2：修改原model_encode kernel，按LOD过滤绘制 #

原model_encode kernel根据当前网格的LOD，读取对应过滤后的实例数据和计数，生成正确的绘制命令：

kernel void model_encode(
    uint meshId [[thread_position_in_grid]],
    device ICBContainer *icbContainer [[buffer(BUF_ICB)]],
    constant ICBMesh *meshes [[buffer(BUF_MESHES)]],
    constant MTLDrawIndexedPrimitivesIndirectArguments *draw [[buffer(BUF_DRAW)]],
    device LODInstanceData *lodData [[buffer(BUF_LOD_DATA)]],
    constant Uniforms &uniforms [[buffer(BUF_UNIFORMS)]])
{
    ICBMesh mesh = meshes[meshId];
    uint currentLOD = mesh.lod;
    
    // 获取当前LOD的实例数量，无实例则直接返回
    uint instanceCount = lodData->instanceCounts[currentLOD];
    if (instanceCount == 0) return;
    
    // 计算当前LOD实例在filteredInstances中的起始位置
    uint baseOffset = 0;
    for (uint i = 0; i < currentLOD; i++) {
        baseOffset += lodData->instanceCounts[i];
    }
    
    render_command cmd(icbContainer->icb, meshId);
    cmd.set_vertex_buffer(&uniforms, BUF_UNIFORMS);
    cmd.set_vertex_buffer(mesh.vertices, BUF_VERTEX);
    cmd.set_vertex_buffer(mesh.uv, BUF_UV);
    cmd.set_vertex_buffer(mesh.tangents, BUF_TANGENT);
    cmd.set_vertex_buffer(mesh.bitangents, BUF_BITANGENT);
    cmd.set_fragment_buffer(mesh.materials, BUF_MATERIAL);
    // 设置过滤后的实例缓冲区，指定起始偏移
    cmd.set_vertex_buffer(lodData->filteredInstances + baseOffset, BUF_INSTANCES);
    
    MTLDrawIndexedPrimitivesIndirectArguments drawArgs = draw[meshId];
    cmd.draw_indexed_primitives(
        primitive_type::triangle,
        drawArgs.indexCount,
        mesh.indices + drawArgs.indexStart,
        instanceCount,
        drawArgs.baseVertex,
        0); // baseInstance设为0，过滤后的实例是连续的
}

关键注意事项 #

• 缓冲区预分配：filteredInstances缓冲区需预分配足够大的空间，大小等于最大可能的实例总数乘以InstanceData的大小，避免溢出。
• 原子操作：必须用原子操作处理写入偏移和计数，防止多线程竞争导致数据错误。
• LOD计算逻辑：可根据实际需求调整距离阈值，也可加入屏幕空间大小等更精准的判断条件。
• 性能优化：若实例数量极大，可考虑分块处理，或使用Metal的threadgroup内存临时存储，减少全局内存的原子操作次数。

内容的提问来源于stack exchange，提问作者Chapak