Bitcoin Forum
December 11, 2016, 06:07:09 AM *
News: Latest stable version of Bitcoin Core: 0.13.1  [Torrent].
 
   Home   Help Search Donate Login Register  
Pages: « 1 2 3 [4] 5 6 »  All
  Print  
Author Topic: DiaKGCN kernel for CGMINER + Phoenix 2 (79XX / 78XX / 77XX / GCN) - 2012-05-25  (Read 26109 times)
d3m0n1q_733rz
Sr. Member
****
Offline Offline

Activity: 378



View Profile WWW
February 08, 2012, 03:31:03 PM
 #61

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 Smiley. 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 Smiley, 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

Funroll_Loops, the theoretically quicker breakfast cereal!
Check out http://www.facebook.com/JupiterICT for all of your computing needs.  If you need it, we can get it.  We have solutions for your computing conundrums.  BTC accepted!  12HWUSguWXRCQKfkPeJygVR1ex5wbg3hAq
Advertised sites are not endorsed by the Bitcoin Forum. They may be unsafe, untrustworthy, or illegal in your jurisdiction. Advertise here.
1481436429
Hero Member
*
Offline Offline

Posts: 1481436429

View Profile Personal Message (Offline)

Ignore
1481436429
Reply with quote  #2

1481436429
Report to moderator
1481436429
Hero Member
*
Offline Offline

Posts: 1481436429

View Profile Personal Message (Offline)

Ignore
1481436429
Reply with quote  #2

1481436429
Report to moderator
1481436429
Hero Member
*
Offline Offline

Posts: 1481436429

View Profile Personal Message (Offline)

Ignore
1481436429
Reply with quote  #2

1481436429
Report to moderator
Diapolo
Hero Member
*****
Offline Offline

Activity: 769



View Profile WWW
February 08, 2012, 03:39:32 PM
 #62

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

Liked my former work for Bitcoin Core? Drop me a donation via:
1PwnvixzVAKnAqp8LCV8iuv7ohzX2pbn5x
bitcoin:1PwnvixzVAKnAqp8LCV8iuv7ohzX2pbn5x?label=Diapolo
d3m0n1q_733rz
Sr. Member
****
Offline Offline

Activity: 378



View Profile WWW
February 08, 2012, 03:44:59 PM
 #63

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.

Funroll_Loops, the theoretically quicker breakfast cereal!
Check out http://www.facebook.com/JupiterICT for all of your computing needs.  If you need it, we can get it.  We have solutions for your computing conundrums.  BTC accepted!  12HWUSguWXRCQKfkPeJygVR1ex5wbg3hAq
Diapolo
Hero Member
*****
Offline Offline

Activity: 769



View Profile WWW
February 08, 2012, 03:50:29 PM
 #64

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 Cheesy.

Dia

Liked my former work for Bitcoin Core? Drop me a donation via:
1PwnvixzVAKnAqp8LCV8iuv7ohzX2pbn5x
bitcoin:1PwnvixzVAKnAqp8LCV8iuv7ohzX2pbn5x?label=Diapolo
d3m0n1q_733rz
Sr. Member
****
Offline Offline

Activity: 378



View Profile WWW
February 08, 2012, 05:10:53 PM
 #65

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 Cheesy.

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...

Funroll_Loops, the theoretically quicker breakfast cereal!
Check out http://www.facebook.com/JupiterICT for all of your computing needs.  If you need it, we can get it.  We have solutions for your computing conundrums.  BTC accepted!  12HWUSguWXRCQKfkPeJygVR1ex5wbg3hAq
d3m0n1q_733rz
Sr. Member
****
Offline Offline

Activity: 378



View Profile WWW
February 12, 2012, 10:43:19 AM
 #66

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.   Grin
Have you gotten the CGminer porting problem solved yet?

Funroll_Loops, the theoretically quicker breakfast cereal!
Check out http://www.facebook.com/JupiterICT for all of your computing needs.  If you need it, we can get it.  We have solutions for your computing conundrums.  BTC accepted!  12HWUSguWXRCQKfkPeJygVR1ex5wbg3hAq
Diapolo
Hero Member
*****
Offline Offline

Activity: 769



View Profile WWW
February 12, 2012, 01:40:55 PM
 #67

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.   Grin
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 Smiley.

Dia

Liked my former work for Bitcoin Core? Drop me a donation via:
1PwnvixzVAKnAqp8LCV8iuv7ohzX2pbn5x
bitcoin:1PwnvixzVAKnAqp8LCV8iuv7ohzX2pbn5x?label=Diapolo
d3m0n1q_733rz
Sr. Member
****
Offline Offline

Activity: 378



View Profile WWW
February 12, 2012, 02:10:33 PM
 #68

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.   Grin
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 Smiley.

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?! 

Funroll_Loops, the theoretically quicker breakfast cereal!
Check out http://www.facebook.com/JupiterICT for all of your computing needs.  If you need it, we can get it.  We have solutions for your computing conundrums.  BTC accepted!  12HWUSguWXRCQKfkPeJygVR1ex5wbg3hAq
e21
Member
**
Offline Offline

Activity: 105


View Profile
February 13, 2012, 09:24:12 PM
 #69

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.   Grin
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 Smiley.

Dia

I look forward to testing this out in a couple days once the HD 7770 is out  Grin
Diapolo
Hero Member
*****
Offline Offline

Activity: 769



View Profile WWW
February 14, 2012, 05:49:37 AM
 #70

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.   Grin
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 Smiley.

Dia

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

7770 should not be GCN, but VLIW4 afaik Wink.

Dia

Liked my former work for Bitcoin Core? Drop me a donation via:
1PwnvixzVAKnAqp8LCV8iuv7ohzX2pbn5x
bitcoin:1PwnvixzVAKnAqp8LCV8iuv7ohzX2pbn5x?label=Diapolo
d3m0n1q_733rz
Sr. Member
****
Offline Offline

Activity: 378



View Profile WWW
February 14, 2012, 08:56:20 AM
 #71

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.   Grin
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 Smiley.

Dia

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

7770 should not be GCN, but VLIW4 afaik Wink.

Dia
I second that.

Funroll_Loops, the theoretically quicker breakfast cereal!
Check out http://www.facebook.com/JupiterICT for all of your computing needs.  If you need it, we can get it.  We have solutions for your computing conundrums.  BTC accepted!  12HWUSguWXRCQKfkPeJygVR1ex5wbg3hAq
Roadhog2k5
Full Member
***
Offline Offline

Activity: 131



View Profile
February 14, 2012, 05:02:21 PM
 #72

7770 is GCN.
Diapolo
Hero Member
*****
Offline Offline

Activity: 769



View Profile WWW
February 14, 2012, 05:24:19 PM
 #73

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

Liked my former work for Bitcoin Core? Drop me a donation via:
1PwnvixzVAKnAqp8LCV8iuv7ohzX2pbn5x
bitcoin:1PwnvixzVAKnAqp8LCV8iuv7ohzX2pbn5x?label=Diapolo
e21
Member
**
Offline Offline

Activity: 105


View Profile
February 14, 2012, 06:07:11 PM
 #74

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  Grin 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
Sr. Member
****
Offline Offline

Activity: 378



View Profile WWW
February 14, 2012, 07:26:59 PM
 #75

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  Grin 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.

Funroll_Loops, the theoretically quicker breakfast cereal!
Check out http://www.facebook.com/JupiterICT for all of your computing needs.  If you need it, we can get it.  We have solutions for your computing conundrums.  BTC accepted!  12HWUSguWXRCQKfkPeJygVR1ex5wbg3hAq
e21
Member
**
Offline Offline

Activity: 105


View Profile
February 17, 2012, 02:12:33 AM
 #76

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  Undecided

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
Sr. Member
****
Offline Offline

Activity: 378



View Profile WWW
February 17, 2012, 05:10:48 AM
 #77

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  Undecided

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.

Funroll_Loops, the theoretically quicker breakfast cereal!
Check out http://www.facebook.com/JupiterICT for all of your computing needs.  If you need it, we can get it.  We have solutions for your computing conundrums.  BTC accepted!  12HWUSguWXRCQKfkPeJygVR1ex5wbg3hAq
Diapolo
Hero Member
*****
Offline Offline

Activity: 769



View Profile WWW
February 17, 2012, 07:23:58 AM
 #78

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  Undecided

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

Liked my former work for Bitcoin Core? Drop me a donation via:
1PwnvixzVAKnAqp8LCV8iuv7ohzX2pbn5x
bitcoin:1PwnvixzVAKnAqp8LCV8iuv7ohzX2pbn5x?label=Diapolo
e21
Member
**
Offline Offline

Activity: 105


View Profile
February 17, 2012, 07:30:05 PM
 #79

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  Undecided

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
Moderator
Legendary
*
Offline Offline

Activity: 2002


Ruu \o/


View Profile WWW
February 20, 2012, 05:25:22 AM
 #80

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.

Primary developer/maintainer for cgminer and ckpool/ckproxy.
Pooled mine at kano.is, solo mine at solo.ckpool.org
-ck
Pages: « 1 2 3 [4] 5 6 »  All
  Print  
 
Jump to:  

Sponsored by , a Bitcoin-accepting VPN.
Powered by MySQL Powered by PHP Powered by SMF 1.1.19 | SMF © 2006-2009, Simple Machines Valid XHTML 1.0! Valid CSS!