CUDA miner

1356769

Comments

  • GenoilGenoil 0xeb9310b185455f863f526dab3d245809f6854b4dMember Posts: 769 ✭✭✭
    edited June 2015
    latest source on my github should compile on Linux without having to tamper with anything. Except I had to make a tiny change to some header file of CUDA 7, because the compiler complained about a "comment in a comment".
  • o0ragman0oo0ragman0o Member, Moderator Posts: 1,291 mod
    you really got to watch that stack on recurrent comments ;)
  • BripsBrips Member Posts: 52
    Thanks Genoil !
    I successfully compile on ubuntu :)
    My test with the GTX 970 :
    - Default ethminer opencl : 17,6MH/s
    - Genoil ethminer opencl : 18,0MH/s
    - Genoil ethminer cuda : 18,4MH/s

    it isn't a big gap, but it's better ! thanks again ;-)
  • GenoilGenoil 0xeb9310b185455f863f526dab3d245809f6854b4dMember Posts: 769 ✭✭✭
    @Brips you will get a bit more if you add the flags --gpu-workgroup-size 128 --gpu-mining-buffers 4 --gpu-batch-size 19
  • BripsBrips Member Posts: 52
    Genoil said:

    @Brips you will get a bit more if you add the flags --gpu-workgroup-size 128 --gpu-mining-buffers 4 --gpu-batch-size 19

    you add these parameters on the command line ? like this :
    ./ethminer -U -M --gpu-workgroup-size 128 --gpu-mining-buffers 4 --gpu-batch-size 19 ?
    i test it but i have a little less 18,3MH/s with cuda
  • GenoilGenoil 0xeb9310b185455f863f526dab3d245809f6854b4dMember Posts: 769 ✭✭✭
    Yes then that's what it is. Had hoped for a bit more on 9xx but the only way to possibly get it faster is to completely turn the implementation of the algorithm upside down, and then still it's a bet...
  • o0ragman0oo0ragman0o Member, Moderator Posts: 1,291 mod
    My Linux build's crashing...

    $ cmake -DETHASHCL=1 -DGUI=0 -DBUNDLE=miner

    ...
    [100%] Building CXX object ethminer/CMakeFiles/ethminer.dir/main.cpp.o
    In file included from /home/o0ragman0o/genoil/cpp-ethereum/ethminer/main.cpp:31:0:
    /home/o0ragman0o/genoil/cpp-ethereum/ethminer/MinerAux.h: In member function ‘void MinerCLI::execute()’:
    /home/o0ragman0o/genoil/cpp-ethereum/ethminer/MinerAux.h:304:4: error: ‘setKernelParameters’ is not a member of ‘dev::eth::Ethash::CUDAMiner {aka dev::eth::Ethash::CPUMiner}’
    ProofOfWork::CUDAMiner::setKernelParameters(gpuMiningBuffers, gpuBatchSize, gpuWorkgroupSize);
    ^
    make[2]: *** [ethminer/CMakeFiles/ethminer.dir/main.cpp.o] Error 1
    make[1]: *** [ethminer/CMakeFiles/ethminer.dir/all] Error 2
    make: *** [all] Error 2
  • GenoilGenoil 0xeb9310b185455f863f526dab3d245809f6854b4dMember Posts: 769 ✭✭✭
    edited June 2015
    @o0ragman0o well that's not good so I should fix that, but it'll go away if you specify -DETHASHCU=1 at cmake.
  • GenoilGenoil 0xeb9310b185455f863f526dab3d245809f6854b4dMember Posts: 769 ✭✭✭
    edited June 2015
    After some extensive testing, I came to the conclusion that my CUDA miner is seriously broken. I had been developing with a very specific Proof of Work and difficulty, and ironically, just for that exact problem it produces the same results as the opencl kernel. So I though I was on the right track. For other Proof of Work though, it occasionally prodcues the same results, but most of the times completely different.

    I'm afraid it's back to the drawing table...

  • GenoilGenoil 0xeb9310b185455f863f526dab3d245809f6854b4dMember Posts: 769 ✭✭✭
    latest revision should be much better..still not 100% identical behavior though..
  • o0ragman0oo0ragman0o Member, Moderator Posts: 1,291 mod
    Still not having any luck with your ethminer... :(

    [ 36%] Building CXX object libdevcore/CMakeFiles/devcore.dir/Guards.cpp.o
    /home/darryl/genoil/cpp-ethereum/libethash-cu/ethash_cu_miner_kernel.cu(196): error: identifier "__ldg" is undefined

    1 error detected in the compilation of "/tmp/tmpxft_00003fc5_00000000-23_ethash_cu_miner_kernel.compute_20.cpp1.ii".
    CMake Error at ethash-cu_generated_ethash_cu_miner_kernel.cu.o.cmake:264 (message):
    Error generating file
    /home/darryl/genoil/cpp-ethereum/build/libethash-cu/CMakeFiles/ethash-cu.dir//./ethash-cu_generated_ethash_cu_miner_kernel.cu.o


    make[2]: *** [libethash-cu/CMakeFiles/ethash-cu.dir/./ethash-cu_generated_ethash_cu_miner_kernel.cu.o] Error 1
    make[1]: *** [libethash-cu/CMakeFiles/ethash-cu.dir/all] Error 2
    make[1]: *** Waiting for unfinished jobs....
  • GenoilGenoil 0xeb9310b185455f863f526dab3d245809f6854b4dMember Posts: 769 ✭✭✭
    @o0ragman0o oh that __ldg is from yesterday, doesn't compile on 2.x. So if you build with -DCOMPUTE=50 (For GTX750) and recompile, you should be fine.
  • o0ragman0oo0ragman0o Member, Moderator Posts: 1,291 mod
    edited July 2015
    ...And after removing the offending "/*" from /usr/local/cuda/include/CL/cl_gl_ext.h:44:4: error: "/*" within comment [-Werror=comment] it compiled and works!

    It's not benchmarking as good as it was on Win 7. -U Only getting 6.6 on a benchmark whereas -G still hitting 6.9. But on actual mining -U is giving me a solid 5Mhs whereas -G only gets 2.4MHs.

    Do you have any idea why mining hashrate is slower than benchmarks?
  • GenoilGenoil 0xeb9310b185455f863f526dab3d245809f6854b4dMember Posts: 769 ✭✭✭
    edited July 2015
    @o0ragman0o it should be able to do 7 (GTX750 non-Ti) as before, but I've made a change in how to optimize thread occupancy. It would be great if you could fing the best values for __launch_bounds__ on line 284 of the kernel. You may try (64, 7), (64, 8), (128, 6), (128, 7). When launching, you can never specify a value for --gpu-workgroup-size larger than the first argument. It defaults to 64.

    The mining hash rate is lower because of some additional code overhead, but it seems quite high on your setup, esp. on OpenCL.
  • o0ragman0oo0ragman0o Member, Moderator Posts: 1,291 mod
    @Genoil, Seems there's a memory leak or something in your CUDA code. After running overnight the workpackages throw:
    Error CUDA mining: an illegal memory access was encountered
  • GenoilGenoil 0xeb9310b185455f863f526dab3d245809f6854b4dMember Posts: 769 ✭✭✭
    edited July 2015
    @o0ragman0o the current version is untested on anything but my own GTX780. I tried to speed up the kernel by completely getting rid of shared memory and replace it with warp shuffles. Took me ages to get it working, only to find out it wasn't any faster than the shared memory approach. And now it doesn't even run on GTX750 :). I'll try to post a version today that allows you to choose between shuffle and shared. Was planning on that anyway.

    scrap that, you were right, it also failed on my GTX780. most recent version should run. btw it won't accept --gpu-workgroup-size 128 for the moment, just 64.
    Post edited by Genoil on
  • GenoilGenoil 0xeb9310b185455f863f526dab3d245809f6854b4dMember Posts: 769 ✭✭✭
    edited July 2015
    made a 14% speed increase on the GTX780 today. now 20.6MH/s @ stock, 24.4MH/s overclocked.
  • o0ragman0oo0ragman0o Member, Moderator Posts: 1,291 mod
    Y'awesome Genoil!
    Up from 6.9
    Benchmarking on platform: { "platform": "CUDA 7.0", "device": "GeForce GTX 750", "version": "Compute 5.0" }
    Preparing DAG...
    Warming up...
    ℹ 11:43:37|cudaminer0 workLoop 0 #00000000… #00000000…
    ℹ 11:43:37|cudaminer0 Initialising miner...
    Using device: GeForce GTX 750(5.0)
    Trial 1... 8388608
    Trial 2... 8388608
    Trial 3... 8475989
    Trial 4... 8388608
    Trial 5... 8388608
    min/mean/max: 8388608/8406084/8475989 H/s
  • GenoilGenoil 0xeb9310b185455f863f526dab3d245809f6854b4dMember Posts: 769 ✭✭✭
    This makes the 750ti (should do 9.1 now) top your charts except the titan and fury, right?
  • o0ragman0oo0ragman0o Member, Moderator Posts: 1,291 mod
    I'm waiting on a 750 ti to be shipped. Maybe @ConradJohnson can give us benchtest on his.
  • eth001eth001 Member Posts: 14
    @Genoil could you help me with the following problem thanks
    [email protected]:~/cpp-ethereum/build$ ethminer/ethminer -U -M --gpu-workgroup-size 128 --gpu-batch-size 20 -t 4
    Benchmarking on platform: { "platform": "CUDA 7.0", "device": "GRID K520", "version": "Compute 3.0" }
    Preparing DAG...
    ℹ 18:35:45|cudaminer0 workLoop 0 #00000000… #00000000…
    ℹ 18:35:45|cudaminer1 workLoop 0 #00000000… #00000000…
    ℹ 18:35:45|cudaminer0 Initialising miner...
    ℹ 18:35:45|cudaminer1 Initialising miner...
    ℹ 18:35:45|cudaminer2 workLoop 0 #00000000… #00000000…
    ℹ 18:35:45|cudaminer2 Initialising miner...
    ℹ 18:35:45|cudaminer3 workLoop 0 #00000000… #00000000…
    ℹ 18:35:45|cudaminer3 Initialising miner...
    Warming up...
    Using device: GRID K520(3.0)
    Using device: GRID K520(3.0)
    Using device: GRID K520(3.0)
    Using device: GRID K520(3.0)
    ✘ 18:35:46|cudaminer2 Error CUDA mining: invalid device symbol
    ✘ 18:35:46|cudaminer1 Error CUDA mining: invalid device symbol
    ✘ 18:35:46|cudaminer0 Error CUDA mining: invalid device symbol
    ✘ 18:35:46|cudaminer3 Error CUDA mining: invalid device symbol
    Trial 1... 0
    Trial 2... 0
    Trial 3... 0
    Trial 4... 0
    Trial 5... 0
    min/mean/max: 0/0/0 H/s
    inner mean: 0 H/s
    Phoning home to find world ranking...
    Error phoning home. ET is sad.
  • PascalVerstPascalVerst Member Posts: 6
    edited July 2015
    I haven't been able to mine for 10 ten days so I decided to rebuild, can you provide some info on building your cuda miner?

    I've downloaded ethereum sln + your sln, but I am build errors.
    The right ethash doesn't seem to be included:
    Error 15 error C2039: 'startCUDA' : is not a member of 'dev::eth::GenericFarm' C:\cpp-ethereum\ethminer\MinerAux.h 404 1 eth
  • eth001eth001 Member Posts: 14
    I think the difficult is too high now , cannot get profile
  • CryptoFutureCryptoFuture Member Posts: 44
    I'm getting invalid device symbol also
  • mooneyjmooneyj Member Posts: 7
    Any tips for getting an nvidia 660ti mining on ubuntu15.04? The gpu's max memory allocation is only 500mb. Can i fit the DAG in there in chunks? I'm using the ubuntu ppa version of cpp-ethereum. Perhaps i need to get on the develop branch & build myself?

    Heres the error: 0 H/s :(


    [email protected]:~$ ethminer -G -F http://127.0.0.1:8545
    [OPENCL]:Found suitable OpenCL device [GeForce GTX 660 Ti] with 2147155968 bytes of GPU memory
    miner 23:35:04|ethminer Getting work package...

    ℹ 23:35:09|ethminer Loading full DAG of seedhash: #290decd9…
    miner 23:35:10|ethminer Getting work package...
    miner 23:35:10|ethminer Grabbing DAG for #5a46dd85…
    ℹ 23:35:14|ethminer Full DAG loaded
    miner 23:35:21|ethminer Got work package:
    miner 23:35:21|ethminer Header-hash: c21c407788fcbfac56d7157526e7a6455ec69e265ec74fe4d2d60f0803c1ac2c
    miner 23:35:21|ethminer Seedhash: 5a46dd85298b0beff65ccf3006c4b5d2f7fa6ca8a5885a7f2132714fe7a48602
    miner 23:35:21|ethminer Target: 000000002bf83e1cdec8746c78083e1637e163db9a5e141cd3a72a4c5c6c9216
    miner 23:35:21|ethminer Mining on PoWhash #c21c4077… : 0 H/s = 0 hashes / 0.5 s
    miner 23:35:21|ethminer Got work package:
    miner 23:35:21|ethminer Header-hash: 36b08af8fcfc0c408110b59cd1102cfa781c5b2c26a80e1ee382f0f869e0acd3
    miner 23:35:21|ethminer Seedhash: 5a46dd85298b0beff65ccf3006c4b5d2f7fa6ca8a5885a7f2132714fe7a48602
    miner 23:35:21|ethminer Target: 000000002bf2bfc4ea19045deb590d852d521f098972df68b7edcaf4b110ccb4
    miner 23:35:22|ethminer Mining on PoWhash #36b08af8… : 0 H/s = 0 hashes / 0.5 s
    miner 23:35:22|ethminer Mining on PoWhash #36b08af8… : 0 H/s = 0 hashes / 1 s
    miner 23:35:23|ethminer Mining on PoWhash #36b08af8… : 0 H/s = 0 hashes / 1.501 s
    miner 23:35:23|ethminer Mining on PoWhash #36b08af8… : 0 H/s = 0 hashes / 2.002 s
    miner 23:35:24|ethminer Mining on PoWhash #36b08af8… : 0 H/s = 0 hashes / 2.503 s
    miner 23:35:24|ethminer Mining on PoWhash #36b08af8… : 0 H/s = 0 hashes / 3.004 s
    miner 23:35:25|ethminer Mining on PoWhash #36b08af8… : 0 H/s = 0 hashes / 3.504 s


    Output from geth gpuinfo:

    =============================================
    ============ OpenCL Device Info =============
    =============================================
    Platform id 0
    Platform Name NVIDIA CUDA
    Platform Vendor NVIDIA Corporation
    Platform Version OpenCL 1.1 CUDA 7.0.35
    Platform Extensions cl_khr_byte_addressable_store cl_khr_icd cl_khr_gl_sharing cl_nv_compiler_options cl_nv_device_attribute_query cl_nv_pragma_unroll cl_nv_copy_opts
    Platform Profile FULL_PROFILE

    Device OpenCL id 0
    Device id for mining 0
    Device Name GeForce GTX 660 Ti
    Vendor NVIDIA Corporation
    Version OpenCL 1.1 CUDA
    Driver version 346.59
    Address bits 32
    Max clock freq 980
    Global mem size 2147155968
    Max constant buffer size 65536
    Max mem alloc size 536788992
    Max compute units 7
    Max work group size 1024
    Max work item sizes [1024 1024 64]
    =============================================
    Found 1 devices. Benchmark first GPU: geth gpubench 0
    Mine using all GPUs: geth --minegpu 0
    You may also need to enable chunking: --gpuchunks
  • o0ragman0oo0ragman0o Member, Moderator Posts: 1,291 mod
    Genoil said:

    This makes the 750ti (should do 9.1 now) top your charts except the titan and fury, right?

    Using device: GeForce GTX 750 Ti(5.0)
    Trial 1... 11097429
    Trial 2... 11097429
    Trial 3... 11097429
    Trial 4... 11010048
    Trial 5... 11097429
    min/mean/max: 11010048/11079952/11097429 H/s
    inner mean: 11097429 H/s

    And I picked it up for less than 100 bucks! That means it's as price and power efficient as the Titans. Actual mining is at 6.1Mhs

    I tried OpenCL benchtest and mining also but got...
    Using platform: NVIDIA CUDA
    Using device: GeForce GTX 750 Ti(OpenCL 1.1 CUDA)
    :396:10: error: redefinition of 'compute_hash'
    hash32_t compute_hash(
    ^
    :354:10: note: previous definition is here
    hash32_t compute_hash(
    ^
    ✘ 09:16:22|gpuminer0 Error GPU mining: clEnqueueWriteBuffer ( -38 )

  • mikkaaamikkaaa Member Posts: 34
    Will this Genoils cuda miner be compatible with the Frontier release? Just trying to map out the mining options...
  • o0ragman0oo0ragman0o Member, Moderator Posts: 1,291 mod
    Ok, using official ethminer:

    [OPENCL]:Found suitable OpenCL device [GeForce GTX 750 Ti] with 2146762752 bytes of GPU memory
    Benchmarking on platform: { "platform": "NVIDIA CUDA", "device": "GeForce GTX 750 Ti", "version": "OpenCL 1.1 CUDA" }
    Preparing DAG...
    Warming up...
    ℹ 09:39:14|gpuminer0 workLoop 0 #00000000… #00000000…
    ℹ 09:39:14|gpuminer0 Initialising miner...
    [OPENCL]:Using platform: NVIDIA CUDA
    [OPENCL]:Using device: GeForce GTX 750 Ti(OpenCL 1.1 CUDA)
    [OPENCL]:Printing program log
    [OPENCL]:

    [OPENCL]:Created one big buffer for the DAG
    [OPENCL]:Loading single big chunk kernels
    [OPENCL]:Creating buffer for header.
    [OPENCL]:Mapping one big chunk.
    [OPENCL]:Creating mining buffer 0
    [OPENCL]:Creating mining buffer 1
    Trial 1... 9087658
    Trial 2... 9175040
    Trial 3... 9087658
    Trial 4... 9175040
    Trial 5... 9087658
    min/mean/max: 9087658/9122610/9175040 H/s
    inner mean: 3058346 H/s
    but actual mining is pretty disappointing:
      ℹ  09:32:44|ethminer  Mining on PoWhash #690a1a93… : 2441468 H/s = 18350080 hashes / 7.516 s
  • rugbyowlrugbyowl Member Posts: 8
    This may be a stupid question, but what do I do after making cpp-ethereum? There isn't a folder labeled ethminer on github anymore, so I'm confused what command I execute to run my miner. I'm running Arch linux if it matters.

    Thanks
  • mooneyjmooneyj Member Posts: 7
    Is the cudaminer branch compatible with frontier? I'm running a g2.x8 instance @ 35.7MH/S on frontier release cpp ethminer. 46MH/s sounds nice :)
Sign In or Register to comment.