DiaKGCN kernel for CGMINER + Phoenix 2 (79XX / 78XX / 77XX / GCN) - 2012-05-25

[d3m0n1q_733rz]

Hey, look into removing the base and uu int by using #if GOFFSET=false as uu and base are not used when GOFFSET is.  They're just two more instructions and register entries.

I had that in mind, because if GOFFSET is active, base is simply unused and can be removed from the code . Thanks for sharing your idea!

Dia

Just remember to use #ifndef GOFFSET with the constant that sets base to uu.
I've also added GOFFSET to Phatk2!  *Happy dance*

It's tested and works.

Nice work , got your version faster on your machine?

Dia

I'm still playing around with it.  I think I'll take a nap and see how well it does.  I'm hoping that it'll work better, but I'm still moving things around.  I'm having trouble keeping Vectors8 from spilling over into memory though.
As for your GCN version, I haven't messed with it as much lately.  I'll get back to it soon.  I'm just wanting to even things out and see what works best with certain ideas.

Also, in your code here:

Code:

#else
	#ifdef GOFFSET
		u nonce = (uint)get_global_id(0);
	#else
		u nonce = ((uint)get_group_id(0) * (uint)WORKSIZExVECSIZE) + (uint)get_local_id(0) + base;

Should we add (u) to the GOFFSET nonce?

Code:

#else
	#ifdef GOFFSET
		u nonce = (uint)get_global_id(0) + (u);
	#else
		u nonce = ((uint)get_group_id(0) * (uint)WORKSIZExVECSIZE) + (uint)get_local_id(0) + base;

  Like this

[1714979855]

1714979855

[1714979855]

1714979855

[1714979855]

1714979855

[Diapolo]

No, (u) is not a variable, it's only used as a type-cast in front of variables, where u is replaced by uint, uint2, uint4 or uint8 (depends on the used vector width).

Example:

#define u uint
ulong Test_ulong = 17
uint Test_uint = (u)Test_ulong -> replaced with uint Test_uint = (uint)Test_ulong

Dia

[d3m0n1q_733rz]

Right, so not adding uint doesn't cause an offset?
I ran into something with the Phatk2 kernel which added uint as PreVal0 and PreVal4 that were required otherwise the offset caused problems (as far as I can tell anyway).  So I just wanted to be sure.

[Diapolo]

Right, so not adding uint doesn't cause an offset?
I ran into something with the Phatk2 kernel which added uint as PreVal0 and PreVal4 that were required otherwise the offset caused problems (as far as I can tell anyway).  So I just wanted to be sure.

You simply can't add (u) or uint, it has no value and should generate a syntax error in the compiler.
A variable can be of type uint (= unsigned integer -> 4 Byte storage for unsigned integer values).
Perhaps you could create an own thread where we can discuss your version as our discussion here bloats this thread .

Dia

[d3m0n1q_733rz]

Right, so not adding uint doesn't cause an offset?
I ran into something with the Phatk2 kernel which added uint as PreVal0 and PreVal4 that were required otherwise the offset caused problems (as far as I can tell anyway).  So I just wanted to be sure.

You simply can't add (u) or uint, it has no value and should generate a syntax error in the compiler.
A variable can be of type uint (= unsigned integer -> 4 Byte storage for unsigned integer values).
Perhaps you could create an own thread where we can discuss your version as our discussion here bloats this thread .

Dia

*Shrugs*  It works.  But okay, I'll toss together a forum later on.  Right now, I'm groggy and am going to take a nap.  And I used (u)(0)...which equals 0 so...zzz...I'm an idiot...zzz...

[d3m0n1q_733rz]

So, Dia, how goes the kernel programming?  I ask because I'm curious; not because I wanted to bump your thread or anything like that.   
Have you gotten the CGminer porting problem solved yet?

[Diapolo]

So, Dia, how goes the kernel programming?  I ask because I'm curious; not because I wanted to bump your thread or anything like that.   
Have you gotten the CGminer porting problem solved yet?

Interest is low, most people don't seem to care about it ... but I'm still working on DiaKGCN .

Dia

[d3m0n1q_733rz]

So, Dia, how goes the kernel programming?  I ask because I'm curious; not because I wanted to bump your thread or anything like that.   
Have you gotten the CGminer porting problem solved yet?

Interest is low, most people don't seem to care about it ... but I'm still working on DiaKGCN .

Dia

The problem with appeal might be in the $560 card required to make full use of it.  But you've included some really fine ideas and work into this kernel.  Honestly, I wish I could help more.  All I can do is suggest different methods of achieving the same thing.  Like when you created another variable to remove const u base when used with GOFFSET, I suggested, "Why not just throw an #ifndef GOFFSET statement into the void search?"
Your methods are really genius.  Heck, I have a difficult time following your code just because I have to look up some of the methods you use.  So keep it up!  I'm not getting as much help with the Phatk2 mod as I had hoped and I've only made it as far as moving around some math and changing the nonce method, but I want to see what these cards are really capable of if we push them to their limits and filter out all the crap.
BTW, I sort of borrowed your direct nonce to miner idea.  I'm actually a little surprised that I didn't think of doing the same thing.  >_<  Who really needs nonce anyway?

Something to ponder and hopefully never attempt:  What would happen if we could write the entire SHA-256 algorithm in a single and really long line of code?  @_@  Would it even be possible?! 

[e21]

So, Dia, how goes the kernel programming?  I ask because I'm curious; not because I wanted to bump your thread or anything like that.   
Have you gotten the CGminer porting problem solved yet?

Interest is low, most people don't seem to care about it ... but I'm still working on DiaKGCN .

Dia

I look forward to testing this out in a couple days once the HD 7770 is out 

[Diapolo]

So, Dia, how goes the kernel programming?  I ask because I'm curious; not because I wanted to bump your thread or anything like that.   
Have you gotten the CGminer porting problem solved yet?

Interest is low, most people don't seem to care about it ... but I'm still working on DiaKGCN .

Dia

I look forward to testing this out in a couple days once the HD 7770 is out 

7770 should not be GCN, but VLIW4 afaik .

Dia

[d3m0n1q_733rz]

So, Dia, how goes the kernel programming?  I ask because I'm curious; not because I wanted to bump your thread or anything like that.   
Have you gotten the CGminer porting problem solved yet?

Interest is low, most people don't seem to care about it ... but I'm still working on DiaKGCN .

Dia

I look forward to testing this out in a couple days once the HD 7770 is out 

7770 should not be GCN, but VLIW4 afaik .

Dia

I second that.

[Roadhog2k5]

7770 is GCN.

[Diapolo]

7770 is GCN.

If 7770 is Cape Verde, I made a mistake with my assumption ... thought 79XX = highend (GCN), 78XX =midrange (GCN) and 77XX = lowend with VLIW4.

Dia

[e21]

7770 is GCN.

If 7770 is Cape Verde, I made a mistake with my assumption ... thought 79XX = highend (GCN), 78XX =midrange (GCN) and 77XX = lowend with VLIW4.

Dia

No worries  Nope; HD 7750 and up are GCN, everything 76xx and under is 40nm VLIW5 I believe (so as to be compatible with CrossFire using Llano A series APUs)

[d3m0n1q_733rz]

7770 is GCN.

If 7770 is Cape Verde, I made a mistake with my assumption ... thought 79XX = highend (GCN), 78XX =midrange (GCN) and 77XX = lowend with VLIW4.

Dia

No worries  Nope; HD 7750 and up are GCN, everything 76xx and under is 40nm VLIW5 I believe (so as to be compatible with CrossFire using Llano A series APUs)

http://lenzfire.com/2011/12/entire-gcn-lineup-hd-7000-series-specs-and-price-revealed-60538/
Verified!
Though, you might be better off getting the 7790 for the price.

[e21]

Hi Dia!

Tried your kernel with the new 7770.. BFI_INT crashes phoenix, so I had to tweak your kernel so it wasn't enabled by default. Getting 148.5MHs at 1000MHz, using VECTORS, this seems to be about the same as I get with phatk unfortunately 

I tried to get your kernel to work with Phoenix 2.0 rc1 but when I specify kernel = diakgcn in the config file and try to run it, I get 'failed to load kernel diakgcn'. Any chance of updating your kernel to work with the new phoenix? I get 161MHs using phatk2 and phoenix 2.0rc1 when autoconfigure = enabled, which is the highest I have gotten from this card at stock speed.

Thanks!

[d3m0n1q_733rz]

Hi Dia!

Tried your kernel with the new 7770.. BFI_INT crashes phoenix, so I had to tweak your kernel so it wasn't enabled by default. Getting 148.5MHs at 1000MHz, using VECTORS, this seems to be about the same as I get with phatk unfortunately 

I tried to get your kernel to work with Phoenix 2.0 rc1 but when I specify kernel = diakgcn in the config file and try to run it, I get 'failed to load kernel diakgcn'. Any chance of updating your kernel to work with the new phoenix? I get 161MHs using phatk2 and phoenix 2.0rc1 when autoconfigure = enabled, which is the highest I have gotten from this card at stock speed.

Thanks!

If you want to disable BFI_INT, the easiest and quickest method is BFI_INT=false.  As for diakgcn, you might play around with it a bit with different vector types, GOFFSET=false and different worksizes.  The same can be said for phatk2.  But don't knock it.  It's some really wonderful work.

[Diapolo]

Hi Dia!

Tried your kernel with the new 7770.. BFI_INT crashes phoenix, so I had to tweak your kernel so it wasn't enabled by default. Getting 148.5MHs at 1000MHz, using VECTORS, this seems to be about the same as I get with phatk unfortunately 

I tried to get your kernel to work with Phoenix 2.0 rc1 but when I specify kernel = diakgcn in the config file and try to run it, I get 'failed to load kernel diakgcn'. Any chance of updating your kernel to work with the new phoenix? I get 161MHs using phatk2 and phoenix 2.0rc1 when autoconfigure = enabled, which is the highest I have gotten from this card at stock speed.

Thanks!

Can you tell me, what the device name is for your 7770? It should be displayed, if you use verbose mode with Phoenix (-v as switch in Phoenix 1 and verbose = true for Phoenix 2 in it's config). I don't supply any updates to the Phoenix 1.x version for Phoenix 2 there was a preview release. If you are interested I can put together a package with the last kernel and init for Phoenix 2 and send it to you for testing? That would be no public official release, as I have not done any documentation for the last changes to build a proper changelog.

Oh and did you use VECTORS2, as VECTORS is no supported switch for DiaKGCN!

Dia

[e21]

Hi Dia!

Tried your kernel with the new 7770.. BFI_INT crashes phoenix, so I had to tweak your kernel so it wasn't enabled by default. Getting 148.5MHs at 1000MHz, using VECTORS, this seems to be about the same as I get with phatk unfortunately 

I tried to get your kernel to work with Phoenix 2.0 rc1 but when I specify kernel = diakgcn in the config file and try to run it, I get 'failed to load kernel diakgcn'. Any chance of updating your kernel to work with the new phoenix? I get 161MHs using phatk2 and phoenix 2.0rc1 when autoconfigure = enabled, which is the highest I have gotten from this card at stock speed.

Thanks!

Can you tell me, what the device name is for your 7770? It should be displayed, if you use verbose mode with Phoenix (-v as switch in Phoenix 1 and verbose = true for Phoenix 2 in it's config). I don't supply any updates to the Phoenix 1.x version for Phoenix 2 there was a preview release. If you are interested I can put together a package with the last kernel and init for Phoenix 2 and send it to you for testing? That would be no public official release, as I have not done any documentation for the last changes to build a proper changelog.

Oh and did you use VECTORS2, as VECTORS is no supported switch for DiaKGCN!

Dia

Thanks for getting back to me Dia,

I would love it if you could send me the latest build to use with Phoenix 2!

Here is the output of Phoenix 1.7.5 using -v and VECTORS2 (VECTORS functions same as VECTORS2 so I guess it just defaults to VECTORS2) with your kernel, (7770 is Capeverde, not sure if that the name you were looking for):


C:\Bitcoin>phoenix.exe -u http://username.worker:password@api2.bitcoin.cz:
8332/ -k diakgcn -v AGGRESSION=9 WORKSIZE=128 VECTORS2 DEVICE=3 BFI_INT F3
[17/02/2012 11:22:30] using PyOpenCL version 0.92
[17/02/2012 11:22:30] checked nonces per kernel execution: 33554432
[17/02/2012 11:22:30] using VECTORS2, resulting global worksize is: 16777216
[17/02/2012 11:22:30] using local worksize of 128 (HW max. is 256)
[17/02/2012 11:22:30] BFI_INT patching not supported on Capeverde
[17/02/2012 11:22:30] OpenCL >= 1.1 supported, using global offset

[17/02/2012 11:22:31] Phoenix v1.7.5 starting...
[17/02/2012 11:22:32] Connected to server
[17/02/2012 11:22:32] Server gave new work; passing to WorkQueue
[17/02/2012 11:22:32] New block (WorkQueue)
[17/02/2012 11:22:32] Currently on block: 167217
[17/02/2012 11:22:32] Server gave new work; passing to WorkQueue
[17/02/2012 11:22:56] positive nonce (lo - 1): 2985823925
[17/02/2012 11:22:57] Result 0000000068f61663... accepted
[17/02/2012 11:23:05] Server gave new work; passing to WorkQueue
[17/02/2012 11:23:27] positive nonce (hi - 1): 3409286864
[17/02/2012 11:23:28] Result 00000000e79cdf84... accepted
[17/02/2012 11:23:33] positive nonce (lo - 1): 4227931077
[17/02/2012 11:23:34] Result 00000000d6000f90... accepted
[17/02/2012 11:23:34] Server gave new work; passing to WorkQueue
[17/02/2012 11:23:35] positive nonce (lo - 1): 171201403
[17/02/2012 11:23:36] Result 00000000b0162afe... accepted
[17/02/2012 11:24:03] Server gave new work; passing to WorkQueue
[17/02/2012 11:24:15] positive nonce (hi - 1): 1829993980
[17/02/2012 11:24:15] Result 000000006309b5fd... accepted
[17/02/2012 11:24:31] Server gave new work; passing to WorkQueue
[17/02/2012 11:25:00] Server gave new work; passing to WorkQueue
[17/02/2012 11:25:16] positive nonce (hi - 1): 1405120782
[17/02/2012 11:25:17] Result 000000005a4c086c... accepted
[150.81 Mhash/sec] [6 Accepted] [0 Rejected] [RPC (+LP)]

150MH/s is while clocked at 1240 MHZ

[-ck]

Diakgcn now working in cgminer 2.2.7 (included). Specify it with -k diakgcn. No you cannot insert a phoenix orientated diakgcn kernel into cgminer.