rename ec_pmul_random_init to ec_pmul_init
This commit is contained in:
parent
2cf36a36f6
commit
f3ab474f24
@ -11,7 +11,7 @@ Fixed-point multiplication uses precomputed multiples of the secp256k1 generator
|
||||
|
||||
These precomputed values are stored in `test/ecdsa_constants.h` as `G1_1_G1SECP256K1.SIG_AFF[]` and loaded into device constant memory `ECDSACONST.d_mul_table[]` during `initialize()`.
|
||||
|
||||
**The problem**: The batch kernel `arith::fixedPMulByCombinedDAA` expects R1 to contain these precomputed multiples, but the current test calls `ec_pmul_random_init()` which overwrites R1 with input points instead.
|
||||
**The problem**: The batch kernel `arith::fixedPMulByCombinedDAA` expects R1 to contain these precomputed multiples, but the current test calls `ec_pmul_init()` which overwrites R1 with input points instead.
|
||||
|
||||
## Two Fixed-Point Implementations
|
||||
|
||||
@ -50,9 +50,9 @@ __global__ void fixedPMulByCombinedDAA(typename EC::Base *R0,
|
||||
|
||||
```cpp
|
||||
// test/ecdsa_ec_fixed_pmul.cu
|
||||
solver.ec_pmul_random_init(RANDOM_S, RANDOM_KEY_X, RANDOM_KEY_Y, count);
|
||||
solver.ec_pmul_init(RANDOM_S, RANDOM_KEY_X, RANDOM_KEY_Y, count);
|
||||
↓
|
||||
// include/gecc/ecdsa/gsv.h:ec_pmul_random_init()
|
||||
// include/gecc/ecdsa/gsv.h:ec_pmul_init()
|
||||
processScalarPoint<<<>>>(..., R1, ...); // ← Fills R1 with input points (WRONG!)
|
||||
↓
|
||||
solver.ecdsa_ec_pmul(MAX_SM_NUMS, 256, false); // false = fixed-point
|
||||
|
||||
@ -4,7 +4,7 @@
|
||||
The fixed-point multiplication test in `test/ecdsa_ec_fixed_pmul.cu` was producing incorrect results.
|
||||
|
||||
## Root Cause
|
||||
The test was calling `ec_pmul_random_init()` which populated R1 with arbitrary input points, but the batch kernel `fixedPMulByCombinedDAA` expected R1 to contain precomputed multiples of the generator G.
|
||||
The test was calling `ec_pmul_init()` which populated R1 with arbitrary input points, but the batch kernel `fixedPMulByCombinedDAA` expected R1 to contain precomputed multiples of the generator G.
|
||||
|
||||
## Solution
|
||||
Created a new test kernel that directly uses the `fixed_point_mult()` device function, which correctly accesses the precomputed table from device constant memory (`ECDSACONST.d_mul_table[]`).
|
||||
@ -21,7 +21,7 @@ Created a new test kernel that directly uses the `fixed_point_mult()` device fun
|
||||
- Stores results properly
|
||||
|
||||
3. **Rewrote correctness test** `test_ecdsa_ec_fixed_pmul_correctness()`:
|
||||
- Allocates memory directly (no `ec_pmul_random_init()`)
|
||||
- Allocates memory directly (no `ec_pmul_init()`)
|
||||
- Calls the new test kernel
|
||||
- Reads results and prints them
|
||||
|
||||
|
||||
@ -1529,8 +1529,8 @@ template <typename BaseField, typename BaseOrder, typename EC, const ECDSAConsta
|
||||
}
|
||||
|
||||
|
||||
// batch EC PMUL Breakdown Test
|
||||
void ec_pmul_random_init(const u64 s[][MAX_LIMBS], const u64 key_x[][MAX_LIMBS], const u64 key_y[][MAX_LIMBS], u32 count) {
|
||||
// batch EC PMUL initialization
|
||||
void ec_pmul_init(const u64 s[][MAX_LIMBS], const u64 key_x[][MAX_LIMBS], const u64 key_y[][MAX_LIMBS], u32 count) {
|
||||
verify_count = count;
|
||||
cudaMallocManaged(&verify_s, Order::SIZE * count); //25MB
|
||||
cudaMallocManaged(&verify_t, Order::SIZE * count);
|
||||
|
||||
@ -30,7 +30,7 @@ void test_ecdsa_ec_fixed_pmul() {
|
||||
printf("--------------------------- %u (%d << %d) --------------------------\n", count, MAX_SM_NUMS, i);
|
||||
|
||||
// solver.verify_init(R, S, E, KEY_X, KEY_Y, count);
|
||||
solver.ec_pmul_random_init(RANDOM_S, RANDOM_KEY_X, RANDOM_KEY_Y, count);
|
||||
solver.ec_pmul_init(RANDOM_S, RANDOM_KEY_X, RANDOM_KEY_Y, count);
|
||||
// warm up
|
||||
solver.ecdsa_ec_pmul(MAX_SM_NUMS<<2, 256, false);
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
@ -23,7 +23,7 @@ void test_ecdsa_ec_unknown_pmul() {
|
||||
count = MAX_SM_NUMS * (1<<i); //1<<18 ~ 1<<23
|
||||
printf("--------------------------- %u (%d << %d) --------------------------\n", count, MAX_SM_NUMS, i);
|
||||
|
||||
solver.ec_pmul_random_init(RANDOM_S, RANDOM_KEY_X, RANDOM_KEY_Y, count);
|
||||
solver.ec_pmul_init(RANDOM_S, RANDOM_KEY_X, RANDOM_KEY_Y, count);
|
||||
// warm up
|
||||
solver.ecdsa_ec_pmul(MAX_SM_NUMS<<2, 256, true);
|
||||
cudaDeviceSynchronize();
|
||||
@ -68,7 +68,7 @@ void test_ecdsa_ec_unknown_pmul_correctness() {
|
||||
printf("Testing %u point multiplications\n\n", count);
|
||||
|
||||
// Initialize with test data
|
||||
solver.ec_pmul_random_init(RANDOM_S, RANDOM_KEY_X, RANDOM_KEY_Y, count);
|
||||
solver.ec_pmul_init(RANDOM_S, RANDOM_KEY_X, RANDOM_KEY_Y, count);
|
||||
|
||||
// Print inputs (first 3 samples)
|
||||
printf("Input scalars (s):\n");
|
||||
|
||||
@ -24,7 +24,7 @@ void test_modinv_in_data_parallel() {
|
||||
printf("--------------------------- %u (~1<< %d) --------------------------\n", count, ((int)log2(MAX_SM_NUMS)) + i);
|
||||
|
||||
// solver.verify_init(R, S, E, KEY_X, KEY_Y, count);
|
||||
solver.ec_pmul_random_init(RANDOM_S, RANDOM_KEY_X, RANDOM_KEY_Y, count);
|
||||
solver.ec_pmul_init(RANDOM_S, RANDOM_KEY_X, RANDOM_KEY_Y, count);
|
||||
// warm up
|
||||
solver.batch_modinv_MTA(MAX_SM_NUMS<<2, 256);
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
Loading…
Reference in New Issue
Block a user