j"1LAnZuoQdcKCkpDBKQMCgziGMoPC4VQUckMÂ’J<div class="post">Hello again, thanks for the code <img alt="Smiley" border="0" src="/static/img/emoticons/smiley.gif"/><br/><br/>but it still doesn't work <img alt="Sad" border="0" src="/static/img/emoticons/sad.gif"/><br/>Nvidia doesn't like reversing numbers it seems <img alt="Angry" border="0" src="/static/img/emoticons/angry.gif"/><br/>however, if I comments anyone of the four affectations, it builds <img alt="Grin" border="0" src="/static/img/emoticons/grin.gif"/><br/><br/>After trying by decomposing the reversing in many ways, I abandoned this path to bypass the problem.<br/><br/>I made a first version running with byte reversing and comparaisons made in the python script, and a second version with a function that make the comparison directly between a word and a reversed word in the OpenCl code. For now this version works only on big endian memory mapping (I hope it's what is used in the GPUs).<br/><br/>Note: I added somme print to the screen to show some informations when running the script.<br/>I also let some comments in the code.<br/><br/>The two versions build and run without error, but I'm not sure that operations are correct.<br/>As expected the second version is more effective (about 14800 khashs/sec with the 1st version, and 15200 khashs/sec with the 2nd version).<br/><br/>Version 1 : reversing and comparison in the python script<br/><div class="codeheader">Code:</div><div class="code">--- F:/Desktop/m0mchil-poclbm-cf33815/poclbm.py Fri Oct 15 08:48:48 2010<br/>+++ F:/Desktop/m0mchil-poclbm-cf33815/poclbm_nv1.py Mon Oct 18 23:36:46 2010<br/>@@ -21,6 +21,9 @@<br/> t2=(rot(a, 30)^rot(a, 19)^rot(a, 10))+((a&b)|(c&(a|b)))<br/> return (uint32(d + t1), uint32(t1+t2))<br/> <br/>+def bytereverse(x):<br/>+ return (((x & 0x000000ffL) << 24) | ((x & 0x0000ff00L) << 8) | ((x & 0x00ff0000L) >> 8) | ((x & 0xff000000L) >> 24))<br/>+<br/> def sysWrite(format, args=()):<br/> sys.stdout.write('
' + format % args)<br/> sys.stdout.flush()<br/>@@ -50,16 +53,40 @@<br/> context = cl.Context([devices[options.device]], None, None)<br/> else:<br/> print 'No device specified, you may use -d to specify one of the following
'<br/>+ # create_some_context should ask for the platform and device to use<br/> context = cl.create_some_context()<br/>+<br/>+device = context.devices[0]<br/>+print "==============================================================="<br/>+print "Platform name:", platform.name<br/>+print "Platform profile:", platform.profile<br/>+print "Platform vendor:", platform.vendor<br/>+print "Platform version:", platform.version<br/>+print "---------------------------------------------------------------"<br/>+print "Device name:", device.name<br/>+print "Device type:", cl.device_type.to_string(device.type)<br/>+print "Device memory: ", device.global_mem_size//1024//1024, 'MB'<br/>+print "Device max clock speed:", device.max_clock_frequency, 'MHz'<br/>+print "Device compute units:", device.max_compute_units<br/>+print "==============================================================="<br/>+<br/> queue = cl.CommandQueue(context)<br/> <br/>-kernelFile = open('btc_miner.cl', 'r')<br/>+kernelFile = open('btc_miner_nv1.cl', 'r')<br/> miner = cl.Program(context, kernelFile.read()).build()<br/> kernelFile.close()<br/> <br/>+print "==============================================================="<br/>+print "build OK"<br/>+print "==============================================================="<br/>+<br/> if (options.worksize == -1):<br/> options.worksize = miner.search.get_work_group_info(cl.kernel_work_group_info.WORK_GROUP_SIZE, context.devices[0])<br/> <br/>+print "==============================================================="<br/>+print "worksize = ", options.worksize<br/>+print "==============================================================="<br/>+<br/> frames = options.frames<br/> frame = float(1)/frames<br/> window = frame/30<br/>@@ -75,7 +102,7 @@<br/> work['extraNonce'] = 0<br/> work['block'] = ''<br/> <br/>-output = np.zeros(2, np.uint32)<br/>+output = np.zeros(6, np.uint32)<br/> <br/> while True:<br/> try:<br/>@@ -117,10 +144,12 @@<br/> output_buf = cl.Buffer(context, mf.WRITE_ONLY | mf.USE_HOST_PTR, hostbuf=output)<br/> <br/> rate = start = time()<br/>+ success = 0<br/>+ nonce = 0<br/> while True:<br/>- if (output[0]):<br/>- work['block'] = work['block'][:152] + pack('I', long(output[1])).encode('hex') + work['block'][160:]<br/>- sysWriteLn('found: %s, %s', (output[1], datetime.now().strftime("%d/%m/%Y %H:%M")))<br/>+ if (success):<br/>+ work['block'] = work['block'][:152] + pack('I', long(nonce)).encode('hex') + work['block'][160:]<br/>+ sysWriteLn('found: %s, %s', (nonce, datetime.now().strftime("%d/%m/%Y %H:%M")))<br/> break<br/> <br/> if (time() - start > options.askrate or base + globalThreads == 0x7FFFFFFF):<br/>@@ -131,8 +160,20 @@<br/> base = 0x7FFFFFFF - globalThreads<br/> <br/> kernelStart = time()<br/>- miner.search(queue, (globalThreads, ), (options.worksize, ), block2[0], block2[1], block2[2], state[0], state[1], state[2], state[3], state[4], state[5], state[6], state[7], state2[1], state2[2], state2[3], state2[5], state2[6], state2[7], target[6], pack('I', base), output_buf)<br/>+ miner.search(queue, (globalThreads, ), (options.worksize, ), block2[0], block2[1], block2[2], state[0], state[1], state[2], state[3], state[4], state[5], state[6], state[7], state2[1], state2[2], state2[3], state2[5], state2[6], state2[7], pack('I', base), output_buf)<br/> cl.enqueue_read_buffer(queue, output_buf, output).wait()<br/>+ #target = target[6]<br/>+ #H.x = output[0] ; H.y = output[1]<br/>+ #G.x = output[2] ; G.y = output[3]<br/>+ #nonce.x = output[4] ; nonce.y = output[5]<br/>+ if (output[0] == 0 and bytereverse(output[2])<=target[6]):<br/>+ success = 1<br/>+ nonce = output[4]<br/>+ <br/>+ if (output[1] == 0 and bytereverse(output[3])<=target[6]):<br/>+ success = 1<br/>+ nonce = output[5]<br/>+ <br/> kernelTime = time() - kernelStart<br/> <br/> if (kernelTime < lower):<br/>@@ -142,4 +183,4 @@<br/> <br/> if (time() - rate > options.rate):<br/> rate = time()<br/>- sysWrite('%s khash/s', int((base / (time() - start)) / 500))<br/>\ No newline at end of file<br/>+ sysWrite('%s : %s khash/s', (datetime.now().strftime("%d/%m/%Y %H:%M:%S"),int((base / (time() - start)) / 500)))<br/>\ No newline at end of file<br/></div><div class="codeheader">Code:</div><div class="code">--- F:/Desktop/m0mchil-poclbm-cf33815/btc_miner.cl Fri Oct 15 08:48:48 2010<br/>+++ F:/Desktop/m0mchil-poclbm-cf33815/btc_miner_nv1.cl Mon Oct 18 22:01:08 2010<br/>@@ -1,5 +1,4 @@<br/>-#define bytereverse(x) ( ((x) << 24) | (((x) << 8) & 0x00ff0000) | (((x) >> 8) & 0x0000ff00) | ((x) >> 24) )<br/>-#define rot(x, y) rotate(x, (uint)y)<br/>+#define rot(x, y) rotate(x, (uint2)y)<br/> #define R(x) (work[x] = (rot(work[x-2],15)^rot(work[x-2],13)^((work[x-2])>>10)) + work[x-7] + (rot(work[x-15],25)^rot(work[x-15],14)^((work[x-15])>>3)) + work[x-16])<br/> #define sharound(a,b,c,d,e,f,g,h,x,K) {h=(h+(rot(e, 26)^rot(e, 21)^rot(e, 7))+(g^(e&(f^g)))+K+x); t1=(rot(a, 30)^rot(a, 19)^rot(a, 10))+((a&b)|(c&(a|b))); d+=h; h+=t1;}<br/> <br/>@@ -8,7 +7,6 @@<br/> const uint state4, const uint state5, const uint state6, const uint state7,<br/> const uint B1, const uint C1, const uint D1,<br/> const uint F1, const uint G1, const uint H1,<br/>- const uint target,<br/> const uint base,<br/> __global uint * output)<br/> {<br/>@@ -209,17 +207,12 @@<br/> G+=0x1f83d9ab;<br/> H+=0x5be0cd19;<br/> <br/>- if((H.x==0) && (bytereverse(G.x)<=target))<br/>- {<br/>- output[0] = 1;<br/>- output[1] = nonce.x;<br/>- }<br/>-<br/>- if((H.y==0) && (bytereverse(G.y)<=target))<br/>- {<br/>- output[0] = 1;<br/>- output[1] = nonce.y;<br/>- }<br/>+ output[0] = H.x;<br/>+ output[1] = H.y;<br/>+ output[2] = G.x;<br/>+ output[3] = G.y;<br/>+ output[4] = nonce.x;<br/>+ output[5] = nonce.y;<br/> }<br/> <br/> // end<br/>\ No newline at end of file<br/></div><br/>Version 2 : function that make the comparison between inversed and non inversed word<br/><div class="codeheader">Code:</div><div class="code">--- F:/Desktop/m0mchil-poclbm-cf33815/poclbm.py Fri Oct 15 08:48:48 2010<br/>+++ F:/Desktop/m0mchil-poclbm-cf33815/poclbm_nv2.py Mon Oct 18 23:53:18 2010<br/>@@ -50,16 +50,40 @@<br/> context = cl.Context([devices[options.device]], None, None)<br/> else:<br/> print 'No device specified, you may use -d to specify one of the following
'<br/>+ # create_some_context should ask for the platform and device to use<br/> context = cl.create_some_context()<br/>+<br/>+device = context.devices[0]<br/>+print "==============================================================="<br/>+print "Platform name:", platform.name<br/>+print "Platform profile:", platform.profile<br/>+print "Platform vendor:", platform.vendor<br/>+print "Platform version:", platform.version<br/>+print "---------------------------------------------------------------"<br/>+print "Device name:", device.name<br/>+print "Device type:", cl.device_type.to_string(device.type)<br/>+print "Device memory: ", device.global_mem_size//1024//1024, 'MB'<br/>+print "Device max clock speed:", device.max_clock_frequency, 'MHz'<br/>+print "Device compute units:", device.max_compute_units<br/>+print "==============================================================="<br/>+<br/> queue = cl.CommandQueue(context)<br/> <br/>-kernelFile = open('btc_miner.cl', 'r')<br/>+kernelFile = open('btc_miner_nv2.cl', 'r')<br/> miner = cl.Program(context, kernelFile.read()).build()<br/> kernelFile.close()<br/> <br/>+print "==============================================================="<br/>+print "build OK"<br/>+print "==============================================================="<br/>+<br/> if (options.worksize == -1):<br/> options.worksize = miner.search.get_work_group_info(cl.kernel_work_group_info.WORK_GROUP_SIZE, context.devices[0])<br/> <br/>+print "==============================================================="<br/>+print "worksize = ", options.worksize<br/>+print "==============================================================="<br/>+<br/> frames = options.frames<br/> frame = float(1)/frames<br/> window = frame/30<br/>@@ -142,4 +166,4 @@<br/> <br/> if (time() - rate > options.rate):<br/> rate = time()<br/>- sysWrite('%s khash/s', int((base / (time() - start)) / 500))<br/>\ No newline at end of file<br/>+ sysWrite('%s : %s khash/s', (datetime.now().strftime("%d/%m/%Y %H:%M:%S"),int((base / (time() - start)) / 500)))<br/>\ No newline at end of file<br/></div><div class="codeheader">Code:</div><div class="code">--- F:/Desktop/m0mchil-poclbm-cf33815/btc_miner.cl Fri Oct 15 08:48:48 2010<br/>+++ F:/Desktop/m0mchil-poclbm-cf33815/btc_miner_nv2.cl Mon Oct 18 23:33:50 2010<br/>@@ -1,8 +1,88 @@<br/>-#define bytereverse(x) ( ((x) << 24) | (((x) << 8) & 0x00ff0000) | (((x) >> 8) & 0x0000ff00) | ((x) >> 24) )<br/>-#define rot(x, y) rotate(x, (uint)y)<br/>+#define rot(x, y) rotate(x, (uint2)y)<br/> #define R(x) (work[x] = (rot(work[x-2],15)^rot(work[x-2],13)^((work[x-2])>>10)) + work[x-7] + (rot(work[x-15],25)^rot(work[x-15],14)^((work[x-15])>>3)) + work[x-16])<br/> #define sharound(a,b,c,d,e,f,g,h,x,K) {h=(h+(rot(e, 26)^rot(e, 21)^rot(e, 7))+(g^(e&(f^g)))+K+x); t1=(rot(a, 30)^rot(a, 19)^rot(a, 10))+((a&b)|(c&(a|b))); d+=h; h+=t1;}<br/> <br/>+// test 1 : recursive function, build error :<br/>+// Error: Code selection failed to select: 0x504c1d8: i8 = NVPTXISD::MoveParam 0x504c150<br/>+/*<br/>+bool recurseReversedCompare(uchar *reversed, uchar *normal, uchar ir, uchar in)<br/>+{<br/>+ if(reversed[ir] > normal[in])<br/>+ {<br/>+ return false;<br/>+ }<br/>+ else if(reversed[ir] < normal[in])<br/>+ {<br/>+ return true;<br/>+ }<br/>+ else if(in > 0)<br/>+ {<br/>+ return recurseReversedCompare(reversed, normal, ir+1, in-1);<br/>+ }<br/>+ else<br/>+ {<br/>+ return true;<br/>+ }<br/>+}<br/>+<br/>+bool reversedSmallerThan(uint reversed, uint normal)<br/>+{<br/>+ uchar *r = (uchar *)&reversed;<br/>+ uchar *n = (uchar *)&normal;<br/>+<br/>+ return recurseReversedCompare(r, n, 0, 3);<br/>+}<br/>+*/<br/>+<br/>+// test 2 : non recursive<br/>+bool reversedSmallerThan(uint reversed, uint normal)<br/>+{<br/>+ uchar *r = (uchar *)&reversed;<br/>+ uchar *n = (uchar *)&normal;<br/>+<br/>+ if(r[0] > n[3])<br/>+ {<br/>+ return false;<br/>+ }<br/>+ else if(r[0] < n[3])<br/>+ {<br/>+ return true;<br/>+ }<br/>+ else<br/>+ {<br/>+ if(r[1] > n[2])<br/>+ {<br/>+ return false;<br/>+ }<br/>+ else if(r[1] < n[2])<br/>+ {<br/>+ return true;<br/>+ }<br/>+ else<br/>+ {<br/>+ if(r[2] > n[1])<br/>+ {<br/>+ return false;<br/>+ }<br/>+ else if(r[2] < n[1])<br/>+ {<br/>+ return true;<br/>+ }<br/>+ else<br/>+ {<br/>+ if(r[3] > n[0])<br/>+ {<br/>+ return false;<br/>+ }<br/>+ else<br/>+ {<br/>+ return true;<br/>+ }<br/>+ }<br/>+ }<br/>+ }<br/>+}<br/>+<br/> __kernel void search( const uint block0, const uint block1, const uint block2,<br/> const uint state0, const uint state1, const uint state2, const uint state3,<br/> const uint state4, const uint state5, const uint state6, const uint state7,<br/>@@ -209,13 +289,13 @@<br/> G+=0x1f83d9ab;<br/> H+=0x5be0cd19;<br/> <br/>- if((H.x==0) && (bytereverse(G.x)<=target))<br/>+ if((H.x==0) && reversedSmallerThan(G.x, target)) // (bytereverse(G.x)<=target)<br/> {<br/> output[0] = 1;<br/> output[1] = nonce.x;<br/> }<br/> <br/>- if((H.y==0) && (bytereverse(G.y)<=target))<br/>+ if((H.y==0) && reversedSmallerThan(G.y, target)) // (bytereverse(G.y)<=target)<br/> {<br/> output[0] = 1;<br/> output[1] = nonce.y;<br/></div><br/>There are probably some optimisations to make <img alt="Smiley" border="0" src="/static/img/emoticons/smiley.gif"/></div> text/html
https://whatsonchain.com/tx/e4464cdff95d20404be55cf5a4dc3063e8deaf4e7c6f3dfb5f5deaf1a34304f1