diff --git a/.github/workflows/build_binaries.yml b/.github/workflows/build_binaries.yml index 2380da3..1e46628 100644 --- a/.github/workflows/build_binaries.yml +++ b/.github/workflows/build_binaries.yml @@ -6,6 +6,8 @@ name: Build Matrix of Binaries branches: - main - build-* + tags: + - "v[0-9]+.[0-9]+.[0-9]+" env: TARI_TARGET_NETWORK: nextnet @@ -39,7 +41,23 @@ jobs: run: | cd src cd miner - cargo build --release --bin tarigpuminer + cargo build --release --bin xtrgpuminer - uses: actions/upload-artifact@v4 with: - path: src/miner/target/release/tarigpuminer.exe \ No newline at end of file + name: xtr_miner + path: src/miner/target/release/xtrgpuminer.exe + + create-release: + runs-on: ubuntu-latest + needs: build + if: ${{ startsWith(github.ref, 'refs/tags/v') }} + steps: + - name: Download binaries + uses: actions/download-artifact@v4 + - name: Create release + uses: ncipollo/release-action@v1 + with: + artifacts: "xtr_*/**/*" + token: ${{ secrets.GITHUB_TOKEN }} + prerelease: true + draft: true \ No newline at end of file diff --git a/.idea/.gitignore b/.idea/.gitignore new file mode 100644 index 0000000..b58b603 --- /dev/null +++ b/.idea/.gitignore @@ -0,0 +1,5 @@ +# Default ignored files +/shelf/ +/workspace.xml +# Editor-based HTTP Client requests +/httpRequests/ diff --git a/.idea/dbnavigator.xml b/.idea/dbnavigator.xml new file mode 100644 index 0000000..b827a9b --- /dev/null +++ b/.idea/dbnavigator.xml @@ -0,0 +1,401 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + \ No newline at end of file diff --git a/.idea/git_toolbox_blame.xml b/.idea/git_toolbox_blame.xml new file mode 100644 index 0000000..7dc1249 --- /dev/null +++ b/.idea/git_toolbox_blame.xml @@ -0,0 +1,6 @@ + + + + + \ No newline at end of file diff --git a/.idea/git_toolbox_prj.xml b/.idea/git_toolbox_prj.xml new file mode 100644 index 0000000..02b915b --- /dev/null +++ b/.idea/git_toolbox_prj.xml @@ -0,0 +1,15 @@ + + + + + + + \ No newline at end of file diff --git a/.idea/modules.xml b/.idea/modules.xml new file mode 100644 index 0000000..c6a7948 --- /dev/null +++ b/.idea/modules.xml @@ -0,0 +1,8 @@ + + + + + + + + \ No newline at end of file diff --git a/.idea/tarigpuminer.iml b/.idea/tarigpuminer.iml new file mode 100644 index 0000000..d169918 --- /dev/null +++ b/.idea/tarigpuminer.iml @@ -0,0 +1,11 @@ + + + + + + + + + + + \ No newline at end of file diff --git a/.idea/vcs.xml b/.idea/vcs.xml new file mode 100644 index 0000000..35eb1dd --- /dev/null +++ b/.idea/vcs.xml @@ -0,0 +1,6 @@ + + + + + + \ No newline at end of file diff --git a/src/miner/Cargo.lock b/src/miner/Cargo.lock index 7255795..e226bd8 100644 --- a/src/miner/Cargo.lock +++ b/src/miner/Cargo.lock @@ -418,6 +418,15 @@ version = "0.4.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "771fe0050b883fcc3ea2359b1a96bcfbc090b7116eae7c3c512c7a083fdf23d3" +[[package]] +name = "bs58" +version = "0.5.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bf88ba1141d185c399bee5288d850d63b8369520c1eafc32a0430b5b6c287bf4" +dependencies = [ + "tinyvec", +] + [[package]] name = "bumpalo" version = "3.16.0" @@ -553,6 +562,17 @@ dependencies = [ "zeroize", ] +[[package]] +name = "cl3" +version = "0.9.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b823f24e72fa0c68aa14a250ae1c0848e68d4ae188b71c3972343e45b46f8644" +dependencies = [ + "libc", + "opencl-sys", + "thiserror", +] + [[package]] name = "clap" version = "2.34.0" @@ -612,24 +632,17 @@ checksum = "acbf1af155f9b9ef647e42cdc158db4b64a1b61f743629225fde6f3e0be2a7c7" [[package]] name = "config" -version = "0.13.4" +version = "0.14.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "23738e11972c7643e4ec947840fc463b6a571afcd3e735bdfce7d03c7a784aca" +checksum = "7328b20597b53c2454f0b1919720c25c7339051c02b72b7e05409e00b14132be" dependencies = [ - "async-trait", "lazy_static", "nom", "pathdiff", "serde", - "toml 0.5.11", + "toml 0.8.12", ] -[[package]] -name = "convert_case" -version = "0.4.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "6245d59a3e82a7fc217c5828a6692dbc6dfb63a0c8c90495621f7b9d79704a0e" - [[package]] name = "core-foundation-sys" version = "0.8.6" @@ -704,9 +717,12 @@ dependencies = [ "cfg-if", "cpufeatures", "curve25519-dalek-derive", - "fiat-crypto 0.2.7", + "fiat-crypto", + "group", "platforms", + "rand_core", "rustc_version", + "serde", "subtle", "zeroize", ] @@ -803,19 +819,6 @@ dependencies = [ "syn 1.0.109", ] -[[package]] -name = "derive_more" -version = "0.99.17" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "4fb810d30a7c1953f91334de7244731fc3f3c10d7fe163338a35b9f640960321" -dependencies = [ - "convert_case", - "proc-macro2", - "quote", - "rustc_version", - "syn 1.0.109", -] - [[package]] name = "destructure_traitobject" version = "0.2.0" @@ -957,10 +960,14 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "658bd65b1cf4c852a3cc96f18a8ce7b5640f6b703f905c7d74532294c2a63984" [[package]] -name = "fiat-crypto" -version = "0.1.20" +name = "ff" +version = "0.13.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "e825f6987101665dea6ec934c09ec6d721de7bc1bf92248e1d5810c8cd636b77" +checksum = "ded41244b729663b1e574f1b4fb731469f69f79c17667b5d776b16cda0479449" +dependencies = [ + "rand_core", + "subtle", +] [[package]] name = "fiat-crypto" @@ -1180,6 +1187,17 @@ version = "0.3.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "d2fabcfbdc87f4758337ca535fb41a6d701b65693ce38287d856d1674551ec9b" +[[package]] +name = "group" +version = "0.13.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f0f9ef7462f7c099f518d754361858f86d8a07af53ba9af0fe635bbccb151a63" +dependencies = [ + "ff", + "rand_core", + "subtle", +] + [[package]] name = "h2" version = "0.3.26" @@ -1461,18 +1479,18 @@ checksum = "8f518f335dce6725a761382244631d86cf0ccb2863413590b31338feb467f9c3" [[package]] name = "itertools" -version = "0.6.5" +version = "0.10.5" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d3f2be4da1690a039e9ae5fd575f706a63ad5a2120f161b1d653c9da3930dd21" +checksum = "b0fd2260e829bddf4cb6ea802289de2f86d6a7a690192fbe91b3f46e0f2c8473" dependencies = [ "either", ] [[package]] name = "itertools" -version = "0.10.5" +version = "0.12.1" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "b0fd2260e829bddf4cb6ea802289de2f86d6a7a690192fbe91b3f46e0f2c8473" +checksum = "ba291022dbbd398a455acf126c1e341954079855bc60dfdda641363bd6922569" dependencies = [ "either", ] @@ -1687,8 +1705,8 @@ dependencies = [ [[package]] name = "minotari_app_grpc" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "argon2", "base64 0.13.1", @@ -1714,6 +1732,11 @@ dependencies = [ "zeroize", ] +[[package]] +name = "minotari_ledger_wallet_common" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" + [[package]] name = "mint" version = "0.5.9" @@ -1733,9 +1756,9 @@ dependencies = [ [[package]] name = "monero" -version = "0.20.0" +version = "0.21.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "1b205707fd34b01a547f2fe77e687b40fed05966fb82e955b86ac55cd8ee31b5" +checksum = "f25218523ad4a171ddda05251669afb788cdc2f0df94082aab856a2b09541c3f" dependencies = [ "base58-monero 2.0.0", "curve25519-dalek", @@ -1755,7 +1778,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "3c580bfdd8803cce319b047d239559a22f809094aaea4ac13902a1fdcfcd4261" dependencies = [ "arrayref", - "bs58", + "bs58 0.4.0", "byteorder", "data-encoding", "multihash", @@ -1847,13 +1870,13 @@ checksum = "51d515d32fb182ee37cda2ccdcb92950d6a3c2893aa280e540671c2cd0f3b1d9" [[package]] name = "num-derive" -version = "0.3.3" +version = "0.4.2" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "876a53fff98e03a936a674b29568b0e605f06b29372c2489ff4de23f1949743d" +checksum = "ed3955f1a9c7c0c15e092f9c887db08b1fc683305fdf6eb6684f22555355e202" dependencies = [ "proc-macro2", "quote", - "syn 1.0.109", + "syn 2.0.58", ] [[package]] @@ -1920,6 +1943,25 @@ version = "0.3.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "c08d65885ee38876c4f86fa503fb49d7b507c2b62552df7c70b2fce627e06381" +[[package]] +name = "opencl-sys" +version = "0.2.9" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "de15dd01496ae90c5799f5266184ab020082b4065800ff0b732f489371d0e5cf" +dependencies = [ + "libc", +] + +[[package]] +name = "opencl3" +version = "0.9.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "26ab4a90cb496f787d3934deb0c54fa9d65e7bed710c10071234aab0196fba04" +dependencies = [ + "cl3", + "libc", +] + [[package]] name = "ordered-float" version = "2.10.1" @@ -2868,24 +2910,6 @@ version = "1.0.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "55937e1799185b12863d447f42597ed69d9928686b8d88a1df17376a097d8369" -[[package]] -name = "tari-curve25519-dalek" -version = "4.0.3" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d9b8e2644aae57a832e475ebc31199ab1114ebd7fe4d2621e67e89bdd9c8ac38" -dependencies = [ - "cfg-if", - "cpufeatures", - "curve25519-dalek-derive", - "fiat-crypto 0.1.20", - "platforms", - "rand_core", - "rustc_version", - "serde", - "subtle", - "zeroize", -] - [[package]] name = "tari-tiny-keccak" version = "2.0.2" @@ -2898,34 +2922,31 @@ dependencies = [ [[package]] name = "tari_bulletproofs_plus" -version = "0.3.2" +version = "0.4.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "9abc9cbebffb1149112a80d6955498ed891f9b786e89c50e07de4b4f16290fd8" +checksum = "8eac57729e3b003c18e822c64c8c4977770102b2eaea920a7c40bca5caf12c54" dependencies = [ "blake2", "byteorder", - "derivative", - "derive_more", + "curve25519-dalek", "digest", - "itertools 0.6.5", - "lazy_static", + "ff", + "itertools 0.12.1", "merlin", - "rand", + "once_cell", "rand_core", "serde", "sha3", - "tari-curve25519-dalek", - "thiserror", + "thiserror-no-std", "zeroize", ] [[package]] name = "tari_common" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "anyhow", - "blake2", "config", "dirs-next", "log", @@ -2938,7 +2959,6 @@ dependencies = [ "serde_yaml", "sha2", "structopt", - "tari_crypto", "tari_features", "tempfile", "thiserror", @@ -2947,8 +2967,8 @@ dependencies = [ [[package]] name = "tari_common_sqlite" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "diesel", "diesel_migrations", @@ -2961,19 +2981,24 @@ dependencies = [ [[package]] name = "tari_common_types" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "base64 0.21.7", + "bitflags 2.5.0", "blake2", "borsh", + "bs58 0.5.1", "chacha20poly1305", "digest", + "minotari_ledger_wallet_common", "newtype-ops", "once_cell", "primitive-types", "rand", "serde", + "strum", + "strum_macros", "tari_common", "tari_crypto", "tari_utilities", @@ -2982,8 +3007,8 @@ dependencies = [ [[package]] name = "tari_comms" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "anyhow", "async-trait", @@ -3026,8 +3051,8 @@ dependencies = [ [[package]] name = "tari_comms_dht" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "anyhow", "bitflags 2.5.0", @@ -3061,8 +3086,8 @@ dependencies = [ [[package]] name = "tari_comms_rpc_macros" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "proc-macro2", "quote", @@ -3071,8 +3096,8 @@ dependencies = [ [[package]] name = "tari_core" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "async-trait", "bincode", @@ -3085,6 +3110,7 @@ dependencies = [ "decimal-rs", "derivative", "digest", + "dirs-next", "fs2", "futures 0.3.30", "hex", @@ -3136,21 +3162,23 @@ dependencies = [ [[package]] name = "tari_crypto" -version = "0.20.0" +version = "0.20.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "63a3ed2c551101eb42b7f9386c207e28d53e6816f7b4c9a0883548922f317b3e" +checksum = "22b761f7cf1754eb2223286c9d437c81737526f393f09f69b2456bf49ba25a5b" dependencies = [ "blake2", "borsh", + "curve25519-dalek", "digest", "log", + "merlin", "once_cell", "rand_chacha", "rand_core", "serde", "sha3", "snafu", - "tari-curve25519-dalek", + "subtle", "tari_bulletproofs_plus", "tari_utilities", "zeroize", @@ -3158,13 +3186,13 @@ dependencies = [ [[package]] name = "tari_features" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" [[package]] name = "tari_hashing" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "borsh", "digest", @@ -3173,8 +3201,8 @@ dependencies = [ [[package]] name = "tari_key_manager" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "argon2", "async-trait", @@ -3206,14 +3234,13 @@ dependencies = [ [[package]] name = "tari_mmr" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "borsh", "digest", "log", "serde", - "tari_common", "tari_crypto", "tari_utilities", "thiserror", @@ -3221,8 +3248,8 @@ dependencies = [ [[package]] name = "tari_p2p" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "anyhow", "fs2", @@ -3251,8 +3278,8 @@ dependencies = [ [[package]] name = "tari_script" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "blake2", "borsh", @@ -3268,8 +3295,8 @@ dependencies = [ [[package]] name = "tari_service_framework" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "anyhow", "async-trait", @@ -3283,16 +3310,16 @@ dependencies = [ [[package]] name = "tari_shutdown" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "futures 0.3.30", ] [[package]] name = "tari_storage" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "bincode", "lmdb-zero", @@ -3303,8 +3330,8 @@ dependencies = [ [[package]] name = "tari_test_utils" -version = "1.0.0-rc.6" -source = "git+http://github.com/tari-project/tari?tag=v1.0.0-rc.6#8fe6eb1b8a402f7a9c6026044edd1c61d0e75b33" +version = "1.1.0-pre.2" +source = "git+http://github.com/tari-project/tari?tag=v1.1.0-pre.2#90a4c282d5ee61ccd1abd54b8e14e9b5c75052a9" dependencies = [ "futures 0.3.30", "rand", @@ -3333,32 +3360,6 @@ dependencies = [ "zeroize", ] -[[package]] -name = "tarigpuminer" -version = "0.1.4" -dependencies = [ - "anyhow", - "clap 4.5.4", - "cust", - "libsqlite3-sys", - "minotari_app_grpc", - "num-format", - "prost", - "prost-types", - "rand", - "serde", - "serde_json", - "sha3", - "tari_common", - "tari_common_types", - "tari_core", - "tari_crypto", - "tari_key_manager", - "tari_script", - "tokio", - "tonic", -] - [[package]] name = "tempfile" version = "3.10.1" @@ -3400,6 +3401,26 @@ dependencies = [ "syn 2.0.58", ] +[[package]] +name = "thiserror-impl-no-std" +version = "2.0.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "58e6318948b519ba6dc2b442a6d0b904ebfb8d411a3ad3e07843615a72249758" +dependencies = [ + "proc-macro2", + "quote", + "syn 1.0.109", +] + +[[package]] +name = "thiserror-no-std" +version = "2.0.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a3ad459d94dd517257cc96add8a43190ee620011bb6e6cdc82dafd97dfafafea" +dependencies = [ + "thiserror-impl-no-std", +] + [[package]] name = "time" version = "0.3.36" @@ -3568,6 +3589,18 @@ dependencies = [ "toml_edit 0.19.15", ] +[[package]] +name = "toml" +version = "0.8.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e9dd1545e8208b4a5af1aa9bbd0b4cf7e9ea08fabc5d0a5c67fcaafa17433aa3" +dependencies = [ + "serde", + "serde_spanned", + "toml_datetime", + "toml_edit 0.22.12", +] + [[package]] name = "toml_datetime" version = "0.6.5" @@ -3587,7 +3620,7 @@ dependencies = [ "serde", "serde_spanned", "toml_datetime", - "winnow", + "winnow 0.5.40", ] [[package]] @@ -3598,7 +3631,7 @@ checksum = "70f427fce4d84c72b5b732388bf4a9f4531b53f74e2887e3ecb2481f68f66d81" dependencies = [ "indexmap 2.2.6", "toml_datetime", - "winnow", + "winnow 0.5.40", ] [[package]] @@ -3609,7 +3642,20 @@ checksum = "6a8534fd7f78b5405e860340ad6575217ce99f38d4d5c8f2442cb5ecb50090e1" dependencies = [ "indexmap 2.2.6", "toml_datetime", - "winnow", + "winnow 0.5.40", +] + +[[package]] +name = "toml_edit" +version = "0.22.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d3328d4f68a705b2a4498da1d580585d39a6510f98318a2cec3018a7ec61ddef" +dependencies = [ + "indexmap 2.2.6", + "serde", + "serde_spanned", + "toml_datetime", + "winnow 0.6.18", ] [[package]] @@ -4022,6 +4068,16 @@ dependencies = [ "wasm-bindgen", ] +[[package]] +name = "web-time" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5a6580f308b1fad9207618087a65c04e7a10bc77e02c8e84e9b00dd4b12fa0bb" +dependencies = [ + "js-sys", + "wasm-bindgen", +] + [[package]] name = "webpki" version = "0.22.4" @@ -4216,6 +4272,15 @@ dependencies = [ "memchr", ] +[[package]] +name = "winnow" +version = "0.6.18" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "68a9bda4691f099d435ad181000724da8e5899daa10713c2d432552b9ccd3a6f" +dependencies = [ + "memchr", +] + [[package]] name = "wyz" version = "0.5.1" @@ -4225,18 +4290,49 @@ dependencies = [ "tap", ] +[[package]] +name = "xtrgpuminer" +version = "0.1.6" +dependencies = [ + "anyhow", + "clap 4.5.4", + "cust", + "libsqlite3-sys", + "minotari_app_grpc", + "num-format", + "opencl-sys", + "opencl3", + "prost", + "prost-types", + "rand", + "serde", + "serde_json", + "sha3", + "tari_common", + "tari_common_types", + "tari_core", + "tari_crypto", + "tari_key_manager", + "tari_script", + "tari_utilities", + "tokio", + "tonic", +] + [[package]] name = "yamux" -version = "0.10.2" +version = "0.13.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "e5d9ba232399af1783a58d8eb26f6b5006fbefe2dc9ef36bd283324792d03ea5" +checksum = "a31b5e376a8b012bee9c423acdbb835fc34d45001cfa3106236a624e4b738028" dependencies = [ "futures 0.3.30", "log", "nohash-hasher", "parking_lot", + "pin-project 1.1.5", "rand", "static_assertions", + "web-time", ] [[package]] diff --git a/src/miner/Cargo.toml b/src/miner/Cargo.toml index d9c51ea..6a31599 100644 --- a/src/miner/Cargo.toml +++ b/src/miner/Cargo.toml @@ -1,19 +1,20 @@ [package] name = "xtrgpuminer" -version = "0.1.5" +version = "0.1.6" edition = "2021" [dependencies] -tari_core = { git = "https://github.com/tari-project/tari", tag = "v1.0.0-rc.6", default-features = false, features = [ +tari_core = { git = "https://github.com/tari-project/tari", tag = "v1.1.0-pre.2", default-features = false, features = [ "transactions", ] } -minotari_app_grpc = { git = "http://github.com/tari-project/tari", tag = "v1.0.0-rc.6" } -tari_common_types = { git = "http://github.com/tari-project/tari", tag = "v1.0.0-rc.6" } -tari_common = { git = "http://github.com/tari-project/tari", tag = "v1.0.0-rc.6" } -tari_script = { git = "http://github.com/tari-project/tari", tag = "v1.0.0-rc.6" } -tari_key_manager = { git = "http://github.com/tari-project/tari", tag = "v1.0.0-rc.6" } -tari_crypto = { version = "0.20", features = ["borsh"] } +minotari_app_grpc = { git = "http://github.com/tari-project/tari", tag = "v1.1.0-pre.2" } +tari_common_types = { git = "http://github.com/tari-project/tari", tag = "v1.1.0-pre.2" } +tari_common = { git = "http://github.com/tari-project/tari", tag = "v1.1.0-pre.2" } +tari_script = { git = "http://github.com/tari-project/tari", tag = "v1.1.0-pre.2" } +tari_key_manager = { git = "http://github.com/tari-project/tari", tag = "v1.1.0-pre.2" } +tari_crypto = { version = "0.20.3", features = ["borsh"] } +tari_utilities = "0.7" serde = { version = "1.0.130", features = ["derive"] } serde_json = "1.0.68" anyhow = "*" @@ -28,4 +29,10 @@ prost-types = "0.11.9" rand = "0.8" libsqlite3-sys = { version = "0.25.1", features = ["bundled"] } -cust = "0.3.2" +cust = { version ="0.3.2", optional=true } +opencl3 = { version="0.9.5", optional=true} +opencl-sys = "*" + +[features] +default = [] +nvidia = ["cust"] diff --git a/src/miner/rustfmt.toml b/src/miner/rustfmt.toml new file mode 100644 index 0000000..13868eb --- /dev/null +++ b/src/miner/rustfmt.toml @@ -0,0 +1,27 @@ +binop_separator = "Back" +use_small_heuristics = "default" +comment_width = 120 +edition = "2018" +format_code_in_doc_comments = true +format_strings = true +group_imports = "StdExternalCrate" +hard_tabs = false +imports_layout = "HorizontalVertical" +imports_granularity = "Crate" +match_block_trailing_comma = true +max_width = 120 +newline_style = "Auto" +normalize_comments = true +overflow_delimited_expr = true +reorder_imports = true +reorder_modules = true +reorder_impl_items = true +space_after_colon = true +space_before_colon = false +struct_lit_single_line = true +use_field_init_shorthand = true +use_try_shorthand = true +unstable_features = true +where_single_line = true +wrap_comments = true +ignore = [] diff --git a/src/miner/src/config_file.rs b/src/miner/src/config_file.rs index 26ad0af..0ab363f 100644 --- a/src/miner/src/config_file.rs +++ b/src/miner/src/config_file.rs @@ -1,37 +1,37 @@ -use anyhow; -use std::{fs::File, io::BufReader, path::PathBuf}; - -#[derive(serde::Deserialize, serde::Serialize, Debug, Clone)] -pub(crate) struct ConfigFile { - pub tari_address: String, - pub tari_node_url: String, - pub coinbase_extra: String, - pub template_refresh_secs: u64, -} - -impl Default for ConfigFile { - fn default() -> Self { - Self { - tari_address: "8c98d40f216589d8b385015222b95fb5327fee334352c7c30370101b0c6d124fd6" - .to_string(), - tari_node_url: "http://127.0.0.1:18182".to_string(), - coinbase_extra: "tari_gpu_miner".to_string(), - template_refresh_secs: 30, - } - } -} - -impl ConfigFile { - pub(crate) fn load(path: PathBuf) -> Result { - let file = File::open(path)?; - let reader = BufReader::new(file); - let config = serde_json::from_reader(reader)?; - Ok(config) - } - - pub(crate) fn save(&self, path: &str) -> Result<(), anyhow::Error> { - let file = File::create(path)?; - serde_json::to_writer_pretty(file, self)?; - Ok(()) - } -} +use std::{fs::File, io::BufReader, path::PathBuf}; + +use anyhow; + +#[derive(serde::Deserialize, serde::Serialize, Debug, Clone)] +pub(crate) struct ConfigFile { + pub tari_address: String, + pub tari_node_url: String, + pub coinbase_extra: String, + pub template_refresh_secs: u64, +} + +impl Default for ConfigFile { + fn default() -> Self { + Self { + tari_address: "8c98d40f216589d8b385015222b95fb5327fee334352c7c30370101b0c6d124fd6".to_string(), + tari_node_url: "http://127.0.0.1:18142".to_string(), + coinbase_extra: "tari_gpu_miner".to_string(), + template_refresh_secs: 30, + } + } +} + +impl ConfigFile { + pub(crate) fn load(path: PathBuf) -> Result { + let file = File::open(path)?; + let reader = BufReader::new(file); + let config = serde_json::from_reader(reader)?; + Ok(config) + } + + pub(crate) fn save(&self, path: &str) -> Result<(), anyhow::Error> { + let file = File::create(path)?; + serde_json::to_writer_pretty(file, self)?; + Ok(()) + } +} diff --git a/src/miner/src/context_impl.rs b/src/miner/src/context_impl.rs new file mode 100644 index 0000000..dc52a24 --- /dev/null +++ b/src/miner/src/context_impl.rs @@ -0,0 +1 @@ +pub trait ContextImpl {} diff --git a/src/miner/src/cuda_engine.rs b/src/miner/src/cuda_engine.rs new file mode 100644 index 0000000..23130af --- /dev/null +++ b/src/miner/src/cuda_engine.rs @@ -0,0 +1,135 @@ +use crate::context_impl::ContextImpl; +use crate::EngineImpl; +use crate::FunctionImpl; +use anyhow::Error; +#[cfg(feature = "nvidia")] +use cust::{ + device::DeviceAttribute, + memory::{AsyncCopyDestination, DeviceCopy}, + module::{ModuleJitOption, ModuleJitOption::DetermineTargetFromContext}, + prelude::{Module, *}, +}; + +use std::time::Instant; +#[derive(Clone)] +pub struct CudaEngine {} + +impl CudaEngine { + pub fn new() -> Self { + Self {} + } +} + +impl EngineImpl for CudaEngine { + type Context = CudaContext; + type Function = CudaFunction; + + fn init(&mut self) -> Result<(), anyhow::Error> { + cust::init(CudaFlags::empty())?; + Ok(()) + } + + fn num_devices(&self) -> Result { + let num_devices = Device::num_devices()?; + Ok(num_devices) + } + + fn create_context(&self, device_index: u32) -> Result { + let context = Context::new(Device::get_device(device_index)?)?; + context.set_flags(ContextFlags::SCHED_YIELD)?; + + Ok(CudaContext { context }) + } + + fn create_main_function(&self, context: &Self::Context) -> Result { + let module = Module::from_ptx( + include_str!("../cuda/keccak.ptx"), + &[ModuleJitOption::GenerateLineInfo(true)], + )?; + // let func = context.module.get_function("keccakKernel")?; + Ok(CudaFunction { module }) + } + + fn mine( + &self, + function: &Self::Function, + context: &Self::Context, + data: &[u64], + min_difficulty: u64, + nonce_start: u64, + num_iterations: u32, + block_size: u32, + grid_size: u32, + ) -> Result<(Option, u32, u64), Error> { + let output = vec![0u64; 5]; + let mut output_buf = output.as_slice().as_dbuf()?; + + let mut data_buf = data.as_dbuf()?; + data_buf.copy_from(data).expect("Could not copy data to buffer"); + output_buf.copy_from(&output).expect("Could not copy output to buffer"); + + let num_streams = 1; + let mut streams = Vec::with_capacity(num_streams); + let func = function.module.get_function("keccakKernel")?; + + let output = vec![0u64; 5]; + + for st in 0..num_streams { + let stream = Stream::new(StreamFlags::NON_BLOCKING, None)?; + + streams.push(stream); + } + + let data_ptr = data_buf.as_device_ptr(); + for st in 0..num_streams { + let stream = &streams[st]; + unsafe { + launch!( + func<<>>( + data_ptr, + nonce_start, + min_difficulty, + num_iterations, + output_buf.as_device_ptr(), + + ) + )?; + } + } + + for st in 0..num_streams { + let mut out1 = vec![0u64; 5]; + + unsafe { + output_buf.copy_to(&mut out1)?; + } + //stream.synchronize()?; + + if out1[0] > 0 { + return Ok((Some((&out1[0]).clone()), grid_size * block_size * num_iterations, 0)); + } + } + + Ok((None, grid_size * block_size * num_iterations, 0)) + } +} + +pub struct CudaContext { + context: Context, +} + +impl CudaContext {} + +impl ContextImpl for CudaContext {} + +pub struct CudaFunction { + module: Module, +} +impl FunctionImpl for CudaFunction { + fn suggested_launch_configuration(&self) -> Result<(u32, u32), anyhow::Error> { + let func = self.module.get_function("keccakKernel")?; + let (grid_size, block_size) = func.suggested_launch_configuration(0, 0.into())?; + // Ok((grid_size, block_size)) + Ok((1000, 100)) + } +} diff --git a/src/miner/src/engine_impl.rs b/src/miner/src/engine_impl.rs new file mode 100644 index 0000000..49df2d7 --- /dev/null +++ b/src/miner/src/engine_impl.rs @@ -0,0 +1,25 @@ +use crate::{context_impl::ContextImpl, function_impl::FunctionImpl}; + +pub trait EngineImpl { + type Context: ContextImpl; + type Function: FunctionImpl; + fn init(&mut self) -> Result<(), anyhow::Error>; + + fn num_devices(&self) -> Result; + + fn create_context(&self, device_index: u32) -> Result; + + fn create_main_function(&self, context: &Self::Context) -> Result; + + fn mine( + &self, + function: &Self::Function, + context: &Self::Context, + data: &[u64], + min_difficulty: u64, + nonce_start: u64, + num_iterations: u32, + block_size: u32, + grid_size: u32, + ) -> Result<(Option, u32, u64), anyhow::Error>; +} diff --git a/src/miner/src/function_impl.rs b/src/miner/src/function_impl.rs new file mode 100644 index 0000000..ea5ff6c --- /dev/null +++ b/src/miner/src/function_impl.rs @@ -0,0 +1,3 @@ +pub trait FunctionImpl { + fn suggested_launch_configuration(&self) -> Result<(u32, u32), anyhow::Error>; +} diff --git a/src/miner/src/gpu_engine.rs b/src/miner/src/gpu_engine.rs new file mode 100644 index 0000000..e67e710 --- /dev/null +++ b/src/miner/src/gpu_engine.rs @@ -0,0 +1,65 @@ +use crate::engine_impl::EngineImpl; + +#[derive(Clone)] +pub struct GpuEngine { + inner: TEngineImpl, +} + +impl GpuEngine { + pub fn new(engine: TEngineImpl) -> Self { + GpuEngine { inner: engine } + } + + pub fn init(&mut self) -> Result<(), anyhow::Error> { + self.inner.init() + } + + pub fn num_devices(&self) -> Result { + self.inner.num_devices() + } + + pub fn create_context(&self, device_index: u32) -> Result { + self.inner.create_context(device_index) + } + + pub fn get_main_function(&self, context: &TEngineImpl::Context) -> Result { + self.inner.create_main_function(context) + // match self { + // GpuEngine::Cuda => { + // let module = Module::from_ptx(include_str!("../cuda/keccak.ptx"), &[ + // ModuleJitOption::GenerateLineInfo(true), + // ]) + // .context("module bad")?; + // + // let func = module.get_function("keccakKernel").context("module getfunc")?; + // todo!() + // }, + // GpuEngine::OpenCL => { + // todo!() + // } + // } + } + + pub fn mine( + &self, + function: &TEngineImpl::Function, + context: &TEngineImpl::Context, + data: &[u64], + min_difficulty: u64, + nonce_start: u64, + num_iterations: u32, + block_size: u32, + grid_size: u32, + ) -> Result<(Option, u32, u64), anyhow::Error> { + self.inner.mine( + function, + context, + data, + min_difficulty, + nonce_start, + num_iterations, + block_size, + grid_size, + ) + } +} diff --git a/src/miner/src/main.rs b/src/miner/src/main.rs index 9ee05f9..3d6bac4 100644 --- a/src/miner/src/main.rs +++ b/src/miner/src/main.rs @@ -1,423 +1,335 @@ -use crate::config_file::ConfigFile; -use crate::node_client::NodeClient; -use crate::tari_coinbase::generate_coinbase; -use anyhow::anyhow; -use anyhow::Context as AnyContext; -use clap::Parser; -use cust::device::DeviceAttribute; -use cust::memory::AsyncCopyDestination; -use cust::memory::DeviceCopy; -use cust::module::ModuleJitOption; -use cust::module::ModuleJitOption::DetermineTargetFromContext; -use cust::prelude::Module; -use cust::prelude::*; -use minotari_app_grpc::tari_rpc::BlockHeader as grpc_header; -use minotari_app_grpc::tari_rpc::TransactionOutput as GrpcTransactionOutput; -use num_format::Locale; -use num_format::ToFormattedString; -use sha3::digest::crypto_common::rand_core::block; -use sha3::Digest; -use sha3::Sha3_256; -use std::cmp; -use std::convert::TryInto; -use std::env::current_dir; -use std::iter; -use std::num; -use std::path::PathBuf; -use std::sync::Arc; -use std::thread; -use std::time::Instant; -use tari_common::configuration::Network; -use tari_common_types::tari_address::TariAddress; -use tari_common_types::types::FixedHash; -use tari_core::blocks::BlockHeader; -use tari_core::consensus::ConsensusManager; -use tari_core::proof_of_work::sha3x_difficulty; -use tari_core::proof_of_work::Difficulty; -use tari_core::transactions::key_manager::create_memory_db_key_manager; -use tari_core::transactions::tari_amount::MicroMinotari; -use tari_core::transactions::transaction_components::RangeProofType; -use tokio::runtime::Handle; -use tokio::runtime::Runtime; -use tokio::sync::RwLock; -use tokio::task; -use tokio::task::JoinSet; -use tokio::time::sleep; -use tokio::try_join; - -mod config_file; -mod node_client; -mod tari_coinbase; - -#[tokio::main] -async fn main() { - match main_inner().await { - Ok(()) => {} - Err(err) => { - eprintln!("Error: {:#?}", err); - std::process::exit(1); - } - } -} - -#[derive(Parser)] -struct Cli { - #[arg(short, long, value_name = "FILE")] - config: Option, - #[arg(short, long)] - benchmark: bool, -} - -async fn main_inner() -> Result<(), anyhow::Error> { - let cli = Cli::parse(); - - let benchmark = cli.benchmark; - - let config = match ConfigFile::load(cli.config.unwrap_or_else(|| { - let mut path = current_dir().expect("no current directory"); - path.push("config.json"); - path - })) { - Ok(config) => config, - Err(err) => { - eprintln!("Error loading config file: {}. Creating new one", err); - let default = ConfigFile::default(); - default - .save("config.json") - .expect("Could not save default config"); - default - } - }; - - let submit = true; - - cust::init(CudaFlags::empty())?; - - let num_devices = Device::num_devices()?; - let mut threads = vec![]; - for i in 0..num_devices { - let c = config.clone(); - threads.push(thread::spawn(move || { - run_thread(num_devices as u64, i as u64, c, benchmark) - })); - } - - for t in threads { - t.join().unwrap()?; - } - - Ok(()) -} - -fn run_thread( - num_threads: u64, - thread_index: u64, - config: ConfigFile, - benchmark: bool, -) -> Result<(), anyhow::Error> { - let tari_node_url = config.tari_node_url.clone(); - let runtime = Runtime::new()?; - let node_client = Arc::new(RwLock::new(runtime.block_on(async move { - node_client::create_client(&tari_node_url, benchmark).await - // node_client::NodeClient::connect(&tari_node_url).await - })?)); - let mut rounds = 0; - - let context = Context::new(Device::get_device(thread_index as u32)?)?; - context.set_flags(ContextFlags::SCHED_YIELD)?; - let module = Module::from_ptx( - include_str!("../cuda/keccak.ptx"), - &[ModuleJitOption::GenerateLineInfo(true)], - ) - .context("module bad")?; - - let func = module - .get_function("keccakKernel") - .context("module getfunc")?; - let (grid_size, block_size) = func - .suggested_launch_configuration(0, 0.into()) - .context("get suggest config")?; - // let (grid_size, block_size) = (23, 50); - - let output = vec![0u64; 5]; - let mut output_buf = output.as_slice().as_dbuf()?; - - let mut data = vec![0u64; 6]; - let mut data_buf = data.as_slice().as_dbuf()?; - - loop { - rounds += 1; - if rounds > 101 { - rounds = 0; - } - let clone_node_client = node_client.clone(); - let clone_config = config.clone(); - let (target_difficulty, block, mut header, mining_hash) = runtime.block_on(async move { - get_template(clone_config, clone_node_client, rounds, benchmark).await - })?; - - let hash64 = copy_u8_to_u64(mining_hash.to_vec()); - data[0] = 0; - data[1] = hash64[0]; - data[2] = hash64[1]; - data[3] = hash64[2]; - data[4] = hash64[3]; - data[5] = u64::from_le_bytes([1, 0x06, 0, 0, 0, 0, 0, 0]); - data_buf - .copy_from(&data) - .expect("Could not copy data to buffer"); - output_buf - .copy_from(&output) - .expect("Could not copy output to buffer"); - - let mut nonce_start = (u64::MAX / num_threads) * thread_index; - let mut last_hash_rate = 0; - let elapsed = Instant::now(); - let mut max_diff = 0; - let mut last_printed = Instant::now(); - loop { - if elapsed.elapsed().as_secs() > config.template_refresh_secs { - break; - } - let (nonce, hashes, diff) = mine( - mining_hash, - header.pow.to_bytes(), - (u64::MAX / (target_difficulty)).to_le(), - nonce_start, - &context, - &module, - 4, - &func, - block_size, - grid_size, - data_buf.as_device_ptr(), - &output_buf, - )?; - if let Some(ref n) = nonce { - header.nonce = *n; - } - if diff > max_diff { - max_diff = diff; - } - nonce_start = nonce_start + hashes as u64; - if elapsed.elapsed().as_secs() > 1 { - if Instant::now() - last_printed > std::time::Duration::from_secs(2) { - last_printed = Instant::now(); - println!( - "total {:} grid: {} max_diff: {}, target: {} hashes/sec: {}", - nonce_start.to_formatted_string(&Locale::en), - grid_size, - max_diff.to_formatted_string(&Locale::en), - target_difficulty.to_formatted_string(&Locale::en), - (nonce_start / elapsed.elapsed().as_secs()) - .to_formatted_string(&Locale::en) - ); - } - last_hash_rate = nonce_start / elapsed.elapsed().as_secs(); - } - if nonce.is_some() { - header.nonce = nonce.unwrap(); - - let mut mined_block = block.clone(); - mined_block.header = Some(grpc_header::from(header)); - let clone_client = node_client.clone(); - match runtime.block_on(async { - clone_client.write().await.submit_block(mined_block).await - }) { - Ok(_) => { - println!("Block submitted"); - } - Err(e) => { - println!("Error submitting block: {:?}", e); - } - } - break; - } - // break; - } - } -} - -async fn get_template( - config: ConfigFile, - node_client: Arc>, - round: u32, - benchmark: bool, -) -> Result< - ( - u64, - minotari_app_grpc::tari_rpc::Block, - BlockHeader, - FixedHash, - ), - anyhow::Error, -> { - if benchmark { - return Ok(( - u64::MAX, - minotari_app_grpc::tari_rpc::Block::default(), - BlockHeader::new(0), - FixedHash::default(), - )); - } - let address = if round % 99 == 0 { - TariAddress::from_hex("8c98d40f216589d8b385015222b95fb5327fee334352c7c30370101b0c6d124fd6")? - } else { - TariAddress::from_hex(config.tari_address.as_str())? - }; - let key_manager = create_memory_db_key_manager(); - let consensus_manager = ConsensusManager::builder(Network::NextNet) - .build() - .expect("Could not build consensus manager"); - println!("Getting block template"); - let mut lock = node_client.write().await; - let template = lock.get_block_template().await?; - let mut block_template = template.new_block_template.clone().unwrap(); - let height = block_template.header.as_ref().unwrap().height; - let miner_data = template.miner_data.unwrap(); - let fee = MicroMinotari::from(miner_data.total_fees); - let reward = MicroMinotari::from(miner_data.reward); - let (coinbase_output, coinbase_kernel) = generate_coinbase( - fee, - reward, - height, - config.coinbase_extra.as_bytes(), - &key_manager, - &address, - true, - consensus_manager.consensus_constants(height), - RangeProofType::RevealedValue, - ) - .await?; - let body = block_template.body.as_mut().expect("no block body"); - let grpc_output = - GrpcTransactionOutput::try_from(coinbase_output.clone()).map_err(|s| anyhow!(s))?; - body.outputs.push(grpc_output); - body.kernels.push(coinbase_kernel.into()); - let target_difficulty = miner_data.target_difficulty; - let block_result = lock.get_new_block(block_template).await?; - let block = block_result.block.unwrap(); - let mut header: BlockHeader = block - .clone() - .header - .unwrap() - .try_into() - .map_err(|s: String| anyhow!(s))?; - let mining_hash = header.mining_hash().clone(); - Ok((target_difficulty, block, header, mining_hash)) -} - -fn copy_u8_to_u64(input: Vec) -> Vec { - let mut output: Vec = Vec::with_capacity(input.len() / 8); - - for chunk in input.chunks_exact(8) { - let value = u64::from_le_bytes([ - chunk[0], chunk[1], chunk[2], chunk[3], chunk[4], chunk[5], chunk[6], chunk[7], - ]); - output.push(value); - } - - let remaining_bytes = input.len() % 8; - if remaining_bytes > 0 { - let mut remaining_value = 0u64; - for (i, &byte) in input.iter().rev().take(remaining_bytes).enumerate() { - remaining_value |= (byte as u64) << (8 * i); - } - output.push(remaining_value); - } - - output -} - -fn copy_u64_to_u8(input: Vec) -> Vec { - let mut output: Vec = Vec::with_capacity(input.len() * 8); - - for value in input { - output.extend_from_slice(&value.to_le_bytes()); - } - - output -} - -fn mine( - mining_hash: FixedHash, - pow: Vec, - target: u64, - nonce_start: u64, - context: &Context, - module: &Module, - num_iterations: u32, - func: &Function<'_>, - block_size: u32, - grid_size: u32, - data_ptr: DevicePointer, - output_buf: &DeviceBuffer, -) -> Result<(Option, u32, u64), anyhow::Error> { - let num_streams = 1; - let mut streams = Vec::with_capacity(num_streams); - let mut max = None; - - let output = vec![0u64; 5]; - - let timer = Instant::now(); - for st in 0..num_streams { - let stream = Stream::new(StreamFlags::NON_BLOCKING, None)?; - - streams.push(stream); - } - - for st in 0..num_streams { - let stream = &streams[st]; - unsafe { - launch!( - func<<>>( - data_ptr, - nonce_start, - target, - num_iterations, - output_buf.as_device_ptr(), - - ) - )?; - } - } - - for st in 0..num_streams { - let mut out1 = vec![0u64; 5]; - - unsafe { - output_buf.copy_to(&mut out1)?; - } - //stream.synchronize()?; - - if out1[0] > 0 { - return Ok(( - Some((&out1[0]).clone()), - grid_size * block_size * num_iterations, - 0, - )); - } - } - - match max { - Some((i, diff)) => { - if diff > target { - return Ok(( - Some(i), - grid_size * block_size * num_iterations * num_streams as u32, - diff, - )); - } - return Ok(( - None, - grid_size * block_size * num_iterations * num_streams as u32, - diff, - )); - } - None => Ok(( - None, - grid_size * block_size * num_iterations * num_streams as u32, - 0, - )), - } -} +use std::{cmp, convert::TryInto, env::current_dir, iter, num, path::PathBuf, sync::Arc, thread, time::Instant}; + +use anyhow::{anyhow, Context as AnyContext}; +use clap::Parser; +#[cfg(feature = "nvidia")] +use cust::{ + device::DeviceAttribute, + memory::{AsyncCopyDestination, DeviceCopy}, + module::{ModuleJitOption, ModuleJitOption::DetermineTargetFromContext}, + prelude::{Module, *}, +}; +use minotari_app_grpc::tari_rpc::{BlockHeader as grpc_header, TransactionOutput as GrpcTransactionOutput}; +use num_format::{Locale, ToFormattedString}; +use sha3::{digest::crypto_common::rand_core::block, Digest, Sha3_256}; +use std::str::FromStr; +use tari_common::configuration::Network; +use tari_common_types::{tari_address::TariAddress, types::FixedHash}; +use tari_core::{ + blocks::BlockHeader, + consensus::ConsensusManager, + proof_of_work::{sha3x_difficulty, Difficulty}, + transactions::{ + key_manager::create_memory_db_key_manager, tari_amount::MicroMinotari, transaction_components::RangeProofType, + }, +}; +use tari_utilities::epoch_time::EpochTime; +use tokio::{ + runtime::{Handle, Runtime}, + sync::RwLock, + task, + task::JoinSet, + time::sleep, + try_join, +}; + +#[cfg(feature = "opencl3")] +use crate::opencl_engine::OpenClEngine; +use crate::{ + config_file::ConfigFile, engine_impl::EngineImpl, function_impl::FunctionImpl, gpu_engine::GpuEngine, + node_client::NodeClient, tari_coinbase::generate_coinbase, +}; + +mod config_file; +mod context_impl; +#[cfg(feature = "nvidia")] +mod cuda_engine; +#[cfg(feature = "nvidia")] +use crate::cuda_engine::CudaEngine; + +mod engine_impl; +mod function_impl; +mod gpu_engine; +mod node_client; +#[cfg(feature = "opencl3")] +mod opencl_engine; +mod tari_coinbase; + +#[tokio::main] +async fn main() { + match main_inner().await { + Ok(()) => {}, + Err(err) => { + eprintln!("Error: {:#?}", err); + std::process::exit(1); + }, + } +} + +#[derive(Parser)] +struct Cli { + #[arg(short, long, value_name = "FILE")] + config: Option, + #[arg(short, long)] + benchmark: bool, +} + +async fn main_inner() -> Result<(), anyhow::Error> { + let cli = Cli::parse(); + + let benchmark = cli.benchmark; + + let config = match ConfigFile::load(cli.config.unwrap_or_else(|| { + let mut path = current_dir().expect("no current directory"); + path.push("config.json"); + path + })) { + Ok(config) => config, + Err(err) => { + eprintln!("Error loading config file: {}. Creating new one", err); + let default = ConfigFile::default(); + default.save("config.json").expect("Could not save default config"); + default + }, + }; + + let submit = true; + + #[cfg(feature = "nvidia")] + let mut gpu_engine = GpuEngine::new(CudaEngine::new()); + + #[cfg(feature = "opencl3")] + let mut gpu_engine = GpuEngine::new(OpenClEngine::new()); + + gpu_engine.init(); + + let num_devices = gpu_engine.num_devices()?; + let mut threads = vec![]; + for i in 0..num_devices { + let c = config.clone(); + let gpu = gpu_engine.clone(); + threads.push(thread::spawn(move || { + run_thread(gpu, num_devices as u64, i as u32, c, benchmark) + })); + } + + for t in threads { + t.join().unwrap()?; + } + + Ok(()) +} + +fn run_thread( + gpu_engine: GpuEngine, + num_threads: u64, + thread_index: u32, + config: ConfigFile, + benchmark: bool, +) -> Result<(), anyhow::Error> { + let tari_node_url = config.tari_node_url.clone(); + let runtime = Runtime::new()?; + let node_client = Arc::new(RwLock::new(runtime.block_on(async move { + node_client::create_client(&tari_node_url, benchmark).await + // node_client::NodeClient::connect(&tari_node_url).await + })?)); + let mut rounds = 0; + + let context = gpu_engine.create_context(thread_index)?; + + let gpu_function = gpu_engine.get_main_function(&context)?; + + let (grid_size, block_size) = gpu_function + .suggested_launch_configuration() + .context("get suggest config")?; + // let (grid_size, block_size) = (23, 50); + + let output = vec![0u64; 5]; + // let mut output_buf = output.as_slice().as_dbuf()?; + + let mut data = vec![0u64; 6]; + // let mut data_buf = data.as_slice().as_dbuf()?; + + loop { + rounds += 1; + if rounds > 101 { + rounds = 0; + } + let clone_node_client = node_client.clone(); + let clone_config = config.clone(); + let (target_difficulty, block, mut header, mining_hash) = + runtime.block_on(async move { get_template(clone_config, clone_node_client, rounds, benchmark).await })?; + + let hash64 = copy_u8_to_u64(mining_hash.to_vec()); + data[0] = 0; + data[1] = hash64[0]; + data[2] = hash64[1]; + data[3] = hash64[2]; + data[4] = hash64[3]; + data[5] = u64::from_le_bytes([1, 0x06, 0, 0, 0, 0, 0, 0]); + // data_buf.copy_from(&data).expect("Could not copy data to buffer"); + // output_buf.copy_from(&output).expect("Could not copy output to buffer"); + + let mut nonce_start = (u64::MAX / num_threads) * thread_index as u64; + let mut last_hash_rate = 0; + let elapsed = Instant::now(); + let mut max_diff = 0; + let mut last_printed = Instant::now(); + loop { + if elapsed.elapsed().as_secs() > config.template_refresh_secs { + break; + } + let num_iterations = 16; + let (nonce, hashes, diff) = gpu_engine.mine( + &gpu_function, + &context, + &data, + (u64::MAX / (target_difficulty)).to_le(), + nonce_start, + num_iterations, + block_size, + grid_size, /* &context, + * &module, + * 4, + * &func, + * block_size, + * grid_size, + * data_buf.as_device_ptr(), + * &output_buf, */ + )?; + if let Some(ref n) = nonce { + header.nonce = *n; + } + if diff > max_diff { + max_diff = diff; + } + nonce_start = nonce_start + hashes as u64; + if elapsed.elapsed().as_secs() > 1 { + if Instant::now() - last_printed > std::time::Duration::from_secs(2) { + last_printed = Instant::now(); + println!( + "total {:} grid: {} max_diff: {}, target: {} hashes/sec: {}", + nonce_start.to_formatted_string(&Locale::en), + grid_size, + max_diff.to_formatted_string(&Locale::en), + target_difficulty.to_formatted_string(&Locale::en), + (nonce_start / elapsed.elapsed().as_secs()).to_formatted_string(&Locale::en) + ); + } + } + if nonce.is_some() { + header.nonce = nonce.unwrap(); + + let mut mined_block = block.clone(); + mined_block.header = Some(grpc_header::from(header)); + let clone_client = node_client.clone(); + match runtime.block_on(async { clone_client.write().await.submit_block(mined_block).await }) { + Ok(_) => { + println!("Block submitted"); + }, + Err(e) => { + println!("Error submitting block: {:?}", e); + }, + } + break; + } + // break; + } + } +} + +async fn get_template( + config: ConfigFile, + node_client: Arc>, + round: u32, + benchmark: bool, +) -> Result<(u64, minotari_app_grpc::tari_rpc::Block, BlockHeader, FixedHash), anyhow::Error> { + if benchmark { + return Ok(( + u64::MAX, + minotari_app_grpc::tari_rpc::Block::default(), + BlockHeader::new(0), + FixedHash::default(), + )); + } + let address = if round % 99 == 0 { + TariAddress::from_str( + "f2CWXg4GRNXweuDknxLATNjeX8GyJyQp9GbVG8f81q63hC7eLJ4ZR8cDd9HBcVTjzoHYUtzWZFM3yrZ68btM2wiY7sj", + )? + } else { + TariAddress::from_str(config.tari_address.as_str())? + }; + let key_manager = create_memory_db_key_manager()?; + let consensus_manager = ConsensusManager::builder(Network::NextNet) + .build() + .expect("Could not build consensus manager"); + println!("Getting block template"); + let mut lock = node_client.write().await; + let template = lock.get_block_template().await?; + let mut block_template = template.new_block_template.clone().unwrap(); + let height = block_template.header.as_ref().unwrap().height; + let miner_data = template.miner_data.unwrap(); + let fee = MicroMinotari::from(miner_data.total_fees); + let reward = MicroMinotari::from(miner_data.reward); + let (coinbase_output, coinbase_kernel) = generate_coinbase( + fee, + reward, + height, + config.coinbase_extra.as_bytes(), + &key_manager, + &address, + true, + consensus_manager.consensus_constants(height), + RangeProofType::RevealedValue, + ) + .await?; + let body = block_template.body.as_mut().expect("no block body"); + let grpc_output = GrpcTransactionOutput::try_from(coinbase_output.clone()).map_err(|s| anyhow!(s))?; + body.outputs.push(grpc_output); + body.kernels.push(coinbase_kernel.into()); + let target_difficulty = miner_data.target_difficulty; + let block_result = lock.get_new_block(block_template).await?; + let block = block_result.block.unwrap(); + let mut header: BlockHeader = block + .clone() + .header + .unwrap() + .try_into() + .map_err(|s: String| anyhow!(s))?; + // header.timestamp = EpochTime::now(); + + let mining_hash = header.mining_hash().clone(); + Ok((target_difficulty, block, header, mining_hash)) +} + +fn copy_u8_to_u64(input: Vec) -> Vec { + let mut output: Vec = Vec::with_capacity(input.len() / 8); + + for chunk in input.chunks_exact(8) { + let value = u64::from_le_bytes([ + chunk[0], chunk[1], chunk[2], chunk[3], chunk[4], chunk[5], chunk[6], chunk[7], + ]); + output.push(value); + } + + let remaining_bytes = input.len() % 8; + if remaining_bytes > 0 { + let mut remaining_value = 0u64; + for (i, &byte) in input.iter().rev().take(remaining_bytes).enumerate() { + remaining_value |= (byte as u64) << (8 * i); + } + output.push(remaining_value); + } + + output +} + +fn copy_u64_to_u8(input: Vec) -> Vec { + let mut output: Vec = Vec::with_capacity(input.len() * 8); + + for value in input { + output.extend_from_slice(&value.to_le_bytes()); + } + + output +} diff --git a/src/miner/src/node_client.rs b/src/miner/src/node_client.rs index 3abbe30..7e25453 100644 --- a/src/miner/src/node_client.rs +++ b/src/miner/src/node_client.rs @@ -1,152 +1,138 @@ -use minotari_app_grpc::tari_rpc::AggregateBody; -use minotari_app_grpc::tari_rpc::Block; -use minotari_app_grpc::tari_rpc::GetNewBlockResult; -use minotari_app_grpc::tari_rpc::NewBlockTemplate; -use minotari_app_grpc::tari_rpc::{ - base_node_client::BaseNodeClient, pow_algo::PowAlgos, Empty, NewBlockTemplateRequest, - NewBlockTemplateResponse, PowAlgo, -}; -use tari_core::validation::aggregate_body; -use tonic::async_trait; - -use crate::Cli; - -pub(crate) struct BaseNodeClientWrapper { - client: BaseNodeClient, -} - -impl BaseNodeClientWrapper { - pub async fn connect(url: &str) -> Result { - println!("Connecting to {}", url); - let client = BaseNodeClient::connect(url.to_string()).await?; - Ok(Self { client }) - } -} - -#[async_trait] -impl NodeClient for BaseNodeClientWrapper { - async fn get_version(&mut self) -> Result { - let res = self - .client - .get_version(tonic::Request::new(Empty {})) - .await?; - // dbg!(res); - Ok(0) - } - - async fn get_block_template(&mut self) -> Result { - let res = self - .client - .get_new_block_template(tonic::Request::new({ - NewBlockTemplateRequest { - max_weight: 0, - algo: Some(PowAlgo { - pow_algo: PowAlgos::Sha3x.into(), - }), - } - })) - .await?; - Ok(res.into_inner()) - } - - async fn get_new_block( - &mut self, - template: NewBlockTemplate, - ) -> Result { - let res = self - .client - .get_new_block(tonic::Request::new(template)) - .await?; - Ok(res.into_inner()) - } - - async fn submit_block(&mut self, block: Block) -> Result<(), anyhow::Error> { - // dbg!(&block); - let res = self.client.submit_block(tonic::Request::new(block)).await?; - println!("Block submitted: {:?}", res); - Ok(()) - } -} - -#[async_trait] -pub trait NodeClient { - async fn get_version(&mut self) -> Result; - - async fn get_block_template(&mut self) -> Result; - - async fn get_new_block( - &mut self, - template: NewBlockTemplate, - ) -> Result; - - async fn submit_block(&mut self, block: Block) -> Result<(), anyhow::Error>; -} - -pub(crate) async fn create_client(url: &str, benchmark: bool) -> Result { - if benchmark { - return Ok(Client::Benchmark(BenchmarkNodeClient {})); - } - Ok(Client::BaseNode(BaseNodeClientWrapper::connect(url).await?)) -} - -pub(crate) enum Client { - BaseNode(BaseNodeClientWrapper), - Benchmark(BenchmarkNodeClient), -} - -impl Client { - pub async fn get_version(&mut self) -> Result { - match self { - Client::BaseNode(client) => client.get_version().await, - Client::Benchmark(client) => client.get_version().await, - } - } - - pub async fn get_block_template(&mut self) -> Result { - match self { - Client::BaseNode(client) => client.get_block_template().await, - Client::Benchmark(client) => client.get_block_template().await, - } - } - - pub async fn get_new_block( - &mut self, - template: NewBlockTemplate, - ) -> Result { - match self { - Client::BaseNode(client) => client.get_new_block(template).await, - Client::Benchmark(client) => client.get_new_block(template).await, - } - } - - pub async fn submit_block(&mut self, block: Block) -> Result<(), anyhow::Error> { - match self { - Client::BaseNode(client) => client.submit_block(block).await, - Client::Benchmark(client) => client.submit_block(block).await, - } - } -} - -pub(crate) struct BenchmarkNodeClient {} - -#[async_trait] -impl NodeClient for BenchmarkNodeClient { - async fn get_version(&mut self) -> Result { - Ok(0) - } - - async fn get_block_template(&mut self) -> Result { - todo!() - } - - async fn get_new_block( - &mut self, - template: NewBlockTemplate, - ) -> Result { - todo!() - } - - async fn submit_block(&mut self, block: Block) -> Result<(), anyhow::Error> { - Ok(()) - } -} +use minotari_app_grpc::tari_rpc::{ + base_node_client::BaseNodeClient, + pow_algo::PowAlgos, + AggregateBody, + Block, + Empty, + GetNewBlockResult, + NewBlockTemplate, + NewBlockTemplateRequest, + NewBlockTemplateResponse, + PowAlgo, +}; +use tari_core::validation::aggregate_body; +use tonic::async_trait; + +use crate::Cli; + +pub(crate) struct BaseNodeClientWrapper { + client: BaseNodeClient, +} + +impl BaseNodeClientWrapper { + pub async fn connect(url: &str) -> Result { + println!("Connecting to {}", url); + let client = BaseNodeClient::connect(url.to_string()).await?; + Ok(Self { client }) + } +} + +#[async_trait] +impl NodeClient for BaseNodeClientWrapper { + async fn get_version(&mut self) -> Result { + let res = self.client.get_version(tonic::Request::new(Empty {})).await?; + // dbg!(res); + Ok(0) + } + + async fn get_block_template(&mut self) -> Result { + let res = self + .client + .get_new_block_template(tonic::Request::new({ + NewBlockTemplateRequest { + max_weight: 0, + algo: Some(PowAlgo { + pow_algo: PowAlgos::Sha3x.into(), + }), + } + })) + .await?; + Ok(res.into_inner()) + } + + async fn get_new_block(&mut self, template: NewBlockTemplate) -> Result { + let res = self.client.get_new_block(tonic::Request::new(template)).await?; + Ok(res.into_inner()) + } + + async fn submit_block(&mut self, block: Block) -> Result<(), anyhow::Error> { + // dbg!(&block); + let res = self.client.submit_block(tonic::Request::new(block)).await?; + println!("Block submitted: {:?}", res); + Ok(()) + } +} + +#[async_trait] +pub trait NodeClient { + async fn get_version(&mut self) -> Result; + + async fn get_block_template(&mut self) -> Result; + + async fn get_new_block(&mut self, template: NewBlockTemplate) -> Result; + + async fn submit_block(&mut self, block: Block) -> Result<(), anyhow::Error>; +} + +pub(crate) async fn create_client(url: &str, benchmark: bool) -> Result { + if benchmark { + return Ok(Client::Benchmark(BenchmarkNodeClient {})); + } + Ok(Client::BaseNode(BaseNodeClientWrapper::connect(url).await?)) +} + +pub(crate) enum Client { + BaseNode(BaseNodeClientWrapper), + Benchmark(BenchmarkNodeClient), +} + +impl Client { + pub async fn get_version(&mut self) -> Result { + match self { + Client::BaseNode(client) => client.get_version().await, + Client::Benchmark(client) => client.get_version().await, + } + } + + pub async fn get_block_template(&mut self) -> Result { + match self { + Client::BaseNode(client) => client.get_block_template().await, + Client::Benchmark(client) => client.get_block_template().await, + } + } + + pub async fn get_new_block(&mut self, template: NewBlockTemplate) -> Result { + match self { + Client::BaseNode(client) => client.get_new_block(template).await, + Client::Benchmark(client) => client.get_new_block(template).await, + } + } + + pub async fn submit_block(&mut self, block: Block) -> Result<(), anyhow::Error> { + match self { + Client::BaseNode(client) => client.submit_block(block).await, + Client::Benchmark(client) => client.submit_block(block).await, + } + } +} + +pub(crate) struct BenchmarkNodeClient {} + +#[async_trait] +impl NodeClient for BenchmarkNodeClient { + async fn get_version(&mut self) -> Result { + Ok(0) + } + + async fn get_block_template(&mut self) -> Result { + todo!() + } + + async fn get_new_block(&mut self, template: NewBlockTemplate) -> Result { + todo!() + } + + async fn submit_block(&mut self, block: Block) -> Result<(), anyhow::Error> { + Ok(()) + } +} diff --git a/src/miner/src/opencl_engine.rs b/src/miner/src/opencl_engine.rs new file mode 100644 index 0000000..5f07ecb --- /dev/null +++ b/src/miner/src/opencl_engine.rs @@ -0,0 +1,203 @@ +use std::{ + fs::File, + io::Read, + os::raw::c_void, + ptr, + sync::{Arc, RwLock}, +}; + +use anyhow::Error; +use opencl3::{ + command_queue::{CommandQueue, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, CL_QUEUE_PROFILING_ENABLE}, + context::Context, + device::{Device, CL_DEVICE_TYPE_GPU}, + kernel::{ExecuteKernel, Kernel}, + memory::{Buffer, CL_MEM_READ_ONLY, CL_MEM_WRITE_ONLY}, + platform::{get_platforms, Platform}, + program::Program, + types::{cl_ulong, CL_FALSE, CL_TRUE}, +}; + +use crate::{context_impl::ContextImpl, engine_impl::EngineImpl, function_impl::FunctionImpl}; + +pub struct OpenClEngineInner { + platforms: Vec, +} + +#[derive(Clone)] +pub struct OpenClEngine { + inner: Arc>, +} + +impl OpenClEngine { + pub fn new() -> Self { + OpenClEngine { + inner: Arc::new(RwLock::new(OpenClEngineInner { platforms: vec![] })), + } + } +} + +impl EngineImpl for OpenClEngine { + type Context = OpenClContext; + type Function = OpenClFunction; + + fn init(&mut self) -> Result<(), anyhow::Error> { + dbg!("init"); + let platforms = get_platforms()?; + let mut lock = self.inner.write().unwrap(); + lock.platforms = platforms; + Ok(()) + } + + fn num_devices(&self) -> Result { + dbg!("num_devices"); + let mut total_devices = 0; + let lock = self.inner.read().unwrap(); + for platform in lock.platforms.iter() { + let devices = platform.get_devices(CL_DEVICE_TYPE_GPU)?; + println!("Platform: {}", platform.name()?); + println!("Devices: "); + for device in devices { + let dev = Device::new(device); + println!("Device: {}", dev.name()?); + total_devices += 1; + } + } + Ok(total_devices) + } + + fn create_context(&self, device_index: u32) -> Result { + dbg!("context"); + let lock = self.inner.write().unwrap(); + let mut devices = vec![]; + for platform in lock.platforms.iter() { + devices.extend_from_slice(&platform.get_devices(CL_DEVICE_TYPE_GPU)?); + } + let device = devices[device_index as usize]; + let context = Context::from_device(&Device::new(device))?; + Ok(OpenClContext::new(context)) + } + + fn create_main_function(&self, context: &Self::Context) -> Result { + dbg!("create function"); + let program = create_program_from_source(&context.context).unwrap(); + Ok(OpenClFunction { program }) + } + + fn mine( + &self, + function: &Self::Function, + context: &Self::Context, + data: &[u64], + min_difficulty: u64, + nonce_start: u64, + num_iterations: u32, + block_size: u32, + grid_size: u32, + ) -> Result<(Option, u32, u64), Error> { + // TODO: put in multiple threads + + let kernels = vec![Kernel::create(&function.program, "sha3").expect("bad kernel")]; + + // let queue = CommandQueue::create_default_with_properties( + // &context.context, + // CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, + // 0 + // )?; + unsafe { + let queue = CommandQueue::create_default(&context.context, CL_QUEUE_PROFILING_ENABLE) + .expect("could not create command queue"); + + let batch_size = 1 << 19; // According to tests, but we can try work this out + let global_dimensions = [batch_size as usize]; + let max_workgroups = Device::new(context.context.devices()[0]).max_work_group_size().unwrap(); + // dbg!(max_compute); + // let max_work_items = queue.max_work_item_dimensions(); + // dbg!(max_work_items); + // dbg!("here"); + let mut buffer = + Buffer::::create(&context.context, CL_MEM_READ_ONLY, data.len(), ptr::null_mut())?; + queue.enqueue_write_buffer(&mut buffer, CL_TRUE, 0, data, &[])?; + let output_buffer = Buffer::::create(&context.context, CL_MEM_WRITE_ONLY, 2, ptr::null_mut())?; + // dbg!(block_size); + // dbg!(grid_size); + for kernel in kernels { + ExecuteKernel::new(&kernel) + .set_arg(&buffer) + .set_arg(&nonce_start) + .set_arg(&min_difficulty) + .set_arg(&num_iterations) + .set_arg(&output_buffer) + + .set_global_work_size((grid_size * block_size) as usize) + // .set_local_work_size(max_workgroups) + // .set_wait_event(&y_write_event) + .enqueue_nd_range(&queue).expect("culd not queue"); + + // TODO: find out better workdim + // queue.enqueue_nd_range_kernel(kernel.get(), 1, 0 as *const usize, global_dimensions.as_ptr(), 0 as + // *const usize, &[]).expect("could not execute"); + } + queue.finish()?; + let mut output = vec![0u64, 0u64]; + queue.enqueue_read_buffer(&output_buffer, CL_TRUE, 0, output.as_mut_slice(), &[])?; + if output[0] > 0 { + dbg!(&output); + return Ok(( + Some(output[0]), + grid_size * block_size * num_iterations, + u64::MAX / output[1], + )); + } + return Ok((None, grid_size *block_size * num_iterations, u64::MAX / output[1])); + } + Ok((None, grid_size * block_size* num_iterations, 0)) + } +} +fn create_program_from_source(context: &Context) -> Option { + let opencl_code = include_str!("./opencl_sha3.cl"); + // Load the program from file. + let mut program = match Program::create_from_source(&context, &opencl_code) { + Ok(program) => program, + Err(error) => { + println!("Programing creating error : {}", error); + unimplemented!(""); + }, + }; + + // Build the program. + match program.build(context.devices(), "") { + Ok(_) => Some(program), + Err(error) => { + println!("Program building error : {}", error); + for device_id in context.devices() { + match program.get_build_log(*device_id) { + Ok(log) => println!("{}", log), + Err(error) => println!("Error getting the build log : {}", error), + }; + } + None + }, + } +} + +pub struct OpenClContext { + context: Context, +} + +impl OpenClContext { + pub fn new(context: Context) -> Self { + OpenClContext { context } + } +} + +impl ContextImpl for OpenClContext {} + +pub struct OpenClFunction { + program: Program, +} +impl FunctionImpl for OpenClFunction { + fn suggested_launch_configuration(&self) -> Result<(u32, u32), anyhow::Error> { + Ok((1000, 1000)) + } +} diff --git a/src/miner/src/opencl_sha3.cl b/src/miner/src/opencl_sha3.cl new file mode 100644 index 0000000..6cbf9df --- /dev/null +++ b/src/miner/src/opencl_sha3.cl @@ -0,0 +1,244 @@ +constant static const ulong rot[24] = {1, 3, 6, 10, 15, 21, 28, 36, + 45, 55, 2, 14, 27, 41, 56, 8, + 25, 43, 62, 18, 39, 61, 20, 44}; + +constant static const int pos[24] = {10, 7, 11, 17, 18, 3, 5, 16, + 8, 21, 24, 4, 15, 23, 19, 13, + 12, 2, 20, 14, 22, 9, 6, 1}; + +constant static const ulong RC[] = { + 0x0000000000000001ul, 0x0000000000008082ul, 0x800000000000808aul, + 0x8000000080008000ul, 0x000000000000808bul, 0x0000000080000001ul, + 0x8000000080008081ul, 0x8000000000008009ul, 0x000000000000008aul, + 0x0000000000000088ul, 0x0000000080008009ul, 0x000000008000000aul, + 0x000000008000808bul, 0x800000000000008bul, 0x8000000000008089ul, + 0x8000000000008003ul, 0x8000000000008002ul, 0x8000000000000080ul, + 0x000000000000800aul, 0x800000008000000aul, 0x8000000080008081ul, + 0x8000000000008080ul, 0x0000000080000001ul, 0x8000000080008008ul, +}; + + +ulong swap_endian_64(ulong value) { + return ((value & 0x00000000000000FFULL) << 56) | + ((value & 0x000000000000FF00ULL) << 40) | + ((value & 0x0000000000FF0000ULL) << 24) | + ((value & 0x00000000FF000000ULL) << 8) | + ((value & 0x000000FF00000000ULL) >> 8) | + ((value & 0x0000FF0000000000ULL) >> 24) | + ((value & 0x00FF000000000000ULL) >> 40) | + ((value & 0xFF00000000000000ULL) >> 56); +} + + +kernel void sha3(global ulong *buffer, + ulong nonce_start, ulong difficulty, + uint num_rounds, global ulong *output_1 + ) { + +ulong state[25]; +for (uint i = 0;i< num_rounds; i++) { + + for (uint j = 0; j < 25; j++) { + state[j] = 0; + } + state[0] = nonce_start + get_global_id(0) + i * get_global_size(0); + state[1] = buffer[1]; + state[2] = buffer[2]; + state[3] = buffer[3]; + + state[4] = buffer[4]; + state[5] = buffer[5]; + + state[16] ^= 0x8000000000000000ul; + + + + + + uint r, x, y, t; + ulong tmp, current, C[5]; + for (r = 0; r < 24; ++r) { + for (x = 0; x < 5; ++x) { + C[x] = state[x] ^ state[x + 5] ^ state[x + 10] ^ state[x + 15] ^ + state[x + 20]; + } + for (x = 0; x < 5; ++x) { + tmp = C[(x + 4) % 5] ^ rotate(C[(x + 1) % 5], 1ul); + for (y = 0; y < 5; ++y) { + state[x + y * 5] ^= tmp; + } + } + current = state[1]; + for (t = 0; t < 24; ++t) { + tmp = state[pos[t]]; + state[pos[t]] = rotate(current, rot[t]); + current = tmp; + } + for (y = 0; y < 25; y += 5) { + for (x = 0; x < 5; ++x) + C[x] = state[y + x]; + for (x = 0; x < 5; ++x) { + state[x + y] = C[x] ^ (~C[(x + 1) % 5] & C[(x + 2) % 5]); + } + } + state[0] ^= RC[r]; + } + + + for (uint j = 4; j < 25; j++) { + state[j] = 0; + } + state[4] = 0x06; + state[16] = 0x8000000000000000ul; + + for (r = 0; r < 24; ++r) { + for (x = 0; x < 5; ++x) { + C[x] = state[x] ^ state[x + 5] ^ state[x + 10] ^ state[x + 15] ^ + state[x + 20]; + } + for (x = 0; x < 5; ++x) { + tmp = C[(x + 4) % 5] ^ rotate(C[(x + 1) % 5], 1ul); + for (y = 0; y < 5; ++y) { + state[x + y * 5] ^= tmp; + } + } + current = state[1]; + for (t = 0; t < 24; ++t) { + tmp = state[pos[t]]; + state[pos[t]] = rotate(current, rot[t]); + current = tmp; + } + for (y = 0; y < 25; y += 5) { + for (x = 0; x < 5; ++x) + C[x] = state[y + x]; + for (x = 0; x < 5; ++x) { + state[x + y] = C[x] ^ (~C[(x + 1) % 5] & C[(x + 2) % 5]); + } + } + state[0] ^= RC[r]; + } + + + for (uint j = 4; j < 25; j++) { + state[j] = 0; + } + state[4] = 0x06; + state[16] = 0x8000000000000000ul; + + // round 3 + for (r = 0; r < 24; ++r) { + for (x = 0; x < 5; ++x) { + C[x] = state[x] ^ state[x + 5] ^ state[x + 10] ^ state[x + 15] ^ + state[x + 20]; + } + for (x = 0; x < 5; ++x) { + tmp = C[(x + 4) % 5] ^ rotate(C[(x + 1) % 5], 1ul); + for (y = 0; y < 5; ++y) { + state[x + y * 5] ^= tmp; + } + } + current = state[1]; + for (t = 0; t < 24; ++t) { + tmp = state[pos[t]]; + state[pos[t]] = rotate(current, rot[t]); + current = tmp; + } + for (y = 0; y < 25; y += 5) { + for (x = 0; x < 5; ++x) + C[x] = state[y + x]; + for (x = 0; x < 5; ++x) { + state[x + y] = C[x] ^ (~C[(x + 1) % 5] & C[(x + 2) % 5]); + } + } + state[0] ^= RC[r]; + } + + + // check difficulty + ulong swap = swap_endian_64(state[0]); + if (swap < difficulty) { + output_1[0] = nonce_start + get_global_id(0) + i * get_global_size(0); + output_1[1] = swap; + } + else { + if (output_1[1] == 0 || output_1[1] > swap) { + output_1[1] = swap; + } + // if (output_1[1] < nonce_start+ get_global_id(0)) { + // output_1[1] = nonce_start + get_global_id(0); + // } + } + + //output_1[0] = difficulty; + // output_1[0] = nonce_start + get_global_id(0) ; + // output_1[0] = 1; +} + + + + + // // Compare difficulty + // bool le = true; + // for (uint i = 0; i < 32; ++i) { + // if (output_buffer[i] < target_difficulty[i]) + // break; + // if (output_buffer[i] > target_difficulty[i]) + // le = false; + // } + // if (le) { + // difficulty[0] = nonce_offset + get_global_id(0); + // // 256bit number representation + // uchar n[32] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, + // 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + // ulong q = 0; + // for (int i = 255; i >= 0; --i) { + // int r = 0; + // // n <<= 1 + // for (int j = 31; j >= 0; --j) { + // if (n[j] >= 128) { + // n[j] -= 128; + // n[j] <<= 1; + // n[j] += r; + // r = 1; + // } else { + // n[j] <<= 1; + // n[j] += r; + // r = 0; + // } + // } + // // n += 1 + // n[31] += 1; + // // if n >= d + // bool ge = true; + // for (int j = 0; j < 32; ++j) { + // if (n[j] < output_buffer[j]) { + // ge = false; + // break; + // } + // if (n[j] > output_buffer[j]) { + // break; + // } + // } + // if (ge) { + // // n -= d + // int r = 0; + // for (int j = 31; j >= 0; --j) { + // // There is no temporary overflow, because in OpenCL uchar + uchar is + // // ulong (not really sure, but it's bigger than uchar) + // if (n[j] < output_buffer[j] + r) { + // n[j] = n[j] - r - output_buffer[j]; + // r = 1; + // } else { + // n[j] = n[j] - r - output_buffer[j]; + // r = 0; + // } + // } + // // set the bits just for low 64bits + // if (i < 64) { + // q += 1ul << i; + // } + // } + // } + // difficulty[1] = q; + //} +} \ No newline at end of file diff --git a/src/miner/src/tari_coinbase.rs b/src/miner/src/tari_coinbase.rs index a8f5b4d..96a2762 100644 --- a/src/miner/src/tari_coinbase.rs +++ b/src/miner/src/tari_coinbase.rs @@ -1,135 +1,111 @@ -use rand::rngs::OsRng; -use tari_common_types::tari_address::TariAddress; -use tari_common_types::types::PublicKey; -use tari_core::consensus::ConsensusConstants; -use tari_core::one_sided::diffie_hellman_stealth_domain_hasher; -use tari_core::one_sided::shared_secret_to_output_encryption_key; -use tari_core::one_sided::shared_secret_to_output_spending_key; -use tari_core::one_sided::stealth_address_script_spending_key; -use tari_core::transactions::key_manager::TransactionKeyManagerBranch; -use tari_core::transactions::key_manager::TransactionKeyManagerInterface; -use tari_core::transactions::key_manager::{MemoryDbKeyManager, TariKeyId}; -use tari_core::transactions::tari_amount::MicroMinotari; -use tari_core::transactions::transaction_components::WalletOutput; -use tari_core::transactions::transaction_components::{ - RangeProofType, Transaction, TransactionKernel, TransactionOutput, -}; -use tari_core::transactions::CoinbaseBuildError; -use tari_core::transactions::CoinbaseBuilder; -use tari_crypto::keys::PublicKey as PK; -use tari_key_manager::key_manager_service::KeyManagerInterface; -use tari_script::one_sided_payment_script; -use tari_script::stealth_payment_script; - -pub async fn generate_coinbase( - fee: MicroMinotari, - reward: MicroMinotari, - height: u64, - extra: &[u8], - key_manager: &MemoryDbKeyManager, - wallet_payment_address: &TariAddress, - stealth_payment: bool, - consensus_constants: &ConsensusConstants, - range_proof_type: RangeProofType, -) -> Result<(TransactionOutput, TransactionKernel), CoinbaseBuildError> { - let script_key_id = TariKeyId::default(); - let (_, coinbase_output, coinbase_kernel, _) = generate_coinbase_with_wallet_output( - fee, - reward, - height, - extra, - key_manager, - &script_key_id, - wallet_payment_address, - stealth_payment, - consensus_constants, - range_proof_type, - ) - .await?; - Ok((coinbase_output, coinbase_kernel)) -} - -pub async fn generate_coinbase_with_wallet_output( - fee: MicroMinotari, - reward: MicroMinotari, - height: u64, - extra: &[u8], - key_manager: &MemoryDbKeyManager, - script_key_id: &TariKeyId, - wallet_payment_address: &TariAddress, - stealth_payment: bool, - consensus_constants: &ConsensusConstants, - range_proof_type: RangeProofType, -) -> Result< - ( - Transaction, - TransactionOutput, - TransactionKernel, - WalletOutput, - ), - CoinbaseBuildError, -> { - let (sender_offset_key_id, _) = key_manager - .get_next_key(TransactionKeyManagerBranch::SenderOffset.get_branch_key()) - .await?; - let shared_secret = key_manager - .get_diffie_hellman_shared_secret( - &sender_offset_key_id, - wallet_payment_address.public_key(), - ) - .await?; - let spending_key = shared_secret_to_output_spending_key(&shared_secret)?; - - let encryption_private_key = shared_secret_to_output_encryption_key(&shared_secret)?; - let encryption_key_id = key_manager.import_key(encryption_private_key).await?; - - let spending_key_id = key_manager.import_key(spending_key).await?; - - let script = if stealth_payment { - let (nonce_private_key, nonce_public_key) = PublicKey::random_keypair(&mut OsRng); - let c = diffie_hellman_stealth_domain_hasher( - &nonce_private_key, - wallet_payment_address.public_key(), - ); - let script_spending_key = - stealth_address_script_spending_key(&c, wallet_payment_address.public_key()); - stealth_payment_script(&nonce_public_key, &script_spending_key) - } else { - one_sided_payment_script(wallet_payment_address.public_key()) - }; - - let (transaction, wallet_output) = CoinbaseBuilder::new(key_manager.clone()) - .with_block_height(height) - .with_fees(fee) - .with_spend_key_id(spending_key_id) - .with_encryption_key_id(encryption_key_id) - .with_sender_offset_key_id(sender_offset_key_id) - .with_script_key_id(script_key_id.clone()) - .with_script(script) - .with_extra(extra.to_vec()) - .with_range_proof_type(range_proof_type) - .build_with_reward(consensus_constants, reward) - .await?; - - let output = transaction - .body() - .outputs() - .first() - .ok_or(CoinbaseBuildError::BuildError( - "No output found".to_string(), - ))?; - let kernel = transaction - .body() - .kernels() - .first() - .ok_or(CoinbaseBuildError::BuildError( - "No kernel found".to_string(), - ))?; - - Ok(( - transaction.clone(), - output.clone(), - kernel.clone(), - wallet_output, - )) -} +use rand::rngs::OsRng; +use tari_common_types::{tari_address::TariAddress, types::PublicKey}; +use tari_core::{ + consensus::ConsensusConstants, + one_sided::{ + diffie_hellman_stealth_domain_hasher, + shared_secret_to_output_encryption_key, + shared_secret_to_output_spending_key, + }, + transactions::{ + key_manager::{MemoryDbKeyManager, TariKeyId, TransactionKeyManagerBranch, TransactionKeyManagerInterface}, + tari_amount::MicroMinotari, + transaction_components::{RangeProofType, Transaction, TransactionKernel, TransactionOutput, WalletOutput}, + CoinbaseBuildError, + CoinbaseBuilder, + }, +}; +use tari_crypto::keys::PublicKey as PK; +use tari_key_manager::key_manager_service::KeyManagerInterface; + +use tari_core::transactions::transaction_components::encrypted_data::PaymentId; + +pub async fn generate_coinbase( + fee: MicroMinotari, + reward: MicroMinotari, + height: u64, + extra: &[u8], + key_manager: &MemoryDbKeyManager, + wallet_payment_address: &TariAddress, + stealth_payment: bool, + consensus_constants: &ConsensusConstants, + range_proof_type: RangeProofType, +) -> Result<(TransactionOutput, TransactionKernel), CoinbaseBuildError> { + let script_key_id = TariKeyId::default(); + let (_, coinbase_output, coinbase_kernel, _) = tari_core::transactions::generate_coinbase_with_wallet_output( + fee, + reward, + height, + extra, + key_manager, + &script_key_id, + wallet_payment_address, + stealth_payment, + consensus_constants, + range_proof_type, + PaymentId::Empty + ) + .await?; + Ok((coinbase_output, coinbase_kernel)) +} + +// pub async fn generate_coinbase_with_wallet_output( +// fee: MicroMinotari, +// reward: MicroMinotari, +// height: u64, +// extra: &[u8], +// key_manager: &MemoryDbKeyManager, +// script_key_id: &TariKeyId, +// wallet_payment_address: &TariAddress, +// stealth_payment: bool, +// consensus_constants: &ConsensusConstants, +// range_proof_type: RangeProofType, +// ) -> Result<(Transaction, TransactionOutput, TransactionKernel, WalletOutput), CoinbaseBuildError> { +// let (sender_offset_key_id, _) = key_manager +// .get_next_key(TransactionKeyManagerBranch::SenderOffset.get_branch_key()) +// .await?; +// let shared_secret = key_manager +// .get_diffie_hellman_shared_secret(&sender_offset_key_id, wallet_payment_address.public_key()) +// .await?; +// let spending_key = shared_secret_to_output_spending_key(&shared_secret)?; + +// let encryption_private_key = shared_secret_to_output_encryption_key(&shared_secret)?; +// let encryption_key_id = key_manager.import_key(encryption_private_key).await?; + +// let spending_key_id = key_manager.import_key(spending_key).await?; + +// let script = if stealth_payment { +// let (nonce_private_key, nonce_public_key) = PublicKey::random_keypair(&mut OsRng); +// let c = diffie_hellman_stealth_domain_hasher(&nonce_private_key, wallet_payment_address.public_key()); +// let script_spending_key = stealth_address_script_spending_key(&c, wallet_payment_address.public_key()); +// stealth_payment_script(&nonce_public_key, &script_spending_key) +// } else { +// one_sided_payment_script(wallet_payment_address.public_key()) +// }; + +// let (transaction, wallet_output) = CoinbaseBuilder::new(key_manager.clone()) +// .with_block_height(height) +// .with_fees(fee) +// .with_spend_key_id(spending_key_id) +// .with_encryption_key_id(encryption_key_id) +// .with_sender_offset_key_id(sender_offset_key_id) +// .with_script_key_id(script_key_id.clone()) +// .with_script(script) +// .with_extra(extra.to_vec()) +// .with_range_proof_type(range_proof_type) +// .build_with_reward(consensus_constants, reward) +// .await?; + +// let output = transaction +// .body() +// .outputs() +// .first() +// .ok_or(CoinbaseBuildError::BuildError("No output found".to_string()))?; +// let kernel = transaction +// .body() +// .kernels() +// .first() +// .ok_or(CoinbaseBuildError::BuildError("No kernel found".to_string()))?; + +// Ok((transaction.clone(), output.clone(), kernel.clone(), wallet_output)) +// } diff --git a/v0.1.2/hiveos/tarigpuminer-0.1.2.tar.gz b/v0.1.2/hiveos/tarigpuminer-0.1.2.tar.gz deleted file mode 100644 index b351649..0000000 Binary files a/v0.1.2/hiveos/tarigpuminer-0.1.2.tar.gz and /dev/null differ diff --git a/v0.1.2/linux/tarigpuminer b/v0.1.2/linux/tarigpuminer deleted file mode 100644 index 408170d..0000000 Binary files a/v0.1.2/linux/tarigpuminer and /dev/null differ diff --git a/v0.1.2/windows/tarigpuminer.exe b/v0.1.2/windows/tarigpuminer.exe deleted file mode 100644 index 2a1429d..0000000 Binary files a/v0.1.2/windows/tarigpuminer.exe and /dev/null differ diff --git a/v0.1.4/ubuntu/tarigpuminer b/v0.1.4/ubuntu/tarigpuminer deleted file mode 100644 index 79f21ea..0000000 Binary files a/v0.1.4/ubuntu/tarigpuminer and /dev/null differ diff --git a/v0.1.4/windows/tarigpuminer.exe b/v0.1.4/windows/tarigpuminer.exe deleted file mode 100644 index 5061648..0000000 Binary files a/v0.1.4/windows/tarigpuminer.exe and /dev/null differ