I must be thick i cannot get the bat file to run on dwarfpool for ether
this is my bat file
@echo off
@set GPU_FORCE_64BIT_PTR 0
@set GPU_MAX_HEAP_SIZE 100
@set GPU_USE_SYNC_OBJECTS 1
@set GPU_MAX_ALLOC_PERCENT 100
@set GPU_SINGLE_ALLOC_PERCENT 100
gatelessgate.exe --gpu-platform 1 -k ethash -o stratum+tcp://eth-eu.dwarfpool.com:8008 u mywalletaddress --xintensity 4620 --worksize 192 --gpu-threads 2 --no-extranonce
pause
Topic: Gateless Gate Sharp 1.3.8: 30Mh/s (Ethash) on RX 480! - page 169. (Read 214410 times)
Very nice zawawa, I appreciate that you are doing this for the community. What is the difference between platform 0 and 1, platform 0 seems to get more hash ~239 from 480 4 GB than platform 1 ~229.
The ROCm OpenCL driver doesn't seem to support inline GCN assembly and I have to use HCC after all.
Grrrrrrrrrr...
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.
Well, the only real reason I would be interested in HCC is that I would be able to use the inline GCN assembly.
If I can convert ROCm (hsaco) binaries into OpenCL binaries, I wouldn't even need HCC.
I am running some experiments, and the results are pretty good so far.
In fact, I am almost done replacing one of the kernels with an LLVM-generated substitute.
It would be such a dirty hack, but, if it works well, who cares?
The ROCm OpenCL driver doesn't seem to support inline GCN assembly and I have to use HCC after all.
Grrrrrrrrrr...
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.
The ROCm OpenCL driver doesn't seem to support inline GCN assembly and I have to use HCC after all.
Grrrrrrrrrr...
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...
The ROCm OpenCL driver doesn't seem to support inline GCN assembly and I have to use HCC after all.
Grrrrrrrrrr...
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.
The ROCm OpenCL driver doesn't seem to support inline GCN assembly and I have to use HCC after all.
Grrrrrrrrrr...
Grrrrrrrrrr...
Hmm... The ROCm driver still spits out my binary.
Let's see if the inline assembly works.
Let's see if the inline assembly works.
Oh wow, the miner actually runs slightly faster with ROCm.
I'm thoroughly impressed so far.
I'm thoroughly impressed so far.
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.
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.
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.
[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?
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.
(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, 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.
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.
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.
BTW your work is truly remarkable. With more experience you will just get better!
I can't wait to see an optimization to Ellesmere.
Yeah yeah mythos open amd drivers lol.
You won't have such issue with NVIDIA but maybe others...
You won't have such issue with NVIDIA but maybe others...
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.
(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.
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...
I'm done building LLVM with GCN support on Windows!
I should create a binary package for this.
This stiff ispretty messed up not straight forward at all.
I should create a binary package for this.
This stiff is
16.40 seems to be working for me. Now back to LLVM...
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.
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.
Jump to: