commit 189878e99676b316362ff8f3ad0b9b611c210770
parent 91cf00dc343e2ed3e870171f8199e8a1ce75689a
Author: quantumish <freifeld.david@gmail.com>
Date: Wed, 29 Jul 2026 01:55:46 -0700
Add slightly parallel SHA256 and old zero knowledge proof thing
Diffstat:
| A | sha.cu | | | 172 | +++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ |
| A | zk-poly.c | | | 195 | +++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ |
2 files changed, 367 insertions(+), 0 deletions(-)
diff --git a/sha.cu b/sha.cu
@@ -0,0 +1,172 @@
+/* Editor's note: a partially parallelized SHA256 implementation.
+ * Written around December 2021 for my high school computer security class.
+ * If I remember correctly the parallelization was done very poorly and there's room
+ * for a lot more optimization.
+ *
+ * The message schedules for the chunks processed by SHA256 are not sequentially
+ * dependent and so this program generates them in parallel and then process them
+ * sequentially when updating the hash.
+ */
+
+#include <cstdint>
+#include <cstring>
+#include <unistd.h>
+#include <stdio.h>
+#include <fcntl.h>
+#include <sys/types.h>
+#include <sys/stat.h>
+#include <cuda_runtime_api.h>
+#include <cuda.h>
+
+// 32-bit bit right rotation
+#define rotr(a,b) (((a) >> (b)) | ((a) << (32-(b))))
+
+// 32-bit byte swap
+#define bswap32(x) ((x>>24)&0xff) | \
+ ((x<<8)&0xff0000) | \
+ ((x>>8)&0xff00) | \
+ ((x<<24)&0xff000000) \
+
+uint64_t k[64] = {
+ 0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, 0x923f82a4, 0xab1c5ed5,
+ 0xd807aa98, 0x12835b01, 0x243185be, 0x550c7dc3, 0x72be5d74, 0x80deb1fe, 0x9bdc06a7, 0xc19bf174,
+ 0xe49b69c1, 0xefbe4786, 0x0fc19dc6, 0x240ca1cc, 0x2de92c6f, 0x4a7484aa, 0x5cb0a9dc, 0x76f988da,
+ 0x983e5152, 0xa831c66d, 0xb00327c8, 0xbf597fc7, 0xc6e00bf3, 0xd5a79147, 0x06ca6351, 0x14292967,
+ 0x27b70a85, 0x2e1b2138, 0x4d2c6dfc, 0x53380d13, 0x650a7354, 0x766a0abb, 0x81c2c92e, 0x92722c85,
+ 0xa2bfe8a1, 0xa81a664b, 0xc24b8b70, 0xc76c51a3, 0xd192e819, 0xd6990624, 0xf40e3585, 0x106aa070,
+ 0x19a4c116, 0x1e376c08, 0x2748774c, 0x34b0bcb5, 0x391c0cb3, 0x4ed8aa4a, 0x5b9cca4f, 0x682e6ff3,
+ 0x748f82ee, 0x78a5636f, 0x84c87814, 0x8cc70208, 0x90befffa, 0xa4506ceb, 0xbef9a3f7, 0xc67178f2,
+};
+
+struct sha_ctx {
+ uint32_t hash[8] = {
+ 0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0xa54ff53a,
+ 0x510e527f, 0x9b05688c, 0x1f83d9ab, 0x5be0cd19,
+ };
+ uint64_t len;
+
+ sha_ctx(uint64_t len);
+ void compress(uint32_t* w);
+ void dump_hash();
+};
+
+sha_ctx::sha_ctx(uint64_t length) :len(length) {}
+
+/*
+ * GPU Kernel responsible for generating message schedules.
+ * Takes:
+ * - `bytes`, a pointer to a buffer containing the input from the file.
+ * - `w`, a pointer to the large buffer to write all the message schedules to.
+ * - `iters`, the number of message schedules to generate per thread.
+ * Does not return anything but instead modifies `w`
+ */
+__global__ void process(const uint8_t* bytes, uint32_t* w, uint32_t iters) {
+ const size_t start_chunk = (blockIdx.y*gridDim.x*blockDim.x)+(blockIdx.x*blockDim.x)+threadIdx.x;
+ for (int off = 0; off < iters; off++) {
+ uint32_t* w_adj = w+(((start_chunk*iters)+off)*64);
+ memcpy(w_adj, bytes+(((start_chunk*iters)+off)*64), 64);
+ for (int i = 0; i < 16; i++) w_adj[i] = bswap32(w_adj[i]);
+ for (int i = 16; i < 64; i++) {
+ uint32_t s0 = (rotr(w_adj[i-15], 7) ^ rotr(w_adj[i-15], 18) ^ (w_adj[i-15] >> 3));
+ uint32_t s1 = (rotr(w_adj[i-2], 17) ^ rotr(w_adj[i-2], 19) ^ (w_adj[i-2] >> 10));
+ w_adj[i] = w_adj[i-16] + s0 + w_adj[i-7] + s1;
+ }
+ }
+}
+
+/*
+ * CPU-side function that updates the hash for a given message schedule.
+ * Takes:
+ * - `w`, a pointer to a single message schedule.
+ * Does not return anything but modifies `sha_ctx::hash`
+ */
+void sha_ctx::compress(uint32_t* w) {
+ uint32_t a[8] = {0};
+ memcpy(a, hash, 8*sizeof(uint32_t));
+ for (int i = 0; i < 64; i++) {
+ uint32_t s1 = (rotr(a[4], 6) ^ rotr(a[4], 11) ^ rotr(a[4], 25));
+ uint32_t ch = (a[4] & a[5]) ^ ((~a[4]) & a[6]);
+ uint32_t temp1 = a[7] + s1 + ch + k[i] + w[i];
+ uint32_t s0 = (rotr(a[0], 2) ^ rotr(a[0], 13) ^ rotr(a[0], 22));
+ uint32_t maj = (a[0] & a[1]) ^ (a[0] & a[2]) ^ (a[1] & a[2]);
+ uint32_t temp2 = s0 + maj;
+ for (int i = 7; i > 0; i--) a[i] = a[i-1];
+ a[4] += temp1;
+ a[0] = temp1 + temp2;
+ }
+ for (int i = 0; i < 8; i++) hash[i] += a[i];
+}
+
+/*
+ * CPU-side function that prints out the current hash.
+ * Returns nothing but writes to stdout.
+ */
+void sha_ctx::dump_hash() {
+ uint32_t temp_hash[8];
+ memcpy(temp_hash, hash, 8*sizeof(uint32_t));
+ for (int i = 0; i < 8; i++) temp_hash[i] = bswap32(temp_hash[i]);
+ auto u8_ptr = reinterpret_cast<uint8_t*>(temp_hash);
+ for (int i = 0; i < 32; i++) printf("%02x", u8_ptr[i]);
+}
+
+int main(int argc, char** argv) {
+ // Open file and initialize sha_ctx
+ int fd = open(argv[1], O_RDONLY | O_NONBLOCK);
+ // Advise kernel we're reading sequentially
+ posix_fadvise(fd, 0, 0, POSIX_FADV_SEQUENTIAL);
+ struct stat stat;
+ fstat(fd, &stat); // Get size of file
+ sha_ctx sha(stat.st_size);
+
+ // Allocate input and message schedule buffers on the FPU
+ const size_t BUFFER_SIZE = 268435456; // 2 GiB
+ uint8_t* buf;
+ cudaMallocManaged(&buf, BUFFER_SIZE);
+ uint32_t* w;
+ cudaMallocManaged(&w, BUFFER_SIZE*4);
+
+ bool padded = false; // Store if we've padded the input yet (for handling edge cases)
+ size_t bytes_read = read(fd, buf, BUFFER_SIZE);
+
+ do {
+ // Maximum size_t (represented by underflow here) is read's return code for error
+ if (bytes_read == (size_t)-1) {
+ printf("Error reading file.");
+ exit(1);
+ }
+ if (!bytes_read) break; // If we're done reading, stop.
+ else if (bytes_read < BUFFER_SIZE) {
+ // Pad buffer in accordance with SHA256's standards
+ size_t buffer_len = 64 * (((bytes_read + 9) / 64) + 1);
+ for (size_t i = bytes_read; i < buffer_len; i++) buf[i] = 0;
+ buf[bytes_read] = 0b10000000;
+ for (int i = 1; i <= 8; i++) buf[buffer_len-i] = sha.len*8 >> (i-1)*8;
+ process<<<dim3{8,8,1}, 64>>>(buf, w, 1024);
+ cudaDeviceSynchronize();
+ for (int i = 0; i < buffer_len/64; i++) sha.compress(w+(i*64));
+ padded=true;
+ } else {
+ process<<<dim3{8,8,1}, 64>>>(buf, w, 1024);
+ cudaDeviceSynchronize();
+ for (int i = 0; i < BUFFER_SIZE/64; i++) sha.compress(w+(i*64));
+ }
+ } while ((bytes_read = read(fd, buf, BUFFER_SIZE)));
+ if (padded == false) {
+ // Pad buffer in accordance with SHA256
+ for (size_t i = 0; i < 64; i++) buf[i] = 0;
+ buf[0] = 0b10000000;
+ for (int i = 1; i <= 8; i++) buf[64-i] = sha.len*8 >> (i-1)*8;
+
+ // Run last chunk on CPU
+ memcpy(w, buf, 64);
+ for (int i = 0; i < 16; i++) w[i] = bswap32(w[i]);
+ for (int i = 16; i < 64; i++) {
+ uint32_t s0 = (rotr(w[i-15], 7) ^ rotr(w[i-15], 18) ^ (w[i-15] >> 3));
+ uint32_t s1 = (rotr(w[i-2], 17) ^ rotr(w[i-2], 19) ^ (w[i-2] >> 10));
+ w[i] = w[i-16] + s0 + w[i-7] + s1;
+ }
+ sha.compress(w);
+ }
+ sha.dump_hash();
+ printf(" %s\n", argv[1]);
+}
diff --git a/zk-poly.c b/zk-poly.c
@@ -0,0 +1,195 @@
+/* A hacky implementation of the "Non-Interactive Zero-Knowledge of a Polynomial" scheme
+ * described in https://doi.org/10.48550/arXiv.1906.07221
+ * Written during late October 2020.
+ */
+
+#include <pthread.h>
+#include <sys/types.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <sys/socket.h>
+#include <arpa/inet.h>
+#include <netinet/in.h>
+#include <string.h>
+#include <math.h>
+#include <time.h>
+#include <gmp.h>
+
+#define PORT 5000
+#define BUFSIZE 1024
+#define EPSILON 0.000001
+
+void* start_verifier(void* args)
+{
+ // Intialize socket with AF_INET IP family and SOCK_DGRAM datagram service, exit if failed
+ int s;
+ if ((s = socket(AF_INET, SOCK_DGRAM, 0)) < 0) {
+ exit(1);
+ }
+ // Establish sockaddr_in struct to pass into bind function
+ struct sockaddr_in addr;
+ memset((char *)&addr, 0, sizeof(addr));
+ addr.sin_family = AF_INET; // Specify address family.
+ addr.sin_addr.s_addr = htonl(INADDR_ANY); // INADDR_ANY just 0.0.0.0, machine IP address
+ addr.sin_port = htons(PORT); // Specify port.
+ if (bind(s, (struct sockaddr *)&addr, sizeof(addr)) < 0) {
+ exit(1);
+ }
+
+ // IP address of client
+ struct sockaddr_in remaddr;
+ socklen_t addrlen = sizeof(remaddr);
+ int recvlen;
+ unsigned char buf[BUFSIZE];
+
+ int connections = 0;
+ int clients[500];
+ int unproven = 1;
+ int g = 6;
+ int degree = 3;
+ int n = 17;
+ int s_1;
+ int a;
+ mpz_t bigs;
+ mpz_init_set_ui(bigs, s_1);
+ mpz_t biga;
+ mpz_init_set_ui(biga, a);
+ mpz_t bigg;
+ mpz_init_set_ui(bigg, g);
+ mpz_t bign;
+ mpz_init_set_ui(bign, n);
+ mpz_t t;
+ while (unproven) {
+ recvlen = recvfrom(s, buf, BUFSIZE, 0, (struct sockaddr *)&remaddr, &addrlen);
+ if (recvlen > BUFSIZE) {
+ sendto(s, "ERR: Too long.", 14, 0, (struct sockaddr *) &remaddr, addrlen);
+ }
+ else if (recvlen > 0) {
+ void* msg = 0x0;
+ buf[recvlen] = 0;
+ printf(" VERIFIER │ Received %d-byte message from %i: \"%s\"\n",
+ recvlen, remaddr.sin_port, buf);
+ if (recvlen==8) {
+ clients[connections] = remaddr.sin_port;
+ connections++;
+ srand(time(0));
+ s_1 = rand() % 500;
+ a = rand() % 500;
+ mpz_init_set_ui(t, (s_1-((int*)buf)[0])*(s_1-((int*)buf)[1]));
+ printf(" VERIFIER │ Chose s = %d and a = %d\n", s_1, a);
+ mpz_t* enc_s = malloc(sizeof(mpz_t)*(degree+1)*2);
+ for (int i = 0; i <= degree; i++) {
+ mpz_init(enc_s[i]);
+ mpz_t temp1;
+
+ mpz_init_set_ui(temp1, (int)pow(s_1, i));
+ mpz_powm_sec(enc_s[i], bigg, temp1, bign);
+ mpz_init(enc_s[degree+1+i]);
+ mpz_t temp2;
+ mpz_init_set_ui(temp2, (int)a*pow(s_1, i));
+ mpz_powm_sec(enc_s[i+1+degree], bigg, temp2, bign);
+ }
+ msg = enc_s;
+ } else if (recvlen == sizeof(mpz_t)*3) {
+ mpz_t tmp1;
+ mpz_init(tmp1);
+ mpz_t tmp2;
+ mpz_init(tmp2);
+ mpz_powm_sec(tmp1, ((mpz_t*)buf)[2], t, bign);
+ mpz_t tmp3;
+ mpz_init(tmp3);
+ mpz_div(tmp3, ((mpz_t*)buf)[0], tmp1);
+ mpz_abs(tmp3, tmp3);
+ mpz_ui_sub(tmp3, 1, tmp3);
+ //gmp_printf("%Fe %Fe (diff %Fe)\n", ((mpz_t*)buf)[0], tmp1, tmp3);
+ if (mpz_cmp_d(tmp3, EPSILON) == -1) printf(" VERIFIER │ Valid roots!\n");
+ else printf(" VERIFIER │ ERR: Invalid roots!\n");
+ mpz_powm_sec(tmp2, ((mpz_t*)buf)[0], biga, bign);
+ mpz_t tmp4;
+ mpz_init(tmp4);
+ mpz_div(tmp4, ((mpz_t*)buf)[1], tmp2);
+ mpz_abs(tmp4, tmp4);
+ mpz_ui_sub(tmp4, 1, tmp4);
+ //gmp_printf("%Fe %Fe (diff %Fe)\n", ((mpz_t*)buf)[1], tmp2, tmp4);
+ if (mpz_cmp_d(tmp4, EPSILON) == -1) printf(" VERIFIER │ Valid form!\n");
+ else printf(" VERIFIER │ ERR: Invalid form!\n");
+ }
+ if (msg != 0x0) {
+ sendto(s, msg, sizeof(mpz_t)*(degree+1)*2, 0, (struct sockaddr *) &remaddr, addrlen);
+ }
+ }
+ }
+
+ return 0x0;
+}
+
+void* start_prover(void* args)
+{
+ // Same socket is needed on client end so initialize all over again.
+ int s;
+ if ((s = socket(AF_INET, SOCK_DGRAM, 0)) < 0) {
+ printf("\n Error : Socket Failed \n");
+ }
+ struct sockaddr_in addr;
+ memset((char *)&addr, 0, sizeof(addr));
+ addr.sin_family = AF_INET; // Specify address family.
+ addr.sin_addr.s_addr = htonl(INADDR_ANY); // INADDR_ANY just 0.0.0.0, machine IP address
+ addr.sin_port = htons(PORT); // Specify port.
+ // Connect to server
+ if(connect(s, (struct sockaddr *)&addr, sizeof(addr)) < 0) {
+ printf("ERR: Connect failed.\n");
+ return 0x0;
+ }
+ // NOTE This can and will not work if flag argument set to 1
+ int roots[2] = {-1,-2};
+ sendto(s, roots, 8, 0, (struct sockaddr*)NULL, sizeof(addr));
+ printf(" Prover │ Informed server of existence.\n");
+ char buf[BUFSIZE];
+ int recvlen;
+ socklen_t len = sizeof(addr);
+ int phase = 0;
+ int constants[4] = {0, 2, 3, 1};
+ int degree = 3;
+ int delta = rand() % 100;
+ while (1==1) {
+ recvlen = recvfrom(s, buf, BUFSIZE, 0, (struct sockaddr *) &addr, &len);
+ if (recvlen > 0) {
+ buf[recvlen] = 0;
+ printf(" Prover │ Received %d-byte message from server: \"%s\"\n", recvlen, buf);
+ mpz_t* enc_ph = malloc(sizeof(mpz_t)*3);
+ mpz_init_set_ui(enc_ph[0], 1);
+ mpz_init_set_ui(enc_ph[1], 1);
+ mpz_init(enc_ph[2]);
+ for (int i = 0; i < degree+1; i++) {
+ mpz_pow_ui(((mpz_t*)buf)[i], ((mpz_t*)buf)[i], abs(constants[i]));
+ mpz_mul(enc_ph[0], enc_ph[0], ((mpz_t*)buf)[i]);
+ mpz_mul(enc_ph[1], enc_ph[1], ((mpz_t*)buf)[i+1+degree]);
+ }
+ mpz_t tmp1;
+ mpz_init(tmp1);
+ mpz_pow_ui(tmp1, ((mpz_t*)buf)[0], -roots[0]);
+ mpz_mul(tmp1, tmp1, ((mpz_t*)buf)[1]);
+ mpz_t tmp2;
+ mpz_init(tmp2);
+ mpz_pow_ui(tmp1, ((mpz_t*)buf)[0], -roots[1]);
+ mpz_mul(tmp1, tmp1, ((mpz_t*)buf)[1]);
+ mpz_t denom;
+ mpz_init(denom);
+ mpz_mul(denom, tmp1, tmp2);
+ mpz_div(enc_ph[2], enc_ph[0], denom);
+ for (int i = 0; i < 3; i++) mpz_pow_ui(enc_ph[i], enc_ph[i], delta);
+ sendto(s, enc_ph, sizeof(mpz_t)*3, 0, (struct sockaddr*)NULL, sizeof(addr));
+ }
+ }
+ return 0x0;
+}
+
+int main()
+{
+ pthread_t verifier;
+ pthread_create(&verifier, NULL, start_verifier, 0x0);
+ pthread_t prover;
+ pthread_create(&prover, NULL, start_prover, 0x0);
+ pthread_join(verifier, NULL);
+}
