MetalEngine demo

Tests/CallUITest/BUILD

@@ -41,7 +41,9 @@ swift_library(
         ":AppResources",
     ],
     deps = [
-        "//submodules/TelegramUI/Components/Calls/CallScreen",
+        "//submodules/Display",
+        "//submodules/MetalEngine",
+        "//submodules/TelegramUI/Components/DustEffect",
     ],
 )
 

Tests/CallUITest/Sources/ViewController.swift

@@ -1,19 +1,47 @@
 import Foundation
 import UIKit
-import CallScreen
+import MetalEngine
+import Display
+import DustEffect
 
 public final class ViewController: UIViewController {
+    private var dustLayer: DustEffectLayer?
+    
     override public func viewDidLoad() {
         super.viewDidLoad()
         
-        let privateCallScreen = PrivateCallScreen(frame: CGRect())
+        self.view.layer.addSublayer(MetalEngine.shared.rootLayer)
-        self.view.addSubview(privateCallScreen)
+        MetalEngine.shared.rootLayer.frame = CGRect(origin: CGPoint(x: 0.0, y: -101.0), size: CGSize(width: 100.0, height: 100.0))
         
-        privateCallScreen.frame = self.view.bounds
+        self.reset()
-        privateCallScreen.update(size: self.view.bounds.size, insets: UIEdgeInsets(top: 44.0, left: 0.0, bottom: 0.0, right: 0.0))
         
-        let context = MetalContext.shared
+        self.view.addGestureRecognizer(UITapGestureRecognizer(target: self, action: #selector(self.imageTap(_:))))
-        self.view.layer.addSublayer(context.rootLayer)
+        
-        context.rootLayer.frame = CGRect(origin: CGPoint(x: 0.0, y: -101.0), size: CGSize(width: 100.0, height: 100.0))
+        self.view.backgroundColor = .white
+        
+        SharedDisplayLinkDriver.shared.updateForegroundState(true)
+    }
+    
+    func reset() {
+        self.dustLayer?.removeFromSuperlayer()
+        
+        let dustLayer = DustEffectLayer()
+        self.dustLayer = dustLayer
+        dustLayer.frame = self.view.bounds
+        
+        self.view.layer.addSublayer(dustLayer)
+    }
+    
+    @objc private func imageTap(_ recognizer: UITapGestureRecognizer) {
+        if case .ended = recognizer.state {
+            guard let dustLayer else {
+                return
+            }
+            
+            let image = UIImage(named: "test")!
+            let itemSize = CGSize(width: 200.0, height: 200.0)
+            let itemFrame = CGRect(origin: CGPoint(x: floor((self.view.bounds.width - itemSize.width) * 0.5), y: floor((self.view.bounds.height - itemSize.height) * 0.5)), size: itemSize)
+            dustLayer.addItem(frame: itemFrame, image: image)
+        }
     }
 }

submodules/MetalEngine/Sources/MetalEngine.swift

@@ -136,14 +136,22 @@ open class MetalEngineSubjectLayer: SimpleLayer {
     }
 }
 
+protocol MetalEngineResource: AnyObject {
+    func free()
+}
+
 public final class PooledTexture {
-    final class Texture {
+    final class Texture: MetalEngineResource {
         let value: MTLTexture
         var isInUse: Bool = false
         
         init(value: MTLTexture) {
             self.value = value
         }
+        
+        public func free() {
+            self.isInUse = false
+        }
     }
     
     public let spec: TextureSpec
@@ -187,6 +195,86 @@ public final class PooledTexture {
     }
 }
 
+public struct BufferSpec: Equatable {
+    public var length: Int
+    
+    public init(length: Int) {
+        self.length = length
+    }
+}
+
+public final class BufferPlaceholder {
+    public let placeholer: Placeholder<MTLBuffer?>
+    public let spec: BufferSpec
+    
+    init(placeholer: Placeholder<MTLBuffer?>, spec: BufferSpec) {
+        self.placeholer = placeholer
+        self.spec = spec
+    }
+}
+
+public final class PooledBuffer {
+    final class Buffer: MetalEngineResource {
+        let value: MTLBuffer
+        var isInUse: Bool = false
+        
+        init(value: MTLBuffer) {
+            self.value = value
+        }
+        
+        public func free() {
+            self.isInUse = false
+        }
+    }
+    
+    public let spec: BufferSpec
+    
+    private let buffers: [Buffer]
+    
+    init(device: MTLDevice, spec: BufferSpec) {
+        self.spec = spec
+        
+        self.buffers = (0 ..< 3).compactMap { _ -> Buffer? in
+            guard let texture = device.makeBuffer(length: spec.length, options: [.storageModePrivate]) else {
+                return nil
+            }
+            return Buffer(value: texture)
+        }
+    }
+    
+    public func get(context: MetalEngineSubjectContext) -> BufferPlaceholder? {
+        #if DEBUG
+        if context.freeResourcesOnCompletion.contains(where: { $0 === self }) {
+            assertionFailure("Trying to get PooledTexture more than once per update cycle")
+        }
+        #endif
+        
+        for buffer in self.buffers {
+            if !buffer.isInUse {
+                buffer.isInUse = true
+                let placeholder = Placeholder<MTLBuffer?>()
+                placeholder.contents = buffer.value
+                context.freeResourcesOnCompletion.append(buffer)
+                return BufferPlaceholder(placeholer: placeholder, spec: self.spec)
+            }
+        }
+        
+        print("PooledBuffer: all textures are in use")
+        return nil
+    }
+}
+
+public final class SharedBuffer {
+    public let buffer: MTLBuffer
+    
+    init?(device: MTLDevice, spec: BufferSpec) {
+        guard let buffer = device.makeBuffer(length: spec.length, options: [.storageModeShared]) else {
+            return nil
+        }
+        self.buffer = buffer
+    }
+}
+
 public final class MetalEngineSubjectContext {
     fileprivate final class ComputeOperation {
         let commands: (MTLCommandBuffer) -> Void
@@ -220,7 +308,7 @@ public final class MetalEngineSubjectContext {
     
     fileprivate var computeOperations: [ComputeOperation] = []
     fileprivate var renderToLayerOperationsGroupedByState: [ObjectIdentifier: [RenderToLayerOperation]] = [:]
-    fileprivate var freeResourcesOnCompletion: [PooledTexture.Texture] = []
+    fileprivate var freeResourcesOnCompletion: [MetalEngineResource] = []
     
     fileprivate init(device: MTLDevice, impl: MetalEngine.Impl) {
         self.device = device
@@ -320,16 +408,6 @@ public final class MetalEngineSubjectContext {
             return commands(commandBuffer, state)
         })
     }
-    
-    public func temporaryTexture(spec: TextureSpec) -> TexturePlaceholder {
-        let descriptor = MTLTextureDescriptor.texture2DDescriptor(pixelFormat: spec.pixelFormat.metalFormat, width: spec.width, height: spec.height, mipmapped: false)
-        descriptor.storageMode = .private
-        descriptor.usage = [.shaderRead, .shaderWrite]
-        
-        let placeholder = Placeholder<MTLTexture?>()
-        placeholder.contents = self.impl.device.makeTexture(descriptor: descriptor)
-        return TexturePlaceholder(placeholer: placeholder, spec: spec)
-    }
 }
 
 public final class MetalEngineSubjectInternalData {
@@ -854,6 +932,8 @@ public final class MetalEngine {
                                     continue
                                 }
                                 
+                                let subRect = surfaceAllocation.effectivePhase.subRect
+                                renderEncoder.setScissorRect(MTLScissorRect(x: Int(subRect.minX), y: Int(subRect.minY), width: Int(subRect.width), height: Int(subRect.height)))
                                 renderToLayerOperation.commands(renderEncoder, RenderLayerPlacement(effectiveRect: surfaceAllocation.effectivePhase.renderingRect))
                             }
                         }
@@ -868,7 +948,7 @@ public final class MetalEngine {
                 commandBuffer.addCompletedHandler { _ in
                     DispatchQueue.main.async {
                         for resource in freeResourcesOnCompletion {
-                            resource.isInUse = false
+                            resource.free()
                         }
                     }
                 }
@@ -916,4 +996,12 @@ public final class MetalEngine {
     public func pooledTexture(spec: TextureSpec) -> PooledTexture {
         return PooledTexture(device: self.device, spec: spec)
     }
+    
+    public func pooledBuffer(spec: BufferSpec) -> PooledBuffer {
+        return PooledBuffer(device: self.device, spec: spec)
+    }
+    
+    public func sharedBuffer(spec: BufferSpec) -> SharedBuffer? {
+        return SharedBuffer(device: self.device, spec: spec)
+    }
 }

submodules/TelegramUI/Components/DustEffect/BUILD

@@ -0,0 +1,63 @@
+load("@build_bazel_rules_swift//swift:swift.bzl", "swift_library")
+
+load(
+    "@build_bazel_rules_apple//apple:resources.bzl",
+    "apple_resource_bundle",
+    "apple_resource_group",
+)
+load("//build-system/bazel-utils:plist_fragment.bzl",
+    "plist_fragment",
+)
+
+filegroup(
+    name = "DustEffectMetalSources",
+    srcs = glob([
+        "Metal/**/*.metal",
+    ]),
+    visibility = ["//visibility:public"],
+)
+
+plist_fragment(
+    name = "DustEffectMetalSourcesBundleInfoPlist",
+    extension = "plist",
+    template =
+    """
+    <key>CFBundleIdentifier</key>
+    <string>org.telegram.DustEffectMetalSources</string>
+    <key>CFBundleDevelopmentRegion</key>
+    <string>en</string>
+    <key>CFBundleName</key>
+    <string>DustEffect</string>
+    """
+)
+
+apple_resource_bundle(
+    name = "DustEffectMetalSourcesBundle",
+    infoplists = [
+        ":DustEffectMetalSourcesBundleInfoPlist",
+    ],
+    resources = [
+        ":DustEffectMetalSources",
+    ],
+)
+
+swift_library(
+    name = "DustEffect",
+    module_name = "DustEffect",
+    srcs = glob([
+        "Sources/**/*.swift",
+    ]),
+    copts = [
+        "-warnings-as-errors",
+    ],
+    data = [
+        ":DustEffectMetalSourcesBundle",
+    ],
+    deps = [
+        "//submodules/Display",
+        "//submodules/MetalEngine",
+    ],
+    visibility = [
+        "//visibility:public",
+    ],
+)

submodules/TelegramUI/Components/DustEffect/Metal/DustEffectShaders.metal

@@ -0,0 +1,129 @@
+#include <metal_stdlib>
+
+#include "loki_header.metal"
+
+using namespace metal;
+
+struct Rectangle {
+    float2 origin;
+    float2 size;
+};
+
+constant static float2 quadVertices[6] = {
+    float2(0.0, 0.0),
+    float2(1.0, 0.0),
+    float2(0.0, 1.0),
+    float2(1.0, 0.0),
+    float2(0.0, 1.0),
+    float2(1.0, 1.0)
+};
+
+struct QuadVertexOut {
+    float4 position [[position]];
+    float2 uv;
+    float alpha;
+};
+
+float2 mapLocalToScreenCoordinates(const device Rectangle &rect, const device float2 &size, float2 position) {
+    float2 result = float2(rect.origin.x + position.x / size.x * rect.size.x, rect.origin.y + position.y / size.y * rect.size.y);
+    result.x = -1.0 + result.x * 2.0;
+    result.y = -1.0 + result.y * 2.0;
+    
+    return result;
+}
+
+struct Particle {
+    packed_float2 offsetFromBasePosition;
+    packed_float2 velocity;
+    float lifetime;
+};
+
+float particleEaseInWindowFunction(float t) {
+    return t;
+}
+
+float particleEaseInValueAt(float fraction, float t) {
+    float windowSize = 0.8;
+
+    float effectiveT = t;
+    float windowStartOffset = -windowSize;
+    float windowEndOffset = 1.0;
+
+    float windowPosition = (1.0 - fraction) * windowStartOffset + fraction * windowEndOffset;
+    float windowT = max(0.0, min(windowSize, effectiveT - windowPosition)) / windowSize;
+    float localT = 1.0 - particleEaseInWindowFunction(windowT);
+
+    return localT;
+}
+
+kernel void dustEffectUpdateParticle(
+    device Particle *particles [[ buffer(0) ]],
+    const device uint2 &size [[ buffer(1) ]],
+    const device float &phase [[ buffer(2) ]],
+    const device float &timeStep [[ buffer(3) ]],
+    uint gid [[ thread_position_in_grid ]]
+) {
+    uint count = size.x * size.y;
+    if (gid >= count) {
+        return;
+    }
+    
+    constexpr float easeInDuration = 0.8;
+    float effectFraction = max(0.0, min(easeInDuration, phase)) / easeInDuration;
+    
+    uint particleX = gid % size.x;
+    float particleXFraction = float(particleX) / float(size.x);
+    float particleFraction = particleEaseInValueAt(effectFraction, particleXFraction);
+    
+    //Loki rng = Loki(gid, uint(phase * timeStep));
+    //float2 offsetNorm = float2(1.0, 1.0) * 10.0 * timeStep;
+    
+    Particle particle = particles[gid];
+    particle.offsetFromBasePosition += (particle.velocity * timeStep) * particleFraction;
+    //particle.velocity += ((-offsetNorm) * 0.5 + float2(rng.rand(), rng.rand()) * offsetNorm) * particleFraction;
+    //particle.velocity = particle.velocity * (1.0 - particleFraction) + particle.velocity * 1.001 * particleFraction;
+    particle.velocity += float2(0.0, timeStep * 120.0) * particleFraction;
+    particle.lifetime = max(0.0, particle.lifetime - timeStep * particleFraction);
+    particles[gid] = particle;
+}
+
+vertex QuadVertexOut dustEffectVertex(
+    const device Rectangle &rect [[ buffer(0) ]],
+    const device float2 &size [[ buffer(1) ]],
+    const device uint2 &particleResolution [[ buffer(2) ]],
+    const device Particle *particles [[ buffer(3) ]],
+    unsigned int vid [[ vertex_id ]],
+    unsigned int particleId [[ instance_id ]]
+) {
+    QuadVertexOut out;
+    
+    float2 quadVertex = quadVertices[vid];
+    
+    uint particleIndexX = particleId % particleResolution.x;
+    uint particleIndexY = particleId / particleResolution.x;
+    
+    Particle particle = particles[particleId];
+    
+    float2 particleSize = size / float2(particleResolution);
+    
+    float2 topLeftPosition = float2(float(particleIndexX) * particleSize.x, float(particleIndexY) * particleSize.y);
+    out.uv = (topLeftPosition + quadVertex * particleSize) / size;
+    
+    topLeftPosition += particle.offsetFromBasePosition;
+    float2 position = topLeftPosition + quadVertex * particleSize;
+    
+    out.position = float4(mapLocalToScreenCoordinates(rect, size, position), 0.0, 1.0);
+    out.alpha = max(0.0, min(0.3, particle.lifetime) / 0.3);
+    
+    return out;
+}
+
+fragment half4 dustEffectFragment(
+    QuadVertexOut in [[stage_in]],
+    texture2d<half, access::sample> inTexture [[ texture(0) ]]
+) {
+    constexpr sampler sampler(coord::normalized, address::clamp_to_edge, filter::linear);
+    
+    half3 color = inTexture.sample(sampler, float2(in.uv.x, 1.0 - in.uv.y)).rgb;
+    return half4(color * in.alpha, in.alpha);
+}

submodules/TelegramUI/Components/DustEffect/Metal/loki.metal

@@ -0,0 +1,54 @@
+#include <metal_stdlib>
+#include "loki_header.metal"
+
+unsigned Loki::TausStep(const unsigned z, const int s1, const int s2, const int s3, const unsigned M)
+{
+    unsigned b=(((z << s1) ^ z) >> s2);
+    return (((z & M) << s3) ^ b);
+}
+
+thread Loki::Loki(const unsigned seed1, const unsigned seed2, const unsigned seed3) {
+    unsigned seed = seed1 * 1099087573UL;
+    unsigned seedb = seed2 * 1099087573UL;
+    unsigned seedc = seed3 * 1099087573UL;
+
+    // Round 1: Randomise seed
+    unsigned z1 = TausStep(seed,13,19,12,429496729UL);
+    unsigned z2 = TausStep(seed,2,25,4,4294967288UL);
+    unsigned z3 = TausStep(seed,3,11,17,429496280UL);
+    unsigned z4 = (1664525*seed + 1013904223UL);
+
+    // Round 2: Randomise seed again using second seed
+    unsigned r1 = (z1^z2^z3^z4^seedb);
+
+    z1 = TausStep(r1,13,19,12,429496729UL);
+    z2 = TausStep(r1,2,25,4,4294967288UL);
+    z3 = TausStep(r1,3,11,17,429496280UL);
+    z4 = (1664525*r1 + 1013904223UL);
+
+    // Round 3: Randomise seed again using third seed
+    r1 = (z1^z2^z3^z4^seedc);
+
+    z1 = TausStep(r1,13,19,12,429496729UL);
+    z2 = TausStep(r1,2,25,4,4294967288UL);
+    z3 = TausStep(r1,3,11,17,429496280UL);
+    z4 = (1664525*r1 + 1013904223UL);
+
+    this->seed = (z1^z2^z3^z4) * 2.3283064365387e-10;
+}
+
+thread float Loki::rand() {
+    unsigned hashed_seed = this->seed * 1099087573UL;
+
+    unsigned z1 = TausStep(hashed_seed,13,19,12,429496729UL);
+    unsigned z2 = TausStep(hashed_seed,2,25,4,4294967288UL);
+    unsigned z3 = TausStep(hashed_seed,3,11,17,429496280UL);
+    unsigned z4 = (1664525*hashed_seed + 1013904223UL);
+
+    thread float old_seed = this->seed;
+
+    this->seed = (z1^z2^z3^z4) * 2.3283064365387e-10;
+
+    return old_seed;
+}
+

submodules/TelegramUI/Components/DustEffect/Metal/loki_header.metal

@@ -0,0 +1,45 @@
+/*
+ * Loki Random Number Generator
+ * Copyright (c) 2017 Youssef Victor All rights reserved.
+ *
+ *      Function                        Result
+ *      ------------------------------------------------------------------
+ *
+ *      TausStep                        Combined Tausworthe Generator or
+ *                                      Linear Feedback Shift Register (LFSR)
+ *                                      random number generator. This is a
+ *                                      helper method for rng, which uses
+ *                                      a hybrid approach combining LFSR with
+ *                                      a Linear Congruential Generator (LCG)
+ *                                      in order to produce random numbers with
+ *                                      periods of well over 2^121
+ *
+ *      rand                            A pseudo-random number based on the
+ *                                      method outlined in "Efficient
+ *                                      pseudo-random number generation
+ *                                      for monte-carlo simulations using
+ *                                      graphic processors" by Siddhant
+ *                                      Mohanty et al 2012.
+ *
+ */
+
+#include <metal_stdlib>
+using namespace metal;
+
+#ifndef LOKI
+#define LOKI
+
+
+class Loki {
+private:
+    thread float seed;
+    unsigned TausStep(const unsigned z, const int s1, const int s2, const int s3, const unsigned M);
+
+public:
+    thread Loki(const unsigned seed1, const unsigned seed2 = 1, const unsigned seed3 = 1);
+
+    thread float rand();
+};
+
+#endif
+

submodules/TelegramUI/Components/DustEffect/Sources/DustEffectLayer.swift

@@ -0,0 +1,274 @@
+import Foundation
+import UIKit
+import Display
+import MetalEngine
+import MetalKit
+
+private final class BundleMarker: NSObject {
+}
+
+private var metalLibraryValue: MTLLibrary?
+func metalLibrary(device: MTLDevice) -> MTLLibrary? {
+    if let metalLibraryValue {
+        return metalLibraryValue
+    }
+    
+    let mainBundle = Bundle(for: BundleMarker.self)
+    guard let path = mainBundle.path(forResource: "DustEffectMetalSourcesBundle", ofType: "bundle") else {
+        return nil
+    }
+    guard let bundle = Bundle(path: path) else {
+        return nil
+    }
+    guard let library = try? device.makeDefaultLibrary(bundle: bundle) else {
+        return nil
+    }
+    
+    metalLibraryValue = library
+    return library
+}
+
+public final class DustEffectLayer: MetalEngineSubjectLayer, MetalEngineSubject {
+    public var internalData: MetalEngineSubjectInternalData?
+    
+    private final class Item {
+        let frame: CGRect
+        let texture: MTLTexture
+        
+        var phase: Float = 0
+        var particleBuffer: SharedBuffer?
+        
+        init?(frame: CGRect, image: UIImage) {
+            self.frame = frame
+            
+            guard let cgImage = image.cgImage, let texture = try? MTKTextureLoader(device: MetalEngine.shared.device).newTexture(cgImage: cgImage, options: [.SRGB: false as NSNumber]) else {
+                return nil
+            }
+            self.texture = texture
+        }
+    }
+    
+    private final class RenderState: RenderToLayerState {
+        let pipelineState: MTLRenderPipelineState
+        
+        init?(device: MTLDevice) {
+            guard let library = metalLibrary(device: device) else {
+                return nil
+            }
+            guard let vertexFunction = library.makeFunction(name: "dustEffectVertex"), let fragmentFunction = library.makeFunction(name: "dustEffectFragment") else {
+                return nil
+            }
+            
+            let pipelineDescriptor = MTLRenderPipelineDescriptor()
+            pipelineDescriptor.vertexFunction = vertexFunction
+            pipelineDescriptor.fragmentFunction = fragmentFunction
+            pipelineDescriptor.colorAttachments[0].pixelFormat = .bgra8Unorm
+            pipelineDescriptor.colorAttachments[0].isBlendingEnabled = true
+            pipelineDescriptor.colorAttachments[0].rgbBlendOperation = .add
+            pipelineDescriptor.colorAttachments[0].alphaBlendOperation = .add
+            pipelineDescriptor.colorAttachments[0].sourceRGBBlendFactor = .one
+            pipelineDescriptor.colorAttachments[0].sourceAlphaBlendFactor = .one
+            pipelineDescriptor.colorAttachments[0].destinationRGBBlendFactor = .oneMinusSourceAlpha
+            pipelineDescriptor.colorAttachments[0].destinationAlphaBlendFactor = .one
+            guard let pipelineState = try? device.makeRenderPipelineState(descriptor: pipelineDescriptor) else {
+                return nil
+            }
+            self.pipelineState = pipelineState
+        }
+    }
+    
+    final class DustComputeState: ComputeState {
+        let computePipelineStateUpdateParticle: MTLComputePipelineState
+        
+        required init?(device: MTLDevice) {
+            guard let library = metalLibrary(device: device) else {
+                return nil
+            }
+            guard let functionDustEffectUpdateParticle = library.makeFunction(name: "dustEffectUpdateParticle") else {
+                return nil
+            }
+            guard let computePipelineStateUpdateParticle = try? device.makeComputePipelineState(function: functionDustEffectUpdateParticle) else {
+                return nil
+            }
+            self.computePipelineStateUpdateParticle = computePipelineStateUpdateParticle
+        }
+    }
+    
+    private var updateLink: SharedDisplayLinkDriver.Link?
+    
+    private var items: [Item] = []
+    
+    override public init() {
+        super.init()
+        
+        self.isOpaque = false
+        self.backgroundColor = nil
+        
+        self.didEnterHierarchy = { [weak self] in
+            guard let self else {
+                return
+            }
+            self.updateNeedsAnimation()
+        }
+        self.didExitHierarchy = { [weak self] in
+            guard let self else {
+                return
+            }
+            self.updateNeedsAnimation()
+        }
+    }
+    
+    override public init(layer: Any) {
+        super.init(layer: layer)
+    }
+    
+    required public init?(coder: NSCoder) {
+        fatalError("init(coder:) has not been implemented")
+    }
+    
+    private func updateItems() {
+        for i in (0 ..< self.items.count).reversed() {
+            self.items[i].phase += 1.0 / 60.0
+            
+            if self.items[i].phase >= 4.0 {
+                self.items.remove(at: i)
+            }
+        }
+        self.updateNeedsAnimation()
+    }
+    
+    private func updateNeedsAnimation() {
+        if !self.items.isEmpty && self.isInHierarchy {
+            if self.updateLink == nil {
+                self.updateLink = SharedDisplayLinkDriver.shared.add { [weak self] in
+                    guard let self else {
+                        return
+                    }
+                    self.updateItems()
+                    self.setNeedsUpdate()
+                }
+            }
+        } else {
+            if self.updateLink != nil {
+                self.updateLink = nil
+            }
+        }
+    }
+    
+    public func addItem(frame: CGRect, image: UIImage) {
+        if let item = Item(frame: frame, image: image) {
+            self.items.append(item)
+            self.updateNeedsAnimation()
+        }
+    }
+    
+    public func update(context: MetalEngineSubjectContext) {
+        if self.bounds.isEmpty {
+            return
+        }
+        
+        let containerSize = self.bounds.size
+        
+        for item in self.items {
+            var itemFrame = item.frame
+            itemFrame.origin.y = containerSize.height - itemFrame.maxY
+            
+            let particleColumnCount = Int(itemFrame.width)
+            let particleRowCount = Int(itemFrame.height)
+            let particleCount = particleColumnCount * particleRowCount
+            
+            if item.particleBuffer == nil {
+                if let particleBuffer = MetalEngine.shared.sharedBuffer(spec: BufferSpec(length: particleCount * 4 * (4 + 1))) {
+                    item.particleBuffer = particleBuffer
+                    
+                    let particles = particleBuffer.buffer.contents().assumingMemoryBound(to: Float.self)
+                    for i in 0 ..< particleCount {
+                        particles[i * 5 + 0] = 0.0;
+                        particles[i * 5 + 1] = 0.0;
+                        
+                        let direction = Float.random(in: 0.0 ..< Float.pi * 2.0)
+                        let velocity = Float.random(in: 0.1 ... 0.2) * 420.0
+                        particles[i * 5 + 2] = cos(direction) * velocity
+                        particles[i * 5 + 3] = sin(direction) * velocity
+                        
+                        particles[i * 5 + 4] = Float.random(in: 0.7 ... 1.5)
+                    }
+                }
+            }
+        }
+        
+        let _ = context.compute(state: DustComputeState.self, commands: { [weak self] commandBuffer, state in
+            guard let self else {
+                return
+            }
+            guard let computeEncoder = commandBuffer.makeComputeCommandEncoder() else {
+                return
+            }
+            
+            for item in self.items {
+                guard let particleBuffer = item.particleBuffer else {
+                    continue
+                }
+                let itemFrame = item.frame
+                let particleColumnCount = Int(itemFrame.width)
+                let particleRowCount = Int(itemFrame.height)
+                
+                let threadgroupSize = MTLSize(width: 32, height: 1, depth: 1)
+                let threadgroupCount = MTLSize(width: (particleRowCount * particleColumnCount + threadgroupSize.width - 1) / threadgroupSize.width, height: 1, depth: 1)
+                
+                computeEncoder.setComputePipelineState(state.computePipelineStateUpdateParticle)
+                computeEncoder.setBuffer(particleBuffer.buffer, offset: 0, index: 0)
+                var particleCount = SIMD2<UInt32>(UInt32(particleColumnCount), UInt32(particleRowCount))
+                computeEncoder.setBytes(&particleCount, length: 4 * 2, index: 1)
+                var phase = item.phase
+                computeEncoder.setBytes(&phase, length: 4, index: 2)
+                var timeStep: Float = 1.0 / 60.0
+                computeEncoder.setBytes(&timeStep, length: 4, index: 3)
+                computeEncoder.dispatchThreadgroups(threadgroupCount, threadsPerThreadgroup: threadgroupSize)
+            }
+            
+            computeEncoder.endEncoding()
+        })
+        
+        context.renderToLayer(spec: RenderLayerSpec(size: RenderSize(width: Int(self.bounds.width * 3.0), height: Int(self.bounds.height * 3.0))), state: RenderState.self, layer: self, commands: { [weak self] encoder, placement in
+            guard let self else {
+                return
+            }
+            
+            for item in self.items {
+                guard let particleBuffer = item.particleBuffer else {
+                    continue
+                }
+                
+                var itemFrame = item.frame
+                itemFrame.origin.y = containerSize.height - itemFrame.maxY
+                
+                let particleColumnCount = Int(itemFrame.width)
+                let particleRowCount = Int(itemFrame.height)
+                let particleCount = particleColumnCount * particleRowCount
+                
+                var effectiveRect = placement.effectiveRect
+                effectiveRect.origin.x += itemFrame.minX / containerSize.width * effectiveRect.width
+                effectiveRect.origin.y += itemFrame.minY / containerSize.height * effectiveRect.height
+                effectiveRect.size.width = itemFrame.width / containerSize.width * effectiveRect.width
+                effectiveRect.size.height = itemFrame.height / containerSize.height * effectiveRect.height
+                
+                var rect = SIMD4<Float>(Float(effectiveRect.minX), Float(effectiveRect.minY), Float(effectiveRect.width), Float(effectiveRect.height))
+                encoder.setVertexBytes(&rect, length: 4 * 4, index: 0)
+                
+                var size = SIMD2<Float>(Float(itemFrame.width), Float(itemFrame.height))
+                encoder.setVertexBytes(&size, length: 4 * 2, index: 1)
+                
+                var particleResolution = SIMD2<UInt32>(UInt32(particleColumnCount), UInt32(particleRowCount))
+                encoder.setVertexBytes(&particleResolution, length: 4 * 2, index: 2)
+                
+                encoder.setVertexBuffer(particleBuffer.buffer, offset: 0, index: 3)
+                
+                encoder.setFragmentTexture(item.texture, index: 0)
+                
+                encoder.drawPrimitives(type: .triangle, vertexStart: 0, vertexCount: 6)
+                encoder.drawPrimitives(type: .triangle, vertexStart: 0, vertexCount: 6, instanceCount: particleCount)
+            }
+        })
+    }
+}