fix: honor reduction_offset and runtime chain_len in the NTLM/MD5 fast-path kernels - #136
Merged
Merged
Conversation
Manifest-driven conversion pipeline for the existing 16,388-part dell3 NetNTLMv1 set, partitioned across dell3 and dell2 by directory range, verified by byte-identical round-trip before each source .rtc is deleted. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…on sort markov_build_sorted() used qsort_r behind #ifdef __APPLE__ (BSD thunk-first signature) / #else (GNU arg-last signature) branches, gated by a top-level #define _GNU_SOURCE. mingw-w64's runtime provides neither variant -- qsort_r is a BSD/glibc extension, not part of the Windows CRT -- so a real Windows cross-build failed to link with "undefined reference to 'qsort_r'" in markov_build_sorted. This was never caught before because the Windows cross-build had only ever been dry-run checked (make -n), never linked. Replace both branches with a single sort_by_freq_desc() helper: an O(n^2) insertion sort descending by frequency, with the same ascending-by-index tie-break as the old comparators (return (int)ia - (int)ib). n is always <= 255 here (charset_len is validated to [1, 255] elsewhere in this file), so O(n^2) is trivial and this removes any dependency on a platform's reentrant qsort variant -- one code path for Linux, macOS, and Windows. Also drops the now-unused #define _GNU_SOURCE. Adds a regression test (group_a2, MS-05/MS-06 in tests/test_markov.c) that specifically exercises a frequency tie in both pos0_freq and a bigram_freq row, asserting the ascending-index tie-break is preserved. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Fixes a real Windows cross-build breakage: mingw-w64 provides no qsort_r variant, and markov.c's #ifdef __APPLE__/#else branches only covered BSD and glibc signatures. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
compute_load_thread_count() called sysconf(_SC_NPROCESSORS_ONLN) unconditionally to size the table-load thread pool. sysconf and _SC_NPROCESSORS_ONLN are POSIX-only and don't exist under mingw-w64, so the Windows cross-build fails with "'_SC_NPROCESSORS_ONLN' undeclared" plus an implicit-declaration warning for sysconf itself. This went unnoticed because the Windows build had only been checked with `make -n` dry runs, never actually linked. Fix: branch on _WIN32 and use GetSystemInfo()/SYSTEM_INFO.dwNumberOfProcessors (windows.h is already included at the top of the file for other Windows code paths) to get the CPU count on Windows, keeping the sysconf path unchanged for Linux/macOS. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Fixes crackalack_lookup.c's compute_load_thread_count() using sysconf, a POSIX-only API unavailable on mingw-w64 -- a second pre-existing Windows cross-build breakage found alongside the qsort_r one. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
cpu_rt_functions.c's fill_plaintext_space_mask() calls mask_keyspace() (mask_parse.c), which wasn't in enumerate_chain's OBJS list. Other backends' linkers apparently tolerated or never exercised the missing symbol; mingw-w64's linker caught it as an undefined reference during a real Windows cross-build. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…lib list Hardcoded -lgcrypt -lm omitted -lbcrypt, which BCryptGenRandom (called from misc.c's get_random(), used by every binary) needs on Windows. Every other target already links via $(LIBS); this was the one exception, caught by a real Windows cross-build. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Fixes the last Windows cross-build blocker: crackalack_plan.exe's link recipe hardcoded -lgcrypt -lm instead of using $(LIBS), omitting -lbcrypt (needed for BCryptGenRandom in misc.c's get_random()). Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The optimized NTLM8 reduce (hash_to_index_ntlm8 / hash_char_to_index_ntlm8) dropped reduction_offset, so it only produced correct chains at table_index 0. Add the parameter, mirroring hash_to_index_netntlmv1_7, and wire it through every consumer: crackalack_ntlm8, precompute_ntlm8, precompute_ntlm8_batch, and false_alarm_check_ntlm8, across all three GPU backends (CL/Metal/CUDA). crackalack_ntlm8 and false_alarm_check_ntlm8 claim the existing g_reduction_offset placeholder at the position the host already binds it (retyping two unsigned long* placeholders to unsigned int* to match). precompute_ntlm8 and precompute_ntlm8_batch instead take g_table_index and derive reduction_offset via TABLE_INDEX_TO_REDUCTION_OFFSET(), matching precompute.cl. precompute_ntlm8_batch's position-3 argument was an unused charset_len placeholder in the kernel body; repurposed it as table_index and updated crackalack_lookup.c's batch dispatch to bind table_index there for the plain-NTLM8 batch path (added a separate branch so the Markov batch kernel, which does read charset_len at that position, is unaffected). At table_index 0 this is a no-op: the three GEN_TESTS sha256 fixtures in crackalack_tests.py still pass unchanged. Precompute test #1 (table_index 16) and lookup test #2 now pass; lookup tests #1/#3/#4/#5 still fail on the known leftover-precalc-file bug (separate, fixed in a later task). Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
batch_precompute_all_hashes() wrote the raw per-hash `hash_output` array straight through to the rcracki.precalc.N disk cache: positions_per_hash (chain_len) entries in ascending-position order. The single-hash precompute_ntlm8 path instead produces chain_len-1 entries in descending-position order (its own "GPU0: 100 94 88 ..." collation). The batch cache file was therefore one entry too long and reversed relative to the per-hash path and the crackalack_tests.py oracle, causing lookup test #3 (which exercises the batch path via 4 hashes) to fail with a precalc hash mismatch even though the underlying reduction_offset math was correct. Reorder/trim the array handed to save_precompute_cache using the same reversal collate_batched_precompute_endpoints() already applies for the in-memory ppi construction, sized to positions_per_hash - 1 and zero-filled (not the ppi's sentinel) to match the raw single-hash cache format exactly. The ppi construction itself is untouched. Verified the 1-hash and 4-hash precompute paths now produce byte-identical rcracki.precalc.0 files for the same hash/table, matching the crackalack_tests.py oracle sha256. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…ed kernels Same defect as NTLM8 (0ddba87): the optimized reduce for each of these paths dropped reduction_offset, so it only produced correct chains at table_index 0. Added the parameter to hash_to_index_*/hash_char_to_index_* for all four paths and wired it through every crackalack_*/precompute_*/false_alarm_check_*/test_* consumer across CL/Metal/CUDA, sourcing it from the same argument slots the host already binds (g_reduction_offset at position 5 for crackalack_*/ false_alarm_check_*, g_table_index at position 7 for precompute_*, matching precompute.cl's TABLE_INDEX_TO_REDUCTION_OFFSET() usage). Several false_alarm_check_* variants had that slot mistyped as unsigned long*/ulong*; retyped to unsigned int* to match the host's actual bind. Relaxed is_ntlm9() by deleting its reduction_offset==0 clause now that the NTLM9 reduce honors the offset, letting 9-char tables at nonzero table_index use the ~12x faster optimized kernel. Updated tests/test_misc.c's IN9-04, which asserted the old (buggy) rejection behavior the brief specifically instructs removing. At table_index 0 this is a no-op: all three GEN_TESTS sha256 fixtures still pass unchanged, both precompute tests pass (including table_index 1024 on the NTLM9 fast path), lookup test #2 passes, and crackalack_unit_tests is back to ALL UNIT TESTS PASS after the test_misc.c update. Lookup #1/#3/#4/#5 still fail on the known leftover-precalc-file bug (Task 3, out of scope here). Left precompute_ntlm10's pre-existing chain_len/plaintext_len_max argument mismatch untouched (unrelated latent bug, no test coverage exists for NTLM10 generation/lookup) -- flagged in the task report for separate follow-up. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Restore the index_filename field to ppi that commit e36d539 dropped, and re-wire it through both precompute paths (cache-hit and cache-save, batch and single-hash) so save_cracked_hash() can unlink the rcracki.precalc.N / .N.index pair once a hash is cracked. Guards against a NULL/empty filename (a hash can crack with no cache entry) and keeps unlink failures non-fatal, matching the original behavior. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
CUDA/precompute.cu and CUDA/false_alarm_check.cu did not exist, so any CUDA lookup falling off the NTLM8/9/10/MD5/NetNTLMv1-7 fast paths (e.g. a loweralpha charset table) died at runtime trying to open a nonexistent kernel file. Add faithful mechanical translations of CL/precompute.cl and CL/false_alarm_check.cl, following the same __global-drop, unsigned long -> unsigned long long, and absolute thread-id (blockIdx.x * blockDim.x + threadIdx.x) conventions already used by CUDA/precompute_ntlm8.cu and CUDA/false_alarm_check_ntlm8.cu. CUDA build/execution is unverified on this machine (no CUDA GPU/toolchain available); macOS build + full test suite pass, confirming no regression to the working backend. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The CUDA backend launches ceil(gws / 256) * 256 threads while OpenCL and
Metal dispatch exactly gws work items, so up to 255 surplus threads run
with a global id >= gws. Several kernels store at their raw global id
before (or without) range-checking:
precompute*.cu: if (target_chain_len < 1) { g_output[gid] = 0; return; }
crackalack*.cu: g_indices[gid] = generate_rainbow_chain(...);
Both output buffers were sized at exactly gws, so those surplus stores
landed past the end of the allocation. In the lookup precompute path the
next two cuMemAllocs are plaintext_space_up_to_index (arg 13) and
plaintext_space_total (arg 14), so the surplus threads zeroed them --
making hash_to_index() a modulo-by-zero, which NVIDIA evaluates to
all-ones. That is why CUDA/precompute.cu returned UINT64_MAX for every
entry with no CUDA error reported, and why the failure was insensitive to
-gws (gws is clamped to output_len first, which is never a multiple of
256 for the test tables' chain_len of 100).
The fast-path precompute kernels had the identical out-of-bounds store all
along; it was latent only because their args 13/14 are unused.
false_alarm_check*.cu is not affected -- its index_pos goes negative for
surplus threads and it returns before any store.
Fix: introduce GPU_LAUNCH_GRANULARITY / GPU_GWS_PAD() in gpu_backend.h
(256 under USE_CUDA, 1 elsewhere, so OpenCL and Metal are bit-for-bit
unchanged), tie cuda_setup.c's block size to the same macro so the two
cannot drift, and size the affected host allocations, device buffers and
transfers with GPU_GWS_PAD(). The consumed element counts are left at
gws, so results and generated tables stay identical across backends.
Verified on Metal: make clean && make macos -j8 && python3
crackalack_tests.py -> ALL TESTS PASS.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…rt buffer precompute_ntlm10 read chain_len from arg 6 (host binds plaintext_len_max there) while ignoring the real chain_len at arg 8, collapsing every NTLM10 lookup's precompute to chain_len==10. precompute_md5_8/md5_9 hardcoded 422000/803000 instead of reading the correctly-typed but unused arg 8. Fixed across CL/Metal/CUDA. Audit of sibling kernels found CL/crackalack_md5_8.cl and CL/crackalack_md5_9.cl (generation, OpenCL only) had the same defect: hardcoded loop bounds ignoring both chain_len and pos_start, while their Metal/CUDA counterparts were already correct. Brought OpenCL in line. Also pad crackalack_sort.c's GPU buffer to GPU_GWS_PAD(n_padded) so CUDA's launch-grid overshoot (surplus threads with gid >= n_padded that pass sort.cu's `if (l <= i) return` guard) can't write past the end of the sort buffer for tables under 256 chains. The sort's k/j loop bounds and gws remain n_padded, unchanged. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…v1.5.8 Apply the explicit signed-cast pattern already used in precompute_ntlm8 to precompute_md5_8/md5_9 across CUDA, Metal, and CL, so target_chain_len is computed in signed arithmetic instead of relying on an implementation-defined unsigned-to-signed conversion for the tail-thread guard. Document the reduction_offset/chain_len compatibility break in CLAUDE.md so pre-fix tables at nonzero table_index (or non-default MD5 chain lengths) are flagged for regeneration. Bump VERSION to v1.5.8.
…uivalence tests Gates is_ntlm8/9/10, is_md5_8/9, and is_netntlmv1_7 in misc.c behind a cached, env-driven kill switch so a single binary can be run through both the optimized and generic kernels for the same table parameters. Adds an oracle-free test group to crackalack_tests.py that runs each optimized path twice (fast path, then RCRT_DISABLE_FAST_PATH=1) against non-default table_index/chain_len parameters and asserts byte-identical rcracki.precalc.0 output, confirming via the kernel-selection log line that each run actually took the path it was supposed to. NTLM9 is implemented but left disabled by default (measured ~14.5 min per precompute run at its required chain_len). Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Task 7: closes the last known coverage gap on this branch. Commit 6d6aca7 fixed batch_precompute_all_hashes() writing its rcracki.precalc.N cache file reversed and one entry too long relative to precompute_hash(); nothing in the suite exercised the batch dispatch (num_hashes >= 2) against the per-hash path directly at a nonzero table_index. Runs the same hash both ways (alone, then as the first of four) against ntlm_ascii-32-95#8-8_32_100x1024_0.rt and asserts the two cache file pairs are byte-identical, locating the batch run's entry for the shared hash via its .index sidecar rather than assuming index 0. Verified non-vacuous via the "Batched precompute (" kernel log line, and verified the test actually fails by temporarily reverting the 6d6aca7 fix (see task-7-report.md). Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…v gating fast_path_disabled() used two non-atomic file-scope statics readable from the multithreaded lookup pipeline, letting one thread observe checked==1 before disabled's write was visible. Switched to pthread_once, which mingw-w64's winpthreads implements natively for the Windows cross-build. Investigated is_markov_ntlm8/9/10 per review finding 2: all three already unconditionally return 0 (commit e383156, unrelated to this kill switch), so none is a live fast path to gate. No code change made there; reasoning documented in the report. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Under the CUDA backend gpu_buffer is CUdeviceptr, an integer type, so the `gpu_buffer foo = NULL` idiom used throughout the host code triggered 305 -Wint-conversion warnings. Add GPU_BUFFER_NULL to gpu_backend.h (0 on CUDA, NULL elsewhere) and use it at every gpu_buffer assignment and comparison. This also silences the lone pointer/integer comparison warning in crackalack_gen.c's first-invocation guard. Also drop a dead `j` in crackalack_lookup.c's main(). No behavior change: 0 was already the sentinel the CUDA macros used. CUDA now builds warning-free (307 -> 0); Metal builds warning-free; the Windows/OpenCL cross-build's 6 remaining warnings are pre-existing and unrelated. Unit tests, CPU tests, and crackalack_tests.py all pass on CUDA (RTX 3080 Ti), including the nonzero-table-index fast-path equivalence tests. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
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
1 participant
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.
Summary
The hand-optimized NTLM8 / NTLM9 / NTLM10 / MD5-8 / MD5-9 kernels omitted
reduction_offsetfrom their hash-to-index reduce, so they only produced correct chains attable_index 0— but the selectors (is_ntlm8,is_ntlm10,is_md5_8,is_md5_9) accepted any table index. That silently corrupted both generation and lookup at nonzero indices. Several adjacent bugs surfaced while fixing it.Any NTLM8 / NTLM10 / MD5-8 / MD5-9 table generated by a pre-fix build at
table_index != 0, or any MD5 table at a non-default chain length, is invalid and must be regenerated. Tables attable_index 0with default chain lengths are unaffected.CL/crackalack_md5_8.clandCL/crackalack_md5_9.clalso restarted each pass at position 0 instead of honoringpos_start(Metal and CUDA were already correct), so OpenCL-generated MD5 tables are affected too.What's fixed
reduction_offsetthreaded through all five optimized kernel families across CUDA, Metal, and OpenCL.chain_lenhonored in the ntlm10/md5_8/md5_9 precompute paths (previously hardcoded 422000/803000; ntlm10 read it from the wrong argument slot), plus a signed cast for the MD5 precompute chain-len guard.is_ntlm9()no longer requiresreduction_offset == 0, so 9-char tables at nonzero indices now use the optimized kernel.precompute/false_alarm_checkkernels.ceil(gws/256)*256threads, so global-id-indexed buffers (including the sort buffer) needed padding via a newGPU_GWS_PAD().gpu_bufferisCUdeviceptr(an integer) under CUDA, so the= NULLidiom emitted 305-Wint-conversionwarnings that buried real ones. AddedGPU_BUFFER_NULLand swept all 306 sites; CUDA now builds warning-free.Tests
RCRT_DISABLE_FAST_PATHkill switch (with apthread_oncerace fixed) plus fast-path/generic equivalence tests that assert byte-identical tables at nonzero table indices: ntlm8@16, ntlm10@128, md5_8@32, md5_9@64.Verification
All three backends built and runtime-tested against this exact tree:
Windows cross-build links; its 6 remaining warnings are pre-existing and unrelated.
🤖 Generated with Claude Code