Bitcoin Forum
December 11, 2017, 10:08:49 PM *
News: Latest stable version of Bitcoin Core: 0.15.1  [Torrent].
 
   Home   Help Search Donate Login Register  
Pages: « 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 [28] 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 ... 85 »
  Print  
Author Topic: Gateless Gate Sharp 1.1.3: zawawa's open-source dual ETH/XMR/PASC/LBC miner  (Read 162083 times)
zawawa
Sr. Member
****
Offline Offline

Activity: 420


Miner Developer


View Profile
February 03, 2017, 06:52:43 PM
 #541

I've been crunching more numbers on the GDS performance, and I think for GPUs with more than 4 memory controllers (like Tahiti & Hawaii) that the best performance may be from splitting the row counters between L2 and GDS. With all the row counters in GDS, the GDS bandwidth should get maxed out before the core/L2 bandwidth limit is reached.


That recommendation is consistent with my somewhat informal observations. Thank you so much!
Just to let you know, I was able to link the object file of the kernel with libclc. The only remaining step is to convert one big LLVM Bit Code (.bc) file into an OpenCL binary. We will see.

Gateless Gate Sharp, an open-source ETH/XMR miner: http://bit.ly/2rJ2x4V
BTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
1513030129
Hero Member
*
Offline Offline

Posts: 1513030129

View Profile Personal Message (Offline)

Ignore
1513030129
Reply with quote  #2

1513030129
Report to moderator
1513030129
Hero Member
*
Offline Offline

Posts: 1513030129

View Profile Personal Message (Offline)

Ignore
1513030129
Reply with quote  #2

1513030129
Report to moderator
1513030129
Hero Member
*
Offline Offline

Posts: 1513030129

View Profile Personal Message (Offline)

Ignore
1513030129
Reply with quote  #2

1513030129
Report to moderator
"With e-currency based on cryptographic proof, without the need to trust a third party middleman, money can be secure and transactions effortless." -- Satoshi
Advertised sites are not endorsed by the Bitcoin Forum. They may be unsafe, untrustworthy, or illegal in your jurisdiction. Advertise here.
1513030129
Hero Member
*
Offline Offline

Posts: 1513030129

View Profile Personal Message (Offline)

Ignore
1513030129
Reply with quote  #2

1513030129
Report to moderator
1513030129
Hero Member
*
Offline Offline

Posts: 1513030129

View Profile Personal Message (Offline)

Ignore
1513030129
Reply with quote  #2

1513030129
Report to moderator
nerdralph
Sr. Member
****
Offline Offline

Activity: 406


View Profile
February 04, 2017, 02:05:21 PM
 #542

Just to let you know, I was able to link the object file of the kernel with libclc. The only remaining step is to convert one big LLVM Bit Code (.bc) file into an OpenCL binary. We will see.

I'd be interested in seeing your complile/link options.  I ran into issues with the linking where it wouldn't find functions like get_global_id(), or when I tried building a different way I was getting duplicate definitions.
It looked like here had been some changes/renaming related to the amdgpu back-end, and I was having a hard time finding documentation on the -x cl and -target options.  I was planning to look through the llvm source but haven't got around to it yet.
zawawa
Sr. Member
****
Offline Offline

Activity: 420


Miner Developer


View Profile
February 04, 2017, 04:16:24 PM
 #543

Just to let you know, I was able to link the object file of the kernel with libclc. The only remaining step is to convert one big LLVM Bit Code (.bc) file into an OpenCL binary. We will see.

I'd be interested in seeing your complile/link options.  I ran into issues with the linking where it wouldn't find functions like get_global_id(), or when I tried building a different way I was getting duplicate definitions.
It looked like here had been some changes/renaming related to the amdgpu back-end, and I was having a hard time finding documentation on the -x cl and -target options.  I was planning to look through the llvm source but haven't got around to it yet.


I will get to that as soon as I am done fixing my Linux installation.
AMD Linux drivers are truly an abomination...

Gateless Gate Sharp, an open-source ETH/XMR miner: http://bit.ly/2rJ2x4V
BTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
nerdralph
Sr. Member
****
Offline Offline

Activity: 406


View Profile
February 04, 2017, 04:30:53 PM
 #544


AMD Linux drivers are truly an abomination...

I had problems with AMDGPU-PRO 16.40, but fglrx has worked great for me on Ubuntu 14.04.  "apt-get install fglrx" just works.   The only other thing you might need to do (if you want clock control) is "aticonfig --initial --adapter=all" and disable gpu-manager.
laik2
Sr. Member
****
Offline Offline

Activity: 392


View Profile
February 04, 2017, 04:32:29 PM
 #545


AMD Linux drivers are truly an abomination...

I had problems with AMDGPU-PRO 16.40, but fglrx has worked great for me on Ubuntu 14.04.  "apt-get install fglrx" just works.   The only other thing you might need to do (if you want clock control) is "aticonfig --initial --adapter=all" and disable gpu-manager.

Use 16.50 with kernel 4.4.(Ubuntu 16.04)
16.60 has a lot of issues, I'm still struggling to get my cards stable. If you need access to linux with already installed "stable" drivers just msg me.
14.04 and fglrx only support pre RX cards.
EDIT: Although I have 14.04 with latest fglrx and is extremely stable on my R9 390s.

ZEC: t1KbbHtXqzSS6qHBaPZDKyWnzxhRjr9oCtW
zawawa
Sr. Member
****
Offline Offline

Activity: 420


Miner Developer


View Profile
February 04, 2017, 07:15:46 PM
 #546

16.40 seems to be working for me. Now back to LLVM...

Gateless Gate Sharp, an open-source ETH/XMR miner: http://bit.ly/2rJ2x4V
BTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
zawawa
Sr. Member
****
Offline Offline

Activity: 420


Miner Developer


View Profile
February 04, 2017, 10:52:45 PM
 #547

I'm done building LLVM with GCN support on Windows!
I should create a binary package for this.
This stiff is pretty messed up not straight forward at all.

Gateless Gate Sharp, an open-source ETH/XMR miner: http://bit.ly/2rJ2x4V
BTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
zawawa
Sr. Member
****
Offline Offline

Activity: 420


Miner Developer


View Profile
February 05, 2017, 06:07:11 AM
 #548

Just to let you know, I was able to link the object file of the kernel with libclc. The only remaining step is to convert one big LLVM Bit Code (.bc) file into an OpenCL binary. We will see.

I'd be interested in seeing your complile/link options.  I ran into issues with the linking where it wouldn't find functions like get_global_id(), or when I tried building a different way I was getting duplicate definitions.
It looked like here had been some changes/renaming related to the amdgpu back-end, and I was having a hard time finding documentation on the -x cl and -target options.  I was planning to look through the llvm source but haven't got around to it yet.


Here they are:
Code:
clang -target amdgcn-amdhsa--polaris11 -IZ:\GitHub\llvm\tools\libclc\generic\include -Xclang -mlink-bitcode-file -Xclang Z:\GitHub\llvm\tools\libclc\built_libs\amdgcn--amdhsa.bc -include clc/clc.h -Dcl_clang_storage_class_specifiers kernel\equihash.cl -D__OPENCL_VERSION__=120 -DWORKSIZE=256 -Wno-typedef-redefinition -o equihashEllesmeregw256l4.asm -S

The resulting file seems to be a valid HSA IL code.
I was also able to generate a binary without "-o equihashEllesmeregw256l4.asm -S", but it does not seem compatible with AMD drivers.
I am getting so close, though...

Gateless Gate Sharp, an open-source ETH/XMR miner: http://bit.ly/2rJ2x4V
BTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
zawawa
Sr. Member
****
Offline Offline

Activity: 420


Miner Developer


View Profile
February 05, 2017, 06:22:55 PM
 #549

Here are my findings regarding GCN support for LLVM.

(1) "-target amdgcn" is broken. LLVM's integrated linker is not capable of generating binaries for AMD's proprietary OpenCL drivers.
(2) "-target amdgcn--amdhsa" produces ROCm binaries.
(3) CLRX could be used to process the assembly output of LLVM's GCN front-end.

This, of course, means more work for me.
I may be able to get some extra help, though. We will see.

Gateless Gate Sharp, an open-source ETH/XMR miner: http://bit.ly/2rJ2x4V
BTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
ioglnx
Sr. Member
****
Offline Offline

Activity: 434

Fighting mob law and inquisition in this forum


View Profile
February 05, 2017, 09:06:08 PM
 #550

Yeah yeah mythos open amd drivers lol.
You won't have such issue with NVIDIA but maybe others...

GTX 1080Ti rocks da house... seriously... this card is a beast³
Owning by now 18x GTX1080Ti :-D @serious love of efficiency
hannusolo
Sr. Member
****
Offline Offline

Activity: 336


Etheal


View Profile
February 05, 2017, 10:07:20 PM
 #551

zawawa, nerdralph have you considered working together on the miner? (on gitter or somewhere?) It might be more efficient that way Cheesy

BTW your work is truly remarkable. With more experience you will just get better! Smiley
I can't wait to see an optimization to Ellesmere.

        ETHEAL           ETHEAL
      ETHEALETHE       ETHEALETHE
    ETHEALETHEALET   ETHEALETHEALET
  ETHEALETHEALETHEA ETHEALETHEALETHEA
 ETHEALETHEALETHEALETHEALETHEALETHEALE
ETHEALETHEALETHEALETHEALETHEALETHEALETH
ETHEALETHEALETHEALETHEALETHEALETHEALETH
ETHEALETHEALETHEALETHEALETHEALETHEALETH
 ETHEALETHEALETHEALETHEALETHEALETHEALE
  ETHEALETHEALETHEALETHEALETHEALETHEA
   ETHEALETHEALETHEALETHEALETHEALETH
    ETHEALETHEALETHEALETHEALETHEALE
      ETHEALETHEALETHEALETHEALETH
        ETHEALETHEALETHEALETHEA
          ETHEALETHEALETHEALE
            ETHEALETHEALETH
              ETHEALETHEA
                ETHEALE
                  ETH
                   E


🅽🅴🆆 Operating System of $7600B+ 🏥 Healthcare industry
🕸 Website  📝 Whitepaper  Video  📢 ANN  💰 Bounty  💬 Telegram
Improving the lives of 1 billion people by making Healthcare transparent.


nerdralph
Sr. Member
****
Offline Offline

Activity: 406


View Profile
February 05, 2017, 11:31:41 PM
 #552

zawawa, nerdralph have you considered working together on the miner? (on gitter or somewhere?) It might be more efficient that way Cheesy

BTW your work is truly remarkable. With more experience you will just get better! Smiley
I can't wait to see an optimization to Ellesmere.

I don't do gitter.  Github's issue tracker and email work fine (that's what JW and I used with Genoi's ethminer).
zawawa seems to have more time to dedicate to development than I do, so if we were trying to work on it together I'd probably slow him down.
nerdralph
Sr. Member
****
Offline Offline

Activity: 406


View Profile
February 05, 2017, 11:41:56 PM
 #553

Here are my findings regarding GCN support for LLVM.

(1) "-target amdgcn" is broken. LLVM's integrated linker is not capable of generating binaries for AMD's proprietary OpenCL drivers.
(2) "-target amdgcn--amdhsa" produces ROCm binaries.
(3) CLRX could be used to process the assembly output of LLVM's GCN front-end.

This, of course, means more work for me.
I may be able to get some extra help, though. We will see.

Interesting.  -target amdgcn worked OK for me to create asm files (.s), while getting a .o to compile from the .s seemed to require -target amdgcn--amdhsa.
I haven't tried, but I suspect the ROCm binaries may work with the AMDGPU-PRO driver.
zawawa
Sr. Member
****
Offline Offline

Activity: 420


Miner Developer


View Profile
February 06, 2017, 03:34:06 AM
 #554

I haven't tried, but I suspect the ROCm binaries may work with the AMDGPU-PRO driver.

Nope, no luck:

Code:
[19:24:15] Error -11: Building Program (clBuildProgram)
[19:24:15] Error: The binary is incorrect or incomplete. Finalization to ISA cou
ldn't be performed.

You know what, I'm now thinking about implementing ROCm support for Gateless Gate.
I mean, if engineers at AMD thinks it is the future, why should I go against the flow?

Gateless Gate Sharp, an open-source ETH/XMR miner: http://bit.ly/2rJ2x4V
BTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
zawawa
Sr. Member
****
Offline Offline

Activity: 420


Miner Developer


View Profile
February 06, 2017, 04:20:24 AM
 #555

Actually, the transition to ROCm seems pretty straight forward, requiring only a minimal amount of work:

https://github.com/RadeonOpenCompute/ROCm

The dream of an open-source GCN compiler may come true at last...
I have been waiting for it forever, you know.

Gateless Gate Sharp, an open-source ETH/XMR miner: http://bit.ly/2rJ2x4V
BTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
zawawa
Sr. Member
****
Offline Offline

Activity: 420


Miner Developer


View Profile
February 06, 2017, 05:06:10 AM
 #556

Oh wow, the miner actually runs slightly faster with ROCm.
I'm thoroughly impressed so far.

Gateless Gate Sharp, an open-source ETH/XMR miner: http://bit.ly/2rJ2x4V
BTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
zawawa
Sr. Member
****
Offline Offline

Activity: 420


Miner Developer


View Profile
February 06, 2017, 05:15:44 AM
 #557

Hmm... The ROCm driver still spits out my binary.
Let's see if the inline assembly works.

Gateless Gate Sharp, an open-source ETH/XMR miner: http://bit.ly/2rJ2x4V
BTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
zawawa
Sr. Member
****
Offline Offline

Activity: 420


Miner Developer


View Profile
February 06, 2017, 06:08:20 AM
 #558

The ROCm OpenCL  driver doesn't seem to support inline GCN assembly and I have to use HCC after all.
Grrrrrrrrrr...

Gateless Gate Sharp, an open-source ETH/XMR miner: http://bit.ly/2rJ2x4V
BTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
jstefanop
Hero Member
*****
Offline Offline

Activity: 838


View Profile
February 06, 2017, 06:22:40 AM
 #559

The ROCm OpenCL  driver doesn't seem to support inline GCN assembly and I have to use HCC after all.
Grrrrrrrrrr...

Yea their OpenCL implementation is pretty barebones...the current release was the first and it was just a developer beta. Will probably be a while before OpenCL matures on that platform.

www.AriseChickun.com 0% fee segwit signaling Litecoin Pool!
FutureBit Moonlander 2 USB Scrypt Stick Miner: https://bitcointalk.org/index.php?topic=2125643.0
LTC:LX5vpxrQE4eLRLPobKwZhw2comkKFCh3p4 - BTC:14w9Lea6kdVzspJk8TQRe7qSYu9LhzJJsh
zawawa
Sr. Member
****
Offline Offline

Activity: 420


Miner Developer


View Profile
February 06, 2017, 07:01:07 AM
 #560

The ROCm OpenCL  driver doesn't seem to support inline GCN assembly and I have to use HCC after all.
Grrrrrrrrrr...

Yea their OpenCL implementation is pretty barebones...the current release was the first and it was just a developer beta. Will probably be a while before OpenCL matures on that platform.

It makes a perfect sense that AMD's commitment to OpenCL seems half-hearted if AMD wants to promote HCC over OpenCL in a long run. I'm not entirely sure if I want to rewrite the kernel, but their ROCm stuff looks pretty neat...

Gateless Gate Sharp, an open-source ETH/XMR miner: http://bit.ly/2rJ2x4V
BTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
Pages: « 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 [28] 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 ... 85 »
  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!