Hi Christian, me again. dunno if im barking up the wrong tree here but if i reduce the amount of cuda threads, i.e:

'case 16: fermi_scrypt_core_kernelA<16><<< grid, threads, 0, stream >>>(d_idata); break;'

Say i set threads to  say 256 (<512). theres a massive increase in speed... But quite a few errors.

Why the errors?

Cuda is still new to me....

http://s22.postimg.org/qdcboxcwh/Capture.jpg

-H 2 is for CPU's that cant deal with the SHA256 workload
basically don't use it unless you have a hashrate that causes cudaminers cpu load to go through the roof.

Still crashed cicc when sm_20, went ok with sm_13. Testing against 2013-12-01, 2013-12-07 spinlock is about 5 - 6 khash/s faster.

With fermi kernel 2013-12-01 with the added offset is around 233, 2013-12-07 is about 228. I tested 2013-12-07 with both your precompiled exe and self built from source, results are identical.

Id like to see feedback from more fermi owners. In particular compute 2.0 fermis and even more particular 448 core gtx560ti.

Well i have a pallit one anyway. Cheapest i could find but that shouldn't cause my results to be totally different to yours. I don't think its a problem with os or drivers because i dual boot 32 bit debian and get the same results there.
Although i cant really see it, maybe my weird ddr3 + fsb + pcie 1.1 abomination is causing it.

I have also noticed a drop in performance since switching to the ulong2 method in fermi. Most noticable in x64. Also, this maybe of no concern of yours as you don't really support cuda 5.5 but, the spinlock kernel crashes the cuda 5.5 cicc compiler when compiling for compute 3.0.

For solo mining,

my coind command: -server -rpcuser=aaa -rpcpassword=123456 -rpcport=8100

my cudaminer command: -o http://127.0.0.1:8100 -O aaa:123456 -C 1 -i 0 -H 1 -l F14x16

I've done a bit of reading on cuda/fermi and bank conflicts. Done a bit of experimentation. I found that the reason i was getting higher performance from my 64 bit builds was due to the extra byte memory alignment offset for 64 bit in the fermi kernel. I found that adding the same offset to the 32 bit kernel brought it up to speed with the 64 bit one. Some reading suggested that 36 byte alignments were optimal for avoiding bank conflicts. I have managed to get up to 234khash/s in linux. I have tested it on a fresh linux install with cuda 5.0 and confirmed that it's not just a quirk with my 5.5 windows setup.

Christian, you say you get 228 on your 560 ti 448 in 32 bit. I only get about 222 with 32 bit in windows or linux but this has pushed it to ~234. I could see yours doing 240 providing mine isn't somehow retarded and i found a way to work around it's retardedness.

Easy way to test,

#define _64BIT_ALIGN 3

whether it be for 64 or 32 bit. can easily be set back to 0 or 1 after if it doesn't really do anything.

I think you should get more than that as i get 230Kh/s from a 560ti 448 core. Maybe try F16x16? oh and -i 0.

Did you try the -H 1 switch? Assuming you have a dual core or more. The 64 bit build seems to use more cpu and may benefit from it more or less depending on cpu speed.

aaaand a new win 64bit exe

http://www.filedropper.com/cudaminer2013-11-20x64cuda-55

Say thanks to cbuchner1, he made this all possible....

Me, well i just pressed a few buttons.

Have fun.

latest git x64, cuda 5.5 windows .exe

http://www.filedropper.com/cudaminer2013-11-14x64cuda-55

valid for 30 days

Probably not. I could only touch about 8kh/s with a lightly overclocked 9500GT. But hey, at least its faster than my 8400GS @ 1.5kh/s.

Is there any more gain to be made with more warps per block?

The increase from 8 to 16 made a few khash/sec difference but i find with -i0 that my system is still responsive. I would expect it to be really laggy if the card was being pushed to its limits.
I would try it myself but my coding skills are still very limited when it comes to cuda and crypto.

function blah uses too much shared memory i think was the error i got

Hint: compile for "compute_10,sm_20" instead of "compute_10,sm_10"

I can but, don't be surprised if it craps out on anything other than a 560ti 448.

http://www.filedropper.com/cudaminerx64

I have compiled the new repo for 64 bit. Meant i had to compile a 64 bit libcurl but now im getting up to 228khash/s but at the same time getting some nasty cpu usage.
This is a 2.4GHz Core2Quad overclocked to 3GHz.

http://s24.postimg.org/yfxb7fa4l/Capture.jpg

Adding an extra sleep before MyStreamSynchronize returns seems to keep it down without impacting performance.

http://s18.postimg.org/mr5rvzk49/Captureb.jpg

I can make two suggestions. Make sure the -C flag is optimal for your card. I get best performance using -C 1. You could try -i 0 but this might make desktop lag unbearable, it depends on the power of your card.

I am actually using  -C 1 -i 0 -l F14x16 -H 1. What are your clock speeds? Here i have stock 732/1900.

I'll try downloading your updated repo and a fresh compile.

*edit* No, still the same.

Hi Christian. I've tried your new repo under windows. Just needed to #undef HAVE_ALLOCA_H and HAVE_SYSLOG_H in cpuminer-config.h to get it to compile under vc2010 and then changed max_warps_per_block() to return 16 in fermi_kernel.h to make use of the extra WARPS_PER_BLOCK.

There is some increase in reported speed. Before, i was using 28x8 with a GTX560Ti 448 core and peaking around 210khash/s. Now using F14x16 im getting up to 218khash/s.

Not much more but still more. Great work, thanks.

Just out to burn some watts... Maybe make a little money on the way.

Good luck.