Topic: [ANN] sgminer v5 - optimized X11/X13/NeoScrypt/Lyra2RE/etc. kernel-switch miner - page 129. (Read 877859 times)

I don't know why GPU-threads was set to 1... I just switched to 2 and went to 130Kh/s which made lower xintensity values usable. Before that I couldn't break 80Kh/s with xintensity below 7. It seems also that xintensity 2 gets me 2 KH/s more than xintensity 3. Worksize 96 still beats any multiple of 32 or 24 by at least 15KH/s.

So here's what I got.

"name" : "WestHash_NeoScrypt",
                "url" : "stratum+tcp://stratum.westhash.com:4341",
                "user" : "X",
                "pass" : "X",
                "algorithm" : "neoscrypt",
                "worksize" : "96",
                "gpu-threads" : "2",
                "xintensity" : "2",
                "thread-concurrency" : "8192"

And The Grand total is... 142KH/s! Thanks a bunch insanid!

Wolf0, if I were to drop the driver down to 13.12 wouldn't most other algos suffer? I thought X13 and X11 do much better on newer drivers?

I wish I could help you with the Scrypt-N HW errors, but I have no experience with it As far as the 270xs: 123 Kh/s isn't terrible, but I know (with FTCs implementation of neoscrypt at least) that it is possible to get higher hash rate out of the 270x; like in the range of 130-140 Kh/s or so from what I recall. I have 7850s and 7870s, which are very similar to the 270x, so maybe try using these settings to get more hashes out of them: gpu-threads2, worksize64, lookup-gap2, thread-concurrency8192, xIntensity3 also you don't need nfactor for neoscrypt, so take that out.

Hey that fixed the problem. It put me up to 97KH/s on each card.
After that I started messing with the intensity and worksize and got up to 123Kh/s per card. Has anyone gotten higher with 270X's? (Core at 1180MHz for now)
{
                "name" : "WestHash_NeoScrypt",
                "url" : "stratum+tcp://stratum.westhash.com:4341",
                "user" : "X",
                "pass" : "X",
                "algorithm" : "neoscrypt",
                "nfactor" : "10",
                "worksize" : "96",
                "gpu-threads" : "1",
                "xintensity" : "9",
                "thread-concurrency" : "8192"
        },

Also, I'm still having issues with Scrypt-N. In 3 minutes each card has about 500 HW errors. Getting about 175KH/s per card.

{
                "name" : "WestHash_Scrypt-N",
                "url" : "stratum+tcp://stratum.westhash.com:4335",
                "user" : "X",
                "pass" : "X",
                "algorithm" : "zuikkis",
                "nfactor" : "11",
                "xintensity" : "4",
                "thread-concurrency" : "5152"
        },

Actually, I just tested good 'ole Scrypt out and it's getting about twice the amount of HW errors at twice the hashrate. I used to mine Scrypt at 4MH/s with my four 270X's but now I'm getting 1.39MH/s. I think that's from a driver change though.

Yay! Made it to 4.4MH/s+ X11 on my 270X

Screenshot (NSFW): https://ottrbutt.com/miner/x11localrigwolf-11182014.png

Hey guys, maybe you can help me out here. I've been struggling to get NeoScrypt to work correctly for a few weeks. I've got 4x270x's set at 1150MHz Core and 1250MHz Memory (tested and 100% stable) and for the life of me I cannot get a single accepted share using Neoscrypt. It says it's hashing away giving me around 79Kh/s per card. But no shares are being accepted. I've experimented with different intensities, worksizes and even nfactor. I should also note that NeoScrypt is giving me roughly one to three HW errors per card every minute.

Scrypt-N is giving me a ton of HW errors as well, like 10 per second though I still get around 420KH/s total.  I'm using sgminer 5.1 dev 2014-11-13 from westhash. Also, please let me know if I've got any "outdated" settings in here, I haven't been messing with the other algos for a few months.

please do.


I fail to see your point here, so I pass.


Cool. We all know that. I used sgminer terminology as I was discussing with someone proficient with that code. Perhaps I could suggest that you go over sgminer codebase and make all appropriate changes there. Most will be grateful and I will buy you lunch if you choose to do so.


You must realize that dispatch of a NDRange call has nothing to do with actual execution of a kernel in a queue. I assume my writing wasn't clear enough so if you let me clarify: I was writing about kernels executing not NDRange calls.


Ok. My idea may be a dead end, I can live with that.


True. Did you have any sucess using HSLS compiler for this purpose? If that's the way to go, let's try.


Huh, this sounds like a description RISK architecture I guess again this is a compiler issue as compiler is supposed to hide this complexity...

May I chime in?
At the cost of being considered nazi, I have to point out there is specific terminology for those: instruction cache (often "I$") and data cache.
They are designed to be efficient in ASIC hardware or FPGA at most. The two problems here are:

• massive I$ overload, because the AMD compiler is too dumb to not unroll stuff (as a side note: HLSL/GLSL compilers usually unroll much more smartly, I currently suspect HLSL compiler might be building a whole tree of possibilities).
• registers must be shuffled across Work Items so most values cannot really be in private memory which brings us to the magic world of LDS layout.
• register pressure: how soon you need the result. To my own surprise it seems GCN 1.0 and 1.1 still cannot dispatch dependent instructions one after another

My understanding regarding barriers is that if one thread hits it then all of them have to in order for kernel to continue execution. So it is a sync mechanism (within a workgroup anyway). I agree that putting "hashes" in local mem would require one small kernel at the beginning and one at the end to copy local data into global mem and this would make sense if concurrent access to global mem from all threads is indeed a bottle neck. You say it's not I just mention it for the sake of completeness.

I completely forgot about gpu-threads, you're right that it keeps gpu busy. This is equivalent of just having 2 command queues that schedule kernels in parallel between them and in a sequential order within one gpu-thread .

BTW why is there a barrier in each kernel anyway? Since currently command queue does not allow for parallel execution of different kernels surely writes of blake should complete before next kernel bmw starts work...

My idea was that we could have some kind of say per thread counter/semaphore associated with "hashes" that would allow a kernel function searchN to start it's work only when this counter/semaphore was of value N. And this counter would be set at the end of each searchN fucntion to N+1. I see now this may not work as it maps poorly with what gpu was designed for.

Anyway I realize I need to spent some time with CodeXL for a while to gain some insight... what beats me is that all hash algos in X11 are designed to have efficient implementation in hardware. So they should be small in code and consume little memory. This kind of thing should be possible to implement directly in thread registers or cl terms "private memory".

No, that's not what I had in mind. Let me explain in detail for the benefit of other in the forum (and to verify my understanding):

The way I see it X(n) algos are much less suitable for GPU mining than Scrypt(n) due to the fact that composition of kernels produces a large code which GPU cannot execute efficiently, mostly due to limited cache size & cache algorithms. So plain vanilla approach to programming a kernel will not be efficient. Example of this is original "darkcoin.cl".

Your approach (so called "darkcoin-mod.cl") was different in the way that you do not schedule whole X11 to be executed an one single go, but in that you have had to split X11 into 11 individual components (in fact hash algos that comprise X11) which execute one after the other thus resolving cache issues as all of the 11 components are small enough to fit into the cache. But now you're scheduling execution of 11 kernels (via 11 "clEnqueueNDRangeKernel" calls).

So each call to the sgminer "opencl_scanhash", function where all the fun is, starts with the job setup, and GPU memory contains block data which needs to be hashed.

Then "blake" hash is executed over that data in all threads. I noticed you use xIntensity of 64 so he would execute 64*2048 "blake" global threads on say R9 280X card. Results of all these threads is stored into global memory. Since all instructions are executed more or less in lock-step (lock-step within a compute unit and possibly out of sync between compute units) there is always a lot of contention for memory access at specific times. Finally GPU waits for all threads to finish via the "barrier" command.

Then the GPU pauses a bit and waits for sgminer to enqueue another kernel which in our case is: "bmw". Global memory contains "blake" hashed block data.

It then proceeds to execute all the threads for "bmw" hash function. First it loads "blake" hashed block data from memory into registers or whatever compiler decides where the variables should be stored, does the thing and saves result (bmw hashed blaked hashed block data) into global memory again.

The whole process repeats for all the kernels (11 of them) and the final hash function checks whether resulting hash of hash of hash ... of block data is below target and if so returns the result which sgminer reads checks - calculates again only using CPU this time if all is ok submits to the pool.
(As we have very few results (hashes below target) per run this last CPU part is not an issue.)

So there are two major points of inefficiency in this approach if we can call it that as this "darkcoin-mod" approach is 40% faster than the original darkcoin.cl:

#1 GPU is not doing much between kernel schedules. This effect is even more problematic as all kernels end with a "barrier" statement thus waiting for all of threads to complete.
#2 kernels are spending time getting and storing data into global memory. Ideally this data should remain in GPU registers but I'm not sure this is possible with different kernels approach.

You say #2 is non issue and possibly I could agree with that. Although one could argue that as all the threads wait for the memory at the beginning for example you cannot take advantage of "waves in flight" as all waves will hit the same obstacle at the beginning of the kernel.
Question for you: why didn't you set "hashes" variable to reside in the local memory?

My thinking is that #1 is a real pain.

One way around it would be to schedule all 11 kernels in parallel on different compute units, some kernels would need more compute units than others due to their relative speed difference as you suggested.

I have another idea:
1. you don't store results from different kernels to the same memory area, but to different area for each algo
2. you can schedule all of them for execution in parallel (well as you currently do in "opencl_scanhash" function except that clCreateCommandQueue should specify "CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE")
3. kernels should be modified to wait for event to triger them. This event would be a signal from the previous kernel that it had finished the work. I'm not sure if this can be done via clEvents as they seem to work on entire queue rather than individual thread.

This would allow for kernels not to have to wait for all of them to finish (no end barrier) as the signaling to triger the next one would be explicit signal.... Reading OpenCl documentation I'm not sure this is possible...

Thank you; it seems Hippie Tech has a bone to pick.

When mining scrypt with 290(x) reference cards there were sweet spots where hash was maximized.  Before Stilt's bios, the reference cards could only mine well with 1250 or 1500 memory and you just adjust the gpu clock to the sweet spot, I think it was usually a 1.4:1 ratio.  After flashing Stilt's bios you could change the memory in steps of 25 mhz and adjust the gpu again.  Maybe there are sweet spots for x11 type coins.  I haven't had time to test much so I have no idea.

I'm running all my cards at 1250 Mhz memory, in same instances I get better hashrates but power consumption are always lower. I guess it's lower memory latency at those clocks...

OK i assume 50W is your idle draw, so that leaves us with 3450kH/58J~=60kH/J
If same ratios were to apply to 290X then 290X should do Cores_2816/1280 * Freq_1018/900 * 3450 kH/s ≃ 8600 kH/s, pull from the wall should be 145W. For X2 those numbers would be double.

This is phenomenal. Although I remember seeing your X2 hash rates to be even higher.

Yes that's what I asked: if "hashes" variable were a local, then only those threads that belong to the same workgroup (64 threads) would have access to this common local resource.

I get the code cache argument, but then this leads to significantly different design, namely different workgroups should run different kernels. This we would have to tie in into sgminer gpu scheduling code.

I'm not as proficient with opencl as I would like to be, but looking at scrypt kernel implementation or SHA256 for that matter (at least from a year ago) all seemed to be rather simple: GPU was treated as a flat SIMD resource with say 2048 threads (7970) and global memory, so we would just schedule 8192 simultaneous threads (-I 13) , or more, so that hopefully we use up all computing resources.

X11 kernel (moded one) looks very different. The way I see it I should somehow take advantage of the fact that work scheduled to different workgroups is in fact executed in MIMD fashion.

Anyway I truly enjoy these conversations and I thank you for it.

Still X2 works like a charm. Currently it runs on stock clocks, 0.9V, 56C, about 150W/core and delivers 5.4MH/s/core....or: 35 kH/J. Are you saying you're doing 60+ kH/J

Yes, I understand. BTW is there an easier way to pass a parameter to "searchX" function other than through global memory? Wouldn't local memory work as long as all threads are in the same local group/compute unit?

Still I wonder: would'n it be better for a single thread to compute the entire hash (all of the 11 functions) rather than having multiple threads evaluating different functions of X11, from a message/argument passing standpoint? (with the former approach you would not need any).

Check the FTC forums - they have a cl for old cards.

Reduce memory usage, reduce power.

I see. Well I'll look into kernel files in the next few days and will post the results if there are any ...
... I see where the main problem is: darkcoin_mod.cl consists of 11 functions one for each algo in X11 (search0-10)  all of which end up writing results of the global memory. This is slow and on top of that all threads compete for writes in the global memory.

Interstingly original darkcoin.cl does not do that and still it is slower. Perhaps other optimizations of darkcoin_mod make more of an impact.