nerdralph
|
|
February 04, 2017, 02:05:21 PM |
|
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.
|
|
|
|
|
|
No Gods or Kings. Only Bitcoin
|
|
|
Advertised sites are not endorsed by the Bitcoin Forum. They may be unsafe, untrustworthy, or illegal in your jurisdiction.
|
|
|
|
zawawa (OP)
Sr. Member
Offline
Activity: 728
Merit: 304
Miner Developer
|
|
February 04, 2017, 04:16:24 PM |
|
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/2rJ2x4VBTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
|
|
|
nerdralph
|
|
February 04, 2017, 04:30:53 PM |
|
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
|
|
February 04, 2017, 04:32:29 PM |
|
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.
|
|
|
|
zawawa (OP)
Sr. Member
Offline
Activity: 728
Merit: 304
Miner Developer
|
|
February 04, 2017, 07:15:46 PM |
|
16.40 seems to be working for me. Now back to LLVM...
|
Gateless Gate Sharp, an open-source ETH/XMR miner: http://bit.ly/2rJ2x4VBTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
|
|
|
zawawa (OP)
Sr. Member
Offline
Activity: 728
Merit: 304
Miner Developer
|
|
February 04, 2017, 10:52:45 PM |
|
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/2rJ2x4VBTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
|
|
|
zawawa (OP)
Sr. Member
Offline
Activity: 728
Merit: 304
Miner Developer
|
|
February 05, 2017, 06:07:11 AM |
|
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: 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/2rJ2x4VBTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
|
|
|
zawawa (OP)
Sr. Member
Offline
Activity: 728
Merit: 304
Miner Developer
|
|
February 05, 2017, 06:22:55 PM |
|
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/2rJ2x4VBTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
|
|
|
ioglnx
Sr. Member
Offline
Activity: 574
Merit: 250
Fighting mob law and inquisition in this forum
|
|
February 05, 2017, 09:06:08 PM |
|
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
|
|
February 05, 2017, 10:07:20 PM |
|
zawawa, nerdralph have you considered working together on the miner? (on gitter or somewhere?) It might be more efficient that way BTW your work is truly remarkable. With more experience you will just get better! I can't wait to see an optimization to Ellesmere.
|
|
|
|
nerdralph
|
|
February 05, 2017, 11:31:41 PM |
|
zawawa, nerdralph have you considered working together on the miner? (on gitter or somewhere?) It might be more efficient that way BTW your work is truly remarkable. With more experience you will just get better! 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
|
|
February 05, 2017, 11:41:56 PM |
|
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 (OP)
Sr. Member
Offline
Activity: 728
Merit: 304
Miner Developer
|
|
February 06, 2017, 03:34:06 AM |
|
I haven't tried, but I suspect the ROCm binaries may work with the AMDGPU-PRO driver.
Nope, no luck: [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/2rJ2x4VBTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
|
|
|
zawawa (OP)
Sr. Member
Offline
Activity: 728
Merit: 304
Miner Developer
|
|
February 06, 2017, 04:20:24 AM |
|
Actually, the transition to ROCm seems pretty straight forward, requiring only a minimal amount of work: https://github.com/RadeonOpenCompute/ROCmThe 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/2rJ2x4VBTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
|
|
|
zawawa (OP)
Sr. Member
Offline
Activity: 728
Merit: 304
Miner Developer
|
|
February 06, 2017, 05:06:10 AM |
|
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/2rJ2x4VBTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
|
|
|
zawawa (OP)
Sr. Member
Offline
Activity: 728
Merit: 304
Miner Developer
|
|
February 06, 2017, 05:15:44 AM |
|
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/2rJ2x4VBTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
|
|
|
zawawa (OP)
Sr. Member
Offline
Activity: 728
Merit: 304
Miner Developer
|
|
February 06, 2017, 06:08:20 AM |
|
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/2rJ2x4VBTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
|
|
|
jstefanop
Legendary
Offline
Activity: 2097
Merit: 1396
|
|
February 06, 2017, 06:22:40 AM |
|
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.
|
|
|
|
zawawa (OP)
Sr. Member
Offline
Activity: 728
Merit: 304
Miner Developer
|
|
February 06, 2017, 07:01:07 AM |
|
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/2rJ2x4VBTC: 1BHwDWVerUTiKxhHPf2ubqKKiBMiKQGomZ
|
|
|
jstefanop
Legendary
Offline
Activity: 2097
Merit: 1396
|
|
February 06, 2017, 03:57:35 PM |
|
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.
|
|
|
|
|