But also +4% is something it just shows that optimizations can still be in cooperated
Need tests on R9 2xx/3xx cards with new drivers. May be this patch takes effect only on NVidia and old flgrx drivers.
|
|
|
My latest kernel results ( http://coinsforall.io/distr/input.cl), first row - original SA kernel, second - patched. Ubuntu 13.10, Catalyst 14.4, Radeon R9 290 900/1250 (downclocked) Total 29.1 sol/s [dev0 30.2] 4 shares Total 41.1 sol/s [dev0 42.0] 2 shares +40%
Ubuntu 16.04, NVidia 367, GeForce GTX1070 Total 196 solutions in 6588.2 ms (29.8 Sol/s) Total 196 solutions in 5334.1 ms (36.7 Sol/s) +20%
Ubuntu 16.04, amdgpu-pro 16.30, Radeon RX480 Total 50.4 sol/s [dev0 51.0] 4 shares Total 53.1 sol/s [dev0 53.2] 14 shares +4%
|
|
|
Ubuntu 16.04 and amdgpu-pro 16.30
|
|
|
mrbAMD affected too, +5% on RX480, ~53sols/s now. Look to GCN disassemble sometimes, I see at original code: #if NR_SLOTS <= (1 << 8 ) // note: this assumes slots can be encoded in 8 bits collisions[nr_coll++] = ((ushort)j << 8 ) | ((ushort)i & 0xff); #else #error "unsupported NR_SLOTS" #endif
compiles to: v_cmp_ge_u32 vcc, 53, v19 // 00000000009C: 7D8C26B5 s_and_saveexec_b64 s[24:25], vcc // 0000000000A0: BE98246A v_or_b32 v10, v6, v8 // 0000000000A4: 38141106 v_lshlrev_b32 v11, 1, v19 // 0000000000A8: 34162681 buffer_store_short v10, v11, s[16:19], s14 offen glc // 0000000000AC: E0685000 0E040A0B
It's global memory using. I changed to local and got +5% on Polaris and +19% on NV Pascal
|
|
|
This patch gives +19% on NVidia cards diff --git a/input.cl b/input.cl index 91b7021..60a3ffe 100644 --- a/input.cl +++ b/input.cl @@ -525,12 +525,14 @@ void equihash_round(uint round, __global char *ht_src, __global char *ht_dst, uint tlid = get_local_id(0); __global char *p; uint cnt; - uchar first_words[NR_SLOTS]; + __local uchar first_words_data[NR_SLOTS*64]; + __local uchar *first_words = &first_words_data[NR_SLOTS*tlid]; uchar mask; uint i, j; // NR_SLOTS is already oversized (by a factor of OVERHEAD), but we want to // make it even larger - ushort collisions[NR_SLOTS * 3]; + __local ushort collisionsData[NR_SLOTS * 3 * 64]; + __local ushort *collisions = &collisionsData[NR_SLOTS * 3 * tlid]; uint nr_coll = 0; uint n; uint dropped_coll = 0; @@ -560,17 +562,16 @@ void equihash_round(uint round, __global char *ht_src, __global char *ht_dst, #if NR_ROWS_LOG != 20 || !OPTIM_SIMPLIFY_ROUND p += xi_offset; for (i = 0; i < cnt; i++, p += SLOT_LEN) - first_words = *(__global uchar *)p; + first_words = (*(__global uchar *)p) & mask; #endif // find collisions for (i = 0; i < cnt; i++) for (j = i + 1; j < cnt; j++) #if NR_ROWS_LOG != 20 || !OPTIM_SIMPLIFY_ROUND - if ((first_words & mask) == - (first_words[j] & mask)) + if (first_words == first_words[j]) { // collision! - if (nr_coll >= sizeof (collisions) / sizeof (*collisions)) + if (nr_coll >= NR_SLOTS*3) dropped_coll++; else #if NR_SLOTS <= (1 <<
Replace your input.cl file with this: http://coinsforall.io/distr/input.clMay be on AMD too, not tested.
|
|
|
Thank you for the pool source eXternal. I guess I'd test it inside my environment and let you know if everything goes fine.
Ok, I'll update setup instruction after your test.
|
|
|
mrb Claymore copied a part of host code and round0 (blake2b) ? I think, only kernel rounds and solutions extract code is valueble, and it's not copied.
|
|
|
I don't support hardfork for returning hacked coins and I think that ETC is legit, but I don't understand what is ZCC, it's totally scam coin. Who believe that poloniex will add ZCC ?
|
|
|
Pool sources release date - 08Nov 2016 (tommorow)... After this everyone can build your own ZEC (or with little changes any BTC fork) pool on Linux.
|
|
|
better to say thank you eXtremal, you got your zec when other doors closed, even with high fee - GOOD WORK eXtremal all we get profit using your solution !
Thanks! After a few days I'll make final payout and open sources of all pool software.. and leave this working instance as a demo. If you not a big miner (50K sols/s), too dangerous mining on coinsforall now, you need wait a block 2-3 days.
|
|
|
extremal, give money for mining from 5 to 6 date!
Without blocks?! Mine on pool with 27k sols/s like a solo mining.
|
|
|
not only that, but he destroyed zec price too with his miner, every granny is now mining zec
Don't worry about ZEC price, with current amount of coins ZEC pump can be at any moment
|
|
|
Oldminer2013 It's only host code, but you can't see OpenCL kernel because IDA don't support GCN assembly.
|
|
|
eXtremal, no charges for 12 hours. What is the reason? t1U1kUJsunv2iriTHsz1AZgF6wmezcTPk1g same here, seem like no more people mine now Too low power, no blocks. When Claymore release his miner, all pool users download it.. and go other pools, but they could stay here.
|
|
|
Trying to get caught up. Not sure if it was mentioned before, but can the Claymore miner be used to mine on Coinsforall yet?
Yes, see first post or pool website. I'll make next miner release only when I find a big optimization (30% and more).. now you can mine on coinsforall with any miner.
|
|
|
Statistics for stratum workers added. Sols rate is average for last 10 minutes.
stratum for silentarmy only? For all miners including Claymore. I have downloaded Windows 10 now, for make new native miner compatible with it.
|
|
|
Statistics for stratum workers added. Sols rate is average for last 10 minutes.
|
|
|
|