From 528679bbf71f9ab487bd2aa3f01478081c65c231 Mon Sep 17 00:00:00 2001 From: Rodrigo Madera Date: Sun, 26 Apr 2026 01:27:49 -0300 Subject: [PATCH 1/6] Increase speed by 20-35% by optimizing mp_mul_mod_word_sub(). --- profanity.cl | 69 ++++++++++++++++++++++++++++++++++------------------ 1 file changed, 45 insertions(+), 24 deletions(-) diff --git a/profanity.cl b/profanity.cl index 75abd70..06709f7 100644 --- a/profanity.cl +++ b/profanity.cl @@ -277,31 +277,52 @@ mp_word mp_mul_word_add_extra(mp_number * const r, const mp_number * const a, co return *extra < cM ? 1 : (*extra == cM ? cA : 0); } -// Multiplies a number with a word, potentially adds modhigher to it, and then subtracts it from en existing number, no extra words, no overflow -// This is a special function only used for modular multiplication +// Multiplies a number with a word, potentially adds modhigher to it, and then subtracts it from +// an existing number, no extra words, no overflow. +// +// This is a special function only used for modular multiplication. +// +// Optimized code (secp256k1 fast reduction) +// by Rodrigo Madera (madera at acm dot org). +// +// Optimization: +// +// pmod = 2^256 - p +// +// p = 0x1000003D1 = 2^32 + 977 +// +// q * pmod = q * (2^256 - p) +// = q * 2^256 - q * p +// +// (r - q * pmod) mod 2^256 == (r + q*p) mod 2^256 +// +// This reduces the amount of bits used giving us 20-35% speed improvements. +// void mp_mul_mod_word_sub(mp_number * const r, const mp_word w, const bool withModHigher) { - // Having these numbers declared here instead of using the global values in __constant address space seems to lead - // to better optimizations by the compiler on my GTX 1070. - mp_number mod = { { 0xfffffc2f, 0xfffffffe, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff} }; - mp_number modhigher = { {0x00000000, 0xfffffc2f, 0xfffffffe, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff} }; - - mp_word cM = 0; // Carry for multiplication - mp_word cS = 0; // Carry for subtraction - mp_word tS = 0; // Temporary storage for subtraction - mp_word tM = 0; // Temporary storage for multiplication - mp_word cA = 0; // Carry for addition of modhigher - - for (mp_word i = 0; i < MP_WORDS; ++i) { - tM = (mod.d[i] * w + cM); - cM = mul_hi(mod.d[i], w) + (tM < cM); - - tM += (withModHigher ? modhigher.d[i] : 0) + cA; - cA = tM < (withModHigher ? modhigher.d[i] : 0) ? 1 : (tM == (withModHigher ? modhigher.d[i] : 0) ? cA : 0); - - tS = r->d[i] - tM - cS; - cS = tS > r->d[i] ? 1 : (tS == r->d[i] ? cS : 0); - - r->d[i] = tS; + const mp_word lo977 = 977u * w; + const mp_word hi977 = mul_hi(977u, w); + + const mp_word p0 = lo977; + const ulong p1_full = (ulong)w + hi977 + (withModHigher ? 0x000003D1u : 0u); + const mp_word p1 = (mp_word)p1_full; + const mp_word p2 = (mp_word)(p1_full >> 32) + (withModHigher ? 1u : 0u); + + ulong s = (ulong)r->d[0] + p0; + r->d[0] = (mp_word)s; + mp_word c = (mp_word)(s >> 32); + + s = (ulong)r->d[1] + p1 + c; + r->d[1] = (mp_word)s; + c = (mp_word)(s >> 32); + + s = (ulong)r->d[2] + p2 + c; + r->d[2] = (mp_word)s; + c = (mp_word)(s >> 32); + + for (mp_word i = 3; i < MP_WORDS; ++i) { + s = (ulong)r->d[i] + c; + r->d[i] = (mp_word)s; + c = (mp_word)(s >> 32); } } From 204cef2569adb72b935302b7338ea16bfd2343e1 Mon Sep 17 00:00:00 2001 From: Denis Date: Sat, 9 May 2026 22:31:37 +0100 Subject: [PATCH 2/6] Fix OpenCL device enumeration error handling --- profanity.cpp | 42 ++++++++++++++++++++++++++++++++---------- 1 file changed, 32 insertions(+), 10 deletions(-) diff --git a/profanity.cpp b/profanity.cpp index a5b0e1d..bd1e611 100644 --- a/profanity.cpp +++ b/profanity.cpp @@ -33,24 +33,46 @@ std::string readFile(const char * const szFilename) return contents.str(); } -std::vector getAllDevices(cl_device_type deviceType = CL_DEVICE_TYPE_GPU) -{ +std::vector getAllDevices(cl_device_type deviceType = CL_DEVICE_TYPE_GPU) { std::vector vDevices; cl_uint platformIdCount = 0; - clGetPlatformIDs (0, NULL, &platformIdCount); + cl_int err = clGetPlatformIDs(0, NULL, &platformIdCount); + if (err != CL_SUCCESS || platformIdCount == 0) { + std::cerr << "warning: no OpenCL platforms found, err = " << err << std::endl; + return vDevices; + } + + std::vector platformIds(platformIdCount); + err = clGetPlatformIDs(platformIdCount, platformIds.data(), NULL); + if (err != CL_SUCCESS) { + std::cerr << "warning: failed to enumerate OpenCL platforms, err = " << err << std::endl; + return vDevices; + } - std::vector platformIds (platformIdCount); - clGetPlatformIDs (platformIdCount, platformIds.data (), NULL); + for (auto it = platformIds.cbegin(); it != platformIds.cend(); ++it) { + cl_uint countDevice = 0; - for( auto it = platformIds.cbegin(); it != platformIds.cend(); ++it ) { - cl_uint countDevice; - clGetDeviceIDs(*it, deviceType, 0, NULL, &countDevice); + err = clGetDeviceIDs(*it, deviceType, 0, NULL, &countDevice); + if (err != CL_SUCCESS || countDevice == 0) { + char platformName[256] = {0}; + clGetPlatformInfo(*it, CL_PLATFORM_NAME, sizeof(platformName), platformName, NULL); + std::cerr << "warning: skipping OpenCL platform without usable GPU devices: " + << platformName << ", err = " << err << std::endl; + continue; + } std::vector deviceIds(countDevice); - clGetDeviceIDs(*it, deviceType, countDevice, deviceIds.data(), &countDevice); + err = clGetDeviceIDs(*it, deviceType, countDevice, deviceIds.data(), &countDevice); + if (err != CL_SUCCESS) { + char platformName[256] = {0}; + clGetPlatformInfo(*it, CL_PLATFORM_NAME, sizeof(platformName), platformName, NULL); + std::cerr << "warning: failed to get GPU devices from platform: " + << platformName << ", err = " << err << std::endl; + continue; + } - std::copy( deviceIds.begin(), deviceIds.end(), std::back_inserter(vDevices) ); + std::copy(deviceIds.begin(), deviceIds.end(), std::back_inserter(vDevices)); } return vDevices; From 2b83d5620af94ade1e41b2c6767c50859631ecfc Mon Sep 17 00:00:00 2001 From: Cursor Agent Date: Wed, 22 Jul 2026 15:48:48 +0000 Subject: [PATCH 3/6] Fix and clarify the fast-reduction comment in mp_mul_mod_word_sub The comment introduced in PR #49 mixed up p and pmod: it defined pmod = 2^256 - p but then called the 33-bit constant 0x1000003D1 "p", and stated the congruence identity in the wrong direction. Rewrite the derivation with consistent naming and state the attribution as "contributed by Rodrigo Madera (@madera)". Co-authored-by: Gleb Alekseev --- profanity.cl | 17 ++++++++--------- 1 file changed, 8 insertions(+), 9 deletions(-) diff --git a/profanity.cl b/profanity.cl index 06709f7..e9a4309 100644 --- a/profanity.cl +++ b/profanity.cl @@ -282,21 +282,20 @@ mp_word mp_mul_word_add_extra(mp_number * const r, const mp_number * const a, co // // This is a special function only used for modular multiplication. // -// Optimized code (secp256k1 fast reduction) -// by Rodrigo Madera (madera at acm dot org). +// Optimization (secp256k1 fast reduction) contributed by Rodrigo Madera (@madera). // -// Optimization: +// The secp256k1 prime has the special form: // -// pmod = 2^256 - p +// p = 2^256 - pmod, where pmod = 2^32 + 977 = 0x1000003D1 // -// p = 0x1000003D1 = 2^32 + 977 +// Therefore, working modulo 2^256: // -// q * pmod = q * (2^256 - p) -// = q * 2^256 - q * p +// q * p = q * (2^256 - pmod) = q * 2^256 - q * pmod == -q * pmod (mod 2^256) // -// (r - q * pmod) mod 2^256 == (r + q*p) mod 2^256 +// (r - q * p) mod 2^256 == (r + q * pmod) mod 2^256 // -// This reduces the amount of bits used giving us 20-35% speed improvements. +// So instead of multiplying q by the full 256-bit p and subtracting, we multiply q by the +// 33-bit pmod and add. This reduces the amount of bits used, giving us 20-35% speed improvements. // void mp_mul_mod_word_sub(mp_number * const r, const mp_word w, const bool withModHigher) { const mp_word lo977 = 977u * w; From 46bf46696892955044960efc9ccf996e5102d5af Mon Sep 17 00:00:00 2001 From: Cursor Agent Date: Wed, 22 Jul 2026 15:48:48 +0000 Subject: [PATCH 4/6] Add OpenCL correctness tests and benchmark for the mp_mod_mul optimization The tests build the real kernel sources (keccak.cl + profanity.cl) plus a harness that keeps a verbatim copy of the pre-optimization mp_mul_mod_word_sub/mp_mod_mul from master, and verify on any OpenCL device: - mp_mul_mod_word_sub old vs new are bit-exact identical, and both match an independent host-side big-integer reference (r - w*p - [withModHigher]*(p<<32)) mod 2^256; - mp_mod_mul old vs new are bit-exact identical on all inputs, and for in-domain inputs (x, y < p) the result is congruent to x*y mod p; - bench_mod_mul times dependency-chained modular multiplications for both variants and cross-checks their outputs. Runs on CPU via PoCL when no GPU is available. On a Xeon with PoCL the optimized variant measures ~x1.7 on isolated mp_mod_mul throughput. Co-authored-by: Gleb Alekseev --- .gitignore | 1 + tests/Makefile | 32 +++ tests/__pycache__/clutil.cpython-312.pyc | Bin 0 -> 3231 bytes tests/bench_mod_mul.cpp | 118 ++++++++++ tests/harness.cl | 129 ++++++++++ tests/test_correctness.cpp | 267 +++++++++++++++++++++ tests/testutil.hpp | 284 +++++++++++++++++++++++ 7 files changed, 831 insertions(+) create mode 100644 tests/Makefile create mode 100644 tests/__pycache__/clutil.cpython-312.pyc create mode 100644 tests/bench_mod_mul.cpp create mode 100644 tests/harness.cl create mode 100644 tests/test_correctness.cpp create mode 100644 tests/testutil.hpp diff --git a/.gitignore b/.gitignore index d58231e..2fcd8e3 100644 --- a/.gitignore +++ b/.gitignore @@ -1,5 +1,6 @@ # Makefile artifacts *.o *.so +*.x64 cache-opencl.* bin diff --git a/tests/Makefile b/tests/Makefile new file mode 100644 index 0000000..dbf27a5 --- /dev/null +++ b/tests/Makefile @@ -0,0 +1,32 @@ +# Builds the mp-math correctness tests and benchmarks. Mirrors the style of the +# main Makefile. Requires an OpenCL runtime; on machines without a GPU, PoCL +# (http://portablecl.org, e.g. `apt install pocl-opencl-icd`) works fine. +# +# Usage: +# make +# ./test_correctness.x64 [num_random_cases] +# ./bench_mod_mul.x64 [global_size] [iterations] [repetitions] + +CC=g++ +CDEFINES= +EXECUTABLES=test_correctness.x64 bench_mod_mul.x64 + +UNAME_S := $(shell uname -s) +ifeq ($(UNAME_S),Darwin) + LDFLAGS=-framework OpenCL + CFLAGS=-std=c++11 -Wall -O2 +else + LDFLAGS=-lOpenCL + CFLAGS=-std=c++11 -Wall -O2 +endif + +all: $(EXECUTABLES) + +test_correctness.x64: test_correctness.cpp testutil.hpp ../types.hpp + $(CC) $(CFLAGS) $(CDEFINES) test_correctness.cpp $(LDFLAGS) -o $@ + +bench_mod_mul.x64: bench_mod_mul.cpp testutil.hpp ../types.hpp + $(CC) $(CFLAGS) $(CDEFINES) bench_mod_mul.cpp $(LDFLAGS) -o $@ + +clean: + rm -rf $(EXECUTABLES) diff --git a/tests/__pycache__/clutil.cpython-312.pyc b/tests/__pycache__/clutil.cpython-312.pyc new file mode 100644 index 0000000000000000000000000000000000000000..ca281ee58fc8abcf2aca5cf94c8fdc8f89a77f7b GIT binary patch literal 3231 zcmbtWO>7&-6`t8$F3BaSKPQT8#K^`$Q&G00O19j_ah&QuN)%BQm8N2vvMbJtBDLhQ zv#URXLU`ywY!y&K3`j--LR@E{pHRKtc%I2&pXQ)1-q!Z?Hx17qf zrxbHmHdVtkWF<}HX*C_2O(}HF=#Zqo`GiIoUk+u|bpI$*jZ-Bh$Mm$3Fw8_cF6%S$ zoEnQMa~&~FzF0F}p4x^xVIOJLFo1PQN+l@ODbt+Q7f*_3GO}VAYHC_j4SSz+Drjl4 zq0iHp3fnng9vVCF0M(j+bmktt_pkRJ{f8e$2h6Pp2xzvaeLx+e*CF^&M9iKB~A7`=eHkbxIDjHckk z7%ihEehMw(AU9@V8Wb%-hcH^4ZdgJF!fpvkJ(0G&6NBTS=tL+qX$e$SNRYFHw31RS zJY%q($@0Td&+Ps<2Ix@?am+3@j5}HaGvRmM9ybL6d_Pl|$>0-PP zFKPwtqnEZxsqW%Vxa93RMA+5x=}&g+8;T2sh2JiH%ftLz#yn!=p}`J*uc6n=|H

zZE`;r1eWf05DsAgNsW_E7=Te}CX=um;i@Thg2z9}Go8RCS(v}mPFsTr#8nMg)n44O zyn;<{)kxAVQPo+w&!LN?*hrA-9Te;sxbAn|LOI)lK;XU^y!lSd(hvU z-i&X=H?<9I`-eNDrRJAP{%ePB!7Cl}Vx5>D`ez;T*A=f7uC2$mu9kckzVx*|@U?z^ z{=ToZfr#JKA zS6{v_G*jmI$8v=g9@`^-g~28}6En`mlK$Wt7~Mz-U`3gB7nb7%ZZ$mT^OwmFe!+M zD|@4L?2`PWpga$8x!f&l3DeZnb~Q~BO1d2i7m*c8m1Vj827EJ2c4dt(o3naa1_h(T z4k2affA9#}fefRd+Y%MST+XN#m(ExsGkx`n<)TVDu3DZq$D?nDCI-S54}6f!KwTw4 zm&L3olqWVYtAa%vw2N>eKV7X}FK0rXFSM)64L^1OCRZfaKA>6ezY8 z+HOv4PHjw;q~{>qk{?KY(-9I}&8z&bKd_4T8k=t}Z9l(}+n(5IDm7k$!s%^Y?cHr` zTA$u{=A-^?e!KT~;)icMX!w4i_mGE=`yO9WDoE>H?*|UOVCYan{`&muPlCp|<@R7d z|1FJ7b~?MkJyrpi=@ifz&a8e$N$19MSm>@(0-;>iVBuH-=s*Vi900n)0o1Yzy=wBDhs>HLWLuI7WS0G3}`dlby#@qDFLlrtUyP|-WYun5DX{(8=^X(7`O;S z34IZ4YYg!0J;DQiD9c~Vl~pNRx4fpN=|#q{k#xD+@~Qv>C?Zfu+JFz(5f>C~UNvYZ z6T?Z`1wHKs@&tewY8ZMdK%(qi*O=fOkUaWY@)uhRtvCC(#1G#1q~-pZw$Cs8A^2JF z3+Z~kf3M-}){CWv=Z_HQ>M01Tc(u!h(g)8m1oZ%d&Ys(%+r76IZk^w0d(d)awSOJ2 zcdZM*8TkfK)N{y#-l157PVgzq9gPw_7LD3>isgcBWtPjw_hgpYlOv|t^|5^2xJiQA z&kMIQ4KE6KhIpd(^8pO7CwSRUc-#^v-wt*6T&1i+l +#include +#include +#include + +#include "testutil.hpp" + +struct BenchResult { + double bestSeconds; + std::vector output; +}; + +static BenchResult benchVariant(const ClSetup & s, const char * const kernelName, + const std::vector & x, const std::vector & y, + const uint32_t iterations, const int reps) +{ + const size_t count = x.size(); + cl_int err; + cl_kernel kernel = clCreateKernel(s.program, kernelName, &err); + clCheck(err, "clCreateKernel"); + + cl_mem yBuf = clCreateBuffer(s.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, count * sizeof(mp_number), (void *)y.data(), &err); + clCheck(err, "clCreateBuffer(y)"); + + BenchResult result; + result.bestSeconds = -1; + result.output.resize(count); + + // reps timed runs + one warmup run (which also produces the output used + // for the old-vs-new cross-check) + for (int rep = 0; rep <= reps; ++rep) { + cl_mem xBuf = clCreateBuffer(s.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, count * sizeof(mp_number), (void *)x.data(), &err); + clCheck(err, "clCreateBuffer(x)"); + clCheck(clSetKernelArg(kernel, 0, sizeof(cl_mem), &xBuf), "clSetKernelArg(0)"); + clCheck(clSetKernelArg(kernel, 1, sizeof(cl_mem), &yBuf), "clSetKernelArg(1)"); + clCheck(clSetKernelArg(kernel, 2, sizeof(uint32_t), &iterations), "clSetKernelArg(2)"); + clCheck(clFinish(s.queue), "clFinish(pre)"); + + const auto t0 = std::chrono::steady_clock::now(); + clCheck(clEnqueueNDRangeKernel(s.queue, kernel, 1, NULL, &count, NULL, 0, NULL, NULL), "clEnqueueNDRangeKernel"); + clCheck(clFinish(s.queue), "clFinish"); + const double seconds = std::chrono::duration(std::chrono::steady_clock::now() - t0).count(); + + if (rep == 0) { + clCheck(clEnqueueReadBuffer(s.queue, xBuf, CL_TRUE, 0, count * sizeof(mp_number), result.output.data(), 0, NULL, NULL), "clEnqueueReadBuffer"); + clCheck(clFinish(s.queue), "clFinish(read)"); + } else if (result.bestSeconds < 0 || seconds < result.bestSeconds) { + result.bestSeconds = seconds; + } + + clReleaseMemObject(xBuf); + } + + clReleaseMemObject(yBuf); + clReleaseKernel(kernel); + return result; +} + +int main(int argc, char * * argv) { + const size_t globalSize = argc > 1 ? (size_t)std::atol(argv[1]) : 4096; + const uint32_t iterations = argc > 2 ? (uint32_t)std::atol(argv[2]) : 2000; + const int reps = argc > 3 ? std::atoi(argv[3]) : 5; + + std::mt19937_64 rng(0xBEEF); + std::vector x(globalSize), y(globalSize); + for (size_t i = 0; i < globalSize; ++i) { + for (int j = 0; j < MP_NWORDS; ++j) { + x[i].d[j] = (uint32_t)rng(); + y[i].d[j] = (uint32_t)rng(); + } + } + + const ClSetup s = clSetup(); + std::printf("Work items: %zu, chained mod-muls per item: %u, repetitions: %d (best time taken)\n", + globalSize, iterations, reps); + + const BenchResult oldResult = benchVariant(s, "bench_mod_mul_old", x, y, iterations, reps); + const BenchResult newResult = benchVariant(s, "bench_mod_mul_new", x, y, iterations, reps); + + size_t mismatches = 0; + for (size_t i = 0; i < globalSize; ++i) { + if (!mpEqual(oldResult.output[i], newResult.output[i])) { + ++mismatches; + } + } + if (mismatches) { + std::printf("CROSS-CHECK FAILED: outputs differ in %zu/%zu work items\n", mismatches, globalSize); + return 1; + } + std::printf("Cross-check: outputs of old and new variants are bit-exact identical (%zu work items x %u chained mod-muls).\n\n", + globalSize, iterations); + + const double totalMuls = (double)globalSize * iterations; + std::printf(" old (master): %9.2f ms (%8.2f Mmul/s)\n", oldResult.bestSeconds * 1e3, totalMuls / oldResult.bestSeconds / 1e6); + std::printf(" new (PR #49): %9.2f ms (%8.2f Mmul/s)\n", newResult.bestSeconds * 1e3, totalMuls / newResult.bestSeconds / 1e6); + std::printf(" speedup: x%.3f (%+.1f%%)\n", oldResult.bestSeconds / newResult.bestSeconds, + (oldResult.bestSeconds / newResult.bestSeconds - 1) * 100); + return 0; +} diff --git a/tests/harness.cl b/tests/harness.cl new file mode 100644 index 0000000..bacae92 --- /dev/null +++ b/tests/harness.cl @@ -0,0 +1,129 @@ +/* harness.cl + * ========== + * Test/benchmark harness appended after the multiprecision section of profanity.cl. + * + * Contains: + * 1. The pre-PR#49 ("old") implementations of mp_mul_mod_word_sub and mp_mod_mul, + * kept verbatim as a reference to test bit-exact equivalence against the + * optimized versions that now live in profanity.cl. + * 2. Thin __kernel wrappers so the host-side tests can invoke both versions. + * 3. Benchmark kernels that run a dependency-chained sequence of modular + * multiplications, so throughput of old vs new can be compared. + */ + +/* ------------------------------------------------------------------------ */ +/* Reference (pre-PR#49) implementations, verbatim from master */ +/* ------------------------------------------------------------------------ */ + +void mp_mul_mod_word_sub_old(mp_number * const r, const mp_word w, const bool withModHigher) { + // Having these numbers declared here instead of using the global values in __constant address space seems to lead + // to better optimizations by the compiler on my GTX 1070. + mp_number mod = { { 0xfffffc2f, 0xfffffffe, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff} }; + mp_number modhigher = { {0x00000000, 0xfffffc2f, 0xfffffffe, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff} }; + + mp_word cM = 0; // Carry for multiplication + mp_word cS = 0; // Carry for subtraction + mp_word tS = 0; // Temporary storage for subtraction + mp_word tM = 0; // Temporary storage for multiplication + mp_word cA = 0; // Carry for addition of modhigher + + for (mp_word i = 0; i < MP_WORDS; ++i) { + tM = (mod.d[i] * w + cM); + cM = mul_hi(mod.d[i], w) + (tM < cM); + + tM += (withModHigher ? modhigher.d[i] : 0) + cA; + cA = tM < (withModHigher ? modhigher.d[i] : 0) ? 1 : (tM == (withModHigher ? modhigher.d[i] : 0) ? cA : 0); + + tS = r->d[i] - tM - cS; + cS = tS > r->d[i] ? 1 : (tS == r->d[i] ? cS : 0); + + r->d[i] = tS; + } +} + +// Copy of mp_mod_mul from profanity.cl that calls the old subtraction routine. +void mp_mod_mul_old(mp_number * const r, const mp_number * const X, const mp_number * const Y) { + mp_number Z = { {0} }; + mp_word extraWord; + + for (int i = MP_WORDS - 1; i >= 0; --i) { + // Z = Z * 2^32 + extraWord = Z.d[7]; Z.d[7] = Z.d[6]; Z.d[6] = Z.d[5]; Z.d[5] = Z.d[4]; Z.d[4] = Z.d[3]; Z.d[3] = Z.d[2]; Z.d[2] = Z.d[1]; Z.d[1] = Z.d[0]; Z.d[0] = 0; + + // Z = Z + X * Y_i + bool overflow = mp_mul_word_add_extra(&Z, X, Y->d[i], &extraWord); + + // Z = Z - qM + mp_mul_mod_word_sub_old(&Z, extraWord, overflow); + } + + *r = Z; +} + +/* ------------------------------------------------------------------------ */ +/* Correctness / equivalence test kernels */ +/* ------------------------------------------------------------------------ */ + +__kernel void k_mul_mod_word_sub_new(__global mp_number * const r, __global const uint * const w, __global const uint * const withModHigher) { + const size_t id = get_global_id(0); + mp_number x = r[id]; + mp_mul_mod_word_sub(&x, w[id], withModHigher[id] != 0); + r[id] = x; +} + +__kernel void k_mul_mod_word_sub_old(__global mp_number * const r, __global const uint * const w, __global const uint * const withModHigher) { + const size_t id = get_global_id(0); + mp_number x = r[id]; + mp_mul_mod_word_sub_old(&x, w[id], withModHigher[id] != 0); + r[id] = x; +} + +__kernel void k_mod_mul_new(__global const mp_number * const X, __global const mp_number * const Y, __global mp_number * const R) { + const size_t id = get_global_id(0); + mp_number x = X[id]; + mp_number y = Y[id]; + mp_number r; + mp_mod_mul(&r, &x, &y); + R[id] = r; +} + +__kernel void k_mod_mul_old(__global const mp_number * const X, __global const mp_number * const Y, __global mp_number * const R) { + const size_t id = get_global_id(0); + mp_number x = X[id]; + mp_number y = Y[id]; + mp_number r; + mp_mod_mul_old(&r, &x, &y); + R[id] = r; +} + +/* ------------------------------------------------------------------------ */ +/* Benchmark kernels */ +/* ------------------------------------------------------------------------ */ +/* Each work item performs `iterations` chained modular multiplications: + * x = x * y (mod p). The data dependency between iterations prevents the + * compiler from removing or reordering the work. The final x is written back + * so nothing is dead code (and so the host can cross-check old vs new). */ + +__kernel void bench_mod_mul_new(__global mp_number * const X, __global const mp_number * const Y, const uint iterations) { + const size_t id = get_global_id(0); + mp_number x = X[id]; + const mp_number y = Y[id]; + + for (uint i = 0; i < iterations; ++i) { + mp_mod_mul(&x, &x, &y); + } + + X[id] = x; +} + +__kernel void bench_mod_mul_old(__global mp_number * const X, __global const mp_number * const Y, const uint iterations) { + const size_t id = get_global_id(0); + mp_number x = X[id]; + const mp_number y = Y[id]; + + for (uint i = 0; i < iterations; ++i) { + mp_mod_mul_old(&x, &x, &y); + } + + X[id] = x; +} diff --git a/tests/test_correctness.cpp b/tests/test_correctness.cpp new file mode 100644 index 0000000..210837e --- /dev/null +++ b/tests/test_correctness.cpp @@ -0,0 +1,267 @@ +/* test_correctness.cpp + * ==================== + * Correctness and equivalence tests for the mp_mul_mod_word_sub optimization (PR #49). + * + * Runs the actual OpenCL kernels from profanity.cl (via tests/harness.cl wrappers) and checks: + * + * 1. mp_mul_mod_word_sub (new, optimized) vs mp_mul_mod_word_sub_old (pre-PR#49 master) + * produce bit-exact identical output. + * 2. Both match an independent host-side big-integer reference: + * r' = (r - w * p - (withModHigher ? (p << 32) : 0)) mod 2^256 + * where p is the secp256k1 prime. + * 3. mp_mod_mul (which uses the new routine) vs mp_mod_mul_old are bit-exact identical + * (checked on all inputs, including out-of-domain ones >= p), and for in-domain + * inputs (field elements X, Y < p, which is how profanity uses it) the result is + * congruent to X * Y (mod p). (mp_mod_mul may return a value that is not fully + * reduced below p -- see "Cutting corners" in profanity.cl -- so congruence is the + * correct property to check.) + * + * Build & run (see tests/Makefile): + * cd tests && make && ./test_correctness.x64 [num_random_cases] + */ + +#include +#include +#include + +#include "testutil.hpp" + +static std::mt19937_64 g_rng(0xC0FFEE); + +static uint32_t randomWord() { + return (uint32_t)g_rng(); +} + +static mp_number randomNumber() { + mp_number n; + for (int i = 0; i < MP_NWORDS; ++i) { + n.d[i] = randomWord(); + } + return n; +} + +/* ------------------------------------------------------------------------ */ +/* mp_mul_mod_word_sub */ +/* ------------------------------------------------------------------------ */ + +struct WordSubCase { + mp_number r; + uint32_t w; + uint32_t withModHigher; +}; + +static std::vector genWordSubCases(const size_t numRandom) { + const mp_number zero = { { 0 } }; + const mp_number allOnes = { { 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff } }; + mp_number pPlusOne = g_p; pPlusOne.d[0] += 1; + mp_number pMinusOne = g_p; pMinusOne.d[0] -= 1; + const mp_number one = { { 1 } }; + const mp_number highBit = { { 0, 0, 0, 0, 0, 0, 0, 0x80000000 } }; + const mp_number pattern = { { 0xffffffff, 0, 0xffffffff, 0, 0xffffffff, 0, 0xffffffff, 0 } }; + + const mp_number edgeR[] = { zero, one, allOnes, g_p, pPlusOne, pMinusOne, highBit, pattern }; + const uint32_t edgeW[] = { 0, 1, 2, 976, 977, 978, 0x3D0, 0x3D1, 0x3D2, + 0x7FFFFFFF, 0x80000000, 0xFFFFFC2E, 0xFFFFFC2F, 0xFFFFFC30, + 0xFFFFFFFE, 0xFFFFFFFF }; + + std::vector cases; + for (const mp_number & r : edgeR) { + for (const uint32_t w : edgeW) { + for (uint32_t wmh = 0; wmh <= 1; ++wmh) { + cases.push_back({ r, w, wmh }); + } + } + } + for (size_t i = 0; i < numRandom; ++i) { + cases.push_back({ randomNumber(), randomWord(), (uint32_t)(g_rng() & 1) }); + } + return cases; +} + +// Runs k_mul_mod_word_sub_{old,new} over all cases in one kernel invocation +static std::vector runWordSubKernel(const ClSetup & s, const char * const kernelName, const std::vector & cases) { + const size_t count = cases.size(); + std::vector r(count); + std::vector w(count), wmh(count); + for (size_t i = 0; i < count; ++i) { + r[i] = cases[i].r; + w[i] = cases[i].w; + wmh[i] = cases[i].withModHigher; + } + + cl_int err; + cl_mem rBuf = clCreateBuffer(s.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, count * sizeof(mp_number), r.data(), &err); + clCheck(err, "clCreateBuffer(r)"); + cl_mem wBuf = clCreateBuffer(s.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, count * sizeof(uint32_t), w.data(), &err); + clCheck(err, "clCreateBuffer(w)"); + cl_mem mBuf = clCreateBuffer(s.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, count * sizeof(uint32_t), wmh.data(), &err); + clCheck(err, "clCreateBuffer(wmh)"); + + cl_kernel kernel = clCreateKernel(s.program, kernelName, &err); + clCheck(err, "clCreateKernel"); + clCheck(clSetKernelArg(kernel, 0, sizeof(cl_mem), &rBuf), "clSetKernelArg(0)"); + clCheck(clSetKernelArg(kernel, 1, sizeof(cl_mem), &wBuf), "clSetKernelArg(1)"); + clCheck(clSetKernelArg(kernel, 2, sizeof(cl_mem), &mBuf), "clSetKernelArg(2)"); + + clCheck(clEnqueueNDRangeKernel(s.queue, kernel, 1, NULL, &count, NULL, 0, NULL, NULL), "clEnqueueNDRangeKernel"); + clCheck(clEnqueueReadBuffer(s.queue, rBuf, CL_TRUE, 0, count * sizeof(mp_number), r.data(), 0, NULL, NULL), "clEnqueueReadBuffer"); + clCheck(clFinish(s.queue), "clFinish"); + + clReleaseKernel(kernel); + clReleaseMemObject(rBuf); + clReleaseMemObject(wBuf); + clReleaseMemObject(mBuf); + return r; +} + +static size_t testMulModWordSub(const ClSetup & s, const size_t numRandom) { + const std::vector cases = genWordSubCases(numRandom); + const std::vector outNew = runWordSubKernel(s, "k_mul_mod_word_sub_new", cases); + const std::vector outOld = runWordSubKernel(s, "k_mul_mod_word_sub_old", cases); + + size_t failures = 0; + for (size_t i = 0; i < cases.size(); ++i) { + if (!mpEqual(outNew[i], outOld[i])) { + if (++failures <= 5) { + std::printf(" EQUIVALENCE MISMATCH case %zu: r=%s w=%#010x wmh=%u\n", + i, mpToHex(cases[i].r).c_str(), cases[i].w, cases[i].withModHigher); + } + continue; + } + const mp_number expect = refMulModWordSub(cases[i].r, cases[i].w, cases[i].withModHigher != 0); + if (!mpEqual(outNew[i], expect)) { + if (++failures <= 5) { + std::printf(" REFERENCE MISMATCH case %zu: r=%s w=%#010x wmh=%u\n got %s\n expected %s\n", + i, mpToHex(cases[i].r).c_str(), cases[i].w, cases[i].withModHigher, + mpToHex(outNew[i]).c_str(), mpToHex(expect).c_str()); + } + } + } + + std::printf("mp_mul_mod_word_sub: %zu cases (old-vs-new bit-exact + big-int reference): %s\n", + cases.size(), failures == 0 ? "OK" : "FAILED"); + return failures; +} + +/* ------------------------------------------------------------------------ */ +/* mp_mod_mul */ +/* ------------------------------------------------------------------------ */ + +struct ModMulCase { + mp_number x; + mp_number y; +}; + +static std::vector genModMulCases(const size_t numRandom) { + const mp_number zero = { { 0 } }; + const mp_number one = { { 1 } }; + const mp_number two = { { 2 } }; + const mp_number allOnes = { { 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff } }; + mp_number pPlusOne = g_p; pPlusOne.d[0] += 1; + mp_number pMinusOne = g_p; pMinusOne.d[0] -= 1; + const mp_number highBit = { { 0, 0, 0, 0, 0, 0, 0, 0x80000000 } }; + const mp_number w977 = { { 977 } }; + const mp_number pmod = { { 0x000003D1, 1 } }; // 2^32 + 977 + + // Out-of-domain values (>= p) are included: mp_mod_mul makes no congruence + // guarantee for them (a pre-existing property of the algorithm, see "Cutting + // corners"), but old and new must still be bit-exact identical on them. + const mp_number edge[] = { zero, one, two, pMinusOne, g_p, pPlusOne, allOnes, highBit, w977, pmod }; + + std::vector cases; + for (const mp_number & x : edge) { + for (const mp_number & y : edge) { + cases.push_back({ x, y }); + } + } + for (size_t i = 0; i < numRandom; ++i) { + cases.push_back({ randomNumber(), randomNumber() }); + } + return cases; +} + +static std::vector runModMulKernel(const ClSetup & s, const char * const kernelName, const std::vector & cases) { + const size_t count = cases.size(); + std::vector x(count), y(count), out(count); + for (size_t i = 0; i < count; ++i) { + x[i] = cases[i].x; + y[i] = cases[i].y; + } + + cl_int err; + cl_mem xBuf = clCreateBuffer(s.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, count * sizeof(mp_number), x.data(), &err); + clCheck(err, "clCreateBuffer(x)"); + cl_mem yBuf = clCreateBuffer(s.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, count * sizeof(mp_number), y.data(), &err); + clCheck(err, "clCreateBuffer(y)"); + cl_mem rBuf = clCreateBuffer(s.context, CL_MEM_WRITE_ONLY, count * sizeof(mp_number), NULL, &err); + clCheck(err, "clCreateBuffer(out)"); + + cl_kernel kernel = clCreateKernel(s.program, kernelName, &err); + clCheck(err, "clCreateKernel"); + clCheck(clSetKernelArg(kernel, 0, sizeof(cl_mem), &xBuf), "clSetKernelArg(0)"); + clCheck(clSetKernelArg(kernel, 1, sizeof(cl_mem), &yBuf), "clSetKernelArg(1)"); + clCheck(clSetKernelArg(kernel, 2, sizeof(cl_mem), &rBuf), "clSetKernelArg(2)"); + + clCheck(clEnqueueNDRangeKernel(s.queue, kernel, 1, NULL, &count, NULL, 0, NULL, NULL), "clEnqueueNDRangeKernel"); + clCheck(clEnqueueReadBuffer(s.queue, rBuf, CL_TRUE, 0, count * sizeof(mp_number), out.data(), 0, NULL, NULL), "clEnqueueReadBuffer"); + clCheck(clFinish(s.queue), "clFinish"); + + clReleaseKernel(kernel); + clReleaseMemObject(xBuf); + clReleaseMemObject(yBuf); + clReleaseMemObject(rBuf); + return out; +} + +static size_t testModMul(const ClSetup & s, const size_t numRandom) { + const std::vector cases = genModMulCases(numRandom); + const std::vector outNew = runModMulKernel(s, "k_mod_mul_new", cases); + const std::vector outOld = runModMulKernel(s, "k_mod_mul_old", cases); + + size_t failures = 0; + size_t congruenceChecked = 0; + for (size_t i = 0; i < cases.size(); ++i) { + if (!mpEqual(outNew[i], outOld[i])) { + if (++failures <= 5) { + std::printf(" EQUIVALENCE MISMATCH case %zu: x=%s y=%s\n", + i, mpToHex(cases[i].x).c_str(), mpToHex(cases[i].y).c_str()); + } + continue; + } + + // Congruence is only guaranteed for in-domain inputs (field elements < p) + if (!mpGte(cases[i].x, g_p) && !mpGte(cases[i].y, g_p)) { + ++congruenceChecked; + const mp_number got = refModP256(outNew[i]); + const mp_number expect = refMulModP(cases[i].x, cases[i].y); + if (!mpEqual(got, expect)) { + if (++failures <= 5) { + std::printf(" CONGRUENCE MISMATCH case %zu: x=%s y=%s\n got%%p %s\n x*y%%p %s\n", + i, mpToHex(cases[i].x).c_str(), mpToHex(cases[i].y).c_str(), + mpToHex(got).c_str(), mpToHex(expect).c_str()); + } + } + } + } + + std::printf("mp_mod_mul: %zu cases old-vs-new bit-exact, %zu in-domain cases congruent mod p: %s\n", + cases.size(), congruenceChecked, failures == 0 ? "OK" : "FAILED"); + return failures; +} + +int main(int argc, char * * argv) { + const size_t numRandom = argc > 1 ? (size_t)std::atol(argv[1]) : 100000; + + const ClSetup s = clSetup(); + + size_t failures = 0; + failures += testMulModWordSub(s, numRandom); + failures += testModMul(s, numRandom / 10 > 1000 ? numRandom / 10 : 1000); + + if (failures) { + std::printf("FAILED: %zu total failures\n", failures); + return 1; + } + std::printf("All tests passed.\n"); + return 0; +} diff --git a/tests/testutil.hpp b/tests/testutil.hpp new file mode 100644 index 0000000..362ecc9 --- /dev/null +++ b/tests/testutil.hpp @@ -0,0 +1,284 @@ +#ifndef HPP_TESTUTIL +#define HPP_TESTUTIL + +/* testutil.hpp + * ============ + * Shared helpers for the mp-math correctness tests and benchmarks: + * - OpenCL boilerplate (context/device selection, program build from + * keccak.cl + profanity.cl + tests/harness.cl, mirroring profanity.cpp) + * - Host-side 256/512-bit reference arithmetic used to verify the kernels + * against an independent implementation. + */ + +#define CL_TARGET_OPENCL_VERSION 120 + +#if defined(__APPLE__) || defined(__MACOSX) +#include +#else +#include +#endif + +#include +#include +#include +#include +#include +#include +#include + +#include "../types.hpp" + +/* ------------------------------------------------------------------------ */ +/* OpenCL helpers */ +/* ------------------------------------------------------------------------ */ + +// The value of PROFANITY_INVERSE_SIZE does not affect the mp_* math functions +// under test; a small value keeps kernel compilation fast on CPU OpenCL +// implementations like PoCL. +static const char * const g_buildOptions = "-D PROFANITY_INVERSE_SIZE=2 -D PROFANITY_MAX_SCORE=40"; + +inline void clCheck(const cl_int err, const char * const what) { + if (err != CL_SUCCESS) { + std::cerr << "OpenCL error " << err << " during: " << what << std::endl; + std::exit(1); + } +} + +// Reads a file, looking in the current directory first and then one level up, +// so the binaries can be run both from the repository root and from tests/. +inline std::string readFile(const std::string & name) { + for (const char * const prefix : { "", "../" }) { + std::ifstream in(prefix + name, std::ios::in | std::ios::binary); + if (in) { + std::string contents; + in.seekg(0, std::ios::end); + contents.resize(in.tellg()); + in.seekg(0, std::ios::beg); + in.read(&contents[0], contents.size()); + return contents; + } + } + + std::cerr << "Could not open " << name << " (run from the repository root or from tests/)" << std::endl; + std::exit(1); +} + +struct ClSetup { + cl_device_id device; + cl_context context; + cl_command_queue queue; + cl_program program; +}; + +inline ClSetup clSetup() { + ClSetup s; + + cl_uint numPlatforms = 0; + clCheck(clGetPlatformIDs(0, NULL, &numPlatforms), "clGetPlatformIDs(count)"); + if (numPlatforms == 0) { + std::cerr << "No OpenCL platforms found" << std::endl; + std::exit(1); + } + + std::vector platforms(numPlatforms); + clCheck(clGetPlatformIDs(numPlatforms, platforms.data(), NULL), "clGetPlatformIDs"); + + // Use the first device of the first platform that has one + s.device = NULL; + for (cl_platform_id platform : platforms) { + cl_uint numDevices = 0; + if (clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, NULL, &numDevices) == CL_SUCCESS && numDevices > 0) { + clCheck(clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 1, &s.device, NULL), "clGetDeviceIDs"); + break; + } + } + if (s.device == NULL) { + std::cerr << "No OpenCL devices found" << std::endl; + std::exit(1); + } + + char deviceName[256] = { 0 }; + clGetDeviceInfo(s.device, CL_DEVICE_NAME, sizeof(deviceName) - 1, deviceName, NULL); + std::cout << "Device: " << deviceName << std::endl; + + cl_int errorCode; + s.context = clCreateContext(NULL, 1, &s.device, NULL, NULL, &errorCode); + clCheck(errorCode, "clCreateContext"); + s.queue = clCreateCommandQueue(s.context, s.device, 0, &errorCode); + clCheck(errorCode, "clCreateCommandQueue"); + + // Same source assembly as profanity.cpp, plus the test harness + const std::string strKeccak = readFile("keccak.cl"); + const std::string strVanity = readFile("profanity.cl"); + const std::string strHarness = readFile("tests/harness.cl"); + const char * szKernels[] = { strKeccak.c_str(), strVanity.c_str(), strHarness.c_str() }; + + s.program = clCreateProgramWithSource(s.context, sizeof(szKernels) / sizeof(char *), szKernels, NULL, &errorCode); + clCheck(errorCode, "clCreateProgramWithSource"); + + if (clBuildProgram(s.program, 1, &s.device, g_buildOptions, NULL, NULL) != CL_SUCCESS) { + size_t sizeLog = 0; + clGetProgramBuildInfo(s.program, s.device, CL_PROGRAM_BUILD_LOG, 0, NULL, &sizeLog); + std::string log(sizeLog, '\0'); + clGetProgramBuildInfo(s.program, s.device, CL_PROGRAM_BUILD_LOG, sizeLog, &log[0], NULL); + std::cerr << "Kernel build failed:" << std::endl << log << std::endl; + std::exit(1); + } + + return s; +} + +/* ------------------------------------------------------------------------ */ +/* Host-side reference arithmetic (independent of the OpenCL code) */ +/* ------------------------------------------------------------------------ */ +/* Numbers use the same representation as mp_number: little-endian 32-bit + * words. All arithmetic below uses plain 64-bit accumulators, no compiler + * extensions. */ + +// secp256k1 prime p = 2^256 - 2^32 - 977 +static const mp_number g_p = { { 0xfffffc2f, 0xfffffffe, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff } }; + +inline bool mpEqual(const mp_number & a, const mp_number & b) { + for (int i = 0; i < MP_NWORDS; ++i) { + if (a.d[i] != b.d[i]) { + return false; + } + } + return true; +} + +// a >= b +inline bool mpGte(const mp_number & a, const mp_number & b) { + for (int i = MP_NWORDS - 1; i >= 0; --i) { + if (a.d[i] != b.d[i]) { + return a.d[i] > b.d[i]; + } + } + return true; +} + +// a -= b (mod 2^256) +inline void mpSub(mp_number & a, const mp_number & b) { + uint64_t borrow = 0; + for (int i = 0; i < MP_NWORDS; ++i) { + const uint64_t t = (uint64_t)a.d[i] - b.d[i] - borrow; + a.d[i] = (uint32_t)t; + borrow = (t >> 32) & 1; + } +} + +inline std::string mpToHex(const mp_number & a) { + static const char digits[] = "0123456789abcdef"; + std::string out = "0x"; + for (int i = MP_NWORDS - 1; i >= 0; --i) { + for (int shift = 28; shift >= 0; shift -= 4) { + out += digits[(a.d[i] >> shift) & 0xF]; + } + } + return out; +} + +// Reference for mp_mul_mod_word_sub: +// r' = (r - w * p - (withModHigher ? (p << 32) : 0)) mod 2^256 +inline mp_number refMulModWordSub(const mp_number & r, const uint32_t w, const bool withModHigher) { + // t = w * p + (withModHigher ? p << 32 : 0), computed mod 2^256 + mp_number t = { { 0 } }; + uint64_t carry = 0; + for (int i = 0; i < MP_NWORDS; ++i) { + uint64_t acc = (uint64_t)g_p.d[i] * w + carry; + if (withModHigher && i > 0) { + acc += g_p.d[i - 1]; + } + t.d[i] = (uint32_t)acc; + carry = acc >> 32; + } + + mp_number out = r; + mpSub(out, t); + return out; +} + +// Full 512-bit product of two 256-bit numbers (schoolbook) +inline void refMul512(const mp_number & a, const mp_number & b, uint32_t out[MP_NWORDS * 2]) { + for (int i = 0; i < MP_NWORDS * 2; ++i) { + out[i] = 0; + } + + for (int i = 0; i < MP_NWORDS; ++i) { + uint64_t carry = 0; + for (int j = 0; j < MP_NWORDS; ++j) { + const uint64_t t = (uint64_t)a.d[i] * b.d[j] + out[i + j] + carry; + out[i + j] = (uint32_t)t; + carry = t >> 32; + } + out[i + MP_NWORDS] = (uint32_t)carry; + } +} + +// Reduces a value of up to 544 bits (17 words) modulo p. +// Uses the same identity as the optimization under test, on the host side: +// v = hi * 2^256 + lo == hi * pmod + lo (mod p), where pmod = 2^32 + 977. +inline mp_number refModP(const uint32_t * const v, const int numWords) { + uint32_t cur[MP_NWORDS * 2 + 1] = { 0 }; + for (int i = 0; i < numWords; ++i) { + cur[i] = v[i]; + } + + for (;;) { + bool hasHigh = false; + for (int i = MP_NWORDS; i < MP_NWORDS * 2 + 1; ++i) { + if (cur[i] != 0) { + hasHigh = true; + break; + } + } + if (!hasHigh) { + break; + } + + // next = lo + hi * 977 + (hi << 32) + uint32_t next[MP_NWORDS * 2 + 1] = { 0 }; + uint64_t carry = 0; + for (int i = 0; i < MP_NWORDS * 2 + 1; ++i) { + uint64_t acc = carry; + if (i < MP_NWORDS) { + acc += cur[i]; // lo + } + const int hiIdx = MP_NWORDS + i; + if (hiIdx < MP_NWORDS * 2 + 1) { + acc += (uint64_t)cur[hiIdx] * 977; // hi * 977 + } + if (i >= 1 && MP_NWORDS + i - 1 < MP_NWORDS * 2 + 1) { + acc += cur[MP_NWORDS + i - 1]; // hi << 32 + } + next[i] = (uint32_t)acc; + carry = acc >> 32; + } + for (int i = 0; i < MP_NWORDS * 2 + 1; ++i) { + cur[i] = next[i]; + } + } + + mp_number out; + for (int i = 0; i < MP_NWORDS; ++i) { + out.d[i] = cur[i]; + } + while (mpGte(out, g_p)) { + mpSub(out, g_p); + } + return out; +} + +// (x * y) mod p via the reference implementation +inline mp_number refMulModP(const mp_number & x, const mp_number & y) { + uint32_t product[MP_NWORDS * 2]; + refMul512(x, y, product); + return refModP(product, MP_NWORDS * 2); +} + +// value mod p for a 256-bit value +inline mp_number refModP256(const mp_number & v) { + return refModP(v.d, MP_NWORDS); +} + +#endif /* HPP_TESTUTIL */ From db2617ac51b7f19316034bd6709a30fbdb8046f4 Mon Sep 17 00:00:00 2001 From: Cursor Agent Date: Wed, 22 Jul 2026 15:48:59 +0000 Subject: [PATCH 5/6] Remove stray Python bytecode cache accidentally committed with tests Co-authored-by: Gleb Alekseev --- .gitignore | 1 + tests/__pycache__/clutil.cpython-312.pyc | Bin 3231 -> 0 bytes 2 files changed, 1 insertion(+) delete mode 100644 tests/__pycache__/clutil.cpython-312.pyc diff --git a/.gitignore b/.gitignore index 2fcd8e3..dc76ad4 100644 --- a/.gitignore +++ b/.gitignore @@ -4,3 +4,4 @@ *.x64 cache-opencl.* bin +__pycache__/ diff --git a/tests/__pycache__/clutil.cpython-312.pyc b/tests/__pycache__/clutil.cpython-312.pyc deleted file mode 100644 index ca281ee58fc8abcf2aca5cf94c8fdc8f89a77f7b..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 3231 zcmbtWO>7&-6`t8$F3BaSKPQT8#K^`$Q&G00O19j_ah&QuN)%BQm8N2vvMbJtBDLhQ zv#URXLU`ywY!y&K3`j--LR@E{pHRKtc%I2&pXQ)1-q!Z?Hx17qf zrxbHmHdVtkWF<}HX*C_2O(}HF=#Zqo`GiIoUk+u|bpI$*jZ-Bh$Mm$3Fw8_cF6%S$ zoEnQMa~&~FzF0F}p4x^xVIOJLFo1PQN+l@ODbt+Q7f*_3GO}VAYHC_j4SSz+Drjl4 zq0iHp3fnng9vVCF0M(j+bmktt_pkRJ{f8e$2h6Pp2xzvaeLx+e*CF^&M9iKB~A7`=eHkbxIDjHckk z7%ihEehMw(AU9@V8Wb%-hcH^4ZdgJF!fpvkJ(0G&6NBTS=tL+qX$e$SNRYFHw31RS zJY%q($@0Td&+Ps<2Ix@?am+3@j5}HaGvRmM9ybL6d_Pl|$>0-PP zFKPwtqnEZxsqW%Vxa93RMA+5x=}&g+8;T2sh2JiH%ftLz#yn!=p}`J*uc6n=|H

zZE`;r1eWf05DsAgNsW_E7=Te}CX=um;i@Thg2z9}Go8RCS(v}mPFsTr#8nMg)n44O zyn;<{)kxAVQPo+w&!LN?*hrA-9Te;sxbAn|LOI)lK;XU^y!lSd(hvU z-i&X=H?<9I`-eNDrRJAP{%ePB!7Cl}Vx5>D`ez;T*A=f7uC2$mu9kckzVx*|@U?z^ z{=ToZfr#JKA zS6{v_G*jmI$8v=g9@`^-g~28}6En`mlK$Wt7~Mz-U`3gB7nb7%ZZ$mT^OwmFe!+M zD|@4L?2`PWpga$8x!f&l3DeZnb~Q~BO1d2i7m*c8m1Vj827EJ2c4dt(o3naa1_h(T z4k2affA9#}fefRd+Y%MST+XN#m(ExsGkx`n<)TVDu3DZq$D?nDCI-S54}6f!KwTw4 zm&L3olqWVYtAa%vw2N>eKV7X}FK0rXFSM)64L^1OCRZfaKA>6ezY8 z+HOv4PHjw;q~{>qk{?KY(-9I}&8z&bKd_4T8k=t}Z9l(}+n(5IDm7k$!s%^Y?cHr` zTA$u{=A-^?e!KT~;)icMX!w4i_mGE=`yO9WDoE>H?*|UOVCYan{`&muPlCp|<@R7d z|1FJ7b~?MkJyrpi=@ifz&a8e$N$19MSm>@(0-;>iVBuH-=s*Vi900n)0o1Yzy=wBDhs>HLWLuI7WS0G3}`dlby#@qDFLlrtUyP|-WYun5DX{(8=^X(7`O;S z34IZ4YYg!0J;DQiD9c~Vl~pNRx4fpN=|#q{k#xD+@~Qv>C?Zfu+JFz(5f>C~UNvYZ z6T?Z`1wHKs@&tewY8ZMdK%(qi*O=fOkUaWY@)uhRtvCC(#1G#1q~-pZw$Cs8A^2JF z3+Z~kf3M-}){CWv=Z_HQ>M01Tc(u!h(g)8m1oZ%d&Ys(%+r76IZk^w0d(d)awSOJ2 zcdZM*8TkfK)N{y#-l157PVgzq9gPw_7LD3>isgcBWtPjw_hgpYlOv|t^|5^2xJiQA z&kMIQ4KE6KhIpd(^8pO7CwSRUc-#^v-wt*6T&1i+l Date: Wed, 22 Jul 2026 13:43:15 -0300 Subject: [PATCH 6/6] make fix and renamings --- Makefile | 2 +- tests/Makefile | 14 +++++++------- .../{bench_mod_mul.cpp => bench_mod_mul_pr49.cpp} | 6 +++--- ...t_correctness.cpp => test_correctness_pr49.cpp} | 6 +++--- 4 files changed, 14 insertions(+), 14 deletions(-) rename tests/{bench_mod_mul.cpp => bench_mod_mul_pr49.cpp} (96%) rename tests/{test_correctness.cpp => test_correctness_pr49.cpp} (98%) diff --git a/Makefile b/Makefile index 3119cd9..b1df3f0 100644 --- a/Makefile +++ b/Makefile @@ -7,7 +7,7 @@ EXECUTABLE=profanity2.x64 UNAME_S := $(shell uname -s) ifeq ($(UNAME_S),Darwin) LDFLAGS=-framework OpenCL - CFLAGS=-c -std=c++11 -Wall -mmmx -O2 + CFLAGS=-c -std=c++11 -Wall -O2 else LDFLAGS=-s -lOpenCL -mcmodel=large CFLAGS=-c -std=c++11 -Wall -mmmx -O2 -mcmodel=large diff --git a/tests/Makefile b/tests/Makefile index dbf27a5..a9b29c7 100644 --- a/tests/Makefile +++ b/tests/Makefile @@ -4,12 +4,12 @@ # # Usage: # make -# ./test_correctness.x64 [num_random_cases] -# ./bench_mod_mul.x64 [global_size] [iterations] [repetitions] +# ./test_correctness_pr49.x64 [num_random_cases] +# ./bench_mod_mul_pr49.x64 [global_size] [iterations] [repetitions] CC=g++ CDEFINES= -EXECUTABLES=test_correctness.x64 bench_mod_mul.x64 +EXECUTABLES=test_correctness_pr49.x64 bench_mod_mul_pr49.x64 UNAME_S := $(shell uname -s) ifeq ($(UNAME_S),Darwin) @@ -22,11 +22,11 @@ endif all: $(EXECUTABLES) -test_correctness.x64: test_correctness.cpp testutil.hpp ../types.hpp - $(CC) $(CFLAGS) $(CDEFINES) test_correctness.cpp $(LDFLAGS) -o $@ +test_correctness_pr49.x64: test_correctness_pr49.cpp testutil.hpp ../types.hpp + $(CC) $(CFLAGS) $(CDEFINES) test_correctness_pr49.cpp $(LDFLAGS) -o $@ -bench_mod_mul.x64: bench_mod_mul.cpp testutil.hpp ../types.hpp - $(CC) $(CFLAGS) $(CDEFINES) bench_mod_mul.cpp $(LDFLAGS) -o $@ +bench_mod_mul_pr49.x64: bench_mod_mul_pr49.cpp testutil.hpp ../types.hpp + $(CC) $(CFLAGS) $(CDEFINES) bench_mod_mul_pr49.cpp $(LDFLAGS) -o $@ clean: rm -rf $(EXECUTABLES) diff --git a/tests/bench_mod_mul.cpp b/tests/bench_mod_mul_pr49.cpp similarity index 96% rename from tests/bench_mod_mul.cpp rename to tests/bench_mod_mul_pr49.cpp index 7483a05..6f2e0ce 100644 --- a/tests/bench_mod_mul.cpp +++ b/tests/bench_mod_mul_pr49.cpp @@ -1,5 +1,5 @@ -/* bench_mod_mul.cpp - * ================= +/* bench_mod_mul_pr49.cpp + * ====================== * Benchmark: mp_mod_mul with the old (pre-PR#49) vs new (optimized) mp_mul_mod_word_sub. * * Each work item performs a dependency-chained sequence of modular multiplications @@ -14,7 +14,7 @@ * exactly the code path the PR optimizes. * * Build & run (see tests/Makefile): - * cd tests && make && ./bench_mod_mul.x64 [global_size] [iterations] [repetitions] + * cd tests && make && ./bench_mod_mul_pr49.x64 [global_size] [iterations] [repetitions] */ #include diff --git a/tests/test_correctness.cpp b/tests/test_correctness_pr49.cpp similarity index 98% rename from tests/test_correctness.cpp rename to tests/test_correctness_pr49.cpp index 210837e..d30d3f8 100644 --- a/tests/test_correctness.cpp +++ b/tests/test_correctness_pr49.cpp @@ -1,5 +1,5 @@ -/* test_correctness.cpp - * ==================== +/* test_correctness_pr49.cpp + * ========================= * Correctness and equivalence tests for the mp_mul_mod_word_sub optimization (PR #49). * * Runs the actual OpenCL kernels from profanity.cl (via tests/harness.cl wrappers) and checks: @@ -17,7 +17,7 @@ * correct property to check.) * * Build & run (see tests/Makefile): - * cd tests && make && ./test_correctness.x64 [num_random_cases] + * cd tests && make && ./test_correctness_pr49.x64 [num_random_cases] */ #include