Topic: [ANN] cudaMiner & ccMiner CUDA based mining applications [Windows/Linux/MacOSX] - page 34. (Read 3426980 times)

will be cleanup to do in the next version so... nice for the final binary size

This one should work:
__inline__ __device__ uint2 ROL2(const uint2 v, const int n) {
      uint2 result;
      if (n <= 32) {
         result.y = ((v.y << (n)) | (v.x >> (32 - n)));
         result.x = ((v.x << (n)) | (v.y >> (32 - n)));
      }
      else {
         result.y = ((v.x << (n - 32)) | (v.y >> (64 - n)));
         result.x = ((v.y << (n - 32)) | (v.x >> (64 - n)));
      }
      return result;
   }

there is a way to do it without a condition statement, but haven't looked into it yet

yes, this code is wrong.

__inline__ __device__ uint2 ROL2(const uint2 v, const int n)
{
   uint2 result;
   result.x = (((v.x) << (n)) | ((v.x) >> (64 - (n))));
   result.y = (((v.y) << (n)) | ((v.y) >> (64 - (n))));
   return result;
}

it is based on post killer groestl code... the only reason it is faster that groestl 512 is that it does less mixing (by definition small_core versus big_core in sph definition). In principle a killer groestl256 (bit slicing) should be faster... but it isn't obvious to write even from the existing code...

I also tried a 64bit implementation with uint2, but it wasn't faster either...

I think I should get uint2 working for compute 3.0, there was a bug (actually it was just wrong ) in the uint2 rotation without funnelshift...

The groestl(256) in lyra implementation is different. Is it faster than the groestl in x11(512). I guess it needs some more work to support 512 rounds?

current status is: "not yet"
actually, I was thinking to implement neoscrypt before I got hired to work on lyra.
But I must admit I kinda lack motivation on this project... (may-be trying to beat Wolf0 will help  )
If there was a bounty for nvidia that would definitely help  (greedy me...)
But for the moment, I am still looking into improvement into lyra... once I am done with that may-be...

I think the third party code forks around here have even more hash rate improvements.

I've reduced my involvement in the mining scene a lot since last summer.

The last thing I did was a GPU miner for CoinShield's CPU channel (together with ChrisH), but that was never made public. Also it looks like CoinShield was quite a dud (i.e. the coin's value is not taking off and the feature rollout of the wallet is really slow)

Check the head on github. I committed the 50KHASH on echo improvement. You should now get around 3MHASH@x11 on the stock clocked windforce black 750ti.

NEOSCRYPT-

DJM34, I now notice that your version of ccminer has blocks of code labelled Neoscrypt.  Would you be able to give a status report on how developed this code is?  Is it possible to mine neoscrypt on GPU with your ccminer?       --scryptr

Christian,

A while back I posted a question regarding performance gains to ccminer versions newer than the v1.0 that I was using. Recently I had to reload my computer from scratch (Xubuntu 14.04) so I compiled v1.2 with Cuda 6.5, driver 340.29. Whereas with two GT640 (GDDR5) mining BTQ I was getting ~1.6Mh/s combined, I'm now getting ~2.1Mh/s. 31% boost for zero added hardware. There will be a little something from me in your stocking soon. Thanks a bunch! 

Z

yes but the number of uint32 operations over a cycle is larger than twice the number of uint64 operation (I think... something like that, a table was posted recently on that on the thread)

In BMW I added a couple of operators. - And SHL/SHR for the Uint2 types. The funnelshift on nvidia is only calculating 32 bit a time so you need 2 instructions to do a full rotate, and then you need code that merge the Upper/Lower registers. In code it meens that a ROL64 4-5 instructions while a shL64 is 1 instruction.
When working in uint2 mode, alot of the merge code 32bit->64 bit is removed. This meens that the ROL64 is down to 2 instructions. For bmw that caused the final routine to use less registers (no memory spills).
But 33% faster was a surprise for me, since the bmw has a lot of SHL/SHL/ + and - wich I suspected would run slower on uint2. (2 instructions instead of 1)

I tried to use it for the whole routine, but for some reason I haven't been able to get it to work...
(actually not sure it will help a lot, the amd is a lot stronger than the nvidia on 64bit calculation, that's why it makes a such a difference on nvidia)

There we are !) good job... i see 2950 with your current git on linux and 2905 on mine, i think i will grab skein one too for the 1.5.1 release... but all your merges with KlausT make your work really hard to follow

We were talking about uint2 possible improvements some days ago with djm on irc... nice you made it :p

For the 970 and 980 is you use compute 5.2, it can be a bit tricky (register allocation is a bit tricky with uint2 and compute 5.2)

It isn't that I don't want you to publish my code, but I prefer to be the first (aside from the VTC team release...) to do it for obvious reason (my own publicity obviously... I guess this one is really obvious )

The way it happened (or my perception of the event): I sent the code to vtc team for the test release/release candidate before going to bed had some run of correction before they decided to go to sleep and when I woke the next morning someone (you) has published my code with no clear reference to the author on another forum where I don't have an account his own version of the code. Sorry by without enough coffee it looks pretty sneaky, and with more coffee it still looks sneaky. (and I am pretty sure you would have reacted in the same way, if someone had done that to you).

Then, I told you to wait (or asked to jk_14 to ask you to wait as I didn't have an account on litecointalk) because it wasn't finished and as it was a test release it is way better if people report problems from only one release and don't start to use code we don't know... (I also told you to wait on irc channel, as I was planning to move the code to your interface )

Then later in the week you sent your code (actually the same, while the RC had already incorporated new parts) to the vtc team with no clear intention since they asked me what they should do with that. That time it was just looking like you were trying to go behind my back on the job I was hired to do...  

That's why I feel a bit pissed off about the whole situation, and I am pretty sure you would have felt the same if you had been in my situation.
I don't think all this was intentional from you, but it was kind of careless...

anyway...

Regarding the difficulty 128 or 256 should do it, just check the pool hashrate and use the one which give the result closest to the real hashrate.

Not sure. Will test this weekend on the 970 and 980.
Note that with the last improvements, and my last echo change (wich is not checked in yet)  a windforce 750ti black edition should do above 3MHASH on stock clocks mining x11.

not bad, what are the numbers on the 9xx serie ? In principle it should be even better however it might require some tuning...

me, i was not paid by vtc team to do it... but somebody was not able to compile it on linux... its also why i imported it.

The next time you dont want i use published code (by VTC not you) tell us before we commit it... not 2 days after. I was planning to release it today, but... the default difficulty setting is still not clear... 128 and 256 seems ok on most pools ive tested.

Sorry for the confusion djm. It yielded better results than miner listed in VTC thread, so I assumed its indeed new

Thanks for the explanation.