bend form A — cuda-secp256k1-batched-mul live on 3090
Per examples/cuda-fanout/plans/form-A-secp256k1-batched-mul.md.
Batched secp256k1 scalar*G via per-thread Jacobian double-and-add
plus per-thread Z-inversion to affine. Field arithmetic uses
FixedPaul/VanitySearch-Bitcrack GPUMath.h verbatim
(commit 66e6f9d, AGPL-3.0, vendored under
vendor/vanity-search-bitcrack/).
Wire: BSCP request / BSCR response distinct from BSHK/BCGB.
"BSCP" u32 op_id u32 n base_xy(64B BE) scalars(n*32B BE)
"BSCR" u32 status u32 n points(n*64B BE x||y)
Validation against coincurve on 3090-ai.foxhop.net byte-identical
across known-small (k in {1,2,3,7,0xdeadbeef,n-1,n,2^128-1})
& random sweeps at n in {32, 1k, 10k, 100k}.
Measured throughput on 3090, kernel-only:
n=10k 2.32 ms 4.31 Mkeys/s
n=100k 15.37 ms 6.51 Mkeys/s
End-to-end over warm TCP daemon from another host:
n=100k 400 ms 250 kkeys/s wall (PCIe + wire serialization
bound; kernel still <16 ms)
Speedup vs coincurve CPU single-thread (~44 kkeys/s host)
~148x at n=100k kernel-only. Day-2 work to add _ModInvGrouped
batched inversion should push toward FixedPaul's 6.9 Gkeys/s
published on 4090.
gpu-worker.lsp: maybe-register-daemon! for cuda-secp256k1-batched-mul,
handle-binary-secp branch in handle-one dispatching on BSCP magic.
Makefile: secp256k1-batch-mul / secp256k1-test / secp256k1-bench
targets. Test harness ships with coincurve preferred, falls back to
python-ecdsa or pure-Python double-and-add for the host oracle.
This commit is contained in:
parent
277250ecfa
commit
766178818c
7 changed files with 2460 additions and 1 deletions
|
|
@ -16,11 +16,24 @@ NVCC ?= nvcc
|
||||||
NVCCFLAGS ?= -O3 -arch=sm_86 -Xcompiler="-O3 -Wall -Wextra"
|
NVCCFLAGS ?= -O3 -arch=sm_86 -Xcompiler="-O3 -Wall -Wextra"
|
||||||
NVCC_PATH ?= $(shell which nvcc 2>/dev/null || echo /usr/local/cuda/bin/nvcc)
|
NVCC_PATH ?= $(shell which nvcc 2>/dev/null || echo /usr/local/cuda/bin/nvcc)
|
||||||
|
|
||||||
all: shake256-fanout cgbn-batch-worker
|
all: shake256-fanout cgbn-batch-worker secp256k1-batch-mul
|
||||||
|
|
||||||
shake256-fanout: shake256-fanout.cu
|
shake256-fanout: shake256-fanout.cu
|
||||||
$(NVCC) $(NVCCFLAGS) -o $@ $<
|
$(NVCC) $(NVCCFLAGS) -o $@ $<
|
||||||
|
|
||||||
|
# bend form A — secp256k1 batched scalar*G via VanitySearch-Bitcrack
|
||||||
|
# vendored GPUMath.h (AGPL-3.0). Per-thread Jacobian double-and-add,
|
||||||
|
# per-thread inversion to affine. See vendor/vanity-search-bitcrack/NOTICE.
|
||||||
|
NVCCFLAGS_SECP ?= -O3 -arch=sm_86 -Xcompiler="-O3 -Wall"
|
||||||
|
secp256k1-batch-mul: secp256k1-batch-mul.cu vendor/vanity-search-bitcrack/GPUMath.h
|
||||||
|
$(NVCC) $(NVCCFLAGS_SECP) -o $@ $<
|
||||||
|
|
||||||
|
secp256k1-test: secp256k1-batch-mul test_secp256k1_known_answers.py
|
||||||
|
python3 -u test_secp256k1_known_answers.py ./secp256k1-batch-mul
|
||||||
|
|
||||||
|
secp256k1-bench: secp256k1-batch-mul test_secp256k1_known_answers.py
|
||||||
|
python3 -u test_secp256k1_known_answers.py ./secp256k1-batch-mul --n 100000
|
||||||
|
|
||||||
# bend form B — CGBN bignum batch worker.
|
# bend form B — CGBN bignum batch worker.
|
||||||
# CGBN_INC must point at a checkout of https://github.com/NVlabs/CGBN/include
|
# CGBN_INC must point at a checkout of https://github.com/NVlabs/CGBN/include
|
||||||
# (header-only consumption; CGBN headers are BSD-3-Clause, our wrapper
|
# (header-only consumption; CGBN headers are BSD-3-Clause, our wrapper
|
||||||
|
|
|
||||||
|
|
@ -42,6 +42,13 @@
|
||||||
(or (get-environment-variable "CGBN_BATCH_WORKER")
|
(or (get-environment-variable "CGBN_BATCH_WORKER")
|
||||||
"./cgbn-batch-worker"))
|
"./cgbn-batch-worker"))
|
||||||
|
|
||||||
|
;; secp256k1-batch-mul — bend form A. Batched scalar*G on secp256k1 via
|
||||||
|
;; vendored VanitySearch-Bitcrack GPUMath.h (AGPL-3.0). Binary wire uses
|
||||||
|
;; BSCP request / BSCR response magic to stay distinct from BSHK/BCGB.
|
||||||
|
(define *binary-secp256k1-batch*
|
||||||
|
(or (get-environment-variable "SECP256K1_BATCH_WORKER")
|
||||||
|
"./secp256k1-batch-mul"))
|
||||||
|
|
||||||
(define (parse-port-arg args)
|
(define (parse-port-arg args)
|
||||||
(let loop ((rest args))
|
(let loop ((rest args))
|
||||||
(cond
|
(cond
|
||||||
|
|
@ -301,6 +308,36 @@
|
||||||
(delete-file in-path) (delete-file out-path)
|
(delete-file in-path) (delete-file out-path)
|
||||||
(wire-send-raw client (string-append "BERR" (cdr status)))))))))
|
(wire-send-raw client (string-append "BERR" (cdr status)))))))))
|
||||||
|
|
||||||
|
;; Binary wire for bend form A (cuda-secp256k1-batched-mul).
|
||||||
|
;; Payload begins with "BSCP"; pass entire blob through to the daemon,
|
||||||
|
;; which expects the same magic + header. Response payload begins
|
||||||
|
;; with "BSCR" on success or "BERR" on failure.
|
||||||
|
(define (handle-binary-secp client payload)
|
||||||
|
(let* ((daemon (cdr (assoc 'cuda-secp256k1-batched-mul *daemons*)))
|
||||||
|
(in-path (gensym-path "/tmp/bend-secp-in" ".bin"))
|
||||||
|
(out-path (gensym-path "/tmp/bend-secp-out" ".bin"))
|
||||||
|
(t-start (current-time-ms)))
|
||||||
|
(write-binary-file in-path payload)
|
||||||
|
(display ";;; bend RECV cuda-secp256k1-batched-mul bytes=")
|
||||||
|
(display (string-length payload))
|
||||||
|
(display " t-ms=") (display t-start) (newline)
|
||||||
|
(let ((status (daemon-process daemon in-path out-path #t)))
|
||||||
|
(let ((wall-ms (- (current-time-ms) t-start)))
|
||||||
|
(cond
|
||||||
|
((eq? status 'ok)
|
||||||
|
(let ((result-blob (read-binary-file out-path)))
|
||||||
|
(delete-file in-path) (delete-file out-path)
|
||||||
|
(display ";;; bend DONE cuda-secp256k1-batched-mul")
|
||||||
|
(display " wall-ms=") (display wall-ms)
|
||||||
|
(display " out-bytes=") (display (string-length result-blob))
|
||||||
|
(newline)
|
||||||
|
(wire-send-raw client result-blob)))
|
||||||
|
(else
|
||||||
|
(display ";;; bend FAIL cuda-secp256k1-batched-mul wall-ms=")
|
||||||
|
(display wall-ms) (newline)
|
||||||
|
(delete-file in-path) (delete-file out-path)
|
||||||
|
(wire-send-raw client (string-append "BERR" (cdr status)))))))))
|
||||||
|
|
||||||
;; Accept one client, handle one request, close. Returns #t to keep
|
;; Accept one client, handle one request, close. Returns #t to keep
|
||||||
;; serving, #f when the server should stop.
|
;; serving, #f when the server should stop.
|
||||||
(define (handle-one server)
|
(define (handle-one server)
|
||||||
|
|
@ -319,6 +356,10 @@
|
||||||
(string=? (substring payload 0 4) "BCGB"))
|
(string=? (substring payload 0 4) "BCGB"))
|
||||||
(handle-binary-cgbn client payload)
|
(handle-binary-cgbn client payload)
|
||||||
(tcp-close client) #t)
|
(tcp-close client) #t)
|
||||||
|
((and (>= (string-length payload) 4)
|
||||||
|
(string=? (substring payload 0 4) "BSCP"))
|
||||||
|
(handle-binary-secp client payload)
|
||||||
|
(tcp-close client) #t)
|
||||||
(else
|
(else
|
||||||
(let* ((req (read-from-string payload))
|
(let* ((req (read-from-string payload))
|
||||||
(resp (handle-request req)))
|
(resp (handle-request req)))
|
||||||
|
|
@ -349,6 +390,7 @@
|
||||||
(set! *worker-port* port)
|
(set! *worker-port* port)
|
||||||
(maybe-register-daemon! 'cuda-shake-fanout *binary-shake-fanout*)
|
(maybe-register-daemon! 'cuda-shake-fanout *binary-shake-fanout*)
|
||||||
(maybe-register-daemon! 'cuda-bignum-cgbn *binary-cgbn-batch*)
|
(maybe-register-daemon! 'cuda-bignum-cgbn *binary-cgbn-batch*)
|
||||||
|
(maybe-register-daemon! 'cuda-secp256k1-batched-mul *binary-secp256k1-batch*)
|
||||||
(let ((server (tcp-listen port)))
|
(let ((server (tcp-listen port)))
|
||||||
(cond
|
(cond
|
||||||
((eq? server #f)
|
((eq? server #f)
|
||||||
|
|
|
||||||
563
examples/cuda-fanout/secp256k1-batch-mul.cu
Normal file
563
examples/cuda-fanout/secp256k1-batch-mul.cu
Normal file
|
|
@ -0,0 +1,563 @@
|
||||||
|
/* secp256k1-batch-mul.cu -- bend form A: batched scalar*base-point on
|
||||||
|
* secp256k1 via CUDA.
|
||||||
|
*
|
||||||
|
* Plan: examples/cuda-fanout/plans/form-A-secp256k1-batched-mul.md
|
||||||
|
* Catalog: CATALOG.md form A (cuda-secp256k1-batched-mul).
|
||||||
|
*
|
||||||
|
* Wire (BSCP input, little-endian u32 headers, big-endian field bytes):
|
||||||
|
* "BSCP" 4 B magic
|
||||||
|
* u32 op_id 0x01 = scalar*base
|
||||||
|
* u32 n number of scalars
|
||||||
|
* u8[32] base_x base point X, big-endian (secp256k1 native)
|
||||||
|
* u8[32] base_y base point Y, big-endian
|
||||||
|
* u8[n*32] scalars big-endian 256-bit ints
|
||||||
|
*
|
||||||
|
* Wire (BSCR output):
|
||||||
|
* "BSCR" 4 B magic on success, "BERR" on failure
|
||||||
|
* u32 status 0 ok / 1 bad-magic / 2 bad-n / 3 cuda-fail /
|
||||||
|
* 4 truncated / 5 bad-op / 6 base-not-on-curve
|
||||||
|
* u32 n echoed scalar count
|
||||||
|
* u8[n*64] points N affine points, each (X32 || Y32) big-endian.
|
||||||
|
* Infinity encoded as all-zero (64 NUL bytes).
|
||||||
|
*
|
||||||
|
* Math: per thread compute k*P via Jacobian double-and-add
|
||||||
|
* (Q = O; for bit msb..lsb: Q = 2*Q; if bit set: Q = Q + P)
|
||||||
|
* Field arithmetic uses GPUMath.h verbatim from FixedPaul/VanitySearch-Bitcrack
|
||||||
|
* (vendored under vendor/vanity-search-bitcrack/, AGPL-3.0, original
|
||||||
|
* Copyright (c) 2019 Jean Luc PONS preserved).
|
||||||
|
*
|
||||||
|
* After the per-thread Jacobian result lands in (X:Y:Z), we convert to
|
||||||
|
* affine via Z^-1, Z^-2, Z^-3. Day-1 ships per-thread _ModInv; Day-2
|
||||||
|
* may swap in _ModInvGrouped across N threads for amortisation.
|
||||||
|
*
|
||||||
|
* Daemon protocol mirrors shake256-fanout / cgbn-batch-worker:
|
||||||
|
* ready ; once at startup
|
||||||
|
* process-bin <in> <out> ; binary request
|
||||||
|
* done <out> ; success
|
||||||
|
* error <reason> ; failure
|
||||||
|
* quit / bye ; shutdown
|
||||||
|
*
|
||||||
|
* CLI: --binary <in> <out> one-shot mode for tests
|
||||||
|
* --daemon stdin loop, used by gpu-worker.lsp
|
||||||
|
*
|
||||||
|
* License: AGPLv3.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#include <cuda_runtime.h>
|
||||||
|
#include <stdio.h>
|
||||||
|
#include <stdlib.h>
|
||||||
|
#include <string.h>
|
||||||
|
#include <stdint.h>
|
||||||
|
#include <time.h>
|
||||||
|
|
||||||
|
#include "vendor/vanity-search-bitcrack/GPUMath.h"
|
||||||
|
|
||||||
|
#define CUDA_CHECK(stmt) do { \
|
||||||
|
cudaError_t err = (stmt); \
|
||||||
|
if (err != cudaSuccess) { \
|
||||||
|
fprintf(stderr, "CUDA error %s at %s:%d: %s\n", \
|
||||||
|
#stmt, __FILE__, __LINE__, cudaGetErrorString(err)); \
|
||||||
|
return -1; \
|
||||||
|
} \
|
||||||
|
} while (0)
|
||||||
|
|
||||||
|
enum {
|
||||||
|
OP_SCALAR_BASE_MUL = 0x01,
|
||||||
|
};
|
||||||
|
|
||||||
|
enum {
|
||||||
|
STATUS_OK = 0,
|
||||||
|
STATUS_BAD_MAGIC = 1,
|
||||||
|
STATUS_BAD_N = 2,
|
||||||
|
STATUS_CUDA_FAIL = 3,
|
||||||
|
STATUS_TRUNCATED = 4,
|
||||||
|
STATUS_BAD_OP = 5,
|
||||||
|
STATUS_NOT_ON_CURVE = 6,
|
||||||
|
};
|
||||||
|
|
||||||
|
/* ------------------------------------------------------------------ */
|
||||||
|
/* Modular helpers we need but GPUMath.h does not provide. */
|
||||||
|
|
||||||
|
/* Add two field elements mod p. r = a + b mod p. a,b assumed < p.
|
||||||
|
* GPUMath.h ships ModSub but no ModAdd, so we synthesise one via a
|
||||||
|
* conditional subtract after a 256-bit addition.
|
||||||
|
*/
|
||||||
|
__device__ __forceinline__ void ModAdd256(uint64_t* r, uint64_t* a, uint64_t* b) {
|
||||||
|
/* a, b assumed < p. Sum fits in 257 bits. Subtract p once if the
|
||||||
|
* 257-bit sum >= p; the conditional is (carry == 1) OR
|
||||||
|
* (carry == 0 AND sum >= p), which equivalently is "subtract-p
|
||||||
|
* produces no borrow once we extend the carry into bit 256".
|
||||||
|
*/
|
||||||
|
uint64_t s[4];
|
||||||
|
uint64_t carry;
|
||||||
|
UADDO(s[0], a[0], b[0]);
|
||||||
|
UADDC(s[1], a[1], b[1]);
|
||||||
|
UADDC(s[2], a[2], b[2]);
|
||||||
|
UADDC(s[3], a[3], b[3]);
|
||||||
|
UADD(carry, 0ULL, 0ULL);
|
||||||
|
/* Trial: s - p */
|
||||||
|
uint64_t t[4]; uint64_t borrow;
|
||||||
|
USUBO(t[0], s[0], 0xFFFFFFFEFFFFFC2FULL);
|
||||||
|
USUBC(t[1], s[1], 0xFFFFFFFFFFFFFFFFULL);
|
||||||
|
USUBC(t[2], s[2], 0xFFFFFFFFFFFFFFFFULL);
|
||||||
|
USUBC(t[3], s[3], 0xFFFFFFFFFFFFFFFFULL);
|
||||||
|
USUB(borrow, 0ULL, 0ULL);
|
||||||
|
/* Take subtracted form if (carry produced) OR (no borrow). */
|
||||||
|
if (carry || !borrow) {
|
||||||
|
r[0] = t[0]; r[1] = t[1]; r[2] = t[2]; r[3] = t[3];
|
||||||
|
} else {
|
||||||
|
r[0] = s[0]; r[1] = s[1]; r[2] = s[2]; r[3] = s[3];
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
/* in-place: r = 2*r mod p */
|
||||||
|
__device__ __forceinline__ void ModDouble256(uint64_t* r) {
|
||||||
|
ModAdd256(r, r, r);
|
||||||
|
}
|
||||||
|
|
||||||
|
/* in-place: r = 3*r mod p */
|
||||||
|
__device__ __forceinline__ void ModTriple256(uint64_t* r, uint64_t* a) {
|
||||||
|
uint64_t t[4];
|
||||||
|
ModAdd256(t, a, a);
|
||||||
|
ModAdd256(r, t, a);
|
||||||
|
}
|
||||||
|
|
||||||
|
/* Convenience wrapper around GPUMath.h's 5-limb _ModInv. The on-curve
|
||||||
|
* inputs are 4 limbs (< p), so we pad to 5 limbs (limb[4]=0) and
|
||||||
|
* accept the in-place result.
|
||||||
|
*/
|
||||||
|
__device__ __forceinline__ void ModInv256(uint64_t* r) {
|
||||||
|
uint64_t v[5];
|
||||||
|
v[0] = r[0]; v[1] = r[1]; v[2] = r[2]; v[3] = r[3]; v[4] = 0;
|
||||||
|
_ModInv(v);
|
||||||
|
r[0] = v[0]; r[1] = v[1]; r[2] = v[2]; r[3] = v[3];
|
||||||
|
}
|
||||||
|
|
||||||
|
/* Convenience wrapper around GPUMath.h's 5-limb _ModMult / _ModSqr.
|
||||||
|
* Limb 4 is set to 0 for the inputs since they are < p; the result
|
||||||
|
* limb 4 is also < 1 because _ModMult does the final fold.
|
||||||
|
*/
|
||||||
|
/* Final-reduce a 4-limb value possibly in [0, 2p) down to [0, p). */
|
||||||
|
__device__ __forceinline__ void cond_sub_p(uint64_t* r) {
|
||||||
|
uint64_t t[4]; uint64_t borrow;
|
||||||
|
USUBO(t[0], r[0], 0xFFFFFFFEFFFFFC2FULL);
|
||||||
|
USUBC(t[1], r[1], 0xFFFFFFFFFFFFFFFFULL);
|
||||||
|
USUBC(t[2], r[2], 0xFFFFFFFFFFFFFFFFULL);
|
||||||
|
USUBC(t[3], r[3], 0xFFFFFFFFFFFFFFFFULL);
|
||||||
|
USUB(borrow, 0ULL, 0ULL);
|
||||||
|
if (!borrow) { r[0]=t[0]; r[1]=t[1]; r[2]=t[2]; r[3]=t[3]; }
|
||||||
|
}
|
||||||
|
|
||||||
|
__device__ __forceinline__ void ModMult4(uint64_t* r, uint64_t* a, uint64_t* b) {
|
||||||
|
/* _ModMult ignores aa[4]/bb[4] but writes only rr[0..3]. */
|
||||||
|
uint64_t aa[5], bb[5], rr[5] = {0,0,0,0,0};
|
||||||
|
aa[0] = a[0]; aa[1] = a[1]; aa[2] = a[2]; aa[3] = a[3]; aa[4] = 0;
|
||||||
|
bb[0] = b[0]; bb[1] = b[1]; bb[2] = b[2]; bb[3] = b[3]; bb[4] = 0;
|
||||||
|
_ModMult(rr, aa, bb);
|
||||||
|
r[0] = rr[0]; r[1] = rr[1]; r[2] = rr[2]; r[3] = rr[3];
|
||||||
|
cond_sub_p(r);
|
||||||
|
}
|
||||||
|
|
||||||
|
__device__ __forceinline__ void ModSqr4(uint64_t* r, uint64_t* a) {
|
||||||
|
uint64_t aa[5], rr[5] = {0,0,0,0,0};
|
||||||
|
aa[0] = a[0]; aa[1] = a[1]; aa[2] = a[2]; aa[3] = a[3]; aa[4] = 0;
|
||||||
|
_ModSqr(rr, aa);
|
||||||
|
r[0] = rr[0]; r[1] = rr[1]; r[2] = rr[2]; r[3] = rr[3];
|
||||||
|
cond_sub_p(r);
|
||||||
|
}
|
||||||
|
|
||||||
|
/* ------------------------------------------------------------------ */
|
||||||
|
/* Jacobian point ops on secp256k1 (a = 0). Coordinates X,Y,Z each
|
||||||
|
* 4-limb little-endian field elements. Identity = (any, any, 0).
|
||||||
|
*/
|
||||||
|
|
||||||
|
__device__ __forceinline__ int is_zero4(uint64_t* a) {
|
||||||
|
return (a[0] | a[1] | a[2] | a[3]) == 0;
|
||||||
|
}
|
||||||
|
|
||||||
|
/* Jacobian double (a=0):
|
||||||
|
* A = X^2; B = Y^2; C = B^2
|
||||||
|
* D = 2*((X+B)^2 - A - C)
|
||||||
|
* E = 3*A; F = E^2
|
||||||
|
* X3 = F - 2*D
|
||||||
|
* Y3 = E*(D - X3) - 8*C
|
||||||
|
* Z3 = 2*Y*Z
|
||||||
|
*/
|
||||||
|
__device__ void jac_double(uint64_t* X3, uint64_t* Y3, uint64_t* Z3,
|
||||||
|
uint64_t* X1, uint64_t* Y1, uint64_t* Z1) {
|
||||||
|
if (is_zero4(Z1)) {
|
||||||
|
X3[0]=X3[1]=X3[2]=X3[3]=0;
|
||||||
|
Y3[0]=Y3[1]=Y3[2]=Y3[3]=0;
|
||||||
|
Z3[0]=Z3[1]=Z3[2]=Z3[3]=0;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
uint64_t A[4], B[4], C[4], D[4], E[4], F[4], XB[4], t[4];
|
||||||
|
ModSqr4(A, X1); /* A = X^2 */
|
||||||
|
ModSqr4(B, Y1); /* B = Y^2 */
|
||||||
|
ModSqr4(C, B); /* C = B^2 */
|
||||||
|
ModAdd256(XB, X1, B); /* XB = X + B */
|
||||||
|
ModSqr4(t, XB); /* t = (X+B)^2 */
|
||||||
|
ModSub256(t, A); /* t -= A */
|
||||||
|
ModSub256(t, C); /* t -= C */
|
||||||
|
ModDouble256(t); /* D = 2*t */
|
||||||
|
D[0]=t[0]; D[1]=t[1]; D[2]=t[2]; D[3]=t[3];
|
||||||
|
ModTriple256(E, A); /* E = 3*A */
|
||||||
|
ModSqr4(F, E); /* F = E^2 */
|
||||||
|
/* X3 = F - 2*D */
|
||||||
|
uint64_t twoD[4];
|
||||||
|
twoD[0]=D[0]; twoD[1]=D[1]; twoD[2]=D[2]; twoD[3]=D[3];
|
||||||
|
ModDouble256(twoD);
|
||||||
|
ModSub256(X3, F, twoD);
|
||||||
|
/* Y3 = E*(D - X3) - 8*C */
|
||||||
|
uint64_t DmX[4], eightC[4], EDX[4];
|
||||||
|
ModSub256(DmX, D, X3);
|
||||||
|
ModMult4(EDX, E, DmX);
|
||||||
|
eightC[0]=C[0]; eightC[1]=C[1]; eightC[2]=C[2]; eightC[3]=C[3];
|
||||||
|
ModDouble256(eightC);
|
||||||
|
ModDouble256(eightC);
|
||||||
|
ModDouble256(eightC);
|
||||||
|
ModSub256(Y3, EDX, eightC);
|
||||||
|
/* Z3 = 2*Y*Z */
|
||||||
|
uint64_t YZ[4];
|
||||||
|
ModMult4(YZ, Y1, Z1);
|
||||||
|
ModAdd256(Z3, YZ, YZ);
|
||||||
|
}
|
||||||
|
|
||||||
|
/* Jacobian + affine addition (a=0). P1 in Jacobian, P2 in affine.
|
||||||
|
* (X2, Y2) affine. Handles P1=O and P1+/-P2 special cases.
|
||||||
|
*/
|
||||||
|
__device__ void jac_add_affine(uint64_t* X3, uint64_t* Y3, uint64_t* Z3,
|
||||||
|
uint64_t* X1, uint64_t* Y1, uint64_t* Z1,
|
||||||
|
uint64_t* X2, uint64_t* Y2) {
|
||||||
|
if (is_zero4(Z1)) {
|
||||||
|
X3[0]=X2[0]; X3[1]=X2[1]; X3[2]=X2[2]; X3[3]=X2[3];
|
||||||
|
Y3[0]=Y2[0]; Y3[1]=Y2[1]; Y3[2]=Y2[2]; Y3[3]=Y2[3];
|
||||||
|
Z3[0]=1; Z3[1]=0; Z3[2]=0; Z3[3]=0;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
uint64_t Z1Z1[4], U2[4], S2[4], H[4], r[4], HH[4], HHH[4], V[4];
|
||||||
|
uint64_t Z1Z1Z1[4], t1[4], t2[4];
|
||||||
|
ModSqr4(Z1Z1, Z1); /* Z1Z1 = Z1^2 */
|
||||||
|
ModMult4(U2, X2, Z1Z1); /* U2 = X2*Z1^2 */
|
||||||
|
ModMult4(Z1Z1Z1, Z1, Z1Z1); /* Z1^3 */
|
||||||
|
ModMult4(S2, Y2, Z1Z1Z1); /* S2 = Y2*Z1^3 */
|
||||||
|
ModSub256(H, U2, X1); /* H = U2 - X1 */
|
||||||
|
ModSub256(r, S2, Y1); /* r = S2 - Y1 */
|
||||||
|
if (is_zero4(H)) {
|
||||||
|
/* Same X. If r == 0 -> point double; else -> infinity */
|
||||||
|
if (is_zero4(r)) {
|
||||||
|
jac_double(X3, Y3, Z3, X1, Y1, Z1);
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
X3[0]=X3[1]=X3[2]=X3[3]=0;
|
||||||
|
Y3[0]=Y3[1]=Y3[2]=Y3[3]=0;
|
||||||
|
Z3[0]=Z3[1]=Z3[2]=Z3[3]=0;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
ModSqr4(HH, H); /* HH = H^2 */
|
||||||
|
ModMult4(HHH, H, HH); /* HHH = H^3 */
|
||||||
|
ModMult4(V, X1, HH); /* V = X1*HH */
|
||||||
|
/* X3 = r^2 - HHH - 2*V */
|
||||||
|
uint64_t r2[4], twoV[4];
|
||||||
|
ModSqr4(r2, r);
|
||||||
|
twoV[0]=V[0]; twoV[1]=V[1]; twoV[2]=V[2]; twoV[3]=V[3];
|
||||||
|
ModDouble256(twoV);
|
||||||
|
ModSub256(t1, r2, HHH);
|
||||||
|
ModSub256(X3, t1, twoV);
|
||||||
|
/* Y3 = r*(V - X3) - Y1*HHH */
|
||||||
|
ModSub256(t1, V, X3);
|
||||||
|
ModMult4(t2, r, t1);
|
||||||
|
ModMult4(t1, Y1, HHH);
|
||||||
|
ModSub256(Y3, t2, t1);
|
||||||
|
/* Z3 = Z1 * H */
|
||||||
|
ModMult4(Z3, Z1, H);
|
||||||
|
}
|
||||||
|
|
||||||
|
/* ------------------------------------------------------------------ */
|
||||||
|
/* Scan a big-endian 32-byte scalar bit-by-bit, top to bottom. */
|
||||||
|
|
||||||
|
__device__ __forceinline__ int scalar_bit_be(const uint8_t* k, int i) {
|
||||||
|
/* bit 0 = LSB. i in [0, 255]. k is BE: k[0] = top byte. */
|
||||||
|
int byte_idx = 31 - (i >> 3);
|
||||||
|
int bit_idx = i & 7;
|
||||||
|
return (k[byte_idx] >> bit_idx) & 1;
|
||||||
|
}
|
||||||
|
|
||||||
|
/* Per-thread: scalar k (big-endian) * base point (Px, Py, affine) ->
|
||||||
|
* Jacobian (X, Y, Z). Output in 4-limb LE.
|
||||||
|
*/
|
||||||
|
__device__ void scalar_mul(uint64_t* outX, uint64_t* outY, uint64_t* outZ,
|
||||||
|
const uint8_t* k_be,
|
||||||
|
uint64_t* Px, uint64_t* Py) {
|
||||||
|
/* Q = O (Z=0) */
|
||||||
|
uint64_t Qx[4] = {0,0,0,0};
|
||||||
|
uint64_t Qy[4] = {0,0,0,0};
|
||||||
|
uint64_t Qz[4] = {0,0,0,0};
|
||||||
|
uint64_t Tx[4], Ty[4], Tz[4];
|
||||||
|
for (int i = 255; i >= 0; i--) {
|
||||||
|
jac_double(Tx, Ty, Tz, Qx, Qy, Qz);
|
||||||
|
Qx[0]=Tx[0]; Qx[1]=Tx[1]; Qx[2]=Tx[2]; Qx[3]=Tx[3];
|
||||||
|
Qy[0]=Ty[0]; Qy[1]=Ty[1]; Qy[2]=Ty[2]; Qy[3]=Ty[3];
|
||||||
|
Qz[0]=Tz[0]; Qz[1]=Tz[1]; Qz[2]=Tz[2]; Qz[3]=Tz[3];
|
||||||
|
if (scalar_bit_be(k_be, i)) {
|
||||||
|
jac_add_affine(Tx, Ty, Tz, Qx, Qy, Qz, Px, Py);
|
||||||
|
Qx[0]=Tx[0]; Qx[1]=Tx[1]; Qx[2]=Tx[2]; Qx[3]=Tx[3];
|
||||||
|
Qy[0]=Ty[0]; Qy[1]=Ty[1]; Qy[2]=Ty[2]; Qy[3]=Ty[3];
|
||||||
|
Qz[0]=Tz[0]; Qz[1]=Tz[1]; Qz[2]=Tz[2]; Qz[3]=Tz[3];
|
||||||
|
}
|
||||||
|
}
|
||||||
|
outX[0]=Qx[0]; outX[1]=Qx[1]; outX[2]=Qx[2]; outX[3]=Qx[3];
|
||||||
|
outY[0]=Qy[0]; outY[1]=Qy[1]; outY[2]=Qy[2]; outY[3]=Qy[3];
|
||||||
|
outZ[0]=Qz[0]; outZ[1]=Qz[1]; outZ[2]=Qz[2]; outZ[3]=Qz[3];
|
||||||
|
}
|
||||||
|
|
||||||
|
/* ------------------------------------------------------------------ */
|
||||||
|
/* Kernel: one thread per scalar. */
|
||||||
|
|
||||||
|
__global__ void secp_mul_batch_kernel(
|
||||||
|
const uint64_t* base_xy, /* 8 limbs total: [Px(4), Py(4)] LE */
|
||||||
|
const uint8_t* scalars, /* n * 32 bytes BE */
|
||||||
|
uint8_t* out_points,/* n * 64 bytes BE: x32 || y32 */
|
||||||
|
uint32_t n)
|
||||||
|
{
|
||||||
|
uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||||
|
if (idx >= n) return;
|
||||||
|
const uint8_t* k_be = scalars + (size_t)idx * 32;
|
||||||
|
|
||||||
|
uint64_t Px[4] = { base_xy[0], base_xy[1], base_xy[2], base_xy[3] };
|
||||||
|
uint64_t Py[4] = { base_xy[4], base_xy[5], base_xy[6], base_xy[7] };
|
||||||
|
|
||||||
|
uint64_t Jx[4], Jy[4], Jz[4];
|
||||||
|
scalar_mul(Jx, Jy, Jz, k_be, Px, Py);
|
||||||
|
|
||||||
|
uint8_t* outp = out_points + (size_t)idx * 64;
|
||||||
|
|
||||||
|
if (is_zero4(Jz)) {
|
||||||
|
/* infinity: emit 64 NUL bytes */
|
||||||
|
for (int i = 0; i < 64; i++) outp[i] = 0;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
/* Affine = (X/Z^2, Y/Z^3). Inverse Z. */
|
||||||
|
uint64_t Zi[4] = { Jz[0], Jz[1], Jz[2], Jz[3] };
|
||||||
|
ModInv256(Zi);
|
||||||
|
uint64_t Zi2[4], Zi3[4];
|
||||||
|
ModSqr4(Zi2, Zi);
|
||||||
|
ModMult4(Zi3, Zi2, Zi);
|
||||||
|
uint64_t Ax[4], Ay[4];
|
||||||
|
ModMult4(Ax, Jx, Zi2);
|
||||||
|
ModMult4(Ay, Jy, Zi3);
|
||||||
|
|
||||||
|
/* Emit BE. limb 0 = LSB; bytes 0..7 of limb 3 = top of value (BE) */
|
||||||
|
for (int li = 0; li < 4; li++) {
|
||||||
|
uint64_t x = Ax[3 - li];
|
||||||
|
for (int b = 0; b < 8; b++) {
|
||||||
|
outp[li*8 + b] = (uint8_t)(x >> (56 - 8*b));
|
||||||
|
}
|
||||||
|
uint64_t y = Ay[3 - li];
|
||||||
|
for (int b = 0; b < 8; b++) {
|
||||||
|
outp[32 + li*8 + b] = (uint8_t)(y >> (56 - 8*b));
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
/* ------------------------------------------------------------------ */
|
||||||
|
/* Host helpers. */
|
||||||
|
|
||||||
|
static int read_file_to_buffer(const char* path, uint8_t** out_buf, size_t* out_len) {
|
||||||
|
FILE* f = fopen(path, "rb");
|
||||||
|
if (!f) { fprintf(stderr, "open %s failed\n", path); return -1; }
|
||||||
|
fseek(f, 0, SEEK_END);
|
||||||
|
long sz = ftell(f);
|
||||||
|
fseek(f, 0, SEEK_SET);
|
||||||
|
uint8_t* buf = (uint8_t*)malloc(sz);
|
||||||
|
if (!buf) { fclose(f); return -1; }
|
||||||
|
if (fread(buf, 1, sz, f) != (size_t)sz) { free(buf); fclose(f); return -1; }
|
||||||
|
fclose(f);
|
||||||
|
*out_buf = buf;
|
||||||
|
*out_len = sz;
|
||||||
|
return 0;
|
||||||
|
}
|
||||||
|
|
||||||
|
static int write_buffer_to_file(const char* path, const uint8_t* buf, size_t len) {
|
||||||
|
FILE* f = fopen(path, "wb");
|
||||||
|
if (!f) { fprintf(stderr, "create %s failed\n", path); return -1; }
|
||||||
|
if (fwrite(buf, 1, len, f) != len) { fclose(f); return -1; }
|
||||||
|
fclose(f);
|
||||||
|
return 0;
|
||||||
|
}
|
||||||
|
|
||||||
|
static uint32_t rd32_le(const uint8_t* p) {
|
||||||
|
return (uint32_t)p[0] | ((uint32_t)p[1] << 8)
|
||||||
|
| ((uint32_t)p[2] << 16) | ((uint32_t)p[3] << 24);
|
||||||
|
}
|
||||||
|
static void wr32_le(uint8_t* p, uint32_t v) {
|
||||||
|
p[0] = v & 0xff; p[1] = (v >> 8) & 0xff;
|
||||||
|
p[2] = (v >> 16) & 0xff; p[3] = (v >> 24) & 0xff;
|
||||||
|
}
|
||||||
|
|
||||||
|
/* Big-endian 32-byte -> 4-limb little-endian u64. */
|
||||||
|
static void be32_to_limbs(uint64_t out[4], const uint8_t* in) {
|
||||||
|
for (int li = 0; li < 4; li++) {
|
||||||
|
uint64_t v = 0;
|
||||||
|
for (int b = 0; b < 8; b++) {
|
||||||
|
v = (v << 8) | in[li*8 + b];
|
||||||
|
}
|
||||||
|
out[3 - li] = v;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
/* Emit BSCR error payload to disk so the daemon protocol can write a
|
||||||
|
* legal output even on input-validation failure.
|
||||||
|
*/
|
||||||
|
static int write_bscr_error(const char* out_path, uint32_t status) {
|
||||||
|
uint8_t buf[12];
|
||||||
|
memcpy(buf, "BSCR", 4);
|
||||||
|
wr32_le(buf + 4, status);
|
||||||
|
wr32_le(buf + 8, 0);
|
||||||
|
return write_buffer_to_file(out_path, buf, sizeof(buf));
|
||||||
|
}
|
||||||
|
|
||||||
|
static int process_one_bin(const char* in_path, const char* out_path) {
|
||||||
|
uint8_t* in_buf = NULL;
|
||||||
|
size_t in_len = 0;
|
||||||
|
if (read_file_to_buffer(in_path, &in_buf, &in_len) != 0) return -1;
|
||||||
|
|
||||||
|
if (in_len < 12 || memcmp(in_buf, "BSCP", 4) != 0) {
|
||||||
|
fprintf(stderr, "bad BSCP magic / short input (%zu bytes)\n", in_len);
|
||||||
|
free(in_buf);
|
||||||
|
write_bscr_error(out_path, STATUS_BAD_MAGIC);
|
||||||
|
return -1;
|
||||||
|
}
|
||||||
|
uint32_t op_id = rd32_le(in_buf + 4);
|
||||||
|
uint32_t n = rd32_le(in_buf + 8);
|
||||||
|
|
||||||
|
if (op_id != OP_SCALAR_BASE_MUL) {
|
||||||
|
fprintf(stderr, "unsupported op_id 0x%02x\n", op_id);
|
||||||
|
free(in_buf);
|
||||||
|
write_bscr_error(out_path, STATUS_BAD_OP);
|
||||||
|
return -1;
|
||||||
|
}
|
||||||
|
|
||||||
|
size_t expected = 12 + 64 + (size_t)n * 32;
|
||||||
|
if (in_len != expected) {
|
||||||
|
fprintf(stderr, "BSCP size %zu != expected %zu (n=%u)\n", in_len, expected, n);
|
||||||
|
free(in_buf);
|
||||||
|
write_bscr_error(out_path, STATUS_TRUNCATED);
|
||||||
|
return -1;
|
||||||
|
}
|
||||||
|
if (n == 0) {
|
||||||
|
free(in_buf);
|
||||||
|
uint8_t hdr[12];
|
||||||
|
memcpy(hdr, "BSCR", 4);
|
||||||
|
wr32_le(hdr + 4, STATUS_OK);
|
||||||
|
wr32_le(hdr + 8, 0);
|
||||||
|
return write_buffer_to_file(out_path, hdr, sizeof(hdr));
|
||||||
|
}
|
||||||
|
|
||||||
|
const uint8_t* base_x_be = in_buf + 12;
|
||||||
|
const uint8_t* base_y_be = in_buf + 12 + 32;
|
||||||
|
const uint8_t* scalars_be = in_buf + 12 + 64;
|
||||||
|
|
||||||
|
uint64_t base_limbs[8];
|
||||||
|
be32_to_limbs(&base_limbs[0], base_x_be);
|
||||||
|
be32_to_limbs(&base_limbs[4], base_y_be);
|
||||||
|
|
||||||
|
/* Device buffers */
|
||||||
|
uint64_t* d_base = NULL;
|
||||||
|
uint8_t* d_scalars = NULL;
|
||||||
|
uint8_t* d_points = NULL;
|
||||||
|
CUDA_CHECK(cudaMalloc(&d_base, 8 * sizeof(uint64_t)));
|
||||||
|
CUDA_CHECK(cudaMalloc(&d_scalars, (size_t)n * 32));
|
||||||
|
CUDA_CHECK(cudaMalloc(&d_points, (size_t)n * 64));
|
||||||
|
|
||||||
|
CUDA_CHECK(cudaMemcpy(d_base, base_limbs, 8 * sizeof(uint64_t),
|
||||||
|
cudaMemcpyHostToDevice));
|
||||||
|
CUDA_CHECK(cudaMemcpy(d_scalars, scalars_be, (size_t)n * 32,
|
||||||
|
cudaMemcpyHostToDevice));
|
||||||
|
|
||||||
|
uint32_t threads_per_block = 128;
|
||||||
|
uint32_t blocks = (n + threads_per_block - 1) / threads_per_block;
|
||||||
|
|
||||||
|
struct timespec t0, t1;
|
||||||
|
clock_gettime(CLOCK_MONOTONIC, &t0);
|
||||||
|
secp_mul_batch_kernel<<<blocks, threads_per_block>>>(d_base, d_scalars, d_points, n);
|
||||||
|
CUDA_CHECK(cudaDeviceSynchronize());
|
||||||
|
clock_gettime(CLOCK_MONOTONIC, &t1);
|
||||||
|
double kernel_ms = (t1.tv_sec - t0.tv_sec) * 1e3
|
||||||
|
+ (t1.tv_nsec - t0.tv_nsec) / 1e6;
|
||||||
|
|
||||||
|
/* Assemble output */
|
||||||
|
size_t out_size = 12 + (size_t)n * 64;
|
||||||
|
uint8_t* out_buf = (uint8_t*)malloc(out_size);
|
||||||
|
if (!out_buf) {
|
||||||
|
cudaFree(d_base); cudaFree(d_scalars); cudaFree(d_points);
|
||||||
|
free(in_buf); return -1;
|
||||||
|
}
|
||||||
|
memcpy(out_buf, "BSCR", 4);
|
||||||
|
wr32_le(out_buf + 4, STATUS_OK);
|
||||||
|
wr32_le(out_buf + 8, n);
|
||||||
|
CUDA_CHECK(cudaMemcpy(out_buf + 12, d_points, (size_t)n * 64,
|
||||||
|
cudaMemcpyDeviceToHost));
|
||||||
|
|
||||||
|
cudaFree(d_base); cudaFree(d_scalars); cudaFree(d_points);
|
||||||
|
free(in_buf);
|
||||||
|
|
||||||
|
int wr = write_buffer_to_file(out_path, out_buf, out_size);
|
||||||
|
free(out_buf);
|
||||||
|
fprintf(stderr, "secp op=0x%02x n=%u kernel=%.2fms %.3f Mkeys/s\n",
|
||||||
|
op_id, n, kernel_ms,
|
||||||
|
(kernel_ms > 0) ? ((double)n / 1000.0) / kernel_ms : 0.0);
|
||||||
|
return wr;
|
||||||
|
}
|
||||||
|
|
||||||
|
/* ------------------------------------------------------------------ */
|
||||||
|
|
||||||
|
int main(int argc, char** argv) {
|
||||||
|
if (argc >= 2 && !strcmp(argv[1], "--daemon")) {
|
||||||
|
int dev_count = 0;
|
||||||
|
cudaGetDeviceCount(&dev_count);
|
||||||
|
if (dev_count == 0) {
|
||||||
|
fprintf(stderr, "no CUDA devices\n");
|
||||||
|
return 1;
|
||||||
|
}
|
||||||
|
cudaFree(0);
|
||||||
|
fprintf(stdout, "ready\n"); fflush(stdout);
|
||||||
|
|
||||||
|
char line[4096];
|
||||||
|
while (fgets(line, sizeof(line), stdin)) {
|
||||||
|
char* nl = strchr(line, '\n');
|
||||||
|
if (nl) *nl = 0;
|
||||||
|
if (!strcmp(line, "quit") || !strcmp(line, "bye")) {
|
||||||
|
fprintf(stdout, "bye\n"); fflush(stdout);
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
if (!strncmp(line, "process-bin ", 12)) {
|
||||||
|
char* args = line + 12;
|
||||||
|
char* sep = strchr(args, ' ');
|
||||||
|
if (!sep) {
|
||||||
|
fprintf(stdout, "error bad-args\n"); fflush(stdout);
|
||||||
|
continue;
|
||||||
|
}
|
||||||
|
*sep = 0;
|
||||||
|
char* in_path = args;
|
||||||
|
char* out_path = sep + 1;
|
||||||
|
int r = process_one_bin(in_path, out_path);
|
||||||
|
fprintf(stdout, r == 0 ? "done %s\n" : "error process-failed\n",
|
||||||
|
out_path);
|
||||||
|
fflush(stdout);
|
||||||
|
} else if (line[0]) {
|
||||||
|
fprintf(stdout, "error unknown-command\n"); fflush(stdout);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
return 0;
|
||||||
|
}
|
||||||
|
if (argc >= 4 && !strcmp(argv[1], "--binary")) {
|
||||||
|
return process_one_bin(argv[2], argv[3]);
|
||||||
|
}
|
||||||
|
fprintf(stderr,
|
||||||
|
"usage: %s --binary <input.bin> <output.bin>\n"
|
||||||
|
" %s --daemon (read commands on stdin)\n",
|
||||||
|
argv[0], argv[0]);
|
||||||
|
return 1;
|
||||||
|
}
|
||||||
250
examples/cuda-fanout/test_secp256k1_known_answers.py
Normal file
250
examples/cuda-fanout/test_secp256k1_known_answers.py
Normal file
|
|
@ -0,0 +1,250 @@
|
||||||
|
"""test_secp256k1_known_answers.py — bend form A validator.
|
||||||
|
|
||||||
|
Pipes BSCP-format scalar*G requests at our secp256k1-batch-mul worker
|
||||||
|
via --binary mode, parses the BSCR response, and compares each output
|
||||||
|
point against a host reference (coincurve, fallback ecdsa).
|
||||||
|
|
||||||
|
Wire format matches secp256k1-batch-mul.cu header doc:
|
||||||
|
BSCP + u32 op_id + u32 n + base_x(BE,32) + base_y(BE,32) + scalars(BE,32*n)
|
||||||
|
BSCR + u32 status + u32 n + points(BE,64*n x||y)
|
||||||
|
|
||||||
|
Usage:
|
||||||
|
python3 -u test_secp256k1_known_answers.py [worker-path]
|
||||||
|
python3 -u test_secp256k1_known_answers.py ./secp256k1-batch-mul --n 1000
|
||||||
|
"""
|
||||||
|
import argparse
|
||||||
|
import os
|
||||||
|
import random
|
||||||
|
import struct
|
||||||
|
import subprocess
|
||||||
|
import sys
|
||||||
|
import tempfile
|
||||||
|
import time
|
||||||
|
|
||||||
|
# secp256k1 curve order and base point.
|
||||||
|
SECP256K1_N = 0xFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFEBAAEDCE6AF48A03BBFD25E8CD0364141
|
||||||
|
SECP256K1_GX = 0x79BE667EF9DCBBAC55A06295CE870B07029BFCDB2DCE28D959F2815B16F81798
|
||||||
|
SECP256K1_GY = 0x483ADA7726A3C4655DA4FBFC0E1108A8FD17B448A68554199C47D08FFB10D4B8
|
||||||
|
|
||||||
|
HAS_COINCURVE = False
|
||||||
|
HAS_ECDSA = False
|
||||||
|
try:
|
||||||
|
from coincurve import PrivateKey, PublicKey
|
||||||
|
HAS_COINCURVE = True
|
||||||
|
except ImportError:
|
||||||
|
try:
|
||||||
|
from ecdsa import SigningKey, SECP256k1, ellipticcurve
|
||||||
|
from ecdsa.ecdsa import generator_secp256k1
|
||||||
|
HAS_ECDSA = True
|
||||||
|
except ImportError:
|
||||||
|
pass
|
||||||
|
|
||||||
|
if not HAS_COINCURVE and not HAS_ECDSA:
|
||||||
|
print("ERROR: install coincurve (pip install coincurve) or python-ecdsa")
|
||||||
|
sys.exit(1)
|
||||||
|
|
||||||
|
|
||||||
|
def host_scalar_mul_G(k):
|
||||||
|
"""Return (x, y) of k*G on secp256k1, or (0, 0) for infinity."""
|
||||||
|
k = k % SECP256K1_N
|
||||||
|
if k == 0:
|
||||||
|
return (0, 0)
|
||||||
|
if HAS_COINCURVE:
|
||||||
|
# coincurve's PrivateKey requires 1 <= k < n; encode public_key.format(False)
|
||||||
|
# gives 0x04 + x32 + y32.
|
||||||
|
priv = PrivateKey.from_int(k)
|
||||||
|
pub = priv.public_key.format(compressed=False)
|
||||||
|
assert pub[0] == 0x04
|
||||||
|
x = int.from_bytes(pub[1:33], "big")
|
||||||
|
y = int.from_bytes(pub[33:65], "big")
|
||||||
|
return (x, y)
|
||||||
|
else:
|
||||||
|
pt = k * generator_secp256k1
|
||||||
|
if pt == ellipticcurve.INFINITY:
|
||||||
|
return (0, 0)
|
||||||
|
return (pt.x(), pt.y())
|
||||||
|
|
||||||
|
|
||||||
|
def host_scalar_mul_point(k, px, py):
|
||||||
|
"""k * (px, py) on secp256k1. Used only when base != G."""
|
||||||
|
if HAS_ECDSA:
|
||||||
|
from ecdsa.ecdsa import generator_secp256k1
|
||||||
|
# Build the affine point on the same curve.
|
||||||
|
curve = generator_secp256k1.curve()
|
||||||
|
pt = ellipticcurve.Point(curve, px, py)
|
||||||
|
result = k * pt
|
||||||
|
if result == ellipticcurve.INFINITY:
|
||||||
|
return (0, 0)
|
||||||
|
return (result.x(), result.y())
|
||||||
|
# coincurve only exposes G-based ops; for non-G base fall through
|
||||||
|
# to a pure-Python double-and-add.
|
||||||
|
return _python_scalar_mul(k, px, py)
|
||||||
|
|
||||||
|
|
||||||
|
_SECP_P = 0xFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFEFFFFFC2F
|
||||||
|
|
||||||
|
def _python_scalar_mul(k, px, py):
|
||||||
|
def inv(a, m): return pow(a, -1, m)
|
||||||
|
def add(P, Q):
|
||||||
|
if P is None: return Q
|
||||||
|
if Q is None: return P
|
||||||
|
if P[0] == Q[0]:
|
||||||
|
if (P[1] + Q[1]) % _SECP_P == 0:
|
||||||
|
return None
|
||||||
|
lam = (3 * P[0] * P[0]) * inv(2 * P[1], _SECP_P) % _SECP_P
|
||||||
|
else:
|
||||||
|
lam = (Q[1] - P[1]) * inv(Q[0] - P[0], _SECP_P) % _SECP_P
|
||||||
|
x3 = (lam * lam - P[0] - Q[0]) % _SECP_P
|
||||||
|
y3 = (lam * (P[0] - x3) - P[1]) % _SECP_P
|
||||||
|
return (x3, y3)
|
||||||
|
R = None
|
||||||
|
P = (px, py)
|
||||||
|
while k:
|
||||||
|
if k & 1: R = add(R, P)
|
||||||
|
P = add(P, P)
|
||||||
|
k >>= 1
|
||||||
|
return (0, 0) if R is None else R
|
||||||
|
|
||||||
|
|
||||||
|
def int_to_be32(v):
|
||||||
|
return int(v % (1 << 256)).to_bytes(32, "big")
|
||||||
|
|
||||||
|
|
||||||
|
def build_bscp(n, base_x, base_y, scalars):
|
||||||
|
parts = [b"BSCP",
|
||||||
|
(0x01).to_bytes(4, "little"),
|
||||||
|
n.to_bytes(4, "little"),
|
||||||
|
int_to_be32(base_x),
|
||||||
|
int_to_be32(base_y)]
|
||||||
|
for s in scalars:
|
||||||
|
parts.append(int_to_be32(s))
|
||||||
|
return b"".join(parts)
|
||||||
|
|
||||||
|
|
||||||
|
def parse_bscr(blob, n):
|
||||||
|
if len(blob) < 12:
|
||||||
|
raise RuntimeError(f"short BSCR: {len(blob)} bytes")
|
||||||
|
magic = blob[:4]
|
||||||
|
status = int.from_bytes(blob[4:8], "little")
|
||||||
|
n_echo = int.from_bytes(blob[8:12], "little")
|
||||||
|
if magic != b"BSCR":
|
||||||
|
raise RuntimeError(f"bad magic {magic!r}, status={status}")
|
||||||
|
if status != 0:
|
||||||
|
raise RuntimeError(f"worker status={status} (nonzero)")
|
||||||
|
if n_echo != n:
|
||||||
|
raise RuntimeError(f"n mismatch: req={n} resp={n_echo}")
|
||||||
|
expected = 12 + n * 64
|
||||||
|
if len(blob) != expected:
|
||||||
|
raise RuntimeError(f"BSCR size {len(blob)} != expected {expected}")
|
||||||
|
out = []
|
||||||
|
for i in range(n):
|
||||||
|
off = 12 + i * 64
|
||||||
|
x = int.from_bytes(blob[off:off+32], "big")
|
||||||
|
y = int.from_bytes(blob[off+32:off+64], "big")
|
||||||
|
out.append((x, y))
|
||||||
|
return out
|
||||||
|
|
||||||
|
|
||||||
|
def run_worker(worker, blob):
|
||||||
|
with tempfile.NamedTemporaryFile(delete=False, suffix=".bscp") as fin:
|
||||||
|
fin.write(blob)
|
||||||
|
in_path = fin.name
|
||||||
|
out_path = in_path.replace(".bscp", ".bscr")
|
||||||
|
try:
|
||||||
|
r = subprocess.run([worker, "--binary", in_path, out_path],
|
||||||
|
capture_output=True, timeout=600)
|
||||||
|
if r.returncode != 0:
|
||||||
|
raise RuntimeError(f"worker exit {r.returncode}: {r.stderr[:400].decode(errors='replace')}")
|
||||||
|
with open(out_path, "rb") as f:
|
||||||
|
return f.read(), r.stderr.decode(errors="replace")
|
||||||
|
finally:
|
||||||
|
for p in (in_path, out_path):
|
||||||
|
try: os.unlink(p)
|
||||||
|
except OSError: pass
|
||||||
|
|
||||||
|
|
||||||
|
def test_known_small(worker):
|
||||||
|
"""Tiny fixed-scalar sanity. Includes 1, 2, 3, n-1, n (->infinity)."""
|
||||||
|
scalars = [1, 2, 3, 7, 0xdeadbeef,
|
||||||
|
SECP256K1_N - 1,
|
||||||
|
SECP256K1_N, # infinity
|
||||||
|
(1 << 128) - 1]
|
||||||
|
n = len(scalars)
|
||||||
|
blob = build_bscp(n, SECP256K1_GX, SECP256K1_GY, scalars)
|
||||||
|
out, _ = run_worker(worker, blob)
|
||||||
|
pts = parse_bscr(out, n)
|
||||||
|
fails = 0
|
||||||
|
for i, k in enumerate(scalars):
|
||||||
|
exp = host_scalar_mul_G(k)
|
||||||
|
got = pts[i]
|
||||||
|
if exp != got:
|
||||||
|
fails += 1
|
||||||
|
if fails <= 3:
|
||||||
|
print(f" FAIL k=0x{k:x}")
|
||||||
|
print(f" exp x=0x{exp[0]:064x} y=0x{exp[1]:064x}")
|
||||||
|
print(f" got x=0x{got[0]:064x} y=0x{got[1]:064x}")
|
||||||
|
label = "known-small"
|
||||||
|
if fails == 0:
|
||||||
|
print(f" PASS {label:25s} n={n}")
|
||||||
|
return True
|
||||||
|
print(f" FAIL {label:25s} n={n} mismatches={fails}/{n}")
|
||||||
|
return False
|
||||||
|
|
||||||
|
|
||||||
|
def test_random_n(worker, n, seed=0xC0FFEE):
|
||||||
|
rng = random.Random(seed ^ n)
|
||||||
|
scalars = [rng.randrange(1, SECP256K1_N) for _ in range(n)]
|
||||||
|
blob = build_bscp(n, SECP256K1_GX, SECP256K1_GY, scalars)
|
||||||
|
t0 = time.monotonic()
|
||||||
|
out, stderr_text = run_worker(worker, blob)
|
||||||
|
wall = time.monotonic() - t0
|
||||||
|
pts = parse_bscr(out, n)
|
||||||
|
fails = 0
|
||||||
|
first = None
|
||||||
|
t_h0 = time.monotonic()
|
||||||
|
for i, k in enumerate(scalars):
|
||||||
|
exp = host_scalar_mul_G(k)
|
||||||
|
if pts[i] != exp:
|
||||||
|
fails += 1
|
||||||
|
if first is None:
|
||||||
|
first = (i, k, exp, pts[i])
|
||||||
|
host_ms = (time.monotonic() - t_h0) * 1000
|
||||||
|
label = f"random n={n}"
|
||||||
|
if fails == 0:
|
||||||
|
thr = n / wall / 1e3
|
||||||
|
print(f" PASS {label:25s} n={n} worker={wall*1000:.1f}ms host={host_ms:.0f}ms kkeys/s={thr:.1f}")
|
||||||
|
print(f" {stderr_text.strip().splitlines()[-1] if stderr_text else ''}")
|
||||||
|
return True
|
||||||
|
i, k, exp, got = first
|
||||||
|
print(f" FAIL {label:25s} n={n} mismatches={fails}/{n}")
|
||||||
|
print(f" first @ i={i}")
|
||||||
|
print(f" k = 0x{k:064x}")
|
||||||
|
print(f" exp = ({exp[0]:#x}, {exp[1]:#x})")
|
||||||
|
print(f" got = ({got[0]:#x}, {got[1]:#x})")
|
||||||
|
return False
|
||||||
|
|
||||||
|
|
||||||
|
def main():
|
||||||
|
ap = argparse.ArgumentParser()
|
||||||
|
ap.add_argument("worker", nargs="?", default="./secp256k1-batch-mul")
|
||||||
|
ap.add_argument("--n", type=int, default=None,
|
||||||
|
help="single N override (default: sweep)")
|
||||||
|
args = ap.parse_args()
|
||||||
|
if not os.path.exists(args.worker):
|
||||||
|
print(f"worker not found: {args.worker}")
|
||||||
|
sys.exit(1)
|
||||||
|
print(f"=== secp256k1-batch-mul validation ({args.worker}) ===")
|
||||||
|
print(f" reference: {'coincurve' if HAS_COINCURVE else 'python-ecdsa'}")
|
||||||
|
ok = test_known_small(args.worker)
|
||||||
|
ns = [args.n] if args.n else [32, 1000, 10000]
|
||||||
|
for n in ns:
|
||||||
|
ok = test_random_n(args.worker, n) and ok
|
||||||
|
if ok:
|
||||||
|
print("\nALL PASS")
|
||||||
|
sys.exit(0)
|
||||||
|
print("\nFAILED")
|
||||||
|
sys.exit(1)
|
||||||
|
|
||||||
|
|
||||||
|
if __name__ == "__main__":
|
||||||
|
main()
|
||||||
891
examples/cuda-fanout/vendor/vanity-search-bitcrack/GPUMath.h
vendored
Normal file
891
examples/cuda-fanout/vendor/vanity-search-bitcrack/GPUMath.h
vendored
Normal file
|
|
@ -0,0 +1,891 @@
|
||||||
|
/*
|
||||||
|
* This file is part of the VanitySearch distribution (https://github.com/JeanLucPons/VanitySearch).
|
||||||
|
* Copyright (c) 2019 Jean Luc PONS.
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, version 3.
|
||||||
|
*
|
||||||
|
* 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
|
||||||
|
* General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------
|
||||||
|
// 256(+64) bits integer CUDA libray for SECPK1
|
||||||
|
// ---------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
// We need 1 extra block for ModInv
|
||||||
|
#define NBBLOCK 5
|
||||||
|
#define BIFULLSIZE 40
|
||||||
|
|
||||||
|
// Assembly directives
|
||||||
|
#define UADDO(c, a, b) asm volatile ("add.cc.u64 %0, %1, %2;" : "=l"(c) : "l"(a), "l"(b) : "memory" );
|
||||||
|
#define UADDC(c, a, b) asm volatile ("addc.cc.u64 %0, %1, %2;" : "=l"(c) : "l"(a), "l"(b) : "memory" );
|
||||||
|
#define UADD(c, a, b) asm volatile ("addc.u64 %0, %1, %2;" : "=l"(c) : "l"(a), "l"(b));
|
||||||
|
|
||||||
|
#define UADDO1(c, a) asm volatile ("add.cc.u64 %0, %0, %1;" : "+l"(c) : "l"(a) : "memory" );
|
||||||
|
#define UADDC1(c, a) asm volatile ("addc.cc.u64 %0, %0, %1;" : "+l"(c) : "l"(a) : "memory" );
|
||||||
|
#define UADD1(c, a) asm volatile ("addc.u64 %0, %0, %1;" : "+l"(c) : "l"(a));
|
||||||
|
|
||||||
|
#define USUBO(c, a, b) asm volatile ("sub.cc.u64 %0, %1, %2;" : "=l"(c) : "l"(a), "l"(b) : "memory" );
|
||||||
|
#define USUBC(c, a, b) asm volatile ("subc.cc.u64 %0, %1, %2;" : "=l"(c) : "l"(a), "l"(b) : "memory" );
|
||||||
|
#define USUB(c, a, b) asm volatile ("subc.u64 %0, %1, %2;" : "=l"(c) : "l"(a), "l"(b));
|
||||||
|
|
||||||
|
#define USUBO1(c, a) asm volatile ("sub.cc.u64 %0, %0, %1;" : "+l"(c) : "l"(a) : "memory" );
|
||||||
|
#define USUBC1(c, a) asm volatile ("subc.cc.u64 %0, %0, %1;" : "+l"(c) : "l"(a) : "memory" );
|
||||||
|
#define USUB1(c, a) asm volatile ("subc.u64 %0, %0, %1;" : "+l"(c) : "l"(a) );
|
||||||
|
|
||||||
|
#define UMULLO(lo,a, b) asm volatile ("mul.lo.u64 %0, %1, %2;" : "=l"(lo) : "l"(a), "l"(b));
|
||||||
|
#define UMULHI(hi,a, b) asm volatile ("mul.hi.u64 %0, %1, %2;" : "=l"(hi) : "l"(a), "l"(b));
|
||||||
|
#define MADDO(r,a,b,c) asm volatile ("mad.hi.cc.u64 %0, %1, %2, %3;" : "=l"(r) : "l"(a), "l"(b), "l"(c) : "memory" );
|
||||||
|
#define MADDC(r,a,b,c) asm volatile ("madc.hi.cc.u64 %0, %1, %2, %3;" : "=l"(r) : "l"(a), "l"(b), "l"(c) : "memory" );
|
||||||
|
#define MADD(r,a,b,c) asm volatile ("madc.hi.u64 %0, %1, %2, %3;" : "=l"(r) : "l"(a), "l"(b), "l"(c));
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
//#define MM64 0xD838091DD2253531ULL
|
||||||
|
//#define MSK62 0x3FFFFFFFFFFFFFFF
|
||||||
|
|
||||||
|
|
||||||
|
__device__ __constant__ uint64_t MM64 = 0xD838091DD2253531ULL;
|
||||||
|
__device__ __constant__ uint64_t MSK62 = 0x3FFFFFFFFFFFFFFFULL;
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
#define _IsPositive(x) (((int64_t)(x[4]))>=0LL)
|
||||||
|
#define _IsNegative(x) (((int64_t)(x[4]))<0LL)
|
||||||
|
#define _IsEqual(a,b) ((a[4] == b[4]) && (a[3] == b[3]) && (a[2] == b[2]) && (a[1] == b[1]) && (a[0] == b[0]))
|
||||||
|
#define _IsZero(a) ((a[4] | a[3] | a[2] | a[1] | a[0]) == 0ULL)
|
||||||
|
#define _IsOne(a) ((a[4] == 0ULL) && (a[3] == 0ULL) && (a[2] == 0ULL) && (a[1] == 0ULL) && (a[0] == 1ULL))
|
||||||
|
|
||||||
|
#define IDX threadIdx.x
|
||||||
|
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
#define AddP(r) { \
|
||||||
|
UADDO1(r[0], 0xFFFFFFFEFFFFFC2FULL); \
|
||||||
|
UADDC1(r[1], 0xFFFFFFFFFFFFFFFFULL); \
|
||||||
|
UADDC1(r[2], 0xFFFFFFFFFFFFFFFFULL); \
|
||||||
|
UADDC1(r[3], 0xFFFFFFFFFFFFFFFFULL); \
|
||||||
|
UADD1(r[4], 0ULL);}
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
#define SubP(r) { \
|
||||||
|
USUBO1(r[0], 0xFFFFFFFEFFFFFC2FULL); \
|
||||||
|
USUBC1(r[1], 0xFFFFFFFFFFFFFFFFULL); \
|
||||||
|
USUBC1(r[2], 0xFFFFFFFFFFFFFFFFULL); \
|
||||||
|
USUBC1(r[3], 0xFFFFFFFFFFFFFFFFULL); \
|
||||||
|
USUB1(r[4], 0ULL);}
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
#define Neg(r) {\
|
||||||
|
USUBO(r[0],0ULL,r[0]); \
|
||||||
|
USUBC(r[1],0ULL,r[1]); \
|
||||||
|
USUBC(r[2],0ULL,r[2]); \
|
||||||
|
USUBC(r[3],0ULL,r[3]); \
|
||||||
|
USUB(r[4],0ULL,r[4]); }
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
#define UMult(r, a, b) {\
|
||||||
|
UMULLO(r[0],a[0],b); \
|
||||||
|
UMULLO(r[1],a[1],b); \
|
||||||
|
MADDO(r[1], a[0],b,r[1]); \
|
||||||
|
UMULLO(r[2],a[2], b); \
|
||||||
|
MADDC(r[2], a[1], b, r[2]); \
|
||||||
|
UMULLO(r[3],a[3], b); \
|
||||||
|
MADDC(r[3], a[2], b, r[3]); \
|
||||||
|
MADD(r[4], a[3], b, 0ULL);}
|
||||||
|
|
||||||
|
|
||||||
|
#define UMultSpecial(r, a) {\
|
||||||
|
uint64_t temp; /* Dichiarazione di temp qui */\
|
||||||
|
r[0] = (a[0] << 32) + (a[0] << 9) + (a[0] << 8) + (a[0] << 7) + (a[0] << 6) + (a[0] << 4) + a[0]; \
|
||||||
|
r[1] = (a[1] << 32) + (a[1] << 9) + (a[1] << 8) + (a[1] << 7) + (a[1] << 6) + (a[1] << 4) + a[1]; \
|
||||||
|
MADDO(r[1], a[0], 0x1000003D1ULL, r[1]); \
|
||||||
|
r[2] = (a[2] << 32) + (a[2] << 9) + (a[2] << 8) + (a[2] << 7) + (a[2] << 6) + (a[2] << 4) + a[2]; \
|
||||||
|
MADDC(r[2], a[1], 0x1000003D1ULL, r[2]); \
|
||||||
|
r[3] = (a[3] << 32) + (a[3] << 9) + (a[3] << 8) + (a[3] << 7) + (a[3] << 6) + (a[3] << 4) + a[3]; \
|
||||||
|
temp = r[3]; \
|
||||||
|
MADDC(r[3], a[2], 0x1000003D1ULL, r[3]); \
|
||||||
|
r[4] = temp + a[3]; \
|
||||||
|
MADD(r[4], a[3], 0x1000003D1ULL, 0ULL); \
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
#define Load(r, a) {\
|
||||||
|
(r)[0] = (a)[0]; \
|
||||||
|
(r)[1] = (a)[1]; \
|
||||||
|
(r)[2] = (a)[2]; \
|
||||||
|
(r)[3] = (a)[3]; \
|
||||||
|
(r)[4] = (a)[4];}
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
#define _LoadI64(r, a) {\
|
||||||
|
(r)[0] = a; \
|
||||||
|
(r)[1] = a>>63; \
|
||||||
|
(r)[2] = (r)[1]; \
|
||||||
|
(r)[3] = (r)[1]; \
|
||||||
|
(r)[4] = (r)[1];}
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
#define Load256(r, a) {\
|
||||||
|
(r)[0] = (a)[0]; \
|
||||||
|
(r)[1] = (a)[1]; \
|
||||||
|
(r)[2] = (a)[2]; \
|
||||||
|
(r)[3] = (a)[3];}
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
#define Load256A(r, a) {\
|
||||||
|
(r)[0] = (a)[IDX]; \
|
||||||
|
(r)[1] = (a)[IDX+blockDim.x]; \
|
||||||
|
(r)[2] = (a)[IDX+2*blockDim.x]; \
|
||||||
|
(r)[3] = (a)[IDX+3*blockDim.x];}
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
#define Store256A(r, a) {\
|
||||||
|
(r)[IDX] = (a)[0]; \
|
||||||
|
(r)[IDX+blockDim.x] = (a)[1]; \
|
||||||
|
(r)[IDX+2*blockDim.x] = (a)[2]; \
|
||||||
|
(r)[IDX+3*blockDim.x] = (a)[3];}
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
__device__ void ShiftR62(uint64_t *r) {
|
||||||
|
|
||||||
|
r[0] = (r[1] << 2) | (r[0] >> 62);
|
||||||
|
r[1] = (r[2] << 2) | (r[1] >> 62);
|
||||||
|
r[2] = (r[3] << 2) | (r[2] >> 62);
|
||||||
|
r[3] = (r[4] << 2) | (r[3] >> 62);
|
||||||
|
// With sign extent
|
||||||
|
r[4] = (int64_t)(r[4]) >> 62;
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
__device__ void IMult(uint64_t* r, uint64_t* a, int64_t b) {
|
||||||
|
|
||||||
|
uint64_t t[NBBLOCK];
|
||||||
|
|
||||||
|
// Make b positive
|
||||||
|
if (b < 0) {
|
||||||
|
b = -b;
|
||||||
|
USUBO(t[0], 0ULL, a[0]);
|
||||||
|
USUBC(t[1], 0ULL, a[1]);
|
||||||
|
USUBC(t[2], 0ULL, a[2]);
|
||||||
|
USUBC(t[3], 0ULL, a[3]);
|
||||||
|
USUB(t[4], 0ULL, a[4]);
|
||||||
|
}
|
||||||
|
else {
|
||||||
|
Load(t, a);
|
||||||
|
}
|
||||||
|
|
||||||
|
UMULLO(r[0], t[0], b);
|
||||||
|
UMULLO(r[1], t[1], b);
|
||||||
|
MADDO(r[1], t[0], b, r[1]);
|
||||||
|
UMULLO(r[2], t[2], b);
|
||||||
|
MADDC(r[2], t[1], b, r[2]);
|
||||||
|
UMULLO(r[3], t[3], b);
|
||||||
|
MADDC(r[3], t[2], b, r[3]);
|
||||||
|
UMULLO(r[4], t[4], b);
|
||||||
|
MADD(r[4], t[3], b, r[4]);
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
__device__ void MulP(uint64_t *r, uint64_t a) {
|
||||||
|
|
||||||
|
uint64_t ah;
|
||||||
|
uint64_t al;
|
||||||
|
|
||||||
|
UMULLO(al, a, 0x1000003D1ULL);
|
||||||
|
UMULHI(ah, a, 0x1000003D1ULL);
|
||||||
|
|
||||||
|
USUBO(r[0], 0ULL, al);
|
||||||
|
USUBC(r[1], 0ULL, ah);
|
||||||
|
USUBC(r[2], 0ULL, 0ULL);
|
||||||
|
USUBC(r[3], 0ULL, 0ULL);
|
||||||
|
USUB(r[4], a, 0ULL);
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
__device__ void ModNeg256(uint64_t* r,uint64_t* a) {
|
||||||
|
|
||||||
|
uint64_t t[4];
|
||||||
|
USUBO(t[0],0ULL,a[0]);
|
||||||
|
USUBC(t[1],0ULL,a[1]);
|
||||||
|
USUBC(t[2],0ULL,a[2]);
|
||||||
|
USUBC(t[3],0ULL,a[3]);
|
||||||
|
UADDO(r[0],t[0],0xFFFFFFFEFFFFFC2FULL);
|
||||||
|
UADDC(r[1],t[1],0xFFFFFFFFFFFFFFFFULL);
|
||||||
|
UADDC(r[2],t[2],0xFFFFFFFFFFFFFFFFULL);
|
||||||
|
UADD(r[3],t[3],0xFFFFFFFFFFFFFFFFULL);
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
__device__ void ModNeg256(uint64_t* r) {
|
||||||
|
|
||||||
|
uint64_t t[4];
|
||||||
|
USUBO(t[0],0ULL,r[0]);
|
||||||
|
USUBC(t[1],0ULL,r[1]);
|
||||||
|
USUBC(t[2],0ULL,r[2]);
|
||||||
|
USUBC(t[3],0ULL,r[3]);
|
||||||
|
UADDO(r[0],t[0],0xFFFFFFFEFFFFFC2FULL);
|
||||||
|
UADDC(r[1],t[1],0xFFFFFFFFFFFFFFFFULL);
|
||||||
|
UADDC(r[2],t[2],0xFFFFFFFFFFFFFFFFULL);
|
||||||
|
UADD(r[3],t[3],0xFFFFFFFFFFFFFFFFULL);
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
__device__ void ModSub256(uint64_t* r, uint64_t* a, uint64_t* b) {
|
||||||
|
uint64_t borrow;
|
||||||
|
uint64_t p[4] = { 0xFFFFFFFEFFFFFC2FULL, 0xFFFFFFFFFFFFFFFFULL,
|
||||||
|
0xFFFFFFFFFFFFFFFFULL, 0xFFFFFFFFFFFFFFFFULL };
|
||||||
|
|
||||||
|
USUBO(r[0], a[0], b[0]);
|
||||||
|
USUBC(r[1], a[1], b[1]);
|
||||||
|
USUBC(r[2], a[2], b[2]);
|
||||||
|
USUBC(r[3], a[3], b[3]);
|
||||||
|
USUB(borrow, 0ULL, 0ULL);
|
||||||
|
|
||||||
|
if (borrow) {
|
||||||
|
UADDO1(r[0], p[0]);
|
||||||
|
UADDC1(r[1], p[1]);
|
||||||
|
UADDC1(r[2], p[2]);
|
||||||
|
UADD1(r[3], p[3]);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
__device__ void ModSub256(uint64_t* r, uint64_t* b) {
|
||||||
|
uint64_t borrow;
|
||||||
|
uint64_t p[4] = { 0xFFFFFFFEFFFFFC2FULL, 0xFFFFFFFFFFFFFFFFULL,
|
||||||
|
0xFFFFFFFFFFFFFFFFULL, 0xFFFFFFFFFFFFFFFFULL };
|
||||||
|
|
||||||
|
USUBO1(r[0], b[0]);
|
||||||
|
USUBC1(r[1], b[1]);
|
||||||
|
USUBC1(r[2], b[2]);
|
||||||
|
USUBC1(r[3], b[3]);
|
||||||
|
USUB(borrow, 0ULL, 0ULL);
|
||||||
|
|
||||||
|
if (borrow) {
|
||||||
|
UADDO1(r[0], p[0]);
|
||||||
|
UADDC1(r[1], p[1]);
|
||||||
|
UADDC1(r[2], p[2]);
|
||||||
|
UADD1(r[3], p[3]);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
__device__ void ModSub256isOdd(uint64_t* a, uint64_t* b, uint8_t* parity) { //no need to compute py, we need only parity
|
||||||
|
|
||||||
|
uint64_t t;
|
||||||
|
uint64_t T[4];
|
||||||
|
|
||||||
|
USUBO(T[0], a[0], b[0]);
|
||||||
|
USUBC(T[1], a[1], b[1]);
|
||||||
|
USUBC(T[2], a[2], b[2]);
|
||||||
|
USUBC(T[3], a[3], b[3]);
|
||||||
|
|
||||||
|
USUB(t, 0ULL, 0ULL); // borrow
|
||||||
|
|
||||||
|
*parity = (T[0] & 1) ^ (t & 1); // LSB of T[0] and LSB of t -> parity od sub
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
// Compute a*b*(mod n)
|
||||||
|
// a and b must be lower than n
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
|
||||||
|
__device__ void _ModMult(uint64_t *r, uint64_t *a, uint64_t *b) {
|
||||||
|
|
||||||
|
uint64_t r512[8];
|
||||||
|
uint64_t t[NBBLOCK];
|
||||||
|
uint64_t ah, al;
|
||||||
|
|
||||||
|
r512[5] = 0;
|
||||||
|
r512[6] = 0;
|
||||||
|
r512[7] = 0;
|
||||||
|
|
||||||
|
// 256*256 multiplier
|
||||||
|
UMult(r512, a, b[0]);
|
||||||
|
UMult(t, a, b[1]);
|
||||||
|
UADDO1(r512[1], t[0]);
|
||||||
|
UADDC1(r512[2], t[1]);
|
||||||
|
UADDC1(r512[3], t[2]);
|
||||||
|
UADDC1(r512[4], t[3]);
|
||||||
|
UADD1(r512[5], t[4]);
|
||||||
|
UMult(t, a, b[2]);
|
||||||
|
UADDO1(r512[2], t[0]);
|
||||||
|
UADDC1(r512[3], t[1]);
|
||||||
|
UADDC1(r512[4], t[2]);
|
||||||
|
UADDC1(r512[5], t[3]);
|
||||||
|
UADD1(r512[6], t[4]);
|
||||||
|
UMult(t, a, b[3]);
|
||||||
|
UADDO1(r512[3], t[0]);
|
||||||
|
UADDC1(r512[4], t[1]);
|
||||||
|
UADDC1(r512[5], t[2]);
|
||||||
|
UADDC1(r512[6], t[3]);
|
||||||
|
UADD1(r512[7], t[4]);
|
||||||
|
|
||||||
|
// Reduce from 512 to 320
|
||||||
|
//UMult(t, (r512 + 4), 0x1000003D1ULL);
|
||||||
|
UMultSpecial(t, (r512 + 4));
|
||||||
|
UADDO1(r512[0], t[0]);
|
||||||
|
UADDC1(r512[1], t[1]);
|
||||||
|
UADDC1(r512[2], t[2]);
|
||||||
|
UADDC1(r512[3], t[3]);
|
||||||
|
|
||||||
|
// Reduce from 320 to 256
|
||||||
|
UADD1(t[4], 0ULL);
|
||||||
|
UMULLO(al, t[4], 0x1000003D1ULL);
|
||||||
|
UMULHI(ah, t[4], 0x1000003D1ULL);
|
||||||
|
UADDO(r[0], r512[0], al);
|
||||||
|
UADDC(r[1], r512[1], ah);
|
||||||
|
UADDC(r[2], r512[2], 0ULL);
|
||||||
|
UADD(r[3], r512[3], 0ULL);
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
__device__ void _ModMult(uint64_t *r, uint64_t *a) {
|
||||||
|
|
||||||
|
uint64_t r512[8];
|
||||||
|
uint64_t t[NBBLOCK];
|
||||||
|
uint64_t ah, al;
|
||||||
|
r512[5] = 0;
|
||||||
|
r512[6] = 0;
|
||||||
|
r512[7] = 0;
|
||||||
|
|
||||||
|
// 256*256 multiplier
|
||||||
|
UMult(r512, a, r[0]);
|
||||||
|
UMult(t, a, r[1]);
|
||||||
|
UADDO1(r512[1], t[0]);
|
||||||
|
UADDC1(r512[2], t[1]);
|
||||||
|
UADDC1(r512[3], t[2]);
|
||||||
|
UADDC1(r512[4], t[3]);
|
||||||
|
UADD1(r512[5], t[4]);
|
||||||
|
UMult(t, a, r[2]);
|
||||||
|
UADDO1(r512[2], t[0]);
|
||||||
|
UADDC1(r512[3], t[1]);
|
||||||
|
UADDC1(r512[4], t[2]);
|
||||||
|
UADDC1(r512[5], t[3]);
|
||||||
|
UADD1(r512[6], t[4]);
|
||||||
|
UMult(t, a, r[3]);
|
||||||
|
UADDO1(r512[3], t[0]);
|
||||||
|
UADDC1(r512[4], t[1]);
|
||||||
|
UADDC1(r512[5], t[2]);
|
||||||
|
UADDC1(r512[6], t[3]);
|
||||||
|
UADD1(r512[7], t[4]);
|
||||||
|
|
||||||
|
// Reduce from 512 to 320
|
||||||
|
UMultSpecial(t, (r512 + 4));
|
||||||
|
UADDO1(r512[0], t[0]);
|
||||||
|
UADDC1(r512[1], t[1]);
|
||||||
|
UADDC1(r512[2], t[2]);
|
||||||
|
UADDC1(r512[3], t[3]);
|
||||||
|
|
||||||
|
// Reduce from 320 to 256
|
||||||
|
UADD1(t[4], 0ULL);
|
||||||
|
UMULLO(al, t[4], 0x1000003D1ULL);
|
||||||
|
UMULHI(ah, t[4], 0x1000003D1ULL);
|
||||||
|
UADDO(r[0], r512[0], al);
|
||||||
|
UADDC(r[1], r512[1], ah);
|
||||||
|
UADDC(r[2], r512[2], 0ULL);
|
||||||
|
UADD(r[3], r512[3], 0ULL);
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
__device__ void _ModSqr(uint64_t *rp, const uint64_t *up) {
|
||||||
|
|
||||||
|
uint64_t r512[8];
|
||||||
|
|
||||||
|
uint64_t u10, u11;
|
||||||
|
|
||||||
|
uint64_t r0;
|
||||||
|
uint64_t r1;
|
||||||
|
uint64_t r3;
|
||||||
|
uint64_t r4;
|
||||||
|
|
||||||
|
uint64_t t1;
|
||||||
|
uint64_t t2;
|
||||||
|
|
||||||
|
|
||||||
|
//k=0
|
||||||
|
UMULLO(r512[0], up[0], up[0]);
|
||||||
|
UMULHI(r1, up[0], up[0]);
|
||||||
|
|
||||||
|
//k=1
|
||||||
|
UMULLO(r3, up[0], up[1]);
|
||||||
|
UMULHI(r4, up[0], up[1]);
|
||||||
|
UADDO1(r3, r3);
|
||||||
|
UADDC1(r4, r4);
|
||||||
|
UADD(t1, 0x0ULL, 0x0ULL);
|
||||||
|
UADDO1(r3, r1);
|
||||||
|
UADDC1(r4, 0x0ULL);
|
||||||
|
UADD1(t1, 0x0ULL);
|
||||||
|
r512[1] = r3;
|
||||||
|
|
||||||
|
//k=2
|
||||||
|
UMULLO(r0, up[0], up[2]);
|
||||||
|
UMULHI(r1, up[0], up[2]);
|
||||||
|
UADDO1(r0, r0);
|
||||||
|
UADDC1(r1, r1);
|
||||||
|
UADD(t2, 0x0ULL, 0x0ULL);
|
||||||
|
UMULLO(u10, up[1], up[1]);
|
||||||
|
UMULHI(u11, up[1], up[1]);
|
||||||
|
UADDO1(r0, u10);
|
||||||
|
UADDC1(r1, u11);
|
||||||
|
UADD1(t2, 0x0ULL);
|
||||||
|
UADDO1(r0, r4);
|
||||||
|
UADDC1(r1, t1);
|
||||||
|
UADD1(t2, 0x0ULL);
|
||||||
|
|
||||||
|
r512[2] = r0;
|
||||||
|
|
||||||
|
//k=3
|
||||||
|
UMULLO(r3, up[0], up[3]);
|
||||||
|
UMULHI(r4, up[0], up[3]);
|
||||||
|
UMULLO(u10, up[1], up[2]);
|
||||||
|
UMULHI(u11, up[1], up[2]);
|
||||||
|
UADDO1(r3, u10);
|
||||||
|
UADDC1(r4, u11);
|
||||||
|
UADD(t1, 0x0ULL, 0x0ULL);
|
||||||
|
t1 += t1;
|
||||||
|
UADDO1(r3, r3);
|
||||||
|
UADDC1(r4, r4);
|
||||||
|
UADD1(t1, 0x0ULL);
|
||||||
|
UADDO1(r3, r1);
|
||||||
|
UADDC1(r4, t2);
|
||||||
|
UADD1(t1, 0x0ULL);
|
||||||
|
|
||||||
|
r512[3] = r3;
|
||||||
|
|
||||||
|
//k=4
|
||||||
|
UMULLO(r0, up[1], up[3]);
|
||||||
|
UMULHI(r1, up[1], up[3]);
|
||||||
|
UADDO1(r0, r0);
|
||||||
|
UADDC1(r1, r1);
|
||||||
|
UADD(t2, 0x0ULL, 0x0ULL);
|
||||||
|
UMULLO(u10, up[2], up[2]);
|
||||||
|
UMULHI(u11, up[2], up[2]);
|
||||||
|
UADDO1(r0, u10);
|
||||||
|
UADDC1(r1, u11);
|
||||||
|
UADD1(t2, 0x0ULL);
|
||||||
|
UADDO1(r0, r4);
|
||||||
|
UADDC1(r1, t1);
|
||||||
|
UADD1(t2, 0x0ULL);
|
||||||
|
|
||||||
|
r512[4] = r0;
|
||||||
|
|
||||||
|
//k=5
|
||||||
|
UMULLO(r3, up[2], up[3]);
|
||||||
|
UMULHI(r4, up[2], up[3]);
|
||||||
|
UADDO1(r3, r3);
|
||||||
|
UADDC1(r4, r4);
|
||||||
|
UADD(t1, 0x0ULL, 0x0ULL);
|
||||||
|
UADDO1(r3, r1);
|
||||||
|
UADDC1(r4, t2);
|
||||||
|
UADD1(t1, 0x0ULL);
|
||||||
|
|
||||||
|
r512[5] = r3;
|
||||||
|
|
||||||
|
//k=6
|
||||||
|
UMULLO(r0, up[3], up[3]);
|
||||||
|
UMULHI(r1, up[3], up[3]);
|
||||||
|
UADDO1(r0, r4);
|
||||||
|
UADD1(r1, t1);
|
||||||
|
r512[6] = r0;
|
||||||
|
|
||||||
|
//k=7
|
||||||
|
r512[7] = r1;
|
||||||
|
|
||||||
|
|
||||||
|
// Reduce from 512 to 320
|
||||||
|
UMULLO(r0, r512[4], 0x1000003D1ULL);
|
||||||
|
UMULLO(r1, r512[5], 0x1000003D1ULL);
|
||||||
|
MADDO(r1, r512[4], 0x1000003D1ULL, r1);
|
||||||
|
UMULLO(t2, r512[6], 0x1000003D1ULL);
|
||||||
|
MADDC(t2, r512[5], 0x1000003D1ULL, t2);
|
||||||
|
UMULLO(r3, r512[7], 0x1000003D1ULL);
|
||||||
|
MADDC(r3, r512[6], 0x1000003D1ULL, r3);
|
||||||
|
MADD(r4, r512[7], 0x1000003D1ULL, 0ULL);
|
||||||
|
|
||||||
|
UADDO1(r512[0], r0);
|
||||||
|
UADDC1(r512[1], r1);
|
||||||
|
UADDC1(r512[2], t2);
|
||||||
|
UADDC1(r512[3], r3);
|
||||||
|
|
||||||
|
// Reduce from 320 to 256
|
||||||
|
UADD1(r4, 0ULL);
|
||||||
|
UMULLO(u10, r4, 0x1000003D1ULL);
|
||||||
|
UMULHI(u11, r4, 0x1000003D1ULL);
|
||||||
|
UADDO(rp[0], r512[0], u10);
|
||||||
|
UADDC(rp[1], r512[1], u11);
|
||||||
|
UADDC(rp[2], r512[2], 0ULL);
|
||||||
|
UADD(rp[3], r512[3], 0ULL);
|
||||||
|
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
|
||||||
|
#define MADDS(r,a,b,c) asm volatile ("madc.hi.s64 %0, %1, %2, %3;" : "=l"(r) : "l"(a), "l"(b), "l"(c));
|
||||||
|
|
||||||
|
#define __sright128(a,b,n) ((a)>>(n))|((b)<<(64-(n)))
|
||||||
|
#define __sleft128(a,b,n) ((b)<<(n))|((a)>>(64-(n)))
|
||||||
|
|
||||||
|
__device__ void ShiftR62(uint64_t dest[5], uint64_t r[5], uint64_t carry) {
|
||||||
|
|
||||||
|
dest[0] = (r[1] << 2) | (r[0] >> 62);
|
||||||
|
dest[1] = (r[2] << 2) | (r[1] >> 62);
|
||||||
|
dest[2] = (r[3] << 2) | (r[2] >> 62);
|
||||||
|
dest[3] = (r[4] << 2) | (r[3] >> 62);
|
||||||
|
dest[4] = (carry << 2) | (r[4] >> 62);
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
__device__ uint64_t IMultC(uint64_t* r, uint64_t* a, int64_t b) {
|
||||||
|
|
||||||
|
uint64_t t[NBBLOCK];
|
||||||
|
uint64_t carry;
|
||||||
|
|
||||||
|
// Make b positive
|
||||||
|
if (b < 0) {
|
||||||
|
b = -b;
|
||||||
|
USUBO(t[0], 0ULL, a[0]);
|
||||||
|
USUBC(t[1], 0ULL, a[1]);
|
||||||
|
USUBC(t[2], 0ULL, a[2]);
|
||||||
|
USUBC(t[3], 0ULL, a[3]);
|
||||||
|
USUB(t[4], 0ULL, a[4]);
|
||||||
|
}
|
||||||
|
else {
|
||||||
|
Load(t, a);
|
||||||
|
}
|
||||||
|
|
||||||
|
UMULLO(r[0], t[0], b);
|
||||||
|
UMULLO(r[1], t[1], b);
|
||||||
|
MADDO(r[1], t[0], b, r[1]);
|
||||||
|
UMULLO(r[2], t[2], b);
|
||||||
|
MADDC(r[2], t[1], b, r[2]);
|
||||||
|
UMULLO(r[3], t[3], b);
|
||||||
|
MADDC(r[3], t[2], b, r[3]);
|
||||||
|
UMULLO(r[4], t[4], b);
|
||||||
|
MADDC(r[4], t[3], b, r[4]);
|
||||||
|
MADDS(carry, t[4], b, 0ULL);
|
||||||
|
|
||||||
|
return carry;
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
__device__ __forceinline__ uint32_t ctz(uint64_t x) {
|
||||||
|
uint32_t n;
|
||||||
|
asm("{\n\t"
|
||||||
|
" .reg .u64 tmp;\n\t"
|
||||||
|
" brev.b64 tmp, %1;\n\t"
|
||||||
|
" clz.b64 %0, tmp;\n\t"
|
||||||
|
"}"
|
||||||
|
: "=r"(n) : "l"(x));
|
||||||
|
return n;
|
||||||
|
}
|
||||||
|
|
||||||
|
// ---------------------------------------------------------------------------------------
|
||||||
|
#define SWAP(tmp,x,y) tmp = x; x = y; y = tmp;
|
||||||
|
|
||||||
|
|
||||||
|
__device__ void _DivStep62(uint64_t u[5], uint64_t v[5],
|
||||||
|
int32_t* pos,
|
||||||
|
int64_t* uu, int64_t* uv,
|
||||||
|
int64_t* vu, int64_t* vv) {
|
||||||
|
|
||||||
|
|
||||||
|
// u' = (uu*u + uv*v) >> bitCount
|
||||||
|
// v' = (vu*u + vv*v) >> bitCount
|
||||||
|
// Do not maintain a matrix for r and s, the number of
|
||||||
|
// 'added P' can be easily calculated
|
||||||
|
|
||||||
|
*uu = 1; *uv = 0;
|
||||||
|
*vu = 0; *vv = 1;
|
||||||
|
|
||||||
|
uint32_t bitCount = 62;
|
||||||
|
uint32_t zeros;
|
||||||
|
uint64_t u0 = u[0];
|
||||||
|
uint64_t v0 = v[0];
|
||||||
|
|
||||||
|
// Extract 64 MSB of u and v
|
||||||
|
// u and v must be positive
|
||||||
|
uint64_t uh, vh;
|
||||||
|
int64_t w, x, y, z;
|
||||||
|
bitCount = 62;
|
||||||
|
|
||||||
|
while (*pos > 0 && (u[*pos] | v[*pos]) == 0) (*pos)--;
|
||||||
|
if (*pos == 0) {
|
||||||
|
|
||||||
|
uh = u[0];
|
||||||
|
vh = v[0];
|
||||||
|
|
||||||
|
}
|
||||||
|
else {
|
||||||
|
|
||||||
|
uint32_t s = __clzll(u[*pos] | v[*pos]);
|
||||||
|
if (s == 0) {
|
||||||
|
uh = u[*pos];
|
||||||
|
vh = v[*pos];
|
||||||
|
}
|
||||||
|
else {
|
||||||
|
uh = __sleft128(u[*pos - 1], u[*pos], s);
|
||||||
|
vh = __sleft128(v[*pos - 1], v[*pos], s);
|
||||||
|
}
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
while (true) {
|
||||||
|
|
||||||
|
// Use a sentinel bit to count zeros only up to bitCount
|
||||||
|
zeros = ctz(v0 | (1ULL << bitCount));
|
||||||
|
|
||||||
|
v0 >>= zeros;
|
||||||
|
vh >>= zeros;
|
||||||
|
*uu <<= zeros;
|
||||||
|
*uv <<= zeros;
|
||||||
|
bitCount -= zeros;
|
||||||
|
|
||||||
|
if (bitCount == 0)
|
||||||
|
break;
|
||||||
|
|
||||||
|
if (vh < uh) {
|
||||||
|
SWAP(w, uh, vh);
|
||||||
|
SWAP(x, u0, v0);
|
||||||
|
SWAP(y, *uu, *vu);
|
||||||
|
SWAP(z, *uv, *vv);
|
||||||
|
}
|
||||||
|
|
||||||
|
vh -= uh;
|
||||||
|
v0 -= u0;
|
||||||
|
*vv -= *uv;
|
||||||
|
*vu -= *uu;
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
__device__ void MatrixVecMulHalf(uint64_t dest[5], uint64_t u[5], uint64_t v[5], int64_t _11, int64_t _12, uint64_t* carry) {
|
||||||
|
|
||||||
|
uint64_t t1[NBBLOCK];
|
||||||
|
uint64_t t2[NBBLOCK];
|
||||||
|
uint64_t c1, c2;
|
||||||
|
|
||||||
|
c1 = IMultC(t1, u, _11);
|
||||||
|
c2 = IMultC(t2, v, _12);
|
||||||
|
|
||||||
|
UADDO(dest[0], t1[0], t2[0]);
|
||||||
|
UADDC(dest[1], t1[1], t2[1]);
|
||||||
|
UADDC(dest[2], t1[2], t2[2]);
|
||||||
|
UADDC(dest[3], t1[3], t2[3]);
|
||||||
|
UADDC(dest[4], t1[4], t2[4]);
|
||||||
|
UADD(*carry, c1, c2);
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
__device__ void MatrixVecMul(uint64_t u[5], uint64_t v[5], int64_t _11, int64_t _12, int64_t _21, int64_t _22) {
|
||||||
|
|
||||||
|
uint64_t t1[NBBLOCK];
|
||||||
|
uint64_t t2[NBBLOCK];
|
||||||
|
uint64_t t3[NBBLOCK];
|
||||||
|
uint64_t t4[NBBLOCK];
|
||||||
|
|
||||||
|
IMult(t1, u, _11);
|
||||||
|
IMult(t2, v, _12);
|
||||||
|
IMult(t3, u, _21);
|
||||||
|
IMult(t4, v, _22);
|
||||||
|
|
||||||
|
UADDO(u[0], t1[0], t2[0]);
|
||||||
|
UADDC(u[1], t1[1], t2[1]);
|
||||||
|
UADDC(u[2], t1[2], t2[2]);
|
||||||
|
UADDC(u[3], t1[3], t2[3]);
|
||||||
|
UADD(u[4], t1[4], t2[4]);
|
||||||
|
|
||||||
|
UADDO(v[0], t3[0], t4[0]);
|
||||||
|
UADDC(v[1], t3[1], t4[1]);
|
||||||
|
UADDC(v[2], t3[2], t4[2]);
|
||||||
|
UADDC(v[3], t3[3], t4[3]);
|
||||||
|
UADD(v[4], t3[4], t4[4]);
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
__device__ uint64_t AddCh(uint64_t r[5], uint64_t a[5], uint64_t carry) {
|
||||||
|
|
||||||
|
uint64_t carryOut;
|
||||||
|
|
||||||
|
UADDO1(r[0], a[0]);
|
||||||
|
UADDC1(r[1], a[1]);
|
||||||
|
UADDC1(r[2], a[2]);
|
||||||
|
UADDC1(r[3], a[3]);
|
||||||
|
UADDC1(r[4], a[4]);
|
||||||
|
UADD(carryOut, carry, 0ULL);
|
||||||
|
|
||||||
|
return carryOut;
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
__device__ __noinline__ void _ModInv(uint64_t* R) {
|
||||||
|
|
||||||
|
// Compute modular inverse of R mop P (using 320bits signed integer)
|
||||||
|
// 0 < this < P , P must be odd
|
||||||
|
// Return 0 if no inverse
|
||||||
|
// See IntMod.cpp for more info.
|
||||||
|
|
||||||
|
int64_t uu, uv, vu, vv;
|
||||||
|
uint64_t mr0, ms0;
|
||||||
|
int32_t pos = NBBLOCK - 1;
|
||||||
|
|
||||||
|
uint64_t u[NBBLOCK];
|
||||||
|
uint64_t v[NBBLOCK];
|
||||||
|
uint64_t r[NBBLOCK];
|
||||||
|
uint64_t s[NBBLOCK];
|
||||||
|
uint64_t tr[NBBLOCK];
|
||||||
|
uint64_t ts[NBBLOCK];
|
||||||
|
uint64_t r0[NBBLOCK];
|
||||||
|
uint64_t s0[NBBLOCK];
|
||||||
|
uint64_t carryR;
|
||||||
|
uint64_t carryS;
|
||||||
|
|
||||||
|
u[0] = 0xFFFFFFFEFFFFFC2F;
|
||||||
|
u[1] = 0xFFFFFFFFFFFFFFFF;
|
||||||
|
u[2] = 0xFFFFFFFFFFFFFFFF;
|
||||||
|
u[3] = 0xFFFFFFFFFFFFFFFF;
|
||||||
|
u[4] = 0;
|
||||||
|
Load(v, R);
|
||||||
|
r[0] = 0; s[0] = 1;
|
||||||
|
r[1] = 0; s[1] = 0;
|
||||||
|
r[2] = 0; s[2] = 0;
|
||||||
|
r[3] = 0; s[3] = 0;
|
||||||
|
r[4] = 0; s[4] = 0;
|
||||||
|
|
||||||
|
// Delayed right shift 62bits
|
||||||
|
|
||||||
|
// DivStep loop -------------------------------
|
||||||
|
|
||||||
|
while (true) {
|
||||||
|
|
||||||
|
_DivStep62(u, v, &pos, &uu, &uv, &vu, &vv);
|
||||||
|
|
||||||
|
MatrixVecMul(u, v, uu, uv, vu, vv);
|
||||||
|
|
||||||
|
if (_IsNegative(u)) {
|
||||||
|
Neg(u);
|
||||||
|
uu = -uu;
|
||||||
|
uv = -uv;
|
||||||
|
}
|
||||||
|
if (_IsNegative(v)) {
|
||||||
|
Neg(v);
|
||||||
|
vu = -vu;
|
||||||
|
vv = -vv;
|
||||||
|
}
|
||||||
|
|
||||||
|
ShiftR62(u);
|
||||||
|
ShiftR62(v);
|
||||||
|
|
||||||
|
// Update r
|
||||||
|
MatrixVecMulHalf(tr, r, s, uu, uv, &carryR);
|
||||||
|
mr0 = (tr[0] * MM64) & MSK62;
|
||||||
|
MulP(r0, mr0);
|
||||||
|
carryR = AddCh(tr, r0, carryR);
|
||||||
|
|
||||||
|
if (_IsZero(v)) {
|
||||||
|
|
||||||
|
ShiftR62(r, tr, carryR);
|
||||||
|
break;
|
||||||
|
|
||||||
|
}
|
||||||
|
else {
|
||||||
|
|
||||||
|
// Update s
|
||||||
|
MatrixVecMulHalf(ts, r, s, vu, vv, &carryS);
|
||||||
|
ms0 = (ts[0] * MM64) & MSK62;
|
||||||
|
MulP(s0, ms0);
|
||||||
|
carryS = AddCh(ts, s0, carryS);
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
ShiftR62(r, tr, carryR);
|
||||||
|
ShiftR62(s, ts, carryS);
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
// u ends with gcd
|
||||||
|
if (!_IsOne(u)) {
|
||||||
|
// No inverse
|
||||||
|
R[0] = 0ULL;
|
||||||
|
R[1] = 0ULL;
|
||||||
|
R[2] = 0ULL;
|
||||||
|
R[3] = 0ULL;
|
||||||
|
R[4] = 0ULL;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
while (_IsNegative(r))
|
||||||
|
AddP(r);
|
||||||
|
while (!_IsNegative(r))
|
||||||
|
SubP(r);
|
||||||
|
AddP(r);
|
||||||
|
|
||||||
|
Load(R, r);
|
||||||
|
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
661
examples/cuda-fanout/vendor/vanity-search-bitcrack/LICENSE.txt
vendored
Normal file
661
examples/cuda-fanout/vendor/vanity-search-bitcrack/LICENSE.txt
vendored
Normal file
|
|
@ -0,0 +1,661 @@
|
||||||
|
GNU AFFERO GENERAL PUBLIC LICENSE
|
||||||
|
Version 3, 19 November 2007
|
||||||
|
|
||||||
|
Copyright (C) 2007 Free Software Foundation, Inc. <https://fsf.org/>
|
||||||
|
Everyone is permitted to copy and distribute verbatim copies
|
||||||
|
of this license document, but changing it is not allowed.
|
||||||
|
|
||||||
|
Preamble
|
||||||
|
|
||||||
|
The GNU Affero General Public License is a free, copyleft license for
|
||||||
|
software and other kinds of works, specifically designed to ensure
|
||||||
|
cooperation with the community in the case of network server software.
|
||||||
|
|
||||||
|
The licenses for most software and other practical works are designed
|
||||||
|
to take away your freedom to share and change the works. By contrast,
|
||||||
|
our General Public Licenses are intended to guarantee your freedom to
|
||||||
|
share and change all versions of a program--to make sure it remains free
|
||||||
|
software for all its users.
|
||||||
|
|
||||||
|
When we speak of free software, we are referring to freedom, not
|
||||||
|
price. Our General Public Licenses are designed to make sure that you
|
||||||
|
have the freedom to distribute copies of free software (and charge for
|
||||||
|
them if you wish), that you receive source code or can get it if you
|
||||||
|
want it, that you can change the software or use pieces of it in new
|
||||||
|
free programs, and that you know you can do these things.
|
||||||
|
|
||||||
|
Developers that use our General Public Licenses protect your rights
|
||||||
|
with two steps: (1) assert copyright on the software, and (2) offer
|
||||||
|
you this License which gives you legal permission to copy, distribute
|
||||||
|
and/or modify the software.
|
||||||
|
|
||||||
|
A secondary benefit of defending all users' freedom is that
|
||||||
|
improvements made in alternate versions of the program, if they
|
||||||
|
receive widespread use, become available for other developers to
|
||||||
|
incorporate. Many developers of free software are heartened and
|
||||||
|
encouraged by the resulting cooperation. However, in the case of
|
||||||
|
software used on network servers, this result may fail to come about.
|
||||||
|
The GNU General Public License permits making a modified version and
|
||||||
|
letting the public access it on a server without ever releasing its
|
||||||
|
source code to the public.
|
||||||
|
|
||||||
|
The GNU Affero General Public License is designed specifically to
|
||||||
|
ensure that, in such cases, the modified source code becomes available
|
||||||
|
to the community. It requires the operator of a network server to
|
||||||
|
provide the source code of the modified version running there to the
|
||||||
|
users of that server. Therefore, public use of a modified version, on
|
||||||
|
a publicly accessible server, gives the public access to the source
|
||||||
|
code of the modified version.
|
||||||
|
|
||||||
|
An older license, called the Affero General Public License and
|
||||||
|
published by Affero, was designed to accomplish similar goals. This is
|
||||||
|
a different license, not a version of the Affero GPL, but Affero has
|
||||||
|
released a new version of the Affero GPL which permits relicensing under
|
||||||
|
this license.
|
||||||
|
|
||||||
|
The precise terms and conditions for copying, distribution and
|
||||||
|
modification follow.
|
||||||
|
|
||||||
|
TERMS AND CONDITIONS
|
||||||
|
|
||||||
|
0. Definitions.
|
||||||
|
|
||||||
|
"This License" refers to version 3 of the GNU Affero General Public License.
|
||||||
|
|
||||||
|
"Copyright" also means copyright-like laws that apply to other kinds of
|
||||||
|
works, such as semiconductor masks.
|
||||||
|
|
||||||
|
"The Program" refers to any copyrightable work licensed under this
|
||||||
|
License. Each licensee is addressed as "you". "Licensees" and
|
||||||
|
"recipients" may be individuals or organizations.
|
||||||
|
|
||||||
|
To "modify" a work means to copy from or adapt all or part of the work
|
||||||
|
in a fashion requiring copyright permission, other than the making of an
|
||||||
|
exact copy. The resulting work is called a "modified version" of the
|
||||||
|
earlier work or a work "based on" the earlier work.
|
||||||
|
|
||||||
|
A "covered work" means either the unmodified Program or a work based
|
||||||
|
on the Program.
|
||||||
|
|
||||||
|
To "propagate" a work means to do anything with it that, without
|
||||||
|
permission, would make you directly or secondarily liable for
|
||||||
|
infringement under applicable copyright law, except executing it on a
|
||||||
|
computer or modifying a private copy. Propagation includes copying,
|
||||||
|
distribution (with or without modification), making available to the
|
||||||
|
public, and in some countries other activities as well.
|
||||||
|
|
||||||
|
To "convey" a work means any kind of propagation that enables other
|
||||||
|
parties to make or receive copies. Mere interaction with a user through
|
||||||
|
a computer network, with no transfer of a copy, is not conveying.
|
||||||
|
|
||||||
|
An interactive user interface displays "Appropriate Legal Notices"
|
||||||
|
to the extent that it includes a convenient and prominently visible
|
||||||
|
feature that (1) displays an appropriate copyright notice, and (2)
|
||||||
|
tells the user that there is no warranty for the work (except to the
|
||||||
|
extent that warranties are provided), that licensees may convey the
|
||||||
|
work under this License, and how to view a copy of this License. If
|
||||||
|
the interface presents a list of user commands or options, such as a
|
||||||
|
menu, a prominent item in the list meets this criterion.
|
||||||
|
|
||||||
|
1. Source Code.
|
||||||
|
|
||||||
|
The "source code" for a work means the preferred form of the work
|
||||||
|
for making modifications to it. "Object code" means any non-source
|
||||||
|
form of a work.
|
||||||
|
|
||||||
|
A "Standard Interface" means an interface that either is an official
|
||||||
|
standard defined by a recognized standards body, or, in the case of
|
||||||
|
interfaces specified for a particular programming language, one that
|
||||||
|
is widely used among developers working in that language.
|
||||||
|
|
||||||
|
The "System Libraries" of an executable work include anything, other
|
||||||
|
than the work as a whole, that (a) is included in the normal form of
|
||||||
|
packaging a Major Component, but which is not part of that Major
|
||||||
|
Component, and (b) serves only to enable use of the work with that
|
||||||
|
Major Component, or to implement a Standard Interface for which an
|
||||||
|
implementation is available to the public in source code form. A
|
||||||
|
"Major Component", in this context, means a major essential component
|
||||||
|
(kernel, window system, and so on) of the specific operating system
|
||||||
|
(if any) on which the executable work runs, or a compiler used to
|
||||||
|
produce the work, or an object code interpreter used to run it.
|
||||||
|
|
||||||
|
The "Corresponding Source" for a work in object code form means all
|
||||||
|
the source code needed to generate, install, and (for an executable
|
||||||
|
work) run the object code and to modify the work, including scripts to
|
||||||
|
control those activities. However, it does not include the work's
|
||||||
|
System Libraries, or general-purpose tools or generally available free
|
||||||
|
programs which are used unmodified in performing those activities but
|
||||||
|
which are not part of the work. For example, Corresponding Source
|
||||||
|
includes interface definition files associated with source files for
|
||||||
|
the work, and the source code for shared libraries and dynamically
|
||||||
|
linked subprograms that the work is specifically designed to require,
|
||||||
|
such as by intimate data communication or control flow between those
|
||||||
|
subprograms and other parts of the work.
|
||||||
|
|
||||||
|
The Corresponding Source need not include anything that users
|
||||||
|
can regenerate automatically from other parts of the Corresponding
|
||||||
|
Source.
|
||||||
|
|
||||||
|
The Corresponding Source for a work in source code form is that
|
||||||
|
same work.
|
||||||
|
|
||||||
|
2. Basic Permissions.
|
||||||
|
|
||||||
|
All rights granted under this License are granted for the term of
|
||||||
|
copyright on the Program, and are irrevocable provided the stated
|
||||||
|
conditions are met. This License explicitly affirms your unlimited
|
||||||
|
permission to run the unmodified Program. The output from running a
|
||||||
|
covered work is covered by this License only if the output, given its
|
||||||
|
content, constitutes a covered work. This License acknowledges your
|
||||||
|
rights of fair use or other equivalent, as provided by copyright law.
|
||||||
|
|
||||||
|
You may make, run and propagate covered works that you do not
|
||||||
|
convey, without conditions so long as your license otherwise remains
|
||||||
|
in force. You may convey covered works to others for the sole purpose
|
||||||
|
of having them make modifications exclusively for you, or provide you
|
||||||
|
with facilities for running those works, provided that you comply with
|
||||||
|
the terms of this License in conveying all material for which you do
|
||||||
|
not control copyright. Those thus making or running the covered works
|
||||||
|
for you must do so exclusively on your behalf, under your direction
|
||||||
|
and control, on terms that prohibit them from making any copies of
|
||||||
|
your copyrighted material outside their relationship with you.
|
||||||
|
|
||||||
|
Conveying under any other circumstances is permitted solely under
|
||||||
|
the conditions stated below. Sublicensing is not allowed; section 10
|
||||||
|
makes it unnecessary.
|
||||||
|
|
||||||
|
3. Protecting Users' Legal Rights From Anti-Circumvention Law.
|
||||||
|
|
||||||
|
No covered work shall be deemed part of an effective technological
|
||||||
|
measure under any applicable law fulfilling obligations under article
|
||||||
|
11 of the WIPO copyright treaty adopted on 20 December 1996, or
|
||||||
|
similar laws prohibiting or restricting circumvention of such
|
||||||
|
measures.
|
||||||
|
|
||||||
|
When you convey a covered work, you waive any legal power to forbid
|
||||||
|
circumvention of technological measures to the extent such circumvention
|
||||||
|
is effected by exercising rights under this License with respect to
|
||||||
|
the covered work, and you disclaim any intention to limit operation or
|
||||||
|
modification of the work as a means of enforcing, against the work's
|
||||||
|
users, your or third parties' legal rights to forbid circumvention of
|
||||||
|
technological measures.
|
||||||
|
|
||||||
|
4. Conveying Verbatim Copies.
|
||||||
|
|
||||||
|
You may convey verbatim copies of the Program's source code as you
|
||||||
|
receive it, in any medium, provided that you conspicuously and
|
||||||
|
appropriately publish on each copy an appropriate copyright notice;
|
||||||
|
keep intact all notices stating that this License and any
|
||||||
|
non-permissive terms added in accord with section 7 apply to the code;
|
||||||
|
keep intact all notices of the absence of any warranty; and give all
|
||||||
|
recipients a copy of this License along with the Program.
|
||||||
|
|
||||||
|
You may charge any price or no price for each copy that you convey,
|
||||||
|
and you may offer support or warranty protection for a fee.
|
||||||
|
|
||||||
|
5. Conveying Modified Source Versions.
|
||||||
|
|
||||||
|
You may convey a work based on the Program, or the modifications to
|
||||||
|
produce it from the Program, in the form of source code under the
|
||||||
|
terms of section 4, provided that you also meet all of these conditions:
|
||||||
|
|
||||||
|
a) The work must carry prominent notices stating that you modified
|
||||||
|
it, and giving a relevant date.
|
||||||
|
|
||||||
|
b) The work must carry prominent notices stating that it is
|
||||||
|
released under this License and any conditions added under section
|
||||||
|
7. This requirement modifies the requirement in section 4 to
|
||||||
|
"keep intact all notices".
|
||||||
|
|
||||||
|
c) You must license the entire work, as a whole, under this
|
||||||
|
License to anyone who comes into possession of a copy. This
|
||||||
|
License will therefore apply, along with any applicable section 7
|
||||||
|
additional terms, to the whole of the work, and all its parts,
|
||||||
|
regardless of how they are packaged. This License gives no
|
||||||
|
permission to license the work in any other way, but it does not
|
||||||
|
invalidate such permission if you have separately received it.
|
||||||
|
|
||||||
|
d) If the work has interactive user interfaces, each must display
|
||||||
|
Appropriate Legal Notices; however, if the Program has interactive
|
||||||
|
interfaces that do not display Appropriate Legal Notices, your
|
||||||
|
work need not make them do so.
|
||||||
|
|
||||||
|
A compilation of a covered work with other separate and independent
|
||||||
|
works, which are not by their nature extensions of the covered work,
|
||||||
|
and which are not combined with it such as to form a larger program,
|
||||||
|
in or on a volume of a storage or distribution medium, is called an
|
||||||
|
"aggregate" if the compilation and its resulting copyright are not
|
||||||
|
used to limit the access or legal rights of the compilation's users
|
||||||
|
beyond what the individual works permit. Inclusion of a covered work
|
||||||
|
in an aggregate does not cause this License to apply to the other
|
||||||
|
parts of the aggregate.
|
||||||
|
|
||||||
|
6. Conveying Non-Source Forms.
|
||||||
|
|
||||||
|
You may convey a covered work in object code form under the terms
|
||||||
|
of sections 4 and 5, provided that you also convey the
|
||||||
|
machine-readable Corresponding Source under the terms of this License,
|
||||||
|
in one of these ways:
|
||||||
|
|
||||||
|
a) Convey the object code in, or embodied in, a physical product
|
||||||
|
(including a physical distribution medium), accompanied by the
|
||||||
|
Corresponding Source fixed on a durable physical medium
|
||||||
|
customarily used for software interchange.
|
||||||
|
|
||||||
|
b) Convey the object code in, or embodied in, a physical product
|
||||||
|
(including a physical distribution medium), accompanied by a
|
||||||
|
written offer, valid for at least three years and valid for as
|
||||||
|
long as you offer spare parts or customer support for that product
|
||||||
|
model, to give anyone who possesses the object code either (1) a
|
||||||
|
copy of the Corresponding Source for all the software in the
|
||||||
|
product that is covered by this License, on a durable physical
|
||||||
|
medium customarily used for software interchange, for a price no
|
||||||
|
more than your reasonable cost of physically performing this
|
||||||
|
conveying of source, or (2) access to copy the
|
||||||
|
Corresponding Source from a network server at no charge.
|
||||||
|
|
||||||
|
c) Convey individual copies of the object code with a copy of the
|
||||||
|
written offer to provide the Corresponding Source. This
|
||||||
|
alternative is allowed only occasionally and noncommercially, and
|
||||||
|
only if you received the object code with such an offer, in accord
|
||||||
|
with subsection 6b.
|
||||||
|
|
||||||
|
d) Convey the object code by offering access from a designated
|
||||||
|
place (gratis or for a charge), and offer equivalent access to the
|
||||||
|
Corresponding Source in the same way through the same place at no
|
||||||
|
further charge. You need not require recipients to copy the
|
||||||
|
Corresponding Source along with the object code. If the place to
|
||||||
|
copy the object code is a network server, the Corresponding Source
|
||||||
|
may be on a different server (operated by you or a third party)
|
||||||
|
that supports equivalent copying facilities, provided you maintain
|
||||||
|
clear directions next to the object code saying where to find the
|
||||||
|
Corresponding Source. Regardless of what server hosts the
|
||||||
|
Corresponding Source, you remain obligated to ensure that it is
|
||||||
|
available for as long as needed to satisfy these requirements.
|
||||||
|
|
||||||
|
e) Convey the object code using peer-to-peer transmission, provided
|
||||||
|
you inform other peers where the object code and Corresponding
|
||||||
|
Source of the work are being offered to the general public at no
|
||||||
|
charge under subsection 6d.
|
||||||
|
|
||||||
|
A separable portion of the object code, whose source code is excluded
|
||||||
|
from the Corresponding Source as a System Library, need not be
|
||||||
|
included in conveying the object code work.
|
||||||
|
|
||||||
|
A "User Product" is either (1) a "consumer product", which means any
|
||||||
|
tangible personal property which is normally used for personal, family,
|
||||||
|
or household purposes, or (2) anything designed or sold for incorporation
|
||||||
|
into a dwelling. In determining whether a product is a consumer product,
|
||||||
|
doubtful cases shall be resolved in favor of coverage. For a particular
|
||||||
|
product received by a particular user, "normally used" refers to a
|
||||||
|
typical or common use of that class of product, regardless of the status
|
||||||
|
of the particular user or of the way in which the particular user
|
||||||
|
actually uses, or expects or is expected to use, the product. A product
|
||||||
|
is a consumer product regardless of whether the product has substantial
|
||||||
|
commercial, industrial or non-consumer uses, unless such uses represent
|
||||||
|
the only significant mode of use of the product.
|
||||||
|
|
||||||
|
"Installation Information" for a User Product means any methods,
|
||||||
|
procedures, authorization keys, or other information required to install
|
||||||
|
and execute modified versions of a covered work in that User Product from
|
||||||
|
a modified version of its Corresponding Source. The information must
|
||||||
|
suffice to ensure that the continued functioning of the modified object
|
||||||
|
code is in no case prevented or interfered with solely because
|
||||||
|
modification has been made.
|
||||||
|
|
||||||
|
If you convey an object code work under this section in, or with, or
|
||||||
|
specifically for use in, a User Product, and the conveying occurs as
|
||||||
|
part of a transaction in which the right of possession and use of the
|
||||||
|
User Product is transferred to the recipient in perpetuity or for a
|
||||||
|
fixed term (regardless of how the transaction is characterized), the
|
||||||
|
Corresponding Source conveyed under this section must be accompanied
|
||||||
|
by the Installation Information. But this requirement does not apply
|
||||||
|
if neither you nor any third party retains the ability to install
|
||||||
|
modified object code on the User Product (for example, the work has
|
||||||
|
been installed in ROM).
|
||||||
|
|
||||||
|
The requirement to provide Installation Information does not include a
|
||||||
|
requirement to continue to provide support service, warranty, or updates
|
||||||
|
for a work that has been modified or installed by the recipient, or for
|
||||||
|
the User Product in which it has been modified or installed. Access to a
|
||||||
|
network may be denied when the modification itself materially and
|
||||||
|
adversely affects the operation of the network or violates the rules and
|
||||||
|
protocols for communication across the network.
|
||||||
|
|
||||||
|
Corresponding Source conveyed, and Installation Information provided,
|
||||||
|
in accord with this section must be in a format that is publicly
|
||||||
|
documented (and with an implementation available to the public in
|
||||||
|
source code form), and must require no special password or key for
|
||||||
|
unpacking, reading or copying.
|
||||||
|
|
||||||
|
7. Additional Terms.
|
||||||
|
|
||||||
|
"Additional permissions" are terms that supplement the terms of this
|
||||||
|
License by making exceptions from one or more of its conditions.
|
||||||
|
Additional permissions that are applicable to the entire Program shall
|
||||||
|
be treated as though they were included in this License, to the extent
|
||||||
|
that they are valid under applicable law. If additional permissions
|
||||||
|
apply only to part of the Program, that part may be used separately
|
||||||
|
under those permissions, but the entire Program remains governed by
|
||||||
|
this License without regard to the additional permissions.
|
||||||
|
|
||||||
|
When you convey a copy of a covered work, you may at your option
|
||||||
|
remove any additional permissions from that copy, or from any part of
|
||||||
|
it. (Additional permissions may be written to require their own
|
||||||
|
removal in certain cases when you modify the work.) You may place
|
||||||
|
additional permissions on material, added by you to a covered work,
|
||||||
|
for which you have or can give appropriate copyright permission.
|
||||||
|
|
||||||
|
Notwithstanding any other provision of this License, for material you
|
||||||
|
add to a covered work, you may (if authorized by the copyright holders of
|
||||||
|
that material) supplement the terms of this License with terms:
|
||||||
|
|
||||||
|
a) Disclaiming warranty or limiting liability differently from the
|
||||||
|
terms of sections 15 and 16 of this License; or
|
||||||
|
|
||||||
|
b) Requiring preservation of specified reasonable legal notices or
|
||||||
|
author attributions in that material or in the Appropriate Legal
|
||||||
|
Notices displayed by works containing it; or
|
||||||
|
|
||||||
|
c) Prohibiting misrepresentation of the origin of that material, or
|
||||||
|
requiring that modified versions of such material be marked in
|
||||||
|
reasonable ways as different from the original version; or
|
||||||
|
|
||||||
|
d) Limiting the use for publicity purposes of names of licensors or
|
||||||
|
authors of the material; or
|
||||||
|
|
||||||
|
e) Declining to grant rights under trademark law for use of some
|
||||||
|
trade names, trademarks, or service marks; or
|
||||||
|
|
||||||
|
f) Requiring indemnification of licensors and authors of that
|
||||||
|
material by anyone who conveys the material (or modified versions of
|
||||||
|
it) with contractual assumptions of liability to the recipient, for
|
||||||
|
any liability that these contractual assumptions directly impose on
|
||||||
|
those licensors and authors.
|
||||||
|
|
||||||
|
All other non-permissive additional terms are considered "further
|
||||||
|
restrictions" within the meaning of section 10. If the Program as you
|
||||||
|
received it, or any part of it, contains a notice stating that it is
|
||||||
|
governed by this License along with a term that is a further
|
||||||
|
restriction, you may remove that term. If a license document contains
|
||||||
|
a further restriction but permits relicensing or conveying under this
|
||||||
|
License, you may add to a covered work material governed by the terms
|
||||||
|
of that license document, provided that the further restriction does
|
||||||
|
not survive such relicensing or conveying.
|
||||||
|
|
||||||
|
If you add terms to a covered work in accord with this section, you
|
||||||
|
must place, in the relevant source files, a statement of the
|
||||||
|
additional terms that apply to those files, or a notice indicating
|
||||||
|
where to find the applicable terms.
|
||||||
|
|
||||||
|
Additional terms, permissive or non-permissive, may be stated in the
|
||||||
|
form of a separately written license, or stated as exceptions;
|
||||||
|
the above requirements apply either way.
|
||||||
|
|
||||||
|
8. Termination.
|
||||||
|
|
||||||
|
You may not propagate or modify a covered work except as expressly
|
||||||
|
provided under this License. Any attempt otherwise to propagate or
|
||||||
|
modify it is void, and will automatically terminate your rights under
|
||||||
|
this License (including any patent licenses granted under the third
|
||||||
|
paragraph of section 11).
|
||||||
|
|
||||||
|
However, if you cease all violation of this License, then your
|
||||||
|
license from a particular copyright holder is reinstated (a)
|
||||||
|
provisionally, unless and until the copyright holder explicitly and
|
||||||
|
finally terminates your license, and (b) permanently, if the copyright
|
||||||
|
holder fails to notify you of the violation by some reasonable means
|
||||||
|
prior to 60 days after the cessation.
|
||||||
|
|
||||||
|
Moreover, your license from a particular copyright holder is
|
||||||
|
reinstated permanently if the copyright holder notifies you of the
|
||||||
|
violation by some reasonable means, this is the first time you have
|
||||||
|
received notice of violation of this License (for any work) from that
|
||||||
|
copyright holder, and you cure the violation prior to 30 days after
|
||||||
|
your receipt of the notice.
|
||||||
|
|
||||||
|
Termination of your rights under this section does not terminate the
|
||||||
|
licenses of parties who have received copies or rights from you under
|
||||||
|
this License. If your rights have been terminated and not permanently
|
||||||
|
reinstated, you do not qualify to receive new licenses for the same
|
||||||
|
material under section 10.
|
||||||
|
|
||||||
|
9. Acceptance Not Required for Having Copies.
|
||||||
|
|
||||||
|
You are not required to accept this License in order to receive or
|
||||||
|
run a copy of the Program. Ancillary propagation of a covered work
|
||||||
|
occurring solely as a consequence of using peer-to-peer transmission
|
||||||
|
to receive a copy likewise does not require acceptance. However,
|
||||||
|
nothing other than this License grants you permission to propagate or
|
||||||
|
modify any covered work. These actions infringe copyright if you do
|
||||||
|
not accept this License. Therefore, by modifying or propagating a
|
||||||
|
covered work, you indicate your acceptance of this License to do so.
|
||||||
|
|
||||||
|
10. Automatic Licensing of Downstream Recipients.
|
||||||
|
|
||||||
|
Each time you convey a covered work, the recipient automatically
|
||||||
|
receives a license from the original licensors, to run, modify and
|
||||||
|
propagate that work, subject to this License. You are not responsible
|
||||||
|
for enforcing compliance by third parties with this License.
|
||||||
|
|
||||||
|
An "entity transaction" is a transaction transferring control of an
|
||||||
|
organization, or substantially all assets of one, or subdividing an
|
||||||
|
organization, or merging organizations. If propagation of a covered
|
||||||
|
work results from an entity transaction, each party to that
|
||||||
|
transaction who receives a copy of the work also receives whatever
|
||||||
|
licenses to the work the party's predecessor in interest had or could
|
||||||
|
give under the previous paragraph, plus a right to possession of the
|
||||||
|
Corresponding Source of the work from the predecessor in interest, if
|
||||||
|
the predecessor has it or can get it with reasonable efforts.
|
||||||
|
|
||||||
|
You may not impose any further restrictions on the exercise of the
|
||||||
|
rights granted or affirmed under this License. For example, you may
|
||||||
|
not impose a license fee, royalty, or other charge for exercise of
|
||||||
|
rights granted under this License, and you may not initiate litigation
|
||||||
|
(including a cross-claim or counterclaim in a lawsuit) alleging that
|
||||||
|
any patent claim is infringed by making, using, selling, offering for
|
||||||
|
sale, or importing the Program or any portion of it.
|
||||||
|
|
||||||
|
11. Patents.
|
||||||
|
|
||||||
|
A "contributor" is a copyright holder who authorizes use under this
|
||||||
|
License of the Program or a work on which the Program is based. The
|
||||||
|
work thus licensed is called the contributor's "contributor version".
|
||||||
|
|
||||||
|
A contributor's "essential patent claims" are all patent claims
|
||||||
|
owned or controlled by the contributor, whether already acquired or
|
||||||
|
hereafter acquired, that would be infringed by some manner, permitted
|
||||||
|
by this License, of making, using, or selling its contributor version,
|
||||||
|
but do not include claims that would be infringed only as a
|
||||||
|
consequence of further modification of the contributor version. For
|
||||||
|
purposes of this definition, "control" includes the right to grant
|
||||||
|
patent sublicenses in a manner consistent with the requirements of
|
||||||
|
this License.
|
||||||
|
|
||||||
|
Each contributor grants you a non-exclusive, worldwide, royalty-free
|
||||||
|
patent license under the contributor's essential patent claims, to
|
||||||
|
make, use, sell, offer for sale, import and otherwise run, modify and
|
||||||
|
propagate the contents of its contributor version.
|
||||||
|
|
||||||
|
In the following three paragraphs, a "patent license" is any express
|
||||||
|
agreement or commitment, however denominated, not to enforce a patent
|
||||||
|
(such as an express permission to practice a patent or covenant not to
|
||||||
|
sue for patent infringement). To "grant" such a patent license to a
|
||||||
|
party means to make such an agreement or commitment not to enforce a
|
||||||
|
patent against the party.
|
||||||
|
|
||||||
|
If you convey a covered work, knowingly relying on a patent license,
|
||||||
|
and the Corresponding Source of the work is not available for anyone
|
||||||
|
to copy, free of charge and under the terms of this License, through a
|
||||||
|
publicly available network server or other readily accessible means,
|
||||||
|
then you must either (1) cause the Corresponding Source to be so
|
||||||
|
available, or (2) arrange to deprive yourself of the benefit of the
|
||||||
|
patent license for this particular work, or (3) arrange, in a manner
|
||||||
|
consistent with the requirements of this License, to extend the patent
|
||||||
|
license to downstream recipients. "Knowingly relying" means you have
|
||||||
|
actual knowledge that, but for the patent license, your conveying the
|
||||||
|
covered work in a country, or your recipient's use of the covered work
|
||||||
|
in a country, would infringe one or more identifiable patents in that
|
||||||
|
country that you have reason to believe are valid.
|
||||||
|
|
||||||
|
If, pursuant to or in connection with a single transaction or
|
||||||
|
arrangement, you convey, or propagate by procuring conveyance of, a
|
||||||
|
covered work, and grant a patent license to some of the parties
|
||||||
|
receiving the covered work authorizing them to use, propagate, modify
|
||||||
|
or convey a specific copy of the covered work, then the patent license
|
||||||
|
you grant is automatically extended to all recipients of the covered
|
||||||
|
work and works based on it.
|
||||||
|
|
||||||
|
A patent license is "discriminatory" if it does not include within
|
||||||
|
the scope of its coverage, prohibits the exercise of, or is
|
||||||
|
conditioned on the non-exercise of one or more of the rights that are
|
||||||
|
specifically granted under this License. You may not convey a covered
|
||||||
|
work if you are a party to an arrangement with a third party that is
|
||||||
|
in the business of distributing software, under which you make payment
|
||||||
|
to the third party based on the extent of your activity of conveying
|
||||||
|
the work, and under which the third party grants, to any of the
|
||||||
|
parties who would receive the covered work from you, a discriminatory
|
||||||
|
patent license (a) in connection with copies of the covered work
|
||||||
|
conveyed by you (or copies made from those copies), or (b) primarily
|
||||||
|
for and in connection with specific products or compilations that
|
||||||
|
contain the covered work, unless you entered into that arrangement,
|
||||||
|
or that patent license was granted, prior to 28 March 2007.
|
||||||
|
|
||||||
|
Nothing in this License shall be construed as excluding or limiting
|
||||||
|
any implied license or other defenses to infringement that may
|
||||||
|
otherwise be available to you under applicable patent law.
|
||||||
|
|
||||||
|
12. No Surrender of Others' Freedom.
|
||||||
|
|
||||||
|
If conditions are imposed on you (whether by court order, agreement or
|
||||||
|
otherwise) that contradict the conditions of this License, they do not
|
||||||
|
excuse you from the conditions of this License. If you cannot convey a
|
||||||
|
covered work so as to satisfy simultaneously your obligations under this
|
||||||
|
License and any other pertinent obligations, then as a consequence you may
|
||||||
|
not convey it at all. For example, if you agree to terms that obligate you
|
||||||
|
to collect a royalty for further conveying from those to whom you convey
|
||||||
|
the Program, the only way you could satisfy both those terms and this
|
||||||
|
License would be to refrain entirely from conveying the Program.
|
||||||
|
|
||||||
|
13. Remote Network Interaction; Use with the GNU General Public License.
|
||||||
|
|
||||||
|
Notwithstanding any other provision of this License, if you modify the
|
||||||
|
Program, your modified version must prominently offer all users
|
||||||
|
interacting with it remotely through a computer network (if your version
|
||||||
|
supports such interaction) an opportunity to receive the Corresponding
|
||||||
|
Source of your version by providing access to the Corresponding Source
|
||||||
|
from a network server at no charge, through some standard or customary
|
||||||
|
means of facilitating copying of software. This Corresponding Source
|
||||||
|
shall include the Corresponding Source for any work covered by version 3
|
||||||
|
of the GNU General Public License that is incorporated pursuant to the
|
||||||
|
following paragraph.
|
||||||
|
|
||||||
|
Notwithstanding any other provision of this License, you have
|
||||||
|
permission to link or combine any covered work with a work licensed
|
||||||
|
under version 3 of the GNU General Public License into a single
|
||||||
|
combined work, and to convey the resulting work. The terms of this
|
||||||
|
License will continue to apply to the part which is the covered work,
|
||||||
|
but the work with which it is combined will remain governed by version
|
||||||
|
3 of the GNU General Public License.
|
||||||
|
|
||||||
|
14. Revised Versions of this License.
|
||||||
|
|
||||||
|
The Free Software Foundation may publish revised and/or new versions of
|
||||||
|
the GNU Affero General Public License from time to time. Such new versions
|
||||||
|
will be similar in spirit to the present version, but may differ in detail to
|
||||||
|
address new problems or concerns.
|
||||||
|
|
||||||
|
Each version is given a distinguishing version number. If the
|
||||||
|
Program specifies that a certain numbered version of the GNU Affero General
|
||||||
|
Public License "or any later version" applies to it, you have the
|
||||||
|
option of following the terms and conditions either of that numbered
|
||||||
|
version or of any later version published by the Free Software
|
||||||
|
Foundation. If the Program does not specify a version number of the
|
||||||
|
GNU Affero General Public License, you may choose any version ever published
|
||||||
|
by the Free Software Foundation.
|
||||||
|
|
||||||
|
If the Program specifies that a proxy can decide which future
|
||||||
|
versions of the GNU Affero General Public License can be used, that proxy's
|
||||||
|
public statement of acceptance of a version permanently authorizes you
|
||||||
|
to choose that version for the Program.
|
||||||
|
|
||||||
|
Later license versions may give you additional or different
|
||||||
|
permissions. However, no additional obligations are imposed on any
|
||||||
|
author or copyright holder as a result of your choosing to follow a
|
||||||
|
later version.
|
||||||
|
|
||||||
|
15. Disclaimer of Warranty.
|
||||||
|
|
||||||
|
THERE IS NO WARRANTY FOR THE PROGRAM, TO THE EXTENT PERMITTED BY
|
||||||
|
APPLICABLE LAW. EXCEPT WHEN OTHERWISE STATED IN WRITING THE COPYRIGHT
|
||||||
|
HOLDERS AND/OR OTHER PARTIES PROVIDE THE PROGRAM "AS IS" WITHOUT WARRANTY
|
||||||
|
OF ANY KIND, EITHER EXPRESSED OR IMPLIED, INCLUDING, BUT NOT LIMITED TO,
|
||||||
|
THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
|
||||||
|
PURPOSE. THE ENTIRE RISK AS TO THE QUALITY AND PERFORMANCE OF THE PROGRAM
|
||||||
|
IS WITH YOU. SHOULD THE PROGRAM PROVE DEFECTIVE, YOU ASSUME THE COST OF
|
||||||
|
ALL NECESSARY SERVICING, REPAIR OR CORRECTION.
|
||||||
|
|
||||||
|
16. Limitation of Liability.
|
||||||
|
|
||||||
|
IN NO EVENT UNLESS REQUIRED BY APPLICABLE LAW OR AGREED TO IN WRITING
|
||||||
|
WILL ANY COPYRIGHT HOLDER, OR ANY OTHER PARTY WHO MODIFIES AND/OR CONVEYS
|
||||||
|
THE PROGRAM AS PERMITTED ABOVE, BE LIABLE TO YOU FOR DAMAGES, INCLUDING ANY
|
||||||
|
GENERAL, SPECIAL, INCIDENTAL OR CONSEQUENTIAL DAMAGES ARISING OUT OF THE
|
||||||
|
USE OR INABILITY TO USE THE PROGRAM (INCLUDING BUT NOT LIMITED TO LOSS OF
|
||||||
|
DATA OR DATA BEING RENDERED INACCURATE OR LOSSES SUSTAINED BY YOU OR THIRD
|
||||||
|
PARTIES OR A FAILURE OF THE PROGRAM TO OPERATE WITH ANY OTHER PROGRAMS),
|
||||||
|
EVEN IF SUCH HOLDER OR OTHER PARTY HAS BEEN ADVISED OF THE POSSIBILITY OF
|
||||||
|
SUCH DAMAGES.
|
||||||
|
|
||||||
|
17. Interpretation of Sections 15 and 16.
|
||||||
|
|
||||||
|
If the disclaimer of warranty and limitation of liability provided
|
||||||
|
above cannot be given local legal effect according to their terms,
|
||||||
|
reviewing courts shall apply local law that most closely approximates
|
||||||
|
an absolute waiver of all civil liability in connection with the
|
||||||
|
Program, unless a warranty or assumption of liability accompanies a
|
||||||
|
copy of the Program in return for a fee.
|
||||||
|
|
||||||
|
END OF TERMS AND CONDITIONS
|
||||||
|
|
||||||
|
How to Apply These Terms to Your New Programs
|
||||||
|
|
||||||
|
If you develop a new program, and you want it to be of the greatest
|
||||||
|
possible use to the public, the best way to achieve this is to make it
|
||||||
|
free software which everyone can redistribute and change under these terms.
|
||||||
|
|
||||||
|
To do so, attach the following notices to the program. It is safest
|
||||||
|
to attach them to the start of each source file to most effectively
|
||||||
|
state the exclusion of warranty; and each file should have at least
|
||||||
|
the "copyright" line and a pointer to where the full notice is found.
|
||||||
|
|
||||||
|
<one line to give the program's name and a brief idea of what it does.>
|
||||||
|
Copyright (C) <year> <name of author>
|
||||||
|
|
||||||
|
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 <https://www.gnu.org/licenses/>.
|
||||||
|
|
||||||
|
Also add information on how to contact you by electronic and paper mail.
|
||||||
|
|
||||||
|
If your software can interact with users remotely through a computer
|
||||||
|
network, you should also make sure that it provides a way for users to
|
||||||
|
get its source. For example, if your program is a web application, its
|
||||||
|
interface could display a "Source" link that leads users to an archive
|
||||||
|
of the code. There are many ways you could offer source, and different
|
||||||
|
solutions will be better for different programs; see section 13 for the
|
||||||
|
specific requirements.
|
||||||
|
|
||||||
|
You should also get your employer (if you work as a programmer) or school,
|
||||||
|
if any, to sign a "copyright disclaimer" for the program, if necessary.
|
||||||
|
For more information on this, and how to apply and follow the GNU AGPL, see
|
||||||
|
<https://www.gnu.org/licenses/>.
|
||||||
39
examples/cuda-fanout/vendor/vanity-search-bitcrack/NOTICE
vendored
Normal file
39
examples/cuda-fanout/vendor/vanity-search-bitcrack/NOTICE
vendored
Normal file
|
|
@ -0,0 +1,39 @@
|
||||||
|
NOTICE -- VanitySearch-Bitcrack vendored headers
|
||||||
|
=================================================
|
||||||
|
|
||||||
|
Upstream source:
|
||||||
|
https://github.com/FixedPaul/VanitySearch-Bitcrack
|
||||||
|
Commit: 66e6f9deda607afc5eaf09731725664f0152fc92
|
||||||
|
License: AGPL-3.0 (LICENSE.txt, vendored alongside)
|
||||||
|
|
||||||
|
The FixedPaul fork in turn descends from JeanLucPons/VanitySearch
|
||||||
|
(https://github.com/JeanLucPons/VanitySearch, GPL-3.0). Each vendored
|
||||||
|
file carries its original "Copyright (c) 2019 Jean Luc PONS" header
|
||||||
|
verbatim per §4 of the license.
|
||||||
|
|
||||||
|
Files vendored
|
||||||
|
--------------
|
||||||
|
|
||||||
|
GPUMath.h
|
||||||
|
secp256k1 field arithmetic for CUDA: 256-bit modular multiply,
|
||||||
|
square, inverse, subtract, negate. Inline PTX. Used as-is by our
|
||||||
|
secp256k1-batch-mul.cu kernel.
|
||||||
|
|
||||||
|
What we do NOT vendor
|
||||||
|
---------------------
|
||||||
|
|
||||||
|
We do not copy GPUEngine.cu, GPUGroup.h, GPUCompute.h, GPUHash.h,
|
||||||
|
GPUBase58.h, GPUWildcard.h, RCGpuUtils.h, or any host-side .cpp file.
|
||||||
|
|
||||||
|
Our kernel computes scalar-times-G via per-thread Jacobian
|
||||||
|
double-and-add with batched Z-inversion at the end — independent of
|
||||||
|
the upstream GRP_SIZE/Shamir keyspace-stride structure. We need only
|
||||||
|
GPUMath.h's field-arithmetic primitives, not the windowed-comb tables
|
||||||
|
or address-lookup pipeline.
|
||||||
|
|
||||||
|
License effect
|
||||||
|
--------------
|
||||||
|
|
||||||
|
Our binary (secp256k1-batch-mul) is a combined work whose vendored
|
||||||
|
portions stay under AGPL-3.0. Our wrapper source (secp256k1-batch-mul.cu,
|
||||||
|
Makefile target, tests) ships under AGPL-3.0 to match.
|
||||||
Loading…
Add table
Add a link
Reference in a new issue