--- title: Compute passes framework: Metal role: collectionGroup role_heading: API Collection platforms: [] path: metal/compute_passes --- # Compute passes Encode a compute pass that runs computations in parallel on a thread grid, processing and manipulating Metal resource data on multiple cores of a GPU. ## Overview Your app can perform large-scale computation or prepare data for a subsequent GPU pass by encoding a compute pass that works on Metal resources in parallel. Compute passes are the part of your Metal pipeline meant for heavy parallelization of tasks requiring fast math, such as ray tracing. Encode commands for your compute pass by creating an [MTLCommandBuffer](mtlcommandbuffer.md) and using it to create a new compute command encoder, using a method like [makeComputeCommandEncoder()](mtlcommandbuffer/makecomputecommandencoder().md) or [makeComputeCommandEncoder(descriptor:)](mtlcommandbuffer/makecomputecommandencoder(descriptor:).md). Add individual dispatches for functions and their data to the compute pass by calling the command encoder’s methods. At the end of assigning data and dispatching a function call to the encoder, create a command that runs in your compute pass with [endEncoding()](mtlcommandencoder/endencoding().md). > **note:** Everything used to set up your compute pass is CPU thread-safe, except for [MTLComputeCommandEncoder](mtlcomputecommandencoder.md). Synchronize [MTLResource](mtlresource.md) instances you share between the CPU and GPU with an [MTLFence](mtlfence.md), an [MTLEvent](mtlevent.md), or a completion callback. For information on dispatching commands to encode, see the [MTLComputeCommandEncoder](mtlcomputecommandencoder.md) reference. Compute passes also support indirect command buffers; for more information, see Dispatching from Indirect Command Buffers. The following two samples demonstrate basic compute passes: - See [Performing calculations on a GPU](performing-calculations-on-a-gpu.md) for an example of configuring and running a compute pass that performs basic parallel math. - See [Combining blit and compute operations in a single pass](combining-blit-and-compute-operations-in-a-single-pass.md) for an example of using a compute pass to modify data for a render pass. ### Kernel arguments and argument tables Compute commands that execute your code on GPU call *kernel functions* in your Metal shader, annotated with `[[kernel]]`. Each kernel has associated argument tables, such as the buffer argument table `[[buffer(n)]]`, used to access data associated with kernel arguments. In addition to annotations describing any argument table, some kernel arguments need information on their address space. For more information, see the following sections of the `Metal Shading Language Specification (PDF):` - For compute kernels, Section 5.1.3 - For function argument tables, Section 5.2 - For address spaces, Section 4 In addition, kernels also use a function table to take advantage of function pointers, allowing them to call visible and intersection functions. Visible functions allow you to use function pointers in kernels, letting you use function stitching and link against Metal dynamic libraries at runtime. Ray tracers use intersection functions on [MTLAccelerationStructure](mtlaccelerationstructure.md) instances to perform quick intersection checks. For more information on function stitching, dynamic libraries, and ray tracing, see: - [Customizing shaders using function pointers and stitching](customizing-shaders-using-function-pointers-and-stitching.md) - [Creating a Metal dynamic library](creating-a-metal-dynamic-library.md) - [Ray tracing with acceleration structures](ray-tracing-with-acceleration-structures.md) For information on per-architecture support for function tables in compute passes and other restrictions, see the `Metal feature set tables (PDF)`. ### Argument buffers and memory residency Compute kernels can access argument data to populate a Metal structure, using an [MTLBuffer](mtlbuffer.md) created by an [MTLArgumentEncoder](mtlargumentencoder.md). For an in-depth discussion of argument buffers, see [Improving CPU performance by using argument buffers](improving-cpu-performance-by-using-argument-buffers.md). Using a resource in an argument buffer requires that it’s resident in GPU memory for the duration of the pass. For a resource to be resident, allocate it with either the [MTLStorageMode.shared](mtlstoragemode/shared.md) or [MTLStorageMode.managed](mtlstoragemode/managed.md) mode. Resources become resident on a per-instance basis by calling methods like [useResource(_:usage:)](mtlcomputecommandencoder/useresource(_:usage:).md) and heaps become resident by calling methods like [useHeap(_:)](mtlcomputecommandencoder/useheap(_:).md). > **important:** For the duration of your compute pass, don’t access any resident resources on the CPU. Doing so in your app can cause GPU memory corruption, such as visual artifacts. When using resident resources, avoid data corruption by using an appropriate [MTLHazardTrackingMode](mtlhazardtrackingmode.md) or by manually managing memory barriers and fences for untracked resources with the methods in Synchronizing Across Command Execution. ### Using tile memory in a compute pass Apple family GPUs offer fast, integrated graphics memory called *tile memory* that’s shared between subsequent passes for fast access to data. Compute passes can reserve this memory space for threadgroup memory or imageblock memory, giving your compute functions the ability to access temporary data at low latency across your shaders. For more information see the following sections of the `Metal Shading Language Specification (PDF):`: - Section 4.4 for information on the `threadgroup` memory space - Section 4.5 for information on the `threadgroup_imageblock` memory space - Section 2.11 for information on imageblocks - Section 5.6 for information on `imageblock` attributes Because tile memory resides on GPU only, you reserve memory on a tile block rather than copy data to it. Use the methods in Encoding Tile Memory Usage to prepare the appropriate block of memory for your kernel. For device support and other tile memory limitations, see `Metal feature set tables (PDF)`. ## Topics ### Essentials - [Performing calculations on a GPU](performing-calculations-on-a-gpu.md) - [Combining blit and compute operations in a single pass](combining-blit-and-compute-operations-in-a-single-pass.md) ### Encoding a compute pass - [Creating threads and threadgroups](creating-threads-and-threadgroups.md) - [Calculating threadgroup and grid sizes](calculating-threadgroup-and-grid-sizes.md) - [MTL4ComputeCommandEncoder](mtl4computecommandencoder.md) - [MTLComputeCommandEncoder](mtlcomputecommandencoder.md) ### Configuring a compute pipeline state - [MTL4ComputePipelineDescriptor](mtl4computepipelinedescriptor.md) - [MTLComputePipelineDescriptor](mtlcomputepipelinedescriptor.md) - [MTLComputePipelineState](mtlcomputepipelinestate.md) - [MTLStageInputOutputDescriptor](mtlstageinputoutputdescriptor.md) - [MTLPipelineBufferDescriptor](mtlpipelinebufferdescriptor.md) - [MTLPipelineBufferDescriptorArray](mtlpipelinebufferdescriptorarray.md) - [MTLPipelineOption](mtlpipelineoption.md) ### Configuring a compute pass - [MTLComputePassDescriptor](mtlcomputepassdescriptor.md) - [MTLDispatchType](mtldispatchtype.md) - [MTLDispatchThreadgroupsIndirectArguments](mtldispatchthreadgroupsindirectarguments.md) - [MTLComputePassSampleBufferAttachmentDescriptor](mtlcomputepasssamplebufferattachmentdescriptor.md) - [MTLComputePassSampleBufferAttachmentDescriptorArray](mtlcomputepasssamplebufferattachmentdescriptorarray.md) ## See Also ### Command encoders - [Render passes](render-passes.md) - [Machine learning passes](machine-learning-passes.md) - [Blit passes](blit-passes.md) - [Indirect command encoding](indirect-command-encoding.md) - [Ray tracing with acceleration structures](ray-tracing-with-acceleration-structures.md)