# 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. ## Discussion 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`](/documentation/Metal/MTLCommandBuffer) and using it to create a new compute command encoder, using a method like [`makeComputeCommandEncoder()`](/documentation/Metal/MTLCommandBuffer/makeComputeCommandEncoder()) or [`makeComputeCommandEncoder(descriptor:)`](/documentation/Metal/MTLCommandBuffer/makeComputeCommandEncoder(descriptor:)). 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()`](/documentation/Metal/MTLCommandEncoder/endEncoding()). > Note: > Everything used to set up your compute pass is CPU thread-safe, except for ``doc://com.apple.metal/documentation/Metal/MTLComputeCommandEncoder``. Synchronize ``doc://com.apple.metal/documentation/Metal/MTLResource`` instances you share between the CPU and GPU with an ``doc://com.apple.metal/documentation/Metal/MTLFence``, an ``doc://com.apple.metal/documentation/Metal/MTLEvent``, or a completion callback. For information on dispatching commands to encode, see the [`MTLComputeCommandEncoder`](/documentation/Metal/MTLComputeCommandEncoder) 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](/documentation/Metal/performing-calculations-on-a-gpu) 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](/documentation/Metal/combining-blit-and-compute-operations-in-a-single-pass) 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):](https://developer.apple.com/metal/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`](/documentation/Metal/MTLAccelerationStructure) 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](/documentation/Metal/customizing-shaders-using-function-pointers-and-stitching) - [Creating a Metal dynamic library](/documentation/Metal/creating-a-metal-dynamic-library) - [Ray tracing with acceleration structures](/documentation/Metal/ray-tracing-with-acceleration-structures) For information on per-architecture support for function tables in compute passes and other restrictions, see the [Metal feature set tables (PDF)](https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf). ### Argument buffers and memory residency Compute kernels can access argument data to populate a Metal structure, using an [`MTLBuffer`](/documentation/Metal/MTLBuffer) created by an [`MTLArgumentEncoder`](/documentation/Metal/MTLArgumentEncoder). For an in-depth discussion of argument buffers, see [Improving CPU performance by using argument buffers](/documentation/Metal/improving-cpu-performance-by-using-argument-buffers). 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`](/documentation/Metal/MTLStorageMode/shared) or [`MTLStorageMode.managed`](/documentation/Metal/MTLStorageMode/managed) mode. Resources become resident on a per-instance basis by calling methods like [`useResource(_:usage:)`](/documentation/Metal/MTLComputeCommandEncoder/useResource(_:usage:)) and heaps become resident by calling methods like [`useHeap(_:)`](/documentation/Metal/MTLComputeCommandEncoder/useHeap(_:)). > 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`](/documentation/Metal/MTLHazardTrackingMode) 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)](https://developer.apple.com/metal/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)](https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf). ## Topics ### Essentials [Performing calculations on a GPU](/documentation/Metal/performing-calculations-on-a-gpu) Use Metal to find GPUs and perform calculations on them. [Combining blit and compute operations in a single pass](/documentation/Metal/combining-blit-and-compute-operations-in-a-single-pass) Run concurrent blit commands and then a compute dispatch in a single pass with a unified compute encoder. ### Encoding a compute pass Encode commands to prepare data and run a kernel in a compute pass. [Creating threads and threadgroups](/documentation/Metal/creating-threads-and-threadgroups) Learn how Metal organizes compute-processing workloads. [Calculating threadgroup and grid sizes](/documentation/Metal/calculating-threadgroup-and-grid-sizes) Calculate the optimum sizes for threadgroups and grids when dispatching compute-processing workloads. [`MTL4ComputeCommandEncoder`](/documentation/Metal/MTL4ComputeCommandEncoder) Encodes computation dispatches, resource copying commands, and acceleration structure building commands for a single pass into a command buffer. [`MTLComputeCommandEncoder`](/documentation/Metal/MTLComputeCommandEncoder) Encodes computation dispatch commands for a single compute pass into a command buffer. ### Configuring a compute pipeline state Define the GPU state for a kernel function call in your compute pass. [`MTL4ComputePipelineDescriptor`](/documentation/Metal/MTL4ComputePipelineDescriptor) Describes a compute pipeline state. [`MTLComputePipelineDescriptor`](/documentation/Metal/MTLComputePipelineDescriptor) An instance describing the desired GPU state for a kernel call in a compute pass. [`MTLComputePipelineState`](/documentation/Metal/MTLComputePipelineState) An interface that represents a GPU pipeline configuration for running kernels in a compute pass. [`MTLStageInputOutputDescriptor`](/documentation/Metal/MTLStageInputOutputDescriptor) A description of the input and output data of a function. [`MTLPipelineBufferDescriptor`](/documentation/Metal/MTLPipelineBufferDescriptor) The mutability options for a buffer that a render or compute pipeline uses. [`MTLPipelineBufferDescriptorArray`](/documentation/Metal/MTLPipelineBufferDescriptorArray) An array of pipeline buffer descriptors. [`MTLPipelineOption`](/documentation/Metal/MTLPipelineOption) Options that determine how Metal prepares the pipeline. ### Configuring a compute pass Define how kernel functions get called throughout your compute pass, and any indirect command buffers to access. [`MTLComputePassDescriptor`](/documentation/Metal/MTLComputePassDescriptor) A description of how to dispatch execution of pass commands and GPU performance sampling. [`MTLDispatchType`](/documentation/Metal/MTLDispatchType) The type of dispatch method to use when calling encoded functions. [`MTLDispatchThreadgroupsIndirectArguments`](/documentation/Metal/MTLDispatchThreadgroupsIndirectArguments) The data layout required for arguments needed to specify the size of threadgroups. [`MTLComputePassSampleBufferAttachmentDescriptor`](/documentation/Metal/MTLComputePassSampleBufferAttachmentDescriptor) A configuration that instructs the GPU where to store counter data from the beginning and end of a compute pass. [`MTLComputePassSampleBufferAttachmentDescriptorArray`](/documentation/Metal/MTLComputePassSampleBufferAttachmentDescriptorArray) A container that stores an array of sample buffer attachments for a compute pass. --- Copyright © 2026 Apple Inc. All rights reserved. | [Terms of Use](https://www.apple.com/legal/internet-services/terms/site.html) | [Privacy Policy](https://www.apple.com/privacy/privacy-policy)