How to access Tier 1 Argument Buffer struct without indexing

Question

According to this example from Apple, Tier 1 Argument Buffers cannot be accessed through pointer indexing (https://developer.apple.com/documentation/metal/buffers/about_argument_buffers). If this is not allowed, how can I index into a particular struct in my argument buffer array?

// Shader.metal
struct MyTexture {
    metal::texture2d<float, metal::access::sample> texture;
};

fragment half4 myFragment(VertexOut vert [[stage_in]],
                          ....,
                          constant int &count [[buffer(4)]],
                          constant MyTexture *textures [[buffer(5)]],
                          ....)
{
   for(int i = 0; i < count; i++) {
       MyTexture resource = textures[i];
       
       float4 color = resource.texture.sample(sampler, pos.xy);
       outputColor = mix(inputColor, color, 0.5); // <-- Causes error
   }
}

The error that I get is from creating the MTLRenderPipelineState with this error message:

Inlining all functions due to use of indirect argument bufferbuffer(5): Argument buffer accessed with non-zero array index.

Answer 1

A workaround for Tier 1 would be to pass a pointer to the instance inside the Argument Buffer, rather than the entire buffer.

Example, look at the use of Material.

// Argument-buffered resource
struct Material {
  metal::sampler sampler [[id(AB_MaterialSampler)]];
  metal::texture2d<float> base_color_texture [[id(AB_MaterialBaseColorTexture)]];
  metal::texture2d<float> normal_map [[id(AB_MaterialNormalMap)]];
  metal::texture2d<float> ao_metallic_roughness_map [[id(AB_MaterialAoMetallicRoughnessMap)]];
  float3 base_color_factor [[id(AB_MaterialBaseColorFactor)]];
  float metallic_factor [[id(AB_MaterialMetallicFactor)]];
  float roughness_factor [[id(AB_MaterialRoughnessFactor)]];
};

// GPU-driven rendering kernel
kernel void icb_frame_kernel(device IcbContainer& icb_container [[buffer(KB_IcbContainer)]],
                             constant VertexUniforms* vertex_uniforms [[buffer(KB_VertexUniforms)]],
                             constant FragmentUniforms* fragment_uniforms [[buffer(KB_FragmentUniforms)]],
                             device Mesh* meshes [[buffer(KB_Meshes)]],
                             constant Transform* transforms [[buffer(KB_Transforms)]],
                             device Material* materials [[buffer(KB_Materials)]],
                             constant ShadowMap* shadow_map [[buffer(KB_ShadowMap)]],
                             constant Ibl* ibl [[buffer(KB_Ibl)]],
                             constant Cubemap* cubemap [[buffer(KB_Cubemap)]],
                             device MTLIndirectCommandBufferExecutionRange& range [[buffer(KB_ExecutionRange)]],
                             const uint instance_id [[thread_position_in_grid]]) {
  device auto& mesh = meshes[instance_id];
  device auto* range_length = reinterpret_cast<device atomic_uint*>(&range.length);
  const auto index = atomic_fetch_add_explicit(range_length, 1, memory_order_relaxed);
  
  render_command cmd(icb_container.icb, index);
  cmd.set_render_pipeline_state(mesh.pipeline_state);
  cmd.set_vertex_buffer(mesh.vertex_buffer, VB_Vertices);
  cmd.set_vertex_buffer(vertex_uniforms, VB_VertexUniforms);
  cmd.set_vertex_buffer(transforms, VB_Transforms);
  cmd.set_fragment_buffer(fragment_uniforms, FB_FragmentUniforms);
  cmd.set_fragment_buffer(transforms, FB_Transforms);
  // Tier 1: use indexed access and pass pointer to instance
  cmd.set_fragment_buffer(&materials[instance_id], FB_Material);
  // Tier 2: pass entire buffer and use indexed access in fragment shader
  cmd.set_fragment_buffer(materials, FB_Material);
  cmd.set_fragment_buffer(shadow_map, FB_ShadowMap);
  cmd.set_fragment_buffer(ibl, FB_Ibl);
  cmd.set_fragment_buffer(cubemap, FB_Cubemap);
  
  if (mesh.is_uint16_index){
    constant auto* index_buffer = static_cast<constant ushort*>(mesh.index_buffer);
    cmd.draw_indexed_primitives(primitive_type::triangle, mesh.index_count, index_buffer, 1, 0, instance_id);
  } else {
    constant auto* index_buffer = static_cast<constant uint*>(mesh.index_buffer);
    cmd.draw_indexed_primitives(primitive_type::triangle, mesh.index_count, index_buffer, 1, 0, instance_id);
  }
}

// Tier 1
fragment half4 pbr_fragment(ProjectedVertex vert [[stage_in]],
                            constant FragmentUniforms& uniforms [[buffer(FB_FragmentUniforms)]],
                            constant Material& material [[buffer(FB_Material)]],
                            constant Ibl& ibl [[buffer(FB_Ibl), function_constant(HAS_IBL)]],
                            constant ShadowMap& shadow_map [[buffer(FB_ShadowMap), function_constant(HAS_SHADOW_MAP)]]
                            ) {
  // Use Material
}

// Tier 2
fragment half4 pbr_fragment(ProjectedVertex vert [[stage_in]],
                            constant FragmentUniforms& uniforms [[buffer(FB_FragmentUniforms)]],
                            device Material* materials [[buffer(FB_Material)]],
                            constant Ibl& ibl [[buffer(FB_Ibl), function_constant(HAS_IBL)]],
                            constant ShadowMap& shadow_map [[buffer(FB_ShadowMap), function_constant(HAS_SHADOW_MAP)]]
                            ) {
  // Use indexed Material
  const auto& material = materials[vert.instance_id];
}

I did not have time to edit the example for brevity, but it should be clear enough.

Side note: the Metal spec recommends to use device address space whenever you use pointer arithmetic (indexed access). See page 61 of the spec.

Answer 2

Short answer: you can't.

The reason you can't is because tier 1 hardware can only emulate argument buffers using regular bind points. With tier2 you can bind any number of textures there, so the driver can't know at bind time how many slots it will need to use, and the hardware itself can't do a dependent read for other GPU objects, like textures and samplers.