From 8f6c5eb2b627f2e59a7a66133d01cc0bda0eafe6 Mon Sep 17 00:00:00 2001 From: Eugene Date: Wed, 26 Aug 2020 19:20:08 +0300 Subject: [PATCH] Style Fixes (#98) * remove warnings * style fixes * update version * introduce several bugfixes, new APIs (#103) * introduce several bugfixes, new APIs * fix alignment * Squashed commit of the following: fix TextureWeightedMix (#104) Add Texture Mix (#100) * add texture mix * naming * update version * update encoder * replace mask reading with mask sampling * fix shared memory size * fix typo in float conversion * Texture Mask Inversion (#106) * add inversion to TextureMask * remove redundant annotation * RGBA to YCbCr (#105) * rgba to ycbcr * update rgba to ycbcr * migrate to gather * fixes * use to matrix multiplication * use matrix transpose * simplify mean calculation Co-authored-by: Andrey Volodin --- Alloy.podspec | 2 +- Alloy/Encoders/EuclideanDistance.swift | 7 +- Alloy/Encoders/LookUpTable.swift | 38 +- Alloy/Encoders/MPSUnaryImageKernels.swift | 25 +- Alloy/Encoders/MaskGuidedBlur.swift | 40 +- Alloy/Encoders/NormalizeKernel.swift | 39 +- Alloy/Encoders/RGBAToYCbCr.swift | 88 +++ Alloy/Encoders/StdMeanNormalization.swift | 39 +- Alloy/Encoders/TextureAddConstant.swift | 37 +- Alloy/Encoders/TextureAffineCrop.swift | 37 +- Alloy/Encoders/TextureCopy.swift | 81 +- .../Encoders/TextureDifferenceHighlight.swift | 54 +- Alloy/Encoders/TextureDivideByConstant.swift | 72 +- Alloy/Encoders/TextureInterpolation.swift | 52 +- Alloy/Encoders/TextureMask.swift | 58 +- Alloy/Encoders/TextureMaskedMix.swift | 65 +- Alloy/Encoders/TextureMax.swift | 43 +- Alloy/Encoders/TextureMean.swift | 43 +- Alloy/Encoders/TextureMin.swift | 43 +- Alloy/Encoders/TextureMultiplyAdd.swift | 50 +- Alloy/Encoders/TextureNormalization.swift | 36 +- Alloy/Encoders/TextureResize.swift | 35 +- Alloy/Encoders/TextureWeightedMix.swift | 8 +- Alloy/Encoders/YCbCrToRGBA.swift | 50 +- Alloy/Float16.swift | 48 +- Alloy/MPS+Functors.swift | 12 +- Alloy/MTLBlitCommandEncoder+Extensions.swift | 45 +- .../MTLComputeCommandEncoder+Extensions.swift | 18 +- Alloy/MTLContext.swift | 2 +- Alloy/MTLOrigin+Extensions.swift | 12 +- .../MTLRenderCommandEncoder+Extensions.swift | 36 +- Alloy/MTLSize+Extensions.swift | 28 + Alloy/MTLSizeAndAlign+Extensions.swift | 17 - Alloy/MTLTexture+Extensions.swift | 49 +- Alloy/Renderers/LinesRenderer.swift | 3 +- Alloy/Renderers/MaskRenderer.swift | 13 +- Alloy/Renderers/PointsRenderer.swift | 6 +- Alloy/Renderers/RectangleRenderer.swift | 6 +- Alloy/Renderers/SimpleGeometryRenderer.swift | 6 +- Alloy/Shaders/Shaders.metal | 706 ++++++++++-------- Demo/AlloyTests/AlloyTests.swift | 30 +- .../Encoders/SwitchDataFormatEncoder.swift | 36 +- Demo/AlloyTests/EuclideanDistanceTests.swift | 18 +- Demo/AlloyTests/TextureCachingTests.swift | 47 +- Demo/AlloyTests/TextureCopyTests.swift | 24 +- Demo/Demo.xcodeproj/project.pbxproj | 16 +- Demo/Demo/ViewController.swift | 6 +- Demo/Podfile.lock | 14 +- 48 files changed, 1213 insertions(+), 1027 deletions(-) create mode 100644 Alloy/Encoders/RGBAToYCbCr.swift create mode 100644 Alloy/MTLSize+Extensions.swift diff --git a/Alloy.podspec b/Alloy.podspec index 2a84e53..0fc74b9 100644 --- a/Alloy.podspec +++ b/Alloy.podspec @@ -1,6 +1,6 @@ Pod::Spec.new do |s| s.name = 'Alloy' - s.version = '0.15.1' + s.version = '0.16.0' s.license = { :type => 'MIT', :file => 'LICENSE' } s.summary = 'Nano helpers for Metal framework' s.homepage = 'https://github.com/s1ddok/Alloy' diff --git a/Alloy/Encoders/EuclideanDistance.swift b/Alloy/Encoders/EuclideanDistance.swift index c18e21a..3ec8151 100644 --- a/Alloy/Encoders/EuclideanDistance.swift +++ b/Alloy/Encoders/EuclideanDistance.swift @@ -67,9 +67,8 @@ final public class EuclideanDistance { let blockSize = BlockSize(width: blockSizeWidth, height: blockSizeHeight) - encoder.set(textures: [textureOne, - textureTwo]) - encoder.set(blockSize, at: 0) + encoder.setTextures(textureOne, textureTwo) + encoder.setValue(blockSize, at: 0) encoder.setBuffer(resultBuffer, offset: 0, index: 1) @@ -77,7 +76,7 @@ final public class EuclideanDistance { let threadgroupMemoryLength = threadgroupSize.width * threadgroupSize.height * 4 - * MemoryLayout.stride + * MemoryLayout>.stride encoder.setThreadgroupMemoryLength(threadgroupMemoryLength, index: 0) diff --git a/Alloy/Encoders/LookUpTable.swift b/Alloy/Encoders/LookUpTable.swift index b68ad6f..4f3b9ed 100644 --- a/Alloy/Encoders/LookUpTable.swift +++ b/Alloy/Encoders/LookUpTable.swift @@ -26,61 +26,59 @@ final public class LookUpTable { // MARK: - Encode - public func callAsFunction(sourceTexture: MTLTexture, - outputTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, lut: MTLTexture, intensity: Float, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTexture: sourceTexture, - outputTexture: outputTexture, + self.encode(source: source, + destination: destination, lut: lut, intensity: intensity, in: commandBuffer) } - public func callAsFunction(sourceTexture: MTLTexture, - outputTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, lut: MTLTexture, intensity: Float, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTexture: sourceTexture, - outputTexture: outputTexture, + self.encode(source: source, + destination: destination, lut: lut, intensity: intensity, using: encoder) } - public func encode(sourceTexture: MTLTexture, - outputTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, lut: MTLTexture, intensity: Float, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Look Up Table" - self.encode(sourceTexture: sourceTexture, - outputTexture: outputTexture, + self.encode(source: source, + destination: destination, lut: lut, intensity: intensity, using: encoder) } } - public func encode(sourceTexture: MTLTexture, - outputTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, lut: MTLTexture, intensity: Float, using encoder: MTLComputeCommandEncoder) { - encoder.set(textures: [sourceTexture, - outputTexture, - lut]) - encoder.set(intensity, at: 0) + encoder.setTextures(source, destination, lut) + encoder.setValue(intensity, at: 0) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: pipelineState, - exactly: outputTexture.size) + exactly: destination.size) } else { encoder.dispatch2d(state: pipelineState, - covering: outputTexture.size) + covering: destination.size) } } diff --git a/Alloy/Encoders/MPSUnaryImageKernels.swift b/Alloy/Encoders/MPSUnaryImageKernels.swift index 938ec2c..1975662 100644 --- a/Alloy/Encoders/MPSUnaryImageKernels.swift +++ b/Alloy/Encoders/MPSUnaryImageKernels.swift @@ -14,22 +14,21 @@ final public class MPSUnaryImageKernels { // MARK: - Encode - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, in: commandBuffer) } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, in commandBuffer: MTLCommandBuffer) { - guard self.kernelQueue - .count != 0 + guard !self.kernelQueue.isEmpty else { return } - let textureDescriptor = sourceTexture.descriptor + let textureDescriptor = source.descriptor textureDescriptor.usage = [.shaderRead, .shaderWrite] textureDescriptor.storageMode = .private // We need only 2 temporary images in the worst case. @@ -43,12 +42,12 @@ final public class MPSUnaryImageKernels { if self.kernelQueue.count == 1 { self.kernelQueue[0] .encode(commandBuffer: commandBuffer, - sourceTexture: sourceTexture, - destinationTexture: destinationTexture) + sourceTexture: source, + destinationTexture: destination) } else { self.kernelQueue[0] .encode(commandBuffer: commandBuffer, - sourceTexture: sourceTexture, + sourceTexture: source, destinationTexture: temporaryImages[0].texture) for i in 1 ..< self.kernelQueue.count - 1 { @@ -63,7 +62,7 @@ final public class MPSUnaryImageKernels { self.kernelQueue[self.kernelQueue.count - 1] .encode(commandBuffer: commandBuffer, sourceTexture: temporaryImages[0].texture, - destinationTexture: destinationTexture) + destinationTexture: destination) } } } diff --git a/Alloy/Encoders/MaskGuidedBlur.swift b/Alloy/Encoders/MaskGuidedBlur.swift index 9055c11..a25ecd8 100644 --- a/Alloy/Encoders/MaskGuidedBlur.swift +++ b/Alloy/Encoders/MaskGuidedBlur.swift @@ -29,24 +29,24 @@ final public class MaskGuidedBlur { // MARK: - Encode - public func callAsFunction(sourceTexture: MTLTexture, - maskTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + mask: MTLTexture, + destination: MTLTexture, sigma: Float, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTexture: sourceTexture, - maskTexture: maskTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + mask: mask, + destination: destination, sigma: sigma, in: commandBuffer) } - public func encode(sourceTexture: MTLTexture, - maskTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + mask: MTLTexture, + destination: MTLTexture, sigma: Float, in commandBuffer: MTLCommandBuffer) { - let temporaryTextureDescriptor = sourceTexture.descriptor + let temporaryTextureDescriptor = source.descriptor temporaryTextureDescriptor.usage = [.shaderRead, .shaderWrite] temporaryTextureDescriptor.storageMode = .private temporaryTextureDescriptor.pixelFormat = .rgba8Unorm @@ -57,30 +57,26 @@ final public class MaskGuidedBlur { textureDescriptor: temporaryTextureDescriptor) defer { temporaryImage.readCount = 0 } - encoder.set(textures: [sourceTexture, - maskTexture, - temporaryImage.texture]) - encoder.set(sigma, at: 0) + encoder.setTextures(source, mask, temporaryImage.texture) + encoder.setValue(sigma, at: 0) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: self.blurRowPassState, - exactly: sourceTexture.size) + exactly: source.size) } else { encoder.dispatch2d(state: self.blurRowPassState, - covering: sourceTexture.size) + covering: source.size) } - encoder.set(textures: [temporaryImage.texture, - maskTexture, - destinationTexture]) - encoder.set(sigma, at: 0) + encoder.setTextures(temporaryImage.texture, mask, destination) + encoder.setValue(sigma, at: 0) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: self.blurColumnPassState, - exactly: sourceTexture.size) + exactly: source.size) } else { encoder.dispatch2d(state: self.blurColumnPassState, - covering: sourceTexture.size) + covering: source.size) } } } diff --git a/Alloy/Encoders/NormalizeKernel.swift b/Alloy/Encoders/NormalizeKernel.swift index 7ce9d13..817bb93 100644 --- a/Alloy/Encoders/NormalizeKernel.swift +++ b/Alloy/Encoders/NormalizeKernel.swift @@ -26,61 +26,60 @@ final public class NormalizeKernel { // MARK: - Encode - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, mean: SIMD3, std: SIMD3, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, mean: mean, std: std, in: commandBuffer) } - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, mean: SIMD3, std: SIMD3, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, mean: mean, std: std, using: encoder) } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, mean: SIMD3, std: SIMD3, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Normalize Kernel" - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, mean: mean, std: std, using: encoder) } } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, mean: SIMD3, std: SIMD3, using encoder: MTLComputeCommandEncoder) { - encoder.set(textures: [sourceTexture, - destinationTexture]) - encoder.set(mean, at: 0) - encoder.set(std, at: 1) + encoder.setTextures(source, destination) + encoder.setValue(mean, at: 0) + encoder.setValue(std, at: 1) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: self.pipelineState, - exactly: destinationTexture.size) + exactly: destination.size) } else { encoder.dispatch2d(state: self.pipelineState, - covering: destinationTexture.size) + covering: destination.size) } } diff --git a/Alloy/Encoders/RGBAToYCbCr.swift b/Alloy/Encoders/RGBAToYCbCr.swift new file mode 100644 index 0000000..65a47b1 --- /dev/null +++ b/Alloy/Encoders/RGBAToYCbCr.swift @@ -0,0 +1,88 @@ +import Metal +import MetalPerformanceShaders + +final public class RGBAToYCbCr { + + // MARK: - Properties + + public let pipelineState: MTLComputePipelineState + private let deviceSupportsNonuniformThreadgroups: Bool + + // MARK: - Life Cycle + + public convenience init(context: MTLContext, + halfSizedCbCr: Bool = true) throws { + try self.init(library: context.library(for: Self.self), + halfSizedCbCr: halfSizedCbCr) + } + + public init(library: MTLLibrary, + halfSizedCbCr: Bool = true) throws { + self.deviceSupportsNonuniformThreadgroups = library.device + .supports(feature: .nonUniformThreadgroups) + let constantValues = MTLFunctionConstantValues() + constantValues.set(self.deviceSupportsNonuniformThreadgroups, + at: 0) + constantValues.set(halfSizedCbCr, + at: 2) + self.pipelineState = try library.computePipelineState(function: Self.functionName, + constants: constantValues) + } + + // MARK: - Encode + + public func callAsFunction(sourceRGBA: MTLTexture, + destinationY: MTLTexture, + destinationCbCr: MTLTexture, + in commandBuffer: MTLCommandBuffer) { + self.encode(sourceRGBA: sourceRGBA, + destinationY: destinationY, + destinationCbCr: destinationCbCr, + in: commandBuffer) + } + + public func callAsFunction(sourceRGBA: MTLTexture, + destinationY: MTLTexture, + destinationCbCr: MTLTexture, + using encoder: MTLComputeCommandEncoder) { + self.encode(sourceRGBA: sourceRGBA, + destinationY: destinationY, + destinationCbCr: destinationCbCr, + using: encoder) + } + + public func encode(sourceRGBA: MTLTexture, + destinationY: MTLTexture, + destinationCbCr: MTLTexture, + in commandBuffer: MTLCommandBuffer) { + commandBuffer.compute { encoder in + encoder.label = "RGBA To YCbCr" + self.encode(sourceRGBA: sourceRGBA, + destinationY: destinationY, + destinationCbCr: destinationCbCr, + using: encoder) + } + } + + private func encode(sourceRGBA: MTLTexture, + destinationY: MTLTexture, + destinationCbCr: MTLTexture, + using encoder: MTLComputeCommandEncoder) { + #if DEBUG + assert(sourceRGBA.size == destinationY.size, + "RGBA and Y textures must have equal sizes.") + #endif + encoder.setTextures(sourceRGBA, + destinationY, + destinationCbCr) + if self.deviceSupportsNonuniformThreadgroups { + encoder.dispatch2d(state: self.pipelineState, + exactly: destinationCbCr.size) + } else { + encoder.dispatch2d(state: self.pipelineState, + covering: destinationCbCr.size) + } + } + + public static let functionName = "rgbaToYCbCr" +} diff --git a/Alloy/Encoders/StdMeanNormalization.swift b/Alloy/Encoders/StdMeanNormalization.swift index 0aee8e5..540fc51 100644 --- a/Alloy/Encoders/StdMeanNormalization.swift +++ b/Alloy/Encoders/StdMeanNormalization.swift @@ -26,61 +26,60 @@ final public class StdMeanNormalization { // MARK: - Encode - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, mean: SIMD3, std: SIMD3, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, mean: mean, std: std, in: commandBuffer) } - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, mean: SIMD3, std: SIMD3, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, mean: mean, std: std, using: encoder) } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, mean: SIMD3, std: SIMD3, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Normalize Kernel" - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, mean: mean, std: std, using: encoder) } } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, mean: SIMD3, std: SIMD3, using encoder: MTLComputeCommandEncoder) { - encoder.set(textures: [sourceTexture, - destinationTexture]) - encoder.set(mean, at: 0) - encoder.set(std, at: 1) + encoder.setTextures(source, destination) + encoder.setValue(mean, at: 0) + encoder.setValue(std, at: 1) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: self.pipelineState, - exactly: destinationTexture.size) + exactly: destination.size) } else { encoder.dispatch2d(state: self.pipelineState, - covering: destinationTexture.size) + covering: destination.size) } } diff --git a/Alloy/Encoders/TextureAddConstant.swift b/Alloy/Encoders/TextureAddConstant.swift index e940be0..78b3b56 100644 --- a/Alloy/Encoders/TextureAddConstant.swift +++ b/Alloy/Encoders/TextureAddConstant.swift @@ -30,53 +30,52 @@ final public class TextureAddConstant { // MARK: - Encode - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, constant: SIMD4, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, constant: constant, in: commandBuffer) } - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, constant: SIMD4, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, constant: constant, using: encoder) } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, constant: SIMD4, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Texture Add Constant" - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, constant: constant, using: encoder) } } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, constant: SIMD4, using encoder: MTLComputeCommandEncoder) { - encoder.set(textures: [sourceTexture, - destinationTexture]) - encoder.set(constant, at: 0) + encoder.setTextures(source, destination) + encoder.setValue(constant, at: 0) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: self.pipelineState, - exactly: sourceTexture.size) + exactly: source.size) } else { encoder.dispatch2d(state: self.pipelineState, - covering: sourceTexture.size) + covering: source.size) } } diff --git a/Alloy/Encoders/TextureAffineCrop.swift b/Alloy/Encoders/TextureAffineCrop.swift index a9dc9ae..c2a41f2 100644 --- a/Alloy/Encoders/TextureAffineCrop.swift +++ b/Alloy/Encoders/TextureAffineCrop.swift @@ -27,53 +27,52 @@ final public class TextureAffineCrop { // MARK: - Encode - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, affineTransform: simd_float3x3, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, affineTransform: affineTransform, in: commandBuffer) } - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, affineTransform: simd_float3x3, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, affineTransform: affineTransform, using: encoder) } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, affineTransform: simd_float3x3, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Texture Affine Crop" - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, affineTransform: affineTransform, using: encoder) } } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, affineTransform: simd_float3x3, using encoder: MTLComputeCommandEncoder) { - encoder.set(textures: [sourceTexture, - destinationTexture]) - encoder.set(affineTransform, at: 0) + encoder.setTextures(source, destination) + encoder.setValue(affineTransform, at: 0) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: self.pipelineState, - exactly: destinationTexture.size) + exactly: destination.size) } else { encoder.dispatch2d(state: self.pipelineState, - covering: destinationTexture.size) + covering: destination.size) } } diff --git a/Alloy/Encoders/TextureCopy.swift b/Alloy/Encoders/TextureCopy.swift index 5621bbd..b9d9d72 100644 --- a/Alloy/Encoders/TextureCopy.swift +++ b/Alloy/Encoders/TextureCopy.swift @@ -29,86 +29,86 @@ final public class TextureCopy { // MARK: - Encode - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, in: commandBuffer) } - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, using: encoder) } public func callAsFunction(region sourceTexureRegion: MTLRegion, - from sourceTexture: MTLTexture, + from source: MTLTexture, to destinationTextureOrigin: MTLOrigin, - of destinationTexture: MTLTexture, + of destination: MTLTexture, in commandBuffer: MTLCommandBuffer) { self.copy(region: sourceTexureRegion, - from: sourceTexture, + from: source, to: destinationTextureOrigin, - of: destinationTexture, + of: destination, in: commandBuffer) } public func callAsFunction(region sourceTexureRegion: MTLRegion, - from sourceTexture: MTLTexture, + from source: MTLTexture, to destinationTextureOrigin: MTLOrigin, - of destinationTexture: MTLTexture, + of destination: MTLTexture, using encoder: MTLComputeCommandEncoder) { self.copy(region: sourceTexureRegion, - from: sourceTexture, + from: source, to: destinationTextureOrigin, - of: destinationTexture, + of: destination, using: encoder) } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Texture Copy" - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, using: encoder) } } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, using encoder: MTLComputeCommandEncoder) { - self.copy(region: sourceTexture.region, - from: sourceTexture, + self.copy(region: source.region, + from: source, to: .zero, - of: destinationTexture, + of: destination, using: encoder) } public func copy(region sourceTexureRegion: MTLRegion, - from sourceTexture: MTLTexture, + from source: MTLTexture, to destinationTextureOrigin: MTLOrigin, - of destinationTexture: MTLTexture, + of destination: MTLTexture, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Texture Copy" self.copy(region: sourceTexureRegion, - from: sourceTexture, + from: source, to: destinationTextureOrigin, - of: destinationTexture, + of: destination, using: encoder) } } public func copy(region sourceTexureRegion: MTLRegion, - from sourceTexture: MTLTexture, + from source: MTLTexture, to destinationTextureOrigin: MTLOrigin, - of destinationTexture: MTLTexture, + of destination: MTLTexture, using encoder: MTLComputeCommandEncoder) { // 1. Calculate read origin correction. let readOriginCorrection = MTLOrigin(x: abs(min(0, sourceTexureRegion.origin.x)), @@ -116,7 +116,7 @@ final public class TextureCopy { z: 0) // 2. Clamp read region to read texture. - guard var readRegion = sourceTexureRegion.clamped(to: sourceTexture.region) + guard var readRegion = sourceTexureRegion.clamped(to: source.region) else { #if DEBUG print("Read region is less or outside of source texture.") @@ -146,9 +146,9 @@ final public class TextureCopy { // 7. Calculate grid size. let gridSize = MTLSize(width: min(readRegion.size.width, - destinationTexture.width - writeOrigin.x), + destination.width - writeOrigin.x), height: min(readRegion.size.height, - destinationTexture.height - writeOrigin.y), + destination.height - writeOrigin.y), depth: 1) guard gridSize.width > 0 @@ -165,20 +165,19 @@ final public class TextureCopy { let writeOffset = SIMD2(x: .init(writeOrigin.x), y: .init(writeOrigin.y)) - encoder.set(textures: [sourceTexture, - destinationTexture]) - encoder.set(readOffset, + encoder.setTextures(source, destination) + encoder.setValue(readOffset, at: 0) - encoder.set(writeOffset, + encoder.setValue(writeOffset, at: 1) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: self.pipelineState, exactly: gridSize) } else { - encoder.set(SIMD2(x: .init(gridSize.width), - y: .init(gridSize.height)), - at: 2) + encoder.setValue(SIMD2(x: .init(gridSize.width), + y: .init(gridSize.height)), + at: 2) encoder.dispatch2d(state: self.pipelineState, covering: gridSize) } diff --git a/Alloy/Encoders/TextureDifferenceHighlight.swift b/Alloy/Encoders/TextureDifferenceHighlight.swift index a5d0539..96e53f7 100644 --- a/Alloy/Encoders/TextureDifferenceHighlight.swift +++ b/Alloy/Encoders/TextureDifferenceHighlight.swift @@ -25,69 +25,67 @@ final public class TextureDifferenceHighlight { // MARK: - Encode - public func callAsFunction(sourceTextureOne: MTLTexture, - sourceTextureTwo: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(sourceOne: MTLTexture, + sourceTwo: MTLTexture, + destination: MTLTexture, color: SIMD4, threshold: Float, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTextureOne: sourceTextureOne, - sourceTextureTwo: sourceTextureTwo, - destinationTexture: destinationTexture, + self.encode(sourceOne: sourceOne, + sourceTwo: sourceTwo, + destination: destination, color: color, threshold: threshold, in: commandBuffer) } - public func callAsFunction(sourceTextureOne: MTLTexture, - sourceTextureTwo: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(sourceOne: MTLTexture, + sourceTwo: MTLTexture, + destination: MTLTexture, color: SIMD4, threshold: Float, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTextureOne: sourceTextureOne, - sourceTextureTwo: sourceTextureTwo, - destinationTexture: destinationTexture, + self.encode(sourceOne: sourceOne, + sourceTwo: sourceTwo, + destination: destination, color: color, threshold: threshold, using: encoder) } - public func encode(sourceTextureOne: MTLTexture, - sourceTextureTwo: MTLTexture, - destinationTexture: MTLTexture, + public func encode(sourceOne: MTLTexture, + sourceTwo: MTLTexture, + destination: MTLTexture, color: SIMD4, threshold: Float, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Texture Difference Highlight" - self.encode(sourceTextureOne: sourceTextureOne, - sourceTextureTwo: sourceTextureTwo, - destinationTexture: destinationTexture, + self.encode(sourceOne: sourceOne, + sourceTwo: sourceTwo, + destination: destination, color: color, threshold: threshold, using: encoder) } } - public func encode(sourceTextureOne: MTLTexture, - sourceTextureTwo: MTLTexture, - destinationTexture: MTLTexture, + public func encode(sourceOne: MTLTexture, + sourceTwo: MTLTexture, + destination: MTLTexture, color: SIMD4, threshold: Float, using encoder: MTLComputeCommandEncoder) { - encoder.set(textures: [sourceTextureOne, - sourceTextureTwo, - destinationTexture]) - encoder.set(color, at: 0) - encoder.set(threshold, at: 1) + encoder.setTextures(sourceOne, sourceTwo, destination) + encoder.setValue(color, at: 0) + encoder.setValue(threshold, at: 1) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: self.pipelineState, - exactly: destinationTexture.size) + exactly: destination.size) } else { encoder.dispatch2d(state: self.pipelineState, - covering: destinationTexture.size) + covering: destination.size) } } diff --git a/Alloy/Encoders/TextureDivideByConstant.swift b/Alloy/Encoders/TextureDivideByConstant.swift index b5a3510..c2b7372 100644 --- a/Alloy/Encoders/TextureDivideByConstant.swift +++ b/Alloy/Encoders/TextureDivideByConstant.swift @@ -30,103 +30,101 @@ final public class TextureDivideByConstant { // MARK: - Encode - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, constant: SIMD4, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, constant: constant, in: commandBuffer) } - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, constant: SIMD4, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, constant: constant, using: encoder) } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, constant: SIMD4, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Texture Divide by Constant" - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, constant: constant, using: encoder) } } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, constant: SIMD4, using encoder: MTLComputeCommandEncoder) { - encoder.set(textures: [sourceTexture, - destinationTexture]) - encoder.set(constant, at: 0) + encoder.setTextures(source, destination) + encoder.setValue(constant, at: 0) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: self.pipelineState, - exactly: sourceTexture.size) + exactly: source.size) } else { encoder.dispatch2d(state: self.pipelineState, - covering: sourceTexture.size) + covering: source.size) } } - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, constant: MTLBuffer, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, constant: constant, in: commandBuffer) } - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, constant: MTLBuffer, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, constant: constant, using: encoder) } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, constant: MTLBuffer, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Texture Divide by Constant" - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, constant: constant, using: encoder) } } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, constant: MTLBuffer, using encoder: MTLComputeCommandEncoder) { - encoder.set(textures: [sourceTexture, - destinationTexture]) + encoder.setTextures(source, destination) encoder.setBuffer(constant, offset: 0, index: 0) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: self.pipelineState, - exactly: sourceTexture.size) + exactly: source.size) } else { encoder.dispatch2d(state: self.pipelineState, - covering: sourceTexture.size) + covering: source.size) } } diff --git a/Alloy/Encoders/TextureInterpolation.swift b/Alloy/Encoders/TextureInterpolation.swift index 44d1225..42c908c 100644 --- a/Alloy/Encoders/TextureInterpolation.swift +++ b/Alloy/Encoders/TextureInterpolation.swift @@ -30,61 +30,59 @@ final public class TextureInterpolation { // MARK: - Encode - public func callAsFunction(sourceTextureOne: MTLTexture, - sourceTextureTwo: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(sourceOne: MTLTexture, + sourceTwo: MTLTexture, + destination: MTLTexture, weight: Float, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTextureOne: sourceTextureOne, - sourceTextureTwo: sourceTextureTwo, - destinationTexture: destinationTexture, + self.encode(sourceOne: sourceOne, + sourceTwo: sourceTwo, + destination: destination, weight: weight, in: commandBuffer) } - public func callAsFunction(sourceTextureOne: MTLTexture, - sourceTextureTwo: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(sourceOne: MTLTexture, + sourceTwo: MTLTexture, + destination: MTLTexture, weight: Float, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTextureOne: sourceTextureOne, - sourceTextureTwo: sourceTextureTwo, - destinationTexture: destinationTexture, + self.encode(sourceOne: sourceOne, + sourceTwo: sourceTwo, + destination: destination, weight: weight, using: encoder) } - public func encode(sourceTextureOne: MTLTexture, - sourceTextureTwo: MTLTexture, - destinationTexture: MTLTexture, + public func encode(sourceOne: MTLTexture, + sourceTwo: MTLTexture, + destination: MTLTexture, weight: Float, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Texture Interpolation" - self.encode(sourceTextureOne: sourceTextureOne, - sourceTextureTwo: sourceTextureTwo, - destinationTexture: destinationTexture, + self.encode(sourceOne: sourceOne, + sourceTwo: sourceTwo, + destination: destination, weight: weight, using: encoder) } } - public func encode(sourceTextureOne: MTLTexture, - sourceTextureTwo: MTLTexture, - destinationTexture: MTLTexture, + public func encode(sourceOne: MTLTexture, + sourceTwo: MTLTexture, + destination: MTLTexture, weight: Float, using encoder: MTLComputeCommandEncoder) { - encoder.set(textures: [sourceTextureOne, - sourceTextureTwo, - destinationTexture]) - encoder.set(weight, at: 0) + encoder.setTextures(sourceOne, sourceTwo, destination) + encoder.setValue(weight, at: 0) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: self.pipelineState, - exactly: destinationTexture.size) + exactly: destination.size) } else { encoder.dispatch2d(state: self.pipelineState, - covering: destinationTexture.size) + covering: destination.size) } } diff --git a/Alloy/Encoders/TextureMask.swift b/Alloy/Encoders/TextureMask.swift index c6c83c1..22505d3 100644 --- a/Alloy/Encoders/TextureMask.swift +++ b/Alloy/Encoders/TextureMask.swift @@ -29,53 +29,59 @@ final public class TextureMask { // MARK: - Encode - public func callAsFunction(sourceTexture: MTLTexture, - maskTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + mask: MTLTexture, + destination: MTLTexture, + isInversed: Bool = false, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTexture: sourceTexture, - maskTexture: maskTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + mask: mask, + destination: destination, + isInversed: isInversed, in: commandBuffer) } - public func callAsFunction(sourceTexture: MTLTexture, - maskTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + mask: MTLTexture, + destination: MTLTexture, + isInversed: Bool = false, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTexture: sourceTexture, - maskTexture: maskTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + mask: mask, + destination: destination, + isInversed: isInversed, using: encoder) } - public func encode(sourceTexture: MTLTexture, - maskTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + mask: MTLTexture, + destination: MTLTexture, + isInversed: Bool = false, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Texture Mask" - self.encode(sourceTexture: sourceTexture, - maskTexture: maskTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + mask: mask, + destination: destination, + isInversed: isInversed, using: encoder) } } - public func encode(sourceTexture: MTLTexture, - maskTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + mask: MTLTexture, + destination: MTLTexture, + isInversed: Bool = false, using encoder: MTLComputeCommandEncoder) { - encoder.set(textures: [sourceTexture, - maskTexture, - destinationTexture]) + encoder.setTextures(source, mask, destination) + encoder.setValue(isInversed, at: 0) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: pipelineState, - exactly: destinationTexture.size) + exactly: destination.size) } else { encoder.dispatch2d(state: pipelineState, - covering: destinationTexture.size) + covering: destination.size) } } diff --git a/Alloy/Encoders/TextureMaskedMix.swift b/Alloy/Encoders/TextureMaskedMix.swift index 70330f1..fc84576 100644 --- a/Alloy/Encoders/TextureMaskedMix.swift +++ b/Alloy/Encoders/TextureMaskedMix.swift @@ -25,60 +25,57 @@ final public class TextureMaskedMix { // MARK: - Encode - public func callAsFunction(sourceTextureOne: MTLTexture, - sourceTextureTwo: MTLTexture, - maskTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(sourceOne: MTLTexture, + sourceTwo: MTLTexture, + mask: MTLTexture, + destination: MTLTexture, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTextureOne: sourceTextureOne, - sourceTextureTwo: sourceTextureTwo, - maskTexture: maskTexture, - destinationTexture: destinationTexture, + self.encode(sourceOne: sourceOne, + sourceTwo: sourceTwo, + mask: mask, + destination: destination, in: commandBuffer) } - public func callAsFunction(sourceTextureOne: MTLTexture, - sourceTextureTwo: MTLTexture, - maskTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(sourceOne: MTLTexture, + sourceTwo: MTLTexture, + mask: MTLTexture, + destination: MTLTexture, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTextureOne: sourceTextureOne, - sourceTextureTwo: sourceTextureTwo, - maskTexture: maskTexture, - destinationTexture: destinationTexture, + self.encode(sourceOne: sourceOne, + sourceTwo: sourceTwo, + mask: mask, + destination: destination, using: encoder) } - public func encode(sourceTextureOne: MTLTexture, - sourceTextureTwo: MTLTexture, - maskTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(sourceOne: MTLTexture, + sourceTwo: MTLTexture, + mask: MTLTexture, + destination: MTLTexture, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Texture Mix" - self.encode(sourceTextureOne: sourceTextureOne, - sourceTextureTwo: sourceTextureTwo, - maskTexture: maskTexture, - destinationTexture: destinationTexture, + self.encode(sourceOne: sourceOne, + sourceTwo: sourceTwo, + mask: mask, + destination: destination, using: encoder) } } - public func encode(sourceTextureOne: MTLTexture, - sourceTextureTwo: MTLTexture, - maskTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(sourceOne: MTLTexture, + sourceTwo: MTLTexture, + mask: MTLTexture, + destination: MTLTexture, using encoder: MTLComputeCommandEncoder) { - encoder.set(textures: [sourceTextureOne, - sourceTextureTwo, - maskTexture, - destinationTexture]) + encoder.setTextures(sourceOne, sourceTwo, mask, destination) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: self.pipelineState, - exactly: destinationTexture.size) + exactly: destination.size) } else { encoder.dispatch2d(state: self.pipelineState, - covering: destinationTexture.size) + covering: destination.size) } } diff --git a/Alloy/Encoders/TextureMax.swift b/Alloy/Encoders/TextureMax.swift index fa3fb4c..f889366 100644 --- a/Alloy/Encoders/TextureMax.swift +++ b/Alloy/Encoders/TextureMax.swift @@ -22,54 +22,53 @@ final public class TextureMax { // MARK: - Encode - public func callAsFunction(sourceTexture: MTLTexture, - resultBuffer: MTLBuffer, + public func callAsFunction(source: MTLTexture, + result: MTLBuffer, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTexture: sourceTexture, - resultBuffer: resultBuffer, + self.encode(source: source, + result: result, in: commandBuffer) } - public func callAsFunction(sourceTexture: MTLTexture, - resultBuffer: MTLBuffer, + public func callAsFunction(source: MTLTexture, + result: MTLBuffer, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTexture: sourceTexture, - resultBuffer: resultBuffer, + self.encode(source: source, + result: result, using: encoder) } - public func encode(sourceTexture: MTLTexture, - resultBuffer: MTLBuffer, + public func encode(source: MTLTexture, + result: MTLBuffer, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Texture Max" - self.encode(sourceTexture: sourceTexture, - resultBuffer: resultBuffer, + self.encode(source: source, + result: result, using: encoder) } } - public func encode(sourceTexture: MTLTexture, - resultBuffer: MTLBuffer, + public func encode(source: MTLTexture, + result: MTLBuffer, using encoder: MTLComputeCommandEncoder) { - let threadgroupSize = MTLSize(width: 8, height: 8, depth: 1).clamped(to: sourceTexture.size) - let blockSizeWidth = (sourceTexture.width + threadgroupSize.width - 1) + let threadgroupSize = MTLSize(width: 8, height: 8, depth: 1).clamped(to: source.size) + let blockSizeWidth = (source.width + threadgroupSize.width - 1) / threadgroupSize.width - let blockSizeHeight = (sourceTexture.height + threadgroupSize.height - 1) + let blockSizeHeight = (source.height + threadgroupSize.height - 1) / threadgroupSize.height let blockSize = BlockSize(width: blockSizeWidth, height: blockSizeHeight) - encoder.set(textures: [sourceTexture]) - encoder.set(blockSize, at: 0) - encoder.setBuffer(resultBuffer, + encoder.setTextures(source) + encoder.setValue(blockSize, at: 0) + encoder.setBuffer(result, offset: 0, index: 1) let threadgroupMemoryLength = threadgroupSize.width * threadgroupSize.height - * 4 - * MemoryLayout.stride + * MemoryLayout>.stride encoder.setThreadgroupMemoryLength(threadgroupMemoryLength, index: 0) diff --git a/Alloy/Encoders/TextureMean.swift b/Alloy/Encoders/TextureMean.swift index 1e1a1e8..515b859 100644 --- a/Alloy/Encoders/TextureMean.swift +++ b/Alloy/Encoders/TextureMean.swift @@ -22,54 +22,53 @@ final public class TextureMean { // MARK: - Encode - public func callAsFunction(sourceTexture: MTLTexture, - resultBuffer: MTLBuffer, + public func callAsFunction(source: MTLTexture, + result: MTLBuffer, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTexture: sourceTexture, - resultBuffer: resultBuffer, + self.encode(source: source, + result: result, in: commandBuffer) } - public func callAsFunction(sourceTexture: MTLTexture, - resultBuffer: MTLBuffer, + public func callAsFunction(source: MTLTexture, + result: MTLBuffer, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTexture: sourceTexture, - resultBuffer: resultBuffer, + self.encode(source: source, + result: result, using: encoder) } - public func encode(sourceTexture: MTLTexture, - resultBuffer: MTLBuffer, + public func encode(source: MTLTexture, + result: MTLBuffer, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Texture Mean" - self.encode(sourceTexture: sourceTexture, - resultBuffer: resultBuffer, + self.encode(source: source, + result: result, using: encoder) } } - public func encode(sourceTexture: MTLTexture, - resultBuffer: MTLBuffer, + public func encode(source: MTLTexture, + result: MTLBuffer, using encoder: MTLComputeCommandEncoder) { - let threadgroupSize = MTLSize(width: 8, height: 8, depth: 1).clamped(to: sourceTexture.size) - let blockSizeWidth = (sourceTexture.width + threadgroupSize.width - 1) + let threadgroupSize = MTLSize(width: 8, height: 8, depth: 1).clamped(to: source.size) + let blockSizeWidth = (source.width + threadgroupSize.width - 1) / threadgroupSize.width - let blockSizeHeight = (sourceTexture.height + threadgroupSize.height - 1) + let blockSizeHeight = (source.height + threadgroupSize.height - 1) / threadgroupSize.height let blockSize = BlockSize(width: blockSizeWidth, height: blockSizeHeight) - encoder.set(textures: [sourceTexture]) - encoder.set(blockSize, at: 0) - encoder.setBuffer(resultBuffer, + encoder.setTextures(source) + encoder.setValue(blockSize, at: 0) + encoder.setBuffer(result, offset: 0, index: 1) let threadgroupMemoryLength = threadgroupSize.width * threadgroupSize.height - * 4 - * MemoryLayout.stride + * MemoryLayout>.stride encoder.setThreadgroupMemoryLength(threadgroupMemoryLength, index: 0) diff --git a/Alloy/Encoders/TextureMin.swift b/Alloy/Encoders/TextureMin.swift index 62a654b..281bded 100644 --- a/Alloy/Encoders/TextureMin.swift +++ b/Alloy/Encoders/TextureMin.swift @@ -22,54 +22,53 @@ final public class TextureMin { // MARK: - Encode - public func callAsFunction(sourceTexture: MTLTexture, - resultBuffer: MTLBuffer, + public func callAsFunction(source: MTLTexture, + result: MTLBuffer, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTexture: sourceTexture, - resultBuffer: resultBuffer, + self.encode(source: source, + result: result, in: commandBuffer) } - public func callAsFunction(sourceTexture: MTLTexture, - resultBuffer: MTLBuffer, + public func callAsFunction(source: MTLTexture, + result: MTLBuffer, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTexture: sourceTexture, - resultBuffer: resultBuffer, + self.encode(source: source, + result: result, using: encoder) } - public func encode(sourceTexture: MTLTexture, - resultBuffer: MTLBuffer, + public func encode(source: MTLTexture, + result: MTLBuffer, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Texture Min" - self.encode(sourceTexture: sourceTexture, - resultBuffer: resultBuffer, + self.encode(source: source, + result: result, using: encoder) } } - public func encode(sourceTexture: MTLTexture, - resultBuffer: MTLBuffer, + public func encode(source: MTLTexture, + result: MTLBuffer, using encoder: MTLComputeCommandEncoder) { - let threadgroupSize = MTLSize(width: 8, height: 8, depth: 1).clamped(to: sourceTexture.size) - let blockSizeWidth = (sourceTexture.width + threadgroupSize.width - 1) + let threadgroupSize = MTLSize(width: 8, height: 8, depth: 1).clamped(to: source.size) + let blockSizeWidth = (source.width + threadgroupSize.width - 1) / threadgroupSize.width - let blockSizeHeight = (sourceTexture.height + threadgroupSize.height - 1) + let blockSizeHeight = (source.height + threadgroupSize.height - 1) / threadgroupSize.height let blockSize = BlockSize(width: blockSizeWidth, height: blockSizeHeight) - encoder.set(textures: [sourceTexture]) - encoder.set(blockSize, at: 0) - encoder.setBuffer(resultBuffer, + encoder.setTextures(source) + encoder.setValue(blockSize, at: 0) + encoder.setBuffer(result, offset: 0, index: 1) let threadgroupMemoryLength = threadgroupSize.width * threadgroupSize.height - * 4 - * MemoryLayout.stride + * MemoryLayout>.stride encoder.setThreadgroupMemoryLength(threadgroupMemoryLength, index: 0) diff --git a/Alloy/Encoders/TextureMultiplyAdd.swift b/Alloy/Encoders/TextureMultiplyAdd.swift index 00a2ce5..a36f424 100644 --- a/Alloy/Encoders/TextureMultiplyAdd.swift +++ b/Alloy/Encoders/TextureMultiplyAdd.swift @@ -30,53 +30,51 @@ final public class TextureMultiplyAdd { // MARK: - Encode - public func callAsFunction(sourceTextureOne: MTLTexture, - sourceTextureTwo: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(sourceOne: MTLTexture, + sourceTwo: MTLTexture, + destination: MTLTexture, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTextureOne: sourceTextureOne, - sourceTextureTwo: sourceTextureTwo, - destinationTexture: destinationTexture, + self.encode(sourceOne: sourceOne, + sourceTwo: sourceTwo, + destination: destination, in: commandBuffer) } - public func callAsFunction(sourceTextureOne: MTLTexture, - sourceTextureTwo: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(sourceOne: MTLTexture, + sourceTwo: MTLTexture, + destination: MTLTexture, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTextureOne: sourceTextureOne, - sourceTextureTwo: sourceTextureTwo, - destinationTexture: destinationTexture, + self.encode(sourceOne: sourceOne, + sourceTwo: sourceTwo, + destination: destination, using: encoder) } - public func encode(sourceTextureOne: MTLTexture, - sourceTextureTwo: MTLTexture, - destinationTexture: MTLTexture, + public func encode(sourceOne: MTLTexture, + sourceTwo: MTLTexture, + destination: MTLTexture, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Texture Multiply Add" - self.encode(sourceTextureOne: sourceTextureOne, - sourceTextureTwo: sourceTextureTwo, - destinationTexture: destinationTexture, + self.encode(sourceOne: sourceOne, + sourceTwo: sourceTwo, + destination: destination, using: encoder) } } - public func encode(sourceTextureOne: MTLTexture, - sourceTextureTwo: MTLTexture, - destinationTexture: MTLTexture, + public func encode(sourceOne: MTLTexture, + sourceTwo: MTLTexture, + destination: MTLTexture, using encoder: MTLComputeCommandEncoder) { - encoder.set(textures: [sourceTextureOne, - sourceTextureTwo, - destinationTexture]) + encoder.setTextures(sourceOne, sourceTwo, destination) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: pipelineState, - exactly: destinationTexture.size) + exactly: destination.size) } else { encoder.dispatch2d(state: pipelineState, - covering: destinationTexture.size) + covering: destination.size) } } diff --git a/Alloy/Encoders/TextureNormalization.swift b/Alloy/Encoders/TextureNormalization.swift index 6de6f6e..3e01d4d 100644 --- a/Alloy/Encoders/TextureNormalization.swift +++ b/Alloy/Encoders/TextureNormalization.swift @@ -23,41 +23,41 @@ final public class TextureNormalization { // MARK: - Encode - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, in: commandBuffer) } - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, using: encoder) } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Texture Normalization" - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, using: encoder) } } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, using encoder: MTLComputeCommandEncoder) { - self.textureMax(sourceTexture: sourceTexture, - resultBuffer: self.intermediateBuffer, + self.textureMax(source: source, + result: self.intermediateBuffer, using: encoder) - self.textureDivide(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.textureDivide(source: source, + destination: destination, constant: self.intermediateBuffer, using: encoder) } diff --git a/Alloy/Encoders/TextureResize.swift b/Alloy/Encoders/TextureResize.swift index 81b08ca..34318c0 100644 --- a/Alloy/Encoders/TextureResize.swift +++ b/Alloy/Encoders/TextureResize.swift @@ -55,48 +55,47 @@ final public class TextureResize { // MARK: - Encode - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, in: commandBuffer) } - public func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func callAsFunction(source: MTLTexture, + destination: MTLTexture, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, using: encoder) } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "Texture Resize" - self.encode(sourceTexture: sourceTexture, - destinationTexture: destinationTexture, + self.encode(source: source, + destination: destination, using: encoder) } } - public func encode(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + public func encode(source: MTLTexture, + destination: MTLTexture, using encoder: MTLComputeCommandEncoder) { - encoder.set(textures: [sourceTexture, - destinationTexture]) + encoder.setTextures(source, destination) encoder.setSamplerState(self.samplerState, index: 0) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: self.pipelineState, - exactly: destinationTexture.size) + exactly: destination.size) } else { encoder.dispatch2d(state: self.pipelineState, - covering: destinationTexture.size) + covering: destination.size) } } diff --git a/Alloy/Encoders/TextureWeightedMix.swift b/Alloy/Encoders/TextureWeightedMix.swift index 4b437f6..aa3324c 100644 --- a/Alloy/Encoders/TextureWeightedMix.swift +++ b/Alloy/Encoders/TextureWeightedMix.swift @@ -69,10 +69,10 @@ final public class TextureWeightedMix { destination: MTLTexture, weight: Float, using encoder: MTLComputeCommandEncoder) { - encoder.set(textures: [sourceOne, - sourceTwo, - destination]) - encoder.set(weight, at: 0) + encoder.setTextures(sourceOne, + sourceTwo, + destination) + encoder.setValue(weight, at: 0) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: self.pipelineState, exactly: destination.size) diff --git a/Alloy/Encoders/YCbCrToRGBA.swift b/Alloy/Encoders/YCbCrToRGBA.swift index 7f69397..2b96d1c 100644 --- a/Alloy/Encoders/YCbCrToRGBA.swift +++ b/Alloy/Encoders/YCbCrToRGBA.swift @@ -25,53 +25,51 @@ final public class YCbCrToRGBA { // MARK: - Encode - public func callAsFunction(sourceYTexture: MTLTexture, - sourceCbCrTexture: MTLTexture, - destinationRGBATexture: MTLTexture, + public func callAsFunction(sourceY: MTLTexture, + sourceCbCr: MTLTexture, + destinationRGBA: MTLTexture, in commandBuffer: MTLCommandBuffer) { - self.encode(sourceYTexture: sourceYTexture, - sourceCbCrTexture: sourceCbCrTexture, - destinationRGBATexture: destinationRGBATexture, + self.encode(sourceY: sourceY, + sourceCbCr: sourceCbCr, + destinationRGBA: destinationRGBA, in: commandBuffer) } - public func callAsFunction(sourceYTexture: MTLTexture, - sourceCbCrTexture: MTLTexture, - destinationRGBATexture: MTLTexture, + public func callAsFunction(sourceY: MTLTexture, + sourceCbCr: MTLTexture, + destinationRGBA: MTLTexture, using encoder: MTLComputeCommandEncoder) { - self.encode(sourceYTexture: sourceYTexture, - sourceCbCrTexture: sourceCbCrTexture, - destinationRGBATexture: destinationRGBATexture, + self.encode(sourceY: sourceY, + sourceCbCr: sourceCbCr, + destinationRGBA: destinationRGBA, using: encoder) } - public func encode(sourceYTexture: MTLTexture, - sourceCbCrTexture: MTLTexture, - destinationRGBATexture: MTLTexture, + public func encode(sourceY: MTLTexture, + sourceCbCr: MTLTexture, + destinationRGBA: MTLTexture, in commandBuffer: MTLCommandBuffer) { commandBuffer.compute { encoder in encoder.label = "YCbCr To RGBA" - self.encode(sourceYTexture: sourceYTexture, - sourceCbCrTexture: sourceCbCrTexture, - destinationRGBATexture: destinationRGBATexture, + self.encode(sourceY: sourceY, + sourceCbCr: sourceCbCr, + destinationRGBA: destinationRGBA, using: encoder) } } - public func encode(sourceYTexture: MTLTexture, - sourceCbCrTexture: MTLTexture, - destinationRGBATexture: MTLTexture, + public func encode(sourceY: MTLTexture, + sourceCbCr: MTLTexture, + destinationRGBA: MTLTexture, using encoder: MTLComputeCommandEncoder) { - encoder.set(textures: [sourceYTexture, - sourceCbCrTexture, - destinationRGBATexture]) + encoder.setTextures(sourceY, sourceCbCr, destinationRGBA) if self.deviceSupportsNonuniformThreadgroups { encoder.dispatch2d(state: self.pipelineState, - exactly: destinationRGBATexture.size) + exactly: destinationRGBA.size) } else { encoder.dispatch2d(state: self.pipelineState, - covering: destinationRGBATexture.size) + covering: destinationRGBA.size) } } diff --git a/Alloy/Float16.swift b/Alloy/Float16.swift index a12f3cd..263b514 100644 --- a/Alloy/Float16.swift +++ b/Alloy/Float16.swift @@ -14,15 +14,24 @@ public typealias Float16 = UInt16 /// - input: A pointer to an array of `Float16`s. /// - count: Number of elements in the array. /// - Returns: An array of regular Swift `Float`s. -public func float16to32(_ input: UnsafeMutableRawPointer, count: Int) -> [Float]? { - var output = [Float](repeating: 0, count: count) - var bufferFloat16 = vImage_Buffer(data: input, height: 1, width: UInt(count), rowBytes: count * 2) - var bufferFloat32 = vImage_Buffer(data: &output, height: 1, width: UInt(count), rowBytes: count * 4) - - if vImageConvert_Planar16FtoPlanarF(&bufferFloat16, &bufferFloat32, 0) != kvImageNoError { - return nil +public func float16to32(_ input: UnsafeMutableRawPointer, + count: Int) -> [Float]? { + var output = [Float](repeating: 0, + count: count) + let status = output.withUnsafeMutableBytes { p -> Int in + var bufferFloat16 = vImage_Buffer(data: input, + height: 1, + width: .init(count), + rowBytes: count * 2) + var bufferFloat32 = vImage_Buffer(data: p.baseAddress, + height: 1, + width: .init(count), + rowBytes: count * 4) + return vImageConvert_Planar16FtoPlanarF(&bufferFloat16, + &bufferFloat32, + 0) } - return output + return status == kvImageNoError ? output : nil } /// Uses vImage to convert an array of Swift floats into a buffer of float16s. @@ -31,13 +40,20 @@ public func float16to32(_ input: UnsafeMutableRawPointer, count: Int) -> [Float] /// - input: A pointer to an array of `Float`s. /// - count: Number of elements in the array. /// - Returns: An array of `Float16`s. -public func float32to16(_ input: UnsafeMutablePointer, count: Int) -> [Float16]? { - var output = [Float16](repeating: 0, count: count) - var bufferFloat32 = vImage_Buffer(data: input, height: 1, width: UInt(count), rowBytes: count * 4) - var bufferFloat16 = vImage_Buffer(data: &output, height: 1, width: UInt(count), rowBytes: count * 2) - - if vImageConvert_PlanarFtoPlanar16F(&bufferFloat32, &bufferFloat16, 0) != kvImageNoError { - return nil +public func float32to16(_ input: UnsafeMutablePointer, + count: Int) -> [Float16]? { + var output = [Float16](repeating: 0, + count: count) + let status = output.withUnsafeMutableBytes { p -> Int in + var bufferFloat32 = vImage_Buffer(data: input, + height: 1, + width: .init(count), + rowBytes: count * 4) + var bufferFloat16 = vImage_Buffer(data: p.baseAddress, + height: 1, + width: .init(count), + rowBytes: count * 2) + return vImageConvert_PlanarFtoPlanar16F(&bufferFloat32, &bufferFloat16, 0) } - return output + return status == kvImageNoError ? output : nil } diff --git a/Alloy/MPS+Functors.swift b/Alloy/MPS+Functors.swift index f2620c9..8f1a4d7 100644 --- a/Alloy/MPS+Functors.swift +++ b/Alloy/MPS+Functors.swift @@ -1,19 +1,19 @@ import MetalPerformanceShaders public extension MPSUnaryImageKernel { - func callAsFunction(sourceTexture: MTLTexture, - destinationTexture: MTLTexture, + func callAsFunction(source: MTLTexture, + destination: MTLTexture, in commandBuffer: MTLCommandBuffer) { self.encode(commandBuffer: commandBuffer, - sourceTexture: sourceTexture, - destinationTexture: destinationTexture) + sourceTexture: source, + destinationTexture: destination) } - func callAsFunction(inPlaceTexture: UnsafeMutablePointer, + func callAsFunction(inPlace: UnsafeMutablePointer, fallbackCopyAllocator: MPSCopyAllocator? = nil, in commandBuffer: MTLCommandBuffer) { self.encode(commandBuffer: commandBuffer, - inPlaceTexture: inPlaceTexture, + inPlaceTexture: inPlace, fallbackCopyAllocator: fallbackCopyAllocator) } diff --git a/Alloy/MTLBlitCommandEncoder+Extensions.swift b/Alloy/MTLBlitCommandEncoder+Extensions.swift index 3947523..06518d3 100644 --- a/Alloy/MTLBlitCommandEncoder+Extensions.swift +++ b/Alloy/MTLBlitCommandEncoder+Extensions.swift @@ -1,28 +1,41 @@ import Metal.MTLBlitCommandEncoder public extension MTLBlitCommandEncoder { - func copy(region: MTLRegion, from texture: MTLTexture, - to targetOrigin: MTLOrigin, of targetTexture: MTLTexture, - sourceSlice: Int = 0, sourceLevel: Int = 0, - destinationSlice: Int = 0, destionationLevel: Int = 0) { - self.copy(from: texture, - sourceSlice: sourceSlice, sourceLevel: sourceLevel, - sourceOrigin: region.origin, sourceSize: region.size, - to: targetTexture, - destinationSlice: destinationSlice, destinationLevel: destionationLevel, + func copy(region: MTLRegion, + from source: MTLTexture, + to targetOrigin: MTLOrigin, + of target: MTLTexture, + sourceSlice: Int = 0, + sourceLevel: Int = 0, + destinationSlice: Int = 0, + destionationLevel: Int = 0) { + self.copy(from: source, + sourceSlice: sourceSlice, + sourceLevel: sourceLevel, + sourceOrigin: region.origin, + sourceSize: region.size, + to: target, + destinationSlice: destinationSlice, + destinationLevel: destionationLevel, destinationOrigin: targetOrigin) } func copy(texture: MTLTexture, - to targetOrigin: MTLOrigin, of targetTexture: MTLTexture, - sourceSlice: Int = 0, sourceLevel: Int = 0, - destinationSlice: Int = 0, destionationLevel: Int = 0) { + to targetOrigin: MTLOrigin, + of target: MTLTexture, + sourceSlice: Int = 0, + sourceLevel: Int = 0, + destinationSlice: Int = 0, + destionationLevel: Int = 0) { let region = texture.region self.copy(from: texture, - sourceSlice: sourceSlice, sourceLevel: sourceLevel, - sourceOrigin: region.origin, sourceSize: region.size, - to: targetTexture, - destinationSlice: destinationSlice, destinationLevel: destionationLevel, + sourceSlice: sourceSlice, + sourceLevel: sourceLevel, + sourceOrigin: region.origin, + sourceSize: region.size, + to: target, + destinationSlice: destinationSlice, + destinationLevel: destionationLevel, destinationOrigin: targetOrigin) } } diff --git a/Alloy/MTLComputeCommandEncoder+Extensions.swift b/Alloy/MTLComputeCommandEncoder+Extensions.swift index eca2412..384ab5d 100644 --- a/Alloy/MTLComputeCommandEncoder+Extensions.swift +++ b/Alloy/MTLComputeCommandEncoder+Extensions.swift @@ -2,21 +2,31 @@ import Metal public extension MTLComputeCommandEncoder { - func set(_ value: T, at index: Int) { + func setValue(_ value: T, at index: Int) { var t = value self.setBytes(&t, length: MemoryLayout.stride, index: index) } - func set(_ value: [T], at index: Int) { + func setValue(_ value: [T], at index: Int) { var t = value self.setBytes(&t, length: MemoryLayout.stride * value.count, index: index) } - func set(textures: [MTLTexture?], startingAt startIndex: Int = 0) { + func setTextures(_ textures: [MTLTexture?], startingAt startIndex: Int = 0) { self.setTextures(textures, range: startIndex..<(startIndex + textures.count)) } - func set(buffers: [MTLBuffer?], offsets: [Int]? = nil, startingAt startIndex: Int = 0) { + func setBuffers(_ buffers: [MTLBuffer?], offsets: [Int]? = nil, startingAt startIndex: Int = 0) { + self.setBuffers(buffers, + offsets: offsets ?? buffers.map { _ in 0 }, + range: startIndex..<(startIndex + buffers.count)) + } + + func setTextures(_ textures: MTLTexture?..., startingAt startIndex: Int = 0) { + self.setTextures(textures, range: startIndex..<(startIndex + textures.count)) + } + + func setBuffers(_ buffers: MTLBuffer?..., offsets: [Int]? = nil, startingAt startIndex: Int = 0) { self.setBuffers(buffers, offsets: offsets ?? buffers.map { _ in 0 }, range: startIndex..<(startIndex + buffers.count)) diff --git a/Alloy/MTLContext.swift b/Alloy/MTLContext.swift index 0341f85..7b81ff0 100644 --- a/Alloy/MTLContext.swift +++ b/Alloy/MTLContext.swift @@ -91,7 +91,7 @@ public final class MTLContext { if isSRGB == nil, let colorSpace = image.colorSpace, - colorSpace.name == CGColorSpace.sRGB + colorSpace.name == CGColorSpace.linearSRGB { isSRGB = true } diff --git a/Alloy/MTLOrigin+Extensions.swift b/Alloy/MTLOrigin+Extensions.swift index b925e2d..dd286b6 100644 --- a/Alloy/MTLOrigin+Extensions.swift +++ b/Alloy/MTLOrigin+Extensions.swift @@ -6,7 +6,7 @@ public extension MTLOrigin { y: value, z: value) } - + func clamped(to size: MTLSize) -> MTLOrigin { return MTLOrigin(x: min(max(self.x, 0), size.width), y: min(max(self.y, 0), size.height), @@ -14,4 +14,14 @@ public extension MTLOrigin { } static let zero = MTLOrigin(repeating: 0) + static func ==(lhs: MTLOrigin, rhs: MTLOrigin) -> Bool { + return lhs.x == rhs.x + && lhs.y == rhs.y + && lhs.z == rhs.z + } + static func !=(lhs: MTLOrigin, rhs: MTLOrigin) -> Bool { + return lhs.x != rhs.x + || lhs.y != rhs.y + || lhs.z != rhs.z + } } diff --git a/Alloy/MTLRenderCommandEncoder+Extensions.swift b/Alloy/MTLRenderCommandEncoder+Extensions.swift index 3b32d6c..2732b07 100644 --- a/Alloy/MTLRenderCommandEncoder+Extensions.swift +++ b/Alloy/MTLRenderCommandEncoder+Extensions.swift @@ -2,41 +2,61 @@ import Metal public extension MTLRenderCommandEncoder { - func set(vertexValue value: T, at index: Int) { + func setVertexValue(_ value: T, at index: Int) { var t = value self.setVertexBytes(&t, length: MemoryLayout.stride, index: index) } - func set(vertexValue value: T, at index: Int) where T: Collection { + func setVertexValue(_ value: T, at index: Int) where T: Collection { var t = value self.setVertexBytes(&t, length: MemoryLayout.stride * value.count, index: index) } - func set(fragmentValue value: T, at index: Int) { + func setFragmentValue(_ value: T, at index: Int) { var t = value self.setFragmentBytes(&t, length: MemoryLayout.stride, index: index) } - func set(fragmentValue value: T, at index: Int) where T: Collection { + func setFragmentValue(_ value: T, at index: Int) where T: Collection { var t = value self.setFragmentBytes(&t, length: MemoryLayout.stride * value.count, index: index) } - func set(vertexTextures textures: [MTLTexture?], startingAt startIndex: Int = 0) { + func setVertexTextures(_ textures: [MTLTexture?], startingAt startIndex: Int = 0) { self.setVertexTextures(textures, range: startIndex..<(startIndex + textures.count)) } - func set(fragmentTextures textures: [MTLTexture?], startingAt startIndex: Int = 0) { + func setVertexTextures(_ textures: MTLTexture?..., startingAt startIndex: Int = 0) { + self.setVertexTextures(textures, range: startIndex..<(startIndex + textures.count)) + } + + func setFragmentTextures(_ textures: [MTLTexture?], startingAt startIndex: Int = 0) { + self.setFragmentTextures(textures, range: startIndex..<(startIndex + textures.count)) + } + + func setFragmentTextures(_ textures: MTLTexture?..., startingAt startIndex: Int = 0) { self.setFragmentTextures(textures, range: startIndex..<(startIndex + textures.count)) } - func set(vertexBuffers buffers: [MTLBuffer?], offsets: [Int]? = nil, startingAt startIndex: Int = 0) { + func setVertexBuffers(_ buffers: [MTLBuffer?], offsets: [Int]? = nil, startingAt startIndex: Int = 0) { + self.setVertexBuffers(buffers, + offsets: offsets ?? buffers.map { _ in 0 }, + range: startIndex..<(startIndex + buffers.count)) + } + + func setVertexBuffers(_ buffers: MTLBuffer?..., offsets: [Int]? = nil, startingAt startIndex: Int = 0) { self.setVertexBuffers(buffers, offsets: offsets ?? buffers.map { _ in 0 }, range: startIndex..<(startIndex + buffers.count)) } - func set(fragmentBuffers buffers: [MTLBuffer?], offsets: [Int]? = nil, startingAt startIndex: Int = 0) { + func setFragmentBuffers(_ buffers: [MTLBuffer?], offsets: [Int]? = nil, startingAt startIndex: Int = 0) { + self.setFragmentBuffers(buffers, + offsets: offsets ?? buffers.map { _ in 0 }, + range: startIndex..<(startIndex + buffers.count)) + } + + func setFragmentBuffers(_ buffers: MTLBuffer?..., offsets: [Int]? = nil, startingAt startIndex: Int = 0) { self.setFragmentBuffers(buffers, offsets: offsets ?? buffers.map { _ in 0 }, range: startIndex..<(startIndex + buffers.count)) diff --git a/Alloy/MTLSize+Extensions.swift b/Alloy/MTLSize+Extensions.swift new file mode 100644 index 0000000..8c2658c --- /dev/null +++ b/Alloy/MTLSize+Extensions.swift @@ -0,0 +1,28 @@ +import Metal + +public extension MTLSize { + init(repeating value: Int) { + self.init(width: value, + height: value, + depth: value) + } + + func clamped(to size: MTLSize) -> MTLSize { + return MTLSize(width: min(max(self.width, 0), size.width), + height: min(max(self.height, 0), size.height), + depth: min(max(self.depth, 0), size.depth)) + } + + static let one = MTLSize(repeating: 1) + static let zero = MTLSize(repeating: 0) + static func ==(lhs: MTLSize, rhs: MTLSize) -> Bool { + return lhs.width == rhs.width + && lhs.height == rhs.height + && lhs.depth == rhs.depth + } + static func !=(lhs: MTLSize, rhs: MTLSize) -> Bool { + return lhs.width != rhs.width + || lhs.height != rhs.height + || lhs.depth != rhs.depth + } +} diff --git a/Alloy/MTLSizeAndAlign+Extensions.swift b/Alloy/MTLSizeAndAlign+Extensions.swift index 0107d00..0721467 100644 --- a/Alloy/MTLSizeAndAlign+Extensions.swift +++ b/Alloy/MTLSizeAndAlign+Extensions.swift @@ -28,20 +28,3 @@ public extension Sequence where Element == MTLTextureDescriptor { } } } - -public extension MTLSize { - init(repeating value: Int) { - self.init(width: value, - height: value, - depth: value) - } - - func clamped(to size: MTLSize) -> MTLSize { - return MTLSize(width: min(max(self.width, 0), size.width), - height: min(max(self.height, 0), size.height), - depth: min(max(self.depth, 0), size.depth)) - } - - static let one = MTLSize(repeating: 1) - static let zero = MTLSize(repeating: 0) -} diff --git a/Alloy/MTLTexture+Extensions.swift b/Alloy/MTLTexture+Extensions.swift index 40cfaaa..edd56ef 100644 --- a/Alloy/MTLTexture+Extensions.swift +++ b/Alloy/MTLTexture+Extensions.swift @@ -20,18 +20,20 @@ public extension MTLTexture { let rowBytes = self.width let length = rowBytes * self.height - let rgbaBytes = [UInt8](repeating: 0, count: length) - self.getBytes(UnsafeMutableRawPointer(mutating: rgbaBytes), - bytesPerRow: rowBytes, - from: self.region, - mipmapLevel: 0) + let rgbaBytes = UnsafeMutableRawPointer.allocate(byteCount: length, + alignment: MemoryLayout.alignment) + defer { rgbaBytes.deallocate() } + self.getBytes(rgbaBytes, + bytesPerRow: rowBytes, + from: self.region, + mipmapLevel: 0) let colorScape = colorSpace ?? CGColorSpaceCreateDeviceGray() let bitmapInfo = CGBitmapInfo(rawValue: self.pixelFormat == .a8Unorm ? CGImageAlphaInfo.alphaOnly.rawValue : CGImageAlphaInfo.none.rawValue) guard let data = CFDataCreate(nil, - rgbaBytes, + rgbaBytes.assumingMemoryBound(to: UInt8.self), length), let dataProvider = CGDataProvider(data: data), let cgImage = CGImage(width: self.width, @@ -52,19 +54,26 @@ public extension MTLTexture { // read texture as byte array let rowBytes = self.width * 4 let length = rowBytes * self.height - let bgraBytes = [UInt8](repeating: 0, count: length) - self.getBytes(UnsafeMutableRawPointer(mutating: bgraBytes), - bytesPerRow: rowBytes, - from: self.region, - mipmapLevel: 0) + + let bgraBytes = UnsafeMutableRawPointer.allocate(byteCount: length, + alignment: MemoryLayout.alignment) + defer { bgraBytes.deallocate() } + + self.getBytes(bgraBytes, + bytesPerRow: rowBytes, + from: self.region, + mipmapLevel: 0) // use Accelerate framework to convert from BGRA to RGBA - var bgraBuffer = vImage_Buffer(data: UnsafeMutableRawPointer(mutating: bgraBytes), + var bgraBuffer = vImage_Buffer(data: bgraBytes, height: vImagePixelCount(self.height), width: vImagePixelCount(self.width), rowBytes: rowBytes) - let rgbaBytes = [UInt8](repeating: 0, count: length) - var rgbaBuffer = vImage_Buffer(data: UnsafeMutableRawPointer(mutating: rgbaBytes), + + let rgbaBytes = UnsafeMutableRawPointer.allocate(byteCount: length, + alignment: MemoryLayout.alignment) + defer { rgbaBytes.deallocate() } + var rgbaBuffer = vImage_Buffer(data: rgbaBytes, height: vImagePixelCount(self.height), width: vImagePixelCount(self.width), rowBytes: rowBytes) @@ -77,7 +86,7 @@ public extension MTLTexture { let colorScape = colorSpace ?? CGColorSpaceCreateDeviceRGB() let bitmapInfo = CGBitmapInfo(rawValue: CGImageAlphaInfo.premultipliedLast.rawValue) guard let data = CFDataCreate(nil, - rgbaBytes, + rgbaBytes.assumingMemoryBound(to: UInt8.self), length), let dataProvider = CGDataProvider(data: data), let cgImage = CGImage(width: self.width, @@ -98,9 +107,11 @@ public extension MTLTexture { let rowBytes = self.width * 4 let length = rowBytes * self.height - let rgbaBytes = [UInt8](repeating: 0, - count: length) - self.getBytes(UnsafeMutableRawPointer(mutating: rgbaBytes), + let rgbaBytes = UnsafeMutableRawPointer.allocate(byteCount: length, + alignment: MemoryLayout.alignment) + defer { rgbaBytes.deallocate() } + + self.getBytes(rgbaBytes, bytesPerRow: rowBytes, from: self.region, mipmapLevel: 0) @@ -108,7 +119,7 @@ public extension MTLTexture { let colorScape = colorSpace ?? CGColorSpaceCreateDeviceRGB() let bitmapInfo = CGBitmapInfo(rawValue: CGImageAlphaInfo.premultipliedLast.rawValue) guard let data = CFDataCreate(nil, - rgbaBytes, + rgbaBytes.assumingMemoryBound(to: UInt8.self), length), let dataProvider = CGDataProvider(data: data), let cgImage = CGImage(width: self.width, diff --git a/Alloy/Renderers/LinesRenderer.swift b/Alloy/Renderers/LinesRenderer.swift index d04aed2..3ac99d1 100644 --- a/Alloy/Renderers/LinesRenderer.swift +++ b/Alloy/Renderers/LinesRenderer.swift @@ -130,8 +130,7 @@ final public class LinesRenderer { renderEncoder.setVertexBuffer(self.linesBuffer, offset: 0, index: 0) - renderEncoder.set(fragmentValue: self.color, - at: 0) + renderEncoder.setFragmentValue(self.color, at: 0) // Draw. renderEncoder.drawPrimitives(type: .triangleStrip, vertexStart: 0, diff --git a/Alloy/Renderers/MaskRenderer.swift b/Alloy/Renderers/MaskRenderer.swift index dc8a4b4..e875b91 100644 --- a/Alloy/Renderers/MaskRenderer.swift +++ b/Alloy/Renderers/MaskRenderer.swift @@ -8,7 +8,7 @@ public class MaskRenderer { /// Mask color. Red in default. public var color: SIMD4 = .init(1, 0, 0, 0.3) /// Texture containig mask information. - public var maskTexture: MTLTexture? = nil + public var mask: MTLTexture? = nil /// Rectrangle described in a normalized coodrinate system. public var normalizedRect: CGRect = .zero @@ -133,14 +133,11 @@ public class MaskRenderer { renderEncoder.setRenderPipelineState(renderPipelineState) // Set any buffers fed into our render pipeline. let rectangle = self.constructRectangle() - renderEncoder.set(vertexValue: rectangle, - at: 0) - renderEncoder.setFragmentTexture(self.maskTexture, + renderEncoder.setVertexValue(rectangle, at: 0) + renderEncoder.setFragmentTexture(self.mask, index: 0) - renderEncoder.set(fragmentValue: self.color, - at: 0) - renderEncoder.set(fragmentValue: isInversed, - at: 1) + renderEncoder.setFragmentValue(self.color, at: 0) + renderEncoder.setFragmentValue(isInversed, at: 1) // Draw. renderEncoder.drawPrimitives(type: .triangleStrip, vertexStart: 0, diff --git a/Alloy/Renderers/PointsRenderer.swift b/Alloy/Renderers/PointsRenderer.swift index d9019bf..b7252db 100644 --- a/Alloy/Renderers/PointsRenderer.swift +++ b/Alloy/Renderers/PointsRenderer.swift @@ -128,10 +128,8 @@ final public class PointsRenderer { renderEncoder.setVertexBuffer(self.pointsPositionsBuffer, offset: 0, index: 0) - renderEncoder.set(vertexValue: self.pointSize, - at: 1) - renderEncoder.set(fragmentValue: self.color, - at: 0) + renderEncoder.setVertexValue(self.pointSize, at: 1) + renderEncoder.setFragmentValue(self.color, at: 0) // Draw. renderEncoder.drawPrimitives(type: .point, vertexStart: 0, diff --git a/Alloy/Renderers/RectangleRenderer.swift b/Alloy/Renderers/RectangleRenderer.swift index 0bda5dc..1f03734 100644 --- a/Alloy/Renderers/RectangleRenderer.swift +++ b/Alloy/Renderers/RectangleRenderer.swift @@ -124,10 +124,8 @@ final public class RectangleRenderer { renderEncoder.setRenderPipelineState(renderPipelineState) // Set any buffers fed into our render pipeline. let rectangle = self.constructRectangle() - renderEncoder.set(vertexValue: rectangle, - at: 0) - renderEncoder.set(fragmentValue: self.color, - at: 0) + renderEncoder.setVertexValue(rectangle, at: 0) + renderEncoder.setFragmentValue(self.color, at: 0) // Draw. renderEncoder.drawPrimitives(type: .triangleStrip, vertexStart: 0, diff --git a/Alloy/Renderers/SimpleGeometryRenderer.swift b/Alloy/Renderers/SimpleGeometryRenderer.swift index 57521e9..fb024ce 100644 --- a/Alloy/Renderers/SimpleGeometryRenderer.swift +++ b/Alloy/Renderers/SimpleGeometryRenderer.swift @@ -87,10 +87,8 @@ final public class SimpleGeometryRenderer { renderEncoder.setVertexBuffer(geometry, offset: 0, index: 0) - renderEncoder.set(vertexValue: matrix, - at: 1) - renderEncoder.set(fragmentValue: color, - at: 0) + renderEncoder.setVertexValue(matrix, at: 1) + renderEncoder.setFragmentValue(color, at: 0) renderEncoder.setTriangleFillMode(fillMode) renderEncoder.setRenderPipelineState(renderPipelineState) renderEncoder.drawIndexedPrimitives(type: type, diff --git a/Alloy/Shaders/Shaders.metal b/Alloy/Shaders/Shaders.metal index ea5d847..3a158d3 100644 --- a/Alloy/Shaders/Shaders.metal +++ b/Alloy/Shaders/Shaders.metal @@ -1,11 +1,3 @@ -// -// Shaders.metal -// AIBeauty -// -// Created by Andrey Volodin on 08.08.2018. -// Copyright © 2018 Andrey Volodin. All rights reserved. -// - #include #include "ColorConversion.h" #include "ShaderStructures.h" @@ -15,7 +7,8 @@ using namespace metal; constant bool deviceSupportsNonuniformThreadgroups [[ function_constant(0) ]]; constant bool deviceDoesntSupportNonuniformThreadgroups = !deviceSupportsNonuniformThreadgroups; -constant float multiplierFC [[function_constant(1)]]; +constant float multiplierFC [[ function_constant(1) ]]; +constant bool halfSizedCbCr [[ function_constant(2) ]]; struct BlockSize { ushort width; @@ -27,33 +20,33 @@ struct BlockSize { // MARK: - Texture Copy template -void textureCopy(texture2d sourceTexture, - texture2d destinationTexture, +void textureCopy(texture2d source, + texture2d destination, constant short2& readOffset, constant short2& writeOffset, constant ushort2& gridSize, const ushort2 position) { - const ushort2 readPosition = ushort2(short2(position) + readOffset); - const ushort2 writePosition = ushort2(short2(position) + writeOffset); checkPosition(position, gridSize, deviceSupportsNonuniformThreadgroups); - const auto resultValue = sourceTexture.read(readPosition); + const auto readPosition = ushort2(short2(position) + readOffset); + const auto writePosition = ushort2(short2(position) + writeOffset); - destinationTexture.write(resultValue, writePosition); + const auto resultValue = source.read(readPosition); + destination.write(resultValue, writePosition); } -#define outerArguments(T) \ -(texture2d sourceTexture [[ texture(0) ]], \ -texture2d destinationTexture [[ texture(1) ]], \ -constant short2& readOffset [[ buffer(0) ]], \ -constant short2& writeOffset [[ buffer(1) ]], \ -constant ushort2& gridSize [[ buffer(2), \ - function_constant(deviceDoesntSupportNonuniformThreadgroups) ]], \ -const ushort2 position [[ thread_position_in_grid ]]) \ +#define outerArguments(T) \ +(texture2d source [[ texture(0) ]], \ +texture2d destination [[ texture(1) ]], \ +constant short2& readOffset [[ buffer(0) ]], \ +constant short2& writeOffset [[ buffer(1) ]], \ +constant ushort2& gridSize [[ buffer(2), \ +function_constant(deviceDoesntSupportNonuniformThreadgroups) ]], \ +const ushort2 position [[ thread_position_in_grid ]]) \ #define innerArguments \ -(sourceTexture, \ -destinationTexture, \ +(source, \ +destination, \ readOffset, \ writeOffset, \ gridSize, \ @@ -66,55 +59,64 @@ generateKernels(textureCopy) // MARK: - Resize Texture -kernel void textureResize(texture2d sourceTexture [[ texture(0) ]], - texture2d destinationTexture [[ texture(1) ]], +kernel void textureResize(texture2d source [[ texture(0) ]], + texture2d destination [[ texture(1) ]], sampler s [[ sampler(0) ]], const ushort2 position [[ thread_position_in_grid ]]) { - - const ushort2 textureSize = ushort2(destinationTexture.get_width(), - destinationTexture.get_height()); + const auto textureSize = ushort2(destination.get_width(), + destination.get_height()); checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); - const float2 normalizedCoord = float2((float(position.x) + 0.5) / textureSize.x, - (float(position.y) + 0.5) / textureSize.y); + const auto positionF = float2(position); + const auto textureSizeF = float2(textureSize); + const auto normalizedPosition = (positionF + 0.5f) / textureSizeF; - auto sampledValue = sourceTexture.sample(s, normalizedCoord); - destinationTexture.write(sampledValue, position); + auto sampledValue = source.sample(s, normalizedPosition); + destination.write(sampledValue, position); } // MARK: - Texture Mask template -void textureMask(texture2d sourceTexture, - texture2d maskTexture, - texture2d destinationTexture, - const ushort2 position [[thread_position_in_grid]]) { - const ushort2 textureSize = ushort2(destinationTexture.get_width(), - destinationTexture.get_height()); +void textureMask(texture2d source, + texture2d mask, + texture2d destination, + constant bool& isInversed, + const ushort2 position) { + const auto textureSize = ushort2(destination.get_width(), + destination.get_height()); checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); - const auto originalPixel = sourceTexture.read(position); + const auto sourceValue = float4(source.read(position)); constexpr sampler s(coord::normalized, address::clamp_to_edge, filter::linear); + const auto positionF = float2(position); + const auto textureSizeF = float2(textureSize); + const auto normalizedPosition = (positionF.x + 0.5f) / textureSizeF; - const auto maskValue = maskTexture.sample(s, (float2(position) + 0.5) / float2(textureSize)); - const auto resultValue = vec(float4(originalPixel) * maskValue.r); + auto maskValue = mask.sample(s, normalizedPosition); + if (isInversed) { + maskValue = 1.0f - maskValue; + } + const auto resultValue = vec(sourceValue * maskValue.r); - destinationTexture.write(resultValue, position); + destination.write(resultValue, position); } -#define outerArguments(T) \ -(texture2d sourceTexture [[ texture(0) ]], \ -texture2d maskTexture [[ texture(1) ]], \ -texture2d destinationTexture [[ texture(2) ]], \ -const ushort2 position [[thread_position_in_grid]]) \ +#define outerArguments(T) \ +(texture2d source [[ texture(0) ]], \ +texture2d mask [[ texture(1) ]], \ +texture2d destination [[ texture(2) ]], \ +constant bool& isInversed [[ buffer(0) ]], \ +const ushort2 position [[thread_position_in_grid]]) \ #define innerArguments \ -(sourceTexture, \ -maskTexture, \ -destinationTexture, \ +(source, \ +mask, \ +destination, \ +isInversed, \ position) \ generateKernels(textureMask) @@ -125,32 +127,32 @@ generateKernels(textureMask) // MARK: - Texture Max template -void textureMax(texture2d sourceTexture, +void textureMax(texture2d source, constant BlockSize& inputBlockSize, device float4& result, threadgroup float4* sharedMemory, const ushort index, const ushort2 position, const ushort2 threadsPerThreadgroup) { - const ushort2 textureSize = ushort2(sourceTexture.get_width(), - sourceTexture.get_height()); + const auto textureSize = ushort2(source.get_width(), + source.get_height()); - ushort2 originalBlockSize = ushort2(inputBlockSize.width, - inputBlockSize.height); - const ushort2 blockStartPosition = position * originalBlockSize; + const auto originalBlockSize = ushort2(inputBlockSize.width, + inputBlockSize.height); + const auto blockStartPosition = position * originalBlockSize; - ushort2 block_size = originalBlockSize; + auto blockSize = originalBlockSize; if (position.x == threadsPerThreadgroup.x || position.y == threadsPerThreadgroup.y) { - const ushort2 readTerritory = blockStartPosition + originalBlockSize; - block_size = originalBlockSize - (readTerritory - textureSize); + const auto readTerritory = blockStartPosition + originalBlockSize; + blockSize = originalBlockSize - (readTerritory - textureSize); } - float4 maxValueInBlock = float4(sourceTexture.read(blockStartPosition)); + auto maxValueInBlock = float4(source.read(blockStartPosition)); - for (ushort x = 0; x < block_size.x; x++) { - for (ushort y = 0; y < block_size.y; y++) { - const ushort2 readPosition = blockStartPosition + ushort2(x, y); - const float4 currentValue = float4(sourceTexture.read(readPosition)); + for (ushort x = 0; x < blockSize.x; x++) { + for (ushort y = 0; y < blockSize.y; y++) { + const auto readPosition = blockStartPosition + ushort2(x, y); + const auto currentValue = float4(source.read(readPosition)); maxValueInBlock = max(maxValueInBlock, currentValue); } } @@ -160,20 +162,17 @@ void textureMax(texture2d sourceTexture, threadgroup_barrier(mem_flags::mem_threadgroup); if (index == 0) { - auto maxValue = sharedMemory[0]; - const ushort threadsInThreadgroup = threadsPerThreadgroup.x * threadsPerThreadgroup.y; + const auto threadsInThreadgroup = threadsPerThreadgroup.x * threadsPerThreadgroup.y; for (ushort i = 1; i < threadsInThreadgroup; i++) { - float4 maxValueInBlock = sharedMemory[i]; - maxValue = max(maxValue, maxValueInBlock); + maxValue = max(maxValue, sharedMemory[i]); } - result = maxValue; } } #define outerArguments(T) \ -(texture2d sourceTexture [[ texture(0) ]], \ +(texture2d source [[ texture(0) ]], \ constant BlockSize& inputBlockSize [[ buffer(0) ]], \ device float4& result [[ buffer(1) ]], \ threadgroup float4* sharedMemory [[ threadgroup(0) ]], \ @@ -182,7 +181,7 @@ const ushort2 position [[ thread_position_in_grid ]], \ const ushort2 threadsPerThreadgroup [[ threads_per_threadgroup ]]) \ #define innerArguments \ -(sourceTexture, \ +(source, \ inputBlockSize, \ result, \ sharedMemory, \ @@ -198,32 +197,32 @@ generateKernels(textureMax) // MARK: - Texture Min template -void textureMin(texture2d sourceTexture, +void textureMin(texture2d source, constant BlockSize& inputBlockSize, device float4& result, threadgroup float4* sharedMemory, const ushort index, const ushort2 position, const ushort2 threadsPerThreadgroup) { - const ushort2 textureSize = ushort2(sourceTexture.get_width(), - sourceTexture.get_height()); + const auto textureSize = ushort2(source.get_width(), + source.get_height()); - ushort2 originalBlockSize = ushort2(inputBlockSize.width, - inputBlockSize.height); - const ushort2 blockStartPosition = position * originalBlockSize; + auto originalBlockSize = ushort2(inputBlockSize.width, + inputBlockSize.height); + const auto blockStartPosition = position * originalBlockSize; - ushort2 blockSize = originalBlockSize; + auto blockSize = originalBlockSize; if (position.x == threadsPerThreadgroup.x || position.y == threadsPerThreadgroup.y) { - const ushort2 readTerritory = blockStartPosition + originalBlockSize; + const auto readTerritory = blockStartPosition + originalBlockSize; blockSize = originalBlockSize - (readTerritory - textureSize); } - float4 minValueInBlock = float4(sourceTexture.read(blockStartPosition)); + auto minValueInBlock = float4(source.read(blockStartPosition)); for (ushort x = 0; x < blockSize.x; x++) { for (ushort y = 0; y < blockSize.y; y++) { - const ushort2 readPosition = blockStartPosition + ushort2(x, y); - const float4 currentValue = float4(sourceTexture.read(readPosition)); + const auto readPosition = blockStartPosition + ushort2(x, y); + const auto currentValue = float4(source.read(readPosition)); minValueInBlock = min(minValueInBlock, currentValue); } } @@ -234,17 +233,16 @@ void textureMin(texture2d sourceTexture, if (index == 0) { float4 minValue = sharedMemory[0]; - const ushort threadsInThreadgroup = threadsPerThreadgroup.x * threadsPerThreadgroup.y; + const auto threadsInThreadgroup = threadsPerThreadgroup.x * threadsPerThreadgroup.y; for (ushort i = 1; i < threadsInThreadgroup; i++) { - float4 minValueInBlock = sharedMemory[i]; - minValue = min(minValue, minValueInBlock); + minValue = min(minValue, sharedMemory[i]); } result = minValue; } } #define outerArguments(T) \ -(texture2d sourceTexture [[ texture(0) ]], \ +(texture2d source [[ texture(0) ]], \ constant BlockSize& inputBlockSize [[ buffer(0) ]], \ device float4& result [[ buffer(1) ]], \ threadgroup float4* sharedMemory [[ threadgroup(0) ]], \ @@ -253,7 +251,7 @@ const ushort2 position [[ thread_position_in_grid ]], \ const ushort2 threadsPerThreadgroup [[ threads_per_threadgroup ]]) \ #define innerArguments \ -(sourceTexture, \ +(source, \ inputBlockSize, \ result, \ sharedMemory, \ @@ -269,32 +267,32 @@ generateKernels(textureMin) // MARK: - Texture Mean template -void textureMean(texture2d sourceTexture, +void textureMean(texture2d source, constant BlockSize& inputBlockSize, device float4& result, threadgroup float4* sharedMemory, const ushort index, const ushort2 position, const ushort2 threadsPerThreadgroup) { - const ushort2 inputTextureSize = ushort2(sourceTexture.get_width(), - sourceTexture.get_height()); + const auto textureSize = ushort2(source.get_width(), + source.get_height()); - ushort2 originalBlockSize = ushort2(inputBlockSize.width, - inputBlockSize.height); - const ushort2 blockStartPosition = position * originalBlockSize; + auto originalBlockSize = ushort2(inputBlockSize.width, + inputBlockSize.height); + const auto blockStartPosition = position * originalBlockSize; - ushort2 blockSize = originalBlockSize; + auto blockSize = originalBlockSize; if (position.x == threadsPerThreadgroup.x || position.y == threadsPerThreadgroup.y) { - const ushort2 readTerritory = blockStartPosition + originalBlockSize; - blockSize = originalBlockSize - (readTerritory - inputTextureSize); + const auto readTerritory = blockStartPosition + originalBlockSize; + blockSize = originalBlockSize - (readTerritory - textureSize); } - float4 totalSumInBlock = float4(0, 0, 0, 0); + auto totalSumInBlock = float4(0); for (ushort x = 0; x < blockSize.x; x++) { for (ushort y = 0; y < blockSize.y; y++) { - const ushort2 read_position = blockStartPosition + ushort2(x, y); - const float4 currentValue = float4(sourceTexture.read(read_position)); + const auto readPosition = blockStartPosition + ushort2(x, y); + const auto currentValue = float4(source.read(readPosition)); totalSumInBlock += currentValue; } } @@ -305,22 +303,21 @@ void textureMean(texture2d sourceTexture, if (index == 0) { - float4 totalSum = sharedMemory[0]; - const ushort threadsInThreadgroup = threadsPerThreadgroup.x * threadsPerThreadgroup.y; + auto totalSum = sharedMemory[0]; + const auto threadsInThreadgroup = threadsPerThreadgroup.x * threadsPerThreadgroup.y; for (ushort i = 1; i < threadsInThreadgroup; i++) { - float4 totalSumInBlock = sharedMemory[i]; - totalSum += totalSumInBlock; + totalSum += sharedMemory[i]; } - half gridSize = inputTextureSize.x * inputTextureSize.y; - float4 meanValue = totalSum / gridSize; + auto gridSize = textureSize.x * textureSize.y; + auto meanValue = totalSum / gridSize; result = meanValue; } } #define outerArguments(T) \ -(texture2d sourceTexture [[ texture(0) ]], \ +(texture2d source [[ texture(0) ]], \ constant BlockSize& inputBlockSize [[ buffer(0) ]], \ device float4& result [[ buffer(1) ]], \ threadgroup float4* sharedMemory [[ threadgroup(0) ]], \ @@ -329,7 +326,7 @@ const ushort2 position [[ thread_position_in_grid ]], \ const ushort2 threadsPerThreadgroup [[ threads_per_threadgroup ]]) \ #define innerArguments \ -(sourceTexture, \ +(source, \ inputBlockSize, \ result, \ sharedMemory, \ @@ -344,101 +341,92 @@ generateKernels(textureMean) // MARK: - Mask Guided Blur -kernel void maskGuidedBlurRowPass(texture2d sourceTexture [[ texture(0) ]], - texture2d maskTexture [[ texture(1) ]], - texture2d destinationTexture [[ texture(2) ]], +kernel void maskGuidedBlurRowPass(texture2d source [[ texture(0) ]], + texture2d mask [[ texture(1) ]], + texture2d destination [[ texture(2) ]], constant float& sigma [[ buffer(0) ]], - ushort2 position [[thread_position_in_grid]]) { - const ushort sourceTextureWidth = sourceTexture.get_width(); - const ushort sourceTextureHeight = sourceTexture.get_height(); + ushort2 position [[ thread_position_in_grid ]]) { + const auto textureSize = ushort2(source.get_width(), + source.get_height()); - if (!deviceSupportsNonuniformThreadgroups) { - if (position.x >= sourceTextureWidth || position.y >= sourceTextureHeight) { - return; - } - } + checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); - constexpr sampler s(filter::linear, coord::normalized); + constexpr sampler s(coord::normalized, + address::clamp_to_edge, + filter::linear); - const float2 srcTid = float2((float(position.x) + 0.5f) / sourceTextureWidth, - (float(position.y) + 0.5f) / sourceTextureHeight); + const auto positionF = float2(position); + const auto textureSizeF = float2(textureSize); + const auto normalizedPosition = (positionF + 0.5f) / textureSizeF; - const float maskValue = maskTexture.sample(s, srcTid).r; + const auto maskValue = mask.sample(s, normalizedPosition).r; - const float sigma_ = (1.0f - maskValue) * sigma; - const int kernelRadius = int(2.0f * sigma_); + const auto sigmaValue = (1.0f - maskValue) * sigma; + const auto kernelRadius = int(2.0f * sigmaValue); - float normalizingConstant = 0.0f; - float3 result = float3(0.0f); + auto normalizingConstant = 0.0f; + auto result = float3(0.0f); - for (int drow = -kernelRadius; drow <= kernelRadius; drow++) { - const float kernelValue = exp(float(-drow * drow) / (2.0f * sigma_ * sigma_ + 1e-5f)); - const uint2 coordinate = uint2(clamp(int(position.x) + drow, 0, sourceTextureWidth - 1), - position.y); - const float2 maskTid = float2(float(coordinate.x) / sourceTextureWidth, - float(coordinate.y) / sourceTextureHeight); - const float maskMultiplier = 1.0f - maskTexture.sample(s, maskTid).r + 1e-5f; - const float totalFactor = kernelValue * maskMultiplier; - normalizingConstant += totalFactor; - result += sourceTexture.read(coordinate).rgb * totalFactor; + for (int row = -kernelRadius; row <= kernelRadius; row++) { + const auto kernelValue = exp(float(-row * row) / (2.0f * sigmaValue * sigmaValue + 1e-5f)); + const auto readPosition = uint2(clamp(position.x + row, 0, position.y - 1), position.y); + const auto readPositionF = float2(readPosition); + const auto normalizedPosition = (readPositionF.x + 0.5f) / textureSizeF; + const auto maskMultiplier = 1.0f - mask.sample(s, normalizedPosition).r + 1e-5f; + const auto totalFactor = kernelValue * maskMultiplier; + normalizingConstant += float(totalFactor); + result += source.read(readPosition).rgb * totalFactor; } result /= normalizingConstant; - destinationTexture.write(float4(result, 1.0f), position); + destination.write(float4(result, 1.0f), position); } -kernel void maskGuidedBlurColumnPass(texture2d sourceTexture [[ texture(0) ]], - texture2d maskTexture [[ texture(1) ]], - texture2d destinationTexture [[ texture(2) ]], +kernel void maskGuidedBlurColumnPass(texture2d source [[ texture(0) ]], + texture2d mask [[ texture(1) ]], + texture2d destination [[ texture(2) ]], constant float& sigma [[ buffer(0) ]], - ushort2 position [[thread_position_in_grid]]) { - const ushort sourceTextureWidth = sourceTexture.get_width(); - const ushort sourceTextureHeight = sourceTexture.get_height(); + uint2 position [[ thread_position_in_grid ]]) { + const auto textureSize = ushort2(source.get_width(), + source.get_height()); - if (!deviceSupportsNonuniformThreadgroups) { - if (position.x >= sourceTextureWidth || position.y >= sourceTextureHeight) { - return; - } - } + checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); - constexpr sampler s(filter::linear, coord::normalized); + constexpr sampler s(coord::normalized, + address::clamp_to_edge, + filter::linear); - const float2 srcTid = float2((float(position.x) + 0.5f) / sourceTextureWidth, - (float(position.y) + 0.5f) / sourceTextureHeight); + const auto textureSizeF = float2(textureSize); + const auto positionF = float2(position); + const auto normalizedPosition = (positionF.x + 0.5f) / textureSizeF; - const float maskValue = maskTexture.sample(s, srcTid).r; + const auto maskValue = mask.sample(s, normalizedPosition).r; - const float sigma_ = (1.0f - maskValue) * sigma; - const int kernelRadius = int(2.0f * sigma_); + const auto sigmaValue = (1.0f - maskValue) * sigma; + const auto kernelRadius = uint(2.0f * sigmaValue); - float normalizingConstant = 0.0f; - float3 result = float3(0.0f); + auto normalizingConstant = 0.0f; + auto result = float3(0.0f); - for (int dcol = -kernelRadius; dcol <= kernelRadius; dcol++) { - const float kernelValue = exp(float(-dcol * dcol) / (2.0f * sigma_ * sigma_ + 1e-5f)); - const uint2 coordinate = uint2(position.x, - clamp(int(position.y) + dcol, 0, sourceTextureHeight - 1)); - const float2 maskTid = float2(float(coordinate.x) / sourceTextureWidth, - float(coordinate.y) / sourceTextureHeight); - const float maskMultiplier = 1.0f - maskTexture.sample(s, maskTid).r + 1e-5f; - const float totalFactor = kernelValue * maskMultiplier; - normalizingConstant += totalFactor; - result += sourceTexture.read(coordinate).rgb * totalFactor; + for (uint column = -kernelRadius; column <= kernelRadius; column++) { + const auto kernelValue = exp(float(-column * column) / (2.0f * sigmaValue * sigmaValue + 1e-5f)); + const auto readPosition = uint2(position.x, clamp(position.y + column, 0u, position.y - 1)); + const auto readPositionF = float2(readPosition); + const auto normalizedPosition = (readPositionF.x + 0.5f) / textureSizeF; + const auto maskMultiplier = 1.0f - mask.sample(s, normalizedPosition).r + 1e-5f; + const auto totalFactor = kernelValue * maskMultiplier; + normalizingConstant += float(totalFactor); + result += source.read(readPosition).rgb * totalFactor; } result /= normalizingConstant; - destinationTexture.write(float4(result, 1.0f), position); + destination.write(float4(result, 1.0f), position); } // MARK: - Euclidean Distance -float euclideanDistance(float4 firstValue, float4 secondValue) { - const float4 diff = firstValue - secondValue; - return sqrt(dot(pow(diff, 2), 1)); -} - template void euclideanDistance(texture2d textureOne, texture2d textureTwo, @@ -448,16 +436,16 @@ void euclideanDistance(texture2d textureOne, const ushort index, const ushort2 position, const ushort2 threadsPerThreadgroup) { - const ushort2 textureSize = ushort2(textureOne.get_width(), - textureOne.get_height()); + const auto textureSize = ushort2(textureOne.get_width(), + textureOne.get_height()); - ushort2 originalBlockSize = ushort2(inputBlockSize.width, - inputBlockSize.height); - const ushort2 blockStartPosition = position * originalBlockSize; + auto originalBlockSize = ushort2(inputBlockSize.width, + inputBlockSize.height); + const auto blockStartPosition = position * originalBlockSize; - ushort2 blockSize = originalBlockSize; + auto blockSize = originalBlockSize; if (position.x == threadsPerThreadgroup.x || position.y == threadsPerThreadgroup.y) { - const ushort2 readTerritory = blockStartPosition + originalBlockSize; + const auto readTerritory = blockStartPosition + originalBlockSize; blockSize = originalBlockSize - (readTerritory - textureSize); } @@ -465,11 +453,10 @@ void euclideanDistance(texture2d textureOne, for (ushort x = 0; x < blockSize.x; x++) { for (ushort y = 0; y < blockSize.y; y++) { - const ushort2 readPosition = blockStartPosition + ushort2(x, y); - const float4 textureOneValue = float4(textureOne.read(readPosition)); - const float4 textureTwoValue = float4(textureTwo.read(readPosition)); - euclideanDistanceSumInBlock += euclideanDistance(textureOneValue, - textureTwoValue); + const auto readPosition = blockStartPosition + ushort2(x, y); + const auto textureOneValue = float4(textureOne.read(readPosition)); + const auto textureTwoValue = float4(textureTwo.read(readPosition)); + euclideanDistanceSumInBlock += sqrt(dot(pow(textureOneValue - textureTwoValue, 2), 1)); } } @@ -478,8 +465,8 @@ void euclideanDistance(texture2d textureOne, threadgroup_barrier(mem_flags::mem_threadgroup); if (index == 0) { - float totalEuclideanDistanceSum = sharedMemory[0]; - const ushort threadsInThreadgroup = threadsPerThreadgroup.x * threadsPerThreadgroup.y; + auto totalEuclideanDistanceSum = sharedMemory[0]; + const auto threadsInThreadgroup = threadsPerThreadgroup.x * threadsPerThreadgroup.y; for (ushort i = 1; i < threadsInThreadgroup; i++) { totalEuclideanDistanceSum += sharedMemory[i]; } @@ -513,32 +500,34 @@ generateKernels(euclideanDistance) #undef outerArguments #undef innerArguments +#undef euclidean + // MARK: - Add Constant template -void addConstant(texture2d sourceTexture, - texture2d destinationTexture, +void addConstant(texture2d source, + texture2d destination, constant float4& constantValue, const ushort2 position) { - const ushort2 textureSize = ushort2(sourceTexture.get_width(), - sourceTexture.get_height()); + const auto textureSize = ushort2(destination.get_width(), + destination.get_height()); checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); - auto sourceTextureValue = sourceTexture.read(position); - auto destinationTextureValue = sourceTextureValue + vec(constantValue); - destinationTexture.write(destinationTextureValue, position); + auto sourceValue = source.read(position); + auto destinationValue = sourceValue + vec(constantValue); + destination.write(destinationValue, position); } -#define outerArguments(T) \ -(texture2d sourceTexture [[ texture(0) ]], \ -texture2d destinationTexture [[ texture(1) ]], \ -constant float4& constantValue [[ buffer(0) ]], \ +#define outerArguments(T) \ +(texture2d source [[ texture(0) ]], \ +texture2d destination [[ texture(1) ]], \ +constant float4& constantValue [[ buffer(0) ]], \ const ushort2 position [[ thread_position_in_grid ]]) #define innerArguments \ -(sourceTexture, \ -destinationTexture, \ +(source, \ +destination, \ constantValue, \ position) @@ -548,28 +537,28 @@ generateKernels(addConstant) #undef innerArguments template -void divideByConstant(texture2d sourceTexture, - texture2d destinationTexture, +void divideByConstant(texture2d source, + texture2d destination, constant float4& constantValue, const ushort2 position) { - const ushort2 textureSize = ushort2(sourceTexture.get_width(), - sourceTexture.get_height()); + const auto textureSize = ushort2(source.get_width(), + source.get_height()); checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); - auto sourceTextureValue = sourceTexture.read(position); - auto destinationTextureValue = sourceTextureValue / vec(constantValue); - destinationTexture.write(destinationTextureValue, position); + auto sourceValue = source.read(position); + auto destinationValue = sourceValue / vec(constantValue); + destination.write(destinationValue, position); } -#define outerArguments(T) \ -(texture2d sourceTexture [[ texture(0) ]], \ -texture2d destinationTexture [[ texture(1) ]], \ -constant float4& constantValue [[ buffer(0) ]], \ +#define outerArguments(T) \ +(texture2d source [[ texture(0) ]], \ +texture2d destination [[ texture(1) ]], \ +constant float4& constantValue [[ buffer(0) ]], \ const ushort2 position [[ thread_position_in_grid ]]) #define innerArguments \ -(sourceTexture, \ -destinationTexture, \ +(source, \ +destination, \ constantValue, \ position) @@ -583,16 +572,23 @@ generateKernels(divideByConstant) kernel void textureMaskedMix(texture2d sourceOne [[ texture(0) ]], texture2d sourceTwo [[ texture(1) ]], - texture2d mask [[ texture(2) ]], + texture2d mask [[ texture(2) ]], texture2d destination [[ texture(3) ]], const ushort2 position [[ thread_position_in_grid ]]) { const ushort2 textureSize = ushort2(destination.get_width(), destination.get_height()); checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); + constexpr sampler s(coord::normalized, + address::clamp_to_edge, + filter::linear); + const auto positionF = float2(position); + const auto textureSizeF = float2(textureSize); + const auto normalizedPosition = (positionF + 0.5f) / textureSizeF; + const auto sourceOneValue = sourceOne.read(position); const auto sourceTwoValue = sourceTwo.read(position); - const auto maskValue = mask.read(position).r; + const auto maskValue = mask.sample(s, normalizedPosition).r; const auto resultValue = mix(sourceOneValue, sourceTwoValue, maskValue); @@ -619,44 +615,43 @@ kernel void textureWeightedMix(texture2d sourceOne [[ textu // MARK: - Texture Multiply Add -kernel void textureMultiplyAdd(texture2d sourceTextureOne [[ texture(0) ]], - texture2d sourceTextureTwo [[ texture(1) ]], - texture2d destinationTexture [[ texture(2) ]], +kernel void textureMultiplyAdd(texture2d sourceOne [[ texture(0) ]], + texture2d sourceTwo [[ texture(1) ]], + texture2d destination [[ texture(2) ]], const ushort2 position [[ thread_position_in_grid ]]) { - const ushort2 textureSize = ushort2(destinationTexture.get_width(), - destinationTexture.get_height()); + const auto textureSize = ushort2(destination.get_width(), + destination.get_height()); checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); - const auto sourceTextureOneValue = sourceTextureOne.read(position); - const auto sourceTextureTwoValue = sourceTextureTwo.read(position); - const auto destinationTextureValue = fma(sourceTextureTwoValue, - multiplierFC, - sourceTextureOneValue); - destinationTexture.write(destinationTextureValue, - position); + const auto sourceOneValue = sourceOne.read(position); + const auto sourceTwoValue = sourceTwo.read(position); + const auto destinationValue = fma(sourceTwoValue, + multiplierFC, + sourceOneValue); + destination.write(destinationValue, + position); } // MARK: - Texture Difference Hightlight -kernel void textureDifferenceHighlight(texture2d sourceTextureOne [[texture(0)]], - texture2d sourceTextureTwo [[texture(1)]], - texture2d destinationTexture [[texture(2)]], +kernel void textureDifferenceHighlight(texture2d sourceOne [[texture(0)]], + texture2d sourceTwo [[texture(1)]], + texture2d destination [[texture(2)]], constant float4& color [[ buffer(0) ]], constant float& threshold [[ buffer(1) ]], ushort2 position [[ thread_position_in_grid ]]) { - const ushort2 textureSize = ushort2(destinationTexture.get_width(), - destinationTexture.get_height()); + const auto textureSize = ushort2(destination.get_width(), + destination.get_height()); checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); - const float4 originalColor = sourceTextureOne.read(position); - const float4 targetColor = sourceTextureTwo.read(position); - const float4 difference = abs(targetColor - originalColor); - const float totalDifference = dot(difference, 1.f); - - destinationTexture.write(mix(targetColor, + const auto originalColor = sourceOne.read(position); + const auto targetColor = sourceTwo.read(position); + const auto difference = abs(targetColor - originalColor); + const auto totalDifference = dot(difference, 1.f); + const auto resultValue = mix(targetColor, color, - step(threshold, totalDifference)), - position); + step(threshold, totalDifference)); + destination.write(resultValue, position); } // MARK: - ML @@ -666,18 +661,15 @@ kernel void normalize(texture2d inputTexture [[ texture(0) ] constant float3& mean [[ buffer(0) ]], constant float3& std [[ buffer(1) ]], uint2 position [[thread_position_in_grid]]) { - const ushort2 textureSize = ushort2(inputTexture.get_width(), + const auto textureSize = ushort2(inputTexture.get_width(), inputTexture.get_height()); - if (!deviceSupportsNonuniformThreadgroups) { - if (position.x >= textureSize.x || position.y >= textureSize.y) { - return; - } - } + checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); + // Read mpsnngraph result value. - const half4 originalValue = inputTexture.read(position); - const half3 meanValue = (half3)mean; - const half3 stdValue = (half3)std; - half4 normalizedValue = originalValue; + const auto originalValue = inputTexture.read(position); + const auto meanValue = (half3)mean; + const auto stdValue = (half3)std; + auto normalizedValue = originalValue; normalizedValue.rgb -= meanValue; normalizedValue.rgb /= stdValue; outputTexture.write(normalizedValue, position); @@ -685,13 +677,13 @@ kernel void normalize(texture2d inputTexture [[ texture(0) ] // MARK: - Rendering -float2 perpendicular(float2 vector) { +inline float2 perpendicular(float2 vector) { return float2(-vector.y, vector.x); } -float2 convertToScreenSpace(float2 vector) { - return float2(-1 + (vector.x * 2), - -1 + ((1 - vector.y) * 2)); +inline float2 convertToScreenSpace(float2 vector) { + return float2(-1.0f + (vector.x * 2.0f), + -1.0f + ((1.0f - vector.y) * 2.0f)); } // MARK: - Rectangle Rendering @@ -708,7 +700,7 @@ vertex VertexOut rectVertex(constant Rectangle& rectangle [[ buffer(0) ]], }; VertexOut out; - float2 position = convertToScreenSpace(positions[vid]); + auto position = convertToScreenSpace(positions[vid]); out.position = float4(position, 0.0, 1.0); return out; @@ -739,27 +731,28 @@ vertex MaskVertexOut maskVertex(constant Rectangle& rectangle [[ buffer(0) ]], Vertex { rectangle.topRight, float2(1.0, 1.0) }, Vertex { rectangle.bottomRight, float2(1.0, 0.0) } }; - - MaskVertexOut out; - float2 position = convertToScreenSpace(vertices[vid].position); - out.position = float4(position, 0.0, 1.0); - out.uv = vertices[vid].uv; + const auto position = convertToScreenSpace(vertices[vid].position); + MaskVertexOut out = { + .position = float4(position, 0.0, 1.0), + .uv = vertices[vid].uv + }; return out; } fragment float4 maskFragment(MaskVertexOut in [[ stage_in ]], - texture2d maskTexture [[ texture(0) ]], + texture2d mask [[ texture(0) ]], constant float4& color [[ buffer(0) ]], constant bool& isInversed [[ buffer(1) ]]) { constexpr sampler s(coord::normalized, address::clamp_to_edge, filter::linear); - float4 maskValue = (float4)maskTexture.sample(s, in.uv).rrrr; + + auto maskValue = mask.sample(s, in.uv).rrrr; if (isInversed) { maskValue = 1.0f - maskValue; } - float4 resultColor = maskValue * color; + auto resultColor = maskValue * color; return resultColor; } @@ -767,16 +760,16 @@ fragment float4 maskFragment(MaskVertexOut in [[ stage_in ]], // MARK: - Lines Rendering vertex VertexOut linesVertex(constant Line *lines [[ buffer(0) ]], - uint vertexId [[vertex_id]], - uint instanceId [[instance_id]]) { + uint vertexId [[ vertex_id ]], + uint instanceId [[ instance_id ]]) { Line line = lines[instanceId]; - float2 startPoint = line.startPoint; - float2 endPoint = line.endPoint; + auto startPoint = line.startPoint; + auto endPoint = line.endPoint; - float2 vector = startPoint - endPoint; - float2 perpendicularVector = perpendicular(normalize(vector)); - float halfWidth = line.width / 2; + auto vector = startPoint - endPoint; + auto perpendicularVector = perpendicular(normalize(vector)); + auto halfWidth = line.width / 2.0f; struct PositionAndOffsetFactor { float2 vertexPosition; @@ -784,17 +777,16 @@ vertex VertexOut linesVertex(constant Line *lines [[ buffer(0) ]], }; const PositionAndOffsetFactor positionsAndOffsetFactors[] = { - PositionAndOffsetFactor { startPoint, -1.0 }, - PositionAndOffsetFactor { endPoint, -1.0 }, - PositionAndOffsetFactor { startPoint, 1.0 }, - PositionAndOffsetFactor { endPoint, 1.0 } + PositionAndOffsetFactor { startPoint, -1.0f }, + PositionAndOffsetFactor { endPoint, -1.0f }, + PositionAndOffsetFactor { startPoint, 1.0f }, + PositionAndOffsetFactor { endPoint, 1.0f } }; - VertexOut out; - const float2 vertexPosition = positionsAndOffsetFactors[vertexId].vertexPosition; - const float offsetFactor = positionsAndOffsetFactors[vertexId].offsetFactor; - float2 position = convertToScreenSpace(vertexPosition + offsetFactor * perpendicularVector * halfWidth); - out.position = float4(position, 0.0, 1.0); + const auto vertexPosition = positionsAndOffsetFactors[vertexId].vertexPosition; + const auto offsetFactor = positionsAndOffsetFactors[vertexId].offsetFactor; + const auto position = convertToScreenSpace(vertexPosition + offsetFactor * perpendicularVector * halfWidth); + VertexOut out = { .position = float4(position, 0.0f, 1.0f) }; return out; } @@ -808,13 +800,14 @@ struct PointVertexOut { vertex PointVertexOut pointVertex(constant float2* pointsPositions [[ buffer(0) ]], constant float& pointSize [[ buffer(1) ]], - uint instanceId [[instance_id]]) { - const float2 pointPosition = pointsPositions[instanceId]; + uint instanceId [[ instance_id ]]) { + const auto pointPosition = pointsPositions[instanceId]; - PointVertexOut out; - float2 position = convertToScreenSpace(pointPosition); - out.position = float4(position, 0, 1); - out.size = pointSize; + const auto position = convertToScreenSpace(pointPosition); + PointVertexOut out = { + .position = float4(position, 0, 1), + .size = pointSize + }; return out; } @@ -822,8 +815,9 @@ vertex PointVertexOut pointVertex(constant float2* pointsPositions [[ buffer(0) fragment float4 pointFragment(PointVertexOut in [[stage_in]], const float2 pointCenter [[ point_coord ]], constant float4& pointColor [[ buffer(0) ]]) { - const float distanceFromCenter = length(2 * (pointCenter - 0.5)); - float4 color = pointColor; + const auto distanceFromCenter = length(2 * (pointCenter - 0.5)); + + auto color = pointColor; color.a = 1.0 - smoothstep(0.9, 1.0, distanceFromCenter); return color; @@ -833,7 +827,6 @@ vertex float4 simpleVertex(constant float4* vertices [[ buffer(0) ]], constant float4x4& matrix [[ buffer(1) ]], uint vid [[vertex_id]]) { const float4 v = vertices[vid]; - return matrix * v; } @@ -848,8 +841,8 @@ kernel void lookUpTable(texture2d source [[ texture(0) ]], texture3d lut [[ texture(2) ]], constant float& intensity [[ buffer(0) ]], uint2 position [[thread_position_in_grid]]) { - const ushort2 textureSize = ushort2(destination.get_width(), - destination.get_height()); + const auto textureSize = ushort2(destination.get_width(), + destination.get_height()); checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); constexpr sampler s(coord::normalized, @@ -857,15 +850,15 @@ kernel void lookUpTable(texture2d source [[ texture(0) ]], filter::linear); // read original color - float4 sourceColor = source.read(position); + auto sourceValue = source.read(position); // use it to sample target color - sourceColor.rgb = mix(sourceColor.rgb, - lut.sample(s, sourceColor.rgb).rgb, + sourceValue.rgb = mix(sourceValue.rgb, + lut.sample(s, sourceValue.rgb).rgb, intensity); // write it to destination texture - destination.write(sourceColor, position); + destination.write(sourceValue, position); } // MARK: Texture affine crop @@ -874,56 +867,56 @@ kernel void textureAffineCrop(texture2d source [[ texture( texture2d destination [[ texture(1) ]], constant float3x3& transform [[ buffer(0) ]], ushort2 position [[thread_position_in_grid]]) { - const ushort2 textureSize = ushort2(destination.get_width(), + const auto textureSize = ushort2(destination.get_width(), destination.get_height()); checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); constexpr sampler s(coord::normalized, address::clamp_to_edge, filter::linear); + const auto positionF = float2(position); + const auto textureSizeF = float2(textureSize); + const auto normalizedPosition = (positionF + 0.5f) / textureSizeF; - const float2 textureSizef = float2(textureSize); - const float2 normalizedPosition = float2(position) / textureSizef; - - const float3 targetPosition = transform * float3(normalizedPosition, 1.0f); + const auto targetPosition = transform * float3(normalizedPosition, 1.0f); // read original color - half4 sourceColor = source.sample(s, targetPosition.xy); + const auto sourceValue = source.sample(s, targetPosition.xy); // write it to destination texture - destination.write(sourceColor, position); + destination.write(sourceValue, position); } // MARK: - Texture Interpolation template -void textureInterpolation(texture2d sourceTextureOne, - texture2d sourceTextureTwo, - texture2d destinationTexture, +void textureInterpolation(texture2d sourceOne, + texture2d sourceTwo, + texture2d destination, constant float& weight, const ushort2 position) { - const ushort2 textureSize = ushort2(destinationTexture.get_width(), - destinationTexture.get_height()); + const auto textureSize = ushort2(destination.get_width(), + destination.get_height()); checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); - const auto sourceValueOne = sourceTextureOne.read(position); - const auto sourceValueTwo = sourceTextureTwo.read(position); + const auto sourceValueOne = sourceOne.read(position); + const auto sourceValueTwo = sourceTwo.read(position); const auto resultValue = sourceValueOne + vec(float4(sourceValueTwo - sourceValueOne) * weight); - destinationTexture.write(resultValue, position); + destination.write(resultValue, position); } -#define outerArguments(T) \ -(texture2d sourceTextureOne [[ texture(0) ]], \ -texture2d sourceTextureTwo [[ texture(1) ]], \ -texture2d destinationTexture [[ texture(2) ]], \ -constant float& weight [[ buffer(0) ]], \ -const ushort2 position [[ thread_position_in_grid ]]) \ +#define outerArguments(T) \ +(texture2d sourceOne [[ texture(0) ]], \ +texture2d sourceTwo [[ texture(1) ]], \ +texture2d destination [[ texture(2) ]], \ +constant float& weight [[ buffer(0) ]], \ +const ushort2 position [[ thread_position_in_grid ]]) \ #define innerArguments \ -(sourceTextureOne, \ -sourceTextureTwo, \ -destinationTexture, \ +(sourceOne, \ +sourceTwo, \ +destination, \ weight, \ position) \ @@ -941,24 +934,79 @@ constant float4x4 ycbcrToRGBTransform = { { -0.7010f, +0.5291f, -0.8860f, +1.0000f } }; -kernel void ycbcrToRGBA(texture2d sourceYTexture [[ texture(0) ]], - texture2d sourceCbCrTexture [[ texture(1) ]], - texture2d destinationRGBATexture [[ texture(2) ]], +kernel void ycbcrToRGBA(texture2d sourceY [[ texture(0) ]], + texture2d sourceCbCr [[ texture(1) ]], + texture2d destinationRGBA [[ texture(2) ]], const ushort2 position [[ thread_position_in_grid ]], const ushort2 totalThreads [[ threads_per_grid ]]) { - const ushort2 textureSize = ushort2(destinationRGBATexture.get_width(), - destinationRGBATexture.get_height()); + const auto textureSize = ushort2(destinationRGBA.get_width(), + destinationRGBA.get_height()); checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); constexpr sampler s(coord::normalized, address::clamp_to_edge, filter::linear); - const auto normalizedPosition = float2(position) / float2(textureSize); + const auto positionF = float2(position); + const auto textureSizeF = float2(textureSize); + const auto normalizedPosition = (positionF + 0.5f) / textureSizeF; - const auto ycbcr = float4(sourceYTexture.sample(s, normalizedPosition).r, - sourceCbCrTexture.sample(s, normalizedPosition).rg, - 1.0); + const auto ycbcr = float4(sourceY.sample(s, normalizedPosition).r, + sourceCbCr.sample(s, normalizedPosition).rg, + 1.0f); const auto destinationValue = ycbcrToRGBTransform * ycbcr; - destinationRGBATexture.write(destinationValue, position); + destinationRGBA.write(destinationValue, position); +} + +// MARK: - RGBA to YCbCr + +constant float4x4 rgbaToYCbCrTransform = { + { +0.2990f, -0.1687f, +0.5000f, +0.0000f }, + { +0.5870f, -0.3313f, -0.4187f, +0.0000f }, + { +0.1140f, +0.5000f, -0.0813f, +0.0000f }, + { -0.0000f, +0.5000f, +0.5000f, +1.0000f } +}; + +kernel void rgbaToYCbCr(texture2d sourceRGBA [[ texture(0) ]], + texture2d destinationY [[ texture(1) ]], + texture2d destinationCbCr [[ texture(2) ]], + const ushort2 position [[ thread_position_in_grid ]]) { + const auto textureSize = ushort2(destinationCbCr.get_width(), + destinationCbCr.get_height()); + checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); + + if (halfSizedCbCr) { + constexpr sampler s(coord::pixel, + address::clamp_to_edge, + filter::nearest); + const auto positionF = float2(position); + const auto gatherPosition = positionF * 2.0f + 1.0f; + + const auto rValuesForQuad = sourceRGBA.gather(s, gatherPosition, + 0, component::x); + const auto gValuesForQuad = sourceRGBA.gather(s, gatherPosition, + 0, component::y); + const auto bValuesForQuad = sourceRGBA.gather(s, gatherPosition, + 0, component::z); + const auto rgbaValues = transpose(float4x4(rValuesForQuad, + gValuesForQuad, + bValuesForQuad, + float4(1.0f))); + const auto ycbcrValues = rgbaToYCbCrTransform * rgbaValues; + + destinationY.write(ycbcrValues[0].r, position * 2 + ushort2(0, 1)); + destinationY.write(ycbcrValues[1].r, position * 2 + ushort2(1, 1)); + destinationY.write(ycbcrValues[2].r, position * 2 + ushort2(1, 0)); + destinationY.write(ycbcrValues[3].r, position * 2 + ushort2(0, 0)); + + const auto cbcrValue = (ycbcrValues * float4(0.25f)).gb; + destinationCbCr.write(float4(cbcrValue, 0.0f), position); + } else { + const auto rgbaValue = sourceRGBA.read(position); + const auto ycbcrValue = rgbaToYCbCrTransform * rgbaValue; + const auto yValue = ycbcrValue.r; + const auto cbcrValue = float4(ycbcrValue.gb, 0.0f); + destinationY.write(yValue, position); + destinationCbCr.write(cbcrValue, position); + } } diff --git a/Demo/AlloyTests/AlloyTests.swift b/Demo/AlloyTests/AlloyTests.swift index dfe58a7..e9adec3 100644 --- a/Demo/AlloyTests/AlloyTests.swift +++ b/Demo/AlloyTests/AlloyTests.swift @@ -63,11 +63,11 @@ class AlloyTests: XCTestCase { func testEvenPerformance() { self.measure { - self.runGPUWork { (encoder, texture, outputTexture) in + self.runGPUWork { (encoder, texture, output) in encoder.setTexture(texture, index: 0) encoder.dispatch2d(state: self.evenInitState, covering: texture.size) - encoder.setTexture(outputTexture, index: 1) + encoder.setTexture(output, index: 1) encoder.dispatch2d(state: self.evenProcessState, covering: texture.size) } } @@ -75,11 +75,11 @@ class AlloyTests: XCTestCase { func testEvenOptimizedPerformance() { self.measure { - self.runGPUWork { (encoder, texture, outputTexture) in + self.runGPUWork { (encoder, texture, output) in encoder.setTexture(texture, index: 0) encoder.dispatch2d(state: self.evenOptimizedInitState, covering: texture.size) - encoder.setTexture(outputTexture, index: 1) + encoder.setTexture(output, index: 1) encoder.dispatch2d(state: self.evenOptimizedProcessState, covering: texture.size) } } @@ -87,11 +87,11 @@ class AlloyTests: XCTestCase { func testExactPerformance() { self.measure { - self.runGPUWork { (encoder, texture, outputTexture) in + self.runGPUWork { (encoder, texture, output) in encoder.setTexture(texture, index: 0) encoder.dispatch2d(state: self.exactInitState, exactly: texture.size) - encoder.setTexture(outputTexture, index: 1) + encoder.setTexture(output, index: 1) encoder.dispatch2d(state: self.exactProcessState, exactly: texture.size) } } @@ -106,22 +106,20 @@ class AlloyTests: XCTestCase { for wd in 0.. 1 { + if original.mipmapLevelCount > 1 { var level: Int = 0 - var width = originalTexture.width - var height = originalTexture.height + var width = original.width + var height = original.height while (width + height > 32) { - guard let originalTextureView = originalTexture.view(level: level), - let decodedTextureView = decodedTexture.view(level: level) + guard let originalTextureView = original.view(level: level), + let decodedTextureView = decoded.view(level: level) else { fatalError("Couldn't create texture view at level \(level)") } euclideanDistance(textureOne: originalTextureView, @@ -136,8 +136,8 @@ final class TextureCachingTests: XCTestCase { } } else { - euclideanDistance(textureOne: originalTexture, - textureTwo: decodedTexture, + euclideanDistance(textureOne: original, + textureTwo: decoded, resultBuffer: resultBuffer, in: commadBuffer) } @@ -154,20 +154,18 @@ final class TextureCachingTests: XCTestCase { generateMipmaps: Bool, in commandBuffer: MTLCommandBuffer) throws -> MTLTexture { var generateMipmaps = generateMipmaps - let resultTexture: MTLTexture + let result: MTLTexture switch pixelFormat.dataFormat { case .normalized: - resultTexture = try self.context - .texture(from: cgImage, - usage: [.shaderRead, .shaderWrite], - generateMipmaps: generateMipmaps) + result = try self.context.texture(from: cgImage, + usage: [.shaderRead, .shaderWrite], + generateMipmaps: generateMipmaps) case .unsignedInteger: generateMipmaps = false - let normalizedTexture = try self.context - .texture(from: cgImage, - usage: [.shaderRead, .shaderWrite]) + let normalized = try self.context.texture(from: cgImage, + usage: [.shaderRead, .shaderWrite]) let unnormalizedTextureDescriptor = MTLTextureDescriptor() unnormalizedTextureDescriptor.width = cgImage.width @@ -176,24 +174,23 @@ final class TextureCachingTests: XCTestCase { unnormalizedTextureDescriptor.usage = [.shaderRead, .shaderWrite] unnormalizedTextureDescriptor.storageMode = .shared - let unnormalizedTexture = try self.context - .texture(descriptor: unnormalizedTextureDescriptor) + let unnormalized = try self.context.texture(descriptor: unnormalizedTextureDescriptor) - self.denormalize(normalizedTexture: normalizedTexture, - unnormalizedTexture: unnormalizedTexture, + self.denormalize(normalized: normalized, + unnormalized: unnormalized, in: commandBuffer) - resultTexture = unnormalizedTexture + result = unnormalized default: throw Error.unsutablePixelFormat } if generateMipmaps { commandBuffer.blit { encoder in - encoder.generateMipmaps(for: resultTexture) + encoder.generateMipmaps(for: result) } } - return resultTexture + return result } } diff --git a/Demo/AlloyTests/TextureCopyTests.swift b/Demo/AlloyTests/TextureCopyTests.swift index 8c79d65..7b5c838 100644 --- a/Demo/AlloyTests/TextureCopyTests.swift +++ b/Demo/AlloyTests/TextureCopyTests.swift @@ -12,9 +12,9 @@ final class TextureCopyTests: XCTestCase { struct TestCase { let sourceRegion: MTLRegion let destinationOrigin: MTLOrigin - let sourceTexture: MTLTexture - let destinationTexture: MTLTexture - let desiredResultTexture: MTLTexture + let source: MTLTexture + let destination: MTLTexture + let desiredResult: MTLTexture init(context: MTLContext, sourceRegion: MTLRegion, @@ -24,12 +24,12 @@ final class TextureCopyTests: XCTestCase { desiredResultImage: CGImage) throws { self.sourceRegion = sourceRegion self.destinationOrigin = destinationOrigin - try self.sourceTexture = context.texture(from: sourceImage, + try self.source = context.texture(from: sourceImage, + usage: [.shaderRead, .shaderWrite]) + try self.destination = context.texture(from: destinationImage, + usage: [.shaderRead, .shaderWrite]) + try self.desiredResult = context.texture(from: desiredResultImage, usage: [.shaderRead, .shaderWrite]) - try self.destinationTexture = context.texture(from: destinationImage, - usage: [.shaderRead, .shaderWrite]) - try self.desiredResultTexture = context.texture(from: desiredResultImage, - usage: [.shaderRead, .shaderWrite]) } } @@ -104,13 +104,13 @@ final class TextureCopyTests: XCTestCase { try self.testCases.forEach { testCase in try self.context.scheduleAndWait { commandBuffer in self.textureCopy(region: testCase.sourceRegion, - from: testCase.sourceTexture, + from: testCase.source, to: testCase.destinationOrigin, - of: testCase.destinationTexture, + of: testCase.destination, in: commandBuffer) - self.euclideanDistanceFloat(textureOne: testCase.destinationTexture, - textureTwo: testCase.desiredResultTexture, + self.euclideanDistanceFloat(textureOne: testCase.destination, + textureTwo: testCase.desiredResult, resultBuffer: resultBuffer, in: commandBuffer) } diff --git a/Demo/Demo.xcodeproj/project.pbxproj b/Demo/Demo.xcodeproj/project.pbxproj index 66dbb83..b2ec996 100644 --- a/Demo/Demo.xcodeproj/project.pbxproj +++ b/Demo/Demo.xcodeproj/project.pbxproj @@ -838,7 +838,7 @@ ); PRODUCT_BUNDLE_IDENTIFIER = me.avolodin.Demo; PRODUCT_NAME = "$(TARGET_NAME)"; - SWIFT_VERSION = 4.2; + SWIFT_VERSION = 5.0; }; name = Debug; }; @@ -858,7 +858,7 @@ ); PRODUCT_BUNDLE_IDENTIFIER = me.avolodin.Demo; PRODUCT_NAME = "$(TARGET_NAME)"; - SWIFT_VERSION = 4.2; + SWIFT_VERSION = 5.0; }; name = Release; }; @@ -883,7 +883,7 @@ ); PRODUCT_BUNDLE_IDENTIFIER = me.avolodin.AlloyTests; PRODUCT_NAME = "$(TARGET_NAME)"; - SWIFT_VERSION = 4.2; + SWIFT_VERSION = 5.0; TEST_HOST = "$(BUILT_PRODUCTS_DIR)/Demo.app/Contents/MacOS/Demo"; }; name = Debug; @@ -909,7 +909,7 @@ ); PRODUCT_BUNDLE_IDENTIFIER = me.avolodin.AlloyTests; PRODUCT_NAME = "$(TARGET_NAME)"; - SWIFT_VERSION = 4.2; + SWIFT_VERSION = 5.0; TEST_HOST = "$(BUILT_PRODUCTS_DIR)/Demo.app/Contents/MacOS/Demo"; }; name = Release; @@ -931,7 +931,7 @@ PRODUCT_BUNDLE_IDENTIFIER = me.avolodin.DemoApp; PRODUCT_NAME = "$(TARGET_NAME)"; SDKROOT = iphoneos; - SWIFT_VERSION = 4.2; + SWIFT_VERSION = 5.0; TARGETED_DEVICE_FAMILY = "1,2"; }; name = Debug; @@ -953,7 +953,7 @@ PRODUCT_BUNDLE_IDENTIFIER = me.avolodin.DemoApp; PRODUCT_NAME = "$(TARGET_NAME)"; SDKROOT = iphoneos; - SWIFT_VERSION = 4.2; + SWIFT_VERSION = 5.0; TARGETED_DEVICE_FAMILY = "1,2"; VALIDATE_PRODUCT = YES; }; @@ -983,7 +983,7 @@ PRODUCT_BUNDLE_IDENTIFIER = me.avolodin.DemoAppTests; PRODUCT_NAME = "$(TARGET_NAME)"; SDKROOT = iphoneos; - SWIFT_VERSION = 4.2; + SWIFT_VERSION = 5.0; TARGETED_DEVICE_FAMILY = "1,2"; TEST_HOST = "$(BUILT_PRODUCTS_DIR)/DemoApp.app/DemoApp"; }; @@ -1013,7 +1013,7 @@ PRODUCT_BUNDLE_IDENTIFIER = me.avolodin.DemoAppTests; PRODUCT_NAME = "$(TARGET_NAME)"; SDKROOT = iphoneos; - SWIFT_VERSION = 4.2; + SWIFT_VERSION = 5.0; TARGETED_DEVICE_FAMILY = "1,2"; TEST_HOST = "$(BUILT_PRODUCTS_DIR)/DemoApp.app/DemoApp"; VALIDATE_PRODUCT = YES; diff --git a/Demo/Demo/ViewController.swift b/Demo/Demo/ViewController.swift index 9c23c92..a2d5bfa 100644 --- a/Demo/Demo/ViewController.swift +++ b/Demo/Demo/ViewController.swift @@ -16,7 +16,7 @@ class ViewController: NSViewController { @IBOutlet weak var imageView: NSImageView! @IBOutlet weak var slider: NSSlider! - let context = try! MTLContext(device: Metal.lowPowerDevice!) + let context = try! MTLContext() var affineCropEncoder: TextureAffineCrop! let image = NSImage(named: "flower")! @@ -55,8 +55,8 @@ class ViewController: NSViewController { try? self.context.scheduleAndWait { buffer in buffer.compute { encoder in - self.affineCropEncoder.encode(sourceTexture: texture, - destinationTexture: cropTexture, + self.affineCropEncoder.encode(source: texture, + destination: cropTexture, affineTransform: simd_float3x3(transform), using: encoder) } diff --git a/Demo/Podfile.lock b/Demo/Podfile.lock index 2f9fbb2..eab2b81 100644 --- a/Demo/Podfile.lock +++ b/Demo/Podfile.lock @@ -1,10 +1,10 @@ PODS: - - Alloy (0.14.1): - - Alloy/Core (= 0.14.1) - - Alloy/Core (0.14.1) - - Alloy/ML (0.14.1): + - Alloy (0.15.0): + - Alloy/Core (= 0.15.0) + - Alloy/Core (0.15.0) + - Alloy/ML (0.15.0): - Alloy/Core - - Alloy/Shaders (0.14.1): + - Alloy/Shaders (0.15.0): - Alloy/Core - SwiftMath (3.2.0) @@ -23,9 +23,9 @@ EXTERNAL SOURCES: :path: "../" SPEC CHECKSUMS: - Alloy: 699cc1a530ca68527ff2edea90584a68e72bd745 + Alloy: 84e03ebb5735cd42203396a4d28f019079369f28 SwiftMath: d010d63f42454f782680c28df987323b08a85caa PODFILE CHECKSUM: 0e625e0798aefa71c5baa3004987245f7e468dd7 -COCOAPODS: 1.9.2 +COCOAPODS: 1.9.3