Bitcoin Forum
November 13, 2018, 07:07:28 PM *
News: Latest Bitcoin Core release: 0.17.0 [Torrent].
 
   Home   Help Search Login Register More  
Poll
Question: Do you want to see improvements in Ethash dual-mining with GGS?
I desperately need it. - 8 (15.7%)
It would be nice. - 11 (21.6%)
It's not worth it anymore. - 32 (62.7%)
Total Voters: 51

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 ... 198 »
  Print  
Author Topic: Gateless Gate Sharp 1.3.8: 30Mh/s (Ethash) on RX 480!  (Read 204145 times)
nerdralph
Sr. Member
****
Offline Offline

Activity: 574
Merit: 251


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

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

Posts: 1542136048

View Profile Personal Message (Offline)

Ignore
1542136048
Reply with quote  #2

1542136048
Report to moderator
1542136048
Hero Member
*
Offline Offline

Posts: 1542136048

View Profile Personal Message (Offline)

Ignore
1542136048
Reply with quote  #2

1542136048
Report to moderator
1542136048
Hero Member
*
Offline Offline

Posts: 1542136048

View Profile Personal Message (Offline)

Ignore
1542136048
Reply with quote  #2

1542136048
Report to moderator
Advertised sites are not endorsed by the Bitcoin Forum. They may be unsafe, untrustworthy, or illegal in your jurisdiction. Advertise here.
1542136048
Hero Member
*
Offline Offline

Posts: 1542136048

View Profile Personal Message (Offline)

Ignore
1542136048
Reply with quote  #2

1542136048
Report to moderator
1542136048
Hero Member
*
Offline Offline

Posts: 1542136048

View Profile Personal Message (Offline)

Ignore
1542136048
Reply with quote  #2

1542136048
Report to moderator
1542136048
Hero Member
*
Offline Offline

Posts: 1542136048

View Profile Personal Message (Offline)

Ignore
1542136048
Reply with quote  #2

1542136048
Report to moderator
zawawa
Sr. Member
****
Offline Offline

Activity: 714
Merit: 300


Miner Developer


View Profile
February 04, 2017, 04:16:24 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.


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: 574
Merit: 251


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


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: 471
Merit: 260



View Profile
February 04, 2017, 04:32:29 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.

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: 714
Merit: 300


Miner Developer


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

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: 714
Merit: 300


Miner Developer


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

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: 714
Merit: 300


Miner Developer


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

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: 714
Merit: 300


Miner Developer


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

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: 574
Merit: 250

Fighting mob law and inquisition in this forum


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

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: 434
Merit: 257


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

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

Activity: 574
Merit: 251


View Profile
February 05, 2017, 11:31:41 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.

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: 574
Merit: 251


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

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: 714
Merit: 300


Miner Developer


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

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: 714
Merit: 300


Miner Developer


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

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: 714
Merit: 300


Miner Developer


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

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: 714
Merit: 300


Miner Developer


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

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: 714
Merit: 300


Miner Developer


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

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

Activity: 1234
Merit: 1119


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

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.

Project Apollo: A Pod Miner Designed for the Home https://bitcointalk.org/index.php?topic=4974036
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: 714
Merit: 300


Miner Developer


View Profile
February 06, 2017, 07:01:07 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.

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

Activity: 1234
Merit: 1119


View Profile
February 06, 2017, 03:57:35 PM
 #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...

Yea but HCC does not make much sense for mining algorithms unless you find a speed up by utilizing inter-gpu memory and data sharing between gpu cores on multi gpu systems. Could be interesting to explore for more advanced Algos like ethash and equihash.

Project Apollo: A Pod Miner Designed for the Home https://bitcointalk.org/index.php?topic=4974036
FutureBit Moonlander 2 USB Scrypt Stick Miner: https://bitcointalk.org/index.php?topic=2125643.0
LTC:LX5vpxrQE4eLRLPobKwZhw2comkKFCh3p4 - BTC:14w9Lea6kdVzspJk8TQRe7qSYu9LhzJJsh
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 ... 198 »
  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!