Native APIs provide ways to emit draws and dispatches that take arguments from buffers on the GPU which is useful when some of the parameters are computed on the GPU.

Native APIs

Metal

In Metal indirect commands are done by passing an indirect buffer and an offset to commands instead of the parameters:

@protocol MTLRenderCommandEncoder <MTLCommandEncoder>
    - (void) drawPrimitives: (MTLPrimitiveType) primitiveType
             indirectBuffer: (id <MTLBuffer>) indirectBuffer
       indirectBufferOffset: (NSUInteger) indirectBufferOffset;

    - (void) drawIndexedPrimitives: (MTLPrimitiveType) primitiveType
                         indexType: (MTLIndexType) indexType
                       indexBuffer: (id <MTLBuffer>) indexBuffer
                 indexBufferOffset: (NSUInteger) indexBufferOffset
                    indirectBuffer: (id <MTLBuffer>) indirectBuffer
              indirectBufferOffset: (NSUInteger) indirectBufferOffset;
@end

@protocol MTLComputeCommandEncoder <MTLCommandEncoder>
    - (void) dispatchThreadgroupsWithIndirectBuffer: (id <MTLBuffer>) indirectBuffer
                               indirectBufferOffset: (NSUInteger) indirectBufferOffset
                              threadsPerThreadgroup: (MTLSize) threadsPerThreadgroup;
@end

The format of the indirect buffers is defined by:

typedef struct {
    uint32_t vertexCount;
    uint32_t instanceCount;
    uint32_t vertexStart;
    uint32_t baseInstance;
} MTLDrawPrimitivesIndirectArguments;

typedef struct {
    uint32_t indexCount;
    uint32_t instanceCount;
    uint32_t indexStart;
    int32_t  baseVertex;
    uint32_t baseInstance;
} MTLDrawIndexedPrimitivesIndirectArguments;

typedef struct {
    uint32_t threadgroupsPerGrid[3];
} MTLDispatchThreadgroupsIndirectArguments;

Vulkan

Vulkan is very similar to Metal except that:

• the primitive type is taken from the last VkCmdBindPipeline
• the index buffer / type and offset are taken from the last VkCmdBindIndexBuffer
• the commands allow sending multiple draws at a fixed stride in the buffer; the maximum number for drawCount is a device limit that has to be at least 1.
• baseInstance must be 0 unless the drawIndirectFirstInstance is enabled.

void vkCmdDrawIndirect(VkCommandBuffer commandBuffer, VkBuffer buffer, VkDeviceSize offset, uint32_t drawCount, uint32_t stride);
void vkCmdDrawIndexedIndirect(VkCommandBuffer commandBuffer, VkBuffer buffer, VkDeviceSize offset, uint32_t drawCount, uint32_t stride);
void vkCmdDispatchIndirect(VkCommandBuffer commandBuffer, VkBuffer buffer, VkDeviceSize offset);

The format of the indirect buffers is defined by:

typedef struct VkDrawIndirectCommand {
    uint32_t    vertexCount;
    uint32_t    instanceCount;
    uint32_t    firstVertex;
    uint32_t    firstInstance;
} VkDrawIndirectCommand;

typedef struct VkDrawIndexedIndirectCommand {
    uint32_t    indexCount;
    uint32_t    instanceCount;
    uint32_t    firstIndex;
    int32_t     vertexOffset;
    uint32_t    firstInstance;
} VkDrawIndexedIndirectCommand;

typedef struct VkDispatchIndirectCommand {
    uint32_t    x;
    uint32_t    y;
    uint32_t    z;
} VkDispatchIndirectCommand;

D3D12

In D3D12 executing commands indirectly is done through "command signature" that are a super-set of what's available in Metal and Vulkan. First a command signature is created that represents the order and layout of commands in the command buffer, then this signature is passed along with the indirect buffer to ID3D12GraphicsCommandList::ExecuteIndirect.

The interesting bits are:

• That the layout of the D3D12_DRAW_ARGUMENTS, D3D12_DRAW_INDEXED_ARGUMENTS and D3D12_DISPATCH_ARGUMENTS structures match the respective structures in Metal and Vulkan.
• A command signature has to be created against a root signature (because there are indirect commands to change root constants and descriptors).
• Command signatures can take the maximum number of commands to process from a "count buffer" which makes command signatures a super-set of "MultiDrawIndirect".

typedef struct D3D12_DRAW_ARGUMENTS {
  UINT VertexCountPerInstance;
  UINT InstanceCount;
  UINT StartVertexLocation;
  UINT StartInstanceLocation;
} D3D12_DRAW_ARGUMENTS;

typedef struct D3D12_DRAW_INDEXED_ARGUMENTS {
  UINT IndexCountPerInstance;
  UINT InstanceCount;
  UINT StartIndexLocation;
  INT  BaseVertexLocation;
  UINT StartInstanceLocation;
} D3D12_DRAW_INDEXED_ARGUMENTS;

typedef struct D3D12_DISPATCH_ARGUMENTS {
  UINT ThreadGroupCountX;
  UINT ThreadGroupCountY;
  UINT ThreadGroupCountZ;
} D3D12_DISPATCH_ARGUMENTS;

Links to relevant parts of the D3D12 documentation:

• Using command signatures: ID3D12GraphicsCommandList::ExecuteIndirect
• Creating command signatures: D3D12_COMMAND_SIGNATURE_DESC and ID3D12Device::CreateCommandSignature
• Specifying a single command in the signature: D3D12_INDIRECT_ARGUMENT_DESC and D3D12_INDIRECT_ARGUMENT_TYPE

Proposed API

All three APIs use the same layout for commands in the indirect buffer, which makes things easier. Because of restrictions of Metal and Vulkan, sending only a single DrawIndirect, DrawIndexedIndirect or DispatchIndirect command at a time should be supported.

The arguments for each of these should be an indirect buffer and offset, and that's it because the rest of the data will be present either in the last bound pipeline state, or in the last bound index buffer. Obviously graphics commands will have to be done in a render pass, and likewise the compute command should be done in a compute pass.

void CommandBuffer::DrawIndirect(Buffer* indirectBuffer, uint32_t indirectOffset);
void CommandBuffer::DrawIndexedIndirect(Buffer* indirectBuffer, uint32_t indirectOffset);
void CommandBuffer::DispatchIndirect(Buffer* indirectBuffer, uint32_t indirectOffset);

// Indirect buffer layout same as D3D12, Metal and Vulkan

Open question: Some Adreno 4XX and 5XX GPUs don't support drawIndirectFirstInstance. Should we require it anyways?