added implementation to poh_verify_many

Author: alexNgr

src/Makefile

@@ -13,8 +13,10 @@ V=debug
 include gpu-common.mk
 
 ECC_TEST_BIN=cuda_ed25519_verify
+POH_VERIFY_MANY_TEST_BIN=cuda_poh_verify_many
 LIB=cuda-crypt
 
+CL_POH_VERIFY_MANY_TEST_BIN=cl_poh_verify_many
 CL_ECC_TEST_BIN=cl_ed25519_verify
 CL_LIB=cl-crypt
 
@@ -25,6 +27,8 @@ CUDA_SHA256_DIR:=cuda-sha256
 
 CUDA_DIR ?= /usr/local/cuda
 
+POH_VERIFY_TEST_DIR=poh-verify-test
+
 CXX ?= g++
 
 CFLAGS_COMMON:=-DENDIAN_NEUTRAL -DLTC_NO_ASM
@@ -39,11 +43,14 @@ CL_CFLAGS_debug:=$(CL_CFLAGS_common) -O0 -g
 CL_CFLAGS:=$(CL_CFLAGS_$V)
 
 all: $(V)/$(ECC_TEST_BIN) \
+     $(V)/$(POH_VERIFY_MANY_TEST_BIN) \
      $(V)/$(CL_ECC_TEST_BIN) \
+	 $(V)/$(CL_POH_VERIFY_MANY_TEST_BIN) \
      $(V)/lib$(LIB).so \
      $(V)/lib$(CL_LIB).so
 
 ECC_DIR:=cuda-ecc-ed25519
+POH_VERIFY_MANY_DIR:=cuda-poh-verify
 
 KEYPAIR_SRCS:=$(addprefix $(ECC_DIR)/,keypair.cu ed25519.h ge.h)
 $V/keypair.o: $(KEYPAIR_SRCS)
@@ -112,9 +119,16 @@ $V/cl_ecc_main.o: $(CL_ECC_DIR)/main.cpp $(ECC_DIR)/ed25519.h
 	@mkdir -p $(@D)
 	$(CXX) $(CL_CFLAGS) -pthread -I$(ECC_DIR) -c $< -o $@
 
+$V/poh_many_main.o: $(POH_VERIFY_TEST_DIR)/main.cpp
+	@mkdir -p $(@D)
+	$(CXX) $(CL_CFLAGS) -pthread -I$(ECC_DIR) -c $< -o $@
+
 $V/$(CL_ECC_TEST_BIN): $V/cl_ecc_main.o $V/lib$(CL_LIB).so
 	$(CXX) $(CL_CFLAGS) -L$(CUDA_DIR)/lib64 -L$V -pthread $< -l$(CL_LIB) -lOpenCL -o $@
 
+$V/$(CL_POH_VERIFY_MANY_TEST_BIN): $V/poh_many_main.o $V/lib$(CL_LIB).so
+	$(CXX) $(CL_CFLAGS) -L$(CUDA_DIR)/lib64 -L$V -pthread $< -l$(CL_LIB) -lOpenCL -o $@
+
 CPU_GPU_OBJS=$(addprefix $V/,verify.o poh_verify.o gpu_ctx.o sign.o seed.o keypair.o)
 
 $V/crypt-dlink.o: $(CPU_GPU_OBJS)
@@ -130,6 +144,10 @@ $V/ecc_main.o: $(addprefix $(ECC_DIR)/,main.cu ed25519.h)
 $V/$(ECC_TEST_BIN): $V/ecc_main.o $V/lib$(LIB).so
 	$(NVCC) $(CFLAGS) -L$V -l$(LIB) $< -o $@
 
+$V/$(POH_VERIFY_MANY_TEST_BIN): $V/poh_many_main.o $V/lib$(LIB).so
+	$(NVCC) $(CFLAGS) -L$V -l$(LIB) $< -o $@
+
+
 .PHONY:clean
 clean:
 	rm -rf $V

src/cuda-poh-verify/poh_verify.cu

@@ -22,8 +22,7 @@ __global__ void poh_verify_kernel(uint8_t* hashes, uint64_t* num_hashes_arr, siz
         sha256_init(&sha_state);
         sha256_process(&sha_state, hash, SHA256_BLOCK_SIZE);
         sha256_done(&sha_state, hash);
-    }
-
+    }  
     memcpy(&hashes[idx * SHA256_BLOCK_SIZE], hash, SHA256_BLOCK_SIZE);
 }
 
@@ -73,6 +72,11 @@ bool poh_init() {
 }
 
 extern "C" {
+
+void poh_verify_many_set_verbose(bool val) {
+    g_verbose = val;
+}
+
 int poh_verify_many(uint8_t* hashes,
                     const uint64_t* num_hashes_arr,
                     size_t num_elems,
@@ -132,7 +136,6 @@ int poh_verify_many(uint8_t* hashes,
     CUDA_CHK(cudaPeekAtLastError());
 
     CUDA_CHK(cudaMemcpyAsync(hashes, cur_ctx->hashes, hashes_size, cudaMemcpyDeviceToHost, stream));
-
     CUDA_CHK(cudaStreamSynchronize(stream));
 
     pthread_mutex_unlock(&cur_ctx->mutex);

src/opencl-ecc-ed25519/verify.cpp

@@ -200,3 +200,55 @@ int cuda_host_unregister(void* ptr)
 {
     return 0;
 }
+
+static int
+get_checked_scalar(unsigned char* scalar, const unsigned char* signature) {
+    // Check if top 4-bits are clear
+    // then scalar is reduced.
+    // if ((signature[31] & 0xf0) == 0) {
+    //     for (int i = 0; i < 32; i++) {
+    //         scalar[i] = signature[i];
+    //     }
+    //     return 0;
+    // }
+
+    // if ((signature[31] >> 7) != 0) {
+    //     return 1;
+    // }
+
+    // scalar32_reduce(scalar);
+    // if (!consttime_equal(scalar, signature)) {
+    //     return 1;
+    // }
+    fprintf(stderr, "get_checked_scalar not implemented.\n");
+    exit(1);
+    return 0;
+
+}
+
+int ed25519_get_checked_scalar(unsigned char* out_scalar, const unsigned char* in_scalar) {
+    return get_checked_scalar(out_scalar, in_scalar);
+}
+
+// Return 0=success if ge unpacks and is not small order
+static int
+check_packed_ge_small_order(const unsigned char* packed_group_element) {
+    // ge_p3 signature_R;
+
+    // fail if ge does not unpack
+    // if (0 != ge_frombytes_negate_vartime(&signature_R, packed_group_element)) {
+    //     return 1;
+    // }
+
+    // // fail if ge is small order
+    // if (0 != ge_is_small_order(&signature_R)) {
+    //     return 1;
+    // }
+    fprintf(stderr, "check_packed_ge_small_order not implemented.\n");
+    exit(1);
+    return 0;
+}
+
+int ed25519_check_packed_ge_small_order(const unsigned char* packed_group_element) {
+    return check_packed_ge_small_order(packed_group_element);
+}

src/opencl-platform/cl_common.h

@@ -16,6 +16,11 @@ using namespace std;
 #define __host__
 #endif
 
+/* this is to silent warnings about opencl version (without this Werror flag will prevent compiling)*/
+#ifndef CL_TARGET_OPENCL_VERSION
+#define CL_TARGET_OPENCL_VERSION 120
+#endif
+
 extern bool g_verbose;
 
 #define LOG(...) if (g_verbose) { printf(__VA_ARGS__); }

src/opencl-platform/cl_init_platform.cpp

@@ -290,9 +290,12 @@ bool cl_check_init(void) {
 
             /* select device based on cli arguments */
             string tmpAttrData = attr_data;
-            
-            // always select last device of type GPU
-            device = device_list[dev];
+
+            // always select first device of first platform            
+            if((dev == 0) && (platf == 0)) {
+                device = device_list[dev];
+                cout << "<----- SELECTED";
+            }
 
             delete[] attr_data;
             cout << endl;

src/opencl-platform/kernels_verify.h

@@ -604,15 +604,14 @@ static __constant ulong32 K[64] = {
 };
 #endif
 
-/* Various logical functions */
 #define Ch(x,y,z)       (z ^ (x & (y ^ z)))
 #define Maj(x,y,z)      (((x | y) & z) | (x & y))
-#define S(x, n)         ROR64c((x),(n))
-#define R(x, n)           (((((ulong)x) & ((ulong)0xFFFFFFFFFFFFFFFFUL))) >> ((ulong)n))
-#define Sigma0(x)       (S(x, 28) ^ S(x, 34) ^ S(x, 39))
-#define Sigma1(x)       (S(x, 14) ^ S(x, 18) ^ S(x, 41))
-#define Gamma0(x)       (S(x, 1) ^ S(x, 8) ^ R(x, 7))
-#define Gamma1(x)       (S(x, 19) ^ S(x, 61) ^ R(x, 6))
+#define S(x, n)         RORc((x),(n))
+#define R(x, n)         (((x)&0xFFFFFFFFUL)>>(n))
+#define Sigma0(x)       (S(x, 2) ^ S(x, 13) ^ S(x, 22))
+#define Sigma1(x)       (S(x, 6) ^ S(x, 11) ^ S(x, 25))
+#define Gamma0(x)       (S(x, 7) ^ S(x, 18) ^ R(x, 3))
+#define Gamma1(x)       (S(x, 17) ^ S(x, 19) ^ R(x, 10))
 #ifndef MIN
    #define MIN(x, y) ( ((x)<(y))?(x):(y) )
 #endif
@@ -642,7 +641,7 @@ static int sha256_compress(hash_state * md, const unsigned char *buf)
 
     /* fill W[16..63] */
     for (i = 16; i < 64; i++) {
-        W[i] += Gamma1(W[i - 2]) + W[i - 7] + Gamma0(W[i - 15]) + W[i - 16];
+        W[i] = Gamma1(W[i - 2]) + W[i - 7] + Gamma0(W[i - 15]) + W[i - 16];
     }
 
     /* Compress */
@@ -738,6 +737,7 @@ static int sha256_compress(hash_state * md, const unsigned char *buf)
     for (i = 0; i < 8; i++) {
         md->sha256.state[i] = md->sha256.state[i] + S[i];
     }
+
     return CRYPT_OK;
 }
 
@@ -4266,7 +4266,6 @@ __kernel void poh_verify_kernel(__global uint8_t* hashes,
     for(int i = 0; i < SHA256_BLOCK_SIZE; i++) {
         hash[i] = hashes[idx * SHA256_BLOCK_SIZE + i];
     }
-    
 
     for (size_t i = 0; i < num_hashes_arr[idx]; i++) {
         hash_state sha_state;