From 212f1110da7d969f783f3ab7e3fac58b181a97f5 Mon Sep 17 00:00:00 2001 From: Matthew Dooley Date: Sat, 16 May 2020 00:12:28 -0400 Subject: [PATCH] init --- .../project.pbxproj | 75 +++++++++-- .../contents.xcworkspacedata | 2 +- .../xcshareddata/swiftpm/Package.resolved | 16 +++ .../Closed-form Fibonacci on GPU.xcscheme | 92 ++++++++++++++ .../xcschemes/xcschememanagement.plist | 8 ++ Closed-form Fibonacci on GPU/.gitignore | 5 + .../MetalFibonacci.swift | 116 ++++++++++++++++++ Closed-form Fibonacci on GPU/fibonacci.metal | 45 +++++++ Closed-form Fibonacci on GPU/main.swift | 114 ++++++++++++++++- 9 files changed, 464 insertions(+), 9 deletions(-) create mode 100644 Closed-form Fibonacci on GPU.xcodeproj/project.xcworkspace/xcshareddata/swiftpm/Package.resolved create mode 100644 Closed-form Fibonacci on GPU.xcodeproj/xcshareddata/xcschemes/Closed-form Fibonacci on GPU.xcscheme create mode 100644 Closed-form Fibonacci on GPU/.gitignore create mode 100644 Closed-form Fibonacci on GPU/MetalFibonacci.swift create mode 100644 Closed-form Fibonacci on GPU/fibonacci.metal diff --git a/Closed-form Fibonacci on GPU.xcodeproj/project.pbxproj b/Closed-form Fibonacci on GPU.xcodeproj/project.pbxproj index 8c43f4b..d2ebe23 100644 --- a/Closed-form Fibonacci on GPU.xcodeproj/project.pbxproj +++ b/Closed-form Fibonacci on GPU.xcodeproj/project.pbxproj @@ -3,11 +3,15 @@ archiveVersion = 1; classes = { }; - objectVersion = 50; + objectVersion = 52; objects = { /* Begin PBXBuildFile section */ - 34896A25246D1F0700165BCE /* main.swift in Sources */ = {isa = PBXBuildFile; fileRef = 34896A24246D1F0700165BCE /* main.swift */; }; + 34432619246F65B800317F3C /* MetalFibonacci.swift in Sources */ = {isa = PBXBuildFile; fileRef = 34432618246F65B800317F3C /* MetalFibonacci.swift */; }; + 34896A2D246D1F6F00165BCE /* Metal.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = 34896A2C246D1F6F00165BCE /* Metal.framework */; }; + 34896A30246D216A00165BCE /* ArgumentParser in Frameworks */ = {isa = PBXBuildFile; productRef = 34896A2F246D216A00165BCE /* ArgumentParser */; }; + 34896A35246D2BF800165BCE /* main.swift in Sources */ = {isa = PBXBuildFile; fileRef = 34896A34246D2BF800165BCE /* main.swift */; }; + 34896A37246D3F3900165BCE /* fibonacci.metal in Sources */ = {isa = PBXBuildFile; fileRef = 34896A36246D3F3900165BCE /* fibonacci.metal */; }; /* End PBXBuildFile section */ /* Begin PBXCopyFilesBuildPhase section */ @@ -23,8 +27,11 @@ /* End PBXCopyFilesBuildPhase section */ /* Begin PBXFileReference section */ + 34432618246F65B800317F3C /* MetalFibonacci.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MetalFibonacci.swift; sourceTree = ""; }; 34896A21246D1F0700165BCE /* Closed-form Fibonacci on GPU */ = {isa = PBXFileReference; explicitFileType = "compiled.mach-o.executable"; includeInIndex = 0; path = "Closed-form Fibonacci on GPU"; sourceTree = BUILT_PRODUCTS_DIR; }; - 34896A24246D1F0700165BCE /* main.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = main.swift; sourceTree = ""; }; + 34896A2C246D1F6F00165BCE /* Metal.framework */ = {isa = PBXFileReference; lastKnownFileType = wrapper.framework; name = Metal.framework; path = System/Library/Frameworks/Metal.framework; sourceTree = SDKROOT; }; + 34896A34246D2BF800165BCE /* main.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = main.swift; sourceTree = ""; }; + 34896A36246D3F3900165BCE /* fibonacci.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = fibonacci.metal; sourceTree = ""; }; /* End PBXFileReference section */ /* Begin PBXFrameworksBuildPhase section */ @@ -32,6 +39,8 @@ isa = PBXFrameworksBuildPhase; buildActionMask = 2147483647; files = ( + 34896A30246D216A00165BCE /* ArgumentParser in Frameworks */, + 34896A2D246D1F6F00165BCE /* Metal.framework in Frameworks */, ); runOnlyForDeploymentPostprocessing = 0; }; @@ -43,6 +52,7 @@ children = ( 34896A23246D1F0700165BCE /* Closed-form Fibonacci on GPU */, 34896A22246D1F0700165BCE /* Products */, + 34896A2B246D1F6F00165BCE /* Frameworks */, ); sourceTree = ""; }; @@ -57,11 +67,21 @@ 34896A23246D1F0700165BCE /* Closed-form Fibonacci on GPU */ = { isa = PBXGroup; children = ( - 34896A24246D1F0700165BCE /* main.swift */, + 34896A34246D2BF800165BCE /* main.swift */, + 34896A36246D3F3900165BCE /* fibonacci.metal */, + 34432618246F65B800317F3C /* MetalFibonacci.swift */, ); path = "Closed-form Fibonacci on GPU"; sourceTree = ""; }; + 34896A2B246D1F6F00165BCE /* Frameworks */ = { + isa = PBXGroup; + children = ( + 34896A2C246D1F6F00165BCE /* Metal.framework */, + ); + name = Frameworks; + sourceTree = ""; + }; /* End PBXGroup section */ /* Begin PBXNativeTarget section */ @@ -78,6 +98,9 @@ dependencies = ( ); name = "Closed-form Fibonacci on GPU"; + packageProductDependencies = ( + 34896A2F246D216A00165BCE /* ArgumentParser */, + ); productName = "Closed-form Fibonacci on GPU"; productReference = 34896A21246D1F0700165BCE /* Closed-form Fibonacci on GPU */; productType = "com.apple.product-type.tool"; @@ -94,6 +117,7 @@ TargetAttributes = { 34896A20246D1F0700165BCE = { CreatedOnToolsVersion = 11.5; + LastSwiftMigration = 1150; }; }; }; @@ -106,6 +130,9 @@ Base, ); mainGroup = 34896A18246D1F0700165BCE; + packageReferences = ( + 34896A2E246D216A00165BCE /* XCRemoteSwiftPackageReference "swift-argument-parser" */, + ); productRefGroup = 34896A22246D1F0700165BCE /* Products */; projectDirPath = ""; projectRoot = ""; @@ -120,7 +147,9 @@ isa = PBXSourcesBuildPhase; buildActionMask = 2147483647; files = ( - 34896A25246D1F0700165BCE /* main.swift in Sources */, + 34896A37246D3F3900165BCE /* fibonacci.metal in Sources */, + 34432619246F65B800317F3C /* MetalFibonacci.swift in Sources */, + 34896A35246D2BF800165BCE /* main.swift in Sources */, ); runOnlyForDeploymentPostprocessing = 0; }; @@ -179,7 +208,7 @@ GCC_WARN_UNUSED_VARIABLE = YES; MACOSX_DEPLOYMENT_TARGET = 10.15; MTL_ENABLE_DEBUG_INFO = INCLUDE_SOURCE; - MTL_FAST_MATH = YES; + MTL_FAST_MATH = NO; ONLY_ACTIVE_ARCH = YES; SDKROOT = macosx; SWIFT_ACTIVE_COMPILATION_CONDITIONS = DEBUG; @@ -233,7 +262,7 @@ GCC_WARN_UNUSED_VARIABLE = YES; MACOSX_DEPLOYMENT_TARGET = 10.15; MTL_ENABLE_DEBUG_INFO = NO; - MTL_FAST_MATH = YES; + MTL_FAST_MATH = NO; SDKROOT = macosx; SWIFT_COMPILATION_MODE = wholemodule; SWIFT_OPTIMIZATION_LEVEL = "-O"; @@ -243,10 +272,17 @@ 34896A29246D1F0700165BCE /* Debug */ = { isa = XCBuildConfiguration; buildSettings = { + CLANG_ENABLE_MODULES = YES; CODE_SIGN_STYLE = Automatic; DEVELOPMENT_TEAM = RUZM7ULTJ2; ENABLE_HARDENED_RUNTIME = YES; + LD_RUNPATH_SEARCH_PATHS = ( + "$(inherited)", + "@executable_path/../Frameworks", + "@loader_path/../Frameworks", + ); PRODUCT_NAME = "$(TARGET_NAME)"; + SWIFT_OPTIMIZATION_LEVEL = "-Onone"; SWIFT_VERSION = 5.0; }; name = Debug; @@ -254,9 +290,15 @@ 34896A2A246D1F0700165BCE /* Release */ = { isa = XCBuildConfiguration; buildSettings = { + CLANG_ENABLE_MODULES = YES; CODE_SIGN_STYLE = Automatic; DEVELOPMENT_TEAM = RUZM7ULTJ2; ENABLE_HARDENED_RUNTIME = YES; + LD_RUNPATH_SEARCH_PATHS = ( + "$(inherited)", + "@executable_path/../Frameworks", + "@loader_path/../Frameworks", + ); PRODUCT_NAME = "$(TARGET_NAME)"; SWIFT_VERSION = 5.0; }; @@ -284,6 +326,25 @@ defaultConfigurationName = Release; }; /* End XCConfigurationList section */ + +/* Begin XCRemoteSwiftPackageReference section */ + 34896A2E246D216A00165BCE /* XCRemoteSwiftPackageReference "swift-argument-parser" */ = { + isa = XCRemoteSwiftPackageReference; + repositoryURL = "https://github.com/apple/swift-argument-parser"; + requirement = { + kind = upToNextMinorVersion; + minimumVersion = 0.0.1; + }; + }; +/* End XCRemoteSwiftPackageReference section */ + +/* Begin XCSwiftPackageProductDependency section */ + 34896A2F246D216A00165BCE /* ArgumentParser */ = { + isa = XCSwiftPackageProductDependency; + package = 34896A2E246D216A00165BCE /* XCRemoteSwiftPackageReference "swift-argument-parser" */; + productName = ArgumentParser; + }; +/* End XCSwiftPackageProductDependency section */ }; rootObject = 34896A19246D1F0700165BCE /* Project object */; } diff --git a/Closed-form Fibonacci on GPU.xcodeproj/project.xcworkspace/contents.xcworkspacedata b/Closed-form Fibonacci on GPU.xcodeproj/project.xcworkspace/contents.xcworkspacedata index 042db99..919434a 100644 --- a/Closed-form Fibonacci on GPU.xcodeproj/project.xcworkspace/contents.xcworkspacedata +++ b/Closed-form Fibonacci on GPU.xcodeproj/project.xcworkspace/contents.xcworkspacedata @@ -2,6 +2,6 @@ + location = "self:"> diff --git a/Closed-form Fibonacci on GPU.xcodeproj/project.xcworkspace/xcshareddata/swiftpm/Package.resolved b/Closed-form Fibonacci on GPU.xcodeproj/project.xcworkspace/xcshareddata/swiftpm/Package.resolved new file mode 100644 index 0000000..61a5ad5 --- /dev/null +++ b/Closed-form Fibonacci on GPU.xcodeproj/project.xcworkspace/xcshareddata/swiftpm/Package.resolved @@ -0,0 +1,16 @@ +{ + "object": { + "pins": [ + { + "package": "swift-argument-parser", + "repositoryURL": "https://github.com/apple/swift-argument-parser", + "state": { + "branch": null, + "revision": "9f04d1ff1afbccd02279338a2c91e5f27c45e93a", + "version": "0.0.5" + } + } + ] + }, + "version": 1 +} diff --git a/Closed-form Fibonacci on GPU.xcodeproj/xcshareddata/xcschemes/Closed-form Fibonacci on GPU.xcscheme b/Closed-form Fibonacci on GPU.xcodeproj/xcshareddata/xcschemes/Closed-form Fibonacci on GPU.xcscheme new file mode 100644 index 0000000..7018a48 --- /dev/null +++ b/Closed-form Fibonacci on GPU.xcodeproj/xcshareddata/xcschemes/Closed-form Fibonacci on GPU.xcscheme @@ -0,0 +1,92 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/Closed-form Fibonacci on GPU.xcodeproj/xcuserdata/Matthew.xcuserdatad/xcschemes/xcschememanagement.plist b/Closed-form Fibonacci on GPU.xcodeproj/xcuserdata/Matthew.xcuserdatad/xcschemes/xcschememanagement.plist index 73f0c23..6bfcdc6 100644 --- a/Closed-form Fibonacci on GPU.xcodeproj/xcuserdata/Matthew.xcuserdatad/xcschemes/xcschememanagement.plist +++ b/Closed-form Fibonacci on GPU.xcodeproj/xcuserdata/Matthew.xcuserdatad/xcschemes/xcschememanagement.plist @@ -10,5 +10,13 @@ 0 + SuppressBuildableAutocreation + + 34896A20246D1F0700165BCE + + primary + + + diff --git a/Closed-form Fibonacci on GPU/.gitignore b/Closed-form Fibonacci on GPU/.gitignore new file mode 100644 index 0000000..95c4320 --- /dev/null +++ b/Closed-form Fibonacci on GPU/.gitignore @@ -0,0 +1,5 @@ +.DS_Store +/.build +/Packages +/*.xcodeproj +xcuserdata/ diff --git a/Closed-form Fibonacci on GPU/MetalFibonacci.swift b/Closed-form Fibonacci on GPU/MetalFibonacci.swift new file mode 100644 index 0000000..bb5d20e --- /dev/null +++ b/Closed-form Fibonacci on GPU/MetalFibonacci.swift @@ -0,0 +1,116 @@ +// +// MetalFibonacci.swift +// Closed-form Fibonacci on GPU +// +// Created by Matthew Dooley on 5/15/20. +// Copyright © 2020 Matthew Dooley. All rights reserved. +// + +import Foundation +import Metal + +class MetalFibonacci +{ + let start: UInt32 + let end: Int + + let length: Int + let intBufferSize: Int = MemoryLayout.stride + let resultBufferSize: Int + + let device: MTLDevice + let fibonacciFunctionPSO: MTLComputePipelineState + let commandQueue: MTLCommandQueue + + var startBuffer: MTLBuffer? + var resultBuffer: MTLBuffer? + + init?(device: MTLDevice, start: Int, end: Int) + { + self.device = device + self.start = UInt32(start) + self.end = end + self.length = end - start + 1 // convert from 0 based counting + self.resultBufferSize = length * MemoryLayout.stride + + let defaultLibrary = self.device.makeDefaultLibrary() + if (defaultLibrary == nil) { + NSLog("Could not find library.") + return nil + } + + let fibonacciFunction = defaultLibrary!.makeFunction(name: "fibonacci") + if (fibonacciFunction == nil) { + NSLog("Could not find function.") + return nil + } + + do { + try self.fibonacciFunctionPSO = self.device.makeComputePipelineState(function: fibonacciFunction!) + } catch { + NSLog("Could not create pipeline") + return nil + } + + self.commandQueue = self.device.makeCommandQueue()! + } + + func prepareData() + { + self.startBuffer = + self.device.makeBuffer(length: self.intBufferSize, + options: MTLResourceOptions.storageModeShared)! + + self.resultBuffer = + self.device.makeBuffer(length: self.resultBufferSize, + options: MTLResourceOptions.storageModeShared)! + let start: UnsafeMutablePointer = + self.startBuffer!.contents().assumingMemoryBound(to: UInt32.self) + + start.pointee = self.start + } + + func encodeFibonacciCommand(_ computeEncoder: MTLComputeCommandEncoder) + { + // Set the work + computeEncoder.setComputePipelineState(self.fibonacciFunctionPSO) + computeEncoder.setBuffer(self.startBuffer, offset: 0, index: 0) + computeEncoder.setBuffer(self.resultBuffer, offset: 0, index: 1) + + let gridSize: MTLSize = MTLSizeMake(self.length, 1, 1) + + // Calculate a threadgroup size. + var threadGroupSizeInt: Int = self.fibonacciFunctionPSO.maxTotalThreadsPerThreadgroup + if (threadGroupSizeInt > self.length) + { + threadGroupSizeInt = self.length + } + let threadGroupSize: MTLSize = MTLSizeMake(threadGroupSizeInt, 1, 1) + + // Encode the compute command. + computeEncoder.dispatchThreads(gridSize, threadsPerThreadgroup: threadGroupSize) + } + + func sendComputeCommand() + { + let commandBuffer: MTLCommandBuffer = self.commandQueue.makeCommandBuffer()! + let computeEncoder: MTLComputeCommandEncoder = commandBuffer.makeComputeCommandEncoder()! + + self.encodeFibonacciCommand(computeEncoder) + + computeEncoder.endEncoding() + commandBuffer.commit() + + commandBuffer.waitUntilCompleted() + } + + func printResults() + { + let result: UnsafeMutablePointer = self.resultBuffer!.contents().assumingMemoryBound(to: UInt32.self) + + for index in 0 ..< self.length + { + print("\(self.start + UInt32(index)): \(result[index])") + } + } +} diff --git a/Closed-form Fibonacci on GPU/fibonacci.metal b/Closed-form Fibonacci on GPU/fibonacci.metal new file mode 100644 index 0000000..e779ba8 --- /dev/null +++ b/Closed-form Fibonacci on GPU/fibonacci.metal @@ -0,0 +1,45 @@ +// +// fibonacci.metal +// Closed-form Fibonacci on GPU +// +// Created by Matthew Dooley on 5/14/20. +// Copyright © 2020 Matthew Dooley. All rights reserved. +// + +#include + +using namespace metal; + +// declare methods +kernel void fibonacci(device unsigned int* start, + device unsigned int* result, + unsigned int index); +unsigned int fibonacci(unsigned int index); +unsigned int binetsFormula(unsigned int n); + +// initialize methods +kernel void fibonacci(device uint* start, + device uint* result, + unsigned int index [[thread_position_in_grid]]) +{ + result[index] = fibonacci(*start + index); +} + +unsigned int fibonacci(unsigned int index) +{ + return binetsFormula(index); + + unsigned int n1 = binetsFormula(index - 1); + unsigned int n2 = binetsFormula(index - 2); + + return n1 + n2; +} + +unsigned int binetsFormula(unsigned int n) +{ + const float sqrt5 = sqrt(5.0); + const float goldenRatio = (1 + sqrt5) / 2; + const float psi = (1 - goldenRatio); + + return (unsigned int)((pow(goldenRatio, n) - pow(psi, n)) / sqrt5); +} diff --git a/Closed-form Fibonacci on GPU/main.swift b/Closed-form Fibonacci on GPU/main.swift index 719ade0..8e8c006 100644 --- a/Closed-form Fibonacci on GPU/main.swift +++ b/Closed-form Fibonacci on GPU/main.swift @@ -5,8 +5,120 @@ // Created by Matthew Dooley on 5/14/20. // Copyright © 2020 Matthew Dooley. All rights reserved. // +// CPU accurate (0, 70] then datatypes start to become annoying +// GPU accurate (0, 32] then datatypes start to become annoying +// import Foundation +import ArgumentParser +import Metal +import Darwin + +// get command line arguments +struct FibonacciOptions: ParsableArguments { + @Flag(help: "Should we use the CPU instead?") + var cpu: Bool + + @Option(name: .shortAndLong, help: "The start position for getting a range.") + var start: Int? + + @Argument(help: "The position to get (or end at).") + var index: Int +} + +// force correct datatypes and any computed values +struct Config { + static let options = FibonacciOptions.parseOrExit() + static let cpu: Bool = options.cpu + static let range: Bool = (options.start != nil) + static let start: Int? = options.start + static let end: Int = options.index + static let index: Int = options.index +} + +if (Config.index < 0) { + print("Error: '' must be >= 0. \(Config.index) < 0") + Darwin.exit(1) +} else if (Config.range && Config.start! < 0) { + print("Error: '' must be >= 0. \(Config.start!) < 0") + Darwin.exit(1) +} else if (Config.range && Config.start! >= Config.index) { + print("Error: '' must be < ''. \(Config.start!) >= \(Config.index)") + Darwin.exit(1) +} + +extension Double { + func toString(decimalPrecision decimal: Int = 9) -> String { + let value = decimal < 0 ? 0 : decimal + var string = String(format: "%.\(value)f", self) + + while string.last == "." { + if string.last == "." { string = String(string.dropLast()); break} + string = String(string.dropLast()) + } + return string + } +} + +func fibonacciRange(_ start: Int, to end: Int) -> [Double] +{ + var output: [Double] = [] + + for i in start...end + { + output.append(fibonacciCPU(i)) + } + + return output +} + +func fibonacciCPU(_ n: Int) -> Double +{ + let n = Double(n) + let n1 = n - 1 + let n2 = n1 - 1 + + let resultn1 = binetsFormula(n1) + let resultn2 = binetsFormula(n2) + + let result = resultn1 + resultn2 + return result +} + +func binetsFormula(_ nth: Double) -> Double +{ + let goldenRatio = ((1 + Double(5).squareRoot()) / 2) + let psi = (1 - goldenRatio) -print("Hello, World!") + return Double( + (pow(goldenRatio, nth) - pow(psi, nth)) + / (Double(5).squareRoot())) +} +if Config.cpu { + // accurate up to and including 70 + print("Calculating with CPU") + + if !Config.range { + print("\(Config.index): \(fibonacciCPU(Config.index).toString(decimalPrecision: 0))") + } else { + var current: Int = Config.start! + for i in fibonacciRange(Config.start!, to: Config.index) + { + print("\(current): \(i.toString(decimalPrecision: 0))") + current += 1 + } + } +} else { + // accurate up to and including 32 + print("Calculating with GPU") + + let device: MTLDevice = MTLCreateSystemDefaultDevice()! + let fibonacci: MetalFibonacci = MetalFibonacci(device: device, + start: Config.start ?? Config.end, + end: Config.end)! + + fibonacci.prepareData() + fibonacci.sendComputeCommand() + fibonacci.printResults() +}