diff --git a/Examples/ExampleMacOS/ExampleMacOS/Logic/Functions.swift b/Examples/ExampleMacOS/ExampleMacOS/Logic/Functions.swift index 6fa961d..dd94600 100644 --- a/Examples/ExampleMacOS/ExampleMacOS/Logic/Functions.swift +++ b/Examples/ExampleMacOS/ExampleMacOS/Logic/Functions.swift @@ -43,3 +43,23 @@ class MyRender { "" } } + +@EMComputeShader3D +class MyCompute3D { + + var intensity: Float = 3 + var tex: MTLTexture? + + var impl: String { + "float2 floatGid = float2(gid.x, gid.y);" + "float2 center = float2(tex.get_width() / 2, tex.get_height() / 2);" + "float dist = distance(center, floatGid);" + "float color = intensity / dist;" + "tex.write(float4(color, color, color, 1), gid);" + } + + var customMetalCode: String { + "" + } +} + diff --git a/LICENSE b/LICENSE index fa20391..261eeb9 100644 --- a/LICENSE +++ b/LICENSE @@ -1,21 +1,201 @@ -MIT License - -Copyright (c) 2024 Yuki Kuwashima - -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in all -copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -SOFTWARE. + Apache License + Version 2.0, January 2004 + http://www.apache.org/licenses/ + + TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION + + 1. Definitions. + + "License" shall mean the terms and conditions for use, reproduction, + and distribution as defined by Sections 1 through 9 of this document. + + "Licensor" shall mean the copyright owner or entity authorized by + the copyright owner that is granting the License. + + "Legal Entity" shall mean the union of the acting entity and all + other entities that control, are controlled by, or are under common + control with that entity. For the purposes of this definition, + "control" means (i) the power, direct or indirect, to cause the + direction or management of such entity, whether by contract or + otherwise, or (ii) ownership of fifty percent (50%) or more of the + outstanding shares, or (iii) beneficial ownership of such entity. + + "You" (or "Your") shall mean an individual or Legal Entity + exercising permissions granted by this License. + + "Source" form shall mean the preferred form for making modifications, + including but not limited to software source code, documentation + source, and configuration files. + + "Object" form shall mean any form resulting from mechanical + transformation or translation of a Source form, including but + not limited to compiled object code, generated documentation, + and conversions to other media types. + + "Work" shall mean the work of authorship, whether in Source or + Object form, made available under the License, as indicated by a + copyright notice that is included in or attached to the work + (an example is provided in the Appendix below). + + "Derivative Works" shall mean any work, whether in Source or Object + form, that is based on (or derived from) the Work and for which the + editorial revisions, annotations, elaborations, or other modifications + represent, as a whole, an original work of authorship. For the purposes + of this License, Derivative Works shall not include works that remain + separable from, or merely link (or bind by name) to the interfaces of, + the Work and Derivative Works thereof. + + "Contribution" shall mean any work of authorship, including + the original version of the Work and any modifications or additions + to that Work or Derivative Works thereof, that is intentionally + submitted to Licensor for inclusion in the Work by the copyright owner + or by an individual or Legal Entity authorized to submit on behalf of + the copyright owner. For the purposes of this definition, "submitted" + means any form of electronic, verbal, or written communication sent + to the Licensor or its representatives, including but not limited to + communication on electronic mailing lists, source code control systems, + and issue tracking systems that are managed by, or on behalf of, the + Licensor for the purpose of discussing and improving the Work, but + excluding communication that is conspicuously marked or otherwise + designated in writing by the copyright owner as "Not a Contribution." + + "Contributor" shall mean Licensor and any individual or Legal Entity + on behalf of whom a Contribution has been received by Licensor and + subsequently incorporated within the Work. + + 2. Grant of Copyright License. Subject to the terms and conditions of + this License, each Contributor hereby grants to You a perpetual, + worldwide, non-exclusive, no-charge, royalty-free, irrevocable + copyright license to reproduce, prepare Derivative Works of, + publicly display, publicly perform, sublicense, and distribute the + Work and such Derivative Works in Source or Object form. + + 3. Grant of Patent License. Subject to the terms and conditions of + this License, each Contributor hereby grants to You a perpetual, + worldwide, non-exclusive, no-charge, royalty-free, irrevocable + (except as stated in this section) patent license to make, have made, + use, offer to sell, sell, import, and otherwise transfer the Work, + where such license applies only to those patent claims licensable + by such Contributor that are necessarily infringed by their + Contribution(s) alone or by combination of their Contribution(s) + with the Work to which such Contribution(s) was submitted. If You + institute patent litigation against any entity (including a + cross-claim or counterclaim in a lawsuit) alleging that the Work + or a Contribution incorporated within the Work constitutes direct + or contributory patent infringement, then any patent licenses + granted to You under this License for that Work shall terminate + as of the date such litigation is filed. + + 4. Redistribution. You may reproduce and distribute copies of the + Work or Derivative Works thereof in any medium, with or without + modifications, and in Source or Object form, provided that You + meet the following conditions: + + (a) You must give any other recipients of the Work or + Derivative Works a copy of this License; and + + (b) You must cause any modified files to carry prominent notices + stating that You changed the files; and + + (c) You must retain, in the Source form of any Derivative Works + that You distribute, all copyright, patent, trademark, and + attribution notices from the Source form of the Work, + excluding those notices that do not pertain to any part of + the Derivative Works; and + + (d) If the Work includes a "NOTICE" text file as part of its + distribution, then any Derivative Works that You distribute must + include a readable copy of the attribution notices contained + within such NOTICE file, excluding those notices that do not + pertain to any part of the Derivative Works, in at least one + of the following places: within a NOTICE text file distributed + as part of the Derivative Works; within the Source form or + documentation, if provided along with the Derivative Works; or, + within a display generated by the Derivative Works, if and + wherever such third-party notices normally appear. The contents + of the NOTICE file are for informational purposes only and + do not modify the License. You may add Your own attribution + notices within Derivative Works that You distribute, alongside + or as an addendum to the NOTICE text from the Work, provided + that such additional attribution notices cannot be construed + as modifying the License. + + You may add Your own copyright statement to Your modifications and + may provide additional or different license terms and conditions + for use, reproduction, or distribution of Your modifications, or + for any such Derivative Works as a whole, provided Your use, + reproduction, and distribution of the Work otherwise complies with + the conditions stated in this License. + + 5. Submission of Contributions. Unless You explicitly state otherwise, + any Contribution intentionally submitted for inclusion in the Work + by You to the Licensor shall be under the terms and conditions of + this License, without any additional terms or conditions. + Notwithstanding the above, nothing herein shall supersede or modify + the terms of any separate license agreement you may have executed + with Licensor regarding such Contributions. + + 6. Trademarks. This License does not grant permission to use the trade + names, trademarks, service marks, or product names of the Licensor, + except as required for reasonable and customary use in describing the + origin of the Work and reproducing the content of the NOTICE file. + + 7. Disclaimer of Warranty. Unless required by applicable law or + agreed to in writing, Licensor provides the Work (and each + Contributor provides its Contributions) on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or + implied, including, without limitation, any warranties or conditions + of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A + PARTICULAR PURPOSE. You are solely responsible for determining the + appropriateness of using or redistributing the Work and assume any + risks associated with Your exercise of permissions under this License. + + 8. Limitation of Liability. In no event and under no legal theory, + whether in tort (including negligence), contract, or otherwise, + unless required by applicable law (such as deliberate and grossly + negligent acts) or agreed to in writing, shall any Contributor be + liable to You for damages, including any direct, indirect, special, + incidental, or consequential damages of any character arising as a + result of this License or out of the use or inability to use the + Work (including but not limited to damages for loss of goodwill, + work stoppage, computer failure or malfunction, or any and all + other commercial damages or losses), even if such Contributor + has been advised of the possibility of such damages. + + 9. Accepting Warranty or Additional Liability. While redistributing + the Work or Derivative Works thereof, You may choose to offer, + and charge a fee for, acceptance of support, warranty, indemnity, + or other liability obligations and/or rights consistent with this + License. However, in accepting such obligations, You may act only + on Your own behalf and on Your sole responsibility, not on behalf + of any other Contributor, and only if You agree to indemnify, + defend, and hold each Contributor harmless for any liability + incurred by, or claims asserted against, such Contributor by reason + of your accepting any such warranty or additional liability. + + END OF TERMS AND CONDITIONS + + APPENDIX: How to apply the Apache License to your work. + + To apply the Apache License to your work, attach the following + boilerplate notice, with the fields enclosed by brackets "[]" + replaced with your own identifying information. (Don't include + the brackets!) The text should be enclosed in the appropriate + comment syntax for the file format. We also recommend that a + file or class name and description of purpose be included on the + same "printed page" as the copyright notice for easier + identification within third-party archives. + + Copyright [yyyy] [name of copyright owner] + + Licensed under the Apache License, Version 2.0 (the "License"); + you may not use this file except in compliance with the License. + You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + + Unless required by applicable law or agreed to in writing, software + distributed under the License is distributed on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. diff --git a/README.md b/README.md index c0d77d5..cab7826 100644 --- a/README.md +++ b/README.md @@ -224,6 +224,7 @@ class MyRender { - Int32 - Float +- Double - Bool - simd_int2 - simd_int3 @@ -234,7 +235,23 @@ class MyRender { - simd_float2x2 - simd_float3x3 - simd_float4x4 -- MTLTexture (only 2d texture) +- simd_double2 +- simd_double3 +- simd_double4 +- simd_double2x2 +- simd_double3x3 +- simd_double4x4 +- MTLTexture (2d) +- MTLTexture (3d) +- IntBuffer (array) +- Int2Buffer (array) +- Int3Buffer (array) +- Int4Buffer (array) +- BoolBuffer (array) +- FloatBuffer (array) +- Float2Buffer (array) +- Float3Buffer (array) +- Float4Buffer (array) ### Supported Platforms - macOS 11.0~ @@ -263,6 +280,25 @@ var customMetalCode: String { } ``` +### Compute Shader for 3D Textures + +```.swift +@EMComputeShader3D +class MyCompute3D { + + @EMTextureArgument(.read, .type3D) + var tex: MTLTexture? + + var impl: String { + "float4 c = tex.read(gid);" + } + + var customMetalCode: String { + "" + } +} +``` + ## Sample Code diff --git a/Sources/EasyMetalShader/Models/ArrayBuffer.swift b/Sources/EasyMetalShader/Models/ArrayBuffer.swift new file mode 100644 index 0000000..2d8f713 --- /dev/null +++ b/Sources/EasyMetalShader/Models/ArrayBuffer.swift @@ -0,0 +1,416 @@ +// +// File.swift +// +// +// Created by Yuki Kuwashima on 2024/08/01. +// + +import MetalKit +import simd + +extension String: LocalizedError { + public var errorDescription: String? { self } +} + +public class BoolBuffer { + + public var buffer: MTLBuffer + public let count: Int + + public init(count: Int) { + self.buffer = ShaderCore.device.makeBuffer(length: MemoryLayout.stride * count)! + self.count = count + } + + @discardableResult + public func setBytes(_ bytes: [Bool]) throws -> Self { + guard bytes.count == count else { + throw "byte count is wrong" + } + self.buffer + .contents() + .copyMemory( + from: bytes, + byteCount: MemoryLayout.stride * count + ) + return self + } + + /// this can be super heavy to call. + public func readBytes() -> [Bool] { + let arrayPointer = buffer.contents().assumingMemoryBound(to: Bool.self) + return Array(UnsafeBufferPointer(start: arrayPointer, count: count)) + } +} + +public class IntBuffer { + + public var buffer: MTLBuffer + public let count: Int + + public init(count: Int) { + self.buffer = ShaderCore.device.makeBuffer(length: MemoryLayout.stride * count)! + self.count = count + } + + @discardableResult + public func setBytes(_ bytes: [Int32]) throws -> Self { + guard bytes.count == count else { + throw "byte count is wrong" + } + self.buffer + .contents() + .copyMemory( + from: bytes, + byteCount: MemoryLayout.stride * count + ) + return self + } + + /// this can be super heavy to call. + public func readBytes() -> [Int32] { + let arrayPointer = buffer.contents().assumingMemoryBound(to: Int32.self) + return Array(UnsafeBufferPointer(start: arrayPointer, count: count)) + } +} + +public class Int2Buffer { + + public var buffer: MTLBuffer + public let count: Int + + public init(count: Int) { + self.buffer = ShaderCore.device.makeBuffer(length: MemoryLayout.stride * count)! + self.count = count + } + + @discardableResult + public func setBytes(_ bytes: [simd_int2]) throws -> Self { + guard bytes.count == count else { + throw "byte count is wrong" + } + self.buffer + .contents() + .copyMemory( + from: bytes, + byteCount: MemoryLayout.stride * count + ) + return self + } + + /// this can be super heavy to call. + public func readBytes() -> [simd_int2] { + let arrayPointer = buffer.contents().assumingMemoryBound(to: simd_int2.self) + return Array(UnsafeBufferPointer(start: arrayPointer, count: count)) + } +} + +public class Int3Buffer { + + public var buffer: MTLBuffer + public let count: Int + + public init(count: Int) { + self.buffer = ShaderCore.device.makeBuffer(length: MemoryLayout.stride * count)! + self.count = count + } + + @discardableResult + public func setBytes(_ bytes: [simd_int3]) throws -> Self { + guard bytes.count == count else { + throw "byte count is wrong" + } + self.buffer + .contents() + .copyMemory( + from: bytes, + byteCount: MemoryLayout.stride * count + ) + return self + } + + /// this can be super heavy to call. + public func readBytes() -> [simd_int3] { + let arrayPointer = buffer.contents().assumingMemoryBound(to: simd_int3.self) + return Array(UnsafeBufferPointer(start: arrayPointer, count: count)) + } +} + +public class Int4Buffer { + + public var buffer: MTLBuffer + public let count: Int + + public init(count: Int) { + self.buffer = ShaderCore.device.makeBuffer(length: MemoryLayout.stride * count)! + self.count = count + } + + @discardableResult + public func setBytes(_ bytes: [simd_int4]) throws -> Self { + guard bytes.count == count else { + throw "byte count is wrong" + } + self.buffer + .contents() + .copyMemory( + from: bytes, + byteCount: MemoryLayout.stride * count + ) + return self + } + + /// this can be super heavy to call. + public func readBytes() -> [simd_int4] { + let arrayPointer = buffer.contents().assumingMemoryBound(to: simd_int4.self) + return Array(UnsafeBufferPointer(start: arrayPointer, count: count)) + } +} + +public class FloatBuffer { + + public var buffer: MTLBuffer + public let count: Int + + public init(count: Int) { + self.buffer = ShaderCore.device.makeBuffer(length: MemoryLayout.stride * count)! + self.count = count + } + + @discardableResult + public func setBytes(_ bytes: [Float]) throws -> Self { + guard bytes.count == count else { + throw "byte count is wrong" + } + self.buffer + .contents() + .copyMemory( + from: bytes, + byteCount: MemoryLayout.stride * count + ) + return self + } + + /// this can be super heavy to call. + public func readBytes() -> [Float] { + let arrayPointer = buffer.contents().assumingMemoryBound(to: Float.self) + return Array(UnsafeBufferPointer(start: arrayPointer, count: count)) + } +} + +public class Float2Buffer { + + public var buffer: MTLBuffer + public let count: Int + + public init(count: Int) { + self.buffer = ShaderCore.device.makeBuffer(length: MemoryLayout.stride * count)! + self.count = count + } + + @discardableResult + public func setBytes(_ bytes: [simd_float2]) throws -> Self { + guard bytes.count == count else { + throw "byte count is wrong" + } + self.buffer + .contents() + .copyMemory( + from: bytes, + byteCount: MemoryLayout.stride * count + ) + return self + } + + /// this can be super heavy to call. + public func readBytes() -> [simd_float2] { + let arrayPointer = buffer.contents().assumingMemoryBound(to: simd_float2.self) + return Array(UnsafeBufferPointer(start: arrayPointer, count: count)) + } +} + +public class Float3Buffer { + + public var buffer: MTLBuffer + public let count: Int + + public init(count: Int) { + self.buffer = ShaderCore.device.makeBuffer(length: MemoryLayout.stride * count)! + self.count = count + } + + @discardableResult + public func setBytes(_ bytes: [simd_float3]) throws -> Self { + guard bytes.count == count else { + throw "byte count is wrong" + } + self.buffer + .contents() + .copyMemory( + from: bytes, + byteCount: MemoryLayout.stride * count + ) + return self + } + + /// this can be super heavy to call. + public func readBytes() -> [simd_float3] { + let arrayPointer = buffer.contents().assumingMemoryBound(to: simd_float3.self) + return Array(UnsafeBufferPointer(start: arrayPointer, count: count)) + } +} + +public class Float4Buffer { + + public var buffer: MTLBuffer + public let count: Int + + public init(count: Int) { + self.buffer = ShaderCore.device.makeBuffer(length: MemoryLayout.stride * count)! + self.count = count + } + + @discardableResult + public func setBytes(_ bytes: [simd_float4]) throws -> Self { + guard bytes.count == count else { + throw "byte count is wrong" + } + self.buffer + .contents() + .copyMemory( + from: bytes, + byteCount: MemoryLayout.stride * count + ) + return self + } + + /// this can be super heavy to call. + public func readBytes() -> [simd_float4] { + let arrayPointer = buffer.contents().assumingMemoryBound(to: simd_float4.self) + return Array(UnsafeBufferPointer(start: arrayPointer, count: count)) + } +} + +public class DoubleBuffer { + + public var buffer: MTLBuffer + public let count: Int + + public init(count: Int) { + self.buffer = ShaderCore.device.makeBuffer(length: MemoryLayout.stride * count)! + self.count = count + } + + @discardableResult + public func setBytes(_ bytes: [Double]) throws -> Self { + guard bytes.count == count else { + throw "byte count is wrong" + } + self.buffer + .contents() + .copyMemory( + from: bytes, + byteCount: MemoryLayout.stride * count + ) + return self + } + + /// this can be super heavy to call. + public func readBytes() -> [Double] { + let arrayPointer = buffer.contents().assumingMemoryBound(to: Double.self) + return Array(UnsafeBufferPointer(start: arrayPointer, count: count)) + } +} + +public class Double2Buffer { + + public var buffer: MTLBuffer + public let count: Int + + public init(count: Int) { + self.buffer = ShaderCore.device.makeBuffer(length: MemoryLayout.stride * count)! + self.count = count + } + + @discardableResult + public func setBytes(_ bytes: [simd_double2]) throws -> Self { + guard bytes.count == count else { + throw "byte count is wrong" + } + self.buffer + .contents() + .copyMemory( + from: bytes, + byteCount: MemoryLayout.stride * count + ) + return self + } + + /// this can be super heavy to call. + public func readBytes() -> [simd_double2] { + let arrayPointer = buffer.contents().assumingMemoryBound(to: simd_double2.self) + return Array(UnsafeBufferPointer(start: arrayPointer, count: count)) + } +} + +public class Double3Buffer { + + public var buffer: MTLBuffer + public let count: Int + + public init(count: Int) { + self.buffer = ShaderCore.device.makeBuffer(length: MemoryLayout.stride * count)! + self.count = count + } + + @discardableResult + public func setBytes(_ bytes: [simd_double3]) throws -> Self { + guard bytes.count == count else { + throw "byte count is wrong" + } + self.buffer + .contents() + .copyMemory( + from: bytes, + byteCount: MemoryLayout.stride * count + ) + return self + } + + /// this can be super heavy to call. + public func readBytes() -> [simd_double3] { + let arrayPointer = buffer.contents().assumingMemoryBound(to: simd_double3.self) + return Array(UnsafeBufferPointer(start: arrayPointer, count: count)) + } +} + +public class Double4Buffer { + + public var buffer: MTLBuffer + public let count: Int + + public init(count: Int) { + self.buffer = ShaderCore.device.makeBuffer(length: MemoryLayout.stride * count)! + self.count = count + } + + @discardableResult + public func setBytes(_ bytes: [simd_double4]) throws -> Self { + guard bytes.count == count else { + throw "byte count is wrong" + } + self.buffer + .contents() + .copyMemory( + from: bytes, + byteCount: MemoryLayout.stride * count + ) + return self + } + + /// this can be super heavy to call. + public func readBytes() -> [simd_double4] { + let arrayPointer = buffer.contents().assumingMemoryBound(to: simd_double4.self) + return Array(UnsafeBufferPointer(start: arrayPointer, count: count)) + } +} diff --git a/Sources/EasyMetalShader/Models/EMMetalArgument.swift b/Sources/EasyMetalShader/Models/EMMetalArgument.swift index 8ea5018..ea0485f 100644 --- a/Sources/EasyMetalShader/Models/EMMetalArgument.swift +++ b/Sources/EasyMetalShader/Models/EMMetalArgument.swift @@ -22,7 +22,27 @@ public enum EMMetalArgument { case float2x2(simd_float2x2) case float3x3(simd_float3x3) case float4x4(simd_float4x4) + case double(Double) + case double2(simd_double2) + case double3(simd_double3) + case double4(simd_double4) + case double2x2(simd_double2x2) + case double3x3(simd_double3x3) + case double4x4(simd_double4x4) case texture2d(MTLTexture!, EMMetalTextureUsage) case texture3d(MTLTexture!, EMMetalTextureUsage) - + case boolBuffer(BoolBuffer) + case intBuffer(IntBuffer) + case int2Buffer(Int2Buffer) + case int3Buffer(Int3Buffer) + case int4Buffer(Int4Buffer) + case floatBuffer(FloatBuffer) + case float2Buffer(Float2Buffer) + case float3Buffer(Float3Buffer) + case float4Buffer(Float4Buffer) + case doubleBuffer(DoubleBuffer) + case double2Buffer(Double2Buffer) + case double3Buffer(Double3Buffer) + case double4Buffer(Double4Buffer) + } diff --git a/Sources/EasyMetalShader/Pipelines/EMMetalComputeFunction.swift b/Sources/EasyMetalShader/Pipelines/EMMetalComputeFunction.swift index a2e5757..c11849e 100644 --- a/Sources/EasyMetalShader/Pipelines/EMMetalComputeFunction.swift +++ b/Sources/EasyMetalShader/Pipelines/EMMetalComputeFunction.swift @@ -72,10 +72,50 @@ extension EMMetalComputeFunction { encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) case .float4x4(let value): encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .double(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .double2(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .double3(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .double4(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .double2x2(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .double3x3(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .double4x4(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) case .texture2d(let value, _): encoder.setTexture(value, index: i+1) case .texture3d(let value, _): encoder.setTexture(value, index: i+1) + case .boolBuffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .intBuffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .int2Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .int3Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .int4Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .floatBuffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .float2Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .float3Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .float4Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .doubleBuffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .double2Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .double3Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .double4Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) case .none: break } @@ -87,4 +127,87 @@ extension EMMetalComputeFunction { threadsPerThreadgroup: size.threadsPerThreadGroup ) } + + public func dispatchArray(_ encoder: MTLComputeCommandEncoder, arrayCount: Int) { + for (i, key) in args.keys.enumerated() { + switch args[key] { + case .bool(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .int(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .int2(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .int3(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .int4(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .float(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .float2(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .float3(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .float4(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .float2x2(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .float3x3(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .float4x4(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .double(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .double2(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .double3(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .double4(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .double2x2(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .double3x3(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .double4x4(let value): + encoder.setBytes([value], length: MemoryLayout.stride, index: i+1) + case .texture2d(let value, _): + encoder.setTexture(value, index: i+1) + case .texture3d(let value, _): + encoder.setTexture(value, index: i+1) + case .boolBuffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .intBuffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .int2Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .int3Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .int4Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .floatBuffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .float2Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .float3Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .float4Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .doubleBuffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .double2Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .double3Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .double4Buffer(let value): + encoder.setBuffer(value.buffer, offset: 0, index: i+1) + case .none: + break + } + } + encoder.setComputePipelineState(computePipelineState) + let size = Self.createDispatchSize(for: computePipelineState, width: arrayCount, height: 1) + encoder.dispatchThreadgroups( + size.threadGroupCount, + threadsPerThreadgroup: size.threadsPerThreadGroup + ) + } } diff --git a/Sources/EasyMetalShader/Pipelines/EMMetalDispatch.swift b/Sources/EasyMetalShader/Pipelines/EMMetalDispatch.swift index b40fb09..7baee88 100644 --- a/Sources/EasyMetalShader/Pipelines/EMMetalDispatch.swift +++ b/Sources/EasyMetalShader/Pipelines/EMMetalDispatch.swift @@ -47,6 +47,9 @@ public class EMMetalDispatch { public func commit() { commandBuffer.commit() + } + + public func waitUntilCompleted() { commandBuffer.waitUntilCompleted() } } diff --git a/Sources/EasyMetalShader/Pipelines/EMMetalRenderFunction.swift b/Sources/EasyMetalShader/Pipelines/EMMetalRenderFunction.swift index 623a199..f2c0f4c 100644 --- a/Sources/EasyMetalShader/Pipelines/EMMetalRenderFunction.swift +++ b/Sources/EasyMetalShader/Pipelines/EMMetalRenderFunction.swift @@ -72,12 +72,72 @@ extension EMMetalRenderFunction { case .float4x4(let value): encoder.setVertexBytes([value], length: MemoryLayout.stride, index: i+1) encoder.setFragmentBytes([value], length: MemoryLayout.stride, index: i+1) + case .double(let value): + encoder.setVertexBytes([value], length: MemoryLayout.stride, index: i+1) + encoder.setFragmentBytes([value], length: MemoryLayout.stride, index: i+1) + case .double2(let value): + encoder.setVertexBytes([value], length: MemoryLayout.stride, index: i+1) + encoder.setFragmentBytes([value], length: MemoryLayout.stride, index: i+1) + case .double3(let value): + encoder.setVertexBytes([value], length: MemoryLayout.stride, index: i+1) + encoder.setFragmentBytes([value], length: MemoryLayout.stride, index: i+1) + case .double4(let value): + encoder.setVertexBytes([value], length: MemoryLayout.stride, index: i+1) + encoder.setFragmentBytes([value], length: MemoryLayout.stride, index: i+1) + case .double2x2(let value): + encoder.setVertexBytes([value], length: MemoryLayout.stride, index: i+1) + encoder.setFragmentBytes([value], length: MemoryLayout.stride, index: i+1) + case .double3x3(let value): + encoder.setVertexBytes([value], length: MemoryLayout.stride, index: i+1) + encoder.setFragmentBytes([value], length: MemoryLayout.stride, index: i+1) + case .double4x4(let value): + encoder.setVertexBytes([value], length: MemoryLayout.stride, index: i+1) + encoder.setFragmentBytes([value], length: MemoryLayout.stride, index: i+1) case .texture2d(let value, _): encoder.setVertexTexture(value, index: i+1) encoder.setFragmentTexture(value, index: i+1) case .texture3d(let value, _): encoder.setVertexTexture(value, index: i+1) encoder.setFragmentTexture(value, index: i+1) + case .boolBuffer(let value): + encoder.setVertexBuffer(value.buffer, offset: 0, index: i+1) + encoder.setFragmentBuffer(value.buffer, offset: 0, index: i+1) + case .intBuffer(let value): + encoder.setVertexBuffer(value.buffer, offset: 0, index: i+1) + encoder.setFragmentBuffer(value.buffer, offset: 0, index: i+1) + case .int2Buffer(let value): + encoder.setVertexBuffer(value.buffer, offset: 0, index: i+1) + encoder.setFragmentBuffer(value.buffer, offset: 0, index: i+1) + case .int3Buffer(let value): + encoder.setVertexBuffer(value.buffer, offset: 0, index: i+1) + encoder.setFragmentBuffer(value.buffer, offset: 0, index: i+1) + case .int4Buffer(let value): + encoder.setVertexBuffer(value.buffer, offset: 0, index: i+1) + encoder.setFragmentBuffer(value.buffer, offset: 0, index: i+1) + case .floatBuffer(let value): + encoder.setVertexBuffer(value.buffer, offset: 0, index: i+1) + encoder.setFragmentBuffer(value.buffer, offset: 0, index: i+1) + case .float2Buffer(let value): + encoder.setVertexBuffer(value.buffer, offset: 0, index: i+1) + encoder.setFragmentBuffer(value.buffer, offset: 0, index: i+1) + case .float3Buffer(let value): + encoder.setVertexBuffer(value.buffer, offset: 0, index: i+1) + encoder.setFragmentBuffer(value.buffer, offset: 0, index: i+1) + case .float4Buffer(let value): + encoder.setVertexBuffer(value.buffer, offset: 0, index: i+1) + encoder.setFragmentBuffer(value.buffer, offset: 0, index: i+1) + case .doubleBuffer(let value): + encoder.setVertexBuffer(value.buffer, offset: 0, index: i+1) + encoder.setFragmentBuffer(value.buffer, offset: 0, index: i+1) + case .double2Buffer(let value): + encoder.setVertexBuffer(value.buffer, offset: 0, index: i+1) + encoder.setFragmentBuffer(value.buffer, offset: 0, index: i+1) + case .double3Buffer(let value): + encoder.setVertexBuffer(value.buffer, offset: 0, index: i+1) + encoder.setFragmentBuffer(value.buffer, offset: 0, index: i+1) + case .double4Buffer(let value): + encoder.setVertexBuffer(value.buffer, offset: 0, index: i+1) + encoder.setFragmentBuffer(value.buffer, offset: 0, index: i+1) case .none: break } diff --git a/Sources/EasyMetalShaderMacro/EMComputeShader/ComputeFunctionStrings/InitCompute.swift b/Sources/EasyMetalShaderMacro/EMComputeShader/ComputeFunctionStrings/InitCompute.swift index eabaa9d..18e9cfe 100644 --- a/Sources/EasyMetalShaderMacro/EMComputeShader/ComputeFunctionStrings/InitCompute.swift +++ b/Sources/EasyMetalShaderMacro/EMComputeShader/ComputeFunctionStrings/InitCompute.swift @@ -50,6 +50,20 @@ public func setup() { functionImpl += "device const float3x3* \\(key)_buf [[buffer(\\(i+1))]]" case .float4x4(_): functionImpl += "device const float4x4* \\(key)_buf [[buffer(\\(i+1))]]" + case .double(_): + functionImpl += "device const double* \\(key)_buf [[buffer(\\(i+1))]]" + case .double2(_): + functionImpl += "device const double2* \\(key)_buf [[buffer(\\(i+1))]]" + case .double3(_): + functionImpl += "device const double3* \\(key)_buf [[buffer(\\(i+1))]]" + case .double4(_): + functionImpl += "device const double4* \\(key)_buf [[buffer(\\(i+1))]]" + case .double2x2(_): + functionImpl += "device const double2x2* \\(key)_buf [[buffer(\\(i+1))]]" + case .double3x3(_): + functionImpl += "device const double3x3* \\(key)_buf [[buffer(\\(i+1))]]" + case .double4x4(_): + functionImpl += "device const double4x4* \\(key)_buf [[buffer(\\(i+1))]]" case .texture2d(_, let usage): switch usage { case .read: @@ -72,6 +86,32 @@ public func setup() { case .sample: functionImpl += "texture3d \\(key) [[texture(\\(i+1))]]" } + case .boolBuffer(_): + functionImpl += "device bool* \\(key) [[buffer(\\(i+1))]]" + case .intBuffer(_): + functionImpl += "device int* \\(key) [[buffer(\\(i+1))]]" + case .int2Buffer(_): + functionImpl += "device int2* \\(key) [[buffer(\\(i+1))]]" + case .int3Buffer(_): + functionImpl += "device int3* \\(key) [[buffer(\\(i+1))]]" + case .int4Buffer(_): + functionImpl += "device int4* \\(key) [[buffer(\\(i+1))]]" + case .floatBuffer(_): + functionImpl += "device float* \\(key) [[buffer(\\(i+1))]]" + case .float2Buffer(_): + functionImpl += "device float2* \\(key) [[buffer(\\(i+1))]]" + case .float3Buffer(_): + functionImpl += "device float3* \\(key) [[buffer(\\(i+1))]]" + case .float4Buffer(_): + functionImpl += "device float4* \\(key) [[buffer(\\(i+1))]]" + case .doubleBuffer(_): + functionImpl += "device double* \\(key) [[buffer(\\(i+1))]]" + case .double2Buffer(_): + functionImpl += "device double2* \\(key) [[buffer(\\(i+1))]]" + case .double3Buffer(_): + functionImpl += "device double3* \\(key) [[buffer(\\(i+1))]]" + case .double4Buffer(_): + functionImpl += "device double4* \\(key) [[buffer(\\(i+1))]]" case .none: break } @@ -106,12 +146,28 @@ public func setup() { functionImpl += "float3x3 \\(key) = \\(key)_buf[0];" case .float4x4(_): functionImpl += "float4x4 \\(key) = \\(key)_buf[0];" + case .double(_): + functionImpl += "double \\(key) = \\(key)_buf[0];" + case .double2(_): + functionImpl += "double2 \\(key) = \\(key)_buf[0];" + case .double3(_): + functionImpl += "double3 \\(key) = \\(key)_buf[0];" + case .double4(_): + functionImpl += "double4 \\(key) = \\(key)_buf[0];" + case .double2x2(_): + functionImpl += "double2x2 \\(key) = \\(key)_buf[0];" + case .double3x3(_): + functionImpl += "double3x3 \\(key) = \\(key)_buf[0];" + case .double4x4(_): + functionImpl += "double4x4 \\(key) = \\(key)_buf[0];" case .texture2d(_, _): break case .texture3d(_, _): break case .none: break + default: + break } } diff --git a/Sources/EasyMetalShaderMacro/EMComputeShader3D/ComputeFunctionStrings3D/InitCompute3D.swift b/Sources/EasyMetalShaderMacro/EMComputeShader3D/ComputeFunctionStrings3D/InitCompute3D.swift index 35fc976..cc1029b 100644 --- a/Sources/EasyMetalShaderMacro/EMComputeShader3D/ComputeFunctionStrings3D/InitCompute3D.swift +++ b/Sources/EasyMetalShaderMacro/EMComputeShader3D/ComputeFunctionStrings3D/InitCompute3D.swift @@ -50,6 +50,20 @@ public func setup() { functionImpl += "device const float3x3* \\(key)_buf [[buffer(\\(i+1))]]" case .float4x4(_): functionImpl += "device const float4x4* \\(key)_buf [[buffer(\\(i+1))]]" + case .double(_): + functionImpl += "device const double* \\(key)_buf [[buffer(\\(i+1))]]" + case .double2(_): + functionImpl += "device const double2* \\(key)_buf [[buffer(\\(i+1))]]" + case .double3(_): + functionImpl += "device const double3* \\(key)_buf [[buffer(\\(i+1))]]" + case .double4(_): + functionImpl += "device const double4* \\(key)_buf [[buffer(\\(i+1))]]" + case .double2x2(_): + functionImpl += "device const double2x2* \\(key)_buf [[buffer(\\(i+1))]]" + case .double3x3(_): + functionImpl += "device const double3x3* \\(key)_buf [[buffer(\\(i+1))]]" + case .double4x4(_): + functionImpl += "device const double4x4* \\(key)_buf [[buffer(\\(i+1))]]" case .texture2d(_, let usage): switch usage { case .read: @@ -72,6 +86,32 @@ public func setup() { case .sample: functionImpl += "texture3d \\(key) [[texture(\\(i+1))]]" } + case .boolBuffer(_): + functionImpl += "device bool* \\(key) [[buffer(\\(i+1))]]" + case .intBuffer(_): + functionImpl += "device int* \\(key) [[buffer(\\(i+1))]]" + case .int2Buffer(_): + functionImpl += "device int2* \\(key) [[buffer(\\(i+1))]]" + case .int3Buffer(_): + functionImpl += "device int3* \\(key) [[buffer(\\(i+1))]]" + case .int4Buffer(_): + functionImpl += "device int4* \\(key) [[buffer(\\(i+1))]]" + case .floatBuffer(_): + functionImpl += "device float* \\(key) [[buffer(\\(i+1))]]" + case .float2Buffer(_): + functionImpl += "device float2* \\(key) [[buffer(\\(i+1))]]" + case .float3Buffer(_): + functionImpl += "device float3* \\(key) [[buffer(\\(i+1))]]" + case .float4Buffer(_): + functionImpl += "device float4* \\(key) [[buffer(\\(i+1))]]" + case .doubleBuffer(_): + functionImpl += "device double* \\(key) [[buffer(\\(i+1))]]" + case .double2Buffer(_): + functionImpl += "device double2* \\(key) [[buffer(\\(i+1))]]" + case .double3Buffer(_): + functionImpl += "device double3* \\(key) [[buffer(\\(i+1))]]" + case .double4Buffer(_): + functionImpl += "device double4* \\(key) [[buffer(\\(i+1))]]" case .none: break } @@ -106,12 +146,28 @@ public func setup() { functionImpl += "float3x3 \\(key) = \\(key)_buf[0];" case .float4x4(_): functionImpl += "float4x4 \\(key) = \\(key)_buf[0];" + case .double(_): + functionImpl += "double \\(key) = \\(key)_buf[0];" + case .double2(_): + functionImpl += "double2 \\(key) = \\(key)_buf[0];" + case .double3(_): + functionImpl += "double3 \\(key) = \\(key)_buf[0];" + case .double4(_): + functionImpl += "double4 \\(key) = \\(key)_buf[0];" + case .double2x2(_): + functionImpl += "double2x2 \\(key) = \\(key)_buf[0];" + case .double3x3(_): + functionImpl += "double3x3 \\(key) = \\(key)_buf[0];" + case .double4x4(_): + functionImpl += "double4x4 \\(key) = \\(key)_buf[0];" case .texture2d(_, _): break case .texture3d(_, _): break case .none: break + default: + break } } diff --git a/Sources/EasyMetalShaderMacro/EMComputeShader3D/EMComputeShader3D+MemberMacro.swift b/Sources/EasyMetalShaderMacro/EMComputeShader3D/EMComputeShader3D+MemberMacro.swift index 48bd979..bb54746 100644 --- a/Sources/EasyMetalShaderMacro/EMComputeShader3D/EMComputeShader3D+MemberMacro.swift +++ b/Sources/EasyMetalShaderMacro/EMComputeShader3D/EMComputeShader3D+MemberMacro.swift @@ -133,7 +133,7 @@ extension EMComputeShader3D: MemberMacro { public var args: [String: EMMetalArgument] = [:] """ - let thisDecl3: DeclSyntax = .init(stringLiteral: ComputeFunctionStrings.initFunc(variableInitStrings: initStringList, gidTypeString: "ushort3")) + let thisDecl3: DeclSyntax = .init(stringLiteral: ComputeFunctionStrings3D.initFunc(variableInitStrings: initStringList, gidTypeString: "ushort3")) var thisDecl4: DeclSyntax = "" diff --git a/Sources/EasyMetalShaderMacro/EMRenderShader/RenderFunctionStrings/InitRender.swift b/Sources/EasyMetalShaderMacro/EMRenderShader/RenderFunctionStrings/InitRender.swift index 1f925a1..b64799c 100644 --- a/Sources/EasyMetalShaderMacro/EMRenderShader/RenderFunctionStrings/InitRender.swift +++ b/Sources/EasyMetalShaderMacro/EMRenderShader/RenderFunctionStrings/InitRender.swift @@ -52,6 +52,20 @@ public func setup(targetPixelFormat: MTLPixelFormat) { functionImpl += "device const float3x3* \\(key)_buf [[buffer(\\(i+1))]]" case .float4x4(_): functionImpl += "device const float4x4* \\(key)_buf [[buffer(\\(i+1))]]" + case .double(_): + functionImpl += "device const double* \\(key)_buf [[buffer(\\(i+1))]]" + case .double2(_): + functionImpl += "device const double2* \\(key)_buf [[buffer(\\(i+1))]]" + case .double3(_): + functionImpl += "device const double3* \\(key)_buf [[buffer(\\(i+1))]]" + case .double4(_): + functionImpl += "device const double4* \\(key)_buf [[buffer(\\(i+1))]]" + case .double2x2(_): + functionImpl += "device const double2x2* \\(key)_buf [[buffer(\\(i+1))]]" + case .double3x3(_): + functionImpl += "device const double3x3* \\(key)_buf [[buffer(\\(i+1))]]" + case .double4x4(_): + functionImpl += "device const double4x4* \\(key)_buf [[buffer(\\(i+1))]]" case .texture2d(_, let usage): switch usage { case .read: @@ -74,6 +88,32 @@ public func setup(targetPixelFormat: MTLPixelFormat) { case .sample: functionImpl += "texture3d \\(key) [[texture(\\(i+1))]]" } + case .boolBuffer(_): + functionImpl += "device bool* \\(key) [[buffer(\\(i+1))]]" + case .intBuffer(_): + functionImpl += "device int* \\(key) [[buffer(\\(i+1))]]" + case .int2Buffer(_): + functionImpl += "device int2* \\(key) [[buffer(\\(i+1))]]" + case .int3Buffer(_): + functionImpl += "device int3* \\(key) [[buffer(\\(i+1))]]" + case .int4Buffer(_): + functionImpl += "device int4* \\(key) [[buffer(\\(i+1))]]" + case .floatBuffer(_): + functionImpl += "device float* \\(key) [[buffer(\\(i+1))]]" + case .float2Buffer(_): + functionImpl += "device float2* \\(key) [[buffer(\\(i+1))]]" + case .float3Buffer(_): + functionImpl += "device float3* \\(key) [[buffer(\\(i+1))]]" + case .float4Buffer(_): + functionImpl += "device float4* \\(key) [[buffer(\\(i+1))]]" + case .doubleBuffer(_): + functionImpl += "device double* \\(key) [[buffer(\\(i+1))]]" + case .double2Buffer(_): + functionImpl += "device double2* \\(key) [[buffer(\\(i+1))]]" + case .double3Buffer(_): + functionImpl += "device double3* \\(key) [[buffer(\\(i+1))]]" + case .double4Buffer(_): + functionImpl += "device double4* \\(key) [[buffer(\\(i+1))]]" case .none: break } @@ -109,12 +149,28 @@ public func setup(targetPixelFormat: MTLPixelFormat) { functionImpl += "float3x3 \\(key) = \\(key)_buf[0];" case .float4x4(_): functionImpl += "float4x4 \\(key) = \\(key)_buf[0];" + case .double(_): + functionImpl += "double \\(key) = \\(key)_buf[0];" + case .double2(_): + functionImpl += "double2 \\(key) = \\(key)_buf[0];" + case .double3(_): + functionImpl += "double3 \\(key) = \\(key)_buf[0];" + case .double4(_): + functionImpl += "double4 \\(key) = \\(key)_buf[0];" + case .double2x2(_): + functionImpl += "double2x2 \\(key) = \\(key)_buf[0];" + case .double3x3(_): + functionImpl += "double3x3 \\(key) = \\(key)_buf[0];" + case .double4x4(_): + functionImpl += "double4x4 \\(key) = \\(key)_buf[0];" case .texture2d(_, _): break case .texture3d(_, _): break case .none: break + default: + break } } @@ -151,6 +207,20 @@ public func setup(targetPixelFormat: MTLPixelFormat) { functionImpl += "device const float3x3* \\(key)_buf [[buffer(\\(i+1))]]" case .float4x4(_): functionImpl += "device const float4x4* \\(key)_buf [[buffer(\\(i+1))]]" + case .double(_): + functionImpl += "device const double* \\(key)_buf [[buffer(\\(i+1))]]" + case .double2(_): + functionImpl += "device const double2* \\(key)_buf [[buffer(\\(i+1))]]" + case .double3(_): + functionImpl += "device const double3* \\(key)_buf [[buffer(\\(i+1))]]" + case .double4(_): + functionImpl += "device const double4* \\(key)_buf [[buffer(\\(i+1))]]" + case .double2x2(_): + functionImpl += "device const double2x2* \\(key)_buf [[buffer(\\(i+1))]]" + case .double3x3(_): + functionImpl += "device const double3x3* \\(key)_buf [[buffer(\\(i+1))]]" + case .double4x4(_): + functionImpl += "device const double4x4* \\(key)_buf [[buffer(\\(i+1))]]" case .texture2d(_, let usage): switch usage { case .read: @@ -173,6 +243,32 @@ public func setup(targetPixelFormat: MTLPixelFormat) { case .sample: functionImpl += "texture3d \\(key) [[texture(\\(i+1))]]" } + case .boolBuffer(_): + functionImpl += "device bool* \\(key) [[buffer(\\(i+1))]]" + case .intBuffer(_): + functionImpl += "device int* \\(key) [[buffer(\\(i+1))]]" + case .int2Buffer(_): + functionImpl += "device int2* \\(key) [[buffer(\\(i+1))]]" + case .int3Buffer(_): + functionImpl += "device int3* \\(key) [[buffer(\\(i+1))]]" + case .int4Buffer(_): + functionImpl += "device int4* \\(key) [[buffer(\\(i+1))]]" + case .floatBuffer(_): + functionImpl += "device float* \\(key) [[buffer(\\(i+1))]]" + case .float2Buffer(_): + functionImpl += "device float2* \\(key) [[buffer(\\(i+1))]]" + case .float3Buffer(_): + functionImpl += "device float3* \\(key) [[buffer(\\(i+1))]]" + case .float4Buffer(_): + functionImpl += "device float4* \\(key) [[buffer(\\(i+1))]]" + case .doubleBuffer(_): + functionImpl += "device double* \\(key) [[buffer(\\(i+1))]]" + case .double2Buffer(_): + functionImpl += "device double2* \\(key) [[buffer(\\(i+1))]]" + case .double3Buffer(_): + functionImpl += "device double3* \\(key) [[buffer(\\(i+1))]]" + case .double4Buffer(_): + functionImpl += "device double4* \\(key) [[buffer(\\(i+1))]]" case .none: break } @@ -209,12 +305,28 @@ public func setup(targetPixelFormat: MTLPixelFormat) { functionImpl += "float3x3 \\(key) = \\(key)_buf[0];" case .float4x4(_): functionImpl += "float4x4 \\(key) = \\(key)_buf[0];" + case .double(_): + functionImpl += "double \\(key) = \\(key)_buf[0];" + case .double2(_): + functionImpl += "double2 \\(key) = \\(key)_buf[0];" + case .double3(_): + functionImpl += "double3 \\(key) = \\(key)_buf[0];" + case .double4(_): + functionImpl += "double4 \\(key) = \\(key)_buf[0];" + case .double2x2(_): + functionImpl += "double2x2 \\(key) = \\(key)_buf[0];" + case .double3x3(_): + functionImpl += "double3x3 \\(key) = \\(key)_buf[0];" + case .double4x4(_): + functionImpl += "double4x4 \\(key) = \\(key)_buf[0];" case .texture2d(_, _): break case .texture3d(_, _): break case .none: break + default: + break } } diff --git a/Sources/EasyMetalShaderMacro/Utils/String+Error.swift b/Sources/EasyMetalShaderMacro/Utils/String+Error.swift index bd83599..3faaffd 100644 --- a/Sources/EasyMetalShaderMacro/Utils/String+Error.swift +++ b/Sources/EasyMetalShaderMacro/Utils/String+Error.swift @@ -7,4 +7,6 @@ import Foundation -extension String: Error {} +extension String: LocalizedError { + public var errorDescription: String? { self } +} diff --git a/Sources/EasyMetalShaderMacro/Utils/Util.swift b/Sources/EasyMetalShaderMacro/Utils/Util.swift index 79d4416..e07c7b0 100644 --- a/Sources/EasyMetalShaderMacro/Utils/Util.swift +++ b/Sources/EasyMetalShaderMacro/Utils/Util.swift @@ -54,7 +54,34 @@ enum Util { return ".float3x3(\(variableName))" case "simd_float4x4": return ".float4x4(\(variableName))" - + + case "BoolBuffer": + return ".boolBuffer(\(variableName))" + case "IntBuffer": + return ".intBuffer(\(variableName))" + case "Int2Buffer": + return ".int2Buffer(\(variableName))" + case "Int3Buffer": + return ".int3Buffer(\(variableName))" + case "Int4Buffer": + return ".int4Buffer(\(variableName))" + case "FloatBuffer": + return ".floatBuffer(\(variableName))" + case "Float2Buffer": + return ".float2Buffer(\(variableName))" + case "Float3Buffer": + return ".float3Buffer(\(variableName))" + case "Float4Buffer": + return ".float4Buffer(\(variableName))" + case "DoubleBuffer": + return ".doubleBuffer(\(variableName))" + case "Double2Buffer": + return ".double2Buffer(\(variableName))" + case "Double3Buffer": + return ".double3Buffer(\(variableName))" + case "Double4Buffer": + return ".double4Buffer(\(variableName))" + default: return nil } diff --git a/Sources/EasyMetalShaderTests/ArrayBufferTests.swift b/Sources/EasyMetalShaderTests/ArrayBufferTests.swift new file mode 100644 index 0000000..15cd25c --- /dev/null +++ b/Sources/EasyMetalShaderTests/ArrayBufferTests.swift @@ -0,0 +1,48 @@ +// +// File.swift +// +// +// Created by Yuki Kuwashima on 2024/08/01. +// + +import XCTest +@testable import EasyMetalShader + +@EMComputeShader +class TestArrayShader { + + var array: BoolBuffer = BoolBuffer(count: 5) + + var impl: String { + "array[0] = true;" + "array[1] = false;" + "array[2] = true;" + "array[3] = false;" + "array[4] = true;" + } + + var customMetalCode: String { + "" + } +} + +class ComputeShaderArrayTests: XCTestCase { + + func testBuild() throws { + + let shader = TestArrayShader() + try shader.array.setBytes([false, false, false, false, false]) + + let dispatch = EMMetalDispatch() + dispatch.compute { encoder in + shader.dispatchArray(encoder, arrayCount: shader.array.count) + } + dispatch.commit() + dispatch.waitUntilCompleted() + + let arrayPointer = shader.array.buffer.contents().assumingMemoryBound(to: Bool.self) + let array = Array(UnsafeBufferPointer(start: arrayPointer, count: shader.array.count)) + + XCTAssertEqual(array, [true, false, true, false, true]) + } +}