Author

Topic: OFFICIAL CGMINER mining software thread for linux/win/osx/mips/arm/r-pi 4.11.0 - page 663. (Read 5805728 times)

member
Activity: 266
Merit: 36
Proofer it will only go to 100% at the temp-overheat value which you have set to 77.

OK, I just removed the gpu-fan parameter and am leaving for the gym.  If I come back and my house is a pile of smoking rubble you will be in BIG TROUBLE!

Ha, at least you will be to worn out to do much about it  Tongue

After about 1.5 hours I see no difference after removing gpu-fan.  I'll report back if I encounter any problems.
donator
Activity: 798
Merit: 500
Proofer it will only go to 100% at the temp-overheat value which you have set to 77.

OK, I just removed the gpu-fan parameter and am leaving for the gym.  If I come back and my house is a pile of smoking rubble you will be in BIG TROUBLE!

Ha, at least you will be to worn out to do much about it  Tongue
member
Activity: 266
Merit: 36
Proofer it will only go to 100% at the temp-overheat value which you have set to 77.

OK, I just removed the gpu-fan parameter and am leaving for the gym.  If I come back and my house is a pile of smoking rubble you will be in BIG TROUBLE!
donator
Activity: 798
Merit: 500
Proofer it will only go to 100% at the temp-overheat value which you have set to 77.
member
Activity: 266
Merit: 36
Must be something other than 5970+Linux:  I've been running 2.2.3 on a 3x5970 rig for several days without problems.  (Well, I did have one DEAD core but I had that occasionally with previous versions as well.)

With p2pool:
Code:
"intensity" : "9",
"gpu-threads" : "1",
"gpu-engine" : "810-810",
"gpu-memclock" : "200",
"gpu-fan" : "85",
"temp-cutoff" : "80",
"temp-overheat" : "77",
"temp-target" : "70",
"temp-hysteresis" : "3",
"auto-fan" : true,
"auto-gpu" : true,

Take this out and see what happens:
Code:
"gpu-fan" : "85",

I believe that will cause one unit's fan -- cores 0-1 -- to go to 100%, which I don't want to see; currently:
Code:
 [P]ool management [G]PU management [S]ettings [D]isplay options [Q]uit
 GPU 0:  72.5C 4593RPM | 364.7/363.8Mh/s | A:1559 R:45 HW:0 U: 1.89/m I: 9
 GPU 1:  76.0C 4593RPM | 365.1/364.2Mh/s | A:1562 R:48 HW:0 U: 1.89/m I: 9
 GPU 2:  68.0C 4438RPM | 365.1/364.2Mh/s | A:1614 R:44 HW:0 U: 1.95/m I: 9
 GPU 3:  63.0C 4438RPM | 365.1/364.2Mh/s | A:1568 R:45 HW:0 U: 1.90/m I: 9
 GPU 4:  65.5C 4493RPM | 365.0/364.1Mh/s | A:1580 R:56 HW:0 U: 1.91/m I: 9
 GPU 5:  68.5C 4493RPM | 365.2/364.1Mh/s | A:1493 R:45 HW:0 U: 1.81/m I: 9

However, if you really believe the experiment may be of use to the community, let me know and I'll do it.
donator
Activity: 798
Merit: 500
Must be something other than 5970+Linux:  I've been running 2.2.3 on a 3x5970 rig for several days without problems.  (Well, I did have one DEAD core but I had that occasionally with previous versions as well.)

With p2pool:
Code:
"intensity" : "9",
"gpu-threads" : "1",
"gpu-engine" : "810-810",
"gpu-memclock" : "200",
"gpu-fan" : "85",
"temp-cutoff" : "80",
"temp-overheat" : "77",
"temp-target" : "70",
"temp-hysteresis" : "3",
"auto-fan" : true,
"auto-gpu" : true,

Take this out and see what happens:
Code:
"gpu-fan" : "85",
member
Activity: 266
Merit: 36

Didn't' forget about this  Tongue I ran 2.1.2 with no errors for 18 hours then started 2.2.3 yesterday with the same flags and got this:
(...)
Weird how las initialised is slightly after start time, but it wasn't disabled for a few hours.  You might be on to something here.

I had to stay on 2.1.2. It's weird just us having reported this problem.
I believe that when people with 5970 rigs running Linux update cgminer, more posts like ours will appear.

Must be something other than 5970+Linux:  I've been running 2.2.3 on a 3x5970 rig for several days without problems.  (Well, I did have one DEAD core but I had that occasionally with previous versions as well.)

With p2pool:
Code:
"intensity" : "9",
"gpu-threads" : "1",
"gpu-engine" : "810-810",
"gpu-memclock" : "200",
"gpu-fan" : "85",
"temp-cutoff" : "80",
"temp-overheat" : "77",
"temp-target" : "70",
"temp-hysteresis" : "3",
"auto-fan" : true,
"auto-gpu" : true,
legendary
Activity: 1988
Merit: 1012
Beyond Imagination
2.2.3 had strange behavior on my fixed-fan-rate rig, 10% lower hash rate on 5870s and crash on 5970
-ck
legendary
Activity: 4088
Merit: 1631
Ruu \o/
Working on the 7970 tuning, I have tried to port both the diapolo and diablo kernels to cgminer. Alas neither of them are actually working yet, so instead I started modifying the existing poclbm kernel in cgminer to improve throughput. This should work on other GPUs as well as the 7970, but I have no idea if it's better or worse than phatk.

When it's released it will get a new date/version number, but I haven't changed the number right now so that people can download it now and give it a try:
https://raw.github.com/ckolivas/cgminer/kernels/poclbm120203.cl

Remember to delete any .bin files and if you're not on a 7970 with the latest cgminer, you'll have to tell it to use that kernel with -k poclbm.

So what's the damage? Well on the 7970 at 1200/1050 clocks, which was getting 694MHash, it's now getting 711Mhash. The 7970 has this unusual behaviour where the hashrate slowly rises for the first 5-10 minutes.
legendary
Activity: 1320
Merit: 1001
Working on i t... workaround for now is to set fan to fixed speed and disable --auto-fan

Thank you, ck. No need to hurry.
I know you are very busy these days.

newbie
Activity: 9
Merit: 0

Didn't' forget about this  Tongue I ran 2.1.2 with no errors for 18 hours then started 2.2.3 yesterday with the same flags and got this:
(...)
Weird how las initialised is slightly after start time, but it wasn't disabled for a few hours.  You might be on to something here.

I had to stay on 2.1.2. It's weird just us having reported this problem.
I believe that when people with 5970 rigs running Linux update cgminer, more posts like ours will appear.


Same here, 2x 6770 on Ubuntu 11.04  ~ 443 MH/s

With 2.2.3 GPU 0 stops working after a few hours. When i restart cgminer it works again for a few hours.
With 2.1.2  both cards are working fine since 30 hrs now...
-ck
legendary
Activity: 4088
Merit: 1631
Ruu \o/

Didn't' forget about this  Tongue I ran 2.1.2 with no errors for 18 hours then started 2.2.3 yesterday with the same flags and got this:
(...)
Weird how las initialised is slightly after start time, but it wasn't disabled for a few hours.  You might be on to something here.

I had to stay on 2.1.2. It's weird just us having reported this problem.
I believe that when people with 5970 rigs running Linux update cgminer, more posts like ours will appear.

Working on i t... workaround for now is to set fan to fixed speed and disable --auto-fan
legendary
Activity: 1320
Merit: 1001

Didn't' forget about this  Tongue I ran 2.1.2 with no errors for 18 hours then started 2.2.3 yesterday with the same flags and got this:
(...)
Weird how las initialised is slightly after start time, but it wasn't disabled for a few hours.  You might be on to something here.

I had to stay on 2.1.2. It's weird just us having reported this problem.
I believe that when people with 5970 rigs running Linux update cgminer, more posts like ours will appear.
hero member
Activity: 772
Merit: 500
Edit: Perhaps we could try my old approach of writing to output in the kernel, because I know that worked for me?

That's the code I used, but uses your NFLAG. It would need to scan the output buffer on host side everytime after a kernel execution, which could lead to higher CPU usage (and needs changes in host code), but saves the IF-clause and another write into output (which saves the kernel quite some instructions, even on GCN).

Code:
u result = (V[7] == 0x136032ed) * nonce;
output[NFLAG & result] = result;

This code would be more like your current code, but uses the approach of comparison and mul to save 0 or a positive nonce in result (and is slower than your current code). But for sure that can't be the problem we are looking for ...

Code:
u result = (V[7] == 0x136032ed) * nonce;
if (result)
output[FOUND] = output[NFLAG & result] = result;

Dia
Writing to output on every iteration isn't going to fix the problem, and I can't see how this would help to be honest. Note that your last code will end up setting output[FOUND] to 0 and would undo anything you wrote to it with other threads  Wink

You are right on the last code, at least I did no commit for it :-P.

Okay, phatk has const u base, which DiaKGCN has too, if GOFFSET is not set, which is currently always true.

phatk:
Code:
base + get_local_id(0) + get_group_id(0) * (WORKSIZE);
diakgcn:
Code:
((uint)get_group_id(0) * (uint)get_local_size(0)) + (uint)get_local_id(0) + base;

I have the (uint) cast, because the returned size_t can have 64 bits on some systems. The brakets I have for group-id * local_size should be unneded, because of * before +. Guess no problem here and as we discussed, output writing should be okay for now.

Dia
hero member
Activity: 772
Merit: 500
Okay, so as I wrote, if Phatk works, then the base-nonces passed to the kernel should be correct for diakgcn. I will check the phatk.cl to be sure. I saw you added a BITALIGN path to diakgcn, that's not using bitalign() or any other OpenCL function, but simply does it's thing directly. What is that for, i'm not sure if that's needed for a GCN kernel anyway Smiley.
Another idea, are you applying a BFI_INT patch on Tahiti (it must not use amd_bytealign())? This is not needed and produces wrong values ... I want that damn thing working Cheesy, I stared at it quite a few hours too ^^.
BITALIGN is to enable amd media ops for platforms that have it.

BFI INT patching does NOT work on Tahiti. It makes a corrupt kernel. SDK2.6 automatically uses the BFI INT instruction anyway so there is no need for this crappy patching.

That's what I said, to be sure that BFI_INT patching is DISABLED for Tahiti, I thought it coule be active, so that would have been a problem Smiley. BITALIGN flag is not needed for DiaKGCN, because amd_bitalign() is use nowhere ... cl_amd_media_ops is only needed for doing BFI_INT patching on non GCN hardware (where amd_bytealign() is "patched" into bfi_int instruction). The amd_bitalign() was used with former SDKs to speedup the rotations, these are now optimized via the OpenCL compiler into bitalign anyway.

Dia
-ck
legendary
Activity: 4088
Merit: 1631
Ruu \o/
Edit: Perhaps we could try my old approach of writing to output in the kernel, because I know that worked for me?

That's the code I used, but uses your NFLAG. It would need to scan the output buffer on host side everytime after a kernel execution, which could lead to higher CPU usage (and needs changes in host code), but saves the IF-clause and another write into output (which saves the kernel quite some instructions, even on GCN).

Code:
u result = (V[7] == 0x136032ed) * nonce;
output[NFLAG & result] = result;

This code would be more like your current code, but uses the approach of comparison and mul to save 0 or a positive nonce in result (and is slower than your current code). But for sure that can't be the problem we are looking for ...

Code:
u result = (V[7] == 0x136032ed) * nonce;
if (result)
output[FOUND] = output[NFLAG & result] = result;

Dia
Writing to output on every iteration isn't going to fix the problem, and I can't see how this would help to be honest. Note that your last code will end up setting output[FOUND] to 0 and would undo anything you wrote to it with other threads  Wink
-ck
legendary
Activity: 4088
Merit: 1631
Ruu \o/
Okay, so as I wrote, if Phatk works, then the base-nonces passed to the kernel should be correct for diakgcn. I will check the phatk.cl to be sure. I saw you added a BITALIGN path to diakgcn, that's not using bitalign() or any other OpenCL function, but simply does it's thing directly. What is that for, i'm not sure if that's needed for a GCN kernel anyway Smiley.
Another idea, are you applying a BFI_INT patch on Tahiti (it must not use amd_bytealign())? This is not needed and produces wrong values ... I want that damn thing working Cheesy, I stared at it quite a few hours too ^^.
BITALIGN is to enable amd media ops for platforms that have it.

BFI INT patching does NOT work on Tahiti. It makes a corrupt kernel. SDK2.6 automatically uses the BFI INT instruction anyway so there is no need for this crappy patching.
hero member
Activity: 772
Merit: 500
Well ... as I said, I have no IDE setup, so currently I can't compile a version for myself. If you don't have the time to fiddle around with my commits, then I really need help in setting up an IDE in Windows. Have you got this in a readme, wiki or can you give me a brief explanation in how to do this? I worked with MS VC++ Express as a hobby some time ago ...

You said local copy, is it a copy of the last version of my fork? As you've observed I am new to this kind of working, but I hope you see my progress Cheesy.

Dia
Compiling this on windows is nothing short of a DISASTER so forget it.

Anyway I fixed up a few things on my Diapolo branch on github. Pull the changes to bring your local tree into sync. Alas I'm still only getting HW errors, so there's clearly something wrong. The return code for giving me a nonce I use works fine, provided I'm testing for the right thing before sending the nonce back. I've stared at it for half a day and can't find what's wrong. I even tried diablo's kernel and encountered exactly the same problem. For some reason I keep thinking it's something to do with confusion about the initial offset of the nonce and what is passed to the kernel.

Okay, so as I wrote, if Phatk works, then the base-nonces passed to the kernel should be correct for diakgcn. I will check the phatk.cl to be sure. I saw you added a BITALIGN path to diakgcn, that's not using bitalign() or any other OpenCL function, but simply does it's thing directly. What is that for, i'm not sure if that's needed for a GCN kernel anyway Smiley.
Another idea, are you applying a BFI_INT patch on Tahiti (it must not use amd_bytealign())? This is not needed and produces wrong values ... I want that damn thing working Cheesy, I stared at it quite a few hours too ^^.

Edit: Perhaps we could try my old approach of writing to output in the kernel, because I know that worked for me?

That's the code I used, but uses your NFLAG. It would need to scan the output buffer on host side everytime after a kernel execution, which could lead to higher CPU usage (and needs changes in host code), but saves the IF-clause and another write into output (which saves the kernel quite some instructions, even on GCN).

Code:
u result = (V[7] == 0x136032ed) * nonce;
output[NFLAG & result] = result;

This code would be more like your current code, but uses the approach of comparison and mul to save 0 or a positive nonce in result (and is slower than your current code). But for sure that can't be the problem we are looking for ...

Code:
u result = (V[7] == 0x136032ed) * nonce;
if (result)
output[FOUND] = output[NFLAG & result] = result;

Dia
-ck
legendary
Activity: 4088
Merit: 1631
Ruu \o/
Well ... as I said, I have no IDE setup, so currently I can't compile a version for myself. If you don't have the time to fiddle around with my commits, then I really need help in setting up an IDE in Windows. Have you got this in a readme, wiki or can you give me a brief explanation in how to do this? I worked with MS VC++ Express as a hobby some time ago ...

You said local copy, is it a copy of the last version of my fork? As you've observed I am new to this kind of working, but I hope you see my progress Cheesy.

Dia
Compiling this on windows is nothing short of a DISASTER so forget it.

Anyway I fixed up a few things on my Diapolo branch on github. Pull the changes to bring your local tree into sync. Alas I'm still only getting HW errors, so there's clearly something wrong. The return code for giving me a nonce I use works fine, provided I'm testing for the right thing before sending the nonce back. I've stared at it for half a day and can't find what's wrong. I even tried diablo's kernel and encountered exactly the same problem. For some reason I keep thinking it's something to do with confusion about the initial offset of the nonce and what is passed to the kernel.
hero member
Activity: 772
Merit: 500
Hey Con,

I looked again through every kernel argument and compared line by line with my Python code. I found 2 small differences and 2 brackets, that are not needed (see last commit https://github.com/Diapolo/cgminer/commit/68e36c657318fbe1e7714be470cf954a1d512333), but I guess they don't fix the persisting problem with false-positive nonces (perhaps you can give it a try - I have no compiler or IDE setup to test it by myself). The argument order is exactly as DiaKGCN awaits it, so that can't be the problem either.

It could be a problem of your changes to the output code in the kernel, a problem with the base-nonces, who are passed to the kernel or something with the output-buffer in the CGMINER host code ... :-/. Where resides the output-buffer processing? As I said my kernel used ulong * natively, which I changed to uint * in one commit of my fork, I guess I need to look at it.

Edit: OMFG, I introduced a bug with one of my former commits, which changed the type of the output buffer from uint * to int * ... fixed that one! It's time for another try Con Cheesy.

Dia
Diapolo... I appreciate the effort you're putting in, and I realise you're new to this collaborative coding and source control management, but probably a good idea to see your code actually compiles before you ask someone to test it. Usually people compile and test their own code before asking someone else to test it for them.

Anyway... I fixed the !(find) in my local copy and it still produces hardware errors.

edit: It doesn't matter what vectors or worksize I try this with.

Well ... as I said, I have no IDE setup, so currently I can't compile a version for myself. If you don't have the time to fiddle around with my commits, then I really need help in setting up an IDE in Windows. Have you got this in a readme, wiki or can you give me a brief explanation in how to do this? I worked with MS VC++ Express as a hobby some time ago ...

You said local copy, is it a copy of the last version of my fork? As you've observed I am new to this kind of working, but I hope you see my progress Cheesy.

Dia
Jump to: