Render advanced 3D graphics and perform data-parallel computations using graphics processors using Metal.

Metal Documentation

Posts under Metal subtopic

Post

Replies

Boosts

Views

Activity

D3DMetal: present hangs permanently when the game window becomes occluded
Environment macOS 27.0.1 (26A434), Apple M4 Pro Apple Game Porting Toolkit 4.0b2 (D3DMetal.framework) via CrossOver 27 D3D12 game, borderless fullscreen window (2540x1650) Symptom When the game window becomes occluded — Mission Control, a system OSD (volume/brightness), or another window being brought to the front — the game never renders another frame. It does not crash: GPU utilization drops to 0%, the process stays alive with near-zero CPU, and it never recovers. Killing and restarting the game is the only way out. D3DMetal log at the moment it happens: [D3DMetal:FATAL:%llX][%s:%d] Failed to get drawable for HWND. Run with D3DM_NO_WINDOW=1 env. Question When CAMetalLayer/nextDrawable has no drawable (window not visible), is there a supported way for D3DMetal to skip frames instead of blocking forever? Is there an environment variable or configuration that makes D3DMetal treat "no drawable" as a recoverable condition? Does D3DM_NO_WINDOW=1 apply here, or is it only for headless/offscreen use? I understand apps should handle a nil drawable and continue, but this path is inside D3DMetal. Any guidance would help.
1
0
347
1d
MetalFX FI on iPhone: ~18 ms in our game and ~11 ms in a triangle demo — optimization guidance?
We are game developers, evaluating MetalFX frame interpolation for an iPhone game. After integrating MTLFXFrameInterpolator into our game, we observed approximately 18 ms per interpolated frame. To investigate the cost, we created a standalone rotating-triangle demo. Even this demo takes approximately 11 ms shortly after launch at 2622 x 1206 output resolution. This baseline cost is already high for our mobile frame budget, so our main question is whether there is meaningful room for optimization while keeping this output resolution and 30-to-60 presentation. The demo runs on iPhone 17 (A19), iOS 26.6, using Xcode 27.1 and a Release build. It uses standalone FI with RGBA16Float color/history/output, Depth32Float depth, and RG16Float motion vectors. The demo's depth and motion inputs are also 2622 x 1206. Presentation and synchronization are adapted from Apple's PresentThread reference. The game and demo are different workloads, so these observations are not a controlled comparison at identical GPU states. We measure the FI interval using GPU timestamp counter samples, with source-readiness and presentation-pacing waits outside the interval. It includes framework-internal synchronization and our profiling barriers; we have not yet completed a same-state comparison with those markers removed. After sustained running, the demo can rise to approximately 30 ms. A same-process trace shows a Maximum-to-Minimum GPU performance-state transition, so performance scaling may contribute to that increase. Our primary concern remains the initial approximately 11 ms cost, even during the Maximum-state period. Is this baseline cost expected on A19? Could you review whether our integration and measurement are correct, and suggest optimizations such as supported texture formats, depth/motion input sizes, or binding a temporal scaler? A recommended iPhone configuration or reference performance range would also help. We have prepared a triangle-only Xcode project with reproduction steps and timing logs, and can provide it for review.
1
0
230
3d
Metal CSM Shadows GPU Culling+Indirect+Vertex Amplification - empty some Shadow Cascades
Hi! I'm porting Cascaded Shadow Maps(CSM) from HLSL Direct3D12/Vulkan I'm using GPU Culling, indirect Call + Vertex Amplification for shadow Casters for creating Shadow map texture. Shadow map texture - Texture2DArray Shadow castes are rendering using Instancing and IndirectDraw, VS shader are checking [amplification_id]] and view mask CS shader are finding shadow caster using InstanceId and view mask for each shadow cascade: // dataIndex - Instance index for featching Instance Matrix // viewMask - for VS [[amplification_id]] uint PackInstanceData(uint viewMask, uint dataIndex) { return (viewMask << 28u) | dataIndex; } kernel void CullObjects(uint2 DTid [[thread_position_in_grid]], constant ArgBuf& argBuf [[buffer(0)]], uint groupIndex [[thread_index_in_threadgroup]]) { threadgroup atomic_uint sharedNumObjects; threadgroup uint sharedInstanceId[32]; threadgroup uint globalBufferIndex; constant auto& cbData = *argBuf.cbData; uint index = DTid.x; if (!groupIndex) { atomic_store_explicit(&sharedNumObjects, 0u, memory_order_relaxed); globalBufferIndex = 0u; } threadgroup_barrier(mem_flags::mem_threadgroup); bool visible = false; uint viewMask = 0; if (index < cbData.numInstanceObjects) { const device CullingObjInfo& cullingObjectInfo = argBuf.cullingObjInfo[index]; // find shadow casters for each shadow cascade for (uint i = 0u; i < NUM_FRUSTUMS; ++i) { constant Frustum& frustum = cbData.frustums[i]; visible = SphereInFrustum(frustum, cullingObjectInfo.boundingSphere); if (visible) { visible = BoundingBoxInFrustum(frustum, cullingObjectInfo.boundingBox.minPoint, cullingObjectInfo.boundingBox.maxPoint); } if (visible) { viewMask |= (1u << i); } } visible = bool(viewMask); } } if (visible) { uint bufferIndex = atomic_fetch_add_explicit(&sharedNumObjects, 1u, memory_order_relaxed); sharedInstanceId[bufferIndex] = PackInstanceData(viewMask, index); } threadgroup_barrier(mem_flags::mem_threadgroup); uint localNumObjects = atomic_load_explicit(&sharedNumObjects, memory_order_relaxed); if (!groupIndex && localNumObjects) { globalBufferIndex = atomic_fetch_add_explicit(&argBuf.numObjects->instanceCount, localNumObjects, memory_order_relaxed); } threadgroup_barrier(mem_flags::mem_threadgroup); if (groupIndex < localNumObjects) { argBuf.instanceId[globalBufferIndex + groupIndex] = sharedInstanceId[groupIndex]; } } VS shader for depth shadow map rendering: // struct PSOutput { float4 pos [[position]]; float2 Depth [[user(locn0)]]; float CullDist [[clip_distance]]; uint layer [[render_target_array_index]]; }; struct VS { float4 pos [[attribute(0)]]; uint InstanceId [[attribute(9)]]; }; struct b { float4x4 lightViewProj[NUM_SM_CASCADES]; }; struct Arg { const device float3x4* InstanceData [[id(0)]]; constant b* cb [[id(1)]]; }; uint UnpackViewInstancingMask(uint packedData) { return packedData >> 28u; } uint UnpackInstanceIndex(uint packedData) { return packedData & 0x0FFFFFFFu; } vertex PSOutput vsMain(VS In [[stage_in]], constant Arg& arg [[buffer(VSArgumentBufferIndex)]], ushort ampID [[amplification_id]]) { constant auto& b0 = *arg.cb; PSOutput output; float4 WorldPos = float4(In.pos.xyz, 1.0f); uint viewMask = UnpackViewInstancingMask(In.InstanceId); output.layer = ampID; // check shadow caster rendering for each cascade using ampID and viewMask if (viewMask & (1u << ampID)) { float3x4 instanceMatrix = arg.InstanceData[UnpackInstanceIndex(In.InstanceId)]; WorldPos = float4((WorldPos * instanceMatrix).xyz, 1.0f); output.CullDist = 1.0f; } else { output.pos = float4(0.0f, 0.0f, 2.0f, 1.0f); // skeep unvisible object for this cascade output.CullDist = -1.0f; return output; } output.pos = b0.lightViewProj[ampID] * WorldPos; output.Depth.xy = output.pos.zz; return output; } I have result left part: CPU Version - separate rendering for each layer for texture2Darray shadow map right part: GPU culling + Indirect + Vertex Amplification Got to fix error ? Why 1 and 3 cascade are empty? 0 cascade has duplicate data from other cascades.
2
0
714
2w
Processing RAW files using CIRAWFilter slows down after opening a few different images
Is there some cleanup required if you are opening multiple different RAW files - I am seeing a slow down each time I open a new file and things grind to a halt after opening 5 or 6 files, one after the other - even allowing some time between opening them. Currently I have a single view with a CAMetalLayer and a single Renderer and CIContent - should I be creating a new CAMetalLayer, render and CIContext each time a new image is opened ? Alternately is there some way to clear the cache or something to maintain performance.
1
0
431
2w
Clarifying Fence Behavior in Metal 4
I have some clarifying questions on how fences work in Metal 4. If two encoders A, B update the same fence and encoder C waits on that fence, is the work encoded by C guaranteed to execute after both A and B have completed, or after either A or B have completed? The following quotes in the waitForFence() documentation seem contradictory. "Encodes a command that instructs the GPU to pause before starting one or more stages of the pass until a pass updates a fence." "When encoding a pass that reuses a fence, wait for other passes to update the fence before repurposing that fence..." Are fences unsignaled when they are waited on? Specifically, can two (or more) encoders wait on a fence that is only updated in a single prior encoder? If an encoder updates a fence, no subsequent encoders in the command queue wait on the fence, and the command queue is committed, is that fence still signaled when used in a subsequent command buffer?
1
0
2.1k
3w
Native Genshin Impact Support on macOS — Feedback ID: FB25106712
Hello, I recently submitted an enhancement request through Feedback Assistant requesting official native macOS support for Genshin Impact. Feedback ID: FB25106712 I am an Apple Developer Beta tester, a Mac user, and a passionate gamer. I regularly test macOS developer beta releases and provide feedback on performance, stability, compatibility, and the overall system experience. I also genuinely enjoy testing games on Mac and seeing what Apple Silicon is capable of. I have personally tested games such as Wuthering Waves, Cyberpunk 2077, and Minecraft on Mac, and my experience with them has been extremely positive. Wuthering Waves in particular runs very smoothly and seamlessly on my Mac. These experiences have made me feel that the capability for excellent gaming on Apple Silicon is already there. The larger limitation, in my experience, is the availability of native support. As a major HoYoverse fan, Genshin Impact is one of the games I would most like to experience and test natively on macOS. At present, Mac users have to rely on third-party workarounds such as PlayCover, compatibility solutions, or cloud gaming rather than having an officially supported native client. I would really like to see Apple work with HoYoverse to bring Genshin Impact natively to macOS, properly optimized for Apple Silicon and Metal. I would also be very happy to test such a release across macOS versions and provide detailed feedback on performance, graphics, stability, thermals, controller support, and the overall gaming experience. Could this enhancement request be brought to the attention of the appropriate Apple team, or is there an appropriate channel through which I can provide further feedback regarding this request? Thank you for your time.
0
0
226
3w
64bit atomics in metal
According to the official Metal Feature Set Tables, the Apple9 family introduces full 64-bit atomics such as atomic_compare_exchange_weak_explicit() and others. However when trying to use it the compiler throws an error "no matching function call...". Example: kernel void test_64bit_cas(device atomic_ulong* atomic_var [[buffer(0)]]) { ulong expected = 0; ulong desired = 1; // This fails to compile despite Apple9 hardware supporting 64-bit buffer atomics atomic_compare_exchange_weak_explicit(atomic_var, &expected, desired, memory_order_relaxed, memory_order_relaxed); } Am i doing something wrong ?
1
0
892
3w
Contribution route for native Metal point-cloud operators used with PyTorch MPS
I maintain mps-pointops, an open-source library implementing farthest-point sampling (FPS), k-nearest-neighbor search, and radius / Ball Query operations with custom Metal kernels for PyTorch MPS tensors. The intended use is 3D point-cloud pipelines whose existing extensions depend on CUDA. I would like to prepare a contribution proposal at the appropriate layer. Apple’s PyTorch/MPS page links to the PyTorch MPS backend contribution guide and directs backend issues to PyTorch’s GitHub tracker with the module: mps label. The current implementation has numerical-contract tests for selection order, ties, radius boundaries, and padding, plus benchmark records from physical M1 and M5 Pro Macs. The repository contains the kernels, Python interfaces, tests, and measurement conditions: https://github.com/gamzerA/mps-pointops My questions are: For specialized operators such as FPS and radius search, is the recommended contribution structure to keep the operator/API in a domain library (for example, a point-cloud or graph library), while proposing only reusable backend primitives to PyTorch MPS? Is there a public Apple-maintained contribution or technical-review channel relevant to this work, beyond the PyTorch route linked from Apple’s page? For an external library that dispatches custom Metal work alongside PyTorch MPS operations, which documented extension/interoperability path should be the long-term target for command-stream ordering and buffer lifetime? I would like to minimize reliance on backend internals when preparing an upstream proposal. Pointers to the appropriate public interface, maintainers, or contribution requirements would help me scope a small, reviewable proposal. I understand that inclusion in an upstream project is subject to its maintainers’ review. Apple documentation referenced: https://developer.apple.com/metal/pytorch/
0
0
183
3w
MTL4FXTemporalDenoisedScaler initialization
I’m trying to use MTL4FXTemporalDenoisedScaler, and I’m seeing a crash during initialization even with a very simple sample app. I created a minimal sample here: https://github.com/tatsuya-ogawa/MetalFXInitExample The exception is: NSException: "-[AGXG16XFamilyHeap baseObject]: unrecognized selector sent to instance ..." What I found is: • This works: descriptor.makeTemporalDenoisedScaler(device: device) • This crashes: descriptor.makeTemporalDenoisedScaler(device: device, compiler: metal4Compiler) So the issue seems to happen only with the Metal4FX version. For testing, I’m using an iPhone 15 Pro. According to the Metal Feature Set Tables, MetalFX denoised upscaling should be supported on Apple9 and later, so I believe the device itself should meet the requirements. Reference: https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf Has anyone seen this before, or knows what might be causing it? I’d appreciate any advice. Thanks.
5
2
1.7k
Sep ’26
Metal FP32 arithmetic rounding-mode and denormal controls for deterministic shaders
Metal already provides a strong precise floating-point contract. With fast math disabled (mathMode = .safe and mathFloatingPointFunctions = .precise, or -fno-fast-math), the Metal Shading Language specification requires correctly rounded FP32 add, subtract, multiply, reciprocal, divide, sqrt, rsqrt, and fma. I am looking for clarification and, if necessary, API support for the two remaining pieces needed for portable bit-exact numerical shaders: Arithmetic rounding mode MSL §8.2 says either round-to-nearest-ties-to-even or round-toward-zero may be supported for floating-point operations. I cannot find a way to select or query the arithmetic rounding mode. The newer MTLCompileOptions.floatingPointConversionRoundingMode appears to apply only to narrowing float-to-float conversions, not arithmetic operations. Do all currently supported Apple GPU families use round-to-nearest-ties-to-even for precise FP32 add/subtract/multiply/divide/sqrt? If so, could that be made a documented guarantee? Otherwise, could Metal expose an arithmetic rounding-mode compile option and a corresponding MTLDevice capability query? Denormal behavior MSL §8.1 and §8.5 permit denormalized FP32 operands and results to be flushed to zero, including with fast math disabled. I cannot find a control or capability query for preserving denormal inputs and results. Do current Apple GPU families support denormal-preserving FP32 arithmetic? Could Metal expose a preserve/flush mode and a MTLDevice query? A convenient end state would be a queryable strict FP32 configuration combining: safe math; precise FP32 functions; contraction disabled when separate rounding points are required; round-to-nearest-ties-to-even arithmetic; preserved FP32 denormal inputs and results; defined signed-zero, infinity, and NaN behavior. My use case is deterministic GPU numerical simulation. A small compute-shader probe can identify effective rounding and denormal behavior on one GPU/OS/compiler combination, but it cannot provide the portable or future-proof contract needed by applications and higher-level APIs such as WebGPU. Related cross-API work: SPIR-V/Vulkan: https://github.com/KhronosGroup/SPIRV-Registry/issues/448 HLSL/DXIL/D3D12: https://github.com/microsoft/hlsl-specs/issues/926 WebGPU/WGSL umbrella issue: https://github.com/gpuweb/gpuweb/issues/2259 Relevant Metal documentation: Metal Shading Language Specification, §§1.6.3 and 8.1–8.5: https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf MTLCompileOptions.mathMode: https://developer.apple.com/documentation/metal/mtlcompileoptions/mathmode MTLCompileOptions.floatingPointConversionRoundingMode: https://developer.apple.com/documentation/metal/mtlcompileoptions/floatingpointconversionroundingmode
2
0
3.2k
Sep ’26
M5 Pro external 5K 165Hz display: Window animations and scrolling UI appear to render at ~60Hz/jitter while cursor remains perfectly smooth
Hello Apple engineers, I’m trying to determine whether what I’m seeing is expected behavior or a software issue with the new M5 Pro platform. System MacBook Pro (M5 Pro) Latest macOS Beta External 5K 165Hz monitor connected via DisplayPort Refresh rate correctly detected as 165Hz What I observe The display itself is clearly running at 165Hz. For example: Mouse cursor movement is extremely smooth. Dragging the desktop by holding an empty area is also perfectly smooth. However: Moving application windows feels much closer to 60Hz. Scrolling in Safari, Chrome and other applications also appears to run at a much lower frame rate than the display refresh rate. Mission Control animations sometimes show similar micro-stutters. This makes the cursor and desktop movement noticeably smoother than normal window animations. ⸻ Troubleshooting already performed Different DisplayPort cables Different timing configurations Different resolutions / HiDPI modes DSC enabled and disabled Refresh rate confirmed at 165Hz Same behavior across multiple applications The issue appears unrelated to the monitor itself because the cursor is clearly rendered at the full refresh rate. ⸻ Additional observation Interestingly, I previously used another external 4K 144Hz HDR monitor and did not notice this behavior. I also found another M5 Pro user reporting nearly the same issue: external 165Hz display smooth cursor window dragging jitter / micro-stuttering At the same time, I haven’t found similar reports from M4 Pro or the base M5 running the same monitor. ⸻ My question Could this be related to the new M5 Pro display pipeline (WindowServer, Display Engine, or DCP)? Is there any known issue regarding high-refresh-rate external displays on the M5 Pro platform? Or is there additional diagnostic logging (WindowServer, DCP, Metal, etc.) that would help identify whether frames are actually being presented at the display refresh rate? I’d be happy to provide: sysdiagnose WindowServer logs Screen recordings Display timing information IORegistry dumps if they would be helpful. Thank you!
5
1
2.1k
Sep ’26
Data corruption when using MTLBlitCommandEncoder.copy from buffer to texture
When we use MTLBlitCommandEncoder.copy to copy from buffer to textures with specific size, we find that the texture content is corruppt. The specific rules we found are: On Apple GPUs (reproduced on Apple M5, macOS 26.3 and iPhone 13 / A15), for ASTC textures uploaded to private storage via the blit encoder, the sampler / texel-fetch mis-addresses the tail of a mip level when ALL of the following hold: that level's texel width is an exact multiple of the 16 KB page width (32 blocks) — 128 texels for ASTC 4x4, 192 for 6x6, 256 for 8x8; the texture width is not divisible by 2^level (a partial block column exists in the chain); that level is taller than one page (>32 block rows). The level's first 128 texel rows (first page row, 4x4) are always read correctly; everything beyond is mis-addressed. Reads landing on unmapped/invalid memory decode as opaque magenta (1,0,1); reads aliasing valid memory show wrong image content. This produced visible purple artifacts in a shipping game (impostor tree atlas, 514x1024 ASTC 4x4, mip2 bottom half). The minimum reproduce code is: ASTCMipTailBugRepro.swift The output is: 514x1024 blit(private) mip2: CORRUPTED (bottom of mip2 mis-addressed) 512x1024 blit(private) mip2: INTACT 514x1024 replaceRegion mip2: INTACT
1
0
1.6k
Sep ’26
Metal 4 and object lifetime
I have a metal kit view and drain the draw method of its delegate like shown below. Let's say I have one or more MTLBuffers with vertex resources bound via the argument table. When is it ok to drop these buffers? As far as I know one cannot schedule a completion handler In Metal 4 and I haven't been able to find any documentation about the lifetime requirements here. Any pointers/ideas appreciated. class RenderCoordinator: NSObject, MTKViewDelegate { public func draw(in view: MTKView) { let commandAllocator: any MTL4CommandAllocator = ... let commandBuffer: any MTL4CommandBuffer = ... let commandQueue: any MTL4CommandQueue = ... guard let drawable = view.currentDrawable else { return } commandBuffer.beginCommandBuffer(allocator: commandAllocator) let state: any MTLRenderPipelineState = ... let encoder: any MTL4RenderCommandEncoder = ... let argTable: any MTL4ArgumentTable = ... encoder.setRenderPipelineState(state) encoder.setArgumentTable(argTable, stages: .vertex) commandBuffer.endCommandBuffer() commandQueue.waitForDrawable(drawable) commandQueue.commit([commandBuffer]) commandQueue.signalDrawable(drawable) drawable.present() } }
2
0
1.6k
Sep ’26
Behaviour of a 0-value `accelerationStructureID`
When constructing a MTLIndirectAccelerationStructureInstanceDescriptor, one specifies an accelerationStructureID. In the equivalents in both Vulkan and DirectX12, one can set it to zero to be an "inactive" instance. However, on Metal, this field does not appear to have any documentation, and thus it is difficult to figure out if there is similar behavior (and no other Metal documentation seems to mention this). Does setting this field to 0 (i.e. null) disable the instance? If not, is there any other way to have an equivalent effect?
5
0
2.6k
Aug ’26
Stress-testing Metal compute pipelines using Autolykos workload characteristics
I have been using a small macOS research project to exercise Metal with a workload that differs from rendering and dense machine-learning kernels. Autolykos v2 is useful for this because it combines a large, height-dependent working set with pseudo-random reads, integer-heavy hashing, sustained execution, and periodic replacement of the dataset. The project happens to be a miner, but my question here is strictly about Metal compute behaviour. On the Apple M4 system used for these measurements, the full dataset contained 216,430,305 elements of 32 bytes each: 6.93 GB, or about 6.45 GiB, held in a .storageModePrivate buffer. For every nonce, the search kernel: computes an index seed, performs 32 pseudo-random 32-byte dataset reads, accumulates the eight 32-bit limbs into a 256-bit sum, and applies a final BLAKE2b compression and target comparison. The normal dispatch uses 128 threads per threadgroup. The wider pipeline also builds the next height's dataset in chunks on a separate command queue while search continues, and keeps two search command buffers in flight. I record command-buffer wall time, gpuStartTime/gpuEndTime, unions of overlapping intervals, and thermal state. To estimate the ceiling imposed by the random gathers, I added a non-consensus microbenchmark. It retains the normal seed calculation, index distribution, all 32 dataset reads, and the complete accumulation, but omits the final BLAKE2b compression. The accumulated result remains observable through a comparison, so the gather loop cannot simply disappear. I expected this stripped kernel to be at least slightly faster. Instead, an order-balanced campaign on an M4 produced: complete search kernel: 3.108 million nonces/s median active throughput gather-only kernel: 2.952 million nonces/s ratio: 105.3% All four same-round ratios were between 103.18% and 105.74%. Each measured run used the full dataset, a 30-second search interval, an excluded warm-up, and a start-temperature gate below 50 °C. Both compute pipeline states reported maxTotalThreadsPerThreadgroup == 1024. My conservative conclusion was not to pursue speculative register-pressure or manual memory-level-parallelism rewrites. The access pattern appears sufficiently dominant, while the supposedly simpler microbenchmark may have changed the compiled pipeline in a way that makes it a poor upper-bound model. My questions are: Can removing the trailing arithmetic legitimately make a memory-latency-heavy Metal kernel slower by changing register allocation, instruction scheduling, or the amount of useful latency hiding? Or would you first suspect a flaw in this kind of gather-only benchmark construction? Also, which Metal GPU counters are the most reliable way to distinguish memory-latency saturation from register-limited occupancy in a long-running compute kernel? I am looking at compute occupancy, buffer and ALU limiters, bandwidth, and cache behaviour, but maxTotalThreadsPerThreadgroup alone is clearly too coarse to explain the result. This is one hardware-specific observation rather than a general claim about Apple GPUs. If useful, I can reduce the workload to a smaller standalone reproducer. The source code, benchmark driver, and complete campaign report are available here: https://github.com/giffeler/ergometal The detailed measurements and validation procedure for this comparison are documented here: https://github.com/giffeler/ergometal/blob/main/Benchmarks/2026-08-15-search-gather-ceiling-ab.md
0
0
786
Aug ’26
Metal-cpp usability issue with MTL::Buffer and MTL::ResidencySet
I know this might be a peeve of mine, but looking into programming a simple Metal4 Compute Shader example, essentially updating the Performing Calculations on a GPU example code to work with CPP and Metal4. I found that MTL::Allocation and MTL::Buffer pointers can't be used interchangeably when you are trying to add allocations to a MTL::ResidencySet, this is forcing you to: Know by heart that they inherit from each other and that you can just cast them (this is a bit suspicious though, it did work for me). Forcefully either C-cast or reinterpret_cast the MTL::Buffer pointer to a MTL::Allocation pointer as the MTL::ResidencySet will only accept that type. I might as well just be plain wrong about how this is used, any tips on correct usage in that case? Is there any expectation to either provide a typecast operator or add inheritance to support the expected behaviour seen in Swift and ObjC, which is just passing the thing? Opened a report with # FB24534953 with some extra information. Bear in mind that the example code uses Premake5, but it can generate an Xcode solution easily.
0
0
765
Aug ’26
Xcode 27.0 b5, macOS 26.6.1, Metal build fails: symbols not found for air64_v28
I've just downloaded the Xcode 27.0 beta 5 on a macOS 26.6.1 machine and tried to build my app (which includes Metal CoreImage kernels). I'm met with a new (to me) error; /Users/…/Developer/…/air-lld:1:1 symbol(s) not found for target 'air64_v28-apple-macosx26.0.0' and from the build log; air-lld: warning: ignoring file '/Applications/Xcode-beta.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk/System/Library/Frameworks/CoreImage.framework/Versions/A/CoreImage.metallib', file AIR version (2.9) is bigger than the one of the target being linked (2.8) air-lld: error: symbol(s) not found for target 'air64_v28-apple-macosx26.0.0' metal: error: air-lld command failed with exit code 1 (use -v to see invocation) I have no MTL_LANGUAGE_REVISION in my build settings. If I add one, with value Metal41 the app builds fine, but crashes at runtime as Metal 4.1 isn't supported on macOS 26. I imagine this is a beta Xcode and or macOS SDK bug, is there a workaround?
1
1
1.3k
Aug ’26
MPS bf16 softmax produces NaN on M5 Max (regression from M4) — breaks all on-device diffusion inference
Metal Performance Shaders produces from bf16/fp16 softmax for large attention tensors on , forcing the entire local generative-AI ecosystem to fall back to fp32. This is a . Environment Minimal reproduction Decomposed softmax on MPS — diffs = x - maxes produces NaN even though every element is identical (result should be all zeros): Real-world impact Running ComfyUI (the dominant local generative-AI UI) on M5 Max: Root cause (per PyTorch MPS maintainers) PyTorch maintainers (@drisspg, @albanD) have traced this to in the MPS/MPSGraph fused kernels. The NaN originates in the x - maxes subtraction inside softmax, then propagates through the attention block and the entire network. PyTorch cannot fix this — it is in the Metal/MPS kernel layer. Not isolated The M4-era fix ("fixed on macOS 15.1") did not survive onto M5, indicating the MPS fused-kernel precision fix was either reverted or not ported to the M5 GPU architecture. Request
1
0
1.1k
Aug ’26
Metal Shader Converter thread safety
Hello Apple! We've got offline shader compilation from HLSL -> Metallib using DXC -> SPIR-V -> metal.exe. This works okay for the most part, but it requires the creation of intermediate files to pass to/from the metal.exe process and we've had some issues with metal.exe sometimes not launching (probably our fault). Then we noticed Metal Shader Converter (MSC) exists and has a DLL - this looks way better since there's no need to launch processes or store intermediate files. However, upon trying to replace metal.exe with it I quickly ran into rampant heap corruption. I was surprised because the docs claim this: Each thread in your program needs to create its own instance of IRCompiler to avoid race conditions. But once I start calling IRCompilerAllocCompileAndLink in parallel all hell breaks loose, whether or not each thread has its own IRCompiler. I figured I must be doing something wrong, so I removed my attempt and compiled DXC locally with the MSC integration and encountered the exact same heap corruption. So I'm inclined to think the library isn't actually thread safe, but I'm wondering if there's something I'm missing? I tried all 3 versions of MSC just in case it was a problem with 3.0, but I got the same result each time. The only way to make it work was to surround compilation with a mutex, which makes its use pointless in our case.
2
0
1.5k
Aug ’26
D3DMetal: present hangs permanently when the game window becomes occluded
Environment macOS 27.0.1 (26A434), Apple M4 Pro Apple Game Porting Toolkit 4.0b2 (D3DMetal.framework) via CrossOver 27 D3D12 game, borderless fullscreen window (2540x1650) Symptom When the game window becomes occluded — Mission Control, a system OSD (volume/brightness), or another window being brought to the front — the game never renders another frame. It does not crash: GPU utilization drops to 0%, the process stays alive with near-zero CPU, and it never recovers. Killing and restarting the game is the only way out. D3DMetal log at the moment it happens: [D3DMetal:FATAL:%llX][%s:%d] Failed to get drawable for HWND. Run with D3DM_NO_WINDOW=1 env. Question When CAMetalLayer/nextDrawable has no drawable (window not visible), is there a supported way for D3DMetal to skip frames instead of blocking forever? Is there an environment variable or configuration that makes D3DMetal treat "no drawable" as a recoverable condition? Does D3DM_NO_WINDOW=1 apply here, or is it only for headless/offscreen use? I understand apps should handle a nil drawable and continue, but this path is inside D3DMetal. Any guidance would help.
Replies
1
Boosts
0
Views
347
Activity
1d
MetalFX FI on iPhone: ~18 ms in our game and ~11 ms in a triangle demo — optimization guidance?
We are game developers, evaluating MetalFX frame interpolation for an iPhone game. After integrating MTLFXFrameInterpolator into our game, we observed approximately 18 ms per interpolated frame. To investigate the cost, we created a standalone rotating-triangle demo. Even this demo takes approximately 11 ms shortly after launch at 2622 x 1206 output resolution. This baseline cost is already high for our mobile frame budget, so our main question is whether there is meaningful room for optimization while keeping this output resolution and 30-to-60 presentation. The demo runs on iPhone 17 (A19), iOS 26.6, using Xcode 27.1 and a Release build. It uses standalone FI with RGBA16Float color/history/output, Depth32Float depth, and RG16Float motion vectors. The demo's depth and motion inputs are also 2622 x 1206. Presentation and synchronization are adapted from Apple's PresentThread reference. The game and demo are different workloads, so these observations are not a controlled comparison at identical GPU states. We measure the FI interval using GPU timestamp counter samples, with source-readiness and presentation-pacing waits outside the interval. It includes framework-internal synchronization and our profiling barriers; we have not yet completed a same-state comparison with those markers removed. After sustained running, the demo can rise to approximately 30 ms. A same-process trace shows a Maximum-to-Minimum GPU performance-state transition, so performance scaling may contribute to that increase. Our primary concern remains the initial approximately 11 ms cost, even during the Maximum-state period. Is this baseline cost expected on A19? Could you review whether our integration and measurement are correct, and suggest optimizations such as supported texture formats, depth/motion input sizes, or binding a temporal scaler? A recommended iPhone configuration or reference performance range would also help. We have prepared a triangle-only Xcode project with reproduction steps and timing logs, and can provide it for review.
Replies
1
Boosts
0
Views
230
Activity
3d
Metal CSM Shadows GPU Culling+Indirect+Vertex Amplification - empty some Shadow Cascades
Hi! I'm porting Cascaded Shadow Maps(CSM) from HLSL Direct3D12/Vulkan I'm using GPU Culling, indirect Call + Vertex Amplification for shadow Casters for creating Shadow map texture. Shadow map texture - Texture2DArray Shadow castes are rendering using Instancing and IndirectDraw, VS shader are checking [amplification_id]] and view mask CS shader are finding shadow caster using InstanceId and view mask for each shadow cascade: // dataIndex - Instance index for featching Instance Matrix // viewMask - for VS [[amplification_id]] uint PackInstanceData(uint viewMask, uint dataIndex) { return (viewMask << 28u) | dataIndex; } kernel void CullObjects(uint2 DTid [[thread_position_in_grid]], constant ArgBuf& argBuf [[buffer(0)]], uint groupIndex [[thread_index_in_threadgroup]]) { threadgroup atomic_uint sharedNumObjects; threadgroup uint sharedInstanceId[32]; threadgroup uint globalBufferIndex; constant auto& cbData = *argBuf.cbData; uint index = DTid.x; if (!groupIndex) { atomic_store_explicit(&sharedNumObjects, 0u, memory_order_relaxed); globalBufferIndex = 0u; } threadgroup_barrier(mem_flags::mem_threadgroup); bool visible = false; uint viewMask = 0; if (index < cbData.numInstanceObjects) { const device CullingObjInfo& cullingObjectInfo = argBuf.cullingObjInfo[index]; // find shadow casters for each shadow cascade for (uint i = 0u; i < NUM_FRUSTUMS; ++i) { constant Frustum& frustum = cbData.frustums[i]; visible = SphereInFrustum(frustum, cullingObjectInfo.boundingSphere); if (visible) { visible = BoundingBoxInFrustum(frustum, cullingObjectInfo.boundingBox.minPoint, cullingObjectInfo.boundingBox.maxPoint); } if (visible) { viewMask |= (1u << i); } } visible = bool(viewMask); } } if (visible) { uint bufferIndex = atomic_fetch_add_explicit(&sharedNumObjects, 1u, memory_order_relaxed); sharedInstanceId[bufferIndex] = PackInstanceData(viewMask, index); } threadgroup_barrier(mem_flags::mem_threadgroup); uint localNumObjects = atomic_load_explicit(&sharedNumObjects, memory_order_relaxed); if (!groupIndex && localNumObjects) { globalBufferIndex = atomic_fetch_add_explicit(&argBuf.numObjects->instanceCount, localNumObjects, memory_order_relaxed); } threadgroup_barrier(mem_flags::mem_threadgroup); if (groupIndex < localNumObjects) { argBuf.instanceId[globalBufferIndex + groupIndex] = sharedInstanceId[groupIndex]; } } VS shader for depth shadow map rendering: // struct PSOutput { float4 pos [[position]]; float2 Depth [[user(locn0)]]; float CullDist [[clip_distance]]; uint layer [[render_target_array_index]]; }; struct VS { float4 pos [[attribute(0)]]; uint InstanceId [[attribute(9)]]; }; struct b { float4x4 lightViewProj[NUM_SM_CASCADES]; }; struct Arg { const device float3x4* InstanceData [[id(0)]]; constant b* cb [[id(1)]]; }; uint UnpackViewInstancingMask(uint packedData) { return packedData >> 28u; } uint UnpackInstanceIndex(uint packedData) { return packedData & 0x0FFFFFFFu; } vertex PSOutput vsMain(VS In [[stage_in]], constant Arg& arg [[buffer(VSArgumentBufferIndex)]], ushort ampID [[amplification_id]]) { constant auto& b0 = *arg.cb; PSOutput output; float4 WorldPos = float4(In.pos.xyz, 1.0f); uint viewMask = UnpackViewInstancingMask(In.InstanceId); output.layer = ampID; // check shadow caster rendering for each cascade using ampID and viewMask if (viewMask & (1u << ampID)) { float3x4 instanceMatrix = arg.InstanceData[UnpackInstanceIndex(In.InstanceId)]; WorldPos = float4((WorldPos * instanceMatrix).xyz, 1.0f); output.CullDist = 1.0f; } else { output.pos = float4(0.0f, 0.0f, 2.0f, 1.0f); // skeep unvisible object for this cascade output.CullDist = -1.0f; return output; } output.pos = b0.lightViewProj[ampID] * WorldPos; output.Depth.xy = output.pos.zz; return output; } I have result left part: CPU Version - separate rendering for each layer for texture2Darray shadow map right part: GPU culling + Indirect + Vertex Amplification Got to fix error ? Why 1 and 3 cascade are empty? 0 cascade has duplicate data from other cascades.
Replies
2
Boosts
0
Views
714
Activity
2w
Processing RAW files using CIRAWFilter slows down after opening a few different images
Is there some cleanup required if you are opening multiple different RAW files - I am seeing a slow down each time I open a new file and things grind to a halt after opening 5 or 6 files, one after the other - even allowing some time between opening them. Currently I have a single view with a CAMetalLayer and a single Renderer and CIContent - should I be creating a new CAMetalLayer, render and CIContext each time a new image is opened ? Alternately is there some way to clear the cache or something to maintain performance.
Replies
1
Boosts
0
Views
431
Activity
2w
Clarifying Fence Behavior in Metal 4
I have some clarifying questions on how fences work in Metal 4. If two encoders A, B update the same fence and encoder C waits on that fence, is the work encoded by C guaranteed to execute after both A and B have completed, or after either A or B have completed? The following quotes in the waitForFence() documentation seem contradictory. "Encodes a command that instructs the GPU to pause before starting one or more stages of the pass until a pass updates a fence." "When encoding a pass that reuses a fence, wait for other passes to update the fence before repurposing that fence..." Are fences unsignaled when they are waited on? Specifically, can two (or more) encoders wait on a fence that is only updated in a single prior encoder? If an encoder updates a fence, no subsequent encoders in the command queue wait on the fence, and the command queue is committed, is that fence still signaled when used in a subsequent command buffer?
Replies
1
Boosts
0
Views
2.1k
Activity
3w
Native Genshin Impact Support on macOS — Feedback ID: FB25106712
Hello, I recently submitted an enhancement request through Feedback Assistant requesting official native macOS support for Genshin Impact. Feedback ID: FB25106712 I am an Apple Developer Beta tester, a Mac user, and a passionate gamer. I regularly test macOS developer beta releases and provide feedback on performance, stability, compatibility, and the overall system experience. I also genuinely enjoy testing games on Mac and seeing what Apple Silicon is capable of. I have personally tested games such as Wuthering Waves, Cyberpunk 2077, and Minecraft on Mac, and my experience with them has been extremely positive. Wuthering Waves in particular runs very smoothly and seamlessly on my Mac. These experiences have made me feel that the capability for excellent gaming on Apple Silicon is already there. The larger limitation, in my experience, is the availability of native support. As a major HoYoverse fan, Genshin Impact is one of the games I would most like to experience and test natively on macOS. At present, Mac users have to rely on third-party workarounds such as PlayCover, compatibility solutions, or cloud gaming rather than having an officially supported native client. I would really like to see Apple work with HoYoverse to bring Genshin Impact natively to macOS, properly optimized for Apple Silicon and Metal. I would also be very happy to test such a release across macOS versions and provide detailed feedback on performance, graphics, stability, thermals, controller support, and the overall gaming experience. Could this enhancement request be brought to the attention of the appropriate Apple team, or is there an appropriate channel through which I can provide further feedback regarding this request? Thank you for your time.
Replies
0
Boosts
0
Views
226
Activity
3w
64bit atomics in metal
According to the official Metal Feature Set Tables, the Apple9 family introduces full 64-bit atomics such as atomic_compare_exchange_weak_explicit() and others. However when trying to use it the compiler throws an error "no matching function call...". Example: kernel void test_64bit_cas(device atomic_ulong* atomic_var [[buffer(0)]]) { ulong expected = 0; ulong desired = 1; // This fails to compile despite Apple9 hardware supporting 64-bit buffer atomics atomic_compare_exchange_weak_explicit(atomic_var, &expected, desired, memory_order_relaxed, memory_order_relaxed); } Am i doing something wrong ?
Replies
1
Boosts
0
Views
892
Activity
3w
Contribution route for native Metal point-cloud operators used with PyTorch MPS
I maintain mps-pointops, an open-source library implementing farthest-point sampling (FPS), k-nearest-neighbor search, and radius / Ball Query operations with custom Metal kernels for PyTorch MPS tensors. The intended use is 3D point-cloud pipelines whose existing extensions depend on CUDA. I would like to prepare a contribution proposal at the appropriate layer. Apple’s PyTorch/MPS page links to the PyTorch MPS backend contribution guide and directs backend issues to PyTorch’s GitHub tracker with the module: mps label. The current implementation has numerical-contract tests for selection order, ties, radius boundaries, and padding, plus benchmark records from physical M1 and M5 Pro Macs. The repository contains the kernels, Python interfaces, tests, and measurement conditions: https://github.com/gamzerA/mps-pointops My questions are: For specialized operators such as FPS and radius search, is the recommended contribution structure to keep the operator/API in a domain library (for example, a point-cloud or graph library), while proposing only reusable backend primitives to PyTorch MPS? Is there a public Apple-maintained contribution or technical-review channel relevant to this work, beyond the PyTorch route linked from Apple’s page? For an external library that dispatches custom Metal work alongside PyTorch MPS operations, which documented extension/interoperability path should be the long-term target for command-stream ordering and buffer lifetime? I would like to minimize reliance on backend internals when preparing an upstream proposal. Pointers to the appropriate public interface, maintainers, or contribution requirements would help me scope a small, reviewable proposal. I understand that inclusion in an upstream project is subject to its maintainers’ review. Apple documentation referenced: https://developer.apple.com/metal/pytorch/
Replies
0
Boosts
0
Views
183
Activity
3w
MTL4FXTemporalDenoisedScaler initialization
I’m trying to use MTL4FXTemporalDenoisedScaler, and I’m seeing a crash during initialization even with a very simple sample app. I created a minimal sample here: https://github.com/tatsuya-ogawa/MetalFXInitExample The exception is: NSException: "-[AGXG16XFamilyHeap baseObject]: unrecognized selector sent to instance ..." What I found is: • This works: descriptor.makeTemporalDenoisedScaler(device: device) • This crashes: descriptor.makeTemporalDenoisedScaler(device: device, compiler: metal4Compiler) So the issue seems to happen only with the Metal4FX version. For testing, I’m using an iPhone 15 Pro. According to the Metal Feature Set Tables, MetalFX denoised upscaling should be supported on Apple9 and later, so I believe the device itself should meet the requirements. Reference: https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf Has anyone seen this before, or knows what might be causing it? I’d appreciate any advice. Thanks.
Replies
5
Boosts
2
Views
1.7k
Activity
Sep ’26
Metal FP32 arithmetic rounding-mode and denormal controls for deterministic shaders
Metal already provides a strong precise floating-point contract. With fast math disabled (mathMode = .safe and mathFloatingPointFunctions = .precise, or -fno-fast-math), the Metal Shading Language specification requires correctly rounded FP32 add, subtract, multiply, reciprocal, divide, sqrt, rsqrt, and fma. I am looking for clarification and, if necessary, API support for the two remaining pieces needed for portable bit-exact numerical shaders: Arithmetic rounding mode MSL §8.2 says either round-to-nearest-ties-to-even or round-toward-zero may be supported for floating-point operations. I cannot find a way to select or query the arithmetic rounding mode. The newer MTLCompileOptions.floatingPointConversionRoundingMode appears to apply only to narrowing float-to-float conversions, not arithmetic operations. Do all currently supported Apple GPU families use round-to-nearest-ties-to-even for precise FP32 add/subtract/multiply/divide/sqrt? If so, could that be made a documented guarantee? Otherwise, could Metal expose an arithmetic rounding-mode compile option and a corresponding MTLDevice capability query? Denormal behavior MSL §8.1 and §8.5 permit denormalized FP32 operands and results to be flushed to zero, including with fast math disabled. I cannot find a control or capability query for preserving denormal inputs and results. Do current Apple GPU families support denormal-preserving FP32 arithmetic? Could Metal expose a preserve/flush mode and a MTLDevice query? A convenient end state would be a queryable strict FP32 configuration combining: safe math; precise FP32 functions; contraction disabled when separate rounding points are required; round-to-nearest-ties-to-even arithmetic; preserved FP32 denormal inputs and results; defined signed-zero, infinity, and NaN behavior. My use case is deterministic GPU numerical simulation. A small compute-shader probe can identify effective rounding and denormal behavior on one GPU/OS/compiler combination, but it cannot provide the portable or future-proof contract needed by applications and higher-level APIs such as WebGPU. Related cross-API work: SPIR-V/Vulkan: https://github.com/KhronosGroup/SPIRV-Registry/issues/448 HLSL/DXIL/D3D12: https://github.com/microsoft/hlsl-specs/issues/926 WebGPU/WGSL umbrella issue: https://github.com/gpuweb/gpuweb/issues/2259 Relevant Metal documentation: Metal Shading Language Specification, §§1.6.3 and 8.1–8.5: https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf MTLCompileOptions.mathMode: https://developer.apple.com/documentation/metal/mtlcompileoptions/mathmode MTLCompileOptions.floatingPointConversionRoundingMode: https://developer.apple.com/documentation/metal/mtlcompileoptions/floatingpointconversionroundingmode
Replies
2
Boosts
0
Views
3.2k
Activity
Sep ’26
M5 Pro external 5K 165Hz display: Window animations and scrolling UI appear to render at ~60Hz/jitter while cursor remains perfectly smooth
Hello Apple engineers, I’m trying to determine whether what I’m seeing is expected behavior or a software issue with the new M5 Pro platform. System MacBook Pro (M5 Pro) Latest macOS Beta External 5K 165Hz monitor connected via DisplayPort Refresh rate correctly detected as 165Hz What I observe The display itself is clearly running at 165Hz. For example: Mouse cursor movement is extremely smooth. Dragging the desktop by holding an empty area is also perfectly smooth. However: Moving application windows feels much closer to 60Hz. Scrolling in Safari, Chrome and other applications also appears to run at a much lower frame rate than the display refresh rate. Mission Control animations sometimes show similar micro-stutters. This makes the cursor and desktop movement noticeably smoother than normal window animations. ⸻ Troubleshooting already performed Different DisplayPort cables Different timing configurations Different resolutions / HiDPI modes DSC enabled and disabled Refresh rate confirmed at 165Hz Same behavior across multiple applications The issue appears unrelated to the monitor itself because the cursor is clearly rendered at the full refresh rate. ⸻ Additional observation Interestingly, I previously used another external 4K 144Hz HDR monitor and did not notice this behavior. I also found another M5 Pro user reporting nearly the same issue: external 165Hz display smooth cursor window dragging jitter / micro-stuttering At the same time, I haven’t found similar reports from M4 Pro or the base M5 running the same monitor. ⸻ My question Could this be related to the new M5 Pro display pipeline (WindowServer, Display Engine, or DCP)? Is there any known issue regarding high-refresh-rate external displays on the M5 Pro platform? Or is there additional diagnostic logging (WindowServer, DCP, Metal, etc.) that would help identify whether frames are actually being presented at the display refresh rate? I’d be happy to provide: sysdiagnose WindowServer logs Screen recordings Display timing information IORegistry dumps if they would be helpful. Thank you!
Replies
5
Boosts
1
Views
2.1k
Activity
Sep ’26
Data corruption when using MTLBlitCommandEncoder.copy from buffer to texture
When we use MTLBlitCommandEncoder.copy to copy from buffer to textures with specific size, we find that the texture content is corruppt. The specific rules we found are: On Apple GPUs (reproduced on Apple M5, macOS 26.3 and iPhone 13 / A15), for ASTC textures uploaded to private storage via the blit encoder, the sampler / texel-fetch mis-addresses the tail of a mip level when ALL of the following hold: that level's texel width is an exact multiple of the 16 KB page width (32 blocks) — 128 texels for ASTC 4x4, 192 for 6x6, 256 for 8x8; the texture width is not divisible by 2^level (a partial block column exists in the chain); that level is taller than one page (>32 block rows). The level's first 128 texel rows (first page row, 4x4) are always read correctly; everything beyond is mis-addressed. Reads landing on unmapped/invalid memory decode as opaque magenta (1,0,1); reads aliasing valid memory show wrong image content. This produced visible purple artifacts in a shipping game (impostor tree atlas, 514x1024 ASTC 4x4, mip2 bottom half). The minimum reproduce code is: ASTCMipTailBugRepro.swift The output is: 514x1024 blit(private) mip2: CORRUPTED (bottom of mip2 mis-addressed) 512x1024 blit(private) mip2: INTACT 514x1024 replaceRegion mip2: INTACT
Replies
1
Boosts
0
Views
1.6k
Activity
Sep ’26
Metal 4 and object lifetime
I have a metal kit view and drain the draw method of its delegate like shown below. Let's say I have one or more MTLBuffers with vertex resources bound via the argument table. When is it ok to drop these buffers? As far as I know one cannot schedule a completion handler In Metal 4 and I haven't been able to find any documentation about the lifetime requirements here. Any pointers/ideas appreciated. class RenderCoordinator: NSObject, MTKViewDelegate { public func draw(in view: MTKView) { let commandAllocator: any MTL4CommandAllocator = ... let commandBuffer: any MTL4CommandBuffer = ... let commandQueue: any MTL4CommandQueue = ... guard let drawable = view.currentDrawable else { return } commandBuffer.beginCommandBuffer(allocator: commandAllocator) let state: any MTLRenderPipelineState = ... let encoder: any MTL4RenderCommandEncoder = ... let argTable: any MTL4ArgumentTable = ... encoder.setRenderPipelineState(state) encoder.setArgumentTable(argTable, stages: .vertex) commandBuffer.endCommandBuffer() commandQueue.waitForDrawable(drawable) commandQueue.commit([commandBuffer]) commandQueue.signalDrawable(drawable) drawable.present() } }
Replies
2
Boosts
0
Views
1.6k
Activity
Sep ’26
Behaviour of a 0-value `accelerationStructureID`
When constructing a MTLIndirectAccelerationStructureInstanceDescriptor, one specifies an accelerationStructureID. In the equivalents in both Vulkan and DirectX12, one can set it to zero to be an "inactive" instance. However, on Metal, this field does not appear to have any documentation, and thus it is difficult to figure out if there is similar behavior (and no other Metal documentation seems to mention this). Does setting this field to 0 (i.e. null) disable the instance? If not, is there any other way to have an equivalent effect?
Replies
5
Boosts
0
Views
2.6k
Activity
Aug ’26
Stress-testing Metal compute pipelines using Autolykos workload characteristics
I have been using a small macOS research project to exercise Metal with a workload that differs from rendering and dense machine-learning kernels. Autolykos v2 is useful for this because it combines a large, height-dependent working set with pseudo-random reads, integer-heavy hashing, sustained execution, and periodic replacement of the dataset. The project happens to be a miner, but my question here is strictly about Metal compute behaviour. On the Apple M4 system used for these measurements, the full dataset contained 216,430,305 elements of 32 bytes each: 6.93 GB, or about 6.45 GiB, held in a .storageModePrivate buffer. For every nonce, the search kernel: computes an index seed, performs 32 pseudo-random 32-byte dataset reads, accumulates the eight 32-bit limbs into a 256-bit sum, and applies a final BLAKE2b compression and target comparison. The normal dispatch uses 128 threads per threadgroup. The wider pipeline also builds the next height's dataset in chunks on a separate command queue while search continues, and keeps two search command buffers in flight. I record command-buffer wall time, gpuStartTime/gpuEndTime, unions of overlapping intervals, and thermal state. To estimate the ceiling imposed by the random gathers, I added a non-consensus microbenchmark. It retains the normal seed calculation, index distribution, all 32 dataset reads, and the complete accumulation, but omits the final BLAKE2b compression. The accumulated result remains observable through a comparison, so the gather loop cannot simply disappear. I expected this stripped kernel to be at least slightly faster. Instead, an order-balanced campaign on an M4 produced: complete search kernel: 3.108 million nonces/s median active throughput gather-only kernel: 2.952 million nonces/s ratio: 105.3% All four same-round ratios were between 103.18% and 105.74%. Each measured run used the full dataset, a 30-second search interval, an excluded warm-up, and a start-temperature gate below 50 °C. Both compute pipeline states reported maxTotalThreadsPerThreadgroup == 1024. My conservative conclusion was not to pursue speculative register-pressure or manual memory-level-parallelism rewrites. The access pattern appears sufficiently dominant, while the supposedly simpler microbenchmark may have changed the compiled pipeline in a way that makes it a poor upper-bound model. My questions are: Can removing the trailing arithmetic legitimately make a memory-latency-heavy Metal kernel slower by changing register allocation, instruction scheduling, or the amount of useful latency hiding? Or would you first suspect a flaw in this kind of gather-only benchmark construction? Also, which Metal GPU counters are the most reliable way to distinguish memory-latency saturation from register-limited occupancy in a long-running compute kernel? I am looking at compute occupancy, buffer and ALU limiters, bandwidth, and cache behaviour, but maxTotalThreadsPerThreadgroup alone is clearly too coarse to explain the result. This is one hardware-specific observation rather than a general claim about Apple GPUs. If useful, I can reduce the workload to a smaller standalone reproducer. The source code, benchmark driver, and complete campaign report are available here: https://github.com/giffeler/ergometal The detailed measurements and validation procedure for this comparison are documented here: https://github.com/giffeler/ergometal/blob/main/Benchmarks/2026-08-15-search-gather-ceiling-ab.md
Replies
0
Boosts
0
Views
786
Activity
Aug ’26
Metal-cpp usability issue with MTL::Buffer and MTL::ResidencySet
I know this might be a peeve of mine, but looking into programming a simple Metal4 Compute Shader example, essentially updating the Performing Calculations on a GPU example code to work with CPP and Metal4. I found that MTL::Allocation and MTL::Buffer pointers can't be used interchangeably when you are trying to add allocations to a MTL::ResidencySet, this is forcing you to: Know by heart that they inherit from each other and that you can just cast them (this is a bit suspicious though, it did work for me). Forcefully either C-cast or reinterpret_cast the MTL::Buffer pointer to a MTL::Allocation pointer as the MTL::ResidencySet will only accept that type. I might as well just be plain wrong about how this is used, any tips on correct usage in that case? Is there any expectation to either provide a typecast operator or add inheritance to support the expected behaviour seen in Swift and ObjC, which is just passing the thing? Opened a report with # FB24534953 with some extra information. Bear in mind that the example code uses Premake5, but it can generate an Xcode solution easily.
Replies
0
Boosts
0
Views
765
Activity
Aug ’26
Xcode 27.0 b5, macOS 26.6.1, Metal build fails: symbols not found for air64_v28
I've just downloaded the Xcode 27.0 beta 5 on a macOS 26.6.1 machine and tried to build my app (which includes Metal CoreImage kernels). I'm met with a new (to me) error; /Users/…/Developer/…/air-lld:1:1 symbol(s) not found for target 'air64_v28-apple-macosx26.0.0' and from the build log; air-lld: warning: ignoring file '/Applications/Xcode-beta.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk/System/Library/Frameworks/CoreImage.framework/Versions/A/CoreImage.metallib', file AIR version (2.9) is bigger than the one of the target being linked (2.8) air-lld: error: symbol(s) not found for target 'air64_v28-apple-macosx26.0.0' metal: error: air-lld command failed with exit code 1 (use -v to see invocation) I have no MTL_LANGUAGE_REVISION in my build settings. If I add one, with value Metal41 the app builds fine, but crashes at runtime as Metal 4.1 isn't supported on macOS 26. I imagine this is a beta Xcode and or macOS SDK bug, is there a workaround?
Replies
1
Boosts
1
Views
1.3k
Activity
Aug ’26
MPS bf16 softmax produces NaN on M5 Max (regression from M4) — breaks all on-device diffusion inference
Metal Performance Shaders produces from bf16/fp16 softmax for large attention tensors on , forcing the entire local generative-AI ecosystem to fall back to fp32. This is a . Environment Minimal reproduction Decomposed softmax on MPS — diffs = x - maxes produces NaN even though every element is identical (result should be all zeros): Real-world impact Running ComfyUI (the dominant local generative-AI UI) on M5 Max: Root cause (per PyTorch MPS maintainers) PyTorch maintainers (@drisspg, @albanD) have traced this to in the MPS/MPSGraph fused kernels. The NaN originates in the x - maxes subtraction inside softmax, then propagates through the attention block and the entire network. PyTorch cannot fix this — it is in the Metal/MPS kernel layer. Not isolated The M4-era fix ("fixed on macOS 15.1") did not survive onto M5, indicating the MPS fused-kernel precision fix was either reverted or not ported to the M5 GPU architecture. Request
Replies
1
Boosts
0
Views
1.1k
Activity
Aug ’26
Metal Shader Converter thread safety
Hello Apple! We've got offline shader compilation from HLSL -> Metallib using DXC -> SPIR-V -> metal.exe. This works okay for the most part, but it requires the creation of intermediate files to pass to/from the metal.exe process and we've had some issues with metal.exe sometimes not launching (probably our fault). Then we noticed Metal Shader Converter (MSC) exists and has a DLL - this looks way better since there's no need to launch processes or store intermediate files. However, upon trying to replace metal.exe with it I quickly ran into rampant heap corruption. I was surprised because the docs claim this: Each thread in your program needs to create its own instance of IRCompiler to avoid race conditions. But once I start calling IRCompilerAllocCompileAndLink in parallel all hell breaks loose, whether or not each thread has its own IRCompiler. I figured I must be doing something wrong, so I removed my attempt and compiled DXC locally with the MSC integration and encountered the exact same heap corruption. So I'm inclined to think the library isn't actually thread safe, but I'm wondering if there's something I'm missing? I tried all 3 versions of MSC just in case it was a problem with 3.0, but I got the same result each time. The only way to make it work was to surround compilation with a mutex, which makes its use pointless in our case.
Replies
2
Boosts
0
Views
1.5k
Activity
Aug ’26