commit b6bf5a2ec75bab67162ca5bacf1eb342d923a5fe Author: tcatm Date: Sat Feb 5 00:44:17 2011 +0100 initial commit diff --git a/COPYING b/COPYING new file mode 100644 index 0000000..d60c31a --- /dev/null +++ b/COPYING @@ -0,0 +1,340 @@ + GNU GENERAL PUBLIC LICENSE + Version 2, June 1991 + + Copyright (C) 1989, 1991 Free Software Foundation, Inc. + 59 Temple Place, Suite 330, Boston, MA 02111-1307 USA + Everyone is permitted to copy and distribute verbatim copies + of this license document, but changing it is not allowed. + + Preamble + + The licenses for most software are designed to take away your +freedom to share and change it. By contrast, the GNU General Public +License is intended to guarantee your freedom to share and change free +software--to make sure the software is free for all its users. This +General Public License applies to most of the Free Software +Foundation's software and to any other program whose authors commit to +using it. (Some other Free Software Foundation software is covered by +the GNU Library General Public License instead.) You can apply it to +your programs, too. + + When we speak of free software, we are referring to freedom, not +price. Our General Public Licenses are designed to make sure that you +have the freedom to distribute copies of free software (and charge for +this service if you wish), that you receive source code or can get it +if you want it, that you can change the software or use pieces of it +in new free programs; and that you know you can do these things. + + To protect your rights, we need to make restrictions that forbid +anyone to deny you these rights or to ask you to surrender the rights. +These restrictions translate to certain responsibilities for you if you +distribute copies of the software, or if you modify it. + + For example, if you distribute copies of such a program, whether +gratis or for a fee, you must give the recipients all the rights that +you have. You must make sure that they, too, receive or can get the +source code. And you must show them these terms so they know their +rights. + + We protect your rights with two steps: (1) copyright the software, and +(2) offer you this license which gives you legal permission to copy, +distribute and/or modify the software. + + Also, for each author's protection and ours, we want to make certain +that everyone understands that there is no warranty for this free +software. If the software is modified by someone else and passed on, we +want its recipients to know that what they have is not the original, so +that any problems introduced by others will not reflect on the original +authors' reputations. + + Finally, any free program is threatened constantly by software +patents. We wish to avoid the danger that redistributors of a free +program will individually obtain patent licenses, in effect making the +program proprietary. To prevent this, we have made it clear that any +patent must be licensed for everyone's free use or not licensed at all. + + The precise terms and conditions for copying, distribution and +modification follow. + + GNU GENERAL PUBLIC LICENSE + TERMS AND CONDITIONS FOR COPYING, DISTRIBUTION AND MODIFICATION + + 0. This License applies to any program or other work which contains +a notice placed by the copyright holder saying it may be distributed +under the terms of this General Public License. The "Program", below, +refers to any such program or work, and a "work based on the Program" +means either the Program or any derivative work under copyright law: +that is to say, a work containing the Program or a portion of it, +either verbatim or with modifications and/or translated into another +language. (Hereinafter, translation is included without limitation in +the term "modification".) Each licensee is addressed as "you". + +Activities other than copying, distribution and modification are not +covered by this License; they are outside its scope. The act of +running the Program is not restricted, and the output from the Program +is covered only if its contents constitute a work based on the +Program (independent of having been made by running the Program). +Whether that is true depends on what the Program does. + + 1. You may copy and distribute verbatim copies of the Program's +source code as you receive it, in any medium, provided that you +conspicuously and appropriately publish on each copy an appropriate +copyright notice and disclaimer of warranty; keep intact all the +notices that refer to this License and to the absence of any warranty; +and give any other recipients of the Program a copy of this License +along with the Program. + +You may charge a fee for the physical act of transferring a copy, and +you may at your option offer warranty protection in exchange for a fee. + + 2. You may modify your copy or copies of the Program or any portion +of it, thus forming a work based on the Program, and copy and +distribute such modifications or work under the terms of Section 1 +above, provided that you also meet all of these conditions: + + a) You must cause the modified files to carry prominent notices + stating that you changed the files and the date of any change. + + b) You must cause any work that you distribute or publish, that in + whole or in part contains or is derived from the Program or any + part thereof, to be licensed as a whole at no charge to all third + parties under the terms of this License. + + c) If the modified program normally reads commands interactively + when run, you must cause it, when started running for such + interactive use in the most ordinary way, to print or display an + announcement including an appropriate copyright notice and a + notice that there is no warranty (or else, saying that you provide + a warranty) and that users may redistribute the program under + these conditions, and telling the user how to view a copy of this + License. (Exception: if the Program itself is interactive but + does not normally print such an announcement, your work based on + the Program is not required to print an announcement.) + +These requirements apply to the modified work as a whole. If +identifiable sections of that work are not derived from the Program, +and can be reasonably considered independent and separate works in +themselves, then this License, and its terms, do not apply to those +sections when you distribute them as separate works. But when you +distribute the same sections as part of a whole which is a work based +on the Program, the distribution of the whole must be on the terms of +this License, whose permissions for other licensees extend to the +entire whole, and thus to each and every part regardless of who wrote it. + +Thus, it is not the intent of this section to claim rights or contest +your rights to work written entirely by you; rather, the intent is to +exercise the right to control the distribution of derivative or +collective works based on the Program. + +In addition, mere aggregation of another work not based on the Program +with the Program (or with a work based on the Program) on a volume of +a storage or distribution medium does not bring the other work under +the scope of this License. + + 3. You may copy and distribute the Program (or a work based on it, +under Section 2) in object code or executable form under the terms of +Sections 1 and 2 above provided that you also do one of the following: + + a) Accompany it with the complete corresponding machine-readable + source code, which must be distributed under the terms of Sections + 1 and 2 above on a medium customarily used for software interchange; or, + + b) Accompany it with a written offer, valid for at least three + years, to give any third party, for a charge no more than your + cost of physically performing source distribution, a complete + machine-readable copy of the corresponding source code, to be + distributed under the terms of Sections 1 and 2 above on a medium + customarily used for software interchange; or, + + c) Accompany it with the information you received as to the offer + to distribute corresponding source code. (This alternative is + allowed only for noncommercial distribution and only if you + received the program in object code or executable form with such + an offer, in accord with Subsection b above.) + +The source code for a work means the preferred form of the work for +making modifications to it. For an executable work, complete source +code means all the source code for all modules it contains, plus any +associated interface definition files, plus the scripts used to +control compilation and installation of the executable. However, as a +special exception, the source code distributed need not include +anything that is normally distributed (in either source or binary +form) with the major components (compiler, kernel, and so on) of the +operating system on which the executable runs, unless that component +itself accompanies the executable. + +If distribution of executable or object code is made by offering +access to copy from a designated place, then offering equivalent +access to copy the source code from the same place counts as +distribution of the source code, even though third parties are not +compelled to copy the source along with the object code. + + 4. You may not copy, modify, sublicense, or distribute the Program +except as expressly provided under this License. Any attempt +otherwise to copy, modify, sublicense or distribute the Program is +void, and will automatically terminate your rights under this License. +However, parties who have received copies, or rights, from you under +this License will not have their licenses terminated so long as such +parties remain in full compliance. + + 5. You are not required to accept this License, since you have not +signed it. However, nothing else grants you permission to modify or +distribute the Program or its derivative works. These actions are +prohibited by law if you do not accept this License. Therefore, by +modifying or distributing the Program (or any work based on the +Program), you indicate your acceptance of this License to do so, and +all its terms and conditions for copying, distributing or modifying +the Program or works based on it. + + 6. Each time you redistribute the Program (or any work based on the +Program), the recipient automatically receives a license from the +original licensor to copy, distribute or modify the Program subject to +these terms and conditions. You may not impose any further +restrictions on the recipients' exercise of the rights granted herein. +You are not responsible for enforcing compliance by third parties to +this License. + + 7. If, as a consequence of a court judgment or allegation of patent +infringement or for any other reason (not limited to patent issues), +conditions are imposed on you (whether by court order, agreement or +otherwise) that contradict the conditions of this License, they do not +excuse you from the conditions of this License. If you cannot +distribute so as to satisfy simultaneously your obligations under this +License and any other pertinent obligations, then as a consequence you +may not distribute the Program at all. For example, if a patent +license would not permit royalty-free redistribution of the Program by +all those who receive copies directly or indirectly through you, then +the only way you could satisfy both it and this License would be to +refrain entirely from distribution of the Program. + +If any portion of this section is held invalid or unenforceable under +any particular circumstance, the balance of the section is intended to +apply and the section as a whole is intended to apply in other +circumstances. + +It is not the purpose of this section to induce you to infringe any +patents or other property right claims or to contest validity of any +such claims; this section has the sole purpose of protecting the +integrity of the free software distribution system, which is +implemented by public license practices. Many people have made +generous contributions to the wide range of software distributed +through that system in reliance on consistent application of that +system; it is up to the author/donor to decide if he or she is willing +to distribute software through any other system and a licensee cannot +impose that choice. + +This section is intended to make thoroughly clear what is believed to +be a consequence of the rest of this License. + + 8. If the distribution and/or use of the Program is restricted in +certain countries either by patents or by copyrighted interfaces, the +original copyright holder who places the Program under this License +may add an explicit geographical distribution limitation excluding +those countries, so that distribution is permitted only in or among +countries not thus excluded. In such case, this License incorporates +the limitation as if written in the body of this License. + + 9. The Free Software Foundation may publish revised and/or new versions +of the General Public License from time to time. Such new versions will +be similar in spirit to the present version, but may differ in detail to +address new problems or concerns. + +Each version is given a distinguishing version number. If the Program +specifies a version number of this License which applies to it and "any +later version", you have the option of following the terms and conditions +either of that version or of any later version published by the Free +Software Foundation. If the Program does not specify a version number of +this License, you may choose any version ever published by the Free Software +Foundation. + + 10. If you wish to incorporate parts of the Program into other free +programs whose distribution conditions are different, write to the author +to ask for permission. For software which is copyrighted by the Free +Software Foundation, write to the Free Software Foundation; we sometimes +make exceptions for this. Our decision will be guided by the two goals +of preserving the free status of all derivatives of our free software and +of promoting the sharing and reuse of software generally. + + NO WARRANTY + + 11. BECAUSE THE PROGRAM IS LICENSED FREE OF CHARGE, THERE IS NO WARRANTY +FOR THE PROGRAM, TO THE EXTENT PERMITTED BY APPLICABLE LAW. EXCEPT WHEN +OTHERWISE STATED IN WRITING THE COPYRIGHT HOLDERS AND/OR OTHER PARTIES +PROVIDE THE PROGRAM "AS IS" WITHOUT WARRANTY OF ANY KIND, EITHER EXPRESSED +OR IMPLIED, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF +MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE. THE ENTIRE RISK AS +TO THE QUALITY AND PERFORMANCE OF THE PROGRAM IS WITH YOU. SHOULD THE +PROGRAM PROVE DEFECTIVE, YOU ASSUME THE COST OF ALL NECESSARY SERVICING, +REPAIR OR CORRECTION. + + 12. IN NO EVENT UNLESS REQUIRED BY APPLICABLE LAW OR AGREED TO IN WRITING +WILL ANY COPYRIGHT HOLDER, OR ANY OTHER PARTY WHO MAY MODIFY AND/OR +REDISTRIBUTE THE PROGRAM AS PERMITTED ABOVE, BE LIABLE TO YOU FOR DAMAGES, +INCLUDING ANY GENERAL, SPECIAL, INCIDENTAL OR CONSEQUENTIAL DAMAGES ARISING +OUT OF THE USE OR INABILITY TO USE THE PROGRAM (INCLUDING BUT NOT LIMITED +TO LOSS OF DATA OR DATA BEING RENDERED INACCURATE OR LOSSES SUSTAINED BY +YOU OR THIRD PARTIES OR A FAILURE OF THE PROGRAM TO OPERATE WITH ANY OTHER +PROGRAMS), EVEN IF SUCH HOLDER OR OTHER PARTY HAS BEEN ADVISED OF THE +POSSIBILITY OF SUCH DAMAGES. + + END OF TERMS AND CONDITIONS + + How to Apply These Terms to Your New Programs + + If you develop a new program, and you want it to be of the greatest +possible use to the public, the best way to achieve this is to make it +free software which everyone can redistribute and change under these terms. + + To do so, attach the following notices to the program. It is safest +to attach them to the start of each source file to most effectively +convey the exclusion of warranty; and each file should have at least +the "copyright" line and a pointer to where the full notice is found. + + + Copyright (C) + + This program is free software; you can redistribute it and/or modify + it under the terms of the GNU General Public License as published by + the Free Software Foundation; either version 2 of the License, or + (at your option) any later version. + + This program is distributed in the hope that it will be useful, + but WITHOUT ANY WARRANTY; without even the implied warranty of + MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + GNU General Public License for more details. + + You should have received a copy of the GNU General Public License + along with this program; if not, write to the Free Software + Foundation, Inc., 59 Temple Place, Suite 330, Boston, MA 02111-1307 USA + + +Also add information on how to contact you by electronic and paper mail. + +If the program is interactive, make it output a short notice like this +when it starts in an interactive mode: + + Gnomovision version 69, Copyright (C) year name of author + Gnomovision comes with ABSOLUTELY NO WARRANTY; for details type `show w'. + This is free software, and you are welcome to redistribute it + under certain conditions; type `show c' for details. + +The hypothetical commands `show w' and `show c' should show the appropriate +parts of the General Public License. Of course, the commands you use may +be called something other than `show w' and `show c'; they could even be +mouse-clicks or menu items--whatever suits your program. + +You should also get your employer (if you work as a programmer) or your +school, if any, to sign a "copyright disclaimer" for the program, if +necessary. Here is a sample; alter the names: + + Yoyodyne, Inc., hereby disclaims all copyright interest in the program + `Gnomovision' (which makes passes at compilers) written by James Hacker. + + , 1 April 1989 + Ty Coon, President of Vice + +This General Public License does not permit incorporating your program into +proprietary programs. If your program is a subroutine library, you may +consider it more useful to permit linking proprietary applications with the +library. If this is what you want to do, use the GNU Library General +Public License instead of this License. diff --git a/LICENSE b/LICENSE new file mode 100644 index 0000000..cbcd60f --- /dev/null +++ b/LICENSE @@ -0,0 +1,3 @@ +oclminer is available under the terms of the GNU Public License version 2. + +See COPYING for details. diff --git a/Makefile b/Makefile new file mode 100644 index 0000000..3821a9c --- /dev/null +++ b/Makefile @@ -0,0 +1,21 @@ +INCLUDEPATHS=-I${ATISTREAMSDKROOT}/include/ -L${ATISTREAMSDKROOT}/ + +LIBS=-lOpenCL -ljansson -lcurl + +DEFS= +DEBUGFLAGS=-g +CFLAGS=-O3 -Wformat $(DEBUGFLAGS) $(DEFS) $(INCLUDEPATHS) +HEADERS= + +OBJS=miner.o ocl.o findnonce.o util.o + +all: oclminer + +%.o: %.c $(HEADERS) + gcc -c $(CFLAGS) -o $@ $< + +oclminer: $(OBJS) + gcc $(CFLAGS) -o $@ $^ $(LIBS) + +clean: + -rm *.o oclminer diff --git a/findnonce.c b/findnonce.c new file mode 100644 index 0000000..04baf8d --- /dev/null +++ b/findnonce.c @@ -0,0 +1,187 @@ +/* + * Copyright 2011 Nils Schneider + * + * This program is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License as published by the Free + * Software Foundation; either version 2 of the License, or (at your option) + * any later version. See COPYING for more details. + */ + +#include +#include + +#include "ocl.h" +#include "findnonce.h" + +int fDebug = 0; + +const uint32_t SHA256_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 +}; + +inline uint32_t ByteReverse(uint32_t value) +{ + __asm__ ("bswap %0" : "=r" (value) : "0" (value)); + return value; +} + +#define rotate(x,y) ((x<>(sizeof(x)*8-y))) +#define rotr(x,y) ((x>>y) | (x<<(sizeof(x)*8-y))) + +#define R(a, b, c, d, e, f, g, h, w, k) \ + h = h + (rotate(e, 26) ^ rotate(e, 21) ^ rotate(e, 7)) + (g ^ (e & (f ^ g))) + k + w; \ + d = d + h; \ + h = h + (rotate(a, 30) ^ rotate(a, 19) ^ rotate(a, 10)) + ((a & b) | (c & (a | b))) + +void precalc_hash(dev_blk_ctx *blk, uint32_t *state, uint32_t *data) { + cl_uint A, B, C, D, E, F, G, H; + cl_uint W[16]; + + A = state[0]; + B = state[1]; + C = state[2]; + D = state[3]; + E = state[4]; + F = state[5]; + G = state[6]; + H = state[7]; + + R(A, B, C, D, E, F, G, H, data[0], SHA256_K[0]); + R(H, A, B, C, D, E, F, G, data[1], SHA256_K[1]); + R(G, H, A, B, C, D, E, F, data[2], SHA256_K[2]); + + blk->cty_a = A; + blk->cty_b = B; + blk->cty_c = C; + blk->cty_d = D; + blk->cty_e = E; + blk->cty_f = F; + blk->cty_g = G; + blk->cty_h = H; + + blk->ctx_a = state[0]; + blk->ctx_b = state[1]; + blk->ctx_c = state[2]; + blk->ctx_d = state[3]; + blk->ctx_e = state[4]; + blk->ctx_f = state[5]; + blk->ctx_g = state[6]; + blk->ctx_h = state[7]; + + blk->merkle = data[0]; + blk->ntime = data[1]; + blk->nbits = data[2]; + + blk->fW0 = data[0] + (rotr(data[1], 7) ^ rotr(data[1], 18) ^ (data[1] >> 3)); + blk->fW1 = data[1] + (rotr(data[2], 7) ^ rotr(data[2], 18) ^ (data[2] >> 3)) + 0x01100000; + blk->fW2 = data[2] + (rotr(blk->fW0, 17) ^ rotr(blk->fW0, 19) ^ (blk->fW0 >> 10)); + blk->fW3 = 0x11002000 + (rotr(blk->fW1, 17) ^ rotr(blk->fW1, 19) ^ (blk->fW1 >> 10)); + blk->fW15 = 0x00000280 + (rotr(blk->fW0, 7) ^ rotr(blk->fW0, 18) ^ (blk->fW0 >> 3)); + blk->fW01r = blk->fW0 + (rotr(blk->fW1, 7) ^ rotr(blk->fW1, 18) ^ (blk->fW1 >> 3)); + + blk->fcty_e = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + 0xe9b5dba5; + blk->fcty_e2 = (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G))); +} + +#define P(t) (W[(t)&0xF] = W[(t-16)&0xF] + (rotate(W[(t-15)&0xF], 25) ^ rotate(W[(t-15)&0xF], 14) ^ (W[(t-15)&0xF] >> 3)) + W[(t-7)&0xF] + (rotate(W[(t-2)&0xF], 15) ^ rotate(W[(t-2)&0xF], 13) ^ (W[(t-2)&0xF] >> 10))) + +#define IR(u) \ + R(A, B, C, D, E, F, G, H, W[u+0], SHA256_K[u+0]); \ + R(H, A, B, C, D, E, F, G, W[u+1], SHA256_K[u+1]); \ + R(G, H, A, B, C, D, E, F, W[u+2], SHA256_K[u+2]); \ + R(F, G, H, A, B, C, D, E, W[u+3], SHA256_K[u+3]); \ + R(E, F, G, H, A, B, C, D, W[u+4], SHA256_K[u+4]); \ + R(D, E, F, G, H, A, B, C, W[u+5], SHA256_K[u+5]); \ + R(C, D, E, F, G, H, A, B, W[u+6], SHA256_K[u+6]); \ + R(B, C, D, E, F, G, H, A, W[u+7], SHA256_K[u+7]) +#define FR(u) \ + R(A, B, C, D, E, F, G, H, P(u+0), SHA256_K[u+0]); \ + R(H, A, B, C, D, E, F, G, P(u+1), SHA256_K[u+1]); \ + R(G, H, A, B, C, D, E, F, P(u+2), SHA256_K[u+2]); \ + R(F, G, H, A, B, C, D, E, P(u+3), SHA256_K[u+3]); \ + R(E, F, G, H, A, B, C, D, P(u+4), SHA256_K[u+4]); \ + R(D, E, F, G, H, A, B, C, P(u+5), SHA256_K[u+5]); \ + R(C, D, E, F, G, H, A, B, P(u+6), SHA256_K[u+6]); \ + R(B, C, D, E, F, G, H, A, P(u+7), SHA256_K[u+7]) + +#define PIR(u) \ + R(F, G, H, A, B, C, D, E, W[u+3], SHA256_K[u+3]); \ + R(E, F, G, H, A, B, C, D, W[u+4], SHA256_K[u+4]); \ + R(D, E, F, G, H, A, B, C, W[u+5], SHA256_K[u+5]); \ + R(C, D, E, F, G, H, A, B, W[u+6], SHA256_K[u+6]); \ + R(B, C, D, E, F, G, H, A, W[u+7], SHA256_K[u+7]) + +#define PFR(u) \ + R(A, B, C, D, E, F, G, H, P(u+0), SHA256_K[u+0]); \ + R(H, A, B, C, D, E, F, G, P(u+1), SHA256_K[u+1]); \ + R(G, H, A, B, C, D, E, F, P(u+2), SHA256_K[u+2]); \ + R(F, G, H, A, B, C, D, E, P(u+3), SHA256_K[u+3]); \ + R(E, F, G, H, A, B, C, D, P(u+4), SHA256_K[u+4]); \ + R(D, E, F, G, H, A, B, C, P(u+5), SHA256_K[u+5]) + +uint32_t postcalc_hash(dev_blk_ctx *blk, uint32_t start, uint32_t end, uint32_t *best_nonce) { + cl_uint A, B, C, D, E, F, G, H; + cl_uint W[16]; + cl_uint nonce; + cl_uint best_g = ~0; + + for(nonce = start; nonce != end; nonce+=1) { + A = blk->cty_a; B = blk->cty_b; + C = blk->cty_c; D = blk->cty_d; + E = blk->cty_e; F = blk->cty_f; + G = blk->cty_g; H = blk->cty_h; + W[0] = blk->merkle; W[1] = blk->ntime; + W[2] = blk->nbits; W[3] = nonce;; + W[4] = 0x80000000; W[5] = 0x00000000; W[6] = 0x00000000; W[7] = 0x00000000; + W[8] = 0x00000000; W[9] = 0x00000000; W[10] = 0x00000000; W[11] = 0x00000000; + W[12] = 0x00000000; W[13] = 0x00000000; W[14] = 0x00000000; W[15] = 0x00000280; + PIR(0); IR(8); + FR(16); FR(24); + FR(32); FR(40); + FR(48); FR(56); + + W[0] = A + blk->ctx_a; W[1] = B + blk->ctx_b; + W[2] = C + blk->ctx_c; W[3] = D + blk->ctx_d; + W[4] = E + blk->ctx_e; W[5] = F + blk->ctx_f; + W[6] = G + blk->ctx_g; W[7] = H + blk->ctx_h; + W[8] = 0x80000000; W[9] = 0x00000000; W[10] = 0x00000000; W[11] = 0x00000000; + W[12] = 0x00000000; W[13] = 0x00000000; W[14] = 0x00000000; W[15] = 0x00000100; + A = 0x6a09e667; B = 0xbb67ae85; + C = 0x3c6ef372; D = 0xa54ff53a; + E = 0x510e527f; F = 0x9b05688c; + G = 0x1f83d9ab; H = 0x5be0cd19; + IR(0); IR(8); + FR(16); FR(24); + FR(32); FR(40); + FR(48); PFR(56); + + if(H == 0xA41F32E7) { + G += 0x1f83d9ab; + G = ByteReverse(G); + + if(G < best_g) { + *best_nonce = nonce; + best_g = G; + } + } + } + + if(best_g == ~0) printf("No best_g found! Error in OpenCL code?\n"); + + return best_g; +} diff --git a/findnonce.h b/findnonce.h new file mode 100644 index 0000000..c7028aa --- /dev/null +++ b/findnonce.h @@ -0,0 +1,17 @@ +#define MAXTHREADS 2000000 + +#include + +typedef struct { + cl_uint ctx_a; cl_uint ctx_b; cl_uint ctx_c; cl_uint ctx_d; + cl_uint ctx_e; cl_uint ctx_f; cl_uint ctx_g; cl_uint ctx_h; + cl_uint cty_a; cl_uint cty_b; cl_uint cty_c; cl_uint cty_d; + cl_uint cty_e; cl_uint cty_f; cl_uint cty_g; cl_uint cty_h; + cl_uint merkle; cl_uint ntime; cl_uint nbits; cl_uint nonce; + cl_uint fW0; cl_uint fW1; cl_uint fW2; cl_uint fW3; cl_uint fW15; + cl_uint fW01r; cl_uint fcty_e; cl_uint fcty_e2; +} dev_blk_ctx; + +extern void precalc_hash(dev_blk_ctx *blk, uint32_t *state, uint32_t *data); +extern uint32_t postcalc_hash(dev_blk_ctx *blk, uint32_t start, uint32_t end, uint32_t *best_nonce); + diff --git a/miner.c b/miner.c new file mode 100644 index 0000000..f056e33 --- /dev/null +++ b/miner.c @@ -0,0 +1,506 @@ + +/* + * Copyright 2010 Jeff Garzik + * 2011 Nils Schneider + * + * This program is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License as published by the Free + * Software Foundation; either version 2 of the License, or (at your option) + * any later version. See COPYING for more details. + */ + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include "miner.h" + +#include "findnonce.h" +#include "ocl.h" + +#define VERSION "0.1" + +#define PROGRAM_NAME "oclminer" +#define DEF_RPC_URL "http://127.0.0.1:8332/" +#define DEF_RPC_USERPASS "rpcuser:rpcpass" + +enum { + STAT_SLEEP_INTERVAL = 100, + STAT_CTR_INTERVAL = 10000000, + FAILURE_INTERVAL = 30, +}; + +int opt_debug = false; +int opt_protocol = false; +int opt_ndevs = false; +static int opt_retries = 10; +static bool program_running = true; +static const bool opt_time = true; +static int opt_n_threads = 1; +static char *rpc_url = DEF_RPC_URL; +static char *userpass = DEF_RPC_USERPASS; + + +struct option_help { + const char *name; + const char *helptext; +}; + +static struct option_help options_help[] = { + { "help", + "(-h) Display this help text" }, + + { "ndevs", + "(-n) Display number of detected GPUs" }, + + { "debug", + "(-D) Enable debug output (default: off)" }, + + { "protocol-dump", + "(-P) Verbose dump of protocol-level activities (default: off)" }, + + { "retries N", + "(-r N) Number of times to retry, if JSON-RPC call fails\n" + "\t(default: 10; use -1 for \"never\")" }, + + { "url URL", + "URL for bitcoin JSON-RPC server " + "(default: " DEF_RPC_URL ")" }, + + { "userpass USERNAME:PASSWORD", + "Username:Password pair for bitcoin JSON-RPC server " + "(default: " DEF_RPC_USERPASS ")" }, +}; + +static struct option options[] = { + { "help", 0, NULL, 'h' }, + { "debug", 0, NULL, 'D' }, + { "protocol-dump", 0, NULL, 'P' }, + { "retries", 1, NULL, 'r' }, + { "url", 1, NULL, 1001 }, + { "userpass", 1, NULL, 1002 }, + { "ndevs", 0, NULL, 'n' }, + { } +}; + +struct work_t { + unsigned char data[128]; + unsigned char hash1[64]; + unsigned char midstate[32]; + unsigned char target[32]; + + unsigned char hash[32]; + uint32_t output[MAXTHREADS]; + uint32_t res_nonce; + bool valid; + uint32_t ready; + dev_blk_ctx blk; +}; + +static bool jobj_binary(const json_t *obj, const char *key, + void *buf, size_t buflen) +{ + const char *hexstr; + json_t *tmp; + + tmp = json_object_get(obj, key); + if (!tmp) { + fprintf(stderr, "JSON key '%s' not found\n", key); + return false; + } + hexstr = json_string_value(tmp); + if (!hexstr) { + fprintf(stderr, "JSON key '%s' is not a string\n", key); + return false; + } + if (!hex2bin(buf, hexstr, buflen)) + return false; + + return true; +} + +static bool work_decode(const json_t *val, struct work_t *work) +{ + if (!jobj_binary(val, "midstate", + work->midstate, sizeof(work->midstate))) { + fprintf(stderr, "JSON inval midstate\n"); + goto err_out; + } + + if (!jobj_binary(val, "data", work->data, sizeof(work->data))) { + fprintf(stderr, "JSON inval data\n"); + goto err_out; + } + + if (!jobj_binary(val, "hash1", work->hash1, sizeof(work->hash1))) { + fprintf(stderr, "JSON inval hash1\n"); + goto err_out; + } + + if (!jobj_binary(val, "target", work->target, sizeof(work->target))) { + fprintf(stderr, "JSON inval target\n"); + goto err_out; + } + + memset(work->hash, 0, sizeof(work->hash)); + + return true; + +err_out: + return false; +} + +static void submit_work(struct work_t *work) +{ + char *hexstr = NULL; + json_t *val, *res; + char s[345]; + + printf("PROOF OF WORK FOUND? submitting...\n"); + + /* build hex string */ + hexstr = bin2hex(work->data, sizeof(work->data)); + if (!hexstr) + goto out; + + /* build JSON-RPC request */ + sprintf(s, + "{\"method\": \"getwork\", \"params\": [ \"%s\" ], \"id\":1}\r\n", + hexstr); + + if (opt_debug) + fprintf(stderr, "DBG: sending RPC call:\n%s", s); + + /* issue JSON-RPC request */ + val = json_rpc_call(rpc_url, userpass, s); + if (!val) { + fprintf(stderr, "submit_work json_rpc_call failed\n"); + goto out; + } + + res = json_object_get(val, "result"); + + printf("PROOF OF WORK RESULT: %s\n", + json_is_true(res) ? "true (yay!!!)" : "false (booooo)"); + + json_decref(val); + +out: + free(hexstr); +} + +static void hashmeter(int thr_id, struct timeval *tv_start, + unsigned long hashes_done) +{ + struct timeval tv_end, diff; + double khashes, secs; + + gettimeofday(&tv_end, NULL); + + timeval_subtract(&diff, &tv_end, tv_start); + + khashes = hashes_done / 1000.0; + secs = (double)diff.tv_sec + ((double)diff.tv_usec / 1000000.0); + + printf("HashMeter(%d): %lu hashes, %.2f khash/sec\n", + thr_id, hashes_done, + khashes / secs); +} + +static bool getwork(struct work_t *work) { + static const char *rpc_req = "{\"method\": \"getwork\", \"params\": [], \"id\":0}\r\n"; + json_t *val; + bool rc; + + /* obtain new work from bitcoin */ + val = json_rpc_call(rpc_url, userpass, rpc_req); + if (!val) { + fprintf(stderr, "json_rpc_call failed, "); + return false; + } + + /* decode result into work state struct */ + rc = work_decode(json_object_get(val, "result"), work); + if (!rc) { + fprintf(stderr, "JSON-decode of work failed, "); + return false; + } + + json_decref(val); + + return true; +} + +static void *miner_thread(void *thr_id_int) +{ + int thr_id = (unsigned long) thr_id_int; + int failures = 0; + + uint32_t res[MAXTHREADS]; + + size_t globalThreads[1]; + size_t localThreads[1]; + + cl_int status; + + _clState *clState; + + char name[32]; + + printf("Init GPU %i\n", thr_id); + clState = initCl(thr_id, name, sizeof(name)); + printf("initCl() finished. Found %s\n", name); + + status = clSetKernelArg(clState->kernel, 0, sizeof(cl_mem), (void *)&clState->inputBuffer); + if(status != CL_SUCCESS) { printf("Error: Setting kernel argument 1.\n"); return false; } + + status = clSetKernelArg(clState->kernel, 1, sizeof(cl_mem), (void *)&clState->outputBuffer); + if(status != CL_SUCCESS) { printf("Error: Setting kernel argument 2.\n"); return false; } + + struct work_t *work; + work = malloc(sizeof(struct work_t)*2); + + work[0].ready = 0; + work[1].ready = 0; + + int frame = 0; + int res_frame = 0; + bool need_work = true; + while (1) { + unsigned long hashes_done; + struct timeval tv_start; + bool rc; + + if (need_work) { + frame++; + frame %= 2; + + if (opt_debug) + fprintf(stderr, "getwork\n"); + + rc = getwork(&work[frame]); + + if (!rc) { + fprintf(stderr, "getwork failed, "); + + if ((opt_retries >= 0) && (++failures > opt_retries)) { + fprintf(stderr, "terminating thread\n"); + return NULL; /* exit thread */ + } + + /* pause, then restart work loop */ + fprintf(stderr, "retry after %d seconds\n", FAILURE_INTERVAL); + sleep(FAILURE_INTERVAL); + continue; + } + + precalc_hash(&work[frame].blk, (uint32_t *)(work[frame].midstate), (uint32_t *)(work[frame].data + 64)); + + work[frame].blk.nonce = 0; + work[frame].valid = true; + work[frame].ready = 0; + + need_work = false; + } + + hashes_done = 0; + gettimeofday(&tv_start, NULL); + + int threads = 102400; + globalThreads[0] = threads; + localThreads[0] = 128; + + printf("%i\n", work[frame].blk.nonce); + + status = clEnqueueWriteBuffer(clState->commandQueue, clState->inputBuffer, CL_TRUE, 0, + sizeof(dev_blk_ctx), (void *)&work[frame].blk, 0, NULL, NULL); + if(status != CL_SUCCESS) { printf("Error: clEnqueueWriteBuffer failed.\n"); return 0; } + + status = clEnqueueNDRangeKernel(clState->commandQueue, clState->kernel, 1, NULL, + globalThreads, localThreads, 0, NULL, NULL); + if(status != CL_SUCCESS) { printf("Error: Enqueueing kernel onto command queue. (clEnqueueNDRangeKernel)\n"); return 0; } + + clFlush(clState->commandQueue); + + hashes_done = 1024 * threads; + + if (work[res_frame].ready) { + rc = false; + + uint32_t bestG = ~0; + uint32_t nonce; + int j; + for(j = 0; j < work[res_frame].ready; j++) { + if(res[j]) { + uint32_t start = (work[res_frame].res_nonce + j)<<10; + uint32_t my_g, my_nonce; + my_g = postcalc_hash(&work[res_frame].blk, start, start + 1024, &my_nonce); + if (opt_debug) + fprintf(stderr, "DEBUG: H0 within %u .. %u, best G = %08x, nonce = %08x\n", start, start + 1024, my_g, my_nonce); + if(my_g < bestG) { + bestG = my_g; + nonce = my_nonce; + if (opt_debug) + fprintf(stderr, "new best\n"); + } + + rc = true; + } + } + + work[res_frame].ready = false; + + uint32_t *target = (uint32_t *)(work[res_frame].target + 24); + + if(rc && bestG <= *target) { + printf("Found solution for %08x: %08x %u\n", *target, bestG, nonce); + + work[res_frame].data[64+12+0] = (nonce>>0) & 0xff; + work[res_frame].data[64+12+1] = (nonce>>8) & 0xff; + work[res_frame].data[64+12+2] = (nonce>>16) & 0xff; + work[res_frame].data[64+12+3] = (nonce>>24) & 0xff; + submit_work(&work[res_frame]); + + need_work = true; + } + } + + clFinish(clState->commandQueue); + + status = clEnqueueReadBuffer(clState->commandQueue, clState->outputBuffer, CL_TRUE, 0, + sizeof(uint32_t) * threads, res, 0, NULL, NULL); + if(status != CL_SUCCESS) { printf("Error: clEnqueueReadBuffer failed. (clEnqueueReadBuffer)\n"); return 0; } + + clFinish(clState->commandQueue); + + res_frame = frame; + work[res_frame].ready = threads; + work[res_frame].res_nonce = work[res_frame].blk.nonce; + + work[frame].blk.nonce += threads; + + if (work[frame].blk.nonce > 4000000 - threads) + need_work = true; + + hashmeter(thr_id, &tv_start, hashes_done); + + failures = 0; + } + + return NULL; +} + +static void show_usage(void) +{ + int i; + + printf("oclminer version %s\n\n", VERSION); + printf("Usage:\tminerd [options]\n\nSupported options:\n"); + for (i = 0; i < ARRAY_SIZE(options_help); i++) { + struct option_help *h; + + h = &options_help[i]; + printf("--%s\n%s\n\n", h->name, h->helptext); + } + + exit(1); +} + +static void parse_arg (int key, char *arg) +{ + int v, i; + + switch(key) { + case 'n': + opt_ndevs = true; + break; + case 'D': + opt_debug = true; + break; + case 'P': + opt_protocol = true; + break; + case 'r': + v = atoi(arg); + if (v < -1 || v > 9999) /* sanity check */ + show_usage(); + + opt_retries = v; + break; + case 1001: /* --url */ + if (strncmp(arg, "http://", 7) && + strncmp(arg, "https://", 8)) + show_usage(); + + rpc_url = arg; + break; + case 1002: /* --userpass */ + if (!strchr(arg, ':')) + show_usage(); + + userpass = arg; + break; + default: + show_usage(); + } +} + +static void parse_cmdline(int argc, char *argv[]) +{ + int key; + + while (1) { + key = getopt_long(argc, argv, "DPh?n", options, NULL); + if (key < 0) + break; + + parse_arg(key, optarg); + } +} + +int main (int argc, char *argv[]) +{ + int i, nDevs; + + /* parse command line */ + parse_cmdline(argc, argv); + + nDevs = clDevicesNum(); + + if (opt_ndevs) { + printf("%i\n", nDevs); + return nDevs; + } + + /* start mining threads */ + for (i = 0; i < opt_n_threads; i++) { + pthread_t t; + + if (pthread_create(&t, NULL, miner_thread, + (void *)(unsigned long) i)) { + fprintf(stderr, "thread %d create failed\n", i); + return 1; + } + + sleep(1); /* don't pound RPC server all at once */ + } + + fprintf(stderr, "%d miner threads started\n", i); + + /* main loop */ + while (program_running) { + sleep(STAT_SLEEP_INTERVAL); + /* do nothing */ + } + + return 0; +} + diff --git a/miner.h b/miner.h new file mode 100644 index 0000000..636ca29 --- /dev/null +++ b/miner.h @@ -0,0 +1,18 @@ +#ifndef __MINER_H__ +#define __MINER_H__ +#include "ocl.h" + +#ifndef ARRAY_SIZE +#define ARRAY_SIZE(arr) (sizeof(arr) / sizeof((arr)[0])) +#endif + +extern int opt_protocol; + +extern json_t *json_rpc_call(const char *url, const char *userpass, + const char *rpc_req); +extern char *bin2hex(unsigned char *p, size_t len); +extern int hex2bin(unsigned char *p, const char *hexstr, size_t len); + +extern int timeval_subtract (struct timeval *result, struct timeval *x, struct timeval *y); + +#endif /* __MINER_H__ */ diff --git a/ocl.c b/ocl.c new file mode 100644 index 0000000..1a4e51c --- /dev/null +++ b/ocl.c @@ -0,0 +1,271 @@ +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "findnonce.h" +#include "ocl.h" + +char *file_contents(const char *filename, int *length) +{ + FILE *f = fopen(filename, "r"); + void *buffer; + + if (!f) { + fprintf(stderr, "Unable to open %s for reading\n", filename); + return NULL; + } + + fseek(f, 0, SEEK_END); + *length = ftell(f); + fseek(f, 0, SEEK_SET); + + buffer = malloc(*length+1); + *length = fread(buffer, 1, *length, f); + fclose(f); + ((char*)buffer)[*length] = '\0'; + + return (char*)buffer; +} + +int clDevicesNum() { + cl_int status = 0; + + cl_uint numPlatforms; + cl_platform_id platform = NULL; + status = clGetPlatformIDs(0, NULL, &numPlatforms); + if(status != CL_SUCCESS) + { + printf("Error: Getting Platforms. (clGetPlatformsIDs)\n"); + return -1; + } + + if(numPlatforms > 0) + { + cl_platform_id* platforms = (cl_platform_id *)malloc(numPlatforms*sizeof(cl_platform_id)); + status = clGetPlatformIDs(numPlatforms, platforms, NULL); + if(status != CL_SUCCESS) + { + printf("Error: Getting Platform Ids. (clGetPlatformsIDs)\n"); + return -1; + } + + unsigned int i; + for(i=0; i < numPlatforms; ++i) + { + char pbuff[100]; + status = clGetPlatformInfo( platforms[i], CL_PLATFORM_VENDOR, sizeof(pbuff), pbuff, NULL); + if(status != CL_SUCCESS) + { + printf("Error: Getting Platform Info. (clGetPlatformInfo)\n"); + free(platforms); + return -1; + } + platform = platforms[i]; + if(!strcmp(pbuff, "Advanced Micro Devices, Inc.")) + { + break; + } + } + free(platforms); + } + + if(platform == NULL) { + perror("NULL platform found!\n"); + return -1; + } + + cl_uint numDevices; + status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &numDevices); + if(status != CL_SUCCESS) + { + printf("Error: Getting Device IDs (num)\n"); + return -1; + } + + return numDevices; +} + +_clState *initCl(int gpu, char *name, size_t nameSize) { + cl_int status = 0; + + _clState *clState = malloc(sizeof(_clState));; + + cl_uint numPlatforms; + cl_platform_id platform = NULL; + status = clGetPlatformIDs(0, NULL, &numPlatforms); + if(status != CL_SUCCESS) + { + printf("Error: Getting Platforms. (clGetPlatformsIDs)\n"); + return NULL; + } + + if(numPlatforms > 0) + { + cl_platform_id* platforms = (cl_platform_id *)malloc(numPlatforms*sizeof(cl_platform_id)); + status = clGetPlatformIDs(numPlatforms, platforms, NULL); + if(status != CL_SUCCESS) + { + printf("Error: Getting Platform Ids. (clGetPlatformsIDs)\n"); + return NULL; + } + + unsigned int i; + for(i=0; i < numPlatforms; ++i) + { + char pbuff[100]; + status = clGetPlatformInfo( platforms[i], CL_PLATFORM_VENDOR, sizeof(pbuff), pbuff, NULL); + if(status != CL_SUCCESS) + { + printf("Error: Getting Platform Info. (clGetPlatformInfo)\n"); + free(platforms); + return NULL; + } + platform = platforms[i]; + if(!strcmp(pbuff, "Advanced Micro Devices, Inc.")) + { + break; + } + } + free(platforms); + } + + if(platform == NULL) { + perror("NULL platform found!\n"); + return NULL; + } + + cl_uint numDevices; + status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &numDevices); + if(status != CL_SUCCESS) + { + printf("Error: Getting Device IDs (num)\n"); + return NULL; + } + + cl_device_id *devices; + if(numDevices > 0 ) { + devices = (cl_device_id *)malloc(numDevices*sizeof(cl_device_id)); + + /* Now, get the device list data */ + + status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, numDevices, devices, NULL); + if(status != CL_SUCCESS) + { + printf("Error: Getting Device IDs (list)\n"); + return NULL; + } + + printf("List of devices:\n"); + + int i; + for(i=0; i= 0 && gpu < numDevices) { + char pbuff[100]; + status = clGetDeviceInfo(devices[gpu], CL_DEVICE_NAME, sizeof(pbuff), pbuff, NULL); + if(status != CL_SUCCESS) + { + printf("Error: Getting Device Info\n"); + return NULL; + } + + printf("Selected %i: %s\n", gpu, pbuff); + strncpy(name, pbuff, nameSize); + } else { + printf("Invalid GPU %i\n", gpu); + return NULL; + } + + } else return NULL; + + cl_context_properties cps[3] = { CL_CONTEXT_PLATFORM, (cl_context_properties)platform, 0 }; + + clState->context = clCreateContextFromType(cps, CL_DEVICE_TYPE_GPU, NULL, NULL, &status); + if(status != CL_SUCCESS) + { + printf("Error: Creating Context. (clCreateContextFromType)\n"); + return NULL; + } + + + ///////////////////////////////////////////////////////////////// + // Load CL file, build CL program object, create CL kernel object + ///////////////////////////////////////////////////////////////// + // + const char * filename = "oclminer.cl"; + int pl; + char *source = file_contents(filename, &pl); + size_t sourceSize[] = {(size_t)pl}; + + clState->program = clCreateProgramWithSource(clState->context, 1, (const char **)&source, sourceSize, &status); + if(status != CL_SUCCESS) + { + printf("Error: Loading Binary into cl_program (clCreateProgramWithBinary)\n"); + return NULL; + } + + /* create a cl program executable for all the devices specified */ + status = clBuildProgram(clState->program, 1, &devices[gpu], NULL, NULL, NULL); + if(status != CL_SUCCESS) + { + printf("Error: Building Program (clBuildProgram)\n"); + size_t logSize; + status = clGetProgramBuildInfo(clState->program, devices[gpu], CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize); + + char *log = malloc(logSize); + status = clGetProgramBuildInfo(clState->program, devices[gpu], CL_PROGRAM_BUILD_LOG, logSize, log, NULL); + printf("%s\n", log); + return NULL; + } + + /* get a kernel object handle for a kernel with the given name */ + clState->kernel = clCreateKernel(clState->program, "oclminer", &status); + if(status != CL_SUCCESS) + { + printf("Error: Creating Kernel from program. (clCreateKernel)\n"); + return NULL; + } + + ///////////////////////////////////////////////////////////////// + // Create an OpenCL command queue + ///////////////////////////////////////////////////////////////// + clState->commandQueue = clCreateCommandQueue( clState->context, devices[gpu], 0, &status); + if(status != CL_SUCCESS) + { + printf("Creating Command Queue. (clCreateCommandQueue)\n"); + return NULL; + } + + clState->inputBuffer = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, sizeof(dev_blk_ctx), NULL, &status); + if(status != CL_SUCCESS) { + printf("Error: clCreateBuffer (inputBuffer)\n"); + return; + } + + clState->outputBuffer = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, sizeof(uint32_t) * MAXTHREADS, NULL, &status); + if(status != CL_SUCCESS) { + printf("Error: clCreateBuffer (outputBuffer)\n"); + return; + } + + return clState; +} + diff --git a/ocl.h b/ocl.h new file mode 100644 index 0000000..fa97ac1 --- /dev/null +++ b/ocl.h @@ -0,0 +1,18 @@ +#ifndef __OCL_H__ +#define __OCL_H__ +#include + +typedef struct { + cl_context context; + cl_kernel kernel; + cl_command_queue commandQueue; + cl_program program; + cl_mem inputBuffer; + cl_mem outputBuffer; +} _clState; + +extern char *file_contents(const char *filename, int *length); +extern int clDevicesNum(); +extern _clState *initCl(int gpu, char *name, size_t nameSize); + +#endif /* __OCL_H__ */ diff --git a/oclminer.cl b/oclminer.cl new file mode 100644 index 0000000..a4b1371 --- /dev/null +++ b/oclminer.cl @@ -0,0 +1,288 @@ +#define rotr(x, n) rotate(x, (uint)(32 - n)) + +#if WORKGROUPSIZE +#define WGS __attribute__((reqd_work_group_size(WORKGROUPSIZE, 1, 1))) +#else +#define WGS +#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 +}; + +typedef struct { + uint ctx_a; uint ctx_b; uint ctx_c; uint ctx_d; + uint ctx_e; uint ctx_f; uint ctx_g; uint ctx_h; + uint cty_a; uint cty_b; uint cty_c; uint cty_d; + uint cty_e; uint cty_f; uint cty_g; uint cty_h; + uint merkle; uint ntime; uint nbits; uint nonce; + uint fW0; uint fW1; uint fW2; uint fW3; uint fW15; + uint fW01r; uint fcty_e; uint fcty_e2; +} dev_blk_ctx; + +__kernel __attribute__((vec_type_hint(uint))) WGS void oclminer( + __constant dev_blk_ctx *ctx, __global uint *output) +{ + const uint fW0 = ctx->fW0; + const uint fW1 = ctx->fW1; + const uint fW2 = ctx->fW2; + const uint fW3 = ctx->fW3; + const uint fW15 = ctx->fW15; + const uint fW01r = ctx->fW01r; + const uint fcty_e = ctx->fcty_e; + const uint fcty_e2 = ctx->fcty_e2; + const uint state0 = ctx->ctx_a; + const uint state1 = ctx->ctx_b; + const uint state2 = ctx->ctx_c; + const uint state3 = ctx->ctx_d; + const uint state4 = ctx->ctx_e; + const uint state5 = ctx->ctx_f; + const uint state6 = ctx->ctx_g; + const uint state7 = ctx->ctx_h; + const uint B1 = ctx->cty_b; + const uint C1 = ctx->cty_c; + const uint D1 = ctx->cty_d; + const uint F1 = ctx->cty_f; + const uint G1 = ctx->cty_g; + const uint H1 = ctx->cty_h; + + uint A, B, C, D, E, F, G, H; + uint W0, W1, W2, W3, W4, W5, W6, W7, W8, W9, W10, W11, W12, W13, W14, W15; + uint it, res = 0; + const uint myid = get_global_id(0); + + const uint tnonce = (ctx->nonce + myid)<<10; + + for(it = 0; it != 1024; it++) { + W3 = it ^ tnonce; + E = fcty_e + W3; A = state0 + E; E = E + fcty_e2; + D = D1 + (rotr(A, 6) ^ rotr(A, 11) ^ rotr(A, 25)) + (C1 ^ (A & (B1 ^ C1))) + K[ 4] + 0x80000000; H = H1 + D; D = D + (rotr(E, 2) ^ rotr(E, 13) ^ rotr(E, 22)) + ((E & F1) | (G1 & (E | F1))); + C = C1 + (rotr(H, 6) ^ rotr(H, 11) ^ rotr(H, 25)) + (B1 ^ (H & (A ^ B1))) + K[ 5]; G = G1 + C; C = C + (rotr(D, 2) ^ rotr(D, 13) ^ rotr(D, 22)) + ((D & E) | (F1 & (D | E))); + B = B1 + (rotr(G, 6) ^ rotr(G, 11) ^ rotr(G, 25)) + (A ^ (G & (H ^ A))) + K[ 6]; F = F1 + B; B = B + (rotr(C, 2) ^ rotr(C, 13) ^ rotr(C, 22)) + ((C & D) | (E & (C | D))); + A = A + (rotr(F, 6) ^ rotr(F, 11) ^ rotr(F, 25)) + (H ^ (F & (G ^ H))) + K[ 7]; E = E + A; A = A + (rotr(B, 2) ^ rotr(B, 13) ^ rotr(B, 22)) + ((B & C) | (D & (B | C))); + H = H + (rotr(E, 6) ^ rotr(E, 11) ^ rotr(E, 25)) + (G ^ (E & (F ^ G))) + K[ 8]; D = D + H; H = H + (rotr(A, 2) ^ rotr(A, 13) ^ rotr(A, 22)) + ((A & B) | (C & (A | B))); + G = G + (rotr(D, 6) ^ rotr(D, 11) ^ rotr(D, 25)) + (F ^ (D & (E ^ F))) + K[ 9]; C = C + G; G = G + (rotr(H, 2) ^ rotr(H, 13) ^ rotr(H, 22)) + ((H & A) | (B & (H | A))); + F = F + (rotr(C, 6) ^ rotr(C, 11) ^ rotr(C, 25)) + (E ^ (C & (D ^ E))) + K[10]; B = B + F; F = F + (rotr(G, 2) ^ rotr(G, 13) ^ rotr(G, 22)) + ((G & H) | (A & (G | H))); + E = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + K[11]; A = A + E; E = E + (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G))); + D = D + (rotr(A, 6) ^ rotr(A, 11) ^ rotr(A, 25)) + (C ^ (A & (B ^ C))) + K[12]; H = H + D; D = D + (rotr(E, 2) ^ rotr(E, 13) ^ rotr(E, 22)) + ((E & F) | (G & (E | F))); + C = C + (rotr(H, 6) ^ rotr(H, 11) ^ rotr(H, 25)) + (B ^ (H & (A ^ B))) + K[13]; G = G + C; C = C + (rotr(D, 2) ^ rotr(D, 13) ^ rotr(D, 22)) + ((D & E) | (F & (D | E))); + B = B + (rotr(G, 6) ^ rotr(G, 11) ^ rotr(G, 25)) + (A ^ (G & (H ^ A))) + K[14]; F = F + B; B = B + (rotr(C, 2) ^ rotr(C, 13) ^ rotr(C, 22)) + ((C & D) | (E & (C | D))); + A = A + (rotr(F, 6) ^ rotr(F, 11) ^ rotr(F, 25)) + (H ^ (F & (G ^ H))) + K[15] + 0x00000280; E = E + A; A = A + (rotr(B, 2) ^ rotr(B, 13) ^ rotr(B, 22)) + ((B & C) | (D & (B | C))); + H = H + (rotr(E, 6) ^ rotr(E, 11) ^ rotr(E, 25)) + (G ^ (E & (F ^ G))) + K[16] + fW0; D = D + H; H = H + (rotr(A, 2) ^ rotr(A, 13) ^ rotr(A, 22)) + ((A & B) | (C & (A | B))); + G = G + (rotr(D, 6) ^ rotr(D, 11) ^ rotr(D, 25)) + (F ^ (D & (E ^ F))) + K[17] + fW1; C = C + G; G = G + (rotr(H, 2) ^ rotr(H, 13) ^ rotr(H, 22)) + ((H & A) | (B & (H | A))); + W2 = (rotr(W3, 7) ^ rotr(W3, 18) ^ (W3 >> 3)) + fW2; + F = F + (rotr(C, 6) ^ rotr(C, 11) ^ rotr(C, 25)) + (E ^ (C & (D ^ E))) + K[18] + W2; B = B + F; F = F + (rotr(G, 2) ^ rotr(G, 13) ^ rotr(G, 22)) + ((G & H) | (A & (G | H))); + W3 = W3 + fW3; + E = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + K[19] + W3; A = A + E; E = E + (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G))); + W4 = (rotr(W2, 17) ^ rotr(W2, 19) ^ (W2 >> 10)) + 0x80000000; + D = D + (rotr(A, 6) ^ rotr(A, 11) ^ rotr(A, 25)) + (C ^ (A & (B ^ C))) + K[20] + W4; H = H + D; D = D + (rotr(E, 2) ^ rotr(E, 13) ^ rotr(E, 22)) + ((E & F) | (G & (E | F))); + W5 = (rotr(W3, 17) ^ rotr(W3, 19) ^ (W3 >> 10)); + C = C + (rotr(H, 6) ^ rotr(H, 11) ^ rotr(H, 25)) + (B ^ (H & (A ^ B))) + K[21] + W5; G = G + C; C = C + (rotr(D, 2) ^ rotr(D, 13) ^ rotr(D, 22)) + ((D & E) | (F & (D | E))); + W6 = (rotr(W4, 17) ^ rotr(W4, 19) ^ (W4 >> 10)) + 0x00000280; + B = B + (rotr(G, 6) ^ rotr(G, 11) ^ rotr(G, 25)) + (A ^ (G & (H ^ A))) + K[22] + W6; F = F + B; B = B + (rotr(C, 2) ^ rotr(C, 13) ^ rotr(C, 22)) + ((C & D) | (E & (C | D))); + W7 = (rotr(W5, 17) ^ rotr(W5, 19) ^ (W5 >> 10)) + fW0; + A = A + (rotr(F, 6) ^ rotr(F, 11) ^ rotr(F, 25)) + (H ^ (F & (G ^ H))) + K[23] + W7; E = E + A; A = A + (rotr(B, 2) ^ rotr(B, 13) ^ rotr(B, 22)) + ((B & C) | (D & (B | C))); + W8 = (rotr(W6, 17) ^ rotr(W6, 19) ^ (W6 >> 10)) + fW1; + H = H + (rotr(E, 6) ^ rotr(E, 11) ^ rotr(E, 25)) + (G ^ (E & (F ^ G))) + K[24] + W8; D = D + H; H = H + (rotr(A, 2) ^ rotr(A, 13) ^ rotr(A, 22)) + ((A & B) | (C & (A | B))); + W9 = W2 + (rotr(W7, 17) ^ rotr(W7, 19) ^ (W7 >> 10)); + G = G + (rotr(D, 6) ^ rotr(D, 11) ^ rotr(D, 25)) + (F ^ (D & (E ^ F))) + K[25] + W9; C = C + G; G = G + (rotr(H, 2) ^ rotr(H, 13) ^ rotr(H, 22)) + ((H & A) | (B & (H | A))); + W10 = W3 + (rotr(W8, 17) ^ rotr(W8, 19) ^ (W8 >> 10)); + F = F + (rotr(C, 6) ^ rotr(C, 11) ^ rotr(C, 25)) + (E ^ (C & (D ^ E))) + K[26] + W10; B = B + F; F = F + (rotr(G, 2) ^ rotr(G, 13) ^ rotr(G, 22)) + ((G & H) | (A & (G | H))); + W11 = W4 + (rotr(W9, 17) ^ rotr(W9, 19) ^ (W9 >> 10)); + E = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + K[27] + W11; A = A + E; E = E + (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G))); + W12 = W5 + (rotr(W10, 17) ^ rotr(W10, 19) ^ (W10 >> 10)); + D = D + (rotr(A, 6) ^ rotr(A, 11) ^ rotr(A, 25)) + (C ^ (A & (B ^ C))) + K[28] + W12; H = H + D; D = D + (rotr(E, 2) ^ rotr(E, 13) ^ rotr(E, 22)) + ((E & F) | (G & (E | F))); + W13 = W6 + (rotr(W11, 17) ^ rotr(W11, 19) ^ (W11 >> 10)); + C = C + (rotr(H, 6) ^ rotr(H, 11) ^ rotr(H, 25)) + (B ^ (H & (A ^ B))) + K[29] + W13; G = G + C; C = C + (rotr(D, 2) ^ rotr(D, 13) ^ rotr(D, 22)) + ((D & E) | (F & (D | E))); + W14 = 0x00a00055 + W7 + (rotr(W12, 17) ^ rotr(W12, 19) ^ (W12 >> 10)); + B = B + (rotr(G, 6) ^ rotr(G, 11) ^ rotr(G, 25)) + (A ^ (G & (H ^ A))) + K[30] + W14; F = F + B; B = B + (rotr(C, 2) ^ rotr(C, 13) ^ rotr(C, 22)) + ((C & D) | (E & (C | D))); + W15 = fW15 + W8 + (rotr(W13, 17) ^ rotr(W13, 19) ^ (W13 >> 10)); + A = A + (rotr(F, 6) ^ rotr(F, 11) ^ rotr(F, 25)) + (H ^ (F & (G ^ H))) + K[31] + W15; E = E + A; A = A + (rotr(B, 2) ^ rotr(B, 13) ^ rotr(B, 22)) + ((B & C) | (D & (B | C))); + W0 = fW01r + W9 + (rotr(W14, 17) ^ rotr(W14, 19) ^ (W14 >> 10)); + H = H + (rotr(E, 6) ^ rotr(E, 11) ^ rotr(E, 25)) + (G ^ (E & (F ^ G))) + K[32] + W0; D = D + H; H = H + (rotr(A, 2) ^ rotr(A, 13) ^ rotr(A, 22)) + ((A & B) | (C & (A | B))); + W1 = fW1 + (rotr(W2, 7) ^ rotr(W2, 18) ^ (W2 >> 3)) + W10 + (rotr(W15, 17) ^ rotr(W15, 19) ^ (W15 >> 10)); + G = G + (rotr(D, 6) ^ rotr(D, 11) ^ rotr(D, 25)) + (F ^ (D & (E ^ F))) + K[33] + W1; C = C + G; G = G + (rotr(H, 2) ^ rotr(H, 13) ^ rotr(H, 22)) + ((H & A) | (B & (H | A))); + W2 = W2 + (rotr(W3, 7) ^ rotr(W3, 18) ^ (W3 >> 3)) + W11 + (rotr(W0, 17) ^ rotr(W0, 19) ^ (W0 >> 10)); + F = F + (rotr(C, 6) ^ rotr(C, 11) ^ rotr(C, 25)) + (E ^ (C & (D ^ E))) + K[34] + W2; B = B + F; F = F + (rotr(G, 2) ^ rotr(G, 13) ^ rotr(G, 22)) + ((G & H) | (A & (G | H))); + W3 = W3 + (rotr(W4, 7) ^ rotr(W4, 18) ^ (W4 >> 3)) + W12 + (rotr(W1, 17) ^ rotr(W1, 19) ^ (W1 >> 10)); + E = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + K[35] + W3; A = A + E; E = E + (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G))); + W4 = W4 + (rotr(W5, 7) ^ rotr(W5, 18) ^ (W5 >> 3)) + W13 + (rotr(W2, 17) ^ rotr(W2, 19) ^ (W2 >> 10)); + D = D + (rotr(A, 6) ^ rotr(A, 11) ^ rotr(A, 25)) + (C ^ (A & (B ^ C))) + K[36] + W4; H = H + D; D = D + (rotr(E, 2) ^ rotr(E, 13) ^ rotr(E, 22)) + ((E & F) | (G & (E | F))); + W5 = W5 + (rotr(W6, 7) ^ rotr(W6, 18) ^ (W6 >> 3)) + W14 + (rotr(W3, 17) ^ rotr(W3, 19) ^ (W3 >> 10)); + C = C + (rotr(H, 6) ^ rotr(H, 11) ^ rotr(H, 25)) + (B ^ (H & (A ^ B))) + K[37] + W5; G = G + C; C = C + (rotr(D, 2) ^ rotr(D, 13) ^ rotr(D, 22)) + ((D & E) | (F & (D | E))); + W6 = W6 + (rotr(W7, 7) ^ rotr(W7, 18) ^ (W7 >> 3)) + W15 + (rotr(W4, 17) ^ rotr(W4, 19) ^ (W4 >> 10)); + B = B + (rotr(G, 6) ^ rotr(G, 11) ^ rotr(G, 25)) + (A ^ (G & (H ^ A))) + K[38] + W6; F = F + B; B = B + (rotr(C, 2) ^ rotr(C, 13) ^ rotr(C, 22)) + ((C & D) | (E & (C | D))); + W7 = W7 + (rotr(W8, 7) ^ rotr(W8, 18) ^ (W8 >> 3)) + W0 + (rotr(W5, 17) ^ rotr(W5, 19) ^ (W5 >> 10)); + A = A + (rotr(F, 6) ^ rotr(F, 11) ^ rotr(F, 25)) + (H ^ (F & (G ^ H))) + K[39] + W7; E = E + A; A = A + (rotr(B, 2) ^ rotr(B, 13) ^ rotr(B, 22)) + ((B & C) | (D & (B | C))); + W8 = W8 + (rotr(W9, 7) ^ rotr(W9, 18) ^ (W9 >> 3)) + W1 + (rotr(W6, 17) ^ rotr(W6, 19) ^ (W6 >> 10)); + H = H + (rotr(E, 6) ^ rotr(E, 11) ^ rotr(E, 25)) + (G ^ (E & (F ^ G))) + K[40] + W8; D = D + H; H = H + (rotr(A, 2) ^ rotr(A, 13) ^ rotr(A, 22)) + ((A & B) | (C & (A | B))); + W9 = W9 + (rotr(W10, 7) ^ rotr(W10, 18) ^ (W10 >> 3)) + W2 + (rotr(W7, 17) ^ rotr(W7, 19) ^ (W7 >> 10)); + G = G + (rotr(D, 6) ^ rotr(D, 11) ^ rotr(D, 25)) + (F ^ (D & (E ^ F))) + K[41] + W9; C = C + G; G = G + (rotr(H, 2) ^ rotr(H, 13) ^ rotr(H, 22)) + ((H & A) | (B & (H | A))); + W10 = W10 + (rotr(W11, 7) ^ rotr(W11, 18) ^ (W11 >> 3)) + W3 + (rotr(W8, 17) ^ rotr(W8, 19) ^ (W8 >> 10)); + F = F + (rotr(C, 6) ^ rotr(C, 11) ^ rotr(C, 25)) + (E ^ (C & (D ^ E))) + K[42] + W10; B = B + F; F = F + (rotr(G, 2) ^ rotr(G, 13) ^ rotr(G, 22)) + ((G & H) | (A & (G | H))); + W11 = W11 + (rotr(W12, 7) ^ rotr(W12, 18) ^ (W12 >> 3)) + W4 + (rotr(W9, 17) ^ rotr(W9, 19) ^ (W9 >> 10)); + E = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + K[43] + W11; A = A + E; E = E + (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G))); + W12 = W12 + (rotr(W13, 7) ^ rotr(W13, 18) ^ (W13 >> 3)) + W5 + (rotr(W10, 17) ^ rotr(W10, 19) ^ (W10 >> 10)); + D = D + (rotr(A, 6) ^ rotr(A, 11) ^ rotr(A, 25)) + (C ^ (A & (B ^ C))) + K[44] + W12; H = H + D; D = D + (rotr(E, 2) ^ rotr(E, 13) ^ rotr(E, 22)) + ((E & F) | (G & (E | F))); + W13 = W13 + (rotr(W14, 7) ^ rotr(W14, 18) ^ (W14 >> 3)) + W6 + (rotr(W11, 17) ^ rotr(W11, 19) ^ (W11 >> 10)); + C = C + (rotr(H, 6) ^ rotr(H, 11) ^ rotr(H, 25)) + (B ^ (H & (A ^ B))) + K[45] + W13; G = G + C; C = C + (rotr(D, 2) ^ rotr(D, 13) ^ rotr(D, 22)) + ((D & E) | (F & (D | E))); + W14 = W14 + (rotr(W15, 7) ^ rotr(W15, 18) ^ (W15 >> 3)) + W7 + (rotr(W12, 17) ^ rotr(W12, 19) ^ (W12 >> 10)); + B = B + (rotr(G, 6) ^ rotr(G, 11) ^ rotr(G, 25)) + (A ^ (G & (H ^ A))) + K[46] + W14; F = F + B; B = B + (rotr(C, 2) ^ rotr(C, 13) ^ rotr(C, 22)) + ((C & D) | (E & (C | D))); + W15 = W15 + (rotr(W0, 7) ^ rotr(W0, 18) ^ (W0 >> 3)) + W8 + (rotr(W13, 17) ^ rotr(W13, 19) ^ (W13 >> 10)); + A = A + (rotr(F, 6) ^ rotr(F, 11) ^ rotr(F, 25)) + (H ^ (F & (G ^ H))) + K[47] + W15; E = E + A; A = A + (rotr(B, 2) ^ rotr(B, 13) ^ rotr(B, 22)) + ((B & C) | (D & (B | C))); + W0 = W0 + (rotr(W1, 7) ^ rotr(W1, 18) ^ (W1 >> 3)) + W9 + (rotr(W14, 17) ^ rotr(W14, 19) ^ (W14 >> 10)); + H = H + (rotr(E, 6) ^ rotr(E, 11) ^ rotr(E, 25)) + (G ^ (E & (F ^ G))) + K[48] + W0; D = D + H; H = H + (rotr(A, 2) ^ rotr(A, 13) ^ rotr(A, 22)) + ((A & B) | (C & (A | B))); + W1 = W1 + (rotr(W2, 7) ^ rotr(W2, 18) ^ (W2 >> 3)) + W10 + (rotr(W15, 17) ^ rotr(W15, 19) ^ (W15 >> 10)); + G = G + (rotr(D, 6) ^ rotr(D, 11) ^ rotr(D, 25)) + (F ^ (D & (E ^ F))) + K[49] + W1; C = C + G; G = G + (rotr(H, 2) ^ rotr(H, 13) ^ rotr(H, 22)) + ((H & A) | (B & (H | A))); + W2 = W2 + (rotr(W3, 7) ^ rotr(W3, 18) ^ (W3 >> 3)) + W11 + (rotr(W0, 17) ^ rotr(W0, 19) ^ (W0 >> 10)); + F = F + (rotr(C, 6) ^ rotr(C, 11) ^ rotr(C, 25)) + (E ^ (C & (D ^ E))) + K[50] + W2; B = B + F; F = F + (rotr(G, 2) ^ rotr(G, 13) ^ rotr(G, 22)) + ((G & H) | (A & (G | H))); + W3 = W3 + (rotr(W4, 7) ^ rotr(W4, 18) ^ (W4 >> 3)) + W12 + (rotr(W1, 17) ^ rotr(W1, 19) ^ (W1 >> 10)); + E = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + K[51] + W3; A = A + E; E = E + (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G))); + W4 = W4 + (rotr(W5, 7) ^ rotr(W5, 18) ^ (W5 >> 3)) + W13 + (rotr(W2, 17) ^ rotr(W2, 19) ^ (W2 >> 10)); + D = D + (rotr(A, 6) ^ rotr(A, 11) ^ rotr(A, 25)) + (C ^ (A & (B ^ C))) + K[52] + W4; H = H + D; D = D + (rotr(E, 2) ^ rotr(E, 13) ^ rotr(E, 22)) + ((E & F) | (G & (E | F))); + W5 = W5 + (rotr(W6, 7) ^ rotr(W6, 18) ^ (W6 >> 3)) + W14 + (rotr(W3, 17) ^ rotr(W3, 19) ^ (W3 >> 10)); + C = C + (rotr(H, 6) ^ rotr(H, 11) ^ rotr(H, 25)) + (B ^ (H & (A ^ B))) + K[53] + W5; G = G + C; C = C + (rotr(D, 2) ^ rotr(D, 13) ^ rotr(D, 22)) + ((D & E) | (F & (D | E))); + W6 = W6 + (rotr(W7, 7) ^ rotr(W7, 18) ^ (W7 >> 3)) + W15 + (rotr(W4, 17) ^ rotr(W4, 19) ^ (W4 >> 10)); + B = B + (rotr(G, 6) ^ rotr(G, 11) ^ rotr(G, 25)) + (A ^ (G & (H ^ A))) + K[54] + W6; F = F + B; B = B + (rotr(C, 2) ^ rotr(C, 13) ^ rotr(C, 22)) + ((C & D) | (E & (C | D))); + W7 = W7 + (rotr(W8, 7) ^ rotr(W8, 18) ^ (W8 >> 3)) + W0 + (rotr(W5, 17) ^ rotr(W5, 19) ^ (W5 >> 10)); + A = A + (rotr(F, 6) ^ rotr(F, 11) ^ rotr(F, 25)) + (H ^ (F & (G ^ H))) + K[55] + W7; E = E + A; A = A + (rotr(B, 2) ^ rotr(B, 13) ^ rotr(B, 22)) + ((B & C) | (D & (B | C))); + W8 = W8 + (rotr(W9, 7) ^ rotr(W9, 18) ^ (W9 >> 3)) + W1 + (rotr(W6, 17) ^ rotr(W6, 19) ^ (W6 >> 10)); + H = H + (rotr(E, 6) ^ rotr(E, 11) ^ rotr(E, 25)) + (G ^ (E & (F ^ G))) + K[56] + W8; D = D + H; H = H + (rotr(A, 2) ^ rotr(A, 13) ^ rotr(A, 22)) + ((A & B) | (C & (A | B))); + W9 = W9 + (rotr(W10, 7) ^ rotr(W10, 18) ^ (W10 >> 3)) + W2 + (rotr(W7, 17) ^ rotr(W7, 19) ^ (W7 >> 10)); + G = G + (rotr(D, 6) ^ rotr(D, 11) ^ rotr(D, 25)) + (F ^ (D & (E ^ F))) + K[57] + W9; C = C + G; G = G + (rotr(H, 2) ^ rotr(H, 13) ^ rotr(H, 22)) + ((H & A) | (B & (H | A))); + W10 = W10 + (rotr(W11, 7) ^ rotr(W11, 18) ^ (W11 >> 3)) + W3 + (rotr(W8, 17) ^ rotr(W8, 19) ^ (W8 >> 10)); + F = F + (rotr(C, 6) ^ rotr(C, 11) ^ rotr(C, 25)) + (E ^ (C & (D ^ E))) + K[58] + W10; B = B + F; F = F + (rotr(G, 2) ^ rotr(G, 13) ^ rotr(G, 22)) + ((G & H) | (A & (G | H))); + W11 = W11 + (rotr(W12, 7) ^ rotr(W12, 18) ^ (W12 >> 3)) + W4 + (rotr(W9, 17) ^ rotr(W9, 19) ^ (W9 >> 10)); + E = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + K[59] + W11; A = A + E; E = E + (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G))); + W12 = W12 + (rotr(W13, 7) ^ rotr(W13, 18) ^ (W13 >> 3)) + W5 + (rotr(W10, 17) ^ rotr(W10, 19) ^ (W10 >> 10)); + D = D + (rotr(A, 6) ^ rotr(A, 11) ^ rotr(A, 25)) + (C ^ (A & (B ^ C))) + K[60] + W12; H = H + D; D = D + (rotr(E, 2) ^ rotr(E, 13) ^ rotr(E, 22)) + ((E & F) | (G & (E | F))); + W13 = W13 + (rotr(W14, 7) ^ rotr(W14, 18) ^ (W14 >> 3)) + W6 + (rotr(W11, 17) ^ rotr(W11, 19) ^ (W11 >> 10)); + C = C + (rotr(H, 6) ^ rotr(H, 11) ^ rotr(H, 25)) + (B ^ (H & (A ^ B))) + K[61] + W13; G = G + C; C = C + (rotr(D, 2) ^ rotr(D, 13) ^ rotr(D, 22)) + ((D & E) | (F & (D | E))); + W14 = W14 + (rotr(W15, 7) ^ rotr(W15, 18) ^ (W15 >> 3)) + W7 + (rotr(W12, 17) ^ rotr(W12, 19) ^ (W12 >> 10)); + B = B + (rotr(G, 6) ^ rotr(G, 11) ^ rotr(G, 25)) + (A ^ (G & (H ^ A))) + K[62] + W14; F = F + B; B = B + (rotr(C, 2) ^ rotr(C, 13) ^ rotr(C, 22)) + ((C & D) | (E & (C | D))); + W15 = W15 + (rotr(W0, 7) ^ rotr(W0, 18) ^ (W0 >> 3)) + W8 + (rotr(W13, 17) ^ rotr(W13, 19) ^ (W13 >> 10)); + A = A + (rotr(F, 6) ^ rotr(F, 11) ^ rotr(F, 25)) + (H ^ (F & (G ^ H))) + K[63] + W15; E = E + A; A = A + (rotr(B, 2) ^ rotr(B, 13) ^ rotr(B, 22)) + ((B & C) | (D & (B | C))); + + W0 = A + state0; W1 = B + state1; + W2 = C + state2; W3 = D + state3; + W4 = E + state4; W5 = F + state5; + W6 = G + state6; W7 = H + state7; + H = 0xb0edbdd0 + K[ 0] + W0; D = 0xa54ff53a + H; H = H + 0x08909ae5; + G = 0x1f83d9ab + (rotr(D, 6) ^ rotr(D, 11) ^ rotr(D, 25)) + (0x9b05688c ^ (D & 0xca0b3af3)) + K[ 1] + W1; C = 0x3c6ef372 + G; G = G + (rotr(H, 2) ^ rotr(H, 13) ^ rotr(H, 22)) + ((H & 0x6a09e667) | (0xbb67ae85 & (H | 0x6a09e667))); + F = 0x9b05688c + (rotr(C, 6) ^ rotr(C, 11) ^ rotr(C, 25)) + (0x510e527f ^ (C & (D ^ 0x510e527f))) + K[ 2] + W2; B = 0xbb67ae85 + F; F = F + (rotr(G, 2) ^ rotr(G, 13) ^ rotr(G, 22)) + ((G & H) | (0x6a09e667 & (G | H))); + E = 0x510e527f + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + K[ 3] + W3; A = 0x6a09e667 + E; E = E + (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G))); + D = D + (rotr(A, 6) ^ rotr(A, 11) ^ rotr(A, 25)) + (C ^ (A & (B ^ C))) + K[ 4] + W4; H = H + D; D = D + (rotr(E, 2) ^ rotr(E, 13) ^ rotr(E, 22)) + ((E & F) | (G & (E | F))); + C = C + (rotr(H, 6) ^ rotr(H, 11) ^ rotr(H, 25)) + (B ^ (H & (A ^ B))) + K[ 5] + W5; G = G + C; C = C + (rotr(D, 2) ^ rotr(D, 13) ^ rotr(D, 22)) + ((D & E) | (F & (D | E))); + B = B + (rotr(G, 6) ^ rotr(G, 11) ^ rotr(G, 25)) + (A ^ (G & (H ^ A))) + K[ 6] + W6; F = F + B; B = B + (rotr(C, 2) ^ rotr(C, 13) ^ rotr(C, 22)) + ((C & D) | (E & (C | D))); + A = A + (rotr(F, 6) ^ rotr(F, 11) ^ rotr(F, 25)) + (H ^ (F & (G ^ H))) + K[ 7] + W7; E = E + A; A = A + (rotr(B, 2) ^ rotr(B, 13) ^ rotr(B, 22)) + ((B & C) | (D & (B | C))); + H = H + (rotr(E, 6) ^ rotr(E, 11) ^ rotr(E, 25)) + (G ^ (E & (F ^ G))) + K[ 8] + 0x80000000; D = D + H; H = H + (rotr(A, 2) ^ rotr(A, 13) ^ rotr(A, 22)) + ((A & B) | (C & (A | B))); + G = G + (rotr(D, 6) ^ rotr(D, 11) ^ rotr(D, 25)) + (F ^ (D & (E ^ F))) + K[ 9]; C = C + G; G = G + (rotr(H, 2) ^ rotr(H, 13) ^ rotr(H, 22)) + ((H & A) | (B & (H | A))); + F = F + (rotr(C, 6) ^ rotr(C, 11) ^ rotr(C, 25)) + (E ^ (C & (D ^ E))) + K[10]; B = B + F; F = F + (rotr(G, 2) ^ rotr(G, 13) ^ rotr(G, 22)) + ((G & H) | (A & (G | H))); + E = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + K[11]; A = A + E; E = E + (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G))); + D = D + (rotr(A, 6) ^ rotr(A, 11) ^ rotr(A, 25)) + (C ^ (A & (B ^ C))) + K[12]; H = H + D; D = D + (rotr(E, 2) ^ rotr(E, 13) ^ rotr(E, 22)) + ((E & F) | (G & (E | F))); + C = C + (rotr(H, 6) ^ rotr(H, 11) ^ rotr(H, 25)) + (B ^ (H & (A ^ B))) + K[13]; G = G + C; C = C + (rotr(D, 2) ^ rotr(D, 13) ^ rotr(D, 22)) + ((D & E) | (F & (D | E))); + B = B + (rotr(G, 6) ^ rotr(G, 11) ^ rotr(G, 25)) + (A ^ (G & (H ^ A))) + K[14]; F = F + B; B = B + (rotr(C, 2) ^ rotr(C, 13) ^ rotr(C, 22)) + ((C & D) | (E & (C | D))); + A = A + (rotr(F, 6) ^ rotr(F, 11) ^ rotr(F, 25)) + (H ^ (F & (G ^ H))) + K[15] + 0x00000100; E = E + A; A = A + (rotr(B, 2) ^ rotr(B, 13) ^ rotr(B, 22)) + ((B & C) | (D & (B | C))); + W0 = W0 + (rotr(W1, 7) ^ rotr(W1, 18) ^ (W1 >> 3)); + H = H + (rotr(E, 6) ^ rotr(E, 11) ^ rotr(E, 25)) + (G ^ (E & (F ^ G))) + K[16] + W0; D = D + H; H = H + (rotr(A, 2) ^ rotr(A, 13) ^ rotr(A, 22)) + ((A & B) | (C & (A | B))); + W1 = W1 + (rotr(W2, 7) ^ rotr(W2, 18) ^ (W2 >> 3)) + 0x00a00000; + G = G + (rotr(D, 6) ^ rotr(D, 11) ^ rotr(D, 25)) + (F ^ (D & (E ^ F))) + K[17] + W1; C = C + G; G = G + (rotr(H, 2) ^ rotr(H, 13) ^ rotr(H, 22)) + ((H & A) | (B & (H | A))); + W2 = W2 + (rotr(W3, 7) ^ rotr(W3, 18) ^ (W3 >> 3)) + (rotr(W0, 17) ^ rotr(W0, 19) ^ (W0 >> 10)); + F = F + (rotr(C, 6) ^ rotr(C, 11) ^ rotr(C, 25)) + (E ^ (C & (D ^ E))) + K[18] + W2; B = B + F; F = F + (rotr(G, 2) ^ rotr(G, 13) ^ rotr(G, 22)) + ((G & H) | (A & (G | H))); + W3 = W3 + (rotr(W4, 7) ^ rotr(W4, 18) ^ (W4 >> 3)) + (rotr(W1, 17) ^ rotr(W1, 19) ^ (W1 >> 10)); + E = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + K[19] + W3; A = A + E; E = E + (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G))); + W4 = W4 + (rotr(W5, 7) ^ rotr(W5, 18) ^ (W5 >> 3)) + (rotr(W2, 17) ^ rotr(W2, 19) ^ (W2 >> 10)); + D = D + (rotr(A, 6) ^ rotr(A, 11) ^ rotr(A, 25)) + (C ^ (A & (B ^ C))) + K[20] + W4; H = H + D; D = D + (rotr(E, 2) ^ rotr(E, 13) ^ rotr(E, 22)) + ((E & F) | (G & (E | F))); + W5 = W5 + (rotr(W6, 7) ^ rotr(W6, 18) ^ (W6 >> 3)) + (rotr(W3, 17) ^ rotr(W3, 19) ^ (W3 >> 10)); + C = C + (rotr(H, 6) ^ rotr(H, 11) ^ rotr(H, 25)) + (B ^ (H & (A ^ B))) + K[21] + W5; G = G + C; C = C + (rotr(D, 2) ^ rotr(D, 13) ^ rotr(D, 22)) + ((D & E) | (F & (D | E))); + W6 = W6 + (rotr(W7, 7) ^ rotr(W7, 18) ^ (W7 >> 3)) + 0x00000100 + (rotr(W4, 17) ^ rotr(W4, 19) ^ (W4 >> 10)); + B = B + (rotr(G, 6) ^ rotr(G, 11) ^ rotr(G, 25)) + (A ^ (G & (H ^ A))) + K[22] + W6; F = F + B; B = B + (rotr(C, 2) ^ rotr(C, 13) ^ rotr(C, 22)) + ((C & D) | (E & (C | D))); + W7 = W7 + 0x11002000 + W0 + (rotr(W5, 17) ^ rotr(W5, 19) ^ (W5 >> 10)); + A = A + (rotr(F, 6) ^ rotr(F, 11) ^ rotr(F, 25)) + (H ^ (F & (G ^ H))) + K[23] + W7; E = E + A; A = A + (rotr(B, 2) ^ rotr(B, 13) ^ rotr(B, 22)) + ((B & C) | (D & (B | C))); + W8 = 0x80000000 + W1 + (rotr(W6, 17) ^ rotr(W6, 19) ^ (W6 >> 10)); + H = H + (rotr(E, 6) ^ rotr(E, 11) ^ rotr(E, 25)) + (G ^ (E & (F ^ G))) + K[24] + W8; D = D + H; H = H + (rotr(A, 2) ^ rotr(A, 13) ^ rotr(A, 22)) + ((A & B) | (C & (A | B))); + W9 = W2 + (rotr(W7, 17) ^ rotr(W7, 19) ^ (W7 >> 10)); + G = G + (rotr(D, 6) ^ rotr(D, 11) ^ rotr(D, 25)) + (F ^ (D & (E ^ F))) + K[25] + W9; C = C + G; G = G + (rotr(H, 2) ^ rotr(H, 13) ^ rotr(H, 22)) + ((H & A) | (B & (H | A))); + W10 = W3 + (rotr(W8, 17) ^ rotr(W8, 19) ^ (W8 >> 10)); + F = F + (rotr(C, 6) ^ rotr(C, 11) ^ rotr(C, 25)) + (E ^ (C & (D ^ E))) + K[26] + W10; B = B + F; F = F + (rotr(G, 2) ^ rotr(G, 13) ^ rotr(G, 22)) + ((G & H) | (A & (G | H))); + W11 = W4 + (rotr(W9, 17) ^ rotr(W9, 19) ^ (W9 >> 10)); + E = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + K[27] + W11; A = A + E; E = E + (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G))); + W12 = W5 + (rotr(W10, 17) ^ rotr(W10, 19) ^ (W10 >> 10)); + D = D + (rotr(A, 6) ^ rotr(A, 11) ^ rotr(A, 25)) + (C ^ (A & (B ^ C))) + K[28] + W12; H = H + D; D = D + (rotr(E, 2) ^ rotr(E, 13) ^ rotr(E, 22)) + ((E & F) | (G & (E | F))); + W13 = W6 + (rotr(W11, 17) ^ rotr(W11, 19) ^ (W11 >> 10)); + C = C + (rotr(H, 6) ^ rotr(H, 11) ^ rotr(H, 25)) + (B ^ (H & (A ^ B))) + K[29] + W13; G = G + C; C = C + (rotr(D, 2) ^ rotr(D, 13) ^ rotr(D, 22)) + ((D & E) | (F & (D | E))); + W14 = 0x00400022 + W7 + (rotr(W12, 17) ^ rotr(W12, 19) ^ (W12 >> 10)); + B = B + (rotr(G, 6) ^ rotr(G, 11) ^ rotr(G, 25)) + (A ^ (G & (H ^ A))) + K[30] + W14; F = F + B; B = B + (rotr(C, 2) ^ rotr(C, 13) ^ rotr(C, 22)) + ((C & D) | (E & (C | D))); + W15 = 0x00000100 + (rotr(W0, 7) ^ rotr(W0, 18) ^ (W0 >> 3)) + W8 + (rotr(W13, 17) ^ rotr(W13, 19) ^ (W13 >> 10)); + A = A + (rotr(F, 6) ^ rotr(F, 11) ^ rotr(F, 25)) + (H ^ (F & (G ^ H))) + K[31] + W15; E = E + A; A = A + (rotr(B, 2) ^ rotr(B, 13) ^ rotr(B, 22)) + ((B & C) | (D & (B | C))); + W0 = W0 + (rotr(W1, 7) ^ rotr(W1, 18) ^ (W1 >> 3)) + W9 + (rotr(W14, 17) ^ rotr(W14, 19) ^ (W14 >> 10)); + H = H + (rotr(E, 6) ^ rotr(E, 11) ^ rotr(E, 25)) + (G ^ (E & (F ^ G))) + K[32] + W0; D = D + H; H = H + (rotr(A, 2) ^ rotr(A, 13) ^ rotr(A, 22)) + ((A & B) | (C & (A | B))); + W1 = W1 + (rotr(W2, 7) ^ rotr(W2, 18) ^ (W2 >> 3)) + W10 + (rotr(W15, 17) ^ rotr(W15, 19) ^ (W15 >> 10)); + G = G + (rotr(D, 6) ^ rotr(D, 11) ^ rotr(D, 25)) + (F ^ (D & (E ^ F))) + K[33] + W1; C = C + G; G = G + (rotr(H, 2) ^ rotr(H, 13) ^ rotr(H, 22)) + ((H & A) | (B & (H | A))); + W2 = W2 + (rotr(W3, 7) ^ rotr(W3, 18) ^ (W3 >> 3)) + W11 + (rotr(W0, 17) ^ rotr(W0, 19) ^ (W0 >> 10)); + F = F + (rotr(C, 6) ^ rotr(C, 11) ^ rotr(C, 25)) + (E ^ (C & (D ^ E))) + K[34] + W2; B = B + F; F = F + (rotr(G, 2) ^ rotr(G, 13) ^ rotr(G, 22)) + ((G & H) | (A & (G | H))); + W3 = W3 + (rotr(W4, 7) ^ rotr(W4, 18) ^ (W4 >> 3)) + W12 + (rotr(W1, 17) ^ rotr(W1, 19) ^ (W1 >> 10)); + E = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + K[35] + W3; A = A + E; E = E + (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G))); + W4 = W4 + (rotr(W5, 7) ^ rotr(W5, 18) ^ (W5 >> 3)) + W13 + (rotr(W2, 17) ^ rotr(W2, 19) ^ (W2 >> 10)); + D = D + (rotr(A, 6) ^ rotr(A, 11) ^ rotr(A, 25)) + (C ^ (A & (B ^ C))) + K[36] + W4; H = H + D; D = D + (rotr(E, 2) ^ rotr(E, 13) ^ rotr(E, 22)) + ((E & F) | (G & (E | F))); + W5 = W5 + (rotr(W6, 7) ^ rotr(W6, 18) ^ (W6 >> 3)) + W14 + (rotr(W3, 17) ^ rotr(W3, 19) ^ (W3 >> 10)); + C = C + (rotr(H, 6) ^ rotr(H, 11) ^ rotr(H, 25)) + (B ^ (H & (A ^ B))) + K[37] + W5; G = G + C; C = C + (rotr(D, 2) ^ rotr(D, 13) ^ rotr(D, 22)) + ((D & E) | (F & (D | E))); + W6 = W6 + (rotr(W7, 7) ^ rotr(W7, 18) ^ (W7 >> 3)) + W15 + (rotr(W4, 17) ^ rotr(W4, 19) ^ (W4 >> 10)); + B = B + (rotr(G, 6) ^ rotr(G, 11) ^ rotr(G, 25)) + (A ^ (G & (H ^ A))) + K[38] + W6; F = F + B; B = B + (rotr(C, 2) ^ rotr(C, 13) ^ rotr(C, 22)) + ((C & D) | (E & (C | D))); + W7 = W7 + (rotr(W8, 7) ^ rotr(W8, 18) ^ (W8 >> 3)) + W0 + (rotr(W5, 17) ^ rotr(W5, 19) ^ (W5 >> 10)); + A = A + (rotr(F, 6) ^ rotr(F, 11) ^ rotr(F, 25)) + (H ^ (F & (G ^ H))) + K[39] + W7; E = E + A; A = A + (rotr(B, 2) ^ rotr(B, 13) ^ rotr(B, 22)) + ((B & C) | (D & (B | C))); + W8 = W8 + (rotr(W9, 7) ^ rotr(W9, 18) ^ (W9 >> 3)) + W1 + (rotr(W6, 17) ^ rotr(W6, 19) ^ (W6 >> 10)); + H = H + (rotr(E, 6) ^ rotr(E, 11) ^ rotr(E, 25)) + (G ^ (E & (F ^ G))) + K[40] + W8; D = D + H; H = H + (rotr(A, 2) ^ rotr(A, 13) ^ rotr(A, 22)) + ((A & B) | (C & (A | B))); + W9 = W9 + (rotr(W10, 7) ^ rotr(W10, 18) ^ (W10 >> 3)) + W2 + (rotr(W7, 17) ^ rotr(W7, 19) ^ (W7 >> 10)); + G = G + (rotr(D, 6) ^ rotr(D, 11) ^ rotr(D, 25)) + (F ^ (D & (E ^ F))) + K[41] + W9; C = C + G; G = G + (rotr(H, 2) ^ rotr(H, 13) ^ rotr(H, 22)) + ((H & A) | (B & (H | A))); + W10 = W10 + (rotr(W11, 7) ^ rotr(W11, 18) ^ (W11 >> 3)) + W3 + (rotr(W8, 17) ^ rotr(W8, 19) ^ (W8 >> 10)); + F = F + (rotr(C, 6) ^ rotr(C, 11) ^ rotr(C, 25)) + (E ^ (C & (D ^ E))) + K[42] + W10; B = B + F; F = F + (rotr(G, 2) ^ rotr(G, 13) ^ rotr(G, 22)) + ((G & H) | (A & (G | H))); + W11 = W11 + (rotr(W12, 7) ^ rotr(W12, 18) ^ (W12 >> 3)) + W4 + (rotr(W9, 17) ^ rotr(W9, 19) ^ (W9 >> 10)); + E = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + K[43] + W11; A = A + E; E = E + (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G))); + W12 = W12 + (rotr(W13, 7) ^ rotr(W13, 18) ^ (W13 >> 3)) + W5 + (rotr(W10, 17) ^ rotr(W10, 19) ^ (W10 >> 10)); + D = D + (rotr(A, 6) ^ rotr(A, 11) ^ rotr(A, 25)) + (C ^ (A & (B ^ C))) + K[44] + W12; H = H + D; D = D + (rotr(E, 2) ^ rotr(E, 13) ^ rotr(E, 22)) + ((E & F) | (G & (E | F))); + W13 = W13 + (rotr(W14, 7) ^ rotr(W14, 18) ^ (W14 >> 3)) + W6 + (rotr(W11, 17) ^ rotr(W11, 19) ^ (W11 >> 10)); + C = C + (rotr(H, 6) ^ rotr(H, 11) ^ rotr(H, 25)) + (B ^ (H & (A ^ B))) + K[45] + W13; G = G + C; C = C + (rotr(D, 2) ^ rotr(D, 13) ^ rotr(D, 22)) + ((D & E) | (F & (D | E))); + W14 = W14 + (rotr(W15, 7) ^ rotr(W15, 18) ^ (W15 >> 3)) + W7 + (rotr(W12, 17) ^ rotr(W12, 19) ^ (W12 >> 10)); + B = B + (rotr(G, 6) ^ rotr(G, 11) ^ rotr(G, 25)) + (A ^ (G & (H ^ A))) + K[46] + W14; F = F + B; B = B + (rotr(C, 2) ^ rotr(C, 13) ^ rotr(C, 22)) + ((C & D) | (E & (C | D))); + W15 = W15 + (rotr(W0, 7) ^ rotr(W0, 18) ^ (W0 >> 3)) + W8 + (rotr(W13, 17) ^ rotr(W13, 19) ^ (W13 >> 10)); + A = A + (rotr(F, 6) ^ rotr(F, 11) ^ rotr(F, 25)) + (H ^ (F & (G ^ H))) + K[47] + W15; E = E + A; A = A + (rotr(B, 2) ^ rotr(B, 13) ^ rotr(B, 22)) + ((B & C) | (D & (B | C))); + W0 = W0 + (rotr(W1, 7) ^ rotr(W1, 18) ^ (W1 >> 3)) + W9 + (rotr(W14, 17) ^ rotr(W14, 19) ^ (W14 >> 10)); + H = H + (rotr(E, 6) ^ rotr(E, 11) ^ rotr(E, 25)) + (G ^ (E & (F ^ G))) + K[48] + W0; D = D + H; H = H + (rotr(A, 2) ^ rotr(A, 13) ^ rotr(A, 22)) + ((A & B) | (C & (A | B))); + W1 = W1 + (rotr(W2, 7) ^ rotr(W2, 18) ^ (W2 >> 3)) + W10 + (rotr(W15, 17) ^ rotr(W15, 19) ^ (W15 >> 10)); + G = G + (rotr(D, 6) ^ rotr(D, 11) ^ rotr(D, 25)) + (F ^ (D & (E ^ F))) + K[49] + W1; C = C + G; G = G + (rotr(H, 2) ^ rotr(H, 13) ^ rotr(H, 22)) + ((H & A) | (B & (H | A))); + W2 = W2 + (rotr(W3, 7) ^ rotr(W3, 18) ^ (W3 >> 3)) + W11 + (rotr(W0, 17) ^ rotr(W0, 19) ^ (W0 >> 10)); + F = F + (rotr(C, 6) ^ rotr(C, 11) ^ rotr(C, 25)) + (E ^ (C & (D ^ E))) + K[50] + W2; B = B + F; F = F + (rotr(G, 2) ^ rotr(G, 13) ^ rotr(G, 22)) + ((G & H) | (A & (G | H))); + W3 = W3 + (rotr(W4, 7) ^ rotr(W4, 18) ^ (W4 >> 3)) + W12 + (rotr(W1, 17) ^ rotr(W1, 19) ^ (W1 >> 10)); + E = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + K[51] + W3; A = A + E; E = E + (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G))); + W4 = W4 + (rotr(W5, 7) ^ rotr(W5, 18) ^ (W5 >> 3)) + W13 + (rotr(W2, 17) ^ rotr(W2, 19) ^ (W2 >> 10)); + D = D + (rotr(A, 6) ^ rotr(A, 11) ^ rotr(A, 25)) + (C ^ (A & (B ^ C))) + K[52] + W4; H = H + D; D = D + (rotr(E, 2) ^ rotr(E, 13) ^ rotr(E, 22)) + ((E & F) | (G & (E | F))); + W5 = W5 + (rotr(W6, 7) ^ rotr(W6, 18) ^ (W6 >> 3)) + W14 + (rotr(W3, 17) ^ rotr(W3, 19) ^ (W3 >> 10)); + C = C + (rotr(H, 6) ^ rotr(H, 11) ^ rotr(H, 25)) + (B ^ (H & (A ^ B))) + K[53] + W5; G = G + C; C = C + (rotr(D, 2) ^ rotr(D, 13) ^ rotr(D, 22)) + ((D & E) | (F & (D | E))); + W6 = W6 + (rotr(W7, 7) ^ rotr(W7, 18) ^ (W7 >> 3)) + W15 + (rotr(W4, 17) ^ rotr(W4, 19) ^ (W4 >> 10)); + B = B + (rotr(G, 6) ^ rotr(G, 11) ^ rotr(G, 25)) + (A ^ (G & (H ^ A))) + K[54] + W6; F = F + B; B = B + (rotr(C, 2) ^ rotr(C, 13) ^ rotr(C, 22)) + ((C & D) | (E & (C | D))); + W7 = W7 + (rotr(W8, 7) ^ rotr(W8, 18) ^ (W8 >> 3)) + W0 + (rotr(W5, 17) ^ rotr(W5, 19) ^ (W5 >> 10)); + A = A + (rotr(F, 6) ^ rotr(F, 11) ^ rotr(F, 25)) + (H ^ (F & (G ^ H))) + K[55] + W7; E = E + A; A = A + (rotr(B, 2) ^ rotr(B, 13) ^ rotr(B, 22)) + ((B & C) | (D & (B | C))); + W8 = W8 + (rotr(W9, 7) ^ rotr(W9, 18) ^ (W9 >> 3)) + W1 + (rotr(W6, 17) ^ rotr(W6, 19) ^ (W6 >> 10)); + H = H + (rotr(E, 6) ^ rotr(E, 11) ^ rotr(E, 25)) + (G ^ (E & (F ^ G))) + K[56] + W8; D = D + H; H = H + (rotr(A, 2) ^ rotr(A, 13) ^ rotr(A, 22)) + ((A & B) | (C & (A | B))); + W9 = W9 + (rotr(W10, 7) ^ rotr(W10, 18) ^ (W10 >> 3)) + W2 + (rotr(W7, 17) ^ rotr(W7, 19) ^ (W7 >> 10)); + G = G + (rotr(D, 6) ^ rotr(D, 11) ^ rotr(D, 25)) + (F ^ (D & (E ^ F))) + K[57] + W9; C = C + G; + W10 = W10 + (rotr(W11, 7) ^ rotr(W11, 18) ^ (W11 >> 3)) + W3 + (rotr(W8, 17) ^ rotr(W8, 19) ^ (W8 >> 10)); + F = F + (rotr(C, 6) ^ rotr(C, 11) ^ rotr(C, 25)) + (E ^ (C & (D ^ E))) + K[58] + W10; B = B + F; + W11 = W11 + (rotr(W12, 7) ^ rotr(W12, 18) ^ (W12 >> 3)) + W4 + (rotr(W9, 17) ^ rotr(W9, 19) ^ (W9 >> 10)); + E = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + K[59] + W11; A = A + E; + W12 = W12 + (rotr(W13, 7) ^ rotr(W13, 18) ^ (W13 >> 3)) + W5 + (rotr(W10, 17) ^ rotr(W10, 19) ^ (W10 >> 10)); + D = D + (rotr(A, 6) ^ rotr(A, 11) ^ rotr(A, 25)) + (C ^ (A & (B ^ C))) + K[60] + W12; H = H + D; + + if(H + 0x5be0cd19 == 0) res = ~0; + } + + output[myid] = res; +} diff --git a/util.c b/util.c new file mode 100644 index 0000000..6088cf1 --- /dev/null +++ b/util.c @@ -0,0 +1,228 @@ + +/* + * Copyright 2010 Jeff Garzik + * + * This program is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License as published by the Free + * Software Foundation; either version 2 of the License, or (at your option) + * any later version. See COPYING for more details. + */ + +#define _GNU_SOURCE + +#include +#include +#include +#include +#include +#include "miner.h" + +struct data_buffer { + void *buf; + size_t len; +}; + +struct upload_buffer { + const void *buf; + size_t len; +}; + +static void databuf_free(struct data_buffer *db) +{ + if (!db) + return; + + free(db->buf); + + memset(db, 0, sizeof(*db)); +} + +static size_t all_data_cb(const void *ptr, size_t size, size_t nmemb, + void *user_data) +{ + struct data_buffer *db = user_data; + size_t len = size * nmemb; + size_t oldlen, newlen; + void *newmem; + static const unsigned char zero; + + oldlen = db->len; + newlen = oldlen + len; + + newmem = realloc(db->buf, newlen + 1); + if (!newmem) + return 0; + + db->buf = newmem; + db->len = newlen; + memcpy(db->buf + oldlen, ptr, len); + memcpy(db->buf + newlen, &zero, 1); /* null terminate */ + + return len; +} + +static size_t upload_data_cb(void *ptr, size_t size, size_t nmemb, + void *user_data) +{ + struct upload_buffer *ub = user_data; + int len = size * nmemb; + + if (len > ub->len) + len = ub->len; + + if (len) { + memcpy(ptr, ub->buf, len); + ub->buf += len; + ub->len -= len; + } + + return len; +} + +json_t *json_rpc_call(const char *url, const char *userpass, const char *rpc_req) +{ + CURL *curl; + json_t *val; + int rc; + struct data_buffer all_data = { }; + struct upload_buffer upload_data; + json_error_t err = { }; + struct curl_slist *headers = NULL; + char len_hdr[64]; + + curl = curl_easy_init(); + if (!curl) + return NULL; + + if (opt_protocol) + curl_easy_setopt(curl, CURLOPT_VERBOSE, 1); + curl_easy_setopt(curl, CURLOPT_URL, url); + curl_easy_setopt(curl, CURLOPT_ENCODING, ""); + curl_easy_setopt(curl, CURLOPT_FAILONERROR, 1); + curl_easy_setopt(curl, CURLOPT_TCP_NODELAY, 1); + curl_easy_setopt(curl, CURLOPT_WRITEFUNCTION, all_data_cb); + curl_easy_setopt(curl, CURLOPT_WRITEDATA, &all_data); + curl_easy_setopt(curl, CURLOPT_READFUNCTION, upload_data_cb); + curl_easy_setopt(curl, CURLOPT_READDATA, &upload_data); + if (userpass) { + curl_easy_setopt(curl, CURLOPT_USERPWD, userpass); + curl_easy_setopt(curl, CURLOPT_HTTPAUTH, CURLAUTH_BASIC); + } + curl_easy_setopt(curl, CURLOPT_POST, 1); + + if (opt_protocol) + printf("JSON protocol request:\n%s\n", rpc_req); + + upload_data.buf = rpc_req; + upload_data.len = strlen(rpc_req); + sprintf(len_hdr, "Content-Length: %lu", + (unsigned long) upload_data.len); + + headers = curl_slist_append(headers, + "Content-type: application/json"); + headers = curl_slist_append(headers, len_hdr); + headers = curl_slist_append(headers, "Expect:"); /* disable Expect hdr*/ + + curl_easy_setopt(curl, CURLOPT_HTTPHEADER, headers); + + rc = curl_easy_perform(curl); + if (rc) + goto err_out; + + val = json_loads(all_data.buf, &err); + if (!val) { + fprintf(stderr, "JSON failed(%d): %s\n", err.line, err.text); + goto err_out; + } + + if (opt_protocol) { + char *s = json_dumps(val, JSON_INDENT(3)); + printf("JSON protocol response:\n%s\n", s); + free(s); + } + + databuf_free(&all_data); + curl_slist_free_all(headers); + curl_easy_cleanup(curl); + return val; + +err_out: + databuf_free(&all_data); + curl_slist_free_all(headers); + curl_easy_cleanup(curl); + return NULL; +} + +char *bin2hex(unsigned char *p, size_t len) +{ + int i; + char *s = malloc((len * 2) + 1); + if (!s) + return NULL; + + for (i = 0; i < len; i++) + sprintf(s + (i * 2), "%02x", (unsigned int) p[i]); + + return s; +} + +int hex2bin(unsigned char *p, const char *hexstr, size_t len) +{ + while (*hexstr && len) { + char hex_byte[3]; + unsigned int v; + + if (!hexstr[1]) { + fprintf(stderr, "hex2bin str truncated\n"); + return 0; + } + + hex_byte[0] = hexstr[0]; + hex_byte[1] = hexstr[1]; + hex_byte[2] = 0; + + if (sscanf(hex_byte, "%x", &v) != 1) { + fprintf(stderr, "hex2bin sscanf '%s' failed\n", + hex_byte); + return 0; + } + + *p = (unsigned char) v; + + p++; + hexstr += 2; + len--; + } + + return len == 0 && *hexstr == 0; +} + +/* Subtract the `struct timeval' values X and Y, + storing the result in RESULT. + Return 1 if the difference is negative, otherwise 0. */ + +int +timeval_subtract ( + struct timeval *result, struct timeval *x, struct timeval *y) +{ + /* Perform the carry for the later subtraction by updating Y. */ + if (x->tv_usec < y->tv_usec) { + int nsec = (y->tv_usec - x->tv_usec) / 1000000 + 1; + y->tv_usec -= 1000000 * nsec; + y->tv_sec += nsec; + } + if (x->tv_usec - y->tv_usec > 1000000) { + int nsec = (x->tv_usec - y->tv_usec) / 1000000; + y->tv_usec += 1000000 * nsec; + y->tv_sec -= nsec; + } + + /* Compute the time remaining to wait. + `tv_usec' is certainly positive. */ + result->tv_sec = x->tv_sec - y->tv_sec; + result->tv_usec = x->tv_usec - y->tv_usec; + + /* Return 1 if result is negative. */ + return x->tv_sec < y->tv_sec; +} +