I use quad_sum to optimize the lighting grid and shadow filter performance. Based on Metal Feature Set Tables, Apple Family 4 should support quad group operations like quad_sum and quad_max. However, on the iPhone X and iPhone 8, during creating pipeline states, we have the following error output: Encountered unlowered function call to air.quad_sum.f32. It works perfectly for iPhone 11 and higher versions. Should I improve my feature-checking logic from Apple Family 4 to Apple Family 5, or do I have other options to fix this unexpected behavior?

I have a test application that draws a large number of simple textured polygons (sprites). Setting CAMetalLayer's displaySyncEnabled to FALSE will cause load on InterruptEventSourceBridge thread in kernel_task. In this case, nanosleep() is used to adjust the amount of METAL commands per unit time so that they are approximately the same. This appears to be a drawing-related thread, but there is no overhead when displaySyncEnabled is TRUE. What are these differences? A specific application is the SDL test program, SDL/test/testsprite.c. https://github.com/libsdl-org/SDL/issues/10475

I have a test application that draws a large number of simple textured polygons (sprites). Setting CAMetalLayer's displaySyncEnabled to FALSE will cause load on InterruptEventSourceBridge thread in kernel_task. (In this case, nanosleep is used to adjust the amount of METAL commands per unit time so that they are approximately the same) This appears to be a drawing-related thread, but there is no overhead when displaySyncEnabled is TRUE. What are these differences? A specific application is the SDL test program, SDL/test/testsprite.c. https://github.com/libsdl-org/SDL/issues/10475

VisionOS 2 beta 5 ，unity text shader errors

We have been having a mysterious crash in our media server app that I've never seen before. After fixing a number of other rare thread safety crashes relating to Metal buffers, this rare crash happens inside a Metal com.Metal.CompletionQueueDispatch? I have no clue what is happening here. It looks to me like Metal is specifically calling abort() for some reason. All of the other threads in the crash log appear to be in a normal state. Thread 70 Crashed:: updateAllMedia Dispatch queue: com.Metal.CompletionQueueDispatch 0 libsystem_kernel.dylib 0x1af572d38 __pthread_kill + 8 1 libsystem_pthread.dylib 0x1af5a7ee0 pthread_kill + 288 2 libsystem_c.dylib 0x1af4e2330 abort + 168 3 libc++abi.dylib 0x1af562b18 abort_message + 132 4 libc++abi.dylib 0x1af552a3c demangling_terminate_handler() + 312 5 libobjc.A.dylib 0x1af4481c8 _objc_terminate() + 160 6 libc++abi.dylib 0x1af561eb4 std::__terminate(void (*)()) + 20 7 libc++abi.dylib 0x1af561e50 std::terminate() + 64 8 libdispatch.dylib 0x1af3e4288 _dispatch_client_callout4 + 40 9 libdispatch.dylib 0x1af40053c _dispatch_mach_msg_invoke + 464 10 libdispatch.dylib 0x1af3eb784 _dispatch_lane_serial_drain + 376 11 libdispatch.dylib 0x1af40125c _dispatch_mach_invoke + 456 12 libdispatch.dylib 0x1af3eb784 _dispatch_lane_serial_drain + 376 13 libdispatch.dylib 0x1af3ec438 _dispatch_lane_invoke + 444 14 libdispatch.dylib 0x1af3eb784 _dispatch_lane_serial_drain + 376 15 libdispatch.dylib 0x1af3ec404 _dispatch_lane_invoke + 392 16 libdispatch.dylib 0x1af3f6c98 _dispatch_workloop_worker_thread + 648 17 libsystem_pthread.dylib 0x1af5a4360 _pthread_wqthread + 288 18 libsystem_pthread.dylib 0x1af5a3080 start_wqthread + 8 Note that the thread name "updateAllMedia" is a misnomer because this thread appears to be a general Metal dispatch queue. I wish there was a debugging option in Metal that called "setThreadName" to name its internal threads.

We have a production Metal app with a complex multithreaded Metal pipeline. When everything is operating smoothly, it works great. Even when extremely overloaded, it does not crash for days at a time. This isn't good enough for our users. Unfortunately, when I have zero visibility into id, I have no way of knowing when metal is "done" with an id. When overloaded, stale metal render passes need to be 'aborted', which results in metal callbacks not being called. for example, these callbacks may not be called after an aborted pass: id<MTLCommandBuffer> m_cmdbuf; [m_cmdbuf addScheduledHandler:^(id <MTLCommandBuffer> cb) { cpr->scheduled = MachAbsoluteTime(); }]; [m_cmdbuf addCompletedHandler:^(id <MTLCommandBuffer> cb) { cpr->completed = MachAbsoluteTime(); }]; For the moment, our workaround is a system which waits a few seconds after we "think" a rendering pass should be done with all its (aborted) resources before releasing buffers. This is not ideal, to say the least. So, in summary, my question is, it would be nice to be able to 'query' an id to know when metal is done with it, so that we know that its safe to release it along with our own internal resources. Is there any such (undocumented) mechanism? I have exhaustively read all existing Metal documentation many times. An idea that I've been toying with... it would be nice to have something akin to Zombie detection running all the time for id only. In OpenGL, it was OK to use a released texture... you may display a bad frame, but not CRASH!. Is there any similar option for id?

How many 32-bit variables can I use concurrently in a single thread of a Metal compute kernel without worrying about the variables getting spilled into the device memory? Alternatively: how many 32-bit registers does a single thread have available for itself? Let's say that each thread of my compute kernel needs to store and work with its own array of N float variables, where N can be 128, 256, 512 or more. To achieve maximum possible performance, I do not want to the local thread variables to get spilled into the slow device memory. I want all N variables to be stored "on-chip", in the thread memory space. To make my question more concrete, let's say there is an array thread float localArray[N]. Assuming an unrealistic hypothetical scenario where localArray is the only variable in the whole kernel, what is the maximum value of N for which no portion of localArray would get spilled into the device memory? I searched in the Metal feature set tables, but I could not find any details.

Our app encountered the following error: Execution of the command buffer was aborted due to an error during execution. Ignored (for causing prior/excessive GPU errors) (00000004:kIOGPUCommandBufferCallbackErrorSubmissionsIgnored)

The Metal feature set tables specifies that beginning with the Apple4 family, the "Maximum threads per threadgroup" is 1024. Given that a single threadgroup is guaranteed to be run on the same GPU shader core, it means that a shader core of any new Apple GPU must be capable of running at least 1024/32 = 32 warps in parallel. From the WWDC session "Scale compute workloads across Apple GPUs (6:17)": For relatively complex kernels, 1K to 2K concurrent threads per shader core is considered a very good occupancy. The cited sentence suggests that a single shader core is capable of running at least 2K (I assume this is meant to be 2048) threads in parallel, so 2048/32 = 64 warps running in parallel. However, I am curious what is the maximum theoretical amount of warps running in parallel on a single shader core (it sounds like it is more than 64). The WWDC session mentions 2K to be only "very good" occupancy. How many threads would be "the best possible" occupancy?

Greetings! I have been battling with a bit of a tough issue. My use case is running a pixelwise regression model on a 2D array of images using CIImageProcessorKernel and a custom Metal Shader. It mostly works great, but the issue that arises is that if the regression calculation in Metal takes too long, an error occurs and the resulting output texture has strange artifacts, for example: The specific error is: Error excuting command buffer = Error Domain=MTLCommandBufferErrorDomain Code=1 "Internal Error (0000000e:Internal Error)" UserInfo={NSLocalizedDescription=Internal Error (0000000e:Internal Error), NSUnderlyingError=0x60000320ca20 {Error Domain=IOGPUCommandQueueErrorDomain Code=14 "(null)"}} (com.apple.CoreImage) There are multiple levels of concurrency: Swift Concurrency calling the Core Image code (which shouldn't have an impact) and of course the Metal command buffer. Is there anyway to ensure the compute command encoder can complete its work? Here is the full implementation of my CIImageProcessorKernel subclass: class ParametricKernel: CIImageProcessorKernel { static let device = MTLCreateSystemDefaultDevice()! override class var outputFormat: CIFormat { return .BGRA8 } override class func formatForInput(at input: Int32) -> CIFormat { return .BGRA8 } override class func process(with inputs: [CIImageProcessorInput]?, arguments: [String : Any]?, output: CIImageProcessorOutput) throws { guard let commandBuffer = output.metalCommandBuffer, let images = arguments?["images"] as? [CGImage], let mask = arguments?["mask"] as? CGImage, let fillTime = arguments?["fillTime"] as? CGFloat, let betaLimit = arguments?["betaLimit"] as? CGFloat, let alphaLimit = arguments?["alphaLimit"] as? CGFloat, let errorScaling = arguments?["errorScaling"] as? CGFloat, let timing = arguments?["timing"], let TTRThreshold = arguments?["ttrthreshold"] as? CGFloat, let input = inputs?.first, let sourceTexture = input.metalTexture, let destinationTexture = output.metalTexture else { return } guard let kernelFunction = device.makeDefaultLibrary()?.makeFunction(name: "parametric") else { return } guard let commandEncoder = commandBuffer.makeComputeCommandEncoder() else { return } let imagesTexture = Texture.textureFromImages(images) let pipelineState = try device.makeComputePipelineState(function: kernelFunction) commandEncoder.setComputePipelineState(pipelineState) commandEncoder.setTexture(imagesTexture, index: 0) let maskTexture = Texture.textureFromImages([mask]) commandEncoder.setTexture(maskTexture, index: 1) commandEncoder.setTexture(destinationTexture, index: 2) var errorScalingFloat = Float(errorScaling) let errorBuffer = device.makeBuffer(bytes: &errorScalingFloat, length: MemoryLayout<Float>.size, options: []) commandEncoder.setBuffer(errorBuffer, offset: 0, index: 1) // Other buffers omitted.... let threadsPerThreadgroup = MTLSizeMake(16, 16, 1) let width = Int(ceil(Float(sourceTexture.width) / Float(threadsPerThreadgroup.width))) let height = Int(ceil(Float(sourceTexture.height) / Float(threadsPerThreadgroup.height))) let threadGroupCount = MTLSizeMake(width, height, 1) commandEncoder.dispatchThreadgroups(threadGroupCount, threadsPerThreadgroup: threadsPerThreadgroup) commandEncoder.endEncoding() } }

I'm trying to create a custom Metal-based visual effect as a UIView to be used inside an existing UIKit-based interface. (An example might be a view that applies a blur effect to what's behind it.) I need to capture the MTLTexture of what's behind the view so that I can feed it to MTLRenderCommandEncoder.setFragmentTexture(_:index:). Can someone show me how or point me to an example? Thanks!

I am trying to convert a ThreeJS project to Metal for the Vision Pro. The issue is ThreeJS doesn't do any color space conversion (when I output a color in a fragment shader and then read it using the digital color meter in SRGB mode I get the same value I inputed in the fragment shader) This is not the case when using metal. When setting up my LayerRenderer I set the colorFormat to rgba16Unorm since it is the only non srgb color format supported on the vision pro apps. However switching between bgra8Unorm_srgb and rgba16Unorm seems to have no affect. when I set up the renderPassDescriptor I use the drawable colorTexture renderPassDescriptor.colorAttachments[0].texture = drawable.colorTextures[0] and when printing its pixel format it seems to be passed from the configuration. If there is anyway to disable this behavior or perform an inverse function of such that I get the original value out from the shader, that would be appreciated.

Guten Tag, my project is simple, first I want draw wired Hexa,-Tetra- and Octahedrons. I draw a cube with Metal but I didn't found rotation, translation and scale. I have searched help , the examples I found are too complicated for me. Mit freundlichen Grüßen VanceRegnet

I am trying to get a little game prototype up and running using Metal using the metal-cpp libraries where I run everything natively at 120Hz with a coupled renderer using Vsync turned on so that I have the absolute physically minimum input to photon latency possible. // Create the metal view SDL_MetalView metal_view = SDL_Metal_CreateView(window); CA::MetalLayer *swap_chain = (CA::MetalLayer *)SDL_Metal_GetLayer(metal_view); // Set up the Metal device MTL::Device *device = MTL::CreateSystemDefaultDevice(); swap_chain->setDevice(device); swap_chain->setPixelFormat(MTL::PixelFormat::PixelFormatBGRA8Unorm); swap_chain->setDisplaySyncEnabled(true); swap_chain->setMaximumDrawableCount(2); I am using SDL3 just for creating the window. Now when I go through my game / render loop - I stall for a long time on getting the next drawable which is understandable - my app runs in about 2-3ms. m_CurrentContext->m_Drawable = m_SwapChain->nextDrawable(); m_CurrentContext->m_CommandBuffer = m_CommandQueue->commandBuffer()->retain(); char frame_label[32]; snprintf(frame_label, sizeof(frame_label), "Frame %d", m_FrameIndex); m_CurrentContext->m_CommandBuffer->setLabel(NS::String::string(frame_label, NS::UTF8StringEncoding)); m_CurrentContext->m_RenderPassDescriptor[ERenderPassTypeNormal] = MTL::RenderPassDescriptor::alloc()->init(); MTL::RenderPassColorAttachmentDescriptor* cd = m_CurrentContext->m_RenderPassDescriptor[ERenderPassTypeNormal]->colorAttachments()->object(0); cd->setTexture(m_CurrentContext->m_Drawable->texture()); cd->setLoadAction(MTL::LoadActionClear); cd->setClearColor(MTL::ClearColor( 0.53f, 0.81f, 0.98f, 1.0f )); cd->setStoreAction(MTL::StoreActionStore); However my ProMotion display does not reliably run at 120Hz when fullscreen and using the direct to display system - it seems to run faster when windowed in composite which is the opposite of what I would expect. The Metal HUD says 120Hz, but the delay to getting the next drawable and looking at what Instruments is saying tells otherwise. When I profile it, the game loop has completed and is sitting there waiting for the next drawable, but the screen does not want to complete in 8.33ms, so the whole thing slows down for no discernible reason. Also as a game developer it is very strange for the command buffer to actually need the drawable texture free to be allowed to encode commands - usually the command buffers and swapping the front and back render buffers are not directly dependent on each other. Usually you only actually need the render buffer texture free when you want to draw to it. I could give myself another drawable, but because I am completing in less than 3ms, all it would do would be to add another frame of latency. I also looked at the FramePacing example and its behaviour is even worse at having high framerate with low latency - the direct to display is always rejected for some reason. Is this just a flaw in the Metal API? Or am I missing something important? I hope someone can help - the behaviour of the display is baffling.

I wanted to try the new logging feature for Metal but could not get it to work. I modified the PerformingCalculationsOnAGPU example by adding os_log_default.log_debug("Hello thread: %d", index); to log the current thread id. But never saw any messages neither in the console nor in Xcode. I also added the -fmetal-enable-logging flag. I am running the Sequoia release candidate 15.0 (24A335) on M1 Max and Xcode 16.0 (16A242). What am I missing?

Hi there, I'm trying to test the "Drawing fully immersive content using Metal" , but when I select Language: Swift, it still shows Objective C code in some sample codes. Please check and update the document Swift Code, thank you.

The following code runs fine when compiled with Swift 5, but crashes when compiled with Swift 6 (stack trace below). In the draw method, commenting out the addCompletedHandler line fixes the problem. I'm testing on iOS 18.0 and see the same behavior in both the simulator and on a device. What's going on here? import Metal import MetalKit import UIKit class ViewController: UIViewController { @IBOutlet var metalView: MTKView! private var commandQueue: MTLCommandQueue? override func viewDidLoad() { super.viewDidLoad() guard let device = MTLCreateSystemDefaultDevice() else { fatalError("expected a Metal device") } self.commandQueue = device.makeCommandQueue() metalView.device = device metalView.enableSetNeedsDisplay = true metalView.isPaused = true metalView.delegate = self } } extension ViewController: MTKViewDelegate { func mtkView(_ view: MTKView, drawableSizeWillChange size: CGSize) {} func draw(in view: MTKView) { guard let commandQueue, let commandBuffer = commandQueue.makeCommandBuffer() else { return } commandBuffer.addCompletedHandler { _ in } // works with Swift 5, crashes with Swift 6 commandBuffer.commit() } } Here's the stack trace: Thread 10 Queue : connection Queue (serial) #0 0x000000010581c3f8 in _dispatch_assert_queue_fail () #1 0x000000010581c384 in dispatch_assert_queue () #2 0x00000002444c63e0 in swift_task_isCurrentExecutorImpl () #3 0x0000000104d71ec4 in closure #1 in ViewController.draw(in:) () #4 0x0000000104d71f58 in thunk for @escaping @callee_guaranteed (@guaranteed MTLCommandBuffer) -> () () #5 0x0000000105ef1950 in __47-[CaptureMTLCommandBuffer _preCommitWithIndex:]_block_invoke_2 () #6 0x00000001c50b35b0 in -[MTLToolsCommandBuffer invokeCompletedHandlers] () #7 0x000000019e94d444 in MTLDispatchListApply () #8 0x000000019e94f558 in -[_MTLCommandBuffer didCompleteWithStartTime:endTime:error:] () #9 0x000000019e95352c in -[_MTLCommandQueue commandBufferDidComplete:startTime:completionTime:error:] () #10 0x0000000226ef50b0 in handleMainConnectionReplies () #11 0x00000001800c9690 in _xpc_connection_call_event_handler () #12 0x00000001800cad90 in _xpc_connection_mach_event () #13 0x000000010581a86c in _dispatch_client_callout4 () #14 0x0000000105837950 in _dispatch_mach_msg_invoke () #15 0x0000000105822870 in _dispatch_lane_serial_drain () #16 0x0000000105838c10 in _dispatch_mach_invoke () #17 0x0000000105822870 in _dispatch_lane_serial_drain () #18 0x00000001058237b0 in _dispatch_lane_invoke () #19 0x00000001058301f0 in _dispatch_root_queue_drain_deferred_wlh () #20 0x000000010582f75c in _dispatch_workloop_worker_thread () #21 0x00000001050abb74 in _pthread_wqthread ()

In my project I need to do the following: In runtime create metal Dynamic library from source. In runtime create metal Executable library from source and Link it with my previous created Dynamic library. Create compute pipeline using those two libraries created above. But I get the following error at the third step: Error Domain=AGXMetalG15X_M1 Code=2 "Undefined symbols: _Z5noisev, referenced from: OnTheFlyKernel " UserInfo={NSLocalizedDescription=Undefined symbols: _Z5noisev, referenced from: OnTheFlyKernel } import Foundation import Metal class MetalShaderCompiler { let device = MTLCreateSystemDefaultDevice()! var pipeline: MTLComputePipelineState! func compileDylib() -> MTLDynamicLibrary { let source = """ #include <metal_stdlib> using namespace metal; half3 noise() { return half3(1, 0, 1); } """ let option = MTLCompileOptions() option.libraryType = .dynamic option.installName = "@executable_path/libFoundation.metallib" let library = try! device.makeLibrary(source: source, options: option) let dylib = try! device.makeDynamicLibrary(library: library) return dylib } func compileExlib(dylib: MTLDynamicLibrary) -> MTLLibrary { let source = """ #include <metal_stdlib> using namespace metal; extern half3 noise(); kernel void OnTheFlyKernel(texture2d<half, access::read> src [[texture(0)]], texture2d<half, access::write> dst [[texture(1)]], ushort2 gid [[thread_position_in_grid]]) { half4 rgba = src.read(gid); rgba.rgb += noise(); dst.write(rgba, gid); } """ let option = MTLCompileOptions() option.libraryType = .executable option.libraries = [dylib] let library = try! self.device.makeLibrary(source: source, options: option) return library } func runtime() { let dylib = self.compileDylib() let exlib = self.compileExlib(dylib: dylib) let pipelineDescriptor = MTLComputePipelineDescriptor() pipelineDescriptor.computeFunction = exlib.makeFunction(name: "OnTheFlyKernel") pipelineDescriptor.preloadedLibraries = [dylib] pipeline = try! device.makeComputePipelineState(descriptor: pipelineDescriptor, options: .bindingInfo, reflection: nil) } }

I use xcode16 and swiftUI for programming on a macos15 system. There is a problem. When I render a picture through mtkview, it is normal when displayed on a regular view. However, when the view is displayed through the .sheet method, the image cannot be displayed. There is no error message from xcode. import Foundation import MetalKit import SwiftUI struct CIImageDisplayView: NSViewRepresentable { typealias NSViewType = MTKView var ciImage: CIImage init(ciImage: CIImage) { self.ciImage = ciImage } func makeNSView(context: Context) -&gt; MTKView { let view = MTKView() view.delegate = context.coordinator view.preferredFramesPerSecond = 60 view.enableSetNeedsDisplay = true view.isPaused = true view.framebufferOnly = false if let defaultDevice = MTLCreateSystemDefaultDevice() { view.device = defaultDevice } view.delegate = context.coordinator return view } func updateNSView(_ nsView: MTKView, context: Context) { } func makeCoordinator() -&gt; RawDisplayRender { RawDisplayRender(ciImage: self.ciImage) } class RawDisplayRender: NSObject, MTKViewDelegate { // MARK: Metal resources var device: MTLDevice! var commandQueue: MTLCommandQueue! // MARK: Core Image resources var context: CIContext! var ciImage: CIImage init(ciImage: CIImage) { self.ciImage = ciImage self.device = MTLCreateSystemDefaultDevice() self.commandQueue = self.device.makeCommandQueue() self.context = CIContext(mtlDevice: self.device) } func mtkView(_ view: MTKView, drawableSizeWillChange size: CGSize) {} func draw(in view: MTKView) { guard let currentDrawable = view.currentDrawable, let commandBuffer = commandQueue.makeCommandBuffer() else { return } let dSize = view.drawableSize let drawImage = self.ciImage let destination = CIRenderDestination(width: Int(dSize.width), height: Int(dSize.height), pixelFormat: view.colorPixelFormat, commandBuffer: commandBuffer, mtlTextureProvider: { () -&gt; MTLTexture in return currentDrawable.texture }) _ = try? self.context.startTask(toClear: destination) _ = try? self.context.startTask(toRender: drawImage, from: drawImage.extent, to: destination, at: CGPoint(x: (dSize.width - drawImage.extent.width) / 2, y: 0)) commandBuffer.present(currentDrawable) commandBuffer.commit() } } } struct ShowCIImageView: View { let cii = CIImage.init(contentsOf: Bundle.main.url(forResource: "9-10", withExtension: "jpg")!)! var body: some View { CIImageDisplayView.init(ciImage: cii).frame(width: 500, height: 500).background(.red) } } struct ContentView: View { @State var showImage = false var body: some View { VStack { Image(systemName: "globe") .imageScale(.large) .foregroundStyle(.tint) Text("Hello, world!") ShowCIImageView() Button { showImage = true } label: { Text("showImage") } } .frame(width: 800, height: 800) .padding() .sheet(isPresented: $showImage) { ShowCIImageView() } } }

hello everyone. I got a texture loading error on visionOS 2.0: Can't create texture(Error Domain=MTKTextureLoaderErrorDomain Code=0 "Pixel format(MTLPixelFormatInvalid) is not valid on this device" UserInfo={NSLocalizedDescription=Pixel format(MTLPixelFormatInvalid) is not valid on this device, MTKTextureLoaderErrorKey=Pixel format(MTLPixelFormatInvalid) is not valid on this device} But this texture can load correctly on visionOS1.3. I don't know what happen between visionOS1.3 and visionOS2.0. The texture is a ktx file which stores cubemap that encoding in astc6x6hdr. And the ktx texture has a glInternalFormat info: GL_COMPRESSED_RGBA_ASTC_6x6. I wonder if visionOS2.0 no longer supports astc6x6hdr cubemap format, or there is something wrong with my assets.