diff --git a/proto-opencl/README.md b/proto-opencl/README.md index 29adb4dc..fc6a3783 100644 --- a/proto-opencl/README.md +++ b/proto-opencl/README.md @@ -19,11 +19,26 @@ Status on 3 October 2026: no AMD device has run this yet. Everything that could and a CPU emulator with 32- and 64-wide sub-groups all give the Mac's 96/96 vectors and cache FNV for the memory-hard pack. The AMD run itself is the next step and the commands for it are below. +## The one-click worker mode (--pack, 4 October 2026) + +`host.c --serve --pack ` serves a pack read at run time (`../proto-cuda/nvrtc/packfile.h`: program.h, seeds.txt, +vectors.h, kernel_bound.cl), whatever pack the exe was built against, so one prebuilt exe serves every hourly program. +The first pair is built and self-tested exactly like a prepared one (cache head, last line and FNV-1a 64, dataset head, +last word and 64 samples, the three vector warps through `igneum_hash_bound` with the pack's seed words), and every +prepared pair is self-tested too; the `ready` line ends in `path prebuilt-generic`. The Windows package ships this as +`igneum-worker-opencl.exe`, built by `../proto-cuda/nvrtc/build-windows.sh` with `IGNEUM_CL_DYNAMIC` (`cl_dynamic.h`: +OpenCL.dll opened with LoadLibrary, every entry point a function pointer, no import library; nothing to install but +the driver) and the Khronos headers fetched by `fetch-redist.sh`. `test-generic.sh` checks the mode here through Apple +OpenCL with the two packs `proto-cuda/nvrtc/emu/test.sh` writes (PASS on 4 October 2026: 192 found lines, prepare and +swap, 15 sampled hashes equal to `igneum-pow hash-bound`). The bench and the compiled-in serve mode are unchanged. + ## Layout ``` proto-opencl/ - host.c C99 host: device list, runtime kernel build, cache + dataset fill, self-tests, vectors, bench, sweep + host.c C99 host: device list, runtime kernel build, cache + dataset fill, self-tests, vectors, bench, sweep, --serve, --pack + cl_dynamic.h Windows one-click build: OpenCL.dll loaded at run time (IGNEUM_CL_DYNAMIC) + test-generic.sh the --pack mode checked here through Apple OpenCL (needs proto-cuda/nvrtc/emu/test.sh's packs) build.sh macOS (-framework OpenCL, or the Khronos ICD loader) and Linux (-lOpenCL) build.bat Windows (MSVC cl.exe + OpenCL.lib) WAVEFRONT.md wave32 vs wave64 on AMD, and why the kernel cannot tell the difference diff --git a/proto-opencl/cl_dynamic.h b/proto-opencl/cl_dynamic.h new file mode 100644 index 00000000..c146499f --- /dev/null +++ b/proto-opencl/cl_dynamic.h @@ -0,0 +1,78 @@ +/* cl_dynamic.h: OpenCL.dll loaded at run time (Windows). 4 October 2026. + * + * The one-click AMD worker (host.c cross-compiled with mingw, build-windows.sh) links no OpenCL import library: + * every cl* entry point host.c calls is a function pointer filled from OpenCL.dll with LoadLibrary/GetProcAddress. + * OpenCL.dll is the Khronos ICD loader that the AMD Adrenalin (and NVIDIA, Intel) driver installs in System32, so + * nothing but the driver is needed on the PC. Included by host.c after when IGNEUM_CL_DYNAMIC is defined; + * the macros below then rename the calls in the rest of the file. The Mac and Linux builds link the library as before. + */ +#ifndef IGNEUM_CL_DYNAMIC_H +#define IGNEUM_CL_DYNAMIC_H +#ifdef _WIN32 +#define WIN32_LEAN_AND_MEAN +#include +#endif +#include + +#define IG_CL_FUNCS(X) \ + X(clGetPlatformIDs) X(clGetPlatformInfo) X(clGetDeviceIDs) X(clGetDeviceInfo) X(clCreateContext) \ + X(clCreateCommandQueue) X(clCreateProgramWithSource) X(clBuildProgram) X(clGetProgramBuildInfo) X(clCreateKernel) \ + X(clSetKernelArg) X(clEnqueueNDRangeKernel) X(clEnqueueReadBuffer) X(clEnqueueWriteBuffer) X(clCreateBuffer) \ + X(clReleaseMemObject) X(clReleaseKernel) X(clReleaseProgram) X(clReleaseCommandQueue) X(clReleaseContext) \ + X(clFinish) X(clWaitForEvents) X(clReleaseEvent) X(clGetEventProfilingInfo) X(clGetKernelWorkGroupInfo) \ + X(clGetExtensionFunctionAddressForPlatform) + +/* One pointer per entry point, typed from the header's declaration (so a signature drift is a compile error). */ +#define IG_CL_DECL(f) static __typeof__(f)* ig_##f = NULL; +IG_CL_FUNCS(IG_CL_DECL) +#undef IG_CL_DECL + +static char ig_cl_error[512]; + +/* Returns 1 when OpenCL.dll and every entry point are there; else 0 with ig_cl_error set. */ +static int ig_cl_load(void) { +#ifdef _WIN32 + HMODULE h = LoadLibraryA("OpenCL.dll"); + if (!h) { + snprintf(ig_cl_error, sizeof(ig_cl_error), "OpenCL.dll is not installed (the GPU driver provides it: AMD Adrenalin, or the NVIDIA or Intel driver); nothing else is needed"); + return 0; + } +#define IG_CL_LOAD(f) ig_##f = (__typeof__(f)*)(void*)GetProcAddress(h, #f); if (!ig_##f) { snprintf(ig_cl_error, sizeof(ig_cl_error), "OpenCL.dll lacks %s", #f); return 0; } + IG_CL_FUNCS(IG_CL_LOAD) +#undef IG_CL_LOAD + return 1; +#else + snprintf(ig_cl_error, sizeof(ig_cl_error), "IGNEUM_CL_DYNAMIC is for the Windows build"); + return 0; +#endif +} + +/* From here on, host.c's calls go through the pointers. */ +#define IG_CL_RENAME(f) +#define clGetPlatformIDs ig_clGetPlatformIDs +#define clGetPlatformInfo ig_clGetPlatformInfo +#define clGetDeviceIDs ig_clGetDeviceIDs +#define clGetDeviceInfo ig_clGetDeviceInfo +#define clCreateContext ig_clCreateContext +#define clCreateCommandQueue ig_clCreateCommandQueue +#define clCreateProgramWithSource ig_clCreateProgramWithSource +#define clBuildProgram ig_clBuildProgram +#define clGetProgramBuildInfo ig_clGetProgramBuildInfo +#define clCreateKernel ig_clCreateKernel +#define clSetKernelArg ig_clSetKernelArg +#define clEnqueueNDRangeKernel ig_clEnqueueNDRangeKernel +#define clEnqueueReadBuffer ig_clEnqueueReadBuffer +#define clEnqueueWriteBuffer ig_clEnqueueWriteBuffer +#define clCreateBuffer ig_clCreateBuffer +#define clReleaseMemObject ig_clReleaseMemObject +#define clReleaseKernel ig_clReleaseKernel +#define clReleaseProgram ig_clReleaseProgram +#define clReleaseCommandQueue ig_clReleaseCommandQueue +#define clReleaseContext ig_clReleaseContext +#define clFinish ig_clFinish +#define clWaitForEvents ig_clWaitForEvents +#define clReleaseEvent ig_clReleaseEvent +#define clGetEventProfilingInfo ig_clGetEventProfilingInfo +#define clGetKernelWorkGroupInfo ig_clGetKernelWorkGroupInfo +#define clGetExtensionFunctionAddressForPlatform ig_clGetExtensionFunctionAddressForPlatform +#endif diff --git a/proto-opencl/host.c b/proto-opencl/host.c index 183cdc70..84435d78 100644 --- a/proto-opencl/host.c +++ b/proto-opencl/host.c @@ -19,6 +19,9 @@ #else #include #endif +#ifdef IGNEUM_CL_DYNAMIC +#include "cl_dynamic.h" /* Windows one-click build: OpenCL.dll loaded at run time, no import library */ +#endif #include #include @@ -37,6 +40,7 @@ #define IGNEUM_NO_CUDA #include "program.h" #include "vectors.h" +#include "../proto-cuda/nvrtc/packfile.h" /* --pack: a pack read at run time (generic serve mode, 4 October 2026) */ #ifndef IGNEUM_DATASET_MODE #define IGNEUM_DATASET_MODE 0 @@ -196,6 +200,8 @@ typedef struct { int noPrepare; // --no-prepare: serve without the prepare command (ready line says "prepare 0"), to test the miner's fallback int kernelGiven; // --kernel was passed const char* vendor; // --vendor S: pick the first GPU whose vendor string contains S (default: first GPU of any vendor) + const char* packDir; // --pack D (serve only): generic mode, the pack is read from D at run time; the compiled-in pack is + // then only the build-time placeholder of the prebuilt exe (4 October 2026) } Options; static int packMib(void) { return (int)(((1ull << IGNEUM_DATASET_LOG2) * 4ull) >> 20); } @@ -220,7 +226,9 @@ static void usage(void) { " --vendor S choose the first GPU whose vendor string contains S (for example \"Advanced Micro Devices\"); fails if none\n" " --serve GPU worker for igneum-miner --worker: reads \"job ...\" lines on stdin, prints found/done lines.\n" " --no-prepare with --serve: no prepare support (the miner then falls back to exit 42 at a seed change).\n" - " Builds the pack's kernel_bound.cl (next to the compiled-in kernel.cl) unless --kernel says otherwise.\n", packMib(), IGNEUM_KERNEL_PATH); + " Builds the pack's kernel_bound.cl (next to the compiled-in kernel.cl) unless --kernel says otherwise.\n" + " --pack D with --serve: serve the pack in directory D (program.h, seeds.txt, vectors.h, kernel_bound.cl), whatever\n" + " pack this exe was built against; it is self-tested against its vectors.h first (the one-click worker)\n", packMib(), IGNEUM_KERNEL_PATH); } static int isPow2(long long v) { return v > 0 && (v & (v - 1)) == 0; } @@ -230,7 +238,7 @@ static Options parseArgs(int argc, char** argv) { Options o; int i; o.datasetMib = 1024; o.batchLog2 = 24; o.batches = 5; o.groupWarps = 1; o.sweep = 0; o.device = -1; - o.exchange = 0; o.list = 0; o.timeWall = -1; o.kernelPath = IGNEUM_KERNEL_PATH; o.extraOpts = ""; o.serve = 0; o.noPrepare = 0; o.kernelGiven = 0; o.vendor = NULL; + o.exchange = 0; o.list = 0; o.timeWall = -1; o.kernelPath = IGNEUM_KERNEL_PATH; o.extraOpts = ""; o.serve = 0; o.noPrepare = 0; o.kernelGiven = 0; o.vendor = NULL; o.packDir = NULL; for (i = 1; i < argc; ++i) { const char* a = argv[i]; int needs = (strcmp(a, "--dataset-mib") == 0 || strcmp(a, "--batch-log2") == 0 || strcmp(a, "--batches") == 0 || @@ -246,6 +254,7 @@ static Options parseArgs(int argc, char** argv) { else if (strcmp(a, "--serve") == 0) o.serve = 1; else if (strcmp(a, "--no-prepare") == 0) o.noPrepare = 1; else if (strcmp(a, "--vendor") == 0) { if (i + 1 >= argc) { usage(); exit(2); } o.vendor = argv[++i]; } + else if (strcmp(a, "--pack") == 0) { if (i + 1 >= argc) { usage(); exit(2); } o.packDir = argv[++i]; } else if (strcmp(a, "--build-opts") == 0) o.extraOpts = argv[++i]; else if (strcmp(a, "--time") == 0) { const char* m = argv[++i]; @@ -394,6 +403,7 @@ typedef struct { size_t subGroupSize; // as queried for a 32-item work-group, 0 if not queried char exchangeNote[512]; char buildOptions[512]; + int groupSize; // work-group size the program was built for (IGNEUM_GROUP = 32 x group-warps) } Device; static const char* exchangeName(int m) { return m == 1 ? "sub_group_shuffle_xor (cl_khr_subgroup_shuffle)" : m == 2 ? "intel_sub_group_shuffle_xor (cl_intel_subgroups)" : "local-memory exchange with barrier"; } @@ -465,6 +475,7 @@ static size_t querySubGroupSize(const Device* dv, const DeviceInfo* di, size_t l static void setupProgram(Device* dv, const DeviceInfo* di, const Options* o, const char* src, size_t srcLen) { int groupSize = 32 * o->groupWarps; int want = 0; + dv->groupSize = groupSize; const char* why = ""; if (strstr(di->extensions, "cl_khr_subgroup_shuffle")) want = 1; else if (strstr(di->extensions, "cl_intel_subgroups")) want = 2; @@ -896,9 +907,19 @@ static int unhexBuf(const char* s, uint8_t* out, size_t cap, size_t* len) { return 1; } +/* Generic serve mode (--pack, 4 October 2026): the sizes and seeds come from the pack directory read at run time, not + * from the compiled-in program.h, so one prebuilt exe serves every pack. The compiled-in values are the defaults. */ +static PfPack gPack; +static int gGeneric = 0; +static uint32_t gServeWords = 1u << IGNEUM_DATASET_LOG2; +#if IGNEUM_DATASET_MODE == 1 +static uint32_t gServeCacheWords = 1u << IGNEUM_CACHE_LOG2_WORDS; +static uint32_t gServeSegments = IGNEUM_CACHE_SEGMENTS; +#endif + #if IGNEUM_DATASET_MODE == 1 /* One resident (program, cache, dataset) triple for a seed pair. The first one is the compiled-in pack on the main - * queue; prepared ones are built from a pack directory on their own queue. */ + * queue (or, with --pack, the directory's pack); prepared ones are built from a pack directory on their own queue. */ typedef struct { char epochHex[65]; char dayHex[512]; @@ -906,7 +927,9 @@ typedef struct { cl_program prog; cl_kernel kHashBound, kCacheFill, kBuild; cl_mem cache, ds; - double buildMs, cacheMs, datasetMs; + double buildMs, cacheMs, datasetMs, checkMs; + char check[1024]; /* the self-test verdict (packfile.h), one line */ + int checked; } ServePair; static void releasePair(ServePair* p) { @@ -927,7 +950,7 @@ typedef struct { char epochHex[65]; char dayHex[512]; char packDir[1024]; - uint32_t words; + uint32_t words, cacheWords, segments; char error[512]; ServePair* result; /* set by the thread on success */ volatile int done; /* 1 when the thread has finished (success or failure) */ @@ -938,6 +961,96 @@ static void prepareFail(PrepareTask* t, const char* what, cl_int err) { snprintf(t->error, sizeof(t->error), "%s (%s)", what, clErrName(err)); } +/* The self-test of packfile.h on a pair: cache head, last line and FNV-1a 64 over the whole cache, dataset head, last + * word and samples, and the vector warps through igneum_hash_bound with the pack's own seed words as init words (that + * is igneum_hash of kernel.cl). The verdict goes to p->check. A pack directory without a readable vectors.h is not + * tested (p->checked stays 0: the miner's CPU re-check still covers every found nonce); a FAIL is an error. */ +static int pairSelfTest(Device* dv, const DeviceInfo* di, cl_command_queue q, ServePair* p, uint32_t words, uint32_t cacheWords, const char* packDir, char* err, size_t errCap) { + PfPack pk; + char perr[256]; + uint32_t cacheHead[16], cacheLast[16], dsHead[16], dsLast = 0; + uint32_t samples[PF_MAX_SAMPLES]; + uint64_t vec[PF_MAX_WARPS * 32]; + uint64_t fnv; + uint32_t* whole; + cl_mem out = NULL, init = NULL; + cl_int e = 0; + int w, i, ok; + size_t local = (size_t)dv->groupSize, g = (size_t)dv->groupSize; /* one work-group; its first 32 lanes are the warp */ + double tb = wallMs(); + (void)di; + p->checked = 0; + perr[0] = 0; + if (!packDir || !packDir[0] || !pf_load(packDir, &pk, perr, sizeof(perr)) || !pk.haveVectors) { + snprintf(p->check, sizeof(p->check), "self-test skipped (%s); the miner's CPU re-check covers every found nonce", (packDir && packDir[0]) ? (perr[0] ? perr : "no vectors.h in the pack") : "no pack directory"); + return 1; + } + whole = (uint32_t*)malloc((size_t)cacheWords * 4u); + if (!whole) { snprintf(err, errCap, "self-test: no host memory for the cache read-back"); return 0; } + e = clEnqueueReadBuffer(q, p->cache, CL_TRUE, 0, (size_t)cacheWords * 4u, whole, 0, NULL, NULL); + if (e != CL_SUCCESS) { free(whole); snprintf(err, errCap, "self-test: cache read-back (%s)", clErrName(e)); return 0; } + memcpy(cacheHead, whole, 64); + memcpy(cacheLast, whole + cacheWords - 16u, 64); + fnv = pf_fnv1a64(whole, (size_t)cacheWords * 4u); + free(whole); + e = clEnqueueReadBuffer(q, p->ds, CL_TRUE, 0, 64, dsHead, 0, NULL, NULL); + if (e == CL_SUCCESS && pk.dsLastIndex < words) e = clEnqueueReadBuffer(q, p->ds, CL_TRUE, (size_t)pk.dsLastIndex * 4u, 4, &dsLast, 0, NULL, NULL); + for (i = 0; i < pk.nSamples && e == CL_SUCCESS; ++i) { samples[i] = 0; if (pk.sampleIdx[i] < words) e = clEnqueueReadBuffer(q, p->ds, CL_TRUE, (size_t)pk.sampleIdx[i] * 4u, 4, &samples[i], 0, NULL, NULL); } + if (e != CL_SUCCESS) { snprintf(err, errCap, "self-test: dataset read-back (%s)", clErrName(e)); return 0; } + out = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, g * sizeof(uint64_t), NULL, &e); + if (e == CL_SUCCESS) init = clCreateBuffer(dv->ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 32, pk.seedw, &e); + for (w = 0; w < pk.vecWarps && e == CL_SUCCESS; ++w) { + cl_uint base = pk.vecBase[w], mask = words - 1u; + e = clSetKernelArg(p->kHashBound, 0, sizeof(cl_mem), &p->ds); + if (e == CL_SUCCESS) e = clSetKernelArg(p->kHashBound, 1, sizeof(cl_mem), &out); + if (e == CL_SUCCESS) e = clSetKernelArg(p->kHashBound, 2, sizeof(cl_uint), &base); + if (e == CL_SUCCESS) e = clSetKernelArg(p->kHashBound, 3, sizeof(cl_uint), &mask); + if (e == CL_SUCCESS) e = clSetKernelArg(p->kHashBound, 4, sizeof(cl_mem), &init); + if (e == CL_SUCCESS) e = clEnqueueNDRangeKernel(q, p->kHashBound, 1, NULL, &g, &local, 0, NULL, NULL); + if (e == CL_SUCCESS) e = clEnqueueReadBuffer(q, out, CL_TRUE, 0, 32 * sizeof(uint64_t), &vec[w * 32], 0, NULL, NULL); + } + if (out) clReleaseMemObject(out); + if (init) clReleaseMemObject(init); + if (e != CL_SUCCESS) { snprintf(err, errCap, "self-test: vector warp (%s)", clErrName(e)); return 0; } + ok = pf_selftest(&pk, cacheHead, cacheLast, fnv, dsHead, dsLast, samples, vec, p->check, sizeof(p->check)); + p->checked = 1; + p->checkMs = wallMs() - tb; + if (!ok) { snprintf(err, errCap, "%s", p->check); return 0; } + return 1; +} + +/* Fills the pair's cache and builds its dataset on queue q (the pair's kernels must exist), then self-tests it. + * Shared by the first pair of --pack and by every prepared pair. Returns 1, or 0 with err set. */ +static int pairBuffers(Device* dv, const DeviceInfo* di, cl_command_queue q, ServePair* p, uint32_t words, uint32_t cacheWords, uint32_t segments, const char* packDir, char* err, size_t errCap) { + cl_int e = 0; + cl_uint nSeg = segments, nItems = words / 16u; + size_t local, g, bytes = (size_t)cacheWords * 4u; + double tb = wallMs(); + p->cache = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, bytes, NULL, &e); + if (e != CL_SUCCESS) { snprintf(err, errCap, "clCreateBuffer cache (%s)", clErrName(e)); return 0; } + local = kernelMaxLocal(dv, p->kCacheFill, di, 256); + g = ((nSeg + local - 1) / local) * local; + e = clSetKernelArg(p->kCacheFill, 0, sizeof(cl_mem), &p->cache); + if (e == CL_SUCCESS) e = clSetKernelArg(p->kCacheFill, 1, sizeof(cl_uint), &nSeg); + if (e == CL_SUCCESS) e = clEnqueueNDRangeKernel(q, p->kCacheFill, 1, NULL, &g, &local, 0, NULL, NULL); + if (e == CL_SUCCESS) e = clFinish(q); + if (e != CL_SUCCESS) { snprintf(err, errCap, "cache fill (%s)", clErrName(e)); return 0; } + p->cacheMs = wallMs() - tb; + tb = wallMs(); + p->ds = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)words * 4u, NULL, &e); + if (e != CL_SUCCESS) { snprintf(err, errCap, "clCreateBuffer dataset (%s)", clErrName(e)); return 0; } + local = kernelMaxLocal(dv, p->kBuild, di, 256); + g = ((nItems + local - 1) / local) * local; + e = clSetKernelArg(p->kBuild, 0, sizeof(cl_mem), &p->ds); + if (e == CL_SUCCESS) e = clSetKernelArg(p->kBuild, 1, sizeof(cl_mem), &p->cache); + if (e == CL_SUCCESS) e = clSetKernelArg(p->kBuild, 2, sizeof(cl_uint), &nItems); + if (e == CL_SUCCESS) e = clEnqueueNDRangeKernel(q, p->kBuild, 1, NULL, &g, &local, 0, NULL, NULL); + if (e == CL_SUCCESS) e = clFinish(q); + if (e != CL_SUCCESS) { snprintf(err, errCap, "dataset build (%s)", clErrName(e)); return 0; } + p->datasetMs = wallMs() - tb; + return pairSelfTest(dv, di, q, p, words, cacheWords, packDir, err, errCap); +} + /* Builds the pair for a prepare request. Runs on its own thread with its own command queue. */ static void prepareRun(PrepareTask* t) { cl_int err = 0; @@ -947,8 +1060,6 @@ static void prepareRun(PrepareTask* t) { ServePair* p = (ServePair*)calloc(1, sizeof(ServePair)); cl_command_queue q = NULL; double tb; - cl_uint nSeg = IGNEUM_CACHE_SEGMENTS, nItems; - size_t local, bytes = (size_t)CACHE_WORDS_HOST * 4u; strncpy(p->epochHex, t->epochHex, 64); p->epochHex[64] = 0; strncpy(p->dayHex, t->dayHex, sizeof(p->dayHex) - 1); snprintf(path, sizeof(path), "%s/kernel_bound.cl", t->packDir); @@ -977,36 +1088,7 @@ static void prepareRun(PrepareTask* t) { if (err != CL_SUCCESS) { prepareFail(t, "clCreateKernel igneum_build", err); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; } q = clCreateCommandQueue(t->dv->ctx, t->di->device, 0, &err); if (err != CL_SUCCESS) { prepareFail(t, "clCreateCommandQueue (prepare)", err); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; } - /* Cache: the same segment count as the compiled-in pack (the dataset schedule is a network constant) */ - tb = wallMs(); - p->cache = clCreateBuffer(t->dv->ctx, CL_MEM_READ_WRITE, bytes, NULL, &err); - if (err != CL_SUCCESS) { prepareFail(t, "clCreateBuffer cache", err); clReleaseCommandQueue(q); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; } - local = kernelMaxLocal(t->dv, p->kCacheFill, t->di, 256); - { - size_t g = ((nSeg + local - 1) / local) * local; - err = clSetKernelArg(p->kCacheFill, 0, sizeof(cl_mem), &p->cache); - if (err == CL_SUCCESS) err = clSetKernelArg(p->kCacheFill, 1, sizeof(cl_uint), &nSeg); - if (err == CL_SUCCESS) err = clEnqueueNDRangeKernel(q, p->kCacheFill, 1, NULL, &g, &local, 0, NULL, NULL); - if (err == CL_SUCCESS) err = clFinish(q); - } - if (err != CL_SUCCESS) { prepareFail(t, "cache fill", err); clReleaseCommandQueue(q); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; } - p->cacheMs = wallMs() - tb; - /* Dataset */ - tb = wallMs(); - nItems = t->words / 16u; - p->ds = clCreateBuffer(t->dv->ctx, CL_MEM_READ_WRITE, (size_t)t->words * 4u, NULL, &err); - if (err != CL_SUCCESS) { prepareFail(t, "clCreateBuffer dataset", err); clReleaseCommandQueue(q); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; } - local = kernelMaxLocal(t->dv, p->kBuild, t->di, 256); - { - size_t g = ((nItems + local - 1) / local) * local; - err = clSetKernelArg(p->kBuild, 0, sizeof(cl_mem), &p->ds); - if (err == CL_SUCCESS) err = clSetKernelArg(p->kBuild, 1, sizeof(cl_mem), &p->cache); - if (err == CL_SUCCESS) err = clSetKernelArg(p->kBuild, 2, sizeof(cl_uint), &nItems); - if (err == CL_SUCCESS) err = clEnqueueNDRangeKernel(q, p->kBuild, 1, NULL, &g, &local, 0, NULL, NULL); - if (err == CL_SUCCESS) err = clFinish(q); - } - if (err != CL_SUCCESS) { prepareFail(t, "dataset build", err); clReleaseCommandQueue(q); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; } - p->datasetMs = wallMs() - tb; + if (!pairBuffers(t->dv, t->di, q, p, t->words, t->cacheWords, t->segments, t->packDir, t->error, sizeof(t->error))) { clReleaseCommandQueue(q); releasePair(p); t->doneAt = wallMs(); t->done = 1; return; } clReleaseCommandQueue(q); t->result = p; t->doneAt = wallMs(); @@ -1029,7 +1111,7 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) { return 2; #else static const uint32_t KEYW[8] = IGNEUM_KEY_INIT; - const uint32_t words = 1u << IGNEUM_DATASET_LOG2; + const uint32_t words = gServeWords; const uint32_t mask = words - 1u; const uint32_t batch = 1u << (o->batchLog2 == 24 ? 22 : o->batchLog2); /* 2^22 nonces per dispatch by default */ size_t groupSize = 32 * (size_t)o->groupWarps; @@ -1045,24 +1127,36 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) { ServePair* old = NULL; /* the previous pair, released after the first job on the new one */ PrepareTask* task = NULL; /* the prepare in flight */ if (!dv->kHashBound) { printf("error 0 the kernel source has no igneum_hash_bound (build from the pack's kernel_bound.cl, or pass --kernel)\n"); fflush(stdout); return 2; } - if (!setupCache(dv, di)) { printf("error 0 cache check failed (device cache differs from the host cache or the pack's FNV)\n"); fflush(stdout); return 1; } cur = (ServePair*)calloc(1, sizeof(ServePair)); - memcpy(cur->sw, SEEDW, 32); memcpy(cur->kw, KEYW, 32); cur->kHashBound = dv->kHashBound; cur->kCacheFill = dv->kCacheFill; cur->kBuild = dv->kBuild; cur->prog = dv->prog; dv->kHashBound = dv->kCacheFill = dv->kBuild = NULL; dv->prog = NULL; /* owned by the pair now */ - cur->cache = gCache; gCache = NULL; - cur->ds = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)words * 4u, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer dataset"); - CL_CHECK(clSetKernelArg(cur->kBuild, 0, sizeof(cl_mem), &cur->ds)); - CL_CHECK(clSetKernelArg(cur->kBuild, 1, sizeof(cl_mem), &cur->cache)); - CL_CHECK(clSetKernelArg(cur->kBuild, 2, sizeof(cl_uint), &nItems)); - clReleaseEvent(launch1D(dv, cur->kBuild, nItems, kernelMaxLocal(dv, cur->kBuild, di, 256))); - CL_CHECK(clFinish(dv->q)); + if (gGeneric) { + /* --pack: the first pair is the directory's pack, built and self-tested exactly like a prepared one */ + char perr[512]; + double t0 = wallMs(); + memcpy(cur->sw, gPack.seedw, 32); memcpy(cur->kw, gPack.keyw, 32); + strncpy(cur->epochHex, gPack.epochHex, 64); cur->epochHex[64] = 0; + strncpy(cur->dayHex, gPack.dayHex, sizeof(cur->dayHex) - 1); + if (!pairBuffers(dv, di, dv->q, cur, words, gServeCacheWords, gServeSegments, o->packDir, perr, sizeof(perr))) { printf("error 0 pack %s: %s\n", o->packDir, perr); fflush(stdout); return 1; } + printf("info first pack %s: cache %.0f dataset %.0f check %.0f ms (%.0f ms in all); %s\n", o->packDir, cur->cacheMs, cur->datasetMs, cur->checkMs, wallMs() - t0, cur->check); + } else { + if (!setupCache(dv, di)) { printf("error 0 cache check failed (device cache differs from the host cache or the pack's FNV)\n"); fflush(stdout); return 1; } + memcpy(cur->sw, SEEDW, 32); memcpy(cur->kw, KEYW, 32); + cur->cache = gCache; gCache = NULL; + cur->ds = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)words * 4u, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer dataset"); + CL_CHECK(clSetKernelArg(cur->kBuild, 0, sizeof(cl_mem), &cur->ds)); + CL_CHECK(clSetKernelArg(cur->kBuild, 1, sizeof(cl_mem), &cur->cache)); + CL_CHECK(clSetKernelArg(cur->kBuild, 2, sizeof(cl_uint), &nItems)); + clReleaseEvent(launch1D(dv, cur->kBuild, nItems, kernelMaxLocal(dv, cur->kBuild, di, 256))); + CL_CHECK(clFinish(dv->q)); + } dOut = clCreateBuffer(dv->ctx, CL_MEM_READ_WRITE, (size_t)batch * sizeof(uint64_t), NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer out"); dInit = clCreateBuffer(dv->ctx, CL_MEM_READ_ONLY, 32, NULL, &err); CL_CHECK_ERR(err, "clCreateBuffer init words"); hOut = (uint64_t*)malloc((size_t)batch * sizeof(uint64_t)); strncpy(devName, di->name, 255); devName[255] = 0; for (k = 0; devName[k]; ++k) if (devName[k] == ' ') devName[k] = '_'; - printf("ready opencl %s platform %s pack %s dataset-log2 %d batch %u exchange %d prepare %d\n", devName, di->platformName, IGNEUM_SEED_STRING, IGNEUM_DATASET_LOG2, batch, dv->exchange, o->noPrepare ? 0 : 1); + printf("ready opencl %s platform %s pack %s dataset-log2 %d batch %u exchange %d prepare %d path %s\n", devName, di->platformName, + gGeneric ? gPack.seedString : IGNEUM_SEED_STRING, log2u32(words), batch, dv->exchange, o->noPrepare ? 0 : 1, gGeneric ? "prebuilt-generic" : "compiled-in"); fflush(stdout); while (fgets(line, sizeof(line), stdin)) { @@ -1091,7 +1185,7 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) { } if (prepared) releasePair(prepared); prepared = p; - printf("prepared %s %s %.1f build %.1f cache %.1f dataset %.1f resident 2 programs 2 datasets\n", p->epochHex, p->dayHex, task->doneAt - task->t0, p->buildMs, p->cacheMs, p->datasetMs); + printf("prepared %s %s %.1f build %.1f cache %.1f dataset %.1f check %.1f resident 2 programs 2 datasets; %s\n", p->epochHex, p->dayHex, task->doneAt - task->t0, p->buildMs, p->cacheMs, p->datasetMs, p->checkMs, p->check); } else { printf("prepare-failed %s %s %s\n", task->epochHex, task->dayHex, task->error); } @@ -1108,7 +1202,7 @@ static int runServe(Device* dv, const DeviceInfo* di, const Options* o) { if (task) { printf("prepare-failed %s %s a prepare is still running\n", f[1], f[2]); fflush(stdout); continue; } if (prepared && strcmp(prepared->epochHex, f[1]) == 0 && strcmp(prepared->dayHex, f[2]) == 0) { printf("prepared %s %s 0 (already resident)\n", f[1], f[2]); fflush(stdout); continue; } task = (PrepareTask*)calloc(1, sizeof(PrepareTask)); - task->dv = dv; task->di = di; task->words = words; task->t0 = wallMs(); + task->dv = dv; task->di = di; task->words = words; task->cacheWords = gServeCacheWords; task->segments = gServeSegments; task->t0 = wallMs(); strncpy(task->epochHex, f[1], 64); strncpy(task->dayHex, f[2], sizeof(task->dayHex) - 1); strncpy(task->packDir, f[3], sizeof(task->packDir) - 1); if (!startPrepareThread(task)) { printf("prepare-failed %s %s cannot start the prepare thread\n", f[1], f[2]); fflush(stdout); free(task); task = NULL; continue; } printf("info prepare started for epoch %.16s day %s from %s (builds in the background)\n", f[1], f[2], f[3]); fflush(stdout); @@ -1211,9 +1305,28 @@ int main(int argc, char** argv) { SizeResult results[5]; int cachePass = 1, overall, anyVec = 0; - printf("igneum-bench-cl pack \"%s\" (test harness: no pool, no network, no wallet)\n", IGNEUM_SEED_STRING); +#ifdef IGNEUM_CL_DYNAMIC + if (!ig_cl_load()) { printf("%s %s\n", o.serve ? "error 0" : "FAIL:", ig_cl_error); fflush(stdout); return 2; } +#endif + if (o.packDir) { + /* Generic serve mode: the pack directory replaces the compiled-in pack (4 October 2026) */ + static char boundPath[1200]; + char perr[512]; + size_t n = strlen(o.packDir); + if (!o.serve) { printf("FAIL: --pack goes with --serve (the bench runs the compiled-in pack)\n"); return 2; } + if (n > 1 && (o.packDir[n - 1] == '/' || o.packDir[n - 1] == '\\')) ((char*)o.packDir)[n - 1] = 0; + if (!pf_load(o.packDir, &gPack, perr, sizeof(perr))) { printf("error 0 pack %s: %s\n", o.packDir, perr); fflush(stdout); return 2; } + gGeneric = 1; + gServeWords = 1u << gPack.datasetLog2; +#if IGNEUM_DATASET_MODE == 1 + gServeCacheWords = 1u << gPack.cacheLog2Words; + gServeSegments = gPack.cacheSegments; +#endif + if (!o.kernelGiven) { snprintf(boundPath, sizeof(boundPath), "%s/kernel_bound.cl", o.packDir); o.kernelPath = boundPath; o.kernelGiven = 1; } + } + printf("igneum-bench-cl pack \"%s\" (test harness: no pool, no network, no wallet)%s\n", gGeneric ? gPack.seedString : IGNEUM_SEED_STRING, gGeneric ? " [--pack: generic serve mode]" : ""); nDev = enumerateDevices(&devs); - if (nDev == 0) { printf("FAIL: no OpenCL platform or device found (is an OpenCL driver / ICD installed?)\n"); return 2; } + if (nDev == 0) { printf("%s no OpenCL platform or device found (is the GPU driver installed? it provides OpenCL)\n", o.serve ? "error 0" : "FAIL:"); fflush(stdout); return 2; } if (o.device >= 0) { if (o.device >= nDev) { printf("FAIL: --device %d out of range (%d devices)\n", o.device, nDev); return 2; } chosen = o.device; diff --git a/proto-opencl/test-generic.sh b/proto-opencl/test-generic.sh new file mode 100755 index 00000000..bf1f2352 --- /dev/null +++ b/proto-opencl/test-generic.sh @@ -0,0 +1,20 @@ +#!/usr/bin/env bash +# Checks the generic serve mode of host.c (--serve --pack , the one-click AMD worker's mode) on this Mac through +# Apple's OpenCL runtime: the exe is built against one pack as a placeholder and then serves two others read at run +# time, with the same protocol script as the CUDA emulation (proto-cuda/nvrtc/emu/serve-check.sh): jobs on pack A, +# a background prepare of pack B with its self-test, the swap, the refusal after the swap, and a sample of found +# hashes against igneum-pow hash-bound. Apple OpenCL is a correctness check only, not an AMD number. +# Usage: ./test-generic.sh (needs proto-cuda/nvrtc/emu/test.sh to have run once: it writes pack-a and pack-b) +set -euo pipefail +cd "$(dirname "$0")" +ROOT="$(cd .. && pwd)" +EMUB="$ROOT/proto-cuda/nvrtc/emu/build" +[ -f "$EMUB/pack-a/seeds.txt" ] && [ -f "$EMUB/pack-b/seeds.txt" ] || { echo "run proto-cuda/nvrtc/emu/test.sh first (it writes $EMUB/pack-a and pack-b)" >&2; exit 1; } +PLACEHOLDER="../proto-cuda/packs/igneum-genesis-mh" # a different pack on purpose: the exe must not depend on it +CC="${CC:-cc}" +"$CC" -std=c99 -O2 -Wall -Wextra -Wno-deprecated-declarations -I "$PLACEHOLDER" -DIGNEUM_KERNEL_PATH="\"$PLACEHOLDER/kernel.cl\"" \ + -o igneum-bench-cl-generic-test host.c -framework OpenCL +echo "built ./igneum-bench-cl-generic-test (placeholder pack $PLACEHOLDER, Apple OpenCL)" +"$ROOT/proto-cuda/nvrtc/emu/serve-check.sh" "$EMUB/pack-a" "$EMUB/pack-b" "$PWD/generic-test.log" \ + ./igneum-bench-cl-generic-test --serve --pack "$EMUB/pack-a" --batch-log2 13 +grep '^ready\|^info first pack\|^prepared' generic-test.log