29:I[49813,["/_next/static/chunks/3z--bwudlglfp.js","/_next/static/chunks/2u04xn-9yg2co.js","/_next/static/chunks/1b5-8ftwqgmtb.js","/_next/static/chunks/3nqo6bdos11l6.js","/_next/static/chunks/0fz0mlrjifmq7.js","/_next/static/chunks/2l43yyot7shmw.js","/_next/static/chunks/29v1ttach42xh.js","/_next/static/chunks/2ziic7alkvd8r.js","/_next/static/chunks/1hplpmn3-d0tx.js"],"CodeWorkspace"] 2a:T2257,/** * Copyright (c) 2011-2026 libbitcoin developers * * This file is part of libbitcoin. * * This program is free software: you can redistribute it and/or modify * it under the terms of the GNU Affero General Public License as published by * the Free Software Foundation, either version 3 of the License, or * (at your option) any later version. * * This program is distributed in the hope that it will be useful, * but WITHOUT ANY WARRANTY; without even the implied warranty of * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the * GNU Affero General Public License for more details. * * You should have received a copy of the GNU Affero General Public License * along with this program. If not, see . */ #include #include "kernels.hpp" namespace libbitcoin { namespace system { namespace secp256k1 { class verifier : public algorithm { public: using algorithm::slice_count; using algorithm::table_words; using algorithm::comb_part_count; /// Entries and words of a comb table part. static constexpr size_t comb_part_size = comb_part_windows * comb_size; static constexpr size_t comb_part_words = comb_part_size * comb_words; /// Point window bits, uniform across the threads of a warp. static constexpr size_t window_bits = 4; static __device__ uint32_t row() NOEXCEPT { return __nvvm_read_ptx_sreg_ctaid_x() * __nvvm_read_ptx_sreg_ntid_x() + __nvvm_read_ptx_sreg_tid_x(); } static __device__ uint8_t verify(const hash_digest& digest, const ec_compressed& key, const ec_signature& signature) NOEXCEPT { affine_t point{}; if (!decode(point, key)) return 0; scalar_t z{}; /* bool */ from_bytes(z, digest); return verify_ecdsa(point, z, array_cast>(signature), array_cast, sizeof(scalar_t)>( signature)) ? 1 : 0; } static __device__ uint8_t verify(const hash_digest& message, const ec_xonly& key, const ec_signature& signature) NOEXCEPT { constexpr auto size = array_count; const auto& r = array_cast(signature); const auto& s = array_cast(signature); const auto challenge = sha256::hash( tagged_midstate<"BIP0340/challenge">, r, key, message); return verify_schnorr(key, challenge, r, s) ? 1 : 0; } /// The output key prefix of the receiver for the point (k = 0), then of /// that key plus each label key [bip352]. static __device__ uint8_t scan(cuda::prefix* prefixes, const ec_compressed& encoded, const cuda::scan_arguments& arguments) NOEXCEPT { affine_t point{}; if (!decode(point, encoded)) return 0; // The device computes in variable time, a server-side optimization // that is opted out of by not compiling CUDA. jacobian_t sum{}; const scalars_t ks{ arguments.scan }; const auto uncomputed = multiply_windows(sum, {}, point, ks); if (f::any(uncomputed)) multiply_complete(sum, {}, point, arguments.scan); if (f::any(sum.infinity)) return 0; // The tagged hash of the shared point and k = 0, as four zero bytes. affine_t shared{}; ec_compressed key{}; data_array data{}; to_affine(shared, sum); to_bytes(key, shared); for (size_t byte{}; byte < key.size(); ++byte) data[byte] = key[byte]; scalar_t tweak{}; constexpr auto midstate = tagged_midstate<"BIP0352/SharedSecret">; const auto hash = sha256::hash(midstate, data); if (!from_bytes(tweak, hash) || is_zero_scalar(tweak)) return 0; uint64_t faults{}; sum = {}; sum.infinity = max_uint64; add_comb(sum, tweak, faults); add_point(sum, arguments.spend, faults); if (is_nonzero(faults)) multiply_complete(sum, tweak, arguments.spend, { 1 }); if (f::any(sum.infinity)) return 0; write(prefixes[0], sum); for (size_t index{}; index < arguments.label_count; ++index) { jacobian_t labeled{}; add_complete(labeled, sum, arguments.labels[index]); if (f::any(labeled.infinity)) return 0; write(prefixes[add1(index)], labeled); } return 1; } /// The point, hash * sum [bip352]. static __device__ uint8_t compute(ec_compressed& out, const ec_compressed& sum, const ec_secret& hash) NOEXCEPT { affine_t point{}; if (!decode(point, sum)) return 0; scalar_t k{}; if (!from_bytes(k, hash) || is_zero_scalar(k)) return 0; jacobian_t product{}; const scalars_t ks{ k }; const auto uncomputed = multiply_windows( product, {}, point, ks); if (f::any(uncomputed)) multiply_complete(product, {}, point, k); if (f::any(product.infinity)) return 0; affine_t normal{}; ec_compressed key{}; to_affine(normal, product); to_bytes(key, normal); for (size_t byte{}; byte < key.size(); ++byte) out[byte] = key[byte]; return 1; } private: // Key rows are unaligned (33 bytes), so x is decoded from a copy. static __device__ bool decode(affine_t& out, const ec_compressed& key) NOEXCEPT { const auto sign = key.front(); if (sign != ec_even_sign && sign != ec_odd_sign) return false; alignas(sizeof(uint64_t)) bytes_t x_bytes{}; for (size_t byte{}; byte < x_bytes.size(); ++byte) x_bytes[byte] = key[add1(byte)]; field_t x{}; const auto odd = sign == ec_odd_sign ? max_uint64 : 0_u64; return from_bytes(x, x_bytes) && is_nonzero(lift(out, x, odd)); } static __device__ void write(cuda::prefix& out, const jacobian_t& point) NOEXCEPT { affine_t normal{}; bytes_t x{}; to_affine(normal, point); to_bytes(x, normal.x); for (size_t byte{}; byte < out.size(); ++byte) out[byte] = x[byte]; } }; extern "C" { __device__ uint64_t generator_table[verifier::slice_count] [verifier::table_words]; __device__ uint64_t comb_table[verifier::comb_part_count] [verifier::comb_part_words]; } __device__ const std_array generator_slices { generator_table[0], generator_table[1], generator_table[2], generator_table[3], generator_table[4], generator_table[5], generator_table[6], generator_table[7], generator_table[8], generator_table[9], generator_table[10], generator_table[11], generator_table[12], generator_table[13], generator_table[14], generator_table[15] }; __device__ const std_array comb_parts { comb_table[0], comb_table[1], comb_table[2] }; } // namespace secp256k1 } // namespace system } // namespace libbitcoin using namespace libbitcoin::system::secp256k1; extern "C" __global__ void verify_ecdsa(cuda::ecdsa_arguments arguments) { const auto row = verifier::row(); if (row < arguments.count) arguments.results[row] = verifier::verify(arguments.digests[row], arguments.keys[row], arguments.signatures[row]); } extern "C" __global__ void verify_schnorr(cuda::schnorr_arguments arguments) { const auto row = verifier::row(); if (row < arguments.count) arguments.results[row] = verifier::verify(arguments.messages[row], arguments.keys[row], arguments.signatures[row]); } extern "C" __global__ void scan_silent(cuda::scan_arguments arguments) { const auto row = verifier::row(); if (row < arguments.count) arguments.valid[row] = verifier::scan( &arguments.prefixes[row * (arguments.label_count + 1u)], arguments.points[row], arguments); } extern "C" __global__ void compute_silent( cuda::silent_arguments arguments) { const auto row = verifier::row(); if (row < arguments.count) arguments.valid[row] = verifier::compute( arguments.points[row], arguments.sums[row], arguments.hashes[row]); } 12:["$","$L29",null,{"project":{"id":"cmuidb72x000001s1932ifphv","owner":"libbitcoin","repo":"libbitcoin-system","displayName":"libbitcoin/libbitcoin-system","defaultBranch":"master"},"commit":{"sha":"83147c18e3a80ab2949c2c8a7e3486950d7cfeff","message":"Merge pull request #2109 from evoskuil/master\n\nPerf/chain parse allocations (cherry-picked from PR2090)","authorLogin":"evoskuil","authorName":"Eric Voskuil","date":"2026-10-22T16:03:52.000Z","htmlUrl":"https://github.com/libbitcoin/libbitcoin-system/commit/83147c18e3a80ab2949c2c8a7e3486950d7cfeff","isHead":true,"headSha":"83147c18e3a80ab2949c2c8a7e3486950d7cfeff"},"gitRef":"master","path":"src/crypto/secp256k1/cuda/kernels.cu","target":{"kind":"file","file":{"path":"src/crypto/secp256k1/cuda/kernels.cu","size":8791,"content":"$2a","tooLarge":false,"binary":false,"htmlUrl":"https://github.com/libbitcoin/libbitcoin-system/blob/83147c18e3a80ab2949c2c8a7e3486950d7cfeff/src/crypto/secp256k1/cuda/kernels.cu","mutants":[],"mutantsAtOtherRevisions":0}},"treeCache":{"":[{"name":".github","path":".github","type":"dir","size":null},{"name":".vscode","path":".vscode","type":"dir","size":null},{"name":"builds","path":"builds","type":"dir","size":null},{"name":"doc","path":"doc","type":"dir","size":null},{"name":"examples","path":"examples","type":"dir","size":null},{"name":"include","path":"include","type":"dir","size":null},{"name":"src","path":"src","type":"dir","size":null},{"name":"test","path":"test","type":"dir","size":null},{"name":"tools","path":"tools","type":"dir","size":null},{"name":".gitignore","path":".gitignore","type":"file","size":1037},{"name":"COPYING","path":"COPYING","type":"file","size":33899},{"name":"README.md","path":"README.md","type":"file","size":9727}],"src":[{"name":"chain","path":"src/chain","type":"dir","size":null},{"name":"config","path":"src/config","type":"dir","size":null},{"name":"crypto","path":"src/crypto","type":"dir","size":null},{"name":"data","path":"src/data","type":"dir","size":null},{"name":"endian","path":"src/endian","type":"dir","size":null},{"name":"error","path":"src/error","type":"dir","size":null},{"name":"filter","path":"src/filter","type":"dir","size":null},{"name":"hash","path":"src/hash","type":"dir","size":null},{"name":"math","path":"src/math","type":"dir","size":null},{"name":"radix","path":"src/radix","type":"dir","size":null},{"name":"serial","path":"src/serial","type":"dir","size":null},{"name":"stream","path":"src/stream","type":"dir","size":null},{"name":"unicode","path":"src/unicode","type":"dir","size":null},{"name":"wallet","path":"src/wallet","type":"dir","size":null},{"name":"words","path":"src/words","type":"dir","size":null},{"name":"x509","path":"src/x509","type":"dir","size":null},{"name":"arena.cpp","path":"src/arena.cpp","type":"file","size":2137},{"name":"define.cpp","path":"src/define.cpp","type":"file","size":5690},{"name":"execution.cpp","path":"src/execution.cpp","type":"file","size":2110},{"name":"settings.cpp","path":"src/settings.cpp","type":"file","size":26844}],"src/crypto":[{"name":"secp256k1","path":"src/crypto/secp256k1","type":"dir","size":null},{"name":"aes_gcm.cpp","path":"src/crypto/aes_gcm.cpp","type":"file","size":1091},{"name":"chacha20_poly1305.cpp","path":"src/crypto/chacha20_poly1305.cpp","type":"file","size":4952},{"name":"chacha20.cpp","path":"src/crypto/chacha20.cpp","type":"file","size":14408},{"name":"der_parser.cpp","path":"src/crypto/der_parser.cpp","type":"file","size":7018},{"name":"fschacha20_poly1305.cpp","path":"src/crypto/fschacha20_poly1305.cpp","type":"file","size":2403},{"name":"fschacha20.cpp","path":"src/crypto/fschacha20.cpp","type":"file","size":1729},{"name":"maybe_random.cpp","path":"src/crypto/maybe_random.cpp","type":"file","size":2517},{"name":"poly1305.cpp","path":"src/crypto/poly1305.cpp","type":"file","size":18744},{"name":"salsa20.cpp","path":"src/crypto/salsa20.cpp","type":"file","size":6080},{"name":"x25519.cpp","path":"src/crypto/x25519.cpp","type":"file","size":8765},{"name":"xsalsa20_poly1305.cpp","path":"src/crypto/xsalsa20_poly1305.cpp","type":"file","size":2948}],"src/crypto/secp256k1":[{"name":"batch","path":"src/crypto/secp256k1/batch","type":"dir","size":null},{"name":"cuda","path":"src/crypto/secp256k1/cuda","type":"dir","size":null},{"name":"generator","path":"src/crypto/secp256k1/generator","type":"dir","size":null},{"name":"algorithm.cpp","path":"src/crypto/secp256k1/algorithm.cpp","type":"file","size":2256},{"name":"ec_context.cpp","path":"src/crypto/secp256k1/ec_context.cpp","type":"file","size":2234},{"name":"ec_context.hpp","path":"src/crypto/secp256k1/ec_context.hpp","type":"file","size":1854},{"name":"ecdsa.cpp","path":"src/crypto/secp256k1/ecdsa.cpp","type":"file","size":6838},{"name":"ellswift.cpp","path":"src/crypto/secp256k1/ellswift.cpp","type":"file","size":2417},{"name":"interface.cpp","path":"src/crypto/secp256k1/interface.cpp","type":"file","size":35523},{"name":"interface.hpp","path":"src/crypto/secp256k1/interface.hpp","type":"file","size":9136},{"name":"keys.cpp","path":"src/crypto/secp256k1/keys.cpp","type":"file","size":4263},{"name":"math.cpp","path":"src/crypto/secp256k1/math.cpp","type":"file","size":6449},{"name":"recover.cpp","path":"src/crypto/secp256k1/recover.cpp","type":"file","size":3462},{"name":"schnorr.cpp","path":"src/crypto/secp256k1/schnorr.cpp","type":"file","size":3384},{"name":"serialize.hpp","path":"src/crypto/secp256k1/serialize.hpp","type":"file","size":2008}],"src/crypto/secp256k1/cuda":[{"name":"context.hpp","path":"src/crypto/secp256k1/cuda/context.hpp","type":"file","size":26316},{"name":"driver.cpp","path":"src/crypto/secp256k1/cuda/driver.cpp","type":"file","size":3153},{"name":"driver.hpp","path":"src/crypto/secp256k1/cuda/driver.hpp","type":"file","size":2609},{"name":"kernels.cu","path":"src/crypto/secp256k1/cuda/kernels.cu","type":"file","size":8791},{"name":"kernels.hpp","path":"src/crypto/secp256k1/cuda/kernels.hpp","type":"file","size":2822},{"name":"layout.hpp","path":"src/crypto/secp256k1/cuda/layout.hpp","type":"file","size":5194},{"name":"library.hpp","path":"src/crypto/secp256k1/cuda/library.hpp","type":"file","size":6300},{"name":"ptx.cpp","path":"src/crypto/secp256k1/cuda/ptx.cpp","type":"file","size":1833830},{"name":"ptx.hpp","path":"src/crypto/secp256k1/cuda/ptx.hpp","type":"file","size":1331},{"name":"turns.hpp","path":"src/crypto/secp256k1/cuda/turns.hpp","type":"file","size":2945}]},"revisions":[{"commitSha":"9fa0f9a07f31c5cc9ebe5efa05e6d4a6d12ba183","branch":null,"commitMessage":"Test script pattern length guards and the sixteen key multisig builder boundary.","commitDate":"2026-09-29T18:05:00.000Z","mutantCount":129},{"commitSha":"98db5bfd7c7b65532e42309c7b71419d82ec5ea8","branch":null,"commitMessage":"Merge pull request #2017 from evoskuil/master","commitDate":"2026-09-28T17:24:16.000Z","mutantCount":230},{"commitSha":"827fb29d0fecbdf7a157e05b1bea59bc916787aa","branch":null,"commitMessage":"Test multisig sighash cache reset across distinct subscripts","commitDate":"2026-09-26T12:24:47.000Z","mutantCount":1},{"commitSha":"8e7a9fdcb4ca72bc5bfe9961815aaa9bc33b64a0","branch":null,"commitMessage":"Remove dead include.","commitDate":"2026-09-25T06:52:28.000Z","mutantCount":90}],"signedIn":false,"pullRequest":null}]