Optimize gpu sigverify

ci/buildkite.yml

@@ -1,6 +1,6 @@
 steps:
   - command: "ci/build.sh"
     name: "build"
-    timeout_in_minutes: 30
+    timeout_in_minutes: 45
     agents:
       - "queue=cuda"

src/cuda-ecc-ed25519/common.cu

@@ -23,4 +23,19 @@ static uint64_t __host__ __device__ load_4(const unsigned char *in) {
     return result;
 }
 
+static uint64_t __host__ __device__ load_7(const unsigned char *in) {
+    uint64_t result;
+
+    result = (uint64_t) in[0];
+    result |= ((uint64_t) in[1]) << 8;
+    result |= ((uint64_t) in[2]) << 16;
+    result |= ((uint64_t) in[3]) << 24;
+    result |= ((uint64_t) in[4]) << 32;
+    result |= ((uint64_t) in[5]) << 40;
+    result |= ((uint64_t) in[6]) << 48;
+
+    return result;
+}
+
+
 #endif

src/cuda-ecc-ed25519/ed25519.h

@@ -75,6 +75,10 @@ bool ED25519_DECLSPEC ed25519_init();
 int cuda_host_register(void* ptr, size_t size, unsigned int flags);
 int cuda_host_unregister(void* ptr);
 
+int ED25519_DECLSPEC ed25519_get_checked_scalar(unsigned char* out_scalar, const unsigned char* in_scalar);
+
+int ED25519_DECLSPEC ed25519_check_packed_ge_small_order(const unsigned char* packed_group_element);
+
 #ifdef __cplusplus
 }
 #endif

src/cuda-ecc-ed25519/fe.cu

@@ -43,6 +43,28 @@ void __device__ __host__ fe_1(fe h) {
 }
 
 
+int __host__ __device__ fe_is_1(fe h) {
+    if (h[0] != 1) {
+        return 0;
+    }
+    for (int i = 1; i < 9; i++) {
+        if (h[i] != 0) {
+            return 0;
+        }
+    }
+    return 1;
+}
+
+int __host__ __device__ fe_is_0(fe h) {
+    for (int i = 0; i < 9; i++) {
+        if (h[i] != 0) {
+            return 0;
+        }
+    }
+    return 1;
+}
+
+
 
 /*
     h = f + g

src/cuda-ecc-ed25519/fe.h

@@ -18,6 +18,8 @@ typedef int32_t fe[10];
 
 void __host__ __device__ fe_0(fe h);
 void __device__ __host__ fe_1(fe h);
+int __device__ __host__ fe_is_0(fe h);
+int __device__ __host__ fe_is_1(fe h);
 
 void __device__ __host__ fe_frombytes(fe h, const unsigned char *s);
 void __device__ __host__ fe_tobytes(unsigned char *s, const fe h);

src/cuda-ecc-ed25519/ge.cu

@@ -180,6 +180,29 @@ int __device__ __host__ ge_frombytes_negate_vartime(ge_p3 *h, const unsigned cha
     return 0;
 }
 
+// x = 1, y = 0, z = 0, t = 1
+int __host__ __device__ ge_is_identity(ge_p3* p) {
+    return (fe_is_0(p->X) &&
+            fe_is_1(p->Y) &&
+            fe_is_1(p->Z) &&
+            fe_is_0(p->T)) ? 1 : 0;
+}
+
+int __host__ __device__ ge_is_small_order(ge_p3* p) {
+    ge_p1p1 r;
+    ge_p2 s;
+    ge_p3 q;
+
+    // calculate q = p * 2*3
+    ge_p3_dbl(&r, p);
+    ge_p1p1_to_p2(&s, &r);
+    ge_p2_dbl(&r, &s);
+    ge_p1p1_to_p2(&s, &r);
+    ge_p2_dbl(&r, &s);
+    ge_p1p1_to_p3(&q, &r);
+
+    return ge_is_identity(&q);
+}
 
 /*
 r = p + q

src/cuda-ecc-ed25519/int128.h

@@ -0,0 +1,37 @@
+#ifndef INT128_H
+#define INT128_H
+
+struct uint128_t {
+  uint64_t low;
+  uint64_t high;
+};
+
+static __device__ __host__ uint128_t mul_128(uint64_t a, uint64_t b) {
+  uint128_t result;
+#ifdef __CUDA_ARCH__
+  result.low = a * b;
+  result.high = __mul64hi(a, b);
+#elif __x86_64__
+  asm( "mulq %3\n\t"
+      : "=a" (result.low), "=d" (result.high)
+      : "%0" (a), "rm" (b));
+#endif
+  return result;
+}
+
+static __device__ __host__ uint128_t add_128(uint128_t a, uint128_t b) {
+  uint128_t result;
+#ifdef __CUDA_ARCH__
+  asm( "add.cc.u64    %0, %2, %4;\n\t"
+       "addc.u64      %1, %3, %5;\n\t"
+       : "=l" (result.low), "=l" (result.high)
+       : "l" (a.low), "l" (a.high),
+       "l" (b.low), "l" (b.high));
+#else
+  result.low = a.low + b.low;
+  result.high = a.high + b.high + (result.low < a.low);
+#endif
+  return result;
+}
+
+#endif

src/cuda-ecc-ed25519/license.txt

@@ -14,3 +14,34 @@ applications, and to alter it and redistribute it freely, subject to the followi
    being the original software.
 
 3. This notice may not be removed or altered from any source distribution.
+
+================================
+
+Copyright (c) 2017-2019 isis agora lovecruft. All rights reserved.
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions are
+met:
+
+1. Redistributions of source code must retain the above copyright
+notice, this list of conditions and the following disclaimer.
+
+2. Redistributions in binary form must reproduce the above copyright
+notice, this list of conditions and the following disclaimer in the
+documentation and/or other materials provided with the distribution.
+
+3. Neither the name of the copyright holder nor the names of its
+contributors may be used to endorse or promote products derived from
+this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS
+IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A
+PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
+HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED
+TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.