Gateless Gate Sharp 1.3.8: 30Mh/s (Ethash) on RX 480!

[nerdralph]

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.

[1713564617]

1713564617

[1713564617]

1713564617

[zawawa]

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

[nerdralph]

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]

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]

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

[zawawa]

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.

[zawawa]

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

[zawawa]

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.

[ioglnx]

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

[hannusolo]

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]

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]

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]

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?

[zawawa]

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.

[zawawa]

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

[zawawa]

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

[zawawa]

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

[jstefanop]

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]

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

[jstefanop]

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.