Skip to content

HTTPS clone URL

Subversion checkout URL

You can clone with
or
.
Download ZIP
Browse files

Screwed up with Git big time, had to redo repo.

Fixed several problems with RPC. Version bump to 1.6.3.

Signed-off-by: jedi95 <jedi95@gmail.com>
  • Loading branch information...
commit 30f838629811dfce0dc1ad0cd04e7a8ca3362a5c 0 parents
@jedi95 authored
2  .gitignore
@@ -0,0 +1,2 @@
+*.pyc
+*.py~
147 ConsoleLogger.py
@@ -0,0 +1,147 @@
+# Copyright (C) 2011 by jedi95 <jedi95@gmail.com> and
+# CFSworks <CFSworks@gmail.com>
+#
+# 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.
+
+import sys
+from time import time
+from datetime import datetime
+
+def formatNumber(n):
+ """Format a positive integer in a more readable fashion."""
+ if n < 0:
+ raise ValueError('can only format positive integers')
+ prefixes = 'KMGTP'
+ whole = str(int(n))
+ decimal = ''
+ i = 0
+ while len(whole) > 3:
+ if i + 1 < len(prefixes):
+ decimal = '.%s' % whole[-3:-1]
+ whole = whole[:-3]
+ i += 1
+ else:
+ break
+ return '%s%s %s' % (whole, decimal, prefixes[i])
+
+class ConsoleLogger(object):
+ """This class will handle printing messages to the console."""
+
+ TIME_FORMAT = '[%d/%m/%Y %H:%M:%S]'
+
+ UPDATE_TIME = 1.0
+
+ def __init__(self, miner, verbose=False):
+ self.verbose = verbose
+ self.miner = miner
+ self.lastUpdate = time() - 1
+ self.rate = 0
+ self.accepted = 0
+ self.invalid = 0
+ self.lineLength = 0
+ self.connectionType = None
+
+ def reportRate(self, rate, update=True):
+ """Used to tell the logger the current Khash/sec."""
+ self.rate = rate
+ if update:
+ self.updateStatus()
+
+ def reportType(self, type):
+ self.connectionType = type
+
+ def reportBlock(self, block):
+ self.log('Currently on block: ' + str(block))
+
+ def reportFound(self, hash, accepted):
+ if accepted:
+ self.accepted += 1
+ else:
+ self.invalid += 1
+
+ hexHash = hash[::-1]
+ hexHash = hexHash[:8].encode('hex')
+ if self.verbose:
+ self.log('Result %s... %s' % (hexHash,
+ 'accepted' if accepted else 'rejected'))
+ else:
+ self.log('Result: %s %s' % (hexHash[8:],
+ 'accepted' if accepted else 'rejected'))
+
+ def reportMsg(self, message):
+ self.log(('MSG: ' + message), True, True)
+
+ def reportConnected(self, connected):
+ if connected:
+ self.log('Connected to server')
+ else:
+ self.log('Disconnected from server')
+
+ def reportConnectionFailed(self):
+ self.log('Failed to connect, retrying...')
+
+ def reportDebug(self, message):
+ if self.verbose:
+ self.log(message)
+
+ def updateStatus(self, force=False):
+ #only update if last update was more than a second ago
+ dt = time() - self.lastUpdate
+ if force or dt > self.UPDATE_TIME:
+ rate = self.rate if (not self.miner.idle) else 0
+ type = " [" + str(self.connectionType) + "]" if self.connectionType is not None else ''
+ status = (
+ "[" + formatNumber(rate) + "hash/sec] "
+ "[" + str(self.accepted) + " Accepted] "
+ "[" + str(self.invalid) + " Rejected]" + type)
+ self.say(status)
+ self.lastUpdate = time()
+
+ def say(self, message, newLine=False, hideTimestamp=False):
+ #add new line if requested
+ if newLine:
+ message += '\n'
+ if hideTimestamp:
+ timestamp = ''
+ else:
+ timestamp = datetime.now().strftime(self.TIME_FORMAT) + ' '
+
+ message = timestamp + message
+
+ #erase the previous line
+ if self.lineLength > 0:
+ sys.stdout.write('\b \b' * self.lineLength)
+ sys.stdout.write(' ' * self.lineLength)
+ sys.stdout.write('\b \b' * self.lineLength)
+
+ #print the line
+ sys.stdout.write(message)
+ sys.stdout.flush()
+
+ #cache the current line length
+ if newLine:
+ self.lineLength = 0
+ else:
+ self.lineLength = len(message)
+
+ def log(self, message, update=True, hideTimestamp=False):
+ self.say(message, True, hideTimestamp)
+ if update:
+ self.updateStatus(True)
+
248 KernelInterface.py
@@ -0,0 +1,248 @@
+# Copyright (C) 2011 by jedi95 <jedi95@gmail.com> and
+# CFSworks <CFSworks@gmail.com>
+#
+# 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.
+
+import os
+from struct import pack, unpack
+from hashlib import sha256
+from twisted.internet import defer, reactor
+
+# I'm using this as a sentinel value to indicate that an option has no default;
+# it must be specified.
+REQUIRED = object()
+
+class KernelOption(object):
+ """This works like a property, and is used in defining easy option tables
+ for kernels.
+ """
+
+ def __init__(self, name, type, help=None, default=REQUIRED,
+ advanced=False, **kwargs):
+ self.localValues = {}
+ self.name = name
+ self.type = type
+ self.help = help
+ self.default = default
+ self.advanced = advanced
+
+ def __get__(self, instance, owner):
+ if instance in self.localValues:
+ return self.localValues[instance]
+ else:
+ return instance.interface._getOption(
+ self.name, self.type, self.default)
+
+ def __set__(self, instance, value):
+ self.localValues[instance] = value
+
+class CoreInterface(object):
+ """An internal class provided for kernels to use when reporting info for
+ one core.
+
+ Only KernelInterface should create this.
+ """
+
+ def __init__(self, kernelInterface):
+ self.kernelInterface = kernelInterface
+ self.averageSamples = []
+ self.kernelInterface.miner._addCore(self)
+
+ def updateRate(self, rate):
+ """Called by a kernel core to report its current rate."""
+
+ numSamples = self.kernelInterface.miner.options.getAvgSamples()
+
+ self.averageSamples.append(rate)
+ self.averageSamples = self.averageSamples[-numSamples:]
+
+ self.kernelInterface.miner.updateAverage()
+
+ def getRate(self):
+ """Retrieve the average rate for this core."""
+
+ if not self.averageSamples:
+ return 0
+
+ return sum(self.averageSamples)/len(self.averageSamples)
+
+ def getKernelInterface(self):
+ return self.kernelInterface
+
+class KernelInterface(object):
+ """This is an object passed to kernels as an API back to the Phoenix
+ framework.
+ """
+
+ def __init__(self, miner):
+ self.miner = miner
+ self._core = None
+
+ def _getOption(self, name, type, default):
+ """KernelOption uses this to read the actual value of the option."""
+ if not name in self.miner.options.kernelOptions:
+ if default == REQUIRED:
+ self.fatal('Required option %s not provided!' % name)
+ else:
+ return default
+
+ givenOption = self.miner.options.kernelOptions[name]
+ if type == bool:
+ # The following are considered true
+ return givenOption is None or \
+ givenOption.lower() in ('t', 'true', 'on', '1', 'y', 'yes')
+
+ try:
+ return type(givenOption)
+ except (TypeError, ValueError):
+ self.fatal('Option %s expects a value of type %s!' % (name, type))
+
+ def getRevision(self):
+ """Return the Phoenix core revision, so that kernels can require a
+ minimum revision before operating (such as if they rely on a certain
+ feature added in a certain revision)
+ """
+
+ return self.miner.REVISION
+
+ def setWorkFactor(self, workFactor):
+ """Deprecated. Kernels are now responsible for requesting optimal size
+ work"""
+
+ def setMeta(self, var, value):
+ """Set metadata for this kernel."""
+
+ self.miner.connection.setMeta(var, value)
+
+ def fetchRange(self, size=None):
+ """Fetch a range from the WorkQueue, optionally specifying a size
+ (in nonces) to include in the range.
+ """
+
+ if size is None:
+ return self.miner.queue.fetchRange()
+ else:
+ return self.miner.queue.fetchRange(size)
+
+ def addStaleCallback(self, callback):
+ """Register a new function to be called, with no arguments, whenever
+ a new block comes out that would render all previous work stale,
+ requiring a kernel to switch immediately.
+ """
+
+ # This should be implemented in a better way in the future...
+ if callback not in self.miner.queue.staleCallbacks:
+ self.miner.queue.staleCallbacks.append(callback)
+
+ def removeStaleCallback(self, callback):
+ """Undo an addStaleCallback."""
+
+ # Likewise.
+ if callback in self.miner.queue.staleCallbacks:
+ self.miner.queue.staleCallbacks.remove(callback)
+
+ def addCore(self):
+ """Return a CoreInterface for a new core."""
+ return CoreInterface(self)
+
+ def checkTarget(self, hash, target):
+ """Utility function that the kernel can use to see if a nonce meets a
+ target before sending it back to the core.
+
+ Since the target is checked before submission anyway, this is mostly
+ intended to be used in hardware sanity-checks.
+ """
+
+ # This for loop compares the bytes of the target and hash in reverse
+ # order, because both are 256-bit little endian.
+ for t,h in zip(target[::-1], hash[::-1]):
+ if ord(t) > ord(h):
+ return True
+ elif ord(t) < ord(h):
+ return False
+ return True
+
+ def calculateHash(self, nr, nonce):
+ """Given a NonceRange and a nonce, calculate the SHA-256 hash of the
+ solution. The resulting hash is returned as a string, which may be
+ compared with the target as a 256-bit little endian unsigned integer.
+ """
+ # Sometimes kernels send weird nonces down the pipe. We can assume they
+ # accidentally set bits outside of the 32-bit space. If the resulting
+ # nonce is invalid, it will be caught anyway...
+ nonce &= 0xFFFFFFFF
+
+ staticDataUnpacked = unpack('<' + 'I'*19, nr.unit.data[:76])
+ staticData = pack('>' + 'I'*19, *staticDataUnpacked)
+ hashInput = pack('>76sI', staticData, nonce)
+ return sha256(sha256(hashInput).digest()).digest()
+
+ def foundNonce(self, nr, nonce):
+ """Called by kernels when they may have found a nonce."""
+
+ # Sometimes kernels send weird nonces down the pipe. We can assume they
+ # accidentally set bits outside of the 32-bit space. If the resulting
+ # nonce is invalid, it will be caught anyway...
+ nonce &= 0xFFFFFFFF
+
+ # Check if the block has changed while this NonceRange was being
+ # processed by the kernel. If so, don't send it to the server.
+ if self.miner.queue.isRangeStale(nr):
+ return False
+
+ # Check if the hash meets the full difficulty before sending.
+ hash = self.calculateHash(nr, nonce)
+
+ if self.checkTarget(hash, nr.unit.target):
+ formattedResult = pack('<76sI', nr.unit.data[:76], nonce)
+ d = self.miner.connection.sendResult(formattedResult)
+ def callback(accepted):
+ self.miner.logger.reportFound(hash, accepted)
+ d.addCallback(callback)
+ return True
+ else:
+ self.miner.logger.reportDebug("Result didn't meet full "
+ "difficulty, not sending")
+ return False
+
+ def debug(self, msg):
+ """Log information as debug so that it can be viewed only when -v is
+ enabled.
+ """
+ self.miner.logger.reportDebug(msg)
+
+ def log(self, msg, withTimestamp=True, withIdentifier=True):
+ """Log some general kernel information to the console."""
+ self.miner.logger.log(msg, True, not withTimestamp)
+
+ def error(self, msg=None):
+ """The kernel has an issue that requires user attention."""
+ if msg is not None:
+ self.miner.logger.log('Kernel error: ' + msg)
+
+ def fatal(self, msg=None):
+ """The kernel has an issue that is preventing it from continuing to
+ operate.
+ """
+ if msg is not None:
+ self.miner.logger.log('FATAL kernel error: ' + msg, False)
+ if reactor.running:
+ reactor.stop()
+ os._exit(0)
+
199 Miner.py
@@ -0,0 +1,199 @@
+# Copyright (C) 2011 by jedi95 <jedi95@gmail.com> and
+# CFSworks <CFSworks@gmail.com>
+#
+# 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.
+
+import platform
+from time import time
+from twisted.internet import reactor
+from minerutil.MMPProtocol import MMPClient
+from KernelInterface import KernelInterface
+
+#The main managing class for the miner itself.
+class Miner(object):
+
+ # This must be manually set for Git
+ VER = (1, 6, 3)
+ REVISION = reduce(lambda x,y: x*100+y, VER)
+ VERSION = 'v%s' % '.'.join(str(x) for x in VER)
+
+ def __init__(self):
+ self.logger = None
+ self.options = None
+ self.connection = None
+ self.kernel = None
+ self.queue = None
+ self.idle = True
+ self.cores = []
+ self.backup = False
+ self.failures = 0
+ self.lastMetaRate = 0.0
+ self.lastRateUpdate = time()
+
+ # Connection callbacks...
+ def onFailure(self):
+ self.logger.reportConnectionFailed()
+
+ #handle failover if url2 has been specified
+ if self.options.url2 is not None:
+ self.failoverCheck()
+
+ def onConnect(self):
+ self.logger.reportConnected(True)
+ def onDisconnect(self):
+ self.logger.reportConnected(False)
+ def onBlock(self, block):
+ self.logger.reportBlock(block)
+ def onMsg(self, msg):
+ self.logger.reportMsg(msg)
+ def onWork(self, work):
+ self.logger.reportDebug('Server gave new work; passing to WorkQueue')
+ self.queue.storeWork(work)
+ def onLongpoll(self, lp):
+ self.logger.reportType('RPC' + (' (+LP)' if lp else ''))
+ def onPush(self, ignored):
+ self.logger.log('LP: New work pushed')
+ def onLog(self, message):
+ self.logger.log(message)
+ def onDebug(self, message):
+ self.logger.reportDebug(message)
+
+ def failoverCheck(self):
+ if self.backup:
+ if (self.failures >= 1):
+ #disconnect and set connection to none
+ self.connection.disconnect()
+ self.connection = None
+
+ #log
+ self.logger.log("Backup server failed,")
+ self.logger.log("attempting to return to primary server.")
+
+ #reset failure count and return to primary server
+ self.failures = 0
+ self.backup = False
+ self.connection = self.options.makeConnection(self)
+ self.connection.connect()
+ else:
+ self.failures += 1
+ else:
+ #The main pool must fail 3 times before moving to the backup pool
+ if (self.failures >= 2):
+ #disconnect and set connection to none
+ self.connection.disconnect()
+ self.connection = None
+
+ #log
+ self.logger.log("Primary server failed too many times,")
+ self.logger.log("attempting to connect to backup server.")
+
+ #reset failure count and connect to backup server
+ self.failures = 0
+ self.backup = True
+ self.connection = self.options.makeConnection(self, True)
+ self.connection.connect()
+ else:
+ self.failures += 1
+
+ #since the main pool may fail from time to time, decrement the
+ #failure count after 5 minutes so we don't end up moving to the
+ #back pool when it isn't nessesary
+ def decrementFailures():
+ if self.failures > 1 and (not self.backup):
+ self.failures -= 1
+ reactor.callLater(300, decrementFailures)
+
+ def start(self, options):
+ #Configures the Miner via the options specified and begins mining.
+
+ self.options = options
+ self.logger = self.options.makeLogger(self, self)
+ self.connection = self.options.makeConnection(self)
+ self.kernel = self.options.makeKernel(KernelInterface(self))
+ self.queue = self.options.makeQueue(self)
+
+ #log a message to let the user know that phoenix is starting
+ self.logger.log("Phoenix %s starting..." % self.VERSION)
+
+ #this will need to be changed to add new protocols
+ if isinstance(self.connection, MMPClient):
+ self.logger.reportType('MMP')
+ else:
+ self.logger.reportType('RPC')
+
+ self.applyMeta()
+
+ # Go!
+ self.connection.connect()
+ self.kernel.start()
+ reactor.addSystemEventTrigger('before', 'shutdown', self.kernel.stop)
+
+ def applyMeta(self):
+ #Applies any static metafields to the connection, such as version,
+ #kernel, hardware, etc.
+
+ # It's important to note here that the name is already put in place by
+ # the Options's makeConnection function, since the Options knows the
+ # user's desired name for this miner anyway.
+
+ self.connection.setVersion(
+ 'phoenix', 'Phoenix Miner', self.VERSION)
+ system = platform.system() + ' ' + platform.version()
+ self.connection.setMeta('os', system)
+
+ #called by CoreInterface to add cores for total hashrate calculation
+ def _addCore(self, core):
+ self.cores.append(core)
+
+ #used by WorkQueue to report when the miner is idle
+ def reportIdle(self, idle):
+
+ #if idle status has changed force an update
+ if self.idle != idle:
+ if idle:
+ self.idle = idle
+ self.logger.log("Warning: work queue empty, miner is idle")
+ self.logger.reportRate(0, True)
+ self.connection.setMeta('rate', 0)
+ self.lastMetaRate = time()
+ self.idleFixer()
+ else:
+ self.idle = idle
+ self.logger.updateStatus(True)
+
+ #request work from the protocol every 15 seconds while idle
+ def idleFixer(self):
+ if self.idle:
+ self.connection.requestWork()
+ reactor.callLater(15, self.idleFixer)
+
+ def updateAverage(self):
+ #Query all mining cores for their Khash/sec rate and sum.
+
+ total = 0
+ if not self.idle:
+ for core in self.cores:
+ total += core.getRate()
+
+ self.logger.reportRate(total)
+
+ # Let's not spam the server with rate messages.
+ if self.lastMetaRate+30 < time():
+ self.connection.setMeta('rate', total)
+ self.lastMetaRate = time()
194 QueueReader.py
@@ -0,0 +1,194 @@
+# Copyright (C) 2011 by jedi95 <jedi95@gmail.com> and
+# CFSworks <CFSworks@gmail.com>
+#
+# 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.
+
+from time import time
+from Queue import Queue, Empty
+from twisted.internet import reactor, defer
+
+from KernelInterface import CoreInterface
+
+class QueueReader(object):
+ """A QueueReader is a very efficient WorkQueue reader that keeps the next
+ nonce range available at all times. The benefit is that threaded mining
+ kernels waste no time getting the next range, since this class will have it
+ completely requested and preprocessed for the next iteration.
+
+ The QueueReader is iterable, so a dedicated mining thread needs only to do
+ for ... in self.qr:
+ """
+
+ SAMPLES = 3
+
+ def __init__(self, core, preprocessor=None, workSizeCallback=None):
+ if not isinstance(core, CoreInterface):
+ # Older kernels used to pass the KernelInterface, and not a
+ # CoreInterface. This is deprecated. We'll go ahead and take care
+ # of that, though...
+ core = core.addCore()
+ self.core = core
+ self.interface = core.getKernelInterface()
+ self.preprocessor = preprocessor
+ self.workSizeCallback = workSizeCallback
+
+ if self.preprocessor is not None:
+ if not callable(self.preprocessor):
+ raise TypeError('the given preprocessor must be callable')
+ if self.workSizeCallback is not None:
+ if not callable(self.workSizeCallback):
+ raise TypeError('the given workSizeCallback must be callable')
+
+ # This shuttles work to the dedicated thread.
+ self.dataQueue = Queue()
+
+ # Used in averaging the last execution times.
+ self.executionTimeSamples = []
+ self.averageExecutionTime = None
+
+ # This gets changed by _updateWorkSize.
+ self.executionSize = None
+
+ # Statistics accessed by the dedicated thread.
+ self.currentData = None
+ self.startedAt = None
+
+ def start(self):
+ """Called by the kernel when it's actually starting."""
+ self._updateWorkSize(None, None)
+ self._requestMore()
+ # We need to know when the current NonceRange in the dataQueue is old.
+ self.interface.addStaleCallback(self._staleCallback)
+
+ def stop(self):
+ """Called by the kernel when it's told to stop. This also brings down
+ the loop running in the mining thread.
+ """
+ # Tell the other thread to exit cleanly.
+ while not self.dataQueue.empty():
+ try:
+ self.dataQueue.get(False)
+ except Empty:
+ pass
+ self.dataQueue.put(StopIteration())
+
+ def _ranExecution(self, dt, nr):
+ """An internal function called after an execution completes, with the
+ time it took. Used to keep track of the time so kernels can use it to
+ tune their execution times.
+ """
+
+ if dt > 0:
+ self.core.updateRate(int(nr.size/dt/1000))
+
+ self.executionTimeSamples.append(dt)
+ self.executionTimeSamples = self.executionTimeSamples[-self.SAMPLES:]
+
+ if len(self.executionTimeSamples) == self.SAMPLES:
+ averageExecutionTime = (sum(self.executionTimeSamples) /
+ len(self.executionTimeSamples))
+
+ self._updateWorkSize(averageExecutionTime, nr.size)
+
+ def _updateWorkSize(self, time, size):
+ """An internal function that tunes the executionSize to that specified
+ by the workSizeCallback; which is in turn passed the average of the
+ last execution times.
+ """
+ if self.workSizeCallback:
+ self.executionSize = self.workSizeCallback(time, size)
+
+ def _requestMore(self):
+ """This is used to start the process of making a new item available in
+ the dataQueue, so the dedicated thread doesn't have to block.
+ """
+
+ # This should only run if there's no ready-to-go work in the queue.
+ if not self.dataQueue.empty():
+ return
+
+ if self.executionSize is None:
+ d = self.interface.fetchRange()
+ else:
+ d = self.interface.fetchRange(self.executionSize)
+
+ def preprocess(nr):
+ # If preprocessing is not necessary, just tuplize right away.
+ if not self.preprocessor:
+ return (nr, nr)
+
+ d2 = defer.maybeDeferred(self.preprocessor, nr)
+
+ # Tuplize the preprocessed result.
+ def callback(x):
+ return (x, nr)
+ d2.addCallback(callback)
+ return d2
+ d.addCallback(preprocess)
+
+ d.addCallback(self.dataQueue.put_nowait)
+
+ def _staleCallback(self):
+ """Called when the WorkQueue gets new work, rendering whatever is in
+ dataQueue old.
+ """
+
+ #only clear queue and request more if no work present, since that
+ #meas a request for more work is already in progress
+ if not self.dataQueue.empty():
+ # Out with the old...
+ while not self.dataQueue.empty():
+ try:
+ self.dataQueue.get(False)
+ except Empty: continue
+ # ...in with the new.
+ self._requestMore()
+
+ def __iter__(self):
+ return self
+ def next(self):
+ """Since QueueReader is iterable, this is the function that runs the
+ for-loop and dispatches work to the thread.
+
+ This should be the only thread that executes outside of the Twisted
+ main thread.
+ """
+
+ # If we just completed a range, we should tell the main thread.
+ now = time()
+ if self.currentData:
+ dt = now - self.startedAt
+ # self.currentData[1] is the un-preprocessed NonceRange.
+ reactor.callFromThread(self._ranExecution, dt, self.currentData[1])
+ self.startedAt = now
+
+ # Block for more data from the main thread. In 99% of cases, though,
+ # there should already be something here.
+ # Note that this comes back with either a tuple, or a StopIteration()
+ self.currentData = self.dataQueue.get(True)
+
+ # Does the main thread want us to shut down, or pass some more data?
+ if isinstance(self.currentData, StopIteration):
+ raise self.currentData
+
+ # We just took the only item in the queue. It needs to be restocked.
+ reactor.callFromThread(self._requestMore)
+
+ # currentData is actually a tuple, with item 0 intended for the kernel.
+ return self.currentData[0]
211 WorkQueue.py
@@ -0,0 +1,211 @@
+# Copyright (C) 2011 by jedi95 <jedi95@gmail.com> and
+# CFSworks <CFSworks@gmail.com>
+#
+# 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.
+
+from minerutil.Midstate import calculateMidstate
+from twisted.internet import defer
+from collections import deque
+
+"""A WorkUnit is a single unit containing 2^32 nonces. A single getWork
+request returns a WorkUnit.
+"""
+class WorkUnit(object):
+ data = None
+ target = None
+ midstate = None
+ nonces = None
+ base = None
+ work = None
+
+"""A NonceRange is a range of nonces from a WorkUnit, to be dispatched in a
+single execution of a mining kernel. The size of the NonceRange can be
+adjusted to tune the performance of the kernel, but will always
+be a multiple of 256.
+
+This class doesn't actually do anything, it's just a well-defined container
+that kernels can pull information out of.
+"""
+class NonceRange(object):
+
+ def __init__(self, unit, base, size):
+ self.unit = unit # The WorkUnit this NonceRange comes from.
+ self.base = base # The base nonce.
+ self.size = size # How many nonces this NonceRange says to test.
+
+
+class WorkQueue(object):
+ """A WorkQueue contains WorkUnits and dispatches NonceRanges when requested
+ by the miner. WorkQueues dispatch deffereds when they runs out of nonces.
+ """
+
+ def __init__(self, miner, options):
+
+ self.miner = miner
+ self.queueSize = options.getQueueSize()
+ self.logger = options.makeLogger(self, miner)
+
+ self.queue = deque('', self.queueSize)
+ self.deferredQueue = deque()
+ self.currentUnit = None
+ self.block = ''
+ self.lastBlock = None
+ self.test = False
+
+ # This is set externally. Not the best practice, but it can be changed
+ # in the future.
+ self.staleCallbacks = []
+
+ # Called by foundNonce to check if a NonceRange is stale before submitting
+ def isRangeStale(self, nr):
+ return (nr.unit.data[4:36] != self.block)
+
+ def storeWork(self, wu):
+
+ #check if this work matches the previous block
+ if self.lastBlock is not None and (wu.data[4:36] == self.lastBlock):
+ self.logger.reportDebug('Server gave work from the previous '
+ 'block, ignoring.')
+ #if the queue is too short request more work
+ if (len(self.queue)) < (self.queueSize):
+ self.miner.connection.requestWork()
+ return
+
+ #create a WorkUnit
+ work = WorkUnit()
+ work.data = wu.data
+ work.target = wu.target
+ work.midstate = calculateMidstate(work.data[:64])
+ work.nonces = 2 ** wu.mask
+ work.base = 0
+
+ #check if there is a new block, if so reset queue
+ newBlock = (wu.data[4:36] != self.block)
+ if newBlock:
+ self.queue.clear()
+ self.currentUnit = None
+ self.lastBlock = self.block
+ self.block = wu.data[4:36]
+ self.logger.reportDebug("New block (WorkQueue)")
+
+ #clear the idle flag since we just added work to queue
+ self.miner.reportIdle(False)
+
+ #add new WorkUnit to queue
+ if work.data and work.target and work.midstate and work.nonces:
+ self.queue.append(work)
+
+ #if the queue is too short request more work
+ if (len(self.queue)) < (self.queueSize):
+ self.miner.connection.requestWork()
+
+ #if there is a new block notify kernels that their work is now stale
+ if newBlock:
+ for callback in self.staleCallbacks:
+ callback()
+
+ #check if there are deferred NonceRange requests pending
+ #since requests to fetch a NonceRange can add additional deferreds to
+ #the queue, cache the size beforehand to avoid infinite loops.
+ for i in range(len(self.deferredQueue)):
+ df, size = self.deferredQueue.popleft()
+ d = self.fetchRange(size)
+ d.chainDeferred(df)
+
+ #gets the next WorkUnit from queue
+ def getNext(self):
+
+ #check if the queue will fall below desired size
+ if (len(self.queue) - 1) < (self.queueSize):
+ self.miner.connection.requestWork()
+
+ #return next WorkUnit
+ return self.queue.popleft()
+
+ def getRangeFromUnit(self, size):
+
+ #get remaining nonces
+ noncesLeft = self.currentUnit.nonces - self.currentUnit.base
+
+ #if there are enough nonces to fill the full reqest
+ if noncesLeft >= size:
+ nr = NonceRange(self.currentUnit, self.currentUnit.base, size)
+
+ #check if this uses up the rest of the WorkUnit
+ if size >= noncesLeft:
+ self.currentUnit = None
+ else:
+ self.currentUnit.base += size
+
+ #otherwise send whatever is left
+ else:
+ nr = NonceRange(
+ self.currentUnit, self.currentUnit.base, noncesLeft)
+ self.currentUnit = None
+
+ #return the range
+ return nr
+
+ def fetchRange(self, size=0x10000):
+
+ #make sure size is not too large
+ size = min(size, 0x100000000)
+
+ #size must be a multiple of 256
+ size = 256 * int(size / 256) #rounded up
+
+ #make sure size is not too small
+ size = max(size, 256)
+
+ #check if the current unit exists
+ if self.currentUnit is not None:
+
+ #get a nonce range
+ nr = self.getRangeFromUnit(size)
+
+ #return the range
+ return defer.succeed(nr)
+
+ #if there is no current unit
+ else:
+ #if there is another unit in queue
+ if len(self.queue) >= 1:
+
+ #get the next unit from queue
+ self.currentUnit = self.getNext()
+
+ #get a nonce range
+ nr = self.getRangeFromUnit(size)
+
+ #return the range
+ return defer.succeed(nr)
+
+ #if the queue is empty
+ else:
+
+ #request more work
+ self.miner.connection.requestWork()
+
+ #report that the miner is idle
+ self.miner.reportIdle(True)
+
+ #set up and return deferred
+ df = defer.Deferred()
+ self.deferredQueue.append((df, size))
+ return df
122 kernels/phatk/BFIPatcher.py
@@ -0,0 +1,122 @@
+# Copyright (C) 2011 by jedi95 <jedi95@gmail.com> and
+# CFSworks <CFSworks@gmail.com>
+#
+# 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.
+
+import struct
+
+class PatchError(Exception): pass
+
+class BFIPatcher(object):
+ """Patches .ELF files compiled for Evergreen GPUs; changes the microcode
+ so that any BYTE_ALIGN_INT instructions become BFI_INT.
+ """
+
+ def __init__(self, interface):
+ self.interface = interface
+
+ def patch(self, data):
+ """Run the process of patching an ELF."""
+
+ self.interface.debug('Finding inner ELF...')
+ innerPos = self.locateInner(data)
+ self.interface.debug('Patching inner ELF...')
+ inner = data[innerPos:]
+ patched = self.patchInner(inner)
+ self.interface.debug('Patch complete, returning to kernel...')
+ return data[:innerPos] + patched
+
+ def patchInner(self, data):
+ sections = self.readELFSections(data)
+ # We're looking for .text -- there should be two of them.
+ textSections = filter(lambda x: x[0] == '.text', sections)
+ if len(textSections) != 2:
+ self.interface.debug('Inner ELF does not have 2 .text sections!')
+ self.interface.debug('Sections are: %r' % sections)
+ raise PatchError()
+ name, offset, size = textSections[1]
+ before, text2, after = (data[:offset], data[offset:offset+size],
+ data[offset+size:])
+
+ self.interface.debug('Patching instructions...')
+ text2 = self.patchInstructions(text2)
+ return before + text2 + after
+
+ def patchInstructions(self, data):
+ output = ''
+ nPatched = 0
+ for i in xrange(len(data)/8):
+ inst, = struct.unpack('Q', data[i*8:i*8+8])
+ # Is it BYTE_ALIGN_INT?
+ if (inst&0x9003f00002001000) == 0x0001a00000000000:
+ nPatched += 1
+ inst ^= (0x0001a00000000000 ^ 0x0000c00000000000) # BFI_INT
+ output += struct.pack('Q', inst)
+ self.interface.debug('BFI-patched %d instructions...' % nPatched)
+ if nPatched < 60:
+ self.interface.debug('Patch safety threshold not met!')
+ raise PatchError()
+ return output
+
+ def locateInner(self, data):
+ """ATI uses an ELF-in-an-ELF. I don't know why. This function's job is
+ to find it.
+ """
+
+ pos = data.find('\x7fELF', 1)
+ if pos == -1 or data.find('\x7fELF', pos+1) != -1: # More than 1 is bad
+ self.interface.debug('Inner ELF not located!')
+ raise PatchError()
+ return pos
+
+ def readELFSections(self, data):
+ try:
+ (ident1, ident2, type, machine, version, entry, phoff,
+ shoff, flags, ehsize, phentsize, phnum, shentsize, shnum,
+ shstrndx) = struct.unpack('QQHHIIIIIHHHHHH', data[:52])
+
+ if ident1 != 0x64010101464c457f:
+ self.interface.debug('Invalid ELF header!')
+ raise PatchError()
+
+ # No section header?
+ if shoff == 0:
+ return []
+
+ # Find out which section contains the section header names
+ shstr = data[shoff+shstrndx*shentsize:shoff+(shstrndx+1)*shentsize]
+ (nameIdx, type, flags, addr, nameTableOffset, size, link, info,
+ addralign, entsize) = struct.unpack('IIIIIIIIII', shstr)
+
+ # Grab the section header.
+ sh = data[shoff:shoff+shnum*shentsize]
+
+ sections = []
+ for i in xrange(shnum):
+ rawEntry = sh[i*shentsize:(i+1)*shentsize]
+ (nameIdx, type, flags, addr, offset, size, link, info,
+ addralign, entsize) = struct.unpack('IIIIIIIIII', rawEntry)
+ nameOffset = nameTableOffset + nameIdx
+ name = data[nameOffset:data.find('\x00', nameOffset)]
+ sections.append((name, offset, size))
+
+ return sections
+ except struct.error:
+ self.interface.debug('A struct.error occurred while reading ELF!')
+ raise PatchError()
427 kernels/phatk/__init__.py
@@ -0,0 +1,427 @@
+# Copyright (C) 2011 by jedi95 <jedi95@gmail.com> and
+# CFSworks <CFSworks@gmail.com> and
+# Phateus <jesse.moll@gmail.com>
+#
+# 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.
+
+import pyopencl as cl
+import numpy as np
+import os
+import math
+
+from hashlib import md5
+from struct import pack, unpack
+from twisted.internet import reactor
+
+from minerutil.Midstate import calculateMidstate
+from QueueReader import QueueReader
+from KernelInterface import *
+from BFIPatcher import *
+
+class KernelData(object):
+ """This class is a container for all the data required for a single kernel
+ execution.
+ """
+
+ def __init__(self, nonceRange, core, vectors, aggression):
+ # Prepare some raw data, converting it into the form that the OpenCL
+ # function expects.
+ data = np.array(
+ unpack('IIII', nonceRange.unit.data[64:]), dtype=np.uint32)
+
+ # Vectors do twice the work per execution, so calculate accordingly...
+ rateDivisor = 2 if vectors else 1
+
+ # get the number of iterations from the aggression and size
+ self.iterations = int((nonceRange.size / (1 << aggression)))
+ self.iterations = max(1, self.iterations)
+
+ #set the size to pass to the kernel based on iterations and vectors
+ self.size = (nonceRange.size / rateDivisor) / self.iterations
+
+ #compute bases for each iteration
+ self.base = [None] * self.iterations
+ for i in range(self.iterations):
+ self.base[i] = pack('I',
+ (nonceRange.base/rateDivisor) + (i * self.size))
+
+ #set up state and precalculated static data
+ self.state = np.array(
+ unpack('IIIIIIII', nonceRange.unit.midstate), dtype=np.uint32)
+ self.state2 = np.array(unpack('IIIIIIII',
+ calculateMidstate(nonceRange.unit.data[64:80] +
+ '\x00\x00\x00\x80' + '\x00'*40 + '\x80\x02\x00\x00',
+ nonceRange.unit.midstate, 3)), dtype=np.uint32)
+ self.state2 = np.array(
+ list(self.state2)[3:] + list(self.state2)[:3], dtype=np.uint32)
+ self.nr = nonceRange
+
+ # added place for another variable
+ self.f = np.zeros(6, np.uint32)
+ self.calculateF(data)
+
+ # D1 == old D1 + (K[4] + W[4])
+ self.state2[3] = np.uint32(self.state2[3] + 0xb956c25b)
+
+ def calculateF(self, data):
+ rot = lambda x,y: x>>y | x<<(32-y)
+
+ # W16
+ self.f[1] = np.uint32(data[0] + (rot(data[1], 7) ^ rot(data[1], 18) ^
+ (data[1] >> 3)))
+ # W17
+ self.f[2] = np.uint32(data[1] + (rot(data[2], 7) ^ rot(data[2], 18) ^
+ (data[2] >> 3)) + 0x01100000)
+
+ # W2 == old W2 + P1() with W16
+ self.f[0] = np.uint32(data[2] + (
+ rot(self.f[1], 32-13) ^
+ rot(self.f[1], 32-15) ^
+ (self.f[1] >> 10)
+ ))
+
+ # W17_2 == P2(19) + P1() with W17
+ self.f[5] = np.uint32(0x11002000+(
+ rot(self.f[2], 32-13) ^
+ rot(self.f[2], 32-15) ^
+ (self.f[2] >> 10)))
+
+ #2 parts of the first SHA round
+ # T1
+ self.f[4] = np.uint32((rot(self.state2[5], 2) ^
+ rot(self.state2[5], 13) ^ rot(self.state2[5], 22)) +
+ ((self.state2[5] & self.state2[6]) | (self.state2[7] &
+ (self.state2[5] | self.state2[6]))))
+
+ # Preval4
+ self.f[3] = np.uint32(self.state[4] + (rot(self.state2[1], 6) ^
+ rot(self.state2[1], 11) ^ rot(self.state2[1], 25)) +
+ (self.state2[3] ^ (self.state2[1] & (self.state2[2] ^
+ self.state2[3]))) + 0xe9b5dba5)
+
+class MiningKernel(object):
+ """A Phoenix Miner-compatible kernel that uses the poclbm OpenCL kernel."""
+
+ PLATFORM = KernelOption(
+ 'PLATFORM', int, default=None,
+ help='The ID of the OpenCL platform to use')
+ DEVICE = KernelOption(
+ 'DEVICE', int, default=None,
+ help='The ID of the OpenCL device to use')
+ VECTORS = KernelOption(
+ 'VECTORS', bool, default=False, advanced=True,
+ help='Enable vector support in the kernel?')
+ FASTLOOP = KernelOption(
+ 'FASTLOOP', bool, default=True, advanced=True,
+ help='Run iterative mining thread?')
+ AGGRESSION = KernelOption(
+ 'AGGRESSION', int, default=5, advanced=True,
+ help='Exponential factor indicating how much work to run '
+ 'per OpenCL execution')
+ WORKSIZE = KernelOption(
+ 'WORKSIZE', int, default=None, advanced=True,
+ help='The worksize to use when executing CL kernels.')
+ BFI_INT = KernelOption(
+ 'BFI_INT', bool, default=True, advanced=True,
+ help='Use the BFI_INT instruction for AMD/ATI GPUs.')
+ OUTPUT_SIZE = 0x100
+
+ # This must be manually set for Git
+ REVISION = 116
+
+ def __init__(self, interface):
+ platforms = cl.get_platforms()
+
+ # Initialize object attributes and retrieve command-line options...)
+ self.device = None
+ self.kernel = None
+ self.interface = interface
+ self.core = self.interface.addCore()
+ self.defines = ''
+ self.loopExponent = 0
+
+ # Set the initial number of nonces to run per execution
+ # 2^(16 + aggression)
+ self.AGGRESSION += 16
+ self.AGGRESSION = min(32, self.AGGRESSION)
+ self.AGGRESSION = max(16, self.AGGRESSION)
+ self.size = 1 << self.AGGRESSION
+
+ # We need a QueueReader to efficiently provide our dedicated thread
+ # with work.
+ self.qr = QueueReader(self.core, lambda nr: self.preprocess(nr),
+ lambda x,y: self.size * 1 << self.loopExponent)
+
+ # The platform selection must be valid to mine.
+ if self.PLATFORM >= len(platforms) or \
+ (self.PLATFORM is None and len(platforms) > 1):
+ self.interface.log(
+ 'Wrong platform or more than one OpenCL platform found, '
+ 'use PLATFORM=ID to select one of the following\n',
+ False, True)
+
+ for i,p in enumerate(platforms):
+ self.interface.log(' [%d]\t%s' % (i, p.name), False, False)
+
+ # Since the platform is invalid, we can't mine.
+ self.interface.fatal()
+ return
+ elif self.PLATFORM is None:
+ self.PLATFORM = 0
+
+ devices = platforms[self.PLATFORM].get_devices()
+
+ # The device selection must be valid to mine.
+ if self.DEVICE >= len(devices) or \
+ (self.DEVICE is None and len(devices) > 1):
+ self.interface.log(
+ 'No device specified or device not found, '
+ 'use DEVICE=ID to specify one of the following\n',
+ False, True)
+
+ for i,d in enumerate(devices):
+ self.interface.log(' [%d]\t%s' % (i, d.name), False, False)
+
+ # Since the device selection is invalid, we can't mine.
+ self.interface.fatal()
+ return
+ elif self.DEVICE is None:
+ self.DEVICE = 0
+
+ self.device = devices[self.DEVICE]
+
+ # We need the appropriate kernel for this device...
+ try:
+ self.loadKernel(self.device)
+ except Exception:
+ self.interface.fatal("Failed to load OpenCL kernel!")
+ return
+
+ # Initialize a command queue to send commands to the device, and a
+ # buffer to collect results in...
+ self.commandQueue = cl.CommandQueue(self.context)
+ self.output = np.zeros(self.OUTPUT_SIZE+1, np.uint32)
+ self.output_buf = cl.Buffer(
+ self.context, cl.mem_flags.WRITE_ONLY | cl.mem_flags.USE_HOST_PTR,
+ hostbuf=self.output)
+
+ self.applyMeta()
+
+ def applyMeta(self):
+ """Apply any kernel-specific metadata."""
+ self.interface.setMeta('kernel', 'phatk r%s' % self.REVISION)
+ self.interface.setMeta('device', self.device.name.replace('\x00',''))
+ self.interface.setMeta('cores', self.device.max_compute_units)
+
+ def loadKernel(self, device):
+ """Load the kernel and initialize the device."""
+ self.context = cl.Context([device], None, None)
+
+ # These definitions are required for the kernel to function.
+ self.defines += (' -DOUTPUT_SIZE=' + str(self.OUTPUT_SIZE))
+ self.defines += (' -DOUTPUT_MASK=' + str(self.OUTPUT_SIZE - 1))
+
+ # If the user wants to mine with vectors, enable the appropriate code
+ # in the kernel source.
+ if self.VECTORS:
+ self.defines += ' -DVECTORS'
+
+ # Some AMD devices support a special "bitalign" instruction that makes
+ # bitwise rotation (required for SHA-256) much faster.
+ if (device.extensions.find('cl_amd_media_ops') != -1):
+ self.defines += ' -DBITALIGN'
+ #enable the expierimental BFI_INT instruction optimization
+ if self.BFI_INT:
+ self.defines += ' -DBFI_INT'
+ else:
+ #Since phatk and phatk2 will error out on Nvidia GPUs
+ #make sure the user knows that they need to use poclbm
+ self.interface.fatal("GPU not supported! phatk is designed for "
+ "ATI 5xxx and newer only. Try -k poclbm instead.")
+ return
+
+ # Locate and read the OpenCL source code in the kernel's directory.
+ kernelFileDir, pyfile = os.path.split(__file__)
+ kernelFilePath = os.path.join(kernelFileDir, 'kernel.cl')
+ kernelFile = open(kernelFilePath, 'r')
+ kernel = kernelFile.read()
+ kernelFile.close()
+
+ # For fast startup, we cache the compiled OpenCL code. The name of the
+ # cache is determined as the hash of a few important,
+ # compilation-specific pieces of information.
+ m = md5()
+ m.update(device.platform.name)
+ m.update(device.platform.version)
+ m.update(device.name)
+ m.update(self.defines)
+ m.update(kernel)
+ cacheName = '%s.elf' % m.hexdigest()
+
+ fileName = os.path.join(kernelFileDir, cacheName)
+
+ # Finally, the actual work of loading the kernel...
+ try:
+ binary = open(fileName, 'rb')
+ except IOError:
+ binary = None
+
+ try:
+ if binary is None:
+ self.kernel = cl.Program(
+ self.context, kernel).build(self.defines)
+
+ #apply BFI_INT if enabled
+ if self.BFI_INT:
+ #patch the binary output from the compiler
+ patcher = BFIPatcher(self.interface)
+ binaryData = patcher.patch(self.kernel.binaries[0])
+
+ self.interface.debug("Applied BFI_INT patch")
+
+ #reload the kernel with the patched binary
+ self.kernel = cl.Program(
+ self.context, [device],
+ [binaryData]).build(self.defines)
+
+ #write the kernel binaries to file
+ binaryW = open(fileName, 'wb')
+ binaryW.write(self.kernel.binaries[0])
+ binaryW.close()
+ else:
+ binaryData = binary.read()
+ self.kernel = cl.Program(
+ self.context, [device], [binaryData]).build(self.defines)
+
+ except cl.LogicError:
+ self.interface.fatal("Failed to compile OpenCL kernel!")
+ return
+ except PatchError:
+ self.interface.fatal('Failed to apply BFI_INT patch to kernel! '
+ 'Is BFI_INT supported on this hardware?')
+ return
+ finally:
+ if binary: binary.close()
+
+ cl.unload_compiler()
+
+ # If the user didn't specify their own worksize, use the maxium
+ # supported by the device.
+ maxSize = self.kernel.search.get_work_group_info(
+ cl.kernel_work_group_info.WORK_GROUP_SIZE, self.device)
+
+ if self.WORKSIZE is None:
+ self.WORKSIZE = maxSize
+ else:
+ if self.WORKSIZE > maxSize:
+ self.interface.log('Warning: Worksize exceeds the maximum of '
+ + str(maxSize) + ', using default.')
+ if self.WORKSIZE < 1:
+ self.interface.log('Warning: Invalid worksize, using default.')
+
+ self.WORKSIZE = min(self.WORKSIZE, maxSize)
+ self.WORKSIZE = max(self.WORKSIZE, 1)
+ #if the worksize is not a power of 2, round down to the nearest one
+ if (self.WORKSIZE & (self.WORKSIZE - 1)) != 0:
+ self.WORKSIZE = 1 << int(math.floor(math.log(X)/math.log(2)))
+
+ def start(self):
+ """Phoenix wants the kernel to start."""
+
+ self.qr.start()
+ reactor.callInThread(self.mineThread)
+
+ def stop(self):
+ """Phoenix wants this kernel to stop. The kernel is not necessarily
+ reusable, so it's safe to clean up as well.
+ """
+ self.qr.stop()
+
+ def updateIterations(self):
+ # Set up the number of internal iterations to run if FASTLOOP enabled
+ rate = self.core.getRate()
+
+ if not (rate <= 0):
+ #calculate the number of iterations to run
+ EXP = max(0, (math.log(rate)/math.log(2)) - (self.AGGRESSION - 8))
+ #prevent switching between loop exponent sizes constantly
+ if EXP > self.loopExponent + 0.54:
+ EXP = round(EXP)
+ elif EXP < self.loopExponent - 0.65:
+ EXP = round(EXP)
+ else:
+ EXP = self.loopExponent
+
+ self.loopExponent = int(max(0, EXP))
+
+ def preprocess(self, nr):
+ if self.FASTLOOP:
+ self.updateIterations()
+
+ kd = KernelData(nr, self.core, self.VECTORS, self.AGGRESSION)
+ return kd
+
+ def postprocess(self, output, nr):
+ # Scans over a single buffer produced as a result of running the
+ # OpenCL kernel on the device. This is done outside of the mining thread
+ # for efficiency reasons.
+
+ # Iterate over only the first OUTPUT_SIZE items. Exclude the last item
+ # which is a duplicate of the most recently-found nonce.
+ for i in xrange(self.OUTPUT_SIZE):
+ if output[i]:
+ if not self.interface.foundNonce(nr, int(output[i])):
+ hash = self.interface.calculateHash(nr, int(output[i]))
+ if not hash.endswith('\x00\x00\x00\x00'):
+ self.interface.error('Unusual behavior from OpenCL. '
+ 'Hardware problem?')
+
+ def mineThread(self):
+ for data in self.qr:
+ for i in range(data.iterations):
+ # added C1 + K[5]
+ # added W17_2
+ # modified Preval4 = Preval4 + T1
+ # modified T1 = T1 - state0
+ self.kernel.search(
+ self.commandQueue, (data.size, ), (self.WORKSIZE, ),
+ data.state[0], data.state[1], data.state[2], data.state[3],
+ data.state[4], data.state[5], data.state[6], data.state[7],
+ data.state2[1], data.state2[2], np.uint32(data.state2[2] + 0x59f111f1), data.state2[3],
+ data.state2[5], data.state2[6], data.state2[7],
+ data.base[i],
+ data.f[0],
+ data.f[1],data.f[2], data.f[5],
+ (data.f[3] + data.f[4]), (data.state[0] - data.f[4]),
+ self.output_buf)
+ cl.enqueue_read_buffer(
+ self.commandQueue, self.output_buf, self.output)
+ self.commandQueue.finish()
+
+ # The OpenCL code will flag the last item in the output buffer when
+ # it finds a valid nonce. If that's the case, send it to the main
+ # thread for postprocessing and clean the buffer for the next pass.
+ if self.output[self.OUTPUT_SIZE]:
+ reactor.callFromThread(self.postprocess, self.output.copy(),
+ data.nr)
+
+ self.output.fill(0)
+ cl.enqueue_write_buffer(
+ self.commandQueue, self.output_buf, self.output)
+
402 kernels/phatk/kernel.cl
@@ -0,0 +1,402 @@
+// This file is taken and modified from the public-domain poclbm project, and
+// we have therefore decided to keep it public-domain in Phoenix.
+
+// 2011-07-17: further modified by Diapolo and still public-domain
+
+#ifdef VECTORS
+ typedef uint2 u;
+#else
+ typedef uint u;
+#endif
+
+__constant uint K[64] = {
+ 0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, 0x923f82a4, 0xab1c5ed5,
+ 0xd807aa98, 0x12835b01, 0x243185be, 0x550c7dc3, 0x72be5d74, 0x80deb1fe, 0x9bdc06a7, 0xc19bf174,
+ 0xe49b69c1, 0xefbe4786, 0x0fc19dc6, 0x240ca1cc, 0x2de92c6f, 0x4a7484aa, 0x5cb0a9dc, 0x76f988da,
+ 0x983e5152, 0xa831c66d, 0xb00327c8, 0xbf597fc7, 0xc6e00bf3, 0xd5a79147, 0x06ca6351, 0x14292967,
+ 0x27b70a85, 0x2e1b2138, 0x4d2c6dfc, 0x53380d13, 0x650a7354, 0x766a0abb, 0x81c2c92e, 0x92722c85,
+ 0xa2bfe8a1, 0xa81a664b, 0xc24b8b70, 0xc76c51a3, 0xd192e819, 0xd6990624, 0xf40e3585, 0x106aa070,
+ 0x19a4c116, 0x1e376c08, 0x2748774c, 0x34b0bcb5, 0x391c0cb3, 0x4ed8aa4a, 0x5b9cca4f, 0x682e6ff3,
+ 0x748f82ee, 0x78a5636f, 0x84c87814, 0x8cc70208, 0x90befffa, 0xa4506ceb, 0xbef9a3f7, 0xc67178f2
+};
+
+// H[6] = 0x08909ae5U + 0xb0edbdd0 + K[0] == 0xfc08884d
+// H[7] = -0x5be0cd19 - (0x90befffa) K[60] == -0xec9fcd13
+__constant uint H[8] = {
+ 0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0x510e527f, 0x9b05688c, 0x1f83d9ab, 0xfc08884d, 0xec9fcd13
+};
+
+// L = 0xa54ff53a + 0xb0edbdd0 + K[0] == 0x198c7e2a2
+__constant ulong L = 0x198c7e2a2;
+
+// offset for W[] array to reduce it's size (W[0] - W[15] are hard-coded or computed without use of P() calculations)
+#define O 15
+
+#ifdef BITALIGN
+ #pragma OPENCL EXTENSION cl_amd_media_ops : enable
+ #define rot(x, y) amd_bitalign(x, x, (u)(32 - y))
+#else
+ #define rot(x, y) rotate(x, (u)y)
+#endif
+
+#ifdef BFI_INT
+ #define Ch(x, y, z) amd_bytealign(x, y, z)
+#else
+ #define Ch(x, y, z) bitselect(z, y, x)
+#endif
+
+// Ma now uses the Ch function, if BFI_INT is enabled, the optimized Ch version is used
+#define Ma(x, y, z) Ch((z ^ x), y, x)
+
+// Various intermediate calculations for each SHA round
+#define s0(n) (rot(Vals[(128 - n) % 8], 30) ^ rot(Vals[(128 - n) % 8], 19) ^ rot(Vals[(128 - n) % 8], 10))
+#define s1(n) (rot(Vals[(132 - n) % 8], 26) ^ rot(Vals[(132 - n) % 8], 21) ^ rot(Vals[(132 - n) % 8], 7))
+#define ch(n) (Ch(Vals[(132 - n) % 8], Vals[(133 - n) % 8], Vals[(134 - n) % 8]))
+#define ma(n) (Ma(Vals[(129 - n) % 8], Vals[(130 - n) % 8], Vals[(128 - n) % 8]))
+#define t1W(n) (K[n % 64] + Vals[(135 - n) % 8] + W[n - O] + s1(n) + ch(n))
+#define t1(n) (K[n % 64] + Vals[(135 - n) % 8] + s1(n) + ch(n))
+
+// intermediate W calculations
+#define P1(x) (rot(W[x - 2 - O], 15) ^ rot(W[x - 2 - O], 13) ^ (W[x - 2 - O] >> 10U))
+#define P2(x) (rot(W[x - 15 - O], 25) ^ rot(W[x - 15 - O], 14) ^ (W[x - 15 - O] >> 3U))
+#define P3(x) W[x - 7 - O]
+#define P4(x) W[x - 16 - O]
+
+// full W calculation
+#define W(x) (W[x - O] = P4(x) + P3(x) + P2(x) + P1(x))
+
+// SHA round without W calc
+#define sharoundW(n) { Vals[(131 - n) % 8] += t1W(n); Vals[(135 - n) % 8] = t1W(n) + s0(n) + ma(n); }
+#define sharound(n) { Vals[(131 - n) % 8] += t1(n); Vals[(135 - n) % 8] = t1(n) + s0(n) + ma(n); }
+
+// description of modified kernel init variables:
+//
+// C1addK5: C1addK5 = C1 + K[5]: C1addK5 = C1 + 0x59f111f1
+// D1: D1 = D1 + K[4] + W[4]: D1 = D1 + 0xe9b5dba5 + 0x80000000U
+// W2: W2 + W16 in P1(): W2 = P1(18) + P4(18)
+// W17_2: 0x80000000U in P2() = 0x11002000 + W17 in P1(): W17_2 = P1(19) + P2(19)
+// PreValaddT1: PreValaddT1 = PreVal4 + T1
+// T1substate0: T1substate0 = T1 - substate0
+
+__kernel void search( const uint state0, const uint state1, const uint state2, const uint state3,
+ const uint state4, const uint state5, const uint state6, const uint state7,
+ const uint B1, const uint C1, const uint C1addK5, const uint D1,
+ const uint F1, const uint G1, const uint H1,
+ const uint base,
+ const uint W2,
+ const uint W16, const uint W17, const uint W17_2,
+ const uint PreVal4addT1, const uint T1substate0,
+ __global uint * output)
+{
+ u W[124 - O];
+ u Vals[8];
+#ifdef VECTORS
+ u W_3 = ((base + get_global_id(0)) << 1) + (uint2)(0, 1);
+#else
+ u W_3 = base + get_global_id(0);
+#endif
+ u Temp;
+
+ Vals[0] = W_3 + PreVal4addT1 + T1substate0;
+ Vals[1] = B1;
+ Vals[2] = C1;
+
+ Vals[4] = W_3 + PreVal4addT1;
+ Vals[5] = F1;
+ Vals[6] = G1;
+
+ // used in: P2(19) == 285220864 (0x11002000), P4(20)
+ // W[4] = 0x80000000U;
+ // P1(x) is 0 for x == 7, 8, 9, 10, 11, 12, 13, 14, 15, 16
+ // P2(x) is 0 for x == 20, 21, 22, 23, 24, 25, 26, 27, 28, 29
+ // P3(x) is 0 for x == 12, 13, 14, 15, 16, 17, 18, 19, 20, 21
+ // P4(x) is 0 for x == 21, 22, 23, 24, 25, 26, 27, 28, 29, 30
+ // W[x] in sharound(x) is 0 for x == 5, 6, 7, 8, 9, 10, 11, 12, 13, 14
+ // W[14] = W[13] = W[12] = W[11] = W[10] = W[9] = W[8] = W[7] = W[6] = W[5] = 0x00000000U;
+ // used in: P2(30) == 10485845 (0xA00055), P3(22), P4(31)
+ // K[15] + W[15] == 0xc19bf174 + 0x00000280U = 0xc19bf3f4
+ W[15 - O] = 0x00000280U;
+ W[16 - O] = W16;
+ W[17 - O] = W17;
+ // P1(18) + P2(18) + P4(18)
+ W[18 - O] = W2 + (rot(W_3, 25) ^ rot(W_3, 14) ^ (W_3 >> 3U));
+ // P1(19) + P2(19) + P4(19)
+ W[19 - O] = W_3 + W17_2;
+ // P1(20) + P4(20)
+ W[20 - O] = (u)0x80000000U + P1(20);
+ W[21 - O] = P1(21);
+ W[22 - O] = P1(22) + P3(22);
+ W[23 - O] = P1(23) + P3(23);
+ W[24 - O] = P1(24) + P3(24);
+ W[25 - O] = P1(25) + P3(25);
+ W[26 - O] = P1(26) + P3(26);
+ W[27 - O] = P1(27) + P3(27);
+ W[28 - O] = P1(28) + P3(28);
+ W[29 - O] = P1(29) + P3(29);
+ W[30 - O] = (u)0xA00055 + P1(30) + P3(30);
+
+ // Round 4
+ Temp = D1 + ch(4) + s1(4);
+ Vals[7] = Temp + H1;
+ Vals[3] = Temp + ma(4) + s0(4);
+
+ // Round 5
+ Temp = C1addK5 + ch(5) + s1(5);
+ Vals[6] = Temp + G1;
+ Vals[2] = Temp + ma(5) + s0(5);
+
+ // W[] is zero for this rounds
+ sharound(6);
+ sharound(7);
+ sharound(8);
+ sharound(9);
+ sharound(10);
+ sharound(11);
+ sharound(12);
+ sharound(13);
+ sharound(14);
+
+ sharoundW(15);
+ sharoundW(16);
+ sharoundW(17);
+ sharoundW(18);
+ sharoundW(19);
+ sharoundW(20);
+ sharoundW(21);
+ sharoundW(22);
+ sharoundW(23);
+ sharoundW(24);
+ sharoundW(25);
+ sharoundW(26);
+ sharoundW(27);
+ sharoundW(28);
+ sharoundW(29);
+ sharoundW(30);
+
+ W(31);
+ sharoundW(31);
+ W(32);
+ sharoundW(32);
+ W(33);
+ sharoundW(33);
+ W(34);
+ sharoundW(34);
+ W(35);
+ sharoundW(35);
+ W(36);
+ sharoundW(36);
+ W(37);
+ sharoundW(37);
+ W(38);
+ sharoundW(38);
+ W(39);
+ sharoundW(39);
+ W(40);
+ sharoundW(40);
+ W(41);
+ sharoundW(41);
+ W(42);
+ sharoundW(42);
+ W(43);
+ sharoundW(43);
+ W(44);
+ sharoundW(44);
+ W(45);
+ sharoundW(45);
+ W(46);
+ sharoundW(46);
+ W(47);
+ sharoundW(47);
+ W(48);
+ sharoundW(48);
+ W(49);
+ sharoundW(49);
+ W(50);
+ sharoundW(50);
+ W(51);
+ sharoundW(51);
+ W(52);
+ sharoundW(52);
+ W(53);
+ sharoundW(53);
+ W(54);
+ sharoundW(54);
+ W(55);
+ sharoundW(55);
+ W(56);
+ sharoundW(56);
+ W(57);
+ sharoundW(57);
+ W(58);
+ sharoundW(58);
+ W(59);
+ sharoundW(59);
+ W(60);
+ sharoundW(60);
+ W(61);
+ sharoundW(61);
+ W(62);
+ sharoundW(62);
+ W(63);
+ sharoundW(63);
+
+ W[64 - O] = state0 + Vals[0];
+ W[65 - O] = state1 + Vals[1];
+ W[66 - O] = state2 + Vals[2];
+ W[67 - O] = state3 + Vals[3];
+ W[68 - O] = state4 + Vals[4];
+ W[69 - O] = state5 + Vals[5];
+ W[70 - O] = state6 + Vals[6];
+ W[71 - O] = state7 + Vals[7];
+ // used in: P2(87) = 285220864 (0x11002000), P4(88)
+ // K[72] + W[72] ==
+ W[72 - O] = 0x80000000U;
+ // P1(x) is 0 for x == 75, 76, 77, 78, 79, 80
+ // P2(x) is 0 for x == 88, 89, 90, 91, 92, 93
+ // P3(x) is 0 for x == 80, 81, 82, 83, 84, 85
+ // P4(x) is 0 for x == 89, 90, 91, 92, 93, 94
+ // W[x] in sharound(x) is 0 for x == 73, 74, 75, 76, 77, 78
+ // W[78] = W[77] = W[76] = W[75] = W[74] = W[73] = 0x00000000U;
+ // used in: P1(81) = 10485760 (0xA00000), P2(94) = 4194338 (0x400022), P3(86), P4(95)
+ // K[79] + W[79] ==
+ W[79 - O] = 0x00000100U;
+
+ Vals[0] = H[0];
+ Vals[1] = H[1];
+ Vals[2] = H[2];
+ Vals[3] = (u)L + W[64 - O];
+ Vals[4] = H[3];
+ Vals[5] = H[4];
+ Vals[6] = H[5];
+ Vals[7] = H[6] + W[64 - O];
+
+ sharoundW(65);
+ sharoundW(66);
+ sharoundW(67);
+ sharoundW(68);
+ sharoundW(69);
+ sharoundW(70);
+ sharoundW(71);
+ sharoundW(72);
+
+ // W[] is zero for this rounds
+ sharound(73);
+ sharound(74);
+ sharound(75);
+ sharound(76);
+ sharound(77);
+ sharound(78);
+
+ sharoundW(79);
+
+ W[80 - O] = P2(80) + P4(80);
+ W[81 - O] = (u)0xA00000 + P4(81) + P2(81);
+ W[82 - O] = P4(82) + P2(82) + P1(82);
+ W[83 - O] = P4(83) + P2(83) + P1(83);
+ W[84 - O] = P4(84) + P2(84) + P1(84);
+ W[85 - O] = P4(85) + P2(85) + P1(85);
+ W(86);
+
+ sharoundW(80);
+ sharoundW(81);
+ sharoundW(82);
+ sharoundW(83);
+ sharoundW(84);
+ sharoundW(85);
+ sharoundW(86);
+
+ W[87 - O] = (u)0x11002000 + P4(87) + P3(87) + P1(87);
+ sharoundW(87);
+ W[88 - O] = (u)0x80000000U + P3(88) + P1(88);
+ sharoundW(88);
+ W[89 - O] = P3(89) + P1(89);
+ sharoundW(89);
+ W[90 - O] = P3(90) + P1(90);
+ sharoundW(90);
+ W[91 - O] = P3(91) + P1(91);
+ sharoundW(91);
+ W[92 - O] = P3(92) + P1(92);
+ sharoundW(92);
+ W[93 - O] = P3(93) + P1(93);
+ sharoundW(93);
+ W[94 - O] = (u)0x400022 + P3(94) + P1(94);
+ sharoundW(94);
+ W[95 - O] = (u)0x00000100U + P3(95) + P2(95) + P1(95);
+ sharoundW(95);
+
+ W(96);
+ sharoundW(96);
+ W(97);
+ sharoundW(97);
+ W(98);
+ sharoundW(98);
+ W(99);
+ sharoundW(99);
+ W(100);
+ sharoundW(100);
+ W(101);
+ sharoundW(101);
+ W(102);
+ sharoundW(102);
+ W(103);
+ sharoundW(103);
+ W(104);
+ sharoundW(104);
+ W(105);
+ sharoundW(105);
+ W(106);
+ sharoundW(106);
+ W(107);
+ sharoundW(107);
+ W(108);
+ sharoundW(108);
+ W(109);
+ sharoundW(109);
+ W(110);
+ sharoundW(110);
+ W(111);
+ sharoundW(111);
+ W(112);
+ sharoundW(112);
+ W(113);
+ sharoundW(113);
+ W(114);
+ sharoundW(114);
+ W(115);
+ sharoundW(115);
+ W(116);
+ sharoundW(116);
+ W(117);
+ sharoundW(117);
+ W(118);
+ sharoundW(118);
+ W(119);
+ sharoundW(119);
+ W(120);
+ sharoundW(120);
+ W(121);
+ sharoundW(121);
+ W(122);
+ sharoundW(122);
+ W(123);
+ sharoundW(123);
+
+ // Round 124
+ Vals[7] += Vals[3] + P4(124) + P3(124) + P2(124) + P1(124) + s1(124) + ch(124);
+
+#ifdef VECTORS
+ if(Vals[7].x == -H[7])
+ {
+ output[OUTPUT_SIZE] = output[(W[3].x >> 2) & OUTPUT_MASK] = W_3.x;
+ }
+ if(Vals[7].y == -H[7])
+ {
+ output[OUTPUT_SIZE] = output[(W[3].y >> 2) & OUTPUT_MASK] = W_3.y;
+ }
+#else
+ if(Vals[7] == -H[7])
+ {
+ output[OUTPUT_SIZE] = output[(W[3] >> 2) & OUTPUT_MASK] = W_3;
+ }
+#endif
+}
122 kernels/phatk2/BFIPatcher.py
@@ -0,0 +1,122 @@
+# Copyright (C) 2011 by jedi95 <jedi95@gmail.com> and
+# CFSworks <CFSworks@gmail.com>
+#
+# 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.
+
+import struct
+
+class PatchError(Exception): pass
+
+class BFIPatcher(object):
+ """Patches .ELF files compiled for Evergreen GPUs; changes the microcode
+ so that any BYTE_ALIGN_INT instructions become BFI_INT.
+ """
+
+ def __init__(self, interface):
+ self.interface = interface
+
+ def patch(self, data):
+ """Run the process of patching an ELF."""
+
+ self.interface.debug('Finding inner ELF...')
+ innerPos = self.locateInner(data)
+ self.interface.debug('Patching inner ELF...')
+ inner = data[innerPos:]
+ patched = self.patchInner(inner)
+ self.interface.debug('Patch complete, returning to kernel...')
+ return data[:innerPos] + patched
+
+ def patchInner(self, data):
+ sections = self.readELFSections(data)
+ # We're looking for .text -- there should be two of them.
+ textSections = filter(lambda x: x[0] == '.text', sections)
+ if len(textSections) != 2:
+ self.interface.debug('Inner ELF does not have 2 .text sections!')
+ self.interface.debug('Sections are: %r' % sections)
+ raise PatchError()
+ name, offset, size = textSections[1]
+ before, text2, after = (data[:offset], data[offset:offset+size],
+ data[offset+size:])
+
+ self.interface.debug('Patching instructions...')
+ text2 = self.patchInstructions(text2)
+ return before + text2 + after
+
+ def patchInstructions(self, data):
+ output = ''
+ nPatched = 0
+ for i in xrange(len(data)/8):
+ inst, = struct.unpack('Q', data[i*8:i*8+8])
+ # Is it BYTE_ALIGN_INT?
+ if (inst&0x9003f00002001000) == 0x0001a00000000000:
+ nPatched += 1
+ inst ^= (0x0001a00000000000 ^ 0x0000c00000000000) # BFI_INT
+ output += struct.pack('Q', inst)
+ self.interface.debug('BFI-patched %d instructions...' % nPatched)
+ if nPatched < 60:
+ self.interface.debug('Patch safety threshold not met!')
+ raise PatchError()
+ return output
+
+ def locateInner(self, data):
+ """ATI uses an ELF-in-an-ELF. I don't know why. This function's job is
+ to find it.
+ """
+
+ pos = data.find('\x7fELF', 1)
+ if pos == -1 or data.find('\x7fELF', pos+1) != -1: # More than 1 is bad
+ self.interface.debug('Inner ELF not located!')
+ raise PatchError()
+ return pos
+
+ def readELFSections(self, data):
+ try:
+ (ident1, ident2, type, machine, version, entry, phoff,
+ shoff, flags, ehsize, phentsize, phnum, shentsize, shnum,
+ shstrndx) = struct.unpack('QQHHIIIIIHHHHHH', data[:52])
+
+ if ident1 != 0x64010101464c457f:
+ self.interface.debug('Invalid ELF header!')
+ raise PatchError()
+
+ # No section header?
+ if shoff == 0:
+ return []
+
+ # Find out which section contains the section header names
+ shstr = data[shoff+shstrndx*shentsize:shoff+(shstrndx+1)*shentsize]
+ (nameIdx, type, flags, addr, nameTableOffset, size, link, info,
+ addralign, entsize) = struct.unpack('IIIIIIIIII', shstr)
+
+ # Grab the section header.
+ sh = data[shoff:shoff+shnum*shentsize]
+
+ sections = []
+ for i in xrange(shnum):
+ rawEntry = sh[i*shentsize:(i+1)*shentsize]
+ (nameIdx, type, flags, addr, offset, size, link, info,
+ addralign, entsize) = struct.unpack('IIIIIIIIII', rawEntry)
+ nameOffset = nameTableOffset + nameIdx
+ name = data[nameOffset:data.find('\x00', nameOffset)]
+ sections.append((name, offset, size))
+
+ return sections
+ except struct.error:
+ self.interface.debug('A struct.error occurred while reading ELF!')
+ raise PatchError()
437 kernels/phatk2/__init__.py
@@ -0,0 +1,437 @@
+# Copyright (C) 2011 by jedi95 <jedi95@gmail.com> and
+# CFSworks <CFSworks@gmail.com>
+# Phateus <jesse.moll@gmail.com>
+#
+# 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.
+
+import pyopencl as cl
+import numpy as np
+import os
+import math
+
+from time import clock
+from hashlib import md5
+from struct import pack, unpack
+from twisted.internet import reactor
+
+from minerutil.Midstate import calculateMidstate
+from QueueReader import QueueReader
+from KernelInterface import *
+from BFIPatcher import *
+
+class KernelData(object):
+ #This class is a container for all the data required for a single kernel execution.
+
+ def __init__(self, nonceRange, core, rateDivisor, aggression):
+ # Prepare some raw data, converting it into the form that the OpenCL
+ # function expects.
+ data = np.array(
+ unpack('IIII', nonceRange.unit.data[64:]), dtype=np.uint32)
+
+ # get the number of iterations from the aggression and size
+ self.iterations = int((nonceRange.size / (1 << aggression)))
+ self.iterations = max(1, self.iterations)
+
+ #set the size to pass to the kernel based on iterations and vectors
+ self.size = (nonceRange.size / rateDivisor) / self.iterations
+ self.totalsize = nonceRange.size
+ #compute bases for each iteration
+
+ self.base = [None] * self.iterations
+ for i in range(self.iterations):
+ if rateDivisor == 1:
+ self.base[i] = pack('I',
+ ((nonceRange.base) + (i * self.size * rateDivisor)))
+ if rateDivisor == 2:
+ self.base[i] = pack('II',
+ ((nonceRange.base) + (i * self.size * rateDivisor))
+ , (1 + (nonceRange.base) + (i * self.size * rateDivisor)))
+ if rateDivisor == 4:
+ self.base[i] = pack('IIII',
+ ((nonceRange.base) + (i * self.size * rateDivisor))
+ , (1 + (nonceRange.base) + (i * self.size * rateDivisor))
+ , (2 + (nonceRange.base) + (i * self.size * rateDivisor))
+ , (3 + (nonceRange.base) + (i * self.size * rateDivisor))
+ )
+ #set up state and precalculated static data
+ self.state = np.array(
+ unpack('IIIIIIII', nonceRange.unit.midstate), dtype=np.uint32)
+ self.state2 = np.array(unpack('IIIIIIII',
+ calculateMidstate(nonceRange.unit.data[64:80] +
+ '\x00\x00\x00\x80' + '\x00'*40 + '\x80\x02\x00\x00',
+ nonceRange.unit.midstate, 3)), dtype=np.uint32)
+ self.state2 = np.array(
+ list(self.state2)[3:] + list(self.state2)[:3], dtype=np.uint32)
+ self.nr = nonceRange
+
+ self.f = np.zeros(9, np.uint32)
+ self.calculateF(data)
+
+ def calculateF(self, data):
+ rot = lambda x,y: x>>y | x<<(32-y)
+ #W2
+ self.f[0] = np.uint32(data[2])
+
+ #W16
+ W16 = np.uint32(data[0] + (rot(data[1], 7) ^ rot(data[1], 18) ^
+ (data[1] >> 3)))
+ self.f[1] = W16
+ #W17
+ W17 = np.uint32(data[1] + (rot(data[2], 7) ^ rot(data[2], 18) ^
+ (data[2] >> 3)) + 0x01100000)
+ self.f[2] = W17
+
+ #2 parts of the first SHA round
+ PreVal4 = (self.state[4] + (rot(self.state2[1], 6) ^
+ rot(self.state2[1], 11) ^ rot(self.state2[1], 25)) +
+ (self.state2[3] ^ (self.state2[1] & (self.state2[2] ^
+ self.state2[3]))) + 0xe9b5dba5)
+ T1 = ((rot(self.state2[5], 2) ^
+ rot(self.state2[5], 13) ^ rot(self.state2[5], 22)) +
+ ((self.state2[5] & self.state2[6]) | (self.state2[7] &
+ (self.state2[5] | self.state2[6]))))
+ self.f[3] = np.uint32(( PreVal4 + T1))
+ self.f[4] = np.uint32( PreVal4 + self.state[0])
+ self.f[5] = np.uint32(0x00000280 + ((rot(W16, 7) ^
+ rot(W16, 18) ^ (W16 >> 3))))
+ self.f[6] = np.uint32(self.f[1] + ((rot(W17, 7) ^
+ rot(W17, 18) ^ (W17 >> 3))))
+
+ self.f[7] = np.uint32(0x11002000 + (rot(W17, 17) ^ rot(W17, 19) ^
+ (W17 >> 10)))
+ self.f[8] = np.uint32(data[2] + (rot(W16, 17) ^ rot(W16, 19) ^
+ (W16 >> 10)))
+
+
+class MiningKernel(object):
+ #A Phoenix Miner-compatible OpenCL kernel created by Phateus
+
+ PLATFORM = KernelOption(
+ 'PLATFORM', int, default=None,
+ help='The ID of the OpenCL platform to use')
+ DEVICE = KernelOption(
+ 'DEVICE', int, default=None,
+ help='The ID of the OpenCL device to use')
+ VECTORS = KernelOption(
+ 'VECTORS', bool, default=False, advanced=True,
+ help='Enable vector support in the kernel?')
+ VECTORS4 = KernelOption(
+ 'VECTORS4', bool, default=False, advanced=True,
+ help='Enable vector uint4 support in the kernel?')
+ FASTLOOP = KernelOption(
+ 'FASTLOOP', bool, default=True, advanced=True,
+ help='Run iterative mining thread?')
+ AGGRESSION = KernelOption(
+ 'AGGRESSION', int, default=5, advanced=True,
+ help='Exponential factor indicating how much work to run '
+ 'per OpenCL execution')
+ WORKSIZE = KernelOption(
+ 'WORKSIZE', int, default=None, advanced=True,
+ help='The worksize to use when executing CL kernels.')
+ BFI_INT = KernelOption(
+ 'BFI_INT', bool, default=True, advanced=True,
+ help='Use the BFI_INT instruction for AMD/ATI GPUs.')
+ OUTPUT_SIZE = WORKSIZE
+
+ # This must be manually set for Git
+ REVISION = 117
+
+ def __init__(self, interface):
+ platforms = cl.get_platforms()
+
+ # Initialize object attributes and retrieve command-line options...)
+ self.device = None
+ self.kernel = None
+ self.interface = interface
+ self.core = self.interface.addCore()
+ self.defines = ''
+ self.loopExponent = 0
+
+ # Set the initial number of nonces to run per execution
+ # 2^(16 + aggression)
+ self.AGGRESSION += 16
+ self.AGGRESSION = min(32, self.AGGRESSION)
+ self.AGGRESSION = max(16, self.AGGRESSION)
+ self.size = 1 << self.AGGRESSION
+
+ # We need a QueueReader to efficiently provide our dedicated thread
+ # with work.
+ self.qr = QueueReader(self.core, lambda nr: self.preprocess(nr),
+ lambda x,y: self.size * 1 << self.loopExponent)
+
+ # The platform selection must be valid to mine.
+ if self.PLATFORM >= len(platforms) or \
+ (self.PLATFORM is None and len(platforms) > 1):
+ self.interface.log(
+ 'Wrong platform or more than one OpenCL platform found, '
+ 'use PLATFORM=ID to select one of the following\n',
+ False, True)
+
+ for i,p in enumerate(platforms):
+ self.interface.log(' [%d]\t%s' % (i, p.name), False, False)
+
+ # Since the platform is invalid, we can't mine.
+ self.interface.fatal()
+ return
+ elif self.PLATFORM is None:
+ self.PLATFORM = 0
+
+ devices = platforms[self.PLATFORM].get_devices()
+
+ # The device selection must be valid to mine.
+ if self.DEVICE >= len(devices) or \
+ (self.DEVICE is None and len(devices) > 1):
+ self.interface.log(
+ 'No device specified or device not found, '
+ 'use DEVICE=ID to specify one of the following\n',
+ False, True)
+
+ for i,d in enumerate(devices):
+ self.interface.log(' [%d]\t%s' % (i, d.name), False, False)
+
+ # Since the device selection is invalid, we can't mine.
+ self.interface.fatal()
+ return
+ elif self.DEVICE is None:
+ self.DEVICE = 0
+
+ self.device = devices[self.DEVICE]
+
+ # We need the appropriate kernel for this device...
+ try:
+ self.loadKernel(self.device)
+ except Exception:
+ self.interface.fatal("Failed to load OpenCL kernel!")
+ return
+
+ # Initialize a command queue to send commands to the device, and a
+ # buffer to collect results in...
+ self.commandQueue = cl.CommandQueue(self.context)
+ self.output = np.zeros(self.OUTPUT_SIZE+1, np.uint32)
+ self.output_buf = cl.Buffer(
+ self.context, cl.mem_flags.WRITE_ONLY | cl.mem_flags.USE_HOST_PTR,
+ hostbuf=self.output)
+
+ self.applyMeta()
+
+ def applyMeta(self):
+ #Apply any kernel-specific metadata.
+ self.interface.setMeta('kernel', 'phatk2 r%s' % self.REVISION)
+ self.interface.setMeta('device', self.device.name.replace('\x00',''))
+ self.interface.setMeta('cores', self.device.max_compute_units)
+
+ def loadKernel(self, device):
+ #Load the kernel and initialize the device.
+ self.context = cl.Context([device], None, None)
+
+ # If the user didn't specify their own worksize, use 256
+ if self.WORKSIZE is None:
+ self.WORKSIZE = 256
+ else:
+ #if the worksize is not a power of 2, round down to the nearest one
+ if (self.WORKSIZE & (self.WORKSIZE - 1)) != 0: