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.

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?

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.

VisionOS 2 beta 5 ，unity text shader errors

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

Why do I get this error almost immediately on starting my rendering pass? Multiline BlockQuote. 2024-05-29 20:02:22.744035-0500 RoomPlanExampleApp[491:10341] [] <<<< AVPointCloudData >>>> Fig assert: "_dataBuffer" at bail (AVPointCloudData.m:217) - (err=0) 2024-05-29 20:02:22.744455-0500 RoomPlanExampleApp[491:10341] [] <<<< AVPointCloudData >>>> Fig assert: "_dataBuffer" at bail (AVPointCloudData.m:217) - (err=0) 2024-05-29 20:05:54.079981-0500 RoomPlanExampleApp[491:10025] [CAMetalLayer nextDrawable] returning nil because allocation failed. 2024-05-29 20:05:54.080144-0500 RoomPlanExampleApp[491:10341] [] <<<< AVPointCloudData >>>> Fig assert: "_dataBuffer" at bail (AVPointCloudData.m:217) - (err=0)

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?

Is there new API for generating Indirect Commands for the Metal Shader Converter? Is there any example project? I currently use a shader to copy indirect commands. Is there a way to do that with the new Shader Converter pipeline?

I am trying to work out how to enable the Metal HUD on iOS for App Store games? I am aware you can go into Developer settings and enable it… but it only appears for some games like HADES or TestFlight apps. I know the HUD appears for sideloaded games too. With sideloadly, I’ve sideloaded GRID Autosport and Myst - per screenshots. But it’s a very time consuming process, on demand resources usually don’t download… and I’m not sure if its legal. I tried using Xcode and Attach to Process for a game like Resident Evil 7 or just anything… but it doesn’t work. I’ve tried restoring a backup and editing the .GlobalPreferences.plist and info.plist file with “MetalForceHudEnabled Boolean Yes”, in a new row… but nothing. any ideas?

I'm trying to create heat maps for a variety of functions of two variables. My first implementation didn't use Metal and was far too slow so now I'm looking into doing it with Metal. I managed to get a very simple example running but I can't figure out how to pass different functions to the fragment shader. Here's the example: in ContentView.swift: struct ContentView: View { var body: some View { Rectangle() .aspectRatio(contentMode: .fit) .visualEffect { content, gp in let width = Shader.Argument.float(gp.size.width) let height = Shader.Argument.float(gp.size.height) return content.colorEffect( ShaderLibrary.heatMap(width, height) ) } } } in Shader.metal: #include <metal_stdlib> using namespace metal; constant float twoPi = 6.283185005187988; // input in [0,1], output in [0,1] float f(float x) { return (sin(twoPi * x) + 1) / 2; } // inputs in [0,1], output in [0,1] float g(float x, float y) { return f(x) * f(y); } [[ stitchable ]] half4 heatMap(float2 pos, half4 color, float width, float height) { float u = pos.x / width; float v = pos.y / height; float c = g(u, v); return half4(c/2, 1-c, c, 1); } As it is, it works great and is blazing fast... ...but the function I'm heat-mapping is hardcoded in the metal file. I'd like to be able to write different functions in Swift and pass them to the shader from within SwiftUI (ie, from the ContentView, by querying a model to get the function). I tried something like this in the metal file: // (u, v) in [0,1] x [0,1] // w = f(u, v) in [0,1] [[ stitchable ]] half4 heatMap( float2 pos, half4 color, float width, float height, float (*f) (float u, float v), half4 (*c) (float w) ) { float u = pos.x / width; float v = pos.y / height; float w = f(u, v); return c(w); } but I couldn't get Swift and C++ to work together to make sense of the function pointers and and now I'm stuck. Any help is greatly appreciated. Many thanks!

Hello As part of my app, I am using Metal shaders on CustomMaterials created and managed using RealityKit. Using the ECS approach, I have a Shader system that iterates through all my materials every frame and passes a SIMD4 of variables (that I can manage on the swift side) that can be interpreted and used every frame on the Metal side to influence elements of the shader. This does work as intended but is limited to just 4 variables when I need more for my use case. I've experimented with trying multiple simd4 or other approaches for passing these to metal and be useable but I haven't had very much luck. I was hoping for some recommendations on the best scalable approach. Swift: class ShaderSystem: System { static let query = EntityQuery(where: .has(ModelComponent.self)) private var startTime: Date required init(scene: Scene) { startTime = Date() } func update(context: SceneUpdateContext) { let audioLevel = AudioSessionManager.shared.audioLevel let elapsedTime = Float(Date().timeIntervalSince(startTime)) guard let sceneType = SceneManager.shared.currentScenes.keys.first else { return } let sceneTime = SceneComposer.shared.getSceneTime(for: sceneType) let multiplier = ControlManager.shared.getControlValue(parameterName: "elapsedTimeMultiplier") ?? 1.0 for entity in context.scene.performQuery(Self.query) { guard var modelComponent = entity.components[ModelComponent.self] as? ModelComponent else { continue } modelComponent.materials = modelComponent.materials.map { material in guard var customMaterial = material as? CustomMaterial else { return material } // Passing audioLevel, elapsedTime, sceneTime, and multiplier customMaterial.custom.value = SIMD4<Float>(audioLevel, elapsedTime, sceneTime, multiplier) return customMaterial } entity.components[ModelComponent.self] = modelComponent } } } metal: struct CustomMaterialUniforms { float4 custom; }; [[visible]] void fractalShader(realitykit::surface_parameters params) { auto uniforms = params.uniforms(); float4 customValues = uniforms.custom_parameter(); float audioLevel = customValues.x; .... Thank you for the assistance

Here is the test code run in a macOS app (MacOS 15 Beta3). If the excutable path does not contain Chinese character, every thing go as We expect. Otherwise(simply place excutable in a Chinese named directory) , the MTLLibrary We made by newLibraryWithSource: function contains no functions, We just got logs: "Library contains the following functions: {}" "Function 'squareKernel' not found." Note: macOS 14 works fine id<MTLDevice> device = MTLCreateSystemDefaultDevice(); if (!device) { NSLog(@"not support Metal."); } NSString *shaderSource = @ "#include <metal_stdlib>\n" "using namespace metal;\n" "kernel void squareKernel(device float* data [[buffer(0)]], uint gid [[thread_position_in_grid]]) {\n" " data[gid] *= data[gid];\n" "}"; MTLCompileOptions *options = [[MTLCompileOptions alloc] init]; options.languageVersion = MTLLanguageVersion2_0; NSError *error = nil; id<MTLLibrary> library = [device newLibraryWithSource:shaderSource options:options error:&error]; if (error) { NSLog(@"New MTLLibrary error: %@", error); } NSArray<NSString *> *functionNames = [library functionNames]; NSLog(@"Library contains the following functions: %@", functionNames); id<MTLFunction> computeShaderFunction = [library newFunctionWithName:@"squareKernel"]; if (computeShaderFunction) { NSLog(@"Found function 'squareKernel'."); NSError *pipelineError = nil; id<MTLComputePipelineState> pipelineState = [device newComputePipelineStateWithFunction:computeShaderFunction error:&pipelineError]; if (pipelineError) { NSLog(@"Create pipeline state error: %@", pipelineError); } NSLog(@"Create pipeline state succeed!"); } else { NSLog(@"Function 'squareKernel' not found."); }

I'm trying to create a MTLFXTemporalScaler as follows (this is adapted from the sample code): func updateTemporalScaler() { let desc = MTLFXTemporalScalerDescriptor() desc.inputWidth = renderTarget.renderSize.width desc.inputHeight = renderTarget.renderSize.height desc.outputWidth = renderTarget.windowSize.width desc.outputHeight = renderTarget.windowSize.height desc.colorTextureFormat = .bgra8Unorm desc.depthTextureFormat = .depth32Float desc.motionTextureFormat = .rg16Float desc.outputTextureFormat = .bgra8Unorm guard let temporalScaler = desc.makeTemporalScaler(device: device) else { fatalError("The temporal scaler effect is not usable!") } temporalScaler.motionVectorScaleX = Float(renderTarget.renderSize.width) temporalScaler.motionVectorScaleY = Float(renderTarget.renderSize.height) mfxTemporalScaler = temporalScaler } I'm getting the following error the 3rd time the code is called: /AppleInternal/Library/BuildRoots/91a344b1-f985-11ee-b563-fe8bc7981bff/Library/Caches/com.apple.xbs/Sources/MetalPerformanceShadersGraph/mpsgraph/MetalPerformanceShadersGraph/Runtimes/MPSRuntime/Operations/RegionOps/ANRegion.mm:855: failed assertion `ANE intermediate buffer handle not same!' When I copy the code out to a playground, it succeeds when called with the same sequence of descriptors. Does this seem like a bug with MTLFXTemporalScaler?

Unity 2022.3.33f1 For some reason modifying MeshRenderer material shader SetVectorArray doesn't work on IOS, but it works on android and windows builds!! I was working on Fog Of War, where I used SimpleFOW by Revision3 it's very simple FOW shader where it manipulates the alpha based on UVs vertices. This is the FogOfWarShaderControl.cs script using System.Collections; using System.Collections.Generic; using UnityEngine; namespace SimpleFOW { [RequireComponent(typeof(MeshRenderer))] public class FogOfWarShaderControl : MonoBehaviour { public static FogOfWarShaderControl Instance { get; private set; } [Header("Maximum amount of revealing points")] [SerializeField] private uint maximumPoints = 512; [Header("Game Camera")] [SerializeField] private Camera mainCamera; private List<Vector4> points = new List<Vector4>(); private Vector2 meshSize, meshExtents; private Vector4[] sendBuffer; private MeshRenderer meshRenderer; private void Awake() { Instance = this; Init(); } // Initialize required variables public void Init() { meshRenderer = GetComponent<MeshRenderer>(); meshExtents = meshRenderer.bounds.extents; meshSize = meshRenderer.bounds.size; points = new List<Vector4>(); sendBuffer = new Vector4[maximumPoints]; } // Transform world point to UV coordinate of FOW mesh public Vector2 WorldPointToMeshUV(Vector2 wp) { Vector2 toRet = Vector2.zero; toRet.x = (transform.position.x - wp.x + meshExtents.x) / meshSize.x; toRet.y = (transform.position.y - wp.y + meshExtents.y) / meshSize.y; return toRet; } // Show or hide FOW public void SetEnabled(bool on) { meshRenderer.enabled = on; } // Add revealing point to FOW renderer if amount of points is lower than MAX_POINTS public void AddPoint(Vector2 worldPoint) { if (points.Count < maximumPoints) { points.Add(WorldPointToMeshUV(worldPoint)); } } // Remove FOW revealing point public void RemovePoint(Vector2 worldPoint) { if (worldPoint == new Vector2(0, 0)) { return; } if (points.Contains(WorldPointToMeshUV(worldPoint))) { points.Remove(WorldPointToMeshUV(worldPoint)); } } // Send any change to revealing point list to shader for rendering public void SendPoints() { points.ToArray().CopyTo(sendBuffer, 0); meshRenderer.material.SetVectorArray("_PointArray", sendBuffer); meshRenderer.material.SetInt("_PointCount", points.Count); } // Send new range value to shader public void SendRange(float range) { meshRenderer.material.SetFloat("_RadarRange", range); } // Send new scale value to shader public void SendScale(float scale) { meshRenderer.material.SetFloat("_Scale", scale); } } } And this is the FogOfWar.shader Shader "Revision3/FogOfWar" { Properties { _MainTex ("Texture", 2D) = "black" {} _PointCount("Point count", Range(0,512)) = 0 _Scale("Scale", Float) = 1.0 _RadarRange("Range", Float) = .5 _MaxAlpha("Maximum Alpha", Float) = 1.0 } SubShader { Tags { "RenderType"="Transparent" "Queue"="Transparent" } LOD 100 ZWrite Off Blend SrcAlpha OneMinusSrcAlpha Pass { CGPROGRAM #pragma vertex vert #pragma fragment frag // make fog work #pragma multi_compile_fog #include "UnityCG.cginc" struct appdata { float4 vertex : POSITION; float2 uv : TEXCOORD0; }; struct v2f { float2 uv : TEXCOORD0; UNITY_FOG_COORDS(1) float4 vertex : SV_POSITION; }; sampler2D _MainTex; float4 _MainTex_ST; float _RadarRange; uint _PointCount; float _Scale; float _MaxAlpha; float2 _PointArray[512]; v2f vert (appdata v) { v2f o; o.vertex = UnityObjectToClipPos(v.vertex); o.uv = TRANSFORM_TEX(v.uv, _MainTex); return o; } float getDistance(float2 pa[512], float2 uv) { float cdist = 99999.0; for (uint i = 0; i < _PointCount; i++) { cdist = min(cdist, distance(pa[i]*_Scale, uv)); } return cdist; } fixed4 frag (v2f i) : SV_Target { // sample the texture fixed4 col = tex2D(_MainTex, i.uv); i.uv *= _Scale; if (_PointCount > 0) col.w = min(_MaxAlpha, max(0.0f, getDistance(_PointArray, i.uv) - _RadarRange)); else col.w = _MaxAlpha; return col; } ENDCG } } } Now I create a gameobject called FogOfWar as follows And then in Unit.cs script and Building.cs script I add the following logic private Vector3 lastPos; private void Update() { if (lastPos != transform.position) { FogOfWarShaderControl.Instance.RemovePoint(lastPos); FogOfWarShaderControl.Instance.AddPoint(transform.position); lastPos = transform.position; FogOfWarShaderControl.Instance.SendPoints(); } } Now this gives me the effect of FOW as follows on IOS where the result should be as follows on other devices I don't know what causes this to happen only on IOS devices. The logic works fine on android/windows/Linux/editor but not IOS devices. So why metal API doesn't support shader set vector array?!

Hi. I am developing an immersive application in Metal. The problem is that when the application starts for the first time, it shows two system windows: one related to the player's surrounding environment and the other to hand tracking permissions. The latter is the one in focus, and when I select the option to grant permissions, both windows disappear and the transition to the immersive view does not occur. I don't know how to handle this issue.

The title is self-exploratory. I wasn't able to find the CAMetalDisplayLink on the most recent metal-cpp release (metal-cpp_macOS15_iOS18-beta). Are there any plans to include it in the next release?

What is the info property of SwiftUI::Layer? I couldn't find any document or resource about it. It appears in SwiftUI::Layer's definition: struct Layer { metal::texture2d<half> tex; float2 info[5]; /// Samples the layer at `p`, in user-space coordinates, /// interpolating linearly between pixel values. Returns an RGBA /// pixel value, with color components premultipled by alpha (i.e. /// [R*A, G*A, B*A, A]), in the layer's working color space. half4 sample(float2 p) const { p = metal::fma(p.x, info[0], metal::fma(p.y, info[1], info[2])); p = metal::clamp(p, info[3], info[4]); return tex.sample(metal::sampler(metal::filter::linear), p); } };

got the above error when using PyTorch with diffusion model same error comes up in few games when running through D3D metal Anyone has any idea what this is about?

Suppose I want to draw a red rectangle onto my render target using a compute shader. id<MTLComputeCommandEncoder> encoder = [commandBuffer computeCommandEncoder]; [encoder setComputePipelineState:pipelineState]; simd_ushort2 position = simd_make_ushort2(100, 100); simd_ushort2 size = simd_make_ushort2(50, 50); [encoder setBytes:&position length:sizeof(position) atIndex:0]; [encoder setTexture:drawable.texture atIndex:0]; [encoder dispatchThreads:MTLSizeMake(size.x, size.y, 1) threadsPerThreadgroup:MTLSizeMake(32, 32, 1)]; [encoder endEncoding]; #include <metal_stdlib> using namespace metal; kernel void Compute(ushort2 position_in_grid [[thread_position_in_grid]], constant ushort2 &position, texture2d<half, access::write> texture) { texture.write(half4(1, 0, 0, 1), position_in_grid + position); } This works just fine: Now, say for whatever reason I want to start using imageblocks in my compute kernel. First, I set the imageblock size on the CPU side: id<MTLComputeCommandEncoder> encoder = [commandBuffer computeCommandEncoder]; [encoder setComputePipelineState:pipelineState]; MTLSize threadgroupSize = MTLSizeMake(32, 32, 1); [encoder setImageblockWidth:threadgroupSize.width height:threadgroupSize.height]; simd_ushort2 position = simd_make_ushort2(100, 100); simd_ushort2 size = simd_make_ushort2(50, 50); [encoder setBytes:&position length:sizeof(position) atIndex:0]; [encoder setTexture:drawable.texture atIndex:0]; MTLSize gridSize = MTLSizeMake(size.x, size.y, 1); [encoder dispatchThreads:gridSize threadsPerThreadgroup:threadgroupSize]; And then I update the compute kernel to simply declare the imageblock – note I never actually read from or write to it: #include <metal_stdlib> using namespace metal; struct Foo { int foo; }; kernel void Compute(ushort2 position_in_grid [[thread_position_in_grid]], constant ushort2 &position, texture2d<half, access::write> texture, imageblock<Foo> imageblock) { texture.write(half4(1, 0, 0, 1), position_in_grid + position); } And now out of nowhere Metal’s shader validation starts complaining about mismatched texture usage flags: 2024-06-22 00:57:15.663132+1000 TextureUsage[80558:4539093] [GPUDebug] Texture usage flags mismatch executing kernel function "Compute" encoder: "1", dispatch: 0 2024-06-22 00:57:15.672004+1000 TextureUsage[80558:4539093] [GPUDebug] Texture usage flags mismatch executing kernel function "Compute" encoder: "1", dispatch: 0 2024-06-22 00:57:15.682422+1000 TextureUsage[80558:4539093] [GPUDebug] Texture usage flags mismatch executing kernel function "Compute" encoder: "1", dispatch: 0 2024-06-22 00:57:15.687587+1000 TextureUsage[80558:4539093] [GPUDebug] Texture usage flags mismatch executing kernel function "Compute" encoder: "1", dispatch: 0 2024-06-22 00:57:15.698106+1000 TextureUsage[80558:4539093] [GPUDebug] Texture usage flags mismatch executing kernel function "Compute" encoder: "1", dispatch: 0 The texture I’m writing to comes from a CAMetalDrawable whose associated CAMetalLayer has framebufferOnly set to NO. What am I missing?