Updated for Swift 5 and Metal 3
This commit is contained in:
parent
0264cd110e
commit
d3ac3f86f3
|
|
@ -0,0 +1,103 @@
|
||||||
|
|
||||||
|
import MetalKit
|
||||||
|
import PlaygroundSupport
|
||||||
|
|
||||||
|
guard let device = MTLCreateSystemDefaultDevice(),
|
||||||
|
let commandQueue = device.makeCommandQueue() else {
|
||||||
|
fatalError("Metal is not supported on this device")
|
||||||
|
}
|
||||||
|
|
||||||
|
let textureDescriptor = MTLTextureDescriptor.texture2DDescriptor(pixelFormat: .rgba8Unorm,
|
||||||
|
width: 600,
|
||||||
|
height: 600,
|
||||||
|
mipmapped: false)
|
||||||
|
textureDescriptor.usage = [.shaderWrite, .shaderRead]
|
||||||
|
guard let texture = device.makeTexture(descriptor: textureDescriptor) else {
|
||||||
|
fatalError("Failed to create texture")
|
||||||
|
}
|
||||||
|
|
||||||
|
let shaderSource = """
|
||||||
|
#include <metal_stdlib>
|
||||||
|
using namespace metal;
|
||||||
|
|
||||||
|
float dist(float2 point, float2 center, float radius)
|
||||||
|
{
|
||||||
|
return length(point - center) - radius;
|
||||||
|
}
|
||||||
|
|
||||||
|
kernel void compute(texture2d<float, access::write> output [[texture(0)]],
|
||||||
|
uint2 gid [[thread_position_in_grid]])
|
||||||
|
{
|
||||||
|
int width = output.get_width();
|
||||||
|
int height = output.get_height();
|
||||||
|
float2 uv = float2(gid) / float2(width, height);
|
||||||
|
uv = uv * 2.0 - 1.0;
|
||||||
|
float distToCircle = dist(uv, float2(0), 0.5);
|
||||||
|
float distToCircle2 = dist(uv, float2(-0.1, 0.1), 0.5);
|
||||||
|
bool inside = distToCircle2 < 0;
|
||||||
|
output.write(inside ? float4(0) : float4(1, 0.7, 0, 1) * (1 - distToCircle), gid);
|
||||||
|
}
|
||||||
|
"""
|
||||||
|
|
||||||
|
guard let library = try? device.makeLibrary(source: shaderSource, options: nil),
|
||||||
|
let kernelFunction = library.makeFunction(name: "compute"),
|
||||||
|
let computePipelineState = try? device.makeComputePipelineState(function: kernelFunction) else {
|
||||||
|
fatalError("Failed to create compute pipeline")
|
||||||
|
}
|
||||||
|
|
||||||
|
guard let commandBuffer = commandQueue.makeCommandBuffer(),
|
||||||
|
let computeEncoder = commandBuffer.makeComputeCommandEncoder() else {
|
||||||
|
fatalError("Failed to create command buffer or compute encoder")
|
||||||
|
}
|
||||||
|
|
||||||
|
let w = computePipelineState.threadExecutionWidth
|
||||||
|
let h = computePipelineState.maxTotalThreadsPerThreadgroup / w
|
||||||
|
let threadsPerThreadgroup = MTLSize(width: w, height: h, depth: 1)
|
||||||
|
let threadgroupsPerGrid = MTLSize(width: (texture.width + w - 1) / w,
|
||||||
|
height: (texture.height + h - 1) / h,
|
||||||
|
depth: 1)
|
||||||
|
computeEncoder.setComputePipelineState(computePipelineState)
|
||||||
|
computeEncoder.setTexture(texture, index: 0)
|
||||||
|
computeEncoder.dispatchThreadgroups(threadgroupsPerGrid, threadsPerThreadgroup: threadsPerThreadgroup)
|
||||||
|
computeEncoder.endEncoding()
|
||||||
|
commandBuffer.commit()
|
||||||
|
commandBuffer.waitUntilCompleted()
|
||||||
|
|
||||||
|
class MetalView: MTKView {
|
||||||
|
var texture: MTLTexture!
|
||||||
|
|
||||||
|
init(frame frameRect: CGRect, device: MTLDevice?, texture: MTLTexture?) {
|
||||||
|
self.texture = texture
|
||||||
|
super.init(frame: frameRect, device: device)
|
||||||
|
self.isPaused = true
|
||||||
|
self.enableSetNeedsDisplay = true
|
||||||
|
self.drawableSize = CGSize(width: 600, height: 600) // Set the drawable size here
|
||||||
|
}
|
||||||
|
|
||||||
|
required init(coder: NSCoder) { fatalError("init(coder:) has not been implemented") }
|
||||||
|
|
||||||
|
override func draw(_ rect: CGRect) {
|
||||||
|
guard let drawable = self.currentDrawable else { return }
|
||||||
|
if drawable.texture.width != texture.width || drawable.texture.height != texture.height {
|
||||||
|
print("Drawable texture size does not match Metal texture size.")
|
||||||
|
return
|
||||||
|
}
|
||||||
|
guard let commandBuffer = device?.makeCommandQueue()?.makeCommandBuffer(),
|
||||||
|
let blitEncoder = commandBuffer.makeBlitCommandEncoder() else { return }
|
||||||
|
blitEncoder.copy(from: texture,
|
||||||
|
sourceSlice: 0,
|
||||||
|
sourceLevel: 0,
|
||||||
|
sourceOrigin: MTLOrigin(x: 0, y: 0, z: 0),
|
||||||
|
sourceSize: MTLSize(width: texture.width, height: texture.height, depth: 1),
|
||||||
|
to: drawable.texture,
|
||||||
|
destinationSlice: 0,
|
||||||
|
destinationLevel: 0,
|
||||||
|
destinationOrigin: MTLOrigin(x: 0, y: 0, z: 0))
|
||||||
|
blitEncoder.endEncoding()
|
||||||
|
commandBuffer.present(drawable)
|
||||||
|
commandBuffer.commit()
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
let view = MetalView(frame: CGRect(x: 0, y: 0, width: 600, height: 600), device: device, texture: texture)
|
||||||
|
PlaygroundPage.current.liveView = view
|
||||||
|
|
@ -0,0 +1,21 @@
|
||||||
|
|
||||||
|
#include <metal_stdlib>
|
||||||
|
using namespace metal;
|
||||||
|
|
||||||
|
float dist(float2 point, float2 center, float radius)
|
||||||
|
{
|
||||||
|
return length(point - center) - radius;
|
||||||
|
}
|
||||||
|
|
||||||
|
kernel void compute(texture2d<float, access::write> output [[texture(0)]],
|
||||||
|
uint2 gid [[thread_position_in_grid]])
|
||||||
|
{
|
||||||
|
int width = output.get_width();
|
||||||
|
int height = output.get_height();
|
||||||
|
float2 uv = float2(gid) / float2(width, height);
|
||||||
|
uv = uv * 2.0 - 1.0;
|
||||||
|
float distToCircle = dist(uv, float2(0), 0.5);
|
||||||
|
float distToCircle2 = dist(uv, float2(-0.1, 0.1), 0.5);
|
||||||
|
bool inside = distToCircle2 < 0;
|
||||||
|
output.write(inside ? float4(0) : float4(1, 0.7, 0, 1) * (1 - distToCircle), gid);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,48 @@
|
||||||
|
|
||||||
|
import MetalKit
|
||||||
|
|
||||||
|
public class MetalView: NSObject, MTKViewDelegate {
|
||||||
|
|
||||||
|
public var device: MTLDevice!
|
||||||
|
var queue: MTLCommandQueue!
|
||||||
|
var cps: MTLComputePipelineState!
|
||||||
|
|
||||||
|
override public init() {
|
||||||
|
super.init()
|
||||||
|
registerShaders()
|
||||||
|
}
|
||||||
|
|
||||||
|
func registerShaders() {
|
||||||
|
device = MTLCreateSystemDefaultDevice()!
|
||||||
|
queue = device.makeCommandQueue()
|
||||||
|
guard let path = Bundle.main.path(forResource: "Shaders", ofType: "metal") else {
|
||||||
|
print("No shader file found")
|
||||||
|
return
|
||||||
|
}
|
||||||
|
do {
|
||||||
|
let input = try String(contentsOfFile: path, encoding: String.Encoding.utf8)
|
||||||
|
let library = try device.makeLibrary(source: input, options: nil)
|
||||||
|
let kernel = library.makeFunction(name: "compute")!
|
||||||
|
cps = try device.makeComputePipelineState(function: kernel)
|
||||||
|
} catch let e {
|
||||||
|
Swift.print("\(e)")
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
public func mtkView(_ view: MTKView, drawableSizeWillChange size: CGSize) {}
|
||||||
|
|
||||||
|
public func draw(in view: MTKView) {
|
||||||
|
if let drawable = view.currentDrawable,
|
||||||
|
let commandBuffer = queue.makeCommandBuffer(),
|
||||||
|
let commandEncoder = commandBuffer.makeComputeCommandEncoder() {
|
||||||
|
commandEncoder.setComputePipelineState(cps)
|
||||||
|
commandEncoder.setTexture(drawable.texture, index: 0)
|
||||||
|
let threadGroupCount = MTLSizeMake(8, 8, 1)
|
||||||
|
let threadGroups = MTLSizeMake(drawable.texture.width / threadGroupCount.width, drawable.texture.height / threadGroupCount.height, 1)
|
||||||
|
commandEncoder.dispatchThreadgroups(threadGroups, threadsPerThreadgroup: threadGroupCount)
|
||||||
|
commandEncoder.endEncoding()
|
||||||
|
commandBuffer.present(drawable)
|
||||||
|
commandBuffer.commit()
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,4 @@
|
||||||
|
<?xml version="1.0" encoding="UTF-8" standalone="yes"?>
|
||||||
|
<playground version='5.0' target-platform='osx'>
|
||||||
|
<timeline fileName='timeline.xctimeline'/>
|
||||||
|
</playground>
|
||||||
|
|
@ -0,0 +1,7 @@
|
||||||
|
<?xml version="1.0" encoding="UTF-8"?>
|
||||||
|
<Workspace
|
||||||
|
version = "1.0">
|
||||||
|
<FileRef
|
||||||
|
location = "self:">
|
||||||
|
</FileRef>
|
||||||
|
</Workspace>
|
||||||
Binary file not shown.
Binary file not shown.
Loading…
Reference in New Issue