Repository navigation
Restore the OpenCL batch-factoring build (WITH_OPENCL), fix the OpenCL host path, add kernel tests - #100
Merged
Conversation
The OpenCL cofactorisation backend added in 5607dd3 lost its Makefile wiring (-DHAVE_OCL_BATCH_FACTOR and the two source files) when Makefile.new became Makefile in abced9f, and section 4d only probed for OpenCL when BATCH_CUDA was set. - Add WITH_OPENCL=1: OpenCL backend without a CUDA toolkit. - BATCH_CUDA=1 with no CUDA toolkit falls back to OpenCL again. - WITH_OPENCL is mutually exclusive with BATCH_CUDA / CUDA_POLY (both backends define the same gpu_* / device_* symbols). - ocl_xface.c: use the same guard as ocl_xface.h; define CL_DEVICE_PREFERRED_WORK_GROUP_SIZE_MULTIPLE when the SDK headers predate OpenCL 3.0 and fall back to the AMD wavefront query. - gpu_cofactorization_cl.c: stop defining HAVE_LASIEVE_MPQS itself (mpqs_factor is not linked into yafu) and use mpz_get_64/mpz_set_64 instead of the lasieve mpz_*_ull helpers. Builds without the new flags are unchanged. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
uecm_stage1 and its 96-bit counterpart call uprac()/uprac96() about 100 times each in straight-line code. They were translated from CUDA's plain __device__ functions to "static inline", and AMD's OpenCL compiler inlines every call, so gbl_ecm alone grew from ~3k lines of IR to ~278k and the AMDGPU backend needed more than 27 minutes (never finished in our test) to build opencl_tinyecm.cl on a Radeon 780M with ROCm 10.1. Marking the two functions noinline brings the IR down to ~98k lines and the full clBuildProgram time to about 42 seconds on the same machine. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
Runs gbl_ecm, gbl_ecm96 and gbl_pm196 from factor/opencl_tinyecm.cl on a GPU, built exactly as gpu_cofactorization_cl.c builds them, and checks every reported factor with GMP. Inputs are balanced 61/62-bit semiprimes (gbl_ecm), 82..87-bit semiprimes with a 30..34-bit factor (gbl_ecm96) and 96-bit composites whose p-1 is built from distinct primes < 500 (gbl_pm196). A wrong (non-dividing) factor is a hard failure; the miss rate must stay under 2%, like the CPU ECM tests. The module is only compiled with WITH_OPENCL=1 (make WITH_OPENCL=1 test) and passes with a note when no OpenCL GPU is present. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
gpu_cofactorization_cl.c had never been executed by yafu since its port from the CUDA code, and diverged from it in three ways that show up as soon as SIQS with TLP (-forceTLP) uses it: - gpu_cofactorization() lacked the single-side cases (first_side -2..-5, do_2lp/do_3lp/do_pm1); take the CUDA logic. - do_gpu_cofactorization() swapped the r-side/a-side codes for "only one side has unfactored residues" and did not set lpb_2lp/lpb_3lp; take the CUDA logic. - Residues were packed with exactly lp_*_num_words words per relation but read back with a fixed stride of 2 or 3, so a side that mixes 2- and 3-word cofactors shifted everything after the first short one and the 96-bit kernels were fed zero moduli (SIGFPE in mpz_tdiv_r). Pack with a fixed per-side stride and zero-pad. - do_gpu_ecm64/ecm_96b/pm1_96b enqueued zero-sized buffer writes when handed an empty list; OpenCL rejects those with CL_INVALID_VALUE (CUDA's memcpy does not), so return early. With these, siqs() of a 75-digit semiprime with -forceTLP runs through the OpenCL cofactorisation on a Radeon 780M and returns the same factors as the CPU-only build. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
NFS and SIQS give every batch thread its own context. On ROCm, with several threads creating, running and releasing contexts on the same GPU at once, yafu died with "Memory access fault by GPU ... address (nil)": the kernel log showed an instruction-fetch page fault at address 0 followed by a full GPU reset (which also disturbs the display on an iGPU). Single batches, and batches that ran one after the other, were fine. gpu_ctx_init now takes a process-wide lock and gpu_ctx_free releases it, so GPU batches run one at a time; the CPU-side work of the other threads is unaffected. The exact failing mechanism inside ROCm is not established; this removes the concurrency that triggered it. With this, an 85-digit GNFS job using a-side 3LP (-nfs_batch_3lp, 8 threads, 34 96-bit and 68 64-bit GPU batches) runs to completion on a Radeon 780M and returns the correct factors, with no GPU fault logged. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
teruteru128
force-pushed
the
feature/opencl-standalone
branch
from
October 9, 2026 08:01
ca4a668 to
b9d2861
Compare
Contributor
Author
|
Thank you for merge! |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Follow-up to #99 — thanks for the quick reply and for being open to this.
Summary
The OpenCL cofactorisation backend from 5607dd3 had lost its Makefile wiring in abced9f, and its host code had never been run from yafu (it had drifted from the CUDA code). This PR makes it build without CUDA via a new
WITH_OPENCL=1, fixes what broke once it actually ran, and adds tests. Five commits, each self-contained:-DHAVE_OCL_BATCH_FACTORand the two sources;BATCH_CUDA=1with no CUDA toolkit falls back to OpenCL again;WITH_OPENCLis mutually exclusive withBATCH_CUDA/CUDA_POLY(samegpu_*/device_*symbols). Also fixesCL_DEVICE_PREFERRED_WORK_GROUP_SIZE_MULTIPLE(OpenCL 3.0 only; missing from ROCm's 2.2 headers) and theocl_xface.cguard, and stopsgpu_cofactorization_cl.cfrom definingHAVE_LASIEVE_MPQSitself (mpqs_factoris not linked into yafu; it now usestinysiqslike the CUDA file) and from using the lasievempz_*_ullhelpers.static inline; AMD's OpenCL compiler inlines everything, which made one kernel ~278k lines of IR and the firstclBuildProgramnever finished in my test.noinlinebrings the IR to ~98k lines for the whole program and the build finishes quickly (later runs come from the driver cache).test/layer2/test_ocl_ecm.c) — checksgbl_ecm,gbl_ecm96,gbl_pm196against known-structure inputs with GMP. Only built withWITH_OPENCL=1; passes with a note when no OpenCL GPU is present.gpu_cofactorization()anddo_gpu_cofactorization()lacked the single-side cases / had the r/a codes swapped (CUDA logic taken over); residues were packed withlp_*_num_wordswords but read with a fixed stride, so a side mixing 2- and 3-word cofactors fed zero moduli to the 96-bit kernels (SIGFPE); zero-sized buffer writes areCL_INVALID_VALUEin OpenCL.gpu_ctx_initnow takes a process-wide lock released bygpu_ctx_free. I haven't established the exact failing mechanism inside ROCm; this removes the concurrency that triggered it.What is not changed
make -nforyafu,msieve,siqs,ecmandtestproduces exactly the same command lists as upstream/master (so the existing CI, which buildsmake yafuwithout GPU flags, is unaffected).build.vc22is untouched (the projects already list the OpenCL files).HAVE_OCL_POLYMakefile wiring is left for later, as discussed.Testing
Ubuntu, GCC 15, ROCm 10.1 OpenCL runtime, Radeon 780M (gfx1103), on the head of this branch:
make WITH_OPENCL=1 test && ./yafu_test ocl_ecm: passes (3 kernels); the full./yafu_testrun withWITH_OPENCL=1(40 tests) also passes.make WITH_OPENCL=1 yafubuilds.siqs(N) -forceTLPon a 75- and an 85-digit semiprime goes through the OpenCL cofactorisation and returns the correct factors.-nfs_batch_3lp, 8 threads, a customnfs_paramstable chosen only to exercise the GPU path) ran to completion through ~100 GPU batches with the correct factors and no GPU fault logged.Not tested: NVIDIA/CUDA hardware, Windows (MinGW/MSVC), natural-size TLP/NFS jobs. I have no reliable performance numbers (the machine was under other load), so this PR makes no speed claims.
Notes for review
hipccfor gfx1103 and gives bit-identical results to the OpenCL kernels on the same inputs; CUDA preprocesses to exactly the same text). I'll open it as its own PR later.Refs #99.
Disclosure: parts of this work were developed with an AI coding assistant (Claude Code; see the
Co-Authored-Bytrailers). Everything described above was built and run on my machine.🤖 Generated with Claude Code