From 35de73aaee99a4f7fd5e2d496310df9117b0ab15 Mon Sep 17 00:00:00 2001 From: Eugene Bokhan Date: Tue, 27 Oct 2020 12:47:37 +0300 Subject: [PATCH 1/5] update for xcode 12 --- .../Metal/MTLPixelFormat+Extensions.swift | 75 +++++++++++++------ .../Metal/MTLTexture+Extensions.swift | 26 +++---- Sources/Alloy/Core/Float16.swift | 67 +++++++++-------- 3 files changed, 98 insertions(+), 70 deletions(-) diff --git a/Sources/Alloy/Core/Extensions/Metal/MTLPixelFormat+Extensions.swift b/Sources/Alloy/Core/Extensions/Metal/MTLPixelFormat+Extensions.swift index 6ae012b..58bdd4d 100644 --- a/Sources/Alloy/Core/Extensions/Metal/MTLPixelFormat+Extensions.swift +++ b/Sources/Alloy/Core/Extensions/Metal/MTLPixelFormat+Extensions.swift @@ -6,6 +6,7 @@ public extension MTLPixelFormat { } var size: Int? { + #if os(iOS) && !targetEnvironment(macCatalyst) switch self { case .a8Unorm, .r8Unorm, .r8Snorm, .r8Uint, .r8Sint, .stencil8, .r8Unorm_srgb: return 1 @@ -13,6 +14,29 @@ public extension MTLPixelFormat { .r16Sint, .r16Float, .rg8Unorm, .rg8Snorm, .rg8Uint, .rg8Sint, .depth16Unorm, .rg8Unorm_srgb: return 2 + case .r32Uint, .r32Sint, .r32Float, + .rg16Unorm, .rg16Snorm, .rg16Uint, + .rg16Sint, .rg16Float, .rgba8Unorm, + .rgba8Unorm_srgb, .rgba8Snorm, .rgba8Uint, + .rgba8Sint, .bgra8Unorm, .bgra8Unorm_srgb, + .rgb10a2Unorm, .rgb10a2Uint, .rg11b10Float, + .rgb9e5Float, .bgr10a2Unorm, .gbgr422, + .bgrg422, .depth32Float, .bgr10_xr_srgb, .bgr10_xr: return 4 + case .rg32Uint, .rg32Sint, .rg32Float, + .rgba16Unorm, .rgba16Snorm, .rgba16Uint, + .rgba16Sint, .rgba16Float, .depth32Float_stencil8, .x32_stencil8, + .bgra10_xr, .bgra10_xr_srgb: return 8 + case .rgba32Uint, .rgba32Sint, .rgba32Float: return 16 + default: return nil + } + #elseif os(macOS) || (os(iOS) && targetEnvironment(macCatalyst)) + switch self { + case .a8Unorm, .r8Unorm, .r8Snorm, + .r8Uint, .r8Sint, .stencil8: return 1 + case .r16Unorm, .r16Snorm, .r16Uint, + .r16Sint, .r16Float, .rg8Unorm, + .rg8Snorm, .rg8Uint, .rg8Sint, + .depth16Unorm: return 2 case .r32Uint, .r32Sint, .r32Float, .rg16Unorm, .rg16Snorm, .rg16Uint, .rg16Sint, .rg16Float, .rgba8Unorm, @@ -21,19 +45,17 @@ public extension MTLPixelFormat { .rgb10a2Unorm, .rgb10a2Uint, .rg11b10Float, .rgb9e5Float, .bgr10a2Unorm, .gbgr422, .bgrg422, .depth32Float, .depth24Unorm_stencil8, - .x24_stencil8, .bgr10_xr_srgb, .bgr10_xr: return 4 + .x24_stencil8: return 4 case .rg32Uint, .rg32Sint, .rg32Float, .rgba16Unorm, .rgba16Snorm, .rgba16Uint, .rgba16Sint, .rgba16Float, .bc1_rgba, - .bc1_rgba_srgb, .depth32Float_stencil8, .x32_stencil8, - .bgra10_xr, .bgra10_xr_srgb: return 8 + .bc1_rgba_srgb, .depth32Float_stencil8, .x32_stencil8: return 8 case .rgba32Uint, .rgba32Sint, .rgba32Float, .bc2_rgba, .bc2_rgba_srgb, .bc3_rgba, .bc3_rgba_srgb: return 16 - default: - // TODO: Finish bc4-bc7 - return nil + default: return nil } + #endif } var isOrdinary8Bit: Bool { @@ -175,11 +197,18 @@ public extension MTLPixelFormat { var isCompressed: Bool { #if os(iOS) && !targetEnvironment(macCatalyst) - return self.isPVRTC - || self.isEAC - || self.isETC - || self.isASTC - || self.isHDRASTC + if #available(iOS 13.0, *) { + return self.isPVRTC + || self.isEAC + || self.isETC + || self.isASTC + || self.isHDRASTC + } else { + return self.isPVRTC + || self.isEAC + || self.isETC + || self.isASTC + } #elseif os(macOS) || (os(iOS) && targetEnvironment(macCatalyst)) return self.isS3TC || self.isRGTC @@ -210,15 +239,13 @@ public extension MTLPixelFormat { } } + @available(iOS 13.0, *) var isHDRASTC: Bool { switch self { - case .astc_4x4_hdr, - .astc_5x4_hdr, .astc_5x5_hdr, - .astc_6x5_hdr, .astc_6x6_hdr, - .astc_8x5_hdr, .astc_8x6_hdr, .astc_8x8_hdr, - .astc_10x5_hdr, .astc_10x6_hdr, .astc_10x8_hdr, .astc_10x10_hdr, + case .astc_4x4_hdr, .astc_5x4_hdr, .astc_5x5_hdr, .astc_6x5_hdr, .astc_6x6_hdr, .astc_8x5_hdr, + .astc_8x6_hdr, .astc_8x8_hdr, .astc_10x5_hdr, .astc_10x6_hdr, .astc_10x8_hdr, .astc_10x10_hdr, .astc_12x10_hdr, .astc_12x12_hdr: - return true + return true default: return false } } @@ -282,10 +309,16 @@ public extension MTLPixelFormat { } var isDepth: Bool { - switch self { - case .depth16Unorm, .depth32Float: - return true - default: return false + if #available(iOS 13.0, *) { + switch self { + case .depth16Unorm, .depth32Float: return true + default: return false + } + } else { + switch self { + case .depth32Float: return true + default: return false + } } } diff --git a/Sources/Alloy/Core/Extensions/Metal/MTLTexture+Extensions.swift b/Sources/Alloy/Core/Extensions/Metal/MTLTexture+Extensions.swift index edd56ef..af95d64 100644 --- a/Sources/Alloy/Core/Extensions/Metal/MTLTexture+Extensions.swift +++ b/Sources/Alloy/Core/Extensions/Metal/MTLTexture+Extensions.swift @@ -20,8 +20,7 @@ public extension MTLTexture { let rowBytes = self.width let length = rowBytes * self.height - let rgbaBytes = UnsafeMutableRawPointer.allocate(byteCount: length, - alignment: MemoryLayout.alignment) + let rgbaBytes = UnsafeMutablePointer.allocate(capacity: length) defer { rgbaBytes.deallocate() } self.getBytes(rgbaBytes, bytesPerRow: rowBytes, @@ -33,7 +32,7 @@ public extension MTLTexture { ? CGImageAlphaInfo.alphaOnly.rawValue : CGImageAlphaInfo.none.rawValue) guard let data = CFDataCreate(nil, - rgbaBytes.assumingMemoryBound(to: UInt8.self), + rgbaBytes, length), let dataProvider = CGDataProvider(data: data), let cgImage = CGImage(width: self.width, @@ -54,25 +53,20 @@ public extension MTLTexture { // read texture as byte array let rowBytes = self.width * 4 let length = rowBytes * self.height - - let bgraBytes = UnsafeMutableRawPointer.allocate(byteCount: length, - alignment: MemoryLayout.alignment) - defer { bgraBytes.deallocate() } - + let bgraBytes = UnsafeMutablePointer.allocate(capacity: length) + let rgbaBytes = UnsafeMutablePointer.allocate(capacity: length) + defer { bgraBytes.deallocate(); rgbaBytes.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: bgraBytes, height: vImagePixelCount(self.height), width: vImagePixelCount(self.width), rowBytes: rowBytes) - - 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), @@ -86,7 +80,7 @@ public extension MTLTexture { let colorScape = colorSpace ?? CGColorSpaceCreateDeviceRGB() let bitmapInfo = CGBitmapInfo(rawValue: CGImageAlphaInfo.premultipliedLast.rawValue) guard let data = CFDataCreate(nil, - rgbaBytes.assumingMemoryBound(to: UInt8.self), + rgbaBytes, length), let dataProvider = CGDataProvider(data: data), let cgImage = CGImage(width: self.width, @@ -107,10 +101,8 @@ public extension MTLTexture { let rowBytes = self.width * 4 let length = rowBytes * self.height - let rgbaBytes = UnsafeMutableRawPointer.allocate(byteCount: length, - alignment: MemoryLayout.alignment) + let rgbaBytes = UnsafeMutablePointer.allocate(capacity: length) defer { rgbaBytes.deallocate() } - self.getBytes(rgbaBytes, bytesPerRow: rowBytes, from: self.region, @@ -119,7 +111,7 @@ public extension MTLTexture { let colorScape = colorSpace ?? CGColorSpaceCreateDeviceRGB() let bitmapInfo = CGBitmapInfo(rawValue: CGImageAlphaInfo.premultipliedLast.rawValue) guard let data = CFDataCreate(nil, - rgbaBytes.assumingMemoryBound(to: UInt8.self), + rgbaBytes, length), let dataProvider = CGDataProvider(data: data), let cgImage = CGImage(width: self.width, diff --git a/Sources/Alloy/Core/Float16.swift b/Sources/Alloy/Core/Float16.swift index 263b514..3f72858 100644 --- a/Sources/Alloy/Core/Float16.swift +++ b/Sources/Alloy/Core/Float16.swift @@ -14,24 +14,29 @@ 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) - 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) +public func float16to32(_ input: UnsafeMutableRawPointer, count: Int) -> [Float]? { + let output = [Float](repeating: 0, count: count) + var data: UnsafeMutableRawPointer? = nil + output.withUnsafeBytes { + data = UnsafeMutableRawPointer(mutating: $0.baseAddress) } - return status == kvImageNoError ? output : nil + + var bufferFloat16 = vImage_Buffer(data: input, + height: 1, + width: UInt(count), + rowBytes: count * 2) + var bufferFloat32 = vImage_Buffer(data: data, + height: 1, + width: UInt(count), + rowBytes: count * 4) + + if vImageConvert_Planar16FtoPlanarF(&bufferFloat16, + &bufferFloat32, + 0) != kvImageNoError { + return nil + } + + return output } /// Uses vImage to convert an array of Swift floats into a buffer of float16s. @@ -40,20 +45,18 @@ public func float16to32(_ input: UnsafeMutableRawPointer, /// - 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) - 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) +public func float32to16(_ input: UnsafeMutablePointer, count: Int) -> [Float16]? { + let output = [Float16](repeating: 0, count: count) + var data: UnsafeMutableRawPointer? = nil + output.withUnsafeBytes { + data = UnsafeMutableRawPointer(mutating: $0.baseAddress) + } + + var bufferFloat32 = vImage_Buffer(data: input, height: 1, width: UInt(count), rowBytes: count * 4) + var bufferFloat16 = vImage_Buffer(data: data, height: 1, width: UInt(count), rowBytes: count * 2) + + if vImageConvert_PlanarFtoPlanar16F(&bufferFloat32, &bufferFloat16, 0) != kvImageNoError { + return nil } - return status == kvImageNoError ? output : nil + return output } From f10643bc8111ba7173ab5b0caddb17574dbcb929 Mon Sep 17 00:00:00 2001 From: Eugene Bokhan Date: Tue, 27 Oct 2020 13:11:30 +0300 Subject: [PATCH 2/5] fix deprecations --- Sources/Alloy/ML/ONNXConvolutionPadding.swift | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/Sources/Alloy/ML/ONNXConvolutionPadding.swift b/Sources/Alloy/ML/ONNXConvolutionPadding.swift index 7128330..66e4514 100644 --- a/Sources/Alloy/ML/ONNXConvolutionPadding.swift +++ b/Sources/Alloy/ML/ONNXConvolutionPadding.swift @@ -30,9 +30,9 @@ public typealias Scales = (height: Int, width: Int) } required convenience public init?(coder aDecoder: NSCoder) { - guard - let data = aDecoder.decodeData(), - let other = NSKeyedUnarchiver.unarchiveObject(with: data) as? ONNXConvolutionPadding + guard let data = aDecoder.decodeData(), + let other = try? NSKeyedUnarchiver.unarchivedObject(ofClass: ONNXConvolutionPadding.self, + from: data) else { return nil } self.init(kernel: other.kernel, strides: other.strides, @@ -43,7 +43,8 @@ public typealias Scales = (height: Int, width: Int) } public func encode(with aCoder: NSCoder) { - aCoder.encode(NSKeyedArchiver.archivedData(withRootObject: self)) + try? aCoder.encode(NSKeyedArchiver.archivedData(withRootObject: self, + requiringSecureCoding: false)) } public func paddingMethod() -> MPSNNPaddingMethod { From eb5218f6924b23dfa8a55dd4935147ef0a4e89fa Mon Sep 17 00:00:00 2001 From: eugene Date: Thu, 4 Mar 2021 18:12:22 +0300 Subject: [PATCH 3/5] actualize with master --- .gitignore | 1 + Alloy.podspec | 2 +- Package.swift | 6 +- .../Extensions/Foundation/Bundle+Module.swift | 2 +- .../Extensions/Metal/MTLDevice+Features.swift | 19 ++++ .../Metal/MTLTexture+Serialization.swift | 8 +- Sources/Alloy/Encoders/Normalization.swift | 86 +++++++++++++++++++ .../Renderers/BoundingBoxesRenderer.swift | 2 +- Sources/Alloy/Renderers/LinesRenderer.swift | 2 +- Sources/Alloy/Renderers/MaskRenderer.swift | 2 +- .../Alloy/Renderers/RectangleRenderer.swift | 2 +- Sources/Alloy/Shaders/Shaders.metal | 16 ++-- 12 files changed, 128 insertions(+), 20 deletions(-) create mode 100644 Sources/Alloy/Encoders/Normalization.swift diff --git a/.gitignore b/.gitignore index 6c8e705..b82bda9 100644 --- a/.gitignore +++ b/.gitignore @@ -6,3 +6,4 @@ Demo/Demo.xcodeproj/xcuserdata/ Demo/Demo.xcworkspace/xcuserdata/ *gen .swiftpm +.build diff --git a/Alloy.podspec b/Alloy.podspec index 0ddb7cc..1d5a21c 100644 --- a/Alloy.podspec +++ b/Alloy.podspec @@ -1,6 +1,6 @@ Pod::Spec.new do |s| s.name = 'Alloy' - s.version = '0.16.3' + s.version = '0.16.5' s.summary = 'Nano helpers for Metal framework' s.homepage = 'https://github.com/s1ddok/Alloy' diff --git a/Package.swift b/Package.swift index 7227e60..8147c54 100644 --- a/Package.swift +++ b/Package.swift @@ -17,8 +17,7 @@ let package = Package( publicHeadersPath: "."), .target(name: "Alloy", dependencies: [.target(name: "AlloyShadersSharedTypes")], - resources: [.process("Shaders/Shaders.metal")], - swiftSettings: [.define("SWIFT_PM")]), + resources: [.process("Shaders/Shaders.metal")]), .target(name: "AlloyTestsResources", path: "Tests/AlloyTestsResources", resources: [ @@ -27,7 +26,6 @@ let package = Package( ]), .testTarget(name: "AlloyTests", dependencies: ["Alloy", "AlloyTestsResources"], - resources: [.process("Shaders/Shaders.metal")], - swiftSettings: [.define("SWIFT_PM")]) + resources: [.process("Shaders/Shaders.metal")]) ] ) diff --git a/Sources/Alloy/Core/Extensions/Foundation/Bundle+Module.swift b/Sources/Alloy/Core/Extensions/Foundation/Bundle+Module.swift index 95bf9bf..cd155e5 100644 --- a/Sources/Alloy/Core/Extensions/Foundation/Bundle+Module.swift +++ b/Sources/Alloy/Core/Extensions/Foundation/Bundle+Module.swift @@ -1,5 +1,5 @@ import Foundation -#if !SWIFT_PM +#if !SWIFT_PACKAGE public extension Bundle { static var module = Bundle(for: MTLContext.self) } diff --git a/Sources/Alloy/Core/Extensions/Metal/MTLDevice+Features.swift b/Sources/Alloy/Core/Extensions/Metal/MTLDevice+Features.swift index aee4448..5f09708 100644 --- a/Sources/Alloy/Core/Extensions/Metal/MTLDevice+Features.swift +++ b/Sources/Alloy/Core/Extensions/Metal/MTLDevice+Features.swift @@ -2,6 +2,7 @@ import Metal public enum Feature { case nonUniformThreadgroups + case readWriteTextures(MTLPixelFormat) } public extension MTLDevice { @@ -15,6 +16,24 @@ public extension MTLDevice { #elseif os(macOS) return self.supportsFeatureSet(.macOS_GPUFamily1_v3) #endif + + case let .readWriteTextures(pixelFormat): + let tierOneSupportedPixelFormats: Set = [ + .r32Float, .r32Uint, .r32Sint + ] + let tierTwoSupportedPixelFormats: Set = tierOneSupportedPixelFormats.union([ + .rgba32Float, .rgba32Uint, .rgba32Sint, .rgba16Float, + .rgba16Uint, .rgba16Sint, .rgba8Unorm, .rgba8Uint, + .rgba8Sint, .r16Float, .r16Uint, .r16Sint, + .r8Unorm, .r8Uint, .r8Sint + ]) + + switch self.readWriteTextureSupport { + case .tier1: return tierOneSupportedPixelFormats.contains(pixelFormat) + case .tier2: return tierTwoSupportedPixelFormats.contains(pixelFormat) + case .tierNone: return false + @unknown default: return false + } } } } diff --git a/Sources/Alloy/Core/Extensions/Metal/MTLTexture+Serialization.swift b/Sources/Alloy/Core/Extensions/Metal/MTLTexture+Serialization.swift index 3df7d0f..705f8ea 100644 --- a/Sources/Alloy/Core/Extensions/Metal/MTLTexture+Serialization.swift +++ b/Sources/Alloy/Core/Extensions/Metal/MTLTexture+Serialization.swift @@ -38,8 +38,10 @@ public class MTLTextureCodableBox: Codable { textureView.getBytes(pointer.advanced(by: offset), bytesPerRow: bytesPerRow, + bytesPerImage: bytesPerImage, from: textureView.region, - mipmapLevel: 0) + mipmapLevel: 0, + slice: 0) offset += bytesPerImage } @@ -80,8 +82,10 @@ public class MTLTextureCodableBox: Codable { textureView.replace(region: textureView.region, mipmapLevel: 0, + slice: 0, withBytes: pointer.advanced(by: offset), - bytesPerRow: bytesPerRow) + bytesPerRow: bytesPerRow, + bytesPerImage: bytesPerImage) offset += bytesPerImage } diff --git a/Sources/Alloy/Encoders/Normalization.swift b/Sources/Alloy/Encoders/Normalization.swift new file mode 100644 index 0000000..6ccd330 --- /dev/null +++ b/Sources/Alloy/Encoders/Normalization.swift @@ -0,0 +1,86 @@ +import Metal + +final public class Normalization { + + // MARK: - Propertires + + private let pipelineState: MTLComputePipelineState + private let deviceSupportsNonuniformThreadgroups: Bool + + // MARK: - Life Cycle + + public convenience init(context: MTLContext) throws { + try self.init(library: context.library(for: .module)) + } + + public init(library: MTLLibrary) throws { + self.deviceSupportsNonuniformThreadgroups = library.device + .supports(feature: .nonUniformThreadgroups) + let constantValues = MTLFunctionConstantValues() + constantValues.set(self.deviceSupportsNonuniformThreadgroups, + at: 0) + self.pipelineState = try library.computePipelineState(function: Self.functionName, + constants: constantValues) + } + + // MARK: - Encode + + public func callAsFunction(source: MTLTexture, + mean: SIMD3, + std: SIMD3, + destination: MTLTexture, + in commandBuffer: MTLCommandBuffer) { + self.encode(source: source, + mean: mean, + std: std, + destination: destination, + in: commandBuffer) + } + + public func callAsFunction(source: MTLTexture, + mean: SIMD3, + std: SIMD3, + destination: MTLTexture, + using encoder: MTLComputeCommandEncoder) { + self.encode(source: source, + mean: mean, + std: std, + destination: destination, + using: encoder) + } + + + public func encode(source: MTLTexture, + mean: SIMD3, + std: SIMD3, + destination: MTLTexture, + in commandBuffer: MTLCommandBuffer) { + commandBuffer.compute { encoder in + self.encode(source: source, + mean: mean, + std: std, + destination: destination, + using: encoder) + } + } + + public func encode(source: MTLTexture, + mean: SIMD3, + std: SIMD3, + destination: MTLTexture, + using encoder: MTLComputeCommandEncoder) { + encoder.setTextures(source, destination) + encoder.setValue(mean, at: 0) + encoder.setValue(std, at: 1) + + if self.deviceSupportsNonuniformThreadgroups { + encoder.dispatch2d(state: self.pipelineState, + exactly: source.size) + } else { + encoder.dispatch2d(state: self.pipelineState, + covering: source.size) + } + } + + public static let functionName = "normalization" +} diff --git a/Sources/Alloy/Renderers/BoundingBoxesRenderer.swift b/Sources/Alloy/Renderers/BoundingBoxesRenderer.swift index 6a0949f..43d1746 100644 --- a/Sources/Alloy/Renderers/BoundingBoxesRenderer.swift +++ b/Sources/Alloy/Renderers/BoundingBoxesRenderer.swift @@ -1,6 +1,6 @@ import Metal import CoreGraphics -#if SWIFT_PM +#if SWIFT_PACKAGE import AlloyShadersSharedTypes #endif diff --git a/Sources/Alloy/Renderers/LinesRenderer.swift b/Sources/Alloy/Renderers/LinesRenderer.swift index 1665bda..18af8e2 100644 --- a/Sources/Alloy/Renderers/LinesRenderer.swift +++ b/Sources/Alloy/Renderers/LinesRenderer.swift @@ -1,5 +1,5 @@ import Metal -#if SWIFT_PM +#if SWIFT_PACKAGE import AlloyShadersSharedTypes #endif diff --git a/Sources/Alloy/Renderers/MaskRenderer.swift b/Sources/Alloy/Renderers/MaskRenderer.swift index 9e935b4..6e4eb78 100644 --- a/Sources/Alloy/Renderers/MaskRenderer.swift +++ b/Sources/Alloy/Renderers/MaskRenderer.swift @@ -1,7 +1,7 @@ import Metal import simd import CoreGraphics -#if SWIFT_PM +#if SWIFT_PACKAGE import AlloyShadersSharedTypes #endif diff --git a/Sources/Alloy/Renderers/RectangleRenderer.swift b/Sources/Alloy/Renderers/RectangleRenderer.swift index 3224773..2010161 100644 --- a/Sources/Alloy/Renderers/RectangleRenderer.swift +++ b/Sources/Alloy/Renderers/RectangleRenderer.swift @@ -1,6 +1,6 @@ import Metal import CoreGraphics -#if SWIFT_PM +#if SWIFT_PACKAGE import AlloyShadersSharedTypes #endif diff --git a/Sources/Alloy/Shaders/Shaders.metal b/Sources/Alloy/Shaders/Shaders.metal index 581a1c5..35330b5 100644 --- a/Sources/Alloy/Shaders/Shaders.metal +++ b/Sources/Alloy/Shaders/Shaders.metal @@ -94,7 +94,7 @@ void textureMask(texture2d source, filter::linear); const auto positionF = float2(position); const auto textureSizeF = float2(textureSize); - const auto normalizedPosition = (positionF.x + 0.5f) / textureSizeF; + const auto normalizedPosition = (positionF + 0.5f) / textureSizeF; auto maskValue = mask.sample(s, normalizedPosition); if (isInversed) { @@ -656,15 +656,15 @@ kernel void textureDifferenceHighlight(texture2d sourceOne // MARK: - ML -kernel void normalize(texture2d inputTexture [[ texture(0) ]], - texture2d outputTexture [[ texture(1) ]], - constant float3& mean [[ buffer(0) ]], - constant float3& std [[ buffer(1) ]], - uint2 position [[thread_position_in_grid]]) { +kernel void normalization(texture2d inputTexture [[ texture(0) ]], + texture2d outputTexture [[ texture(1) ]], + constant float3& mean [[ buffer(0) ]], + constant float3& std [[ buffer(1) ]], + uint2 position [[thread_position_in_grid]]) { const auto textureSize = ushort2(inputTexture.get_width(), - inputTexture.get_height()); + inputTexture.get_height()); checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); - + // Read mpsnngraph result value. const auto originalValue = inputTexture.read(position); const auto meanValue = (half3)mean; From b2692ea6d4053d65de6fcdc82f7c8166673c3479 Mon Sep 17 00:00:00 2001 From: eugene Date: Thu, 4 Mar 2021 18:14:11 +0300 Subject: [PATCH 4/5] Revert "actualize with master" This reverts commit eb5218f6924b23dfa8a55dd4935147ef0a4e89fa. --- .gitignore | 1 - Alloy.podspec | 2 +- Package.swift | 6 +- .../Extensions/Foundation/Bundle+Module.swift | 2 +- .../Extensions/Metal/MTLDevice+Features.swift | 19 ---- .../Metal/MTLTexture+Serialization.swift | 8 +- Sources/Alloy/Encoders/Normalization.swift | 86 ------------------- .../Renderers/BoundingBoxesRenderer.swift | 2 +- Sources/Alloy/Renderers/LinesRenderer.swift | 2 +- Sources/Alloy/Renderers/MaskRenderer.swift | 2 +- .../Alloy/Renderers/RectangleRenderer.swift | 2 +- Sources/Alloy/Shaders/Shaders.metal | 16 ++-- 12 files changed, 20 insertions(+), 128 deletions(-) delete mode 100644 Sources/Alloy/Encoders/Normalization.swift diff --git a/.gitignore b/.gitignore index b82bda9..6c8e705 100644 --- a/.gitignore +++ b/.gitignore @@ -6,4 +6,3 @@ Demo/Demo.xcodeproj/xcuserdata/ Demo/Demo.xcworkspace/xcuserdata/ *gen .swiftpm -.build diff --git a/Alloy.podspec b/Alloy.podspec index 1d5a21c..0ddb7cc 100644 --- a/Alloy.podspec +++ b/Alloy.podspec @@ -1,6 +1,6 @@ Pod::Spec.new do |s| s.name = 'Alloy' - s.version = '0.16.5' + s.version = '0.16.3' s.summary = 'Nano helpers for Metal framework' s.homepage = 'https://github.com/s1ddok/Alloy' diff --git a/Package.swift b/Package.swift index 8147c54..7227e60 100644 --- a/Package.swift +++ b/Package.swift @@ -17,7 +17,8 @@ let package = Package( publicHeadersPath: "."), .target(name: "Alloy", dependencies: [.target(name: "AlloyShadersSharedTypes")], - resources: [.process("Shaders/Shaders.metal")]), + resources: [.process("Shaders/Shaders.metal")], + swiftSettings: [.define("SWIFT_PM")]), .target(name: "AlloyTestsResources", path: "Tests/AlloyTestsResources", resources: [ @@ -26,6 +27,7 @@ let package = Package( ]), .testTarget(name: "AlloyTests", dependencies: ["Alloy", "AlloyTestsResources"], - resources: [.process("Shaders/Shaders.metal")]) + resources: [.process("Shaders/Shaders.metal")], + swiftSettings: [.define("SWIFT_PM")]) ] ) diff --git a/Sources/Alloy/Core/Extensions/Foundation/Bundle+Module.swift b/Sources/Alloy/Core/Extensions/Foundation/Bundle+Module.swift index cd155e5..95bf9bf 100644 --- a/Sources/Alloy/Core/Extensions/Foundation/Bundle+Module.swift +++ b/Sources/Alloy/Core/Extensions/Foundation/Bundle+Module.swift @@ -1,5 +1,5 @@ import Foundation -#if !SWIFT_PACKAGE +#if !SWIFT_PM public extension Bundle { static var module = Bundle(for: MTLContext.self) } diff --git a/Sources/Alloy/Core/Extensions/Metal/MTLDevice+Features.swift b/Sources/Alloy/Core/Extensions/Metal/MTLDevice+Features.swift index 5f09708..aee4448 100644 --- a/Sources/Alloy/Core/Extensions/Metal/MTLDevice+Features.swift +++ b/Sources/Alloy/Core/Extensions/Metal/MTLDevice+Features.swift @@ -2,7 +2,6 @@ import Metal public enum Feature { case nonUniformThreadgroups - case readWriteTextures(MTLPixelFormat) } public extension MTLDevice { @@ -16,24 +15,6 @@ public extension MTLDevice { #elseif os(macOS) return self.supportsFeatureSet(.macOS_GPUFamily1_v3) #endif - - case let .readWriteTextures(pixelFormat): - let tierOneSupportedPixelFormats: Set = [ - .r32Float, .r32Uint, .r32Sint - ] - let tierTwoSupportedPixelFormats: Set = tierOneSupportedPixelFormats.union([ - .rgba32Float, .rgba32Uint, .rgba32Sint, .rgba16Float, - .rgba16Uint, .rgba16Sint, .rgba8Unorm, .rgba8Uint, - .rgba8Sint, .r16Float, .r16Uint, .r16Sint, - .r8Unorm, .r8Uint, .r8Sint - ]) - - switch self.readWriteTextureSupport { - case .tier1: return tierOneSupportedPixelFormats.contains(pixelFormat) - case .tier2: return tierTwoSupportedPixelFormats.contains(pixelFormat) - case .tierNone: return false - @unknown default: return false - } } } } diff --git a/Sources/Alloy/Core/Extensions/Metal/MTLTexture+Serialization.swift b/Sources/Alloy/Core/Extensions/Metal/MTLTexture+Serialization.swift index 705f8ea..3df7d0f 100644 --- a/Sources/Alloy/Core/Extensions/Metal/MTLTexture+Serialization.swift +++ b/Sources/Alloy/Core/Extensions/Metal/MTLTexture+Serialization.swift @@ -38,10 +38,8 @@ public class MTLTextureCodableBox: Codable { textureView.getBytes(pointer.advanced(by: offset), bytesPerRow: bytesPerRow, - bytesPerImage: bytesPerImage, from: textureView.region, - mipmapLevel: 0, - slice: 0) + mipmapLevel: 0) offset += bytesPerImage } @@ -82,10 +80,8 @@ public class MTLTextureCodableBox: Codable { textureView.replace(region: textureView.region, mipmapLevel: 0, - slice: 0, withBytes: pointer.advanced(by: offset), - bytesPerRow: bytesPerRow, - bytesPerImage: bytesPerImage) + bytesPerRow: bytesPerRow) offset += bytesPerImage } diff --git a/Sources/Alloy/Encoders/Normalization.swift b/Sources/Alloy/Encoders/Normalization.swift deleted file mode 100644 index 6ccd330..0000000 --- a/Sources/Alloy/Encoders/Normalization.swift +++ /dev/null @@ -1,86 +0,0 @@ -import Metal - -final public class Normalization { - - // MARK: - Propertires - - private let pipelineState: MTLComputePipelineState - private let deviceSupportsNonuniformThreadgroups: Bool - - // MARK: - Life Cycle - - public convenience init(context: MTLContext) throws { - try self.init(library: context.library(for: .module)) - } - - public init(library: MTLLibrary) throws { - self.deviceSupportsNonuniformThreadgroups = library.device - .supports(feature: .nonUniformThreadgroups) - let constantValues = MTLFunctionConstantValues() - constantValues.set(self.deviceSupportsNonuniformThreadgroups, - at: 0) - self.pipelineState = try library.computePipelineState(function: Self.functionName, - constants: constantValues) - } - - // MARK: - Encode - - public func callAsFunction(source: MTLTexture, - mean: SIMD3, - std: SIMD3, - destination: MTLTexture, - in commandBuffer: MTLCommandBuffer) { - self.encode(source: source, - mean: mean, - std: std, - destination: destination, - in: commandBuffer) - } - - public func callAsFunction(source: MTLTexture, - mean: SIMD3, - std: SIMD3, - destination: MTLTexture, - using encoder: MTLComputeCommandEncoder) { - self.encode(source: source, - mean: mean, - std: std, - destination: destination, - using: encoder) - } - - - public func encode(source: MTLTexture, - mean: SIMD3, - std: SIMD3, - destination: MTLTexture, - in commandBuffer: MTLCommandBuffer) { - commandBuffer.compute { encoder in - self.encode(source: source, - mean: mean, - std: std, - destination: destination, - using: encoder) - } - } - - public func encode(source: MTLTexture, - mean: SIMD3, - std: SIMD3, - destination: MTLTexture, - using encoder: MTLComputeCommandEncoder) { - encoder.setTextures(source, destination) - encoder.setValue(mean, at: 0) - encoder.setValue(std, at: 1) - - if self.deviceSupportsNonuniformThreadgroups { - encoder.dispatch2d(state: self.pipelineState, - exactly: source.size) - } else { - encoder.dispatch2d(state: self.pipelineState, - covering: source.size) - } - } - - public static let functionName = "normalization" -} diff --git a/Sources/Alloy/Renderers/BoundingBoxesRenderer.swift b/Sources/Alloy/Renderers/BoundingBoxesRenderer.swift index 43d1746..6a0949f 100644 --- a/Sources/Alloy/Renderers/BoundingBoxesRenderer.swift +++ b/Sources/Alloy/Renderers/BoundingBoxesRenderer.swift @@ -1,6 +1,6 @@ import Metal import CoreGraphics -#if SWIFT_PACKAGE +#if SWIFT_PM import AlloyShadersSharedTypes #endif diff --git a/Sources/Alloy/Renderers/LinesRenderer.swift b/Sources/Alloy/Renderers/LinesRenderer.swift index 18af8e2..1665bda 100644 --- a/Sources/Alloy/Renderers/LinesRenderer.swift +++ b/Sources/Alloy/Renderers/LinesRenderer.swift @@ -1,5 +1,5 @@ import Metal -#if SWIFT_PACKAGE +#if SWIFT_PM import AlloyShadersSharedTypes #endif diff --git a/Sources/Alloy/Renderers/MaskRenderer.swift b/Sources/Alloy/Renderers/MaskRenderer.swift index 6e4eb78..9e935b4 100644 --- a/Sources/Alloy/Renderers/MaskRenderer.swift +++ b/Sources/Alloy/Renderers/MaskRenderer.swift @@ -1,7 +1,7 @@ import Metal import simd import CoreGraphics -#if SWIFT_PACKAGE +#if SWIFT_PM import AlloyShadersSharedTypes #endif diff --git a/Sources/Alloy/Renderers/RectangleRenderer.swift b/Sources/Alloy/Renderers/RectangleRenderer.swift index 2010161..3224773 100644 --- a/Sources/Alloy/Renderers/RectangleRenderer.swift +++ b/Sources/Alloy/Renderers/RectangleRenderer.swift @@ -1,6 +1,6 @@ import Metal import CoreGraphics -#if SWIFT_PACKAGE +#if SWIFT_PM import AlloyShadersSharedTypes #endif diff --git a/Sources/Alloy/Shaders/Shaders.metal b/Sources/Alloy/Shaders/Shaders.metal index 35330b5..581a1c5 100644 --- a/Sources/Alloy/Shaders/Shaders.metal +++ b/Sources/Alloy/Shaders/Shaders.metal @@ -94,7 +94,7 @@ void textureMask(texture2d source, filter::linear); const auto positionF = float2(position); const auto textureSizeF = float2(textureSize); - const auto normalizedPosition = (positionF + 0.5f) / textureSizeF; + const auto normalizedPosition = (positionF.x + 0.5f) / textureSizeF; auto maskValue = mask.sample(s, normalizedPosition); if (isInversed) { @@ -656,15 +656,15 @@ kernel void textureDifferenceHighlight(texture2d sourceOne // MARK: - ML -kernel void normalization(texture2d inputTexture [[ texture(0) ]], - texture2d outputTexture [[ texture(1) ]], - constant float3& mean [[ buffer(0) ]], - constant float3& std [[ buffer(1) ]], - uint2 position [[thread_position_in_grid]]) { +kernel void normalize(texture2d inputTexture [[ texture(0) ]], + texture2d outputTexture [[ texture(1) ]], + constant float3& mean [[ buffer(0) ]], + constant float3& std [[ buffer(1) ]], + uint2 position [[thread_position_in_grid]]) { const auto textureSize = ushort2(inputTexture.get_width(), - inputTexture.get_height()); + inputTexture.get_height()); checkPosition(position, textureSize, deviceSupportsNonuniformThreadgroups); - + // Read mpsnngraph result value. const auto originalValue = inputTexture.read(position); const auto meanValue = (half3)mean; From c065328af0f361c0cd6fb43c1008b8bd0da401fa Mon Sep 17 00:00:00 2001 From: eugene Date: Thu, 4 Mar 2021 19:12:15 +0300 Subject: [PATCH 5/5] style fixes --- .../MPSImage+Extensions.swift | 40 ++++++++++++++----- Sources/Alloy/Core/Float16.swift | 26 ++++++++---- 2 files changed, 48 insertions(+), 18 deletions(-) diff --git a/Sources/Alloy/Core/Extensions/MetalPerformanceShaders/MPSImage+Extensions.swift b/Sources/Alloy/Core/Extensions/MetalPerformanceShaders/MPSImage+Extensions.swift index ffd4c03..263dc75 100644 --- a/Sources/Alloy/Core/Extensions/MetalPerformanceShaders/MPSImage+Extensions.swift +++ b/Sources/Alloy/Core/Extensions/MetalPerformanceShaders/MPSImage+Extensions.swift @@ -9,36 +9,53 @@ public extension MPSImage { func toFloatArray() -> [Float]? { switch self.pixelFormat { case .r16Float, .rg16Float, .rgba16Float: - guard var outputFloat16 = self.convert(initial: Float16(0)) else { return nil } - return float16to32(&outputFloat16, count: outputFloat16.count) + if #available(iOS 14.0, macOS 11.0, macCatalyst 14.5, *), + let float16Array: [Swift.Float16] = self.toArray() { + return float16Array.map(Float.init) + } else if var float16Array: [Float16] = self.toArray() { + return float16to32(&float16Array, + count: float16Array.count) + } else { return nil } case .r32Float, .rg32Float, .rgba32Float, .depth32Float: - return self.convert(initial: Float(0)) + return self.toArray() case .invalid: return nil default: return nil } } - private func convert(initial: T) -> [T]? { - guard self.texture.isAccessibleOnCPU else { return nil } + private func toArray() -> [T]? { + guard self.texture.isAccessibleOnCPU + else { return nil } let numSlices = (self.featureChannels + 3) / 4 /// If the number of channels is not a multiple of 4, we may need to add /// padding. For 1 and 2 channels we don't need padding. - let channelsPlusPadding = (self.featureChannels < 3) ? self.featureChannels : numSlices * 4 + let channelsPlusPadding = (self.featureChannels < 3) + ? self.featureChannels + : numSlices * 4 /// How many elements we need to copy over from each pixel in a slice. /// For 1 channel it's just 1 element (R); for 2 channels it is 2 elements /// (R+G), and for any other number of channels it is 4 elements (RGBA). - let numComponents = (self.featureChannels < 3) ? self.featureChannels : 4 + let numComponents = (self.featureChannels < 3) + ? self.featureChannels + : 4 /// Allocate the memory for the array. If batching is used, then we need to /// copy numSlices slices for each image in the batch. - let count = self.width * self.height * channelsPlusPadding * self.numberOfImages - var output = [T](repeating: initial, count: count) + let count = self.width + * self.height + * channelsPlusPadding + * self.numberOfImages - let region = MTLRegion(origin: MTLOrigin(x: 0, y: 0, z: 0), - size: MTLSize(width: self.width, height: self.height, depth: 1)) + var output = [T](repeating: .zero, + count: count) + + let region = MTLRegion(origin: .zero, + size: .init(width: self.width, + height: self.height, + depth: 1)) for i in 0 ..< numSlices * self.numberOfImages { self.texture.getBytes(&(output[self.width * self.height * numComponents * i]), @@ -48,6 +65,7 @@ public extension MPSImage { mipmapLevel: 0, slice: i) } + return output } } diff --git a/Sources/Alloy/Core/Float16.swift b/Sources/Alloy/Core/Float16.swift index 3f72858..3ec3e74 100644 --- a/Sources/Alloy/Core/Float16.swift +++ b/Sources/Alloy/Core/Float16.swift @@ -14,8 +14,10 @@ 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]? { - let output = [Float](repeating: 0, count: count) +public func float16to32(_ input: UnsafeMutableRawPointer, + count: Int) -> [Float]? { + let output = [Float](repeating: 0, + count: count) var data: UnsafeMutableRawPointer? = nil output.withUnsafeBytes { data = UnsafeMutableRawPointer(mutating: $0.baseAddress) @@ -45,17 +47,27 @@ 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]? { - let output = [Float16](repeating: 0, count: count) +public func float32to16(_ input: UnsafeMutablePointer, + count: Int) -> [Float16]? { + let output = [Float16](repeating: 0, + count: count) var data: UnsafeMutableRawPointer? = nil output.withUnsafeBytes { data = UnsafeMutableRawPointer(mutating: $0.baseAddress) } - var bufferFloat32 = vImage_Buffer(data: input, height: 1, width: UInt(count), rowBytes: count * 4) - var bufferFloat16 = vImage_Buffer(data: data, height: 1, width: UInt(count), rowBytes: count * 2) + var bufferFloat32 = vImage_Buffer(data: input, + height: 1, + width: UInt(count), + rowBytes: count * 4) + var bufferFloat16 = vImage_Buffer(data: data, + height: 1, + width: UInt(count), + rowBytes: count * 2) - if vImageConvert_PlanarFtoPlanar16F(&bufferFloat32, &bufferFloat16, 0) != kvImageNoError { + if vImageConvert_PlanarFtoPlanar16F(&bufferFloat32, + &bufferFloat16, + 0) != kvImageNoError { return nil } return output