From 411807fc2ba4e665cfb74c0d3d8dc29e4320163e Mon Sep 17 00:00:00 2001 From: eugene Date: Fri, 12 Feb 2021 18:50:54 +0300 Subject: [PATCH 1/2] add read write texture decl --- Sources/mtlswift/AST/ASTNode.swift | 2 +- Sources/mtlswift/AST/ASTShader.swift | 10 +++- Sources/mtlswift/AST/CustomDeclaration.swift | 11 ++++ .../Encoder Generator/MTLKernelEncoder.swift | 60 ++++++++++++++++--- 4 files changed, 71 insertions(+), 12 deletions(-) diff --git a/Sources/mtlswift/AST/ASTNode.swift b/Sources/mtlswift/AST/ASTNode.swift index 2e5ecb5..6e8c3f4 100644 --- a/Sources/mtlswift/AST/ASTNode.swift +++ b/Sources/mtlswift/AST/ASTNode.swift @@ -21,7 +21,7 @@ public class ASTNode { public init(parsingString inputString: String) throws { let scanner = StringScanner(string: inputString) - + guard let prefix = scanner.readWord() else { throw Errors.parsingError } diff --git a/Sources/mtlswift/AST/ASTShader.swift b/Sources/mtlswift/AST/ASTShader.swift index da39f17..d602825 100644 --- a/Sources/mtlswift/AST/ASTShader.swift +++ b/Sources/mtlswift/AST/ASTShader.swift @@ -58,14 +58,18 @@ public struct ASTShader { var swiftNameLookup: [String: String] = [:] var swiftTypeLookup: [String: String] = [:] + var inPlaceTextures = [String]() for declaration in self.customDeclarations { - if case .swiftParameterName(let oldName, let newName) = declaration { + if case let .swiftParameterName(oldName, newName) = declaration { swiftNameLookup[oldName] = newName } - if case .swiftParameterType(let parameter, let type) = declaration { + if case let .swiftParameterType(parameter, type) = declaration { swiftTypeLookup[parameter] = type } + if case let .inPlaceTexture(textureName) = declaration { + inPlaceTextures.append(textureName) + } } // individual access levels are not supported for now @@ -95,7 +99,7 @@ public struct ASTShader { return MTLKernelEncoder.Parameter(name: name, swiftTypeName: type, - kind: .texture, + kind: inPlaceTextures.contains(p.name) ? .inPlaceTexture : .texture, index: p.index ?? -1, defaultValueString: nil) case .buffer: diff --git a/Sources/mtlswift/AST/CustomDeclaration.swift b/Sources/mtlswift/AST/CustomDeclaration.swift index 27462e9..8e2a2fe 100644 --- a/Sources/mtlswift/AST/CustomDeclaration.swift +++ b/Sources/mtlswift/AST/CustomDeclaration.swift @@ -12,6 +12,14 @@ public enum CustomDeclaration { // TODO: Check it is valid Swift identifier self = .swiftName(name: name) return + } else if scanner.skip(exact: CustomDeclaration.inPlaceTexture) { + guard let name = scanner.readWord() else { + print("ERROR: Failed to parse \(CustomDeclaration.inPlaceTexture), skipping: \(scanner.leftString)") + return nil + } + + self = .inPlaceTexture(name: name) + return } else if scanner.skip(exact: CustomDeclaration.swiftParameterNameDeclaration) { guard let oldName = scanner.readWord(), @@ -186,6 +194,9 @@ public enum CustomDeclaration { public static let accessLevelDeclaration = "accessLevel:" case accessLevel(level: AccessLevel) + + public static let inPlaceTexture = "inPlaceTexture:" + case inPlaceTexture(name: String) public static let swiftNameDeclaration = "swiftName:" case swiftName(name: String) diff --git a/Sources/mtlswift/Encoder Generator/MTLKernelEncoder.swift b/Sources/mtlswift/Encoder Generator/MTLKernelEncoder.swift index 8ace7b0..3d4bc9c 100644 --- a/Sources/mtlswift/Encoder Generator/MTLKernelEncoder.swift +++ b/Sources/mtlswift/Encoder Generator/MTLKernelEncoder.swift @@ -2,7 +2,7 @@ public struct MTLKernelEncoder { public struct Parameter { public enum Kind { - case texture, buffer, sampler, threadgroupMemory + case texture, inPlaceTexture, buffer, sampler, threadgroupMemory } public var name: String public var swiftTypeName: String @@ -45,19 +45,35 @@ public struct MTLKernelEncoder { sourceBuilder.add(line: "\(self.accessLevel.rawValue) let \(bc.name): \(bc.type.swiftTypeDelcaration)") } + let inPlaceTextures = self.parameters.filter { $0.kind == .inPlaceTexture } + let containsInPlaceTextures = !inPlaceTextures.isEmpty + sourceBuilder.blankLine() sourceBuilder.add(line: "\(self.accessLevel.rawValue) let pipelineState: MTLComputePipelineState") + if containsInPlaceTextures { + sourceBuilder.add(line: "\(self.accessLevel.rawValue) let textureCopy: TextureCopy") + } sourceBuilder.blankLine() if self.usedConstants.isEmpty && self.branchingConstant == nil { // MARK: Generate inits - sourceBuilder.add(line: "\(self.accessLevel.rawValue) init(library: MTLLibrary) throws {") + if containsInPlaceTextures { + sourceBuilder.add(line: "\(self.accessLevel.rawValue) init(library: MTLLibrary, textureCopy: TextureCopy) throws {") + } else { + sourceBuilder.add(line: "\(self.accessLevel.rawValue) init(library: MTLLibrary) throws {") + } sourceBuilder.pushLevel() sourceBuilder.add(line: "self.pipelineState = try library.computePipelineState(function: \"\(self.shaderName)\")") } else { let parameterString = ", " + self.usedConstants.map { "\($0.name): \($0.type.swiftTypeDelcaration)" }.joined(separator: ", ") - sourceBuilder.add(line: "\(self.accessLevel.rawValue) init(library: MTLLibrary\(self.usedConstants.isEmpty ? "" : parameterString)) throws {") + + if containsInPlaceTextures { + sourceBuilder.add(line: "\(self.accessLevel.rawValue) init(library: MTLLibrary, textureCopy: TextureCopy\(self.usedConstants.isEmpty ? "" : parameterString)) throws {") + } else { + sourceBuilder.add(line: "\(self.accessLevel.rawValue) init(library: MTLLibrary\(self.usedConstants.isEmpty ? "" : parameterString)) throws {") + } + sourceBuilder.pushLevel() sourceBuilder.add(line: "let constantValues = MTLFunctionConstantValues()") @@ -76,6 +92,10 @@ public struct MTLKernelEncoder { sourceBuilder.add(line: "self.pipelineState = try library.computePipelineState(function: \"\(self.shaderName)\", constants: constantValues)") } + if containsInPlaceTextures { + sourceBuilder.add(line: "self.textureCopy = textureCopy") + } + // MARK: Balancing for init sourceBuilder.popLevel() sourceBuilder.add(line: "}") @@ -143,6 +163,29 @@ public struct MTLKernelEncoder { // Encode in commandBuffer sourceBuilder.add(line: "\(self.accessLevel.rawValue) func encode(\(parameterString)\(gridSizeParameterString)\(threadgroupParameterString)in commandBuffer: MTLCommandBuffer) {") sourceBuilder.pushLevel() + + if containsInPlaceTextures { + for inPlaceTextureParameter in inPlaceTextures { + let name = inPlaceTextureParameter.name + let imageCopyName = "\(name)CopyImage" + let originalTextureName = "\(name)OriginalTexture" + + sourceBuilder.add(line: "var \(name) = \(name)") + sourceBuilder.add(line: "if !self.pipelineState.device.supports(feature: .readWriteTextures(\(name).pixelFormat)) {") + sourceBuilder.pushLevel() + + sourceBuilder.add(line: "let \(originalTextureName) = \(name)") + sourceBuilder.add(line: "let \(imageCopyName) = \(name).matchingTemporaryImage(commandBuffer: commandBuffer)") + sourceBuilder.add(line: "defer { \(imageCopyName).readCount = .zero }") + sourceBuilder.add(line: "\(name) = \(imageCopyName).texture") + sourceBuilder.add(line: "self.textureCopy(source: \(originalTextureName), destination: \(name), in: commandBuffer)") + sourceBuilder.add(line: "}") + sourceBuilder.popLevel() + + sourceBuilder.blankLine() + } + } + sourceBuilder.add(line: "commandBuffer.compute { encoder in") sourceBuilder.pushLevel() sourceBuilder.add(line: "encoder.label = \"\(self.swiftName)\"") @@ -153,7 +196,8 @@ public struct MTLKernelEncoder { sourceBuilder.add(line: "}") // Ecode using encoder - sourceBuilder.add(line: "\(self.accessLevel.rawValue) func encode(\(parameterString)\(gridSizeParameterString)\(threadgroupParameterString)using encoder: MTLComputeCommandEncoder) {") + let accessLevel = containsInPlaceTextures ? "private" : "\(self.accessLevel.rawValue)" + sourceBuilder.add(line: "\(accessLevel) func encode(\(parameterString)\(gridSizeParameterString)\(threadgroupParameterString)using encoder: MTLComputeCommandEncoder) {") } sourceBuilder.pushLevel() sourceBuilder.add(line: threadgroupVariableString) @@ -166,7 +210,7 @@ public struct MTLKernelEncoder { } else { sourceBuilder.add(line: "encoder.setValue(\(parameter.name), at: \(parameter.index))") } - case .texture: + case .texture, .inPlaceTexture: sourceBuilder.add(line: "encoder.setTexture(\(parameter.name), index: \(parameter.index))") case .sampler: sourceBuilder.add(line: "encoder.setSamplerState(\(parameter.name), index: \(parameter.index))") @@ -201,7 +245,7 @@ public struct MTLKernelEncoder { case .even(parameters: .over(let argument)): if let targetParameter = self.parameters.first(where: { $0.name == argument }), - targetParameter.kind == .texture { + (targetParameter.kind == .texture || targetParameter.kind == .inPlaceTexture) { sourceBuilder.add(line: "encoder.dispatch2d(state: self.pipelineState, covering: \(targetParameter.name).size\(threadgroupExpressionString))") } else { fatalError("Could not generate dispatching over parameter \(argument)") @@ -216,7 +260,7 @@ public struct MTLKernelEncoder { case .exact(parameters: .over(let argument)): if let targetParameter = self.parameters.first(where: { $0.name == argument }), - targetParameter.kind == .texture { + (targetParameter.kind == .texture || targetParameter.kind == .inPlaceTexture) { sourceBuilder.add(line: "encoder.dispatch2d(state: self.pipelineState, exactly: \(targetParameter.name).size\(threadgroupExpressionString))") } else { print("Could not generate dispatching over parameter \(argument)") @@ -231,7 +275,7 @@ public struct MTLKernelEncoder { sourceBuilder.add(line: "if self.\(bc.name) { encoder.dispatch2d(state: self.pipelineState, exactly: \(self.swiftName).gridSize\(idx)\(threadgroupExpressionString)) } else { encoder.dispatch2d(state: self.pipelineState, covering: \(self.swiftName).gridSize\(idx)\(threadgroupExpressionString)) }") case .optimal(_, parameters: .over(let argument)): if let targetParameter = self.parameters.first(where: { $0.name == argument }), - targetParameter.kind == .texture { + (targetParameter.kind == .texture || targetParameter.kind == .inPlaceTexture) { let bc = self.branchingConstant! sourceBuilder.add(line: "if self.\(bc.name) { encoder.dispatch2d(state: self.pipelineState, exactly: \(targetParameter.name).size\(threadgroupExpressionString)) } else { encoder.dispatch2d(state: self.pipelineState, covering: \(targetParameter.name).size\(threadgroupExpressionString)) }") } else { print("Could not generate dispatching over parameter \(argument)") } From 91577b551911dde03e9fb3821624e73db81b42f7 Mon Sep 17 00:00:00 2001 From: Eugene Bokhan Date: Fri, 19 Feb 2021 23:41:10 +0300 Subject: [PATCH 2/2] update --- Sources/mtlswift/AST/ASTShader.swift | 13 ++- Sources/mtlswift/AST/CustomDeclaration.swift | 16 ++- .../Encoder Generator/MTLKernelEncoder.swift | 109 +++++++++++++----- 3 files changed, 100 insertions(+), 38 deletions(-) diff --git a/Sources/mtlswift/AST/ASTShader.swift b/Sources/mtlswift/AST/ASTShader.swift index d602825..8bc0b4a 100644 --- a/Sources/mtlswift/AST/ASTShader.swift +++ b/Sources/mtlswift/AST/ASTShader.swift @@ -58,7 +58,7 @@ public struct ASTShader { var swiftNameLookup: [String: String] = [:] var swiftTypeLookup: [String: String] = [:] - var inPlaceTextures = [String]() + var inPlaceTextureNameMappings = [MTLKernelEncoder.InPlaceTextureNameMapping]() for declaration in self.customDeclarations { if case let .swiftParameterName(oldName, newName) = declaration { swiftNameLookup[oldName] = newName @@ -67,8 +67,8 @@ public struct ASTShader { if case let .swiftParameterType(parameter, type) = declaration { swiftTypeLookup[parameter] = type } - if case let .inPlaceTexture(textureName) = declaration { - inPlaceTextures.append(textureName) + if case let .inPlaceTexture(source, destination, inPlace) = declaration { + inPlaceTextureNameMappings.append(.init(source: source, destination: destination, inPlace: inPlace)) } } @@ -96,10 +96,10 @@ public struct ASTShader { print("WARNING: Swift Types are not available for texture parameters, ignoring \(type)") type = "MTLTexture" } - + return MTLKernelEncoder.Parameter(name: name, swiftTypeName: type, - kind: inPlaceTextures.contains(p.name) ? .inPlaceTexture : .texture, + kind: .texture, index: p.index ?? -1, defaultValueString: nil) case .buffer: @@ -176,7 +176,8 @@ public struct ASTShader { encodingVariants: [MTLKernelEncoder.EncodingVariant(dispatchType: type, threadgroupSize: size)], usedConstants: constants, branchingConstant: branchingConstant, - threadgroupMemoryCalculations: threadgroupMemoryCalculatiosn) + threadgroupMemoryCalculations: threadgroupMemoryCalculatiosn, + inPlaceTextureNameMappings: inPlaceTextureNameMappings) } return nil diff --git a/Sources/mtlswift/AST/CustomDeclaration.swift b/Sources/mtlswift/AST/CustomDeclaration.swift index 8e2a2fe..5e43130 100644 --- a/Sources/mtlswift/AST/CustomDeclaration.swift +++ b/Sources/mtlswift/AST/CustomDeclaration.swift @@ -4,7 +4,8 @@ public enum CustomDeclaration { let scanner = StringScanner(string: rawString) if scanner.skip(exact: CustomDeclaration.swiftNameDeclaration) { - guard let name = scanner.readWord() else { + guard let name = scanner.readWord() + else { print("ERROR: Failed to parse \(CustomDeclaration.swiftNameDeclaration), skipping: \(scanner.leftString)") return nil } @@ -13,12 +14,19 @@ public enum CustomDeclaration { self = .swiftName(name: name) return } else if scanner.skip(exact: CustomDeclaration.inPlaceTexture) { - guard let name = scanner.readWord() else { + guard let source = scanner.readWord(), + scanner.skip(exact: ":"), + let destination = scanner.readWord(), + scanner.skip(exact: ":"), + let inPlace = scanner.readWord() + else { print("ERROR: Failed to parse \(CustomDeclaration.inPlaceTexture), skipping: \(scanner.leftString)") return nil } - self = .inPlaceTexture(name: name) + self = .inPlaceTexture(source: source, + destination: destination, + inPlace: inPlace) return } else if scanner.skip(exact: CustomDeclaration.swiftParameterNameDeclaration) { guard @@ -196,7 +204,7 @@ public enum CustomDeclaration { case accessLevel(level: AccessLevel) public static let inPlaceTexture = "inPlaceTexture:" - case inPlaceTexture(name: String) + case inPlaceTexture(source: String, destination: String, inPlace: String) public static let swiftNameDeclaration = "swiftName:" case swiftName(name: String) diff --git a/Sources/mtlswift/Encoder Generator/MTLKernelEncoder.swift b/Sources/mtlswift/Encoder Generator/MTLKernelEncoder.swift index 3d4bc9c..24269f7 100644 --- a/Sources/mtlswift/Encoder Generator/MTLKernelEncoder.swift +++ b/Sources/mtlswift/Encoder Generator/MTLKernelEncoder.swift @@ -1,8 +1,8 @@ public struct MTLKernelEncoder { - public struct Parameter { + public struct Parameter: Hashable { public enum Kind { - case texture, inPlaceTexture, buffer, sampler, threadgroupMemory + case texture, buffer, sampler, threadgroupMemory } public var name: String public var swiftTypeName: String @@ -12,6 +12,12 @@ public struct MTLKernelEncoder { public var defaultValueString: String? = nil } + + public struct InPlaceTextureNameMapping { + public let source: String + public let destination: String + public let inPlace: String + } public enum ThreadgroupMemoryLengthCalculation { case total(index: Int, bytes: Int) @@ -32,6 +38,7 @@ public struct MTLKernelEncoder { public var usedConstants: [ASTFunctionConstant] public var branchingConstant: ASTFunctionConstant? public var threadgroupMemoryCalculations: [ThreadgroupMemoryLengthCalculation] + public var inPlaceTextureNameMappings: [InPlaceTextureNameMapping] public var shaderString: String { let sourceBuilder = SourceStringBuilder() @@ -45,8 +52,7 @@ public struct MTLKernelEncoder { sourceBuilder.add(line: "\(self.accessLevel.rawValue) let \(bc.name): \(bc.type.swiftTypeDelcaration)") } - let inPlaceTextures = self.parameters.filter { $0.kind == .inPlaceTexture } - let containsInPlaceTextures = !inPlaceTextures.isEmpty + let containsInPlaceTextures = !self.inPlaceTextureNameMappings.isEmpty sourceBuilder.blankLine() sourceBuilder.add(line: "\(self.accessLevel.rawValue) let pipelineState: MTLComputePipelineState") @@ -133,11 +139,48 @@ public struct MTLKernelEncoder { sourceBuilder.add(line: "\(self.accessLevel.rawValue) func encode(\(gridSizeParameterString)\(threadgroupParameterString)using encoder: MTLComputeCommandEncoder) {") } else { var parameterString = "" + var filteredParameterString = "" + + let inPlaceParameters = self.inPlaceTextureNameMappings.map { mapping -> Parameter in + return .init(name: mapping.inPlace, + swiftTypeName: "MTLTexture", + kind: .texture, + index: -1) + } + let sourceParameters = self.inPlaceTextureNameMappings.map { mapping -> Parameter in + let index = self.parameters.first(where: { $0.name == mapping.source })?.index ?? 0 + return .init(name: mapping.source, + swiftTypeName: "MTLTexture", + kind: .texture, + index: index) + } + let destinationParameters = self.inPlaceTextureNameMappings.map { mapping -> Parameter in + let index = self.parameters.first(where: { $0.name == mapping.destination })?.index ?? 0 + return .init(name: mapping.destination, + swiftTypeName: "MTLTexture", + kind: .texture, + index: index) + } + let filteredParameters = containsInPlaceTextures + ? inPlaceParameters + Set(self.parameters).subtracting(Set(sourceParameters)) + .subtracting(Set(destinationParameters)) + .map { $0 } + : self.parameters + for parameter in self.parameters { parameterString += "\(parameter.name): \(parameter.swiftTypeName), " } + + for filteredParameter in filteredParameters { + filteredParameterString += "\(filteredParameter.name): \(filteredParameter.swiftTypeName), " + } + + if containsInPlaceTextures { + print(2) + } var parametersBodyString = "" + var filteredParametersBodyString = "" let gridSizeValueString = gridSizeParameterString.isEmpty ? "" : ", gridSize: gridSize" let threadgroupSizeValueString = threadgroupParameterString.isEmpty ? "" : ", threadgroupSize: threadgroupSize" for parameterIndex in 0 ..< self.parameters.count { @@ -145,40 +188,50 @@ public struct MTLKernelEncoder { let parameterSeparator = parameterIndex < self.parameters.count - 1 ? ", " : "" parametersBodyString += parameterName + ": " + parameterName + parameterSeparator } + for parameterIndex in 0 ..< filteredParameters.count { + let parameterName = filteredParameters[parameterIndex].name + let parameterSeparator = parameterIndex < filteredParameters.count - 1 ? ", " : "" + filteredParametersBodyString += parameterName + ": " + parameterName + parameterSeparator + } // Call as function in commandBuffer - sourceBuilder.add(line: "\(self.accessLevel.rawValue) func callAsFunction(\(parameterString)\(gridSizeParameterString)\(threadgroupParameterString)in commandBuffer: MTLCommandBuffer) {") + sourceBuilder.add(line: "\(self.accessLevel.rawValue) func callAsFunction(\(filteredParameterString)\(gridSizeParameterString)\(threadgroupParameterString)in commandBuffer: MTLCommandBuffer) {") sourceBuilder.pushLevel() - sourceBuilder.add(line: "self.encode(\(parametersBodyString)\(gridSizeValueString)\(threadgroupSizeValueString), in: commandBuffer)") + sourceBuilder.add(line: "self.encode(\(filteredParametersBodyString)\(gridSizeValueString)\(threadgroupSizeValueString), in: commandBuffer)") sourceBuilder.popLevel() sourceBuilder.add(line: "}") - // Call as function using encoder - sourceBuilder.add(line: "\(self.accessLevel.rawValue) func callAsFunction(\(parameterString)\(gridSizeParameterString)\(threadgroupParameterString)using encoder: MTLComputeCommandEncoder) {") - sourceBuilder.pushLevel() - sourceBuilder.add(line: "self.encode(\(parametersBodyString)\(gridSizeValueString)\(threadgroupSizeValueString), using: encoder)") - sourceBuilder.popLevel() - sourceBuilder.add(line: "}") + if !containsInPlaceTextures { + // Call as function using encoder + sourceBuilder.add(line: "\(self.accessLevel.rawValue) func callAsFunction(\(filteredParameterString)\(gridSizeParameterString)\(threadgroupParameterString)using encoder: MTLComputeCommandEncoder) {") + sourceBuilder.pushLevel() + sourceBuilder.add(line: "self.encode(\(filteredParametersBodyString)\(gridSizeValueString)\(threadgroupSizeValueString), using: encoder)") + sourceBuilder.popLevel() + sourceBuilder.add(line: "}") + } // Encode in commandBuffer - sourceBuilder.add(line: "\(self.accessLevel.rawValue) func encode(\(parameterString)\(gridSizeParameterString)\(threadgroupParameterString)in commandBuffer: MTLCommandBuffer) {") + sourceBuilder.add(line: "\(self.accessLevel.rawValue) func encode(\(filteredParameterString)\(gridSizeParameterString)\(threadgroupParameterString)in commandBuffer: MTLCommandBuffer) {") sourceBuilder.pushLevel() if containsInPlaceTextures { - for inPlaceTextureParameter in inPlaceTextures { - let name = inPlaceTextureParameter.name - let imageCopyName = "\(name)CopyImage" - let originalTextureName = "\(name)OriginalTexture" - - sourceBuilder.add(line: "var \(name) = \(name)") - sourceBuilder.add(line: "if !self.pipelineState.device.supports(feature: .readWriteTextures(\(name).pixelFormat)) {") + for inPlaceTextureNameMapping in self.inPlaceTextureNameMappings { + let sourceName = inPlaceTextureNameMapping.source + let destinationName = inPlaceTextureNameMapping.destination + let inPlaceName = inPlaceTextureNameMapping.inPlace + let imageCopyName = "\(sourceName)CopyImage" + let originalTextureName = "\(sourceName)OriginalTexture" + + sourceBuilder.add(line: "var \(sourceName) = \(inPlaceName)") + sourceBuilder.add(line: "let \(destinationName) = \(inPlaceName)") + sourceBuilder.add(line: "if !self.pipelineState.device.supports(feature: .readWriteTextures(\(sourceName).pixelFormat)) {") sourceBuilder.pushLevel() - sourceBuilder.add(line: "let \(originalTextureName) = \(name)") - sourceBuilder.add(line: "let \(imageCopyName) = \(name).matchingTemporaryImage(commandBuffer: commandBuffer)") + sourceBuilder.add(line: "let \(originalTextureName) = \(sourceName)") + sourceBuilder.add(line: "let \(imageCopyName) = \(sourceName).matchingTemporaryImage(commandBuffer: commandBuffer)") sourceBuilder.add(line: "defer { \(imageCopyName).readCount = .zero }") - sourceBuilder.add(line: "\(name) = \(imageCopyName).texture") - sourceBuilder.add(line: "self.textureCopy(source: \(originalTextureName), destination: \(name), in: commandBuffer)") + sourceBuilder.add(line: "\(sourceName) = \(imageCopyName).texture") + sourceBuilder.add(line: "self.textureCopy(source: \(originalTextureName), destination: \(sourceName), in: commandBuffer)") sourceBuilder.add(line: "}") sourceBuilder.popLevel() @@ -210,7 +263,7 @@ public struct MTLKernelEncoder { } else { sourceBuilder.add(line: "encoder.setValue(\(parameter.name), at: \(parameter.index))") } - case .texture, .inPlaceTexture: + case .texture: sourceBuilder.add(line: "encoder.setTexture(\(parameter.name), index: \(parameter.index))") case .sampler: sourceBuilder.add(line: "encoder.setSamplerState(\(parameter.name), index: \(parameter.index))") @@ -245,7 +298,7 @@ public struct MTLKernelEncoder { case .even(parameters: .over(let argument)): if let targetParameter = self.parameters.first(where: { $0.name == argument }), - (targetParameter.kind == .texture || targetParameter.kind == .inPlaceTexture) { + (targetParameter.kind == .texture) { sourceBuilder.add(line: "encoder.dispatch2d(state: self.pipelineState, covering: \(targetParameter.name).size\(threadgroupExpressionString))") } else { fatalError("Could not generate dispatching over parameter \(argument)") @@ -260,7 +313,7 @@ public struct MTLKernelEncoder { case .exact(parameters: .over(let argument)): if let targetParameter = self.parameters.first(where: { $0.name == argument }), - (targetParameter.kind == .texture || targetParameter.kind == .inPlaceTexture) { + (targetParameter.kind == .texture) { sourceBuilder.add(line: "encoder.dispatch2d(state: self.pipelineState, exactly: \(targetParameter.name).size\(threadgroupExpressionString))") } else { print("Could not generate dispatching over parameter \(argument)") @@ -275,7 +328,7 @@ public struct MTLKernelEncoder { sourceBuilder.add(line: "if self.\(bc.name) { encoder.dispatch2d(state: self.pipelineState, exactly: \(self.swiftName).gridSize\(idx)\(threadgroupExpressionString)) } else { encoder.dispatch2d(state: self.pipelineState, covering: \(self.swiftName).gridSize\(idx)\(threadgroupExpressionString)) }") case .optimal(_, parameters: .over(let argument)): if let targetParameter = self.parameters.first(where: { $0.name == argument }), - (targetParameter.kind == .texture || targetParameter.kind == .inPlaceTexture) { + (targetParameter.kind == .texture) { let bc = self.branchingConstant! sourceBuilder.add(line: "if self.\(bc.name) { encoder.dispatch2d(state: self.pipelineState, exactly: \(targetParameter.name).size\(threadgroupExpressionString)) } else { encoder.dispatch2d(state: self.pipelineState, covering: \(targetParameter.name).size\(threadgroupExpressionString)) }") } else { print("Could not generate dispatching over parameter \(argument)") }