Go to Top

Swift and Metal GPU Programming on tvOS for the new Apple TV

Apple just announced tvOS for the new Apple TV. The new Apple TV has an A8 CPU/GPU and supports Metal for GPU programming!

This posting gives a basic example how to use Swift with Metal on TVOS on Apple.  The code is available on github.com/memkite/MetalForTVOS.

But first, what is Metal?

Apple describes Metal originally as: “Metal provides the lowest-overhead access to the GPU, enabling you to maximize the graphics and compute potential of your iOS 9, OSX 10.11 and now tvOS app. With a streamlined API, precompiled shaders, and support for efficient multi-threading, Metal can take your game or graphics app to the next level of performance and capability.” – source: https://developer.apple.com/metal/

I consider Metal as an alternative to OpenGL for graphics processing, but also as an alternative to OpenCL and Nvidia’s Cuda
on data-parallel programming for GPUs (also called General Purpose GPU programming – GPGPU).

1. Swift and Metal GPU Programming on tvOS 9.0

This example shows a very simple Metal shader function – doubler – that doubles each element in the input vector

1.1 Summary of steps:

  1. Create a new OSX Cocoa Swift project in Xcode (requires Xcode 7.1 beta or newer)
  2. Add dependency libraries (added Metal and Quartzcore libraries in this case)
  3. Create a Metal Shader method – doubler – in Shaders.metal
  4. Create a (convenience) MetalViewController class in MetalViewController.swift that provides the methods:
    1. setupMetal() – creates metalDevice, metalCommandQueue, metalDefaultLibrary and metalCommandBuffer objects
    2. setupShaderInMetalPipeline(shaderName:String) – returns shader function, computePipelineState and computePipelineErrors objects
    3. createMetalBuffer(var vector:[Float]) – creates a Metal byte buffer from a Swift float array
  5. Create a ViewController  in GameViewController.swift that:
    1. creates an input vector and (empty) output vector of same length
    2. creates  Metal Byte buffers for those vectors
    3. creates a MetalComputeCommandEncoder – command encoder (primary object for using Metal)
    4. adds metal input and output buffers to the command encoder
    5. connects the shader function to the command encoder (indirect via computePipelineState object)
    6. configures number of threads and thread groups to use on the GPU for the command encoder
    7. finalizes configuration of command encoder (i.e. ends encoding)
    8. starts metal job by commit() and waits until it is completed
    9. gets output data from the GPU into an NSData object
    10. converts NSData object into a regular Swift Array
    11. prints output array


Have shown a simple example of using Metal for GPU-programming on tvOS, that is as mentioned currently untested, but compiles. The tvOS app code in Swift resembles iOS very much (as expected), and the Metal part resembles both iOS and OSX. My primary interest in Metal is Deep Learning, e.g. for Image recognition and Text Classification, and it will be fun to see what can be done in that direction on Apple TV. (Disclaimer: this code just compiles and resembles Metal code for iOS and OSX but is not yet tested since I don’t have a new Apple TV yet – if you want to send one => address).

Best regards,

Amund Tveit


Appendix – Code

Appendix 1 – Metal Shader – Shaders.metal

using namespace metal;

kernel void doubler(const device float *inVector [[ buffer(0) ]],
                    device float *outVector [[ buffer(1) ]],
                    uint id [[ thread_position_in_grid ]]) {
    outVector[id] = 2*inVector[id];

Appendix 2 – Game View Controller – GameViewController.swift

import Metal

class GameViewController:MetalViewController {//, MTKViewDelegate {
    override func viewDidLoad() {
        // Create input and output vectors, and corresponding metal buffer
        let N = 100
        let inputVector = createInputVector(N)
        let outputVector = doublerWithMetal(inputVector)
        print("input = \(inputVector)")
        print("output = \(outputVector)")
    func doublerWithMetal(inputVector:[Float]) -> [Float] {
        // uses metal to calculate double array
        let (_, computePipelineState, _) = setupShaderInMetalPipeline("doubler")
        let inputMetalBuffer = createMetalBuffer(inputVector)
        var outputVector = [Float](count: inputVector.count, repeatedValue: 0.0)
        let outputMetalBuffer = createMetalBuffer(outputVector)
        // Create Metal Compute Command Encoder and add input and output buffers to it
        metalComputeCommandEncoder = metalCommandBuffer.computeCommandEncoder()
        metalComputeCommandEncoder.setBuffer(inputMetalBuffer, offset: 0, atIndex: 0)
        metalComputeCommandEncoder.setBuffer(outputMetalBuffer, offset: 0, atIndex: 1)
        // Set the shader function that Metal will use
        // Find max number of parallel GPU threads (threadExecutionWidth) in computePipelineState
        let threadExecutionWidth = computePipelineState.threadExecutionWidth
        // Set up thread groups on GPU
        let threadsPerGroup = MTLSize(width:threadExecutionWidth,height:1,depth:1)
        let numThreadgroups = MTLSize(width:(inputVector.count+threadExecutionWidth)/threadExecutionWidth, height:1, depth:1)
        metalComputeCommandEncoder.dispatchThreadgroups(numThreadgroups, threadsPerThreadgroup: threadsPerGroup)
        // Finalize configuration
        print("outputVector before job is running: \(outputVector)")
        // Start job
        // Wait for it to finish
        // Get output data from Metal/GPU into Swift
        let data = NSData(bytesNoCopy: outputMetalBuffer.contents(),
            length: outputVector.count*sizeof(Float), freeWhenDone: false)
        data.getBytes(&outputVector, length:inputVector.count * sizeof(Float))
        return outputVector
    func createInputVector(N: Int) -> [Float] {
        var vector = [Float](count: N, repeatedValue: 0.0)
        for (index, _) in vector.enumerate() {
            vector[index] = Float(index)
        return vector

Appendix 3 – MetalViewController – MetalViewController.swift

import Foundation
import UIKit
import Metal
import QuartzCore

class MetalViewController : UIViewController {

    override func viewDidLoad() {
    var metalDevice:MTLDevice!
    var metalCommandQueue:MTLCommandQueue!
    var metalDefaultLibrary:MTLLibrary!
    var metalCommandBuffer:MTLCommandBuffer!
    var metalComputeCommandEncoder:MTLComputeCommandEncoder!
    func setupMetal() {
        // Get access to iPhone or iPad GPU
        metalDevice = MTLCreateSystemDefaultDevice()
        // Queue to handle an ordered list of command buffers
        metalCommandQueue = metalDevice.newCommandQueue()
        // Access to Metal functions that are stored in Shaders.metal file, e.g. sigmoid()
        metalDefaultLibrary = metalDevice.newDefaultLibrary()
        // Buffer for storing encoded commands that are sent to GPU
        metalCommandBuffer = metalCommandQueue.commandBuffer()
    func setupShaderInMetalPipeline(shaderName:String) -> (shader:MTLFunction!,
        computePipelineErrors:NSErrorPointer!)  {
            var shader = metalDefaultLibrary.newFunctionWithName(shaderName)
            var computePipeLineDescriptor = MTLComputePipelineDescriptor()
            computePipeLineDescriptor.computeFunction = shader
            //        var computePipelineErrors = NSErrorPointer()
            //            let computePipelineState:MTLComputePipelineState = metalDevice.newComputePipelineStateWithFunction(shader!, completionHandler: {(})
            var computePipelineErrors = NSErrorPointer()
            var computePipelineState:MTLComputePipelineState? = nil
            do {
                computePipelineState = try metalDevice.newComputePipelineStateWithFunction(shader!)
            } catch {
            return (shader, computePipelineState, computePipelineErrors)   
    func createMetalBuffer(var vector:[Float]) -> MTLBuffer {
        let byteLength = vector.count*sizeof(Float)
        return metalDevice.newBufferWithBytes(&vector, length: byteLength, options: MTLResourceOptions.CPUCacheModeDefaultCache)

, , , ,

About Amund Tveit (@atveit - amund@memkite.com)

Amund Tveit works in Memkite on developing large-scale Deep Learning and Search (Convolutional Neural Network) with Swift and Metal for iOS (see deeplearning.education for a Memkite app video demo). He also maintains the deeplearning.university bibliography (github.com/memkite/DeepLearningBibliography)

Amund previously co-founded Atbrox , a cloud computing/big data service company (partner with Amazon Web Services), also doing some “sweat equity” startup investments in US and Nordic startups. His presentations about Hadoop/Mapreduce Algorithms and Search were among top 3% of all SlideShare presentations in 2013 and his blog posts has been frequently quoted by Big Data Industry Leaders and featured on front pages of YCombinator News and Reddit Programming

He previously worked for Google, where he was tech.lead for Google News for iPhone (mentioned as “Google News Now Looks Beautiful On Your iPhone” on Mashable.com), lead a team measuring and improving Google Services in the Scandinavian Countries (Maps and Search) and worked as a software engineer on infrastructure projects. Other work experience include telecom (IBM Canada) and insurance/finance (Storebrand).

Amund has a PhD in Computer Science. His publications has been cited more than 500 times. He also holds 4 US patents in the areas of search and advertisement technology, and a pending US patent in the area of brain-controlled search with consumer-level EEG devices.

Amund enjoys coding, in particular Python, C++ and Swift (iOS)

2 Responses to "Swift and Metal GPU Programming on tvOS for the new Apple TV"

Leave a Reply to Memkite – Deep Learning for iOS (tested on iPhone 6S), tvOS and OS X developed in Metal and Swift | Memkite Cancel reply

Your email address will not be published. Required fields are marked *

You may use these HTML tags and attributes: <a href="" title=""> <abbr title=""> <acronym title=""> <b> <blockquote cite=""> <cite> <code> <del datetime=""> <em> <i> <q cite=""> <strike> <strong>