Tile-Based GPUs Explained: Apple, Mali, Adreno, PowerVR and TBDR

A tile-based GPU splits the screen into small tiles, sorts every triangle into the tiles it touches, then renders one tile at a time inside fast on-chip memory and writes only the finished pixels out. An immediate mode GPU shades each triangle as it arrives and does its depth and colour work against external memory, so it moves far more bytes per frame.

That difference in where the framebuffer lives is the reason almost every phone, tablet and Apple silicon Mac uses a tiler, and it changes how you should write render passes and compute kernels for them. This guide goes deeper than our overviews of AI chip architecture and how a GPU is made. It is built from vendor documentation from Apple, Arm, Qualcomm, Imagination and Khronos, listed in full under Sources. Where a vendor has not published a number, we say so rather than guess.

A note on the code: the Metal compute example was compiled and run on an Apple M5 Mac while we wrote this, and produced the correct result. The Vulkan and Metal render pass snippets follow vendor documentation but were not executed, because they need a full rendering app around them. We publish no timings of our own; every performance number is either a vendor's published measurement or a worked model with its assumptions stated.

Key takeaways

• Two loops, swapped. An immediate mode renderer (IMR) loops over primitives, then fragments. A tile-based renderer loops over tiles, then over the primitives in each tile, so the colour, depth and stencil working set for one tile fits on chip.
• Two passes per render pass. Tilers first run a geometry or binning pass that writes transformed vertices and per-tile lists to system memory, then a per-tile fragment pass that reads them back. Geometry traffic goes up; framebuffer traffic goes down, usually by much more.
• Deferred means visibility first. Apple and Imagination PowerVR GPUs resolve which opaque surface is in front at each pixel before running fragment shaders, so overdraw from opaque geometry is removed regardless of draw order. Imagination's own capture of a game frame went from 4.7 to 1.2 shaded fragments per pixel.
• Load and store actions are the biggest lever you control. Clear or don't care on load, store only what the next pass reads, and make depth and MSAA attachments transient or memoryless. Arm measured a 600.2 MiB/s read bandwidth difference from a single wrong load action.
• Compute is not automatically tile aware. A compute dispatch between render passes forces tile memory out to DRAM. Apple tile shaders, Vulkan subpasses and Qualcomm's tile shading extensions exist to keep compute inside the tile.
• Mobile uses tiles because DRAM traffic costs power. Arm's own estimate is around 100 mW per GB/s of external DDR traffic on a phone, against an approximate 2.5 W device power budget in its example.

What is the difference between immediate mode and tile-based rendering?

Every rasterising GPU does the same logical work for a draw: vertex shading, triangle setup and clipping, rasterisation, fragment shading, depth and stencil tests, and blending into the colour attachment. The two families differ in the order of that work, and therefore in where the intermediate data lives.

The Arm authored pipeline barriers sample in the Khronos Vulkan Samples repository puts it in two pieces of pseudo code. An immediate mode GPU runs this:

for each ( primitive )
    for each ( fragment )
        execute fragment shader

A tile-based GPU runs this:

for each ( tile )
    for each ( primitive in tile )
        for each ( fragment in primitive in tile )
            execute fragment shader

The sample explains why the first loop is expensive: every shaded fragment has to "touch the working sets which include blending, depth testing and stencil operations", which at high resolution means several read, modify and write operations per fragment against external memory.

What an immediate mode renderer does

 Immediate mode renderer (IMR), one render pass

 draw 1 ─> vertex shade ─> raster ─> fragment shade ─> depth test ─> blend ──┐
 draw 2 ─> vertex shade ─> raster ─> fragment shade ─> depth test ─> blend ──┤
 draw 3 ─> vertex shade ─> raster ─> fragment shade ─> depth test ─> blend ──┤
                                                                            v
                         ┌──────────────────────────────────────────────────────┐
                         │ DRAM: colour, depth, stencil for the WHOLE screen     │
                         │ (read and written per fragment, through L2 caches)    │
                         └──────────────────────────────────────────────────────┘

Draws flow through the pipeline in submission order. A fragment from draw 3 that lands on top of a fragment from draw 1 causes the shader to run twice for that pixel, and both results touch the depth and colour buffers. Early depth testing helps if the application sorts opaque geometry from front to back, but Imagination's guide notes the limits: the sort must be redone as the camera moves, it is too coarse for intersecting objects, and it conflicts with sorting draws by state.

Discrete desktop GPUs are usually described as immediate mode renderers, and they compensate with very high bandwidth memory and large caches. Their vendors do not document their rasterisers as tile-based deferred renderers, so we do not describe their internals. Our post on GPU cost per million tokens shows how much of a data centre GPU's value is memory bandwidth.

What a tile-based renderer does

 Tile-based renderer, one render pass

 PASS 1: geometry / binning (all draws in the pass)
   draws ─> vertex (position) shade ─> clip, cull ─> bin into tiles
                                                         │
                                                         v
          ┌──────────────────────────────────────────────────────────────┐
          │ DRAM: transformed vertices + per-tile primitive lists         │
          │ Apple: "Tiled Vertex Buffer"  PowerVR: "Parameter Buffer"     │
          │ Adreno: "visibility stream"                                   │
          └──────────────────────────────────────────────────────────────┘
                                                         │
 PASS 2: per tile fragment work                          v
   for tile in screen:
     load action ─> raster ─> visibility (HSR / depth) ─> fragment shade ─> blend
                 \___________ all inside on-chip TILE MEMORY ___________/
                                                         │
                                          store action   v
                                                   DRAM framebuffer

Apple's WWDC 2020 session "Harness Apple GPUs with Metal" describes exactly these two phases. In the tiling phase, for the entire render pass, the GPU splits the viewport into tiles, shades all the vertices and bins the transformed primitives into per-tile lists. That output goes into a structure Apple calls the Tiled Vertex Buffer. In the rendering phase, for each tile, the GPU runs the load action, rasterises the primitives, computes visibility, shades the visible pixels and runs the store action.

The point of the second loop is that the colour, depth and stencil data for one tile is small enough to sit in on-chip SRAM. Apple's Metal documentation lists what that buys: tile memory bandwidth "many times faster than device memory", latency many times lower, and energy consumption significantly less. DRAM sees only the load and store actions for the framebuffer.

TBR versus TBDR: what "deferred" adds

Not every tiler is deferred in the same way, and the terms get blurred. A tile-based renderer (TBR) gets the bandwidth win from on-chip tile memory. A tile-based deferred renderer (TBDR) goes one step further: inside each tile, it works out which opaque primitive is visible at each pixel before it runs any fragment shader, so fragments that would later be covered are never shaded at all.

Imagination's guide states the PowerVR version plainly: the architecture "will defer all texturing and shading operations until all objects have been tested for visibility", and its hidden surface removal can remove overdraw entirely for fully opaque renders. Apple uses the TBDR name for its GPUs. Arm and Qualcomm document their own early visibility features, covered below.

Side by side

How binning works: the parameter buffer, visibility streams and partial renders

The first pass is where tilers pay for their framebuffer savings. Each vendor names the intermediate structure differently, but the contents are similar.

Imagination PowerVR. The Tiler stage runs vertex processing, then clip, project and cull, then a block called the Tile Accelerator (TA) determines which tiles each transformed primitive touches. Per-tile lists hold primitive lists that point at the transformed vertex data. Both the tile lists and the vertex data go into the Parameter Buffer (PB), which Imagination says resides in system memory and is mostly managed by the hardware.

Apple. The equivalent is the Tiled Vertex Buffer. Apple's WWDC 2020 session says it stores post-transform vertex data and other internal data, is mostly opaque to the developer, and "may cause a Partial Render if full". A partial render means the GPU splits the render pass to flush that buffer, which forces tile contents out to memory and back in. Very geometry heavy passes can therefore lose the tiler's main advantage.

Qualcomm Adreno. The binning pass marks which bins each triangle is visible in and records this in a visibility stream that is stored in system memory. In the rendering pass, only the primitives visible in the current bin are processed, using GMEM as the local colour and depth buffer, and the finished bin is resolved to system memory. The compiler generates a position-only vertex shader for binning: a version of your vertex shader reduced to the instructions that affect vertex positions. Since the Snapdragon A7x series, Qualcomm says, concurrent binning lets that pass run asynchronously ahead of other work when no dependency prevents it.

Arm Mali. Arm's reference sheet lists an Index-Driven Vertex Shading (IDVS) geometry pipeline on every Mali GPU from Bifrost onwards. Arm's product pages describe the 5th Gen architecture, starting with Mali-G720 and Immortalis-G720, as having a "deferred vertex shading geometry pipeline". Arm's detailed description of DVS is on its developer documentation site, which we could not retrieve for this post, so we only repeat the product page claim that Mali-G720 reduces memory bandwidth usage compared with its predecessor.

What all four share is a cost model. Tiling turns per-fragment framebuffer traffic into per-vertex and per-triangle traffic. Many pixels with moderate geometry strongly favours the tiler; very dense geometry over few pixels, such as a detailed shadow map, does not. Qualcomm's FlexRender, covered below, exists for exactly that case.

How big is a tile, and how much on-chip memory does it use?

Tile size matters because it sets how much per-pixel data fits on chip. Some vendors publish it; others choose it dynamically.

Two worked numbers show why these sizes are chosen.

Apple. A 32 by 32 tile at 128 bytes per pixel is 32 x 32 x 128 = 131,072 bytes, exactly 128 KB. The Feature Set Tables list a maximum implicit imageblock allocation of 128 KB on Apple4 to Apple8 and 256 KB on Apple9 and Apple10, so on Apple7 and Apple8 the largest per-pixel layout at full tile size matches the allocation limit exactly. The tables also halve the tile to 32 by 16 at 4x MSAA, which is consistent with a fixed byte budget shared by more samples per pixel.

Arm. A 16 by 16 tile at 256 bits (32 bytes) per pixel is 8 KB of colour storage. If Mali-G710 also uses 16 by 16 tiles, its 1024 bits (128 bytes) per pixel would be 32 KB per tile. Arm's reference sheet has a footnote that tile storage per pixel can exceed the listed figure, "but with reduced tile size", and the Arm subpasses sample says the same thing about G-buffers that exceed the limit: they work, at the cost of smaller tiles, lower fragment throughput and more bandwidth reading tile lists.

How many tiles in a frame? At a phone resolution of 2400 by 1080 pixels, a 32 by 32 grid needs 75 x 34 = 2,550 tiles per frame, and a 16 by 16 grid needs 150 x 68 = 10,200 tiles. Arm's subpasses sample does the same calculation for a Galaxy S10 class device with a Mali-G76 at 2220 by 1080: 2,397,600 pixels divided by 256 is about 9.4 thousand tiles per frame, or about 281 thousand per second at 30 fps. Arm rounds this to "~9k" and "270k/s", and that is what its physical tiles counter showed. If you round partial tiles up at the edges, the grid is 139 by 68, which is 9,452 tiles.

Hidden surface removal: how tilers avoid shading what you cannot see

Overdraw is the number of times the fragment shader runs per screen pixel. On a tiler, fragment shading is often the most expensive stage left once framebuffer traffic is on chip, so removing overdraw is the second big win.

Apple

Apple's WWDC 2020 session walks through Apple's hidden surface removal (HSR) step by step. Before any fragment shader runs, the on-chip depth buffer tracks the frontmost visible layer at each pixel, storing a depth value and a primitive ID. Rasterising an opaque triangle updates those values without shading anything. Apple describes HSR as "pixel perfect and submission order independent", so two opaque triangles drawn back to front still produce no overdraw.

The deferral has limits that Apple spells out. When a translucent primitive arrives, the HSR block has to flush the pixels it covers, meaning it runs the fragment shaders for the opaque surfaces underneath immediately so the blend has something to blend onto. A shader that uses discard or writes depth also forces the pipeline back to HSR. Apple's guidance is therefore to submit opaque meshes first, then alpha tested, discard or depth feedback meshes, then translucent meshes, and not to interleave them. You do not need to sort opaque meshes among themselves.

Imagination PowerVR

In PowerVR's per-tile Renderer, the Image Synthesis Processor (ISP) fetches only the screen-space position data for the tile, then performs HSR along with depth and stencil tests. A Tag Buffer records which triangle is on top at each pixel. Only after that does the Texture and Shading Processor (TSP) run fragment shaders for visible pixels, followed by alpha test and blend. Imagination's guide includes a measured example from a commercial mobile game: without HSR or early depth, the frame shaded 4.7 fragments per screen pixel on average; on a PowerVR device after HSR, 1.2. That is 1 minus 1.2 / 4.7 = 74.5 percent fewer fragments, which Imagination rounds to 75 percent. The remaining 0.2 above one is from blended geometry that cannot be removed.

The same guide tells developers not to do a depth pre-pass on PowerVR. On an IMR, a depth-only pass followed by a shading pass with depth test set to equal is a common trick to avoid overdraw. On a deferred renderer, the hardware already does this for opaque geometry, so a pre-pass just doubles the geometry work.

Arm Mali

Arm's reference sheet lists hidden surface removal and transaction elimination as supported on every Mali GPU in the sheet, from Midgard Mali-T720 through Valhall Mali-G710. Arm's product pages add "Fragment prepass" for Immortalis-G925 and Mali-G725 alongside "Doubled tiler throughput". Arm's detailed description of fragment prepass was not retrievable for this post, so we do not describe its mechanism.

Transaction elimination is a different kind of saving and is worth knowing. Arm's technology page describes it as comparing each 16 by 16 pixel tile of the current frame with the previous frame using a CRC-based signature, and writing out only the tiles that changed. For user interfaces and games with static regions, unchanged tiles cost no write bandwidth.

Qualcomm Adreno

Adreno builds a low resolution Z buffer (LRZ) during the binning pass and uses it in the rendering pass to reject pixels before testing them against the full-resolution depth buffer. Qualcomm's best practices list many things that disable LRZ for a draw or until the next depth clear: blending, stencil use, writing depth from the shader, discard, alpha to coverage, changing the depth comparison direction, reading from the framebuffer, and masked writes. LRZ stays active in direct rendering mode but, Qualcomm notes, usually gives much less benefit there than with binned rendering.

Worked numbers: how much bandwidth does tiling save?

This section builds a simple model, checks it against Arm's published measurements, and then shows where tiling stops paying. Every computation here was checked in Python.

A deliberately simple IMR model

Take a phone render target of 2400 by 1080 pixels (2,592,000 pixels) at 60 frames per second, with a 4-byte RGBA8 colour buffer and a 4-byte 32-bit depth buffer. Assume an average opaque overdraw of 3, and, to keep the model an upper bound, assume no cache hits and that every fragment passes the depth test.

Per frame on an IMR:

• clearing colour and depth: 2,592,000 x 8 bytes = 20.7 MB
• per fragment depth read, depth write and colour write: 2,592,000 x 3 x 12 bytes = 93.3 MB
• total: 114.0 MB per frame, or 6.84 GB/s at 60 fps

On a tiler with colour clear and store, and depth clear and don't care:

• clears happen in tile memory: 0 bytes to DRAM
• per fragment work happens in tile memory: 0 bytes to DRAM
• store colour once: 2,592,000 x 4 bytes = 10.4 MB per frame, or 0.62 GB/s

That is an 11 to 1 ratio for the framebuffer alone. Real IMRs do much better than this upper bound because of large caches, early depth rejection and framebuffer compression. The point of the model is the shape, not the exact ratio: IMR framebuffer traffic scales with pixels times overdraw times bytes per access, while tiler framebuffer traffic scales with pixels times bytes stored.

The tiler's geometry bill

The tiler has to write and then read back its binned geometry. Vendors do not publish the size of a binned vertex or a tile list entry, so the following numbers are assumptions for illustration. Assume 32 bytes per transformed vertex and 4 bytes per tile list reference, with each triangle touching 1.5 tiles on average.

At moderate geometry, the tiler still moves about a third of the bytes the simple IMR model moves. At ten times the geometry, the binning bill alone is bigger than the IMR's framebuffer bill. This crossover is the honest trade-off of tiling, and the reason Apple warns about partial renders and every vendor asks for level of detail systems that keep triangle counts sensible.

Checking against Arm's measurements

Arm's render passes sample gives a formula for the cost of loading or storing one uncompressed attachment: width x height x bytes per pixel x frames per second. For 2220 by 1080 at 4 bytes and about 60 fps, Arm quotes about 575 MiB/s. Strictly, 2220 x 1080 x 4 x 60 = 575,424,000 bytes per second, which is 575 MB/s or 548.8 MiB/s; the sample's figure is in decimal megabytes. Arm then measured the real difference between LOAD_OP_LOAD and LOAD_OP_CLEAR on a colour attachment: external read bandwidth fell from 1533.9 to 933.7 MiB/s, a difference of 600.2 MiB/s. That is roughly one full attachment read per frame, which is what the formula predicts.

The same repository's MSAA sample measured what happens when a 4x multisampled colour attachment is written out and resolved in a separate pass instead of resolved on chip: about 5 GB/s of extra bandwidth on a high-end Mali-G76 phone. Our own estimate agrees. At 2400 by 1080, a 4x MSAA RGBA8 attachment is 2,592,000 x 4 samples x 4 bytes = 41.5 MB. Writing it and reading it back once per frame at 60 fps is 4.98 GB/s.

What that costs in power

Arm's MSAA sample gives the conversion it uses for a mobile device: external DDR bandwidth "costs around 100 mW per GB/s". At that rate the extra 5 GB/s from a separate resolve pass costs about 500 mW, which Arm describes as 20 percent of an approximate 2.5 W device power budget. In the sample's worst case, where multisampled depth and colour are both written out, the total bandwidth increase was 6.3 GB/s, about 630 mW or 25 percent of that budget. These are Arm's estimates for one class of phone, not universal constants, but they explain the design priority.

G-buffers and subpasses

Deferred shading writes a G-buffer of several attachments in one pass and reads them in a lighting pass. With four RGBA8 G-buffer targets plus 4 bytes of depth, that is 20 bytes per pixel. If the G-buffer goes to DRAM and back, our 2400 by 1080 target moves 2,592,000 x 20 x 2 = 103.7 MB per frame, or 6.22 GB/s at 60 fps. If the two passes are merged into one tile pass, that traffic disappears. Four RGBA8 colour targets are 128 bits per pixel, which fits even the oldest Mali tile budget in Arm's reference sheet.

Arm's subpasses sample measured exactly this on a Mali-G76. With subpasses merged, the physical tiles counter read 262.2 thousand per second; with separate render passes, 614.7 thousand per second. By our arithmetic that is a 57 percent reduction in tiles rendered; Arm's text describes it as approximately 55 percent, and says bandwidth fell by a similar share.

Apple GPUs: TBDR, unified memory and Metal's tile features

Architecture and unified memory

Apple has designed its own GPUs for its A-series and M-series chips, and its Metal documentation describes all of them as TBDR GPUs. Apple's WWDC 2020 session for developers summarises the memory model in two sentences: the CPU and GPU share system memory under a unified memory architecture, and the GPU has a dedicated pool of on-chip memory called tile memory. There is no separate video memory. Apple's press release for M1, dated 10 November 2020, describes the chip as having a unified memory architecture and a GPU with up to eight cores "capable of running nearly 25,000 threads simultaneously", with 2.6 teraflops of throughput. That teraflops number is Apple's peak figure, not a sustained one.

Unified memory has two practical consequences for compute. First, a buffer allocated with shared storage is visible to both the CPU and GPU without a copy across a bus. Second, the GPU's DRAM bandwidth is shared with the CPU and every other block on the chip, which is one more reason tile memory matters.

GPU families and the limits that matter

The Metal Feature Set Tables map chips to GPU families. In the 21 May 2026 edition:

The same tables show the feature timeline that matters for tile work. Imageblocks, tile shaders, raster order groups and imageblock sample coverage control arrive with Apple4 (A11). Memoryless render targets and programmable blending date back to Apple2. SIMD-scoped reduction operations and SIMD-scoped matrix multiply arrive with Apple7 (A14 and M1). Apple10 adds 8x MSAA with a 16 by 16 tile.

SIMD-groups of 32

Metal groups threads in a threadgroup into SIMD-groups, which Apple's WWDC 2022 session "Scale compute workloads across Apple GPUs" says are "also known as waves or warps in other compute dialects". The session states that the pipeline's threadExecutionWidth returns the SIMD width and "on all Apple GPUs, it is equal to 32", and that threadgroups can have up to 1024 threads sharing up to 32 KB of threadgroup memory. Our Apple M5 test reported the same: 32 and 1024. Apple's tables note that 1024 is a theoretical maximum; a specific pipeline can report less.

So a full 1024-thread threadgroup is 32 SIMD-groups of 32 threads, and 32 KB of threadgroup memory holds 8,192 single-precision floats, which is 8 floats per thread at full occupancy of that threadgroup.

Tile memory as an API concept

Apple exposes tile memory more directly than most vendors. Its Metal documentation on tailoring apps for Apple GPUs describes four features available from A11 onwards:

• Imageblocks. Two-dimensional structures in tile memory with a per-pixel layout you define, which persist for the lifetime of a tile across draws and dispatches. A fragment sees only its own pixel; a compute thread can see the whole imageblock.
• Tile shaders. Compute or fragment functions that run as part of a render pass and read and write tile memory, so you can mix render and compute work without a round trip through device memory. Apple's example is a tiled deferred renderer whose G-buffer, light culling and lighting phases share one render pass.
• Raster order groups. An attribute on memory pointers that makes overlapping fragment shader threads access the same pixel in submission order, so fragment shader read, modify and write sequences do not race.
• Imageblock sample coverage control. Control of per-sample coverage when working with MSAA data in imageblocks.

The fourth piece is the memoryless storage mode. Apple's documentation says a memoryless resource "uses tile memory", is available only on Apple family GPUs, and exists only temporarily during a render pass. Depth and multisample attachments that are never stored are the classic use.

Arm Mali and Immortalis: Midgard to the 5th Gen architecture

Arm licenses Mali GPU designs to chip makers, so the same GPU appears in many phones from different brands, and each chip maker chooses the core count and clocks. Immortalis is the name Arm gives its flagship GPUs, and Arm's pages list hardware ray tracing support for them.

Generations

Arm's product pages give no announcement dates, so we omit them. The Arm reference sheet below is dated 2022 and stops at Mali-G710; we found no Arm per-core figures for 5th Gen parts.

Per-core configuration from Arm's reference sheet

The register file trade-off is spelled out for Valhall. Mali-G710 runs 2048 threads per core when a shader uses 0 to 32 work registers, and 1024 threads when it uses 33 to 64. At a warp width of 16, that is 128 warps per core with small shaders and 64 warps with large ones. This is the Mali equivalent of the occupancy calculation CUDA developers do: halving the resident threads halves the warps available to hide memory latency.

To turn per-core figures into a chip figure you multiply by the core count and clock the chip maker chose, which Arm does not publish for a specific phone. For example, a hypothetical 10-core Mali-G710 at 850 MHz would peak at 10 x 128 x 0.85 = 1,088 GFLOPS of FP32. That is an illustration of the arithmetic, not the spec of any shipping device.

Mali features that matter for tilers

• Transaction elimination. CRC signatures per 16 by 16 tile; unchanged tiles are not written back.
• Arm Frame Buffer Compression (AFBC). A lossless compression format used across Arm's multimedia IP. Arm's page says it can reduce system bandwidth and power for transferring image data "by up to 50 percent" in video post processing uses. Arm's Vulkan sample on render passes notes that load and store savings are lower when images are AFBC compressed, because there are fewer bytes to save.
• Tile storage limits for subpass merging. Arm's subpasses sample says Mali merges Vulkan subpasses only when they use at most 128 bits per pixel of tile buffer colour storage, with later GPUs such as Mali-G72 raising that to 256 bits. Exceed it and the driver either shrinks tiles or splits the pass.

Qualcomm Adreno: GMEM, FlexRender and tile shading

GMEM and binning

Qualcomm's Adreno documentation describes a tiled architecture in its own vocabulary. The frame is divided into bins, region sizes are chosen automatically, and each bin is rendered in on-chip high performance Graphics Memory, which Qualcomm calls GMEM. Moving data between GMEM and system memory has its own names: a GMEMLoad brings existing contents into GMEM at the start of a bin, a GMEMStore resolves results out at the end. Qualcomm's advice mirrors the load and store guidance in this post: invalidate early with glInvalidateFramebuffer, clear with glClear to avoid GMEM loads, and in Vulkan use render passes and subpasses so work stays in GMEM.

FlexRender

FlexRender is Qualcomm's name for Adreno's ability to switch, per render target and mid-frame, between binned rendering through GMEM and direct rendering to the framebuffer in system memory. Qualcomm states that the driver and GPU analyse the rendering parameters and select the mode automatically, and that the heuristics are not exposed to developers. The scenarios it lists as generally triggering direct mode are:

• a high ratio of texture samples in vertex shaders to vertices
• a small number of vertices or draws
• use of tessellation or geometry shaders

That matches the worked trade-off earlier. Snapdragon Profiler's Rendering Stages metric shows which mode each surface used. One consequence: Qualcomm notes that Vulkan subpass merging only applies when a surface is rendered in binning mode.

Tile shading extensions

For Adreno 840 and later, Qualcomm says, two Vulkan extensions expose GMEM directly to applications. VK_QCOM_tile_memory_heap lets you allocate images and buffers in GMEM and control whether they persist across render passes. VK_QCOM_tile_shading lets compute shaders run inside a render pass or dynamic rendering scope and access tiled attachments in GMEM, and supports per-tile draws for GPU-driven rendering. Before these extensions, Qualcomm notes, GMEM was available only to the driver and was "entirely off-limits for compute shaders". The Khronos reference page for VK_QCOM_tile_shading lists a last modified date of 13 August 2025.

Qualcomm describes Adreno's shader cores as unified and scalar, but does not publish wave sizes, ALU counts or GMEM capacities per Adreno model in the developer documentation we read, so we do not quote any.

Imagination PowerVR: the original TBDR

Imagination's history page says the company began as VideoLogic in 1985, supplied the PowerVR Series2 GPU in Sega's Dreamcast console, renamed itself Imagination in 1999 to focus on licensing IP, and turned to mobile phones in the mid 2000s. TBDR has been the core of the design throughout.

The current architecture guide splits the GPU into two halves.

 PowerVR TBDR pipeline (from Imagination's architecture guide)

 TILER                                         RENDERER (per tile)
 ─────                                         ───────────────────
 vertex processing                             fetch tile list from Parameter Buffer
   -> clip, project, cull                        -> ISP: fetch positions,
   -> Tile Accelerator (TA): which tiles?           hidden surface removal,
   -> per-tile primitive lists                      depth and stencil test
        │                                        -> Tag Buffer: which triangle
        v                                           is on top per pixel
 ┌───────────────────────────┐                   -> TSP: texture and shade
 │ Parameter Buffer (DRAM):   │ ─────────────>      visible pixels only
 │ tile lists + transformed   │                   -> alpha test, blend
 │ vertex data                │                   -> write tile to framebuffer
 └───────────────────────────┘

Three further details from the guide matter to developers. Colour, depth and stencil read, modify and write happen in on-chip buffers, and only attachments the application chose to preserve are written to system memory. The shader engine is massively multi-threaded and schedules tasks based on data availability. And every PowerVR core is managed by firmware, which handles most high-level graphics events inside the GPU instead of in the CPU driver, to cut latency and CPU overhead. That firmware is one more piece of software that reaches devices through updates, which matters for the supply chain section at the end.

Writing render passes for a tiler

The single most effective change for tilers is to tell the driver what you need loaded at the start of a pass and stored at the end. Both Vulkan and Metal make this explicit.

Vulkan

The Khronos reference page for VkAttachmentLoadOp defines VK_ATTACHMENT_LOAD_OP_CLEAR as clearing the render area to a uniform value and VK_ATTACHMENT_LOAD_OP_DONT_CARE as meaning "the previous contents within the area need not be preserved". On a tiler, either one avoids reading the old attachment from DRAM into tile memory. LOAD_OP_LOAD is the one that costs bandwidth.

The following snippet follows the Khronos reference pages and Arm's render passes sample. It was not executed here.

/* Depth is only needed during this pass: clear on load, drop on store,
   and let a tiler keep it entirely in tile memory. */
VkAttachmentDescription depth = {
    .flags          = 0,
    .format         = VK_FORMAT_D32_SFLOAT,
    .samples        = VK_SAMPLE_COUNT_1_BIT,
    .loadOp         = VK_ATTACHMENT_LOAD_OP_CLEAR,
    .storeOp        = VK_ATTACHMENT_STORE_OP_DONT_CARE,
    .stencilLoadOp  = VK_ATTACHMENT_LOAD_OP_DONT_CARE,
    .stencilStoreOp = VK_ATTACHMENT_STORE_OP_DONT_CARE,
    .initialLayout  = VK_IMAGE_LAYOUT_UNDEFINED,
    .finalLayout    = VK_IMAGE_LAYOUT_DEPTH_STENCIL_ATTACHMENT_OPTIMAL,
};

/* The image itself is transient, so a tiler need not back it with memory. */
VkImageCreateInfo depth_image = {
    .sType       = VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO,
    .imageType   = VK_IMAGE_TYPE_2D,
    .format      = VK_FORMAT_D32_SFLOAT,
    .extent      = { width, height, 1 },
    .mipLevels   = 1,
    .arrayLayers = 1,
    .samples     = VK_SAMPLE_COUNT_1_BIT,
    .tiling      = VK_IMAGE_TILING_OPTIMAL,
    .usage       = VK_IMAGE_USAGE_DEPTH_STENCIL_ATTACHMENT_BIT |
                   VK_IMAGE_USAGE_TRANSIENT_ATTACHMENT_BIT,
    .sharingMode = VK_SHARING_MODE_EXCLUSIVE,
    .initialLayout = VK_IMAGE_LAYOUT_UNDEFINED,
};
/* When allocating, prefer a memory type with
   VK_MEMORY_PROPERTY_LAZILY_ALLOCATED_BIT; if no memory type exposes it,
   fall back to ordinary device-local memory. */

Arm's sample explains the pairing: an image with the transient usage bit, backed by lazily allocated memory, "will not even need physical storage" on a tiler, and its contents must be invalidated with STORE_OP_DONT_CARE at the end of the pass. The same pattern applies to a multisampled colour attachment that is resolved on write-back through pResolveAttachments: keep the multisampled image transient and store only the resolved single-sample image. That is how Arm's MSAA sample avoids the extra 5 GB/s measured above.

Vulkan 1.4 adds another route for tilers. VK_KHR_dynamic_rendering_local_read, which the Khronos reference page shows as promoted to Vulkan 1.4, lets a dynamic rendering pass read attachments written earlier in the same pass at the same pixel, giving dynamic rendering users the same on-chip G-buffer pattern that subpasses gave render pass users.

Metal

The Metal equivalent uses load and store actions on a render pass descriptor and the memoryless storage mode. This snippet follows Apple's documentation; it was not executed here.

import Metal

// A depth buffer that lives only in tile memory.
let depthDesc = MTLTextureDescriptor.texture2DDescriptor(
    pixelFormat: .depth32Float, width: width, height: height, mipmapped: false)
depthDesc.usage = .renderTarget
depthDesc.storageMode = .memoryless          // tile memory only, Apple GPUs only
let depth = device.makeTexture(descriptor: depthDesc)!

let pass = MTLRenderPassDescriptor()
pass.colorAttachments[0].texture     = drawable.texture
pass.colorAttachments[0].loadAction  = .clear      // no read from DRAM
pass.colorAttachments[0].clearColor  = MTLClearColor(red: 0, green: 0, blue: 0, alpha: 1)
pass.colorAttachments[0].storeAction = .store      // the one write we need
pass.depthAttachment.texture         = depth
pass.depthAttachment.loadAction      = .clear
pass.depthAttachment.storeAction     = .dontCare   // required for memoryless

Apple's WWDC 2020 guidance is the same as Arm's and Qualcomm's: "only load the data that you need", clear instead of loading when you do not need the old contents, and only store what you need, such as the main colour attachment.

Why mobile uses tiles, and what that means for compute

The power argument

A phone GPU shares its memory bus with the CPU and the rest of the chip and lives inside a few watts. At Arm's estimate of around 100 mW per GB/s of external DDR traffic, against roughly 2.5 W for the device in its example, a few gigabytes per second of avoidable framebuffer traffic is a large fraction of the budget. Arm's sample text adds the hardware side: the IMR approach needs high-frequency memory with wide interfaces and many pins, which are energy intensive.

Imagination and Apple make the same argument in their own documentation: less external memory traffic means more performance and less power, which matters most on battery powered devices.

Desktop and data centre GPUs make a different trade. They have hundreds of watts to spend and attach GDDR or HBM memory with very wide buses, as described in our post on what an AI accelerator is. Paying for that bandwidth lets them avoid the geometry cost and pass-boundary constraints of binning.

What changes for compute work

Compute kernels on a tiler do not run in the per-tile loop. A dispatch on a separate compute encoder or queue reads and writes device memory like it would on any GPU. That creates three practical rules.

1. A compute pass in the middle of rendering splits the render pass. Everything in tile memory that the compute step needs is stored to DRAM first and loaded back afterwards. Apple's tile shaders, Vulkan subpasses with input attachments, VK_KHR_dynamic_rendering_local_read and Qualcomm's VK_QCOM_tile_shading are all ways to keep that work inside the render pass.
2. Threadgroup or workgroup memory is your on-chip scratchpad. On Apple GPUs, 32 KB per threadgroup from Apple4 onwards. Apple suggests threadgroup memory for unstructured data and imageblocks, which compute kernels can also use, for image data.
3. Unified memory removes copies but not bandwidth limits. On Apple silicon, a buffer in shared storage mode is written by the CPU and read by the GPU with no transfer step, which our example below does. But a memory-bound kernel is still limited by the DRAM bandwidth the whole chip shares.

For a reduction, arithmetic intensity is about 1 floating-point add per 4 bytes read, 0.25 FLOP per byte. That is far below the ratio of peak FLOPS to bandwidth on any modern GPU, so the kernel is memory bound and its job is to read each element once, coalesced, and do the rest on chip. The same reasoning drives the matrix multiply tiling in our AI chip architecture post.

A short Metal compute example

This kernel reduces a large float array to one partial sum per threadgroup, using SIMD-group reductions and a small threadgroup array. It needs SIMD-scoped reduction operations, which Apple's Feature Set Tables list from the Apple7 family (A14 and M1) onwards.

#include <metal_stdlib>
using namespace metal;

// One partial sum per threadgroup. Launch with up to 1024 threads per threadgroup.
kernel void partial_sum(device const float *in       [[buffer(0)]],
                        device float       *partials [[buffer(1)]],
                        constant uint      &n        [[buffer(2)]],
                        uint gid      [[thread_position_in_grid]],
                        uint tgid     [[threadgroup_position_in_grid]],
                        uint sg_index [[simdgroup_index_in_threadgroup]],
                        uint lane     [[thread_index_in_simdgroup]],
                        uint sg_count [[simdgroups_per_threadgroup]])
{
    threadgroup float sg_sums[32];          // 1024 threads / 32 lanes = 32 SIMD-groups

    float v = (gid < n) ? in[gid] : 0.0f;   // guard the last, partial threadgroup
    v = simd_sum(v);                        // reduce across the lanes of this SIMD-group
    if (lane == 0) { sg_sums[sg_index] = v; }
    threadgroup_barrier(mem_flags::mem_threadgroup);

    if (sg_index == 0) {
        float s = (lane < sg_count) ? sg_sums[lane] : 0.0f;
        s = simd_sum(s);
        if (lane == 0) { partials[tgid] = s; }
    }
}

The host code in Swift compiles the kernel at runtime, fills a shared buffer from the CPU and dispatches one thread per element.

import Metal
import Foundation

let device = MTLCreateSystemDefaultDevice()!
let src = try String(contentsOfFile: "reduce.metal", encoding: .utf8)
let library = try device.makeLibrary(source: src, options: nil)
let pipeline = try device.makeComputePipelineState(
    function: library.makeFunction(name: "partial_sum")!)
print("threadExecutionWidth:", pipeline.threadExecutionWidth,
      "maxTotalThreadsPerThreadgroup:", pipeline.maxTotalThreadsPerThreadgroup,
      "hasUnifiedMemory:", device.hasUnifiedMemory)

let n = 1 << 24                                  // 16,777,216 floats = 64 MiB
let input = device.makeBuffer(length: n * MemoryLayout<Float>.stride,
                              options: .storageModeShared)!
let p = input.contents().bindMemory(to: Float.self, capacity: n)
for i in 0..<n { p[i] = Float(i % 7) }           // CPU writes straight into GPU-visible memory

let tgSize = min(1024, pipeline.maxTotalThreadsPerThreadgroup)
let groups = (n + tgSize - 1) / tgSize
let partials = device.makeBuffer(length: groups * MemoryLayout<Float>.stride,
                                 options: .storageModeShared)!
var count = UInt32(n)

let queue = device.makeCommandQueue()!
let cmd = queue.makeCommandBuffer()!
let enc = cmd.makeComputeCommandEncoder()!
enc.setComputePipelineState(pipeline)
enc.setBuffer(input, offset: 0, index: 0)
enc.setBuffer(partials, offset: 0, index: 1)
enc.setBytes(&count, length: MemoryLayout<UInt32>.stride, index: 2)
enc.dispatchThreadgroups(MTLSize(width: groups, height: 1, depth: 1),
                         threadsPerThreadgroup: MTLSize(width: tgSize, height: 1, depth: 1))
enc.endEncoding()
cmd.commit()
cmd.waitUntilCompleted()

let q = partials.contents().bindMemory(to: Float.self, capacity: groups)
var gpu = 0.0
for g in 0..<groups { gpu += Double(q[g]) }
print("threadgroups:", groups, "GPU sum:", gpu)

What happened when we ran it on an Apple M5 Mac with swiftc -O main.swift and the kernel in reduce.metal beside it: the pipeline reported a threadExecutionWidth of 32, a maxTotalThreadsPerThreadgroup of 1024 and unified memory; the dispatch used 16,384 threadgroups; and the GPU sum, 50,331,645, matched a CPU sum of the same data exactly. We did not time it, and we publish no speed numbers for it.

Notes on the design, which carry over to other GPUs:

• Read once, coalesced. Adjacent lanes read adjacent addresses, and each element is read from DRAM exactly once.
• Reduce in registers first. simd_sum combines 32 values without touching threadgroup memory, so only 128 bytes per threadgroup go into threadgroup memory, far below the 32 KB limit, with a single barrier.
• Respect the pipeline's own limits. The host asks for maxTotalThreadsPerThreadgroup instead of assuming 1024, as Apple advises.
• Floating-point order differs. A GPU reduction adds in a different order from a CPU loop, so compare with a tolerance unless, as here, the data is exactly representable.

This kernel is pure compute and does not use tile memory. A tile shader version, which would run inside a render pass and read an imageblock, needs a full render pipeline around it, and Apple's sample code library is the right place to start for that.

Profiling a tiler

Each vendor ships a profiler with tile-level counters: Xcode's GPU tools for Apple, Streamline for Arm Mali (the source of the Arm measurements above, with counters such as physical tiles and external read and write bytes), and Snapdragon Profiler for Adreno, whose Rendering Stages metric shows whether each surface was binned or direct. A check that works everywhere: compute width x height x bytes per pixel x frames per second for every attachment you load or store, as Arm's sample does, and compare the sum against the external bandwidth counter. A counter far above your sum means something is loading or storing that you did not intend, or a pass is being split.

Where Safeguard fits

Safeguard does not design, program or speed up GPUs, and nothing in this post depends on Safeguard. The honest connection is the software around the GPU. A mobile game or graphics app ships a shader toolchain, an engine, Android Gradle and Maven dependencies, JavaScript or Python build tooling, and container images for build and CI farms. Those are software components with published vulnerabilities, and they need an inventory and a patch path like any other dependency. Safeguard generates SBOMs in SPDX and CycloneDX, inventories transitive dependencies across ecosystems including npm, PyPI, Maven and Gradle, Go, Cargo, NuGet, Composer and Bundler, scans container image layers, matches components against vulnerability databases, ranks findings with reachability analysis, and can open fix pull requests when a fixed version exists (remediation ships switched off by default).

It has limits here that you should plan around. GPU kernel drivers, user-mode drivers and the firmware that runs on the GPU itself, which Imagination's guide describes for every PowerVR core, ship from the GPU vendor and the device maker through operating system and vendor updates. Safeguard does not scan or patch GPU firmware or drivers on phones. For those, track the device maker's security bulletins and keep devices updated.

Rendering code is only one layer of a graphics product; the engines, build tools, shader compilers and container images around it pull in open source components that receive security fixes like any other dependency. The fit with GPU hardware itself is indirect, but the dependency hygiene below applies to every repository and image you build.

Common mistakes

1. Loading attachments you are about to overwrite. LOAD_OP_LOAD or Metal's .load on a colour buffer you then fully redraw reads a whole attachment from DRAM per frame. Arm measured 600.2 MiB/s of avoidable reads from this one mistake.
2. Storing depth and multisample attachments. If nothing reads depth after the pass, store it as don't care and make it transient or memoryless. Resolve MSAA on write-back, not in a separate pass.
3. Splitting passes with compute or barriers. A compute dispatch, a blit or a barrier between two render passes that use the same attachments forces a store and a reload. Use subpasses, local read, tile shaders or tile shading extensions where the platform supports them.
4. Interleaving opaque and blended draws. On Apple and PowerVR, a translucent draw in the middle of opaque ones flushes HSR and causes overdraw. Draw opaque first, then alpha tested, then translucent.
5. Doing a depth pre-pass on a deferred renderer. On PowerVR and Apple GPUs the hardware already removes opaque overdraw, so a pre-pass doubles geometry work for no gain. On Adreno and Mali, measure before assuming either way.
6. Using discard and depth writes casually. They defeat HSR on Apple and PowerVR and disable LRZ on Adreno.
7. Ignoring geometry cost. Tilers write binned geometry to DRAM and read it back. Millions of tiny triangles can overflow the parameter buffer (Apple's partial render) or push Adreno into direct mode. Use level of detail and keep triangles at least a few pixels in size.
8. Fat G-buffers. Exceeding per-pixel tile storage (128 or 256 bits on older Mali, up to 1024 bits on Mali-G710, 128 bytes per sample on recent Apple GPUs) shrinks tiles or splits passes.

FAQ

What is tile-based deferred rendering in simple terms?

The GPU first works out which triangles touch each small square of the screen. Then, for each square, it figures out which surface is in front at every pixel, shades only those visible pixels in fast on-chip memory, and writes the finished square to memory once. "Tile-based" is the square-by-square part; "deferred" is the visibility-before-shading part.

Do Apple Silicon Macs use tile-based rendering?

Yes. Apple's Metal documentation says the GPUs in Apple silicon implement tile-based deferred rendering, and the Metal Feature Set Tables list M1 through M5 in the Apple7 to Apple10 families, which support imageblocks, tile shaders and memoryless render targets.

Are desktop NVIDIA and AMD GPUs immediate mode renderers?

They are generally classed as immediate mode, and they rely on high bandwidth memory and large caches. Their vendors have not documented their rasterisers as tile-based deferred renderers, so we do not make claims about internal binning behaviour. Vulkan's load and store actions are still worth getting right on desktop, because they also reduce cache and memory traffic there.

What tile size do mobile GPUs use?

Apple's Metal tables give a maximum of 32 by 32 pixels without MSAA, 32 by 16 with 4x MSAA and 16 by 16 with 8x MSAA on Apple10. Arm's documentation and samples use 16 by 16 pixels for Mali. Qualcomm says Adreno chooses bin sizes automatically, influenced by resolution, MSAA and render target count.

What is FlexRender?

FlexRender is Adreno's ability to choose, per render target, between binned rendering through GMEM and direct rendering to system memory. The driver decides automatically; Qualcomm lists heavy vertex texturing, very few vertices or draws, and tessellation or geometry shaders as common triggers for direct mode.

What is GMEM on Adreno?

Qualcomm's name for the on-chip graphics memory that holds colour and depth for the current bin. On Adreno 840 and later, two Vulkan extensions let applications allocate resources in GMEM and run compute shaders against it.

What is the SIMD width on Apple GPUs, and the warp width on Mali?

Apple says threadExecutionWidth is 32 on all Apple GPUs, and our M5 test agreed. Arm's reference sheet lists 4 for early Bifrost, 8 for Mali-G52 and G76, and 16 for Valhall from Mali-G77 to G710.

Why do mobile GPUs use tiles instead of fast memory like desktop GPUs?

Power and cost. External DRAM traffic is expensive in energy; Arm estimates around 100 mW per GB/s on a phone. Keeping colour, depth and stencil traffic on chip lets a phone GPU render at high resolution within a few watts on a shared LPDDR bus.

Does tile-based rendering help GPU compute or machine learning?

Not automatically. Ordinary compute dispatches read and write device memory on any GPU. Tiling helps compute when the work can run inside a render pass on tile memory, through Apple tile shaders, Vulkan local read or subpasses, or Qualcomm's tile shading extension. For general compute, threadgroup memory and coalesced access matter more than the rasteriser design.

More in this series

Other posts in our GPUs, CPUs, QPUs and lithography series:

• AMD GPU Architecture Explained: GCN, RDNA, CDNA, Infinity Fabric and HIP
• CPU vs GPU vs QPU Explained: How Each Works and Why They Coexist
• CUDA Programming Model Explained: Grids, Blocks, Warps and Memory
• GPU Memory and Interconnects: HBM, GDDR, NVLink, PCIe, CXL
• CUDA vs HIP vs SYCL vs OpenCL vs Metal vs Vulkan vs Triton Compared
• GPU Software Supply Chain: Drivers, CUDA, cuDNN, NCCL and Your SBOM
• How a QPU Works: Qubit Chips, Fridges, Control Pulses and Readout
• How CPUs Work: Pipelines, Out-of-Order, Caches, Branches, SIMD
• How EUV Lithography Machines Work: Tin Plasma, Mirrors and High-NA
• How GPUs Work: Graphics Pipeline, Shader Cores, Ray Tracing and Compute
• Inside an NVIDIA SM: Warp Schedulers, Registers, Tensor Cores
• Intel GPU Architecture: Xe Cores, XMX, Ponte Vecchio and oneAPI SYCL
• NVIDIA GPU Architecture History: From Tesla and Fermi to Blackwell
• NVIDIA vs AMD vs Intel vs Apple GPUs: Architectures Compared
• Optimizing a CUDA Matrix Multiply Step by Step, From Naive to Tensor Cores
• Photolithography Explained: Contact Printing to DUV Immersion and Overlay

Fix it without waiting for upstream

Whatever the vulnerability, most of the risk sits in two gaps. The first is the time between a fix being published and that fix actually running in your production builds. The second, and worse, is the time before any fix exists at all, while a maintainer works out a patch and a release schedule. Waiting through both gaps is a habit, not a requirement. Safeguard gives teams three ways to close them.

1. Autonomous remediation: pull requests, not tickets

When a fixed version exists, Safeguard's Griffin models open a tested, policy gated pull request that moves the affected dependency to the safe version. You can run it on one repository or in bulk across every project that carries the vulnerable version. Remediation covers eight categories, including dependencies, licenses, security posture and findings. Reachability analysis ranks the vulnerable code paths your application actually calls first, so the first pull requests you merge are the ones that matter most. Autonomous remediation ships switched off, with every category set to Manual, so your team decides how much runs on its own and what still needs a human approval.

2. Gold Open Source: the safe version in seconds, free

gold.safeguard.sh is free and needs no login. Search any CVE or package to see affected and fixed versions, CVSS, EPSS, CISA KEV status and a recommended safe version, or use the "Am I affected?" checker. The same data is available through a free read API, README status badges and RSS threat feeds. The safeguard-sh/gold-scan-action GitHub Action can fail a build on a known exploited vulnerability, on any known vulnerability, or never, so a version you have just removed cannot quietly come back in next week's merge.

3. Gold Registry: start from components that are already fixed

registry.safeguard.sh provides hardened, drop in replacements for upstream packages and container images. Every Gold artifact is held continuously to zero exploitable critical or high CVEs, zero malware and a completed license check, with SLSA L3 or higher provenance, a Sigstore signature and a public evidence page you can hand to an auditor. An artifact loses its Gold badge if it regresses. Teams can point their own package directory at the registry so hardened equivalents resolve without changing how developers install anything.

When upstream has no fix yet, or will not ship one, Safeguard first works to get the fix accepted upstream. If the issue is severe or upstream will not move, Safeguard builds a fork, fixes the vulnerability there and publishes a hardened Gold build, so you are not stuck on a maintainer's release schedule. You can also request a hardened, compatibility tested build of a specific package. This fork and patch flow is in early access today, so ask which packages it covers before you plan around it.

One honest limit: when the vulnerable code lives in a vendor appliance, a hosted service or an operating system runtime you do not build yourself, the fix still has to come from that vendor or distribution. The steps above apply to the open source dependencies and container images you build and ship, which is where most teams carry the largest and least tracked share of their exposure.

See how Safeguard would help

The first demo is a real Safeguard CLI session in a repository: it reads a vulnerability finding and the advisory, asks before any command that changes something, edits the dependency, runs the install and tests with your approval, and pushes nothing. The same flow applies to the build tooling repository behind a game or graphics app, for example the Gradle project of an Android client or the scripts that drive a shader build.

Click through it at your own pace, or open it full screen.

The second demo sets up the Safeguard runner on your own machine, approves it with a device code and scans from the terminal. Graphics teams often keep large asset and build trees on dedicated build machines, and a local runner lets you scan the repository where it already lives.

Click through it at your own pace, or open it full screen.

The third demo connects Claude to the Safeguard MCP server with OAuth and lets you choose which tools Claude can use. An engineer can then ask about the dependencies in an engine or tools repository from an assistant they already use, without granting it more access than they choose.

Click through it at your own pace, or open it full screen.

Download the apps at https://safeguard.sh/download and read the CLI documentation at https://docs.safeguard.sh/docs/cli .

Quick check

Sources

• Apple: Tailor your apps for Apple GPUs and tile-based deferred rendering
• Apple: Metal Feature Set Tables (PDF, 21 May 2026 edition)
• Apple: MTLStorageMode.memoryless
• Apple: Calculating threadgroup and grid sizes
• Apple WWDC20: Harness Apple GPUs with Metal
• Apple WWDC20: Optimize Metal Performance for Apple silicon Macs
• Apple WWDC22: Scale compute workloads across Apple GPUs
• Apple Newsroom: Apple unleashes M1 (10 November 2020)
• Arm: Mali GPU reference sheet (PDF, 2022)
• Arm: Mali and Immortalis GPU product overview
• Arm: Mali-G720 product page
• Arm: Mali-G77 product page
• Arm: Mali G1-Ultra product page
• Arm: Transaction Elimination
• Arm: Arm Frame Buffer Compression
• Khronos Vulkan Samples (Arm authored): Render passes
• Khronos Vulkan Samples (Arm authored): Subpasses
• Khronos Vulkan Samples (Arm authored): MSAA
• Khronos Vulkan Samples (Arm authored): Pipeline barriers and tile-based rendering
• Qualcomm: Adreno GPU overview
• Qualcomm: Adreno GPU on mobile, best practices
• Imagination: PowerVR architecture overview for developers
• Imagination: Tile-Based Deferred Rendering
• Imagination: Vertex Processing (Tiler)
• Imagination: Per-Tile Rasterization (Renderer)
• Imagination: Hidden Surface Removal Efficiency
• Imagination: Further TBDR Details
• Imagination: History of PowerVR
• Imagination: Do Not Use Depth Pre-pass
• Khronos: VkAttachmentLoadOp reference page
• Khronos: VK_KHR_dynamic_rendering_local_read reference page
• Khronos: VK_QCOM_tile_shading reference page

Last reviewed: 25 October 2026. GPU families, extension availability and vendor guidance change with each hardware and driver release; check the linked vendor documentation for the device you target.