diff --git a/CHANGELOG.md b/CHANGELOG.md index a499bcc..c018a94 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -1,3 +1,9 @@ +### v0.4.1-dev; 2026-08-22 +``` +complete rewrite of codebase +add OpenCL GPU acceleration for yescrypt / gost-yescrypt with CPU fallback +add ETA status for -w wordlists +``` ### v0.3.1; 2026-08-08 ``` add gost-yescrypt support diff --git a/README.md b/README.md index 64600f7..7daff09 100644 --- a/README.md +++ b/README.md @@ -5,69 +5,94 @@ [![GitHub release](https://img.shields.io/github/release/cyclone-github/yescrypt_crack.svg)](https://github.com/cyclone-github/yescrypt_crack/releases) [![Go Reference](https://pkg.go.dev/badge/github.com/cyclone-github/yescrypt_crack.svg)](https://pkg.go.dev/github.com/cyclone-github/yescrypt_crack) +### Note: v0.4.0-dev adds beta support for OpenCL GPU acceleration. + ``` ./yescrypt_crack.bin -h hash.txt -w wordlist.txt - -------------------------------------------------- | Cyclone's Yescrypt Cracker | | https://github.com/cyclone-github/yescrypt_crack | -------------------------------------------------- Hash file: hash.txt -Total Hashes: 2 -CPU Threads: 16 +Total Hashes: 26 +CPU Threads: 56 Wordlist: wordlist.txt -2026/08/08 11:04:49 Working... -$y$j9T$z7lNWyBfW4ZruGHCsFzDz/$Sz1GtrDDnsf0KfUE8mQHNJqGyG32TDWC287DdU97dz.:cyclone123 -$gy$j9T$zHbAdM1G1TomuQ6P6lRwc.$kTKZeBwUiSlCcxHfPQuzczMUxeNgp0LaK2LbsAK8KO/:cyclone123 -2026/08/08 11:04:49 Finished -2026/08/08 11:04:49 Cracked: 2/2 23.51 h/s 00h:00m:00s +Backend: auto +GPU selection: all +GPU batch: 1280 +2026/08/20 11:54:25 Counting wordlist lines... +2026/08/20 11:54:25 Tuning GPU... +2026/08/20 11:54:26 OpenCL GPU 0: NVIDIA GeForce RTX 4090, 23.5 GiB VRAM, batch cap 1280, self-test passed +2026/08/20 11:54:26 OpenCL GPU 1: NVIDIA GeForce RTX 4090, 23.5 GiB VRAM, batch cap 1280, self-test passed +2026/08/20 11:54:26 Working... +2026/08/20 11:55:26 Cracked: 0/26 7338.51 H/s 01m:00s/-- +2026/08/20 11:56:26 Cracked: 0/26 7391.90 H/s 02m:00s/-- +2026/08/20 11:57:26 Cracked: 0/26 7416.85 H/s 03m/1d:17h:06m +2026/08/20 11:58:26 Cracked: 0/26 7434.64 H/s 04m/1d:20h:21m +2026/08/20 11:59:26 Cracked: 0/26 7441.04 H/s 05m/1d:13h:26m +```` +### Install this dev branch: ``` - -### Install from latest: +go install github.com/cyclone-github/yescrypt_crack@yescrypt_crack_gpu +``` +### Install latest main branch: ``` go install github.com/cyclone-github/yescrypt_crack@main ``` ### Info: -I wrote this tool since yescrypt has become the default /etc/shadow hash for many popular linux distros such as Debian, Ubuntu, RHEL, Fedora, Arch, etc, and due to the very limited hash cracking tools that supported yescrypt. - -Supports both yescrypt and gost-yescrypt. - -Since `yescrypt_crack` is written in pure Go, it easily compiles and runs on just about any OS and architecture such as Intel/ARM, Linux, Windows, Mac. +I originally wrote this tool in 2025 since yescrypt had become the default /etc/shadow hash for many popular Linux distros such as Debian, Ubuntu, RHEL, Fedora, Arch, etc, and due to the very limited hash cracking tools that supported yescrypt at that time. -It is worth noting that JtR can be faster than `yescrypt_crack`, so YMMV. +`yescrypt_crack` supports OpenCL GPU acceleration for yescrypt and gost-yescrypt. GPU mode is used by default when a supported OpenCL GPU is available, and CPU mode is avaialble by using flag `-cpu`. All R&D has been performed on Debian 12/13 Linux. Cross platform support for Windows has been implemented, but is not guaranteed. For best results, Linux is recommended. Mac is not supported from v0.4.0 onward, but is supported on v0.3.1. ### Example hash:plain: ``` -$y$j9T$z7lNWyBfW4ZruGHCsFzDz/$Sz1GtrDDnsf0KfUE8mQHNJqGyG32TDWC287DdU97dz.:cyclone123 -$gy$j9T$Mq.X1lTWvTD74i0lyu6lC0$11n1S4OQtGaonRa4j4y9FsY.DnZUDwBFGKyzH2J/ry.:cyclone123 +$y$j9T$ss392e/1r/sS364AbKyZU1$Z6lrn5PE2YqIDbgeH590BnPC5fP8BFrkRgOhvo8WKVC:yescrypt_crack +$gy$j9T$pFEYYB6sbVC37XBX.uMqK/$xY/CaU3ESVzvfR9YErup5kUn2FiHSOoqgAx1zUA99u1:yescrypt_crack ``` ### Supported options: ``` -w {wordlist} (omit -w to read from stdin) --h {yescrypt_hash} +-h {yescrypt_hash_file} -o {output} (omit -o to write to stdout) --t {cpu threads} +-t {cpu threads} (selects CPU mode) -s {print status every nth sec} +-gpu [all|0,1|list] (default: all GPUs; omit value to use all GPUs) +-cpu (force CPU mode) +-b {gpu batch size} (optional; 0 = auto) -version (version info) -help (usage instructions) +``` + +### Examples: +``` +./yescrypt_crack.bin -h hashes.txt -w wordlist.txt -o found.txt -./yescrypt_crack.bin -h {yescrypt_hash} -w {wordlist} -o {output} -t {cpu threads} -s {print status every nth sec} +./yescrypt_crack.bin -h hashes.txt -w wordlist.txt -gpu -./yescrypt_crack.bin -h yescrypt.txt -w wordlist.txt -o cracked.txt -t 16 -s 10 +./yescrypt_crack.bin -h hashes.txt -w wordlist.txt -gpu 0,1 -cat wordlist | ./yescrypt_crack.bin -h yescrypt.txt +./yescrypt_crack.bin -h hashes.txt -w wordlist.txt -gpu 0,1 -b 1280 -s 10 -./yescrypt_crack.bin -h yescrypt.txt -w wordlist.txt -o output.txt +./yescrypt_crack.bin -h hashes.txt -w wordlist.txt -cpu -t 16 -s 10 + +cat wordlist | ./yescrypt_crack.bin -h hashes.txt + +./yescrypt_crack.bin -gpu list ``` +For backwards compatibility, explicitly specifying `-cpu` selects CPU mode. + ### Credits: -* `yescrypt_crack` tool was written by cyclone in pure Go -* `yescrypt_crack` uses Solar Designer's yescrypt-go implementation: https://github.com/openwall/yescrypt-go -* The yescrypt algo was written by Solar Designer: https://www.openwall.com/yescrypt/ +* `yescrypt_crack` was written by cyclone in Go +* The yescrypt algorithm was designed by Solar Designer: https://www.openwall.com/yescrypt/ +* The CPU yescrypt implementation in `yescrypt_cpu.go` is adapted from `openwall/yescrypt-go`: https://github.com/openwall/yescrypt-go +* The GPU implementation in `yescrypt.cl` is adapted from hashcat's yescrypt OpenCL implementation: https://github.com/hashcat/hashcat/blob/master/OpenCL/inc_hash_yescrypt.cl +* Streebog-256 is provided by `github.com/tarantool/go-gostcrypto/streebog`. +* See `THIRD_PARTY_NOTICES.md` for upstream copyright notices and license terms. ### Changelog: - https://github.com/cyclone-github/yescrypt_crack/blob/main/CHANGELOG.md @@ -77,9 +102,9 @@ cat wordlist | ./yescrypt_crack.bin -h yescrypt.txt - This assumes you have Go and Git installed - `git clone https://github.com/cyclone-github/yescrypt_crack.git` # clone repo - `cd yescrypt_crack` # enter project directory - - `go mod tidy` # download dependencies - - `go build -ldflags="-s -w" .` # compile binary in current directory - - `go install -ldflags="-s -w" .` # compile binary and install to $GOPATH - - `./yescrypt_crack -h {hash file} -w {wordlist file} -t {CPU threads to use (optional)}` # run yescrypt_crack + - `go mod tidy` # download dependencies + - `go build -ldflags="-s -w" .` # compile binary in current directory + - `go install -ldflags="-s -w" .` # compile binary and install to $GOPATH + - `./yescrypt_crack -h {hash file} -w {wordlist file}` # run yescrypt_crack - Compile from source code how-to: - https://github.com/cyclone-github/scripts/blob/main/intro_to_go.txt diff --git a/THIRD_PARTY_NOTICES.md b/THIRD_PARTY_NOTICES.md new file mode 100644 index 0000000..baeb0cb --- /dev/null +++ b/THIRD_PARTY_NOTICES.md @@ -0,0 +1,69 @@ +# Third-Party Notices + +`yescrypt_crack` is licensed under GPLv2. The project also contains or depends on work from the projects below. + +## yescrypt / yescrypt-go + +yescrypt was designed by Solar Designer. Portions of the CPU implementation in `yescrypt_cpu.go` are adapted from `openwall/yescrypt-go`. + +Project: https://github.com/openwall/yescrypt-go + +License: BSD-3-Clause + +Copyright (c) 2009-2020 The Go Authors. All rights reserved. +Copyright (c) 2024 Solar Designer. All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions are met: + +* Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer. +* Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution. +* Neither the name of Google Inc. nor the names of its contributors may be used to endorse or promote products derived from this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +## hashcat + +The GPU implementation in yescrypt.cl is adapted from hashcat's yescrypt OpenCL implementation. + +Relevant implementation: https://github.com/hashcat/hashcat/blob/master/OpenCL/inc_hash_yescrypt.cl + +License: MIT + +Copyright (c) 2015-2025 Jens Steube + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in all +copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +SOFTWARE. + +## go-gostcrypto + +The gost-yescrypt finalization path imports `github.com/tarantool/go-gostcrypto/streebog`. + +Project: https://github.com/tarantool/go-gostcrypto + +License: BSD-2-Clause + +Copyright (c) 2026, Tarantool AUTHORS + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions are met: + +1. Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer. +2. Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. diff --git a/check_hex.go b/check_hex.go index 343c9f9..b0922ac 100644 --- a/check_hex.go +++ b/check_hex.go @@ -43,7 +43,7 @@ func checkForHexBytes(line []byte) ([]byte, []byte, int) { } decodedBytes = make([]byte, hex.DecodedLen(len(cleaned))) - _, err = hex.Decode(decodedBytes, cleaned) + n, err = hex.Decode(decodedBytes, cleaned) if err != nil { return line, line, 1 } diff --git a/go.mod b/go.mod index 8237a14..3d5eb12 100644 --- a/go.mod +++ b/go.mod @@ -2,9 +2,4 @@ module github.com/cyclone-github/yescrypt_crack go 1.26.2 -require ( - github.com/openwall/yescrypt-go v1.0.0 - github.com/tarantool/go-gostcrypto v0.1.0 -) - -require golang.org/x/crypto v0.55.0 // indirect +require github.com/tarantool/go-gostcrypto v0.1.0 diff --git a/go.sum b/go.sum index 94c35f8..222541d 100644 --- a/go.sum +++ b/go.sum @@ -1,6 +1,2 @@ -github.com/openwall/yescrypt-go v1.0.0 h1:jsGk48zkFvtUjGVOhYPGh+CS595JmTRcKnpggK2AON4= -github.com/openwall/yescrypt-go v1.0.0/go.mod h1:e6CWtFizUEOUttaOjeVMiv1lJaJie3mfOtLJ9CCD6sA= github.com/tarantool/go-gostcrypto v0.1.0 h1:sGjZil1nfXzxvBU4TOIZErTXcvCuCLH/pZnwfXr9uaY= github.com/tarantool/go-gostcrypto v0.1.0/go.mod h1:oieRliQHLNWjZ5qrpIamIFNHJ3j7dEvU4evBr6HeEFc= -golang.org/x/crypto v0.55.0 h1:+KWHjbgOaAQ66dh/YlkZKHlz9ZUlq61AFirAR9ntP8M= -golang.org/x/crypto v0.55.0/go.mod h1:uq0V9dE/fzQuJtbnL+2EhWOE63vo164FY8xqEnV9xis= diff --git a/gost.go b/gost.go index 30cc7a6..3feb0f2 100644 --- a/gost.go +++ b/gost.go @@ -1,13 +1,5 @@ package main -import ( - "bytes" - "crypto/hmac" - - "github.com/openwall/yescrypt-go" - "github.com/tarantool/go-gostcrypto/streebog" -) - var cryptBase64DecodeTable = [...]byte{ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 64, 64, 64, 64, 64, 64, 64, @@ -18,6 +10,8 @@ var cryptBase64DecodeTable = [...]byte{ 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, } +const cryptBase64Alphabet = "./0123456789ABCDEFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuvwxyz" + func cryptBase64Value(c byte) int { if c >= '.' && c <= 'z' { return int(cryptBase64DecodeTable[c-'.']) @@ -29,11 +23,9 @@ func cryptBase64Value(c byte) int { // by yescrypt and gost-yescrypt. It is not RFC 4648 base64. func decodeCryptBase64(src []byte) []byte { dst := make([]byte, 0, len(src)*3/4) - for i := 0; i < len(src); { var value uint32 var bits uint32 - for ; bits < 24 && i < len(src); bits += 6 { c := cryptBase64Value(src[i]) if c > 63 { @@ -42,73 +34,39 @@ func decodeCryptBase64(src []byte) []byte { i++ value |= uint32(c) << bits } - if bits < 12 { return nil } - for ; bits >= 8; bits -= 8 { dst = append(dst, byte(value)) value >>= 8 } - if value != 0 { return nil } } - return dst } -func crackGostYescrypt(password, fullHash []byte) bool { - if !bytes.HasPrefix(fullHash, []byte("$gy$")) { - return false - } - - // libxcrypt computes the normal yescrypt result first using the same - // parameters and salt. Convert "$gy$..." to "$y$..." for yescrypt-go. - yescryptSetting := make([]byte, 0, len(fullHash)-1) - yescryptSetting = append(yescryptSetting, '$', 'y', '$') - yescryptSetting = append(yescryptSetting, fullHash[4:]...) - - yescryptHash, err := yescrypt.Hash(password, yescryptSetting) - if err != nil { - return false - } - - yescryptDigestPos := bytes.LastIndexByte(yescryptHash, '$') - if yescryptDigestPos < 0 { - return false - } - yescryptDigest := decodeCryptBase64(yescryptHash[yescryptDigestPos+1:]) - if len(yescryptDigest) != 32 { - return false - } - - // the final '$' separates the gost-yescrypt setting from its digest - // libxcrypt intentionally excludes this '$' from the inner HMAC message - gostDigestPos := bytes.LastIndexByte(fullHash, '$') - if gostDigestPos < 0 { - return false - } - expectedDigest := decodeCryptBase64(fullHash[gostDigestPos+1:]) - if len(expectedDigest) != 32 { - return false +func encodeCryptBase64(src []byte) []byte { + dst := make([]byte, 0, (len(src)*8+5)/6) + for i := 0; i < len(src); { + var value uint32 + bits := 0 + for bits < 24 && i < len(src) { + value |= uint32(src[i]) << bits + i++ + bits += 8 + } + for bits > 0 { + dst = append(dst, cryptBase64Alphabet[value&0x3f]) + value >>= 6 + bits -= 6 + } } + return dst +} - // libxcrypt construction: - // hk = Streebog-256(password) - // interm = HMAC-Streebog-256(hk, gost setting without trailing '$') - // result = HMAC-Streebog-256(interm, raw yescrypt 256-bit result) - hk := streebog.Sum256(password) - - mac := hmac.New(streebog.New256, hk[:]) - _, _ = mac.Write(fullHash[:gostDigestPos]) - interm := mac.Sum(nil) - - mac = hmac.New(streebog.New256, interm) - _, _ = mac.Write(yescryptDigest) - generatedDigest := mac.Sum(nil) - - return hmac.Equal(expectedDigest, generatedDigest) +func crackGostYescrypt(password, fullHash []byte) bool { + return crackHash(password, fullHash) } diff --git a/gpu_opencl.c b/gpu_opencl.c new file mode 100644 index 0000000..2c01b4d --- /dev/null +++ b/gpu_opencl.c @@ -0,0 +1,549 @@ +#include "gpu_opencl.h" + +#include +#include +#include +#include +#include +#include +#include + +/* Minimal OpenCL 1.2 ABI declarations. We intentionally do not include + * CL/cl.h so the binary can be built without OpenCL development headers. */ +typedef int32_t cl_int; +typedef uint32_t cl_uint; +typedef uint64_t cl_ulong; +typedef cl_ulong cl_bitfield; +typedef cl_bitfield cl_device_type; +typedef cl_bitfield cl_mem_flags; +typedef cl_bitfield cl_command_queue_properties; +typedef uint32_t cl_bool; +typedef intptr_t cl_context_properties; +typedef struct _cl_platform_id *cl_platform_id; +typedef struct _cl_device_id *cl_device_id; +typedef struct _cl_context *cl_context; +typedef struct _cl_command_queue *cl_command_queue; +typedef struct _cl_mem *cl_mem; +typedef struct _cl_program *cl_program; +typedef struct _cl_kernel *cl_kernel; +typedef struct _cl_event *cl_event; + +#define CL_SUCCESS 0 +#define CL_DEVICE_NOT_FOUND -1 +#define CL_PLATFORM_NOT_FOUND_KHR -1001 +#define CL_DEVICE_TYPE_GPU (1ULL << 2) +#define CL_DEVICE_NAME 0x102B +#define CL_DEVICE_GLOBAL_MEM_SIZE 0x101F +#define CL_DEVICE_MAX_MEM_ALLOC_SIZE 0x1010 +#define CL_PROGRAM_BUILD_LOG 0x1183 +#define CL_MEM_READ_WRITE (1ULL << 0) +#define CL_MEM_WRITE_ONLY (1ULL << 1) +#define CL_MEM_READ_ONLY (1ULL << 2) +#define CL_FALSE 0 +#define CL_TRUE 1 + +#define MAX_GPU_DEVICES 32 +#define MAX_PW 256ULL +#define MAX_R 32ULL +#define B_STRIDE (128ULL * MAX_R) +#define XY_STRIDE_WORDS (32ULL * MAX_R) +#define S_WORDS 1536ULL +#define ONE_MIB (1024ULL * 1024ULL) + +static void *ocl_lib; + +#define DECL(name, ret, args) static ret (*p_##name) args +DECL(clGetPlatformIDs, cl_int, (cl_uint, cl_platform_id *, cl_uint *)); +DECL(clGetDeviceIDs, cl_int, (cl_platform_id, cl_device_type, cl_uint, cl_device_id *, cl_uint *)); +DECL(clGetDeviceInfo, cl_int, (cl_device_id, cl_uint, size_t, void *, size_t *)); +DECL(clCreateContext, cl_context, (const cl_context_properties *, cl_uint, const cl_device_id *, void (*)(const char *, const void *, size_t, void *), void *, cl_int *)); +DECL(clCreateCommandQueue, cl_command_queue, (cl_context, cl_device_id, cl_command_queue_properties, cl_int *)); +DECL(clCreateProgramWithSource, cl_program, (cl_context, cl_uint, const char **, const size_t *, cl_int *)); +DECL(clBuildProgram, cl_int, (cl_program, cl_uint, const cl_device_id *, const char *, void (*)(cl_program, void *), void *)); +DECL(clGetProgramBuildInfo, cl_int, (cl_program, cl_device_id, cl_uint, size_t, void *, size_t *)); +DECL(clCreateKernel, cl_kernel, (cl_program, const char *, cl_int *)); +DECL(clCreateBuffer, cl_mem, (cl_context, cl_mem_flags, size_t, void *, cl_int *)); +DECL(clSetKernelArg, cl_int, (cl_kernel, cl_uint, size_t, const void *)); +DECL(clEnqueueWriteBuffer, cl_int, (cl_command_queue, cl_mem, cl_bool, size_t, size_t, const void *, cl_uint, const cl_event *, cl_event *)); +DECL(clEnqueueNDRangeKernel, cl_int, (cl_command_queue, cl_kernel, cl_uint, const size_t *, const size_t *, const size_t *, cl_uint, const cl_event *, cl_event *)); +DECL(clEnqueueReadBuffer, cl_int, (cl_command_queue, cl_mem, cl_bool, size_t, size_t, void *, cl_uint, const cl_event *, cl_event *)); +DECL(clFinish, cl_int, (cl_command_queue)); +DECL(clReleaseMemObject, cl_int, (cl_mem)); +DECL(clReleaseKernel, cl_int, (cl_kernel)); +DECL(clReleaseProgram, cl_int, (cl_program)); +DECL(clReleaseCommandQueue, cl_int, (cl_command_queue)); +DECL(clReleaseContext, cl_int, (cl_context)); +#undef DECL + +static int seterr(char *err, size_t errlen, const char *fmt, ...) { + if (err && errlen) { + va_list ap; + va_start(ap, fmt); + vsnprintf(err, errlen, fmt, ap); + va_end(ap); + } + return -1; +} + +static int load_opencl(char *err, size_t errlen) { + if (ocl_lib) return 0; + const char *libs[] = {"libOpenCL.so.1", "libOpenCL.so", NULL}; + for (int i = 0; libs[i]; i++) { + ocl_lib = dlopen(libs[i], RTLD_NOW | RTLD_LOCAL); + if (ocl_lib) break; + } + if (!ocl_lib) return seterr(err, errlen, "OpenCL loader not found: %s", dlerror()); + +#define LOAD(name) do { *(void **)(&p_##name) = dlsym(ocl_lib, #name); if (!p_##name) return seterr(err, errlen, "OpenCL symbol %s missing", #name); } while (0) + LOAD(clGetPlatformIDs); LOAD(clGetDeviceIDs); LOAD(clGetDeviceInfo); + LOAD(clCreateContext); LOAD(clCreateCommandQueue); LOAD(clCreateProgramWithSource); + LOAD(clBuildProgram); LOAD(clGetProgramBuildInfo); LOAD(clCreateKernel); + LOAD(clCreateBuffer); LOAD(clSetKernelArg); LOAD(clEnqueueWriteBuffer); + LOAD(clEnqueueNDRangeKernel); LOAD(clEnqueueReadBuffer); LOAD(clFinish); + LOAD(clReleaseMemObject); LOAD(clReleaseKernel); LOAD(clReleaseProgram); + LOAD(clReleaseCommandQueue); LOAD(clReleaseContext); +#undef LOAD + return 0; +} + +static int enumerate_gpus(cl_device_id *out, int maxout, char *err, size_t errlen) { + if (load_opencl(err, errlen) != 0) return -1; + cl_uint np = 0; + cl_int rc = p_clGetPlatformIDs(0, NULL, &np); + if (rc == CL_PLATFORM_NOT_FOUND_KHR || np == 0) return 0; + if (rc != CL_SUCCESS) return seterr(err, errlen, "clGetPlatformIDs failed: %d", rc); + cl_platform_id *plats = calloc(np, sizeof(*plats)); + if (!plats) return seterr(err, errlen, "out of memory enumerating OpenCL platforms"); + rc = p_clGetPlatformIDs(np, plats, NULL); + if (rc != CL_SUCCESS) { free(plats); return seterr(err, errlen, "clGetPlatformIDs(list) failed: %d", rc); } + int count = 0; + for (cl_uint p = 0; p < np && count < maxout; p++) { + cl_uint nd = 0; + rc = p_clGetDeviceIDs(plats[p], CL_DEVICE_TYPE_GPU, 0, NULL, &nd); + if (rc == CL_DEVICE_NOT_FOUND || nd == 0) continue; + if (rc != CL_SUCCESS) continue; + cl_device_id *devs = calloc(nd, sizeof(*devs)); + if (!devs) continue; + if (p_clGetDeviceIDs(plats[p], CL_DEVICE_TYPE_GPU, nd, devs, NULL) == CL_SUCCESS) { + for (cl_uint d = 0; d < nd && count < maxout; d++) out[count++] = devs[d]; + } + free(devs); + } + free(plats); + return count; +} + +int ycl_opencl_device_count(char *err, size_t errlen) { + cl_device_id devs[MAX_GPU_DEVICES]; + return enumerate_gpus(devs, MAX_GPU_DEVICES, err, errlen); +} + +int ycl_opencl_device_info(int index, char *name, size_t namelen, + uint64_t *global_mem, uint64_t *max_alloc, + char *err, size_t errlen) { + cl_device_id devs[MAX_GPU_DEVICES]; + int n = enumerate_gpus(devs, MAX_GPU_DEVICES, err, errlen); + if (n < 0) return -1; + if (index < 0 || index >= n) return seterr(err, errlen, "OpenCL GPU index %d out of range (found %d)", index, n); + cl_device_id d = devs[index]; + if (name && namelen) { + if (p_clGetDeviceInfo(d, CL_DEVICE_NAME, namelen, name, NULL) != CL_SUCCESS) snprintf(name, namelen, "OpenCL GPU %d", index); + name[namelen - 1] = 0; + } + cl_ulong gm = 0, ma = 0; + if (p_clGetDeviceInfo(d, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(gm), &gm, NULL) != CL_SUCCESS) return seterr(err, errlen, "cannot query GPU global memory"); + if (p_clGetDeviceInfo(d, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(ma), &ma, NULL) != CL_SUCCESS) return seterr(err, errlen, "cannot query GPU max allocation"); + if (global_mem) *global_mem = gm; + if (max_alloc) *max_alloc = ma; + return 0; +} + +struct ycl_gpu_ctx { + int index; + cl_device_id dev; + cl_context ctx; + cl_command_queue queue; + + char *source; + size_t source_len; + cl_program program; + cl_kernel init_kernel; + cl_kernel loop_kernel; + cl_kernel final_kernel; + + cl_ulong global_mem; + cl_ulong max_alloc; + uint32_t N, r, capacity; + atomic_int cancel_requested; + + cl_mem passwords; + cl_mem pw_lens; + cl_mem salt; + cl_mem P; + cl_mem scratch; + cl_mem S; + cl_mem yctx; + cl_mem V[4]; + cl_mem out; +}; + +static void release_buffers(ycl_gpu_ctx *g) { + cl_mem *all[] = { + &g->passwords, &g->pw_lens, &g->salt, + &g->P, &g->scratch, &g->S, &g->yctx, + &g->V[0], &g->V[1], &g->V[2], &g->V[3], &g->out + }; + for (size_t i = 0; i < sizeof(all) / sizeof(all[0]); i++) { + if (*all[i]) { + p_clReleaseMemObject(*all[i]); + *all[i] = NULL; + } + } + g->capacity = 0; +} + +static void release_program(ycl_gpu_ctx *g) { + if (g->init_kernel) { p_clReleaseKernel(g->init_kernel); g->init_kernel = NULL; } + if (g->loop_kernel) { p_clReleaseKernel(g->loop_kernel); g->loop_kernel = NULL; } + if (g->final_kernel) { p_clReleaseKernel(g->final_kernel); g->final_kernel = NULL; } + if (g->program) { p_clReleaseProgram(g->program); g->program = NULL; } + g->N = 0; + g->r = 0; +} + +void ycl_gpu_cancel(ycl_gpu_ctx *g) { + if (g) atomic_store_explicit(&g->cancel_requested, 1, memory_order_relaxed); +} + +void ycl_gpu_reset_cancel(ycl_gpu_ctx *g) { + if (g) atomic_store_explicit(&g->cancel_requested, 0, memory_order_relaxed); +} + +void ycl_gpu_destroy(ycl_gpu_ctx *g) { + if (!g) return; + release_buffers(g); + release_program(g); + if (g->queue) p_clReleaseCommandQueue(g->queue); + if (g->ctx) p_clReleaseContext(g->ctx); + free(g->source); + free(g); +} + +static int build_specialized_program(ycl_gpu_ctx *g, uint32_t N, uint32_t r, + char *err, size_t errlen) { + if (g->program && g->N == N && g->r == r) return 0; + + release_program(g); + + cl_int rc; + const char *sources[1] = {g->source}; + size_t lengths[1] = {g->source_len}; + g->program = p_clCreateProgramWithSource(g->ctx, 1, sources, lengths, &rc); + if (!g->program || rc != CL_SUCCESS) { + release_program(g); + return seterr(err, errlen, "clCreateProgramWithSource failed: %d", rc); + } + + char options[256]; + snprintf(options, sizeof(options), "-cl-std=CL1.2 -DYC_N=%u -DYC_R=%u", N, r); + rc = p_clBuildProgram(g->program, 1, &g->dev, options, NULL, NULL); + if (rc != CL_SUCCESS) { + size_t logsz = 0; + p_clGetProgramBuildInfo(g->program, g->dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &logsz); + char *log = calloc(logsz + 1, 1); + if (log) p_clGetProgramBuildInfo(g->program, g->dev, CL_PROGRAM_BUILD_LOG, logsz, log, NULL); + seterr(err, errlen, "OpenCL kernel build failed (%d) [%s]: %s", rc, options, log ? log : "no build log"); + free(log); + release_program(g); + return -1; + } + + g->init_kernel = p_clCreateKernel(g->program, "yescrypt_init", &rc); + if (!g->init_kernel || rc != CL_SUCCESS) { + seterr(err, errlen, "clCreateKernel(yescrypt_init) failed: %d", rc); + release_program(g); + return -1; + } + g->loop_kernel = p_clCreateKernel(g->program, "yescrypt_loop", &rc); + if (!g->loop_kernel || rc != CL_SUCCESS) { + seterr(err, errlen, "clCreateKernel(yescrypt_loop) failed: %d", rc); + release_program(g); + return -1; + } + g->final_kernel = p_clCreateKernel(g->program, "yescrypt_final", &rc); + if (!g->final_kernel || rc != CL_SUCCESS) { + seterr(err, errlen, "clCreateKernel(yescrypt_final) failed: %d", rc); + release_program(g); + return -1; + } + + g->N = N; + g->r = r; + return 0; +} + +ycl_gpu_ctx *ycl_gpu_create(int index, const char *source, size_t source_len, + char *err, size_t errlen) { + cl_device_id devs[MAX_GPU_DEVICES]; + int n = enumerate_gpus(devs, MAX_GPU_DEVICES, err, errlen); + if (n < 0) return NULL; + if (index < 0 || index >= n) { + seterr(err, errlen, "OpenCL GPU index %d out of range (found %d)", index, n); + return NULL; + } + if (!source || source_len == 0) { + seterr(err, errlen, "empty OpenCL source"); + return NULL; + } + + ycl_gpu_ctx *g = calloc(1, sizeof(*g)); + if (!g) { + seterr(err, errlen, "out of memory creating GPU context"); + return NULL; + } + atomic_init(&g->cancel_requested, 0); + g->index = index; + g->dev = devs[index]; + g->source = malloc(source_len + 1); + if (!g->source) { + seterr(err, errlen, "out of memory copying OpenCL source"); + ycl_gpu_destroy(g); + return NULL; + } + memcpy(g->source, source, source_len); + g->source[source_len] = 0; + g->source_len = source_len; + + p_clGetDeviceInfo(g->dev, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(g->global_mem), &g->global_mem, NULL); + p_clGetDeviceInfo(g->dev, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(g->max_alloc), &g->max_alloc, NULL); + + cl_int rc; + g->ctx = p_clCreateContext(NULL, 1, &g->dev, NULL, NULL, &rc); + if (!g->ctx || rc != CL_SUCCESS) { + seterr(err, errlen, "clCreateContext failed: %d", rc); + ycl_gpu_destroy(g); + return NULL; + } + g->queue = p_clCreateCommandQueue(g->ctx, g->dev, 0, &rc); + if (!g->queue || rc != CL_SUCCESS) { + seterr(err, errlen, "clCreateCommandQueue failed: %d", rc); + ycl_gpu_destroy(g); + return NULL; + } + return g; +} + +static cl_mem mkbuf(ycl_gpu_ctx *g, cl_mem_flags flags, size_t size, + const char *what, char *err, size_t errlen) { + cl_int rc; + if (size == 0) size = 8; + cl_mem m = p_clCreateBuffer(g->ctx, flags, size, NULL, &rc); + if (!m || rc != CL_SUCCESS) { + seterr(err, errlen, "clCreateBuffer(%s, %.2f MiB) failed: %d", + what, (double)size / (1024.0 * 1024.0), rc); + return NULL; + } + return m; +} + +int ycl_gpu_configure(ycl_gpu_ctx *g, uint32_t N, uint32_t r, + uint32_t batch_hint, uint32_t *capacity, + char *err, size_t errlen) { + if (!g || !N || !r || r > MAX_R || (N & (N - 1U))) + return seterr(err, errlen, "invalid yescrypt GPU parameters N=%u r=%u", N, r); + + /* Rebuild only when the yescrypt geometry changes. N/r become compile-time + * constants so the OpenCL compiler can specialize the cooperative inner loop. */ + const int geometry_changed = (g->N != N || g->r != r || !g->program); + if (!geometry_changed && g->capacity && (!batch_hint || batch_hint == g->capacity)) { + if (capacity) *capacity = g->capacity; + return 0; + } + + release_buffers(g); + if (build_specialized_program(g, N, r, err, errlen) != 0) return -1; + + const uint64_t state_bytes = 128ULL * (uint64_t)r; + const uint64_t vbytes = state_bytes * (uint64_t)N; + const uint64_t scratch_bytes = state_bytes < 256ULL ? 256ULL : state_bytes; + const uint64_t ctx_bytes = 48ULL; + const uint64_t workspace = MAX_PW + sizeof(cl_uint) + 64ULL + + state_bytes + scratch_bytes + + S_WORDS * sizeof(cl_ulong) + + ctx_bytes + 32ULL; + + uint64_t reserve = g->global_mem / 10ULL; + if (reserve < 512ULL * ONE_MIB) reserve = 512ULL * ONE_MIB; + if (reserve >= g->global_mem) reserve = g->global_mem / 5ULL; + const uint64_t usable = g->global_mem - reserve; + uint64_t cap_mem = usable / (vbytes + workspace); + + const uint64_t alloc_limit = (g->max_alloc * 95ULL) / 100ULL; + const uint64_t slots_per_segment = vbytes ? alloc_limit / vbytes : 0; + uint64_t cap_alloc = slots_per_segment * 4ULL; + uint64_t cap = cap_mem; + if (cap > cap_alloc) cap = cap_alloc; + if (cap > 4096ULL) cap = 4096ULL; + if (batch_hint) { + if (cap > batch_hint) cap = batch_hint; + } else { + cap = (cap * 95ULL) / 100ULL; + if (cap >= 32ULL) cap = (cap / 32ULL) * 32ULL; + } + if (cap == 0) + return seterr(err, errlen, + "GPU cannot allocate one yescrypt V region: need %.2f MiB/candidate, max allocation %.2f MiB", + (double)vbytes / ONE_MIB, (double)g->max_alloc / ONE_MIB); + if (cap >= 4ULL) cap = (cap / 4ULL) * 4ULL; + if (cap == 0) cap = 1; + + g->capacity = (uint32_t)cap; + const uint64_t c = cap; + + g->passwords = mkbuf(g, CL_MEM_READ_ONLY, (size_t)(c * MAX_PW), "passwords", err, errlen); if (!g->passwords) goto fail; + g->pw_lens = mkbuf(g, CL_MEM_READ_ONLY, (size_t)(c * sizeof(cl_uint)), "pw_lens", err, errlen); if (!g->pw_lens) goto fail; + g->salt = mkbuf(g, CL_MEM_READ_ONLY, 64, "salt", err, errlen); if (!g->salt) goto fail; + g->P = mkbuf(g, CL_MEM_READ_WRITE, (size_t)(c * state_bytes), "P", err, errlen); if (!g->P) goto fail; + g->scratch = mkbuf(g, CL_MEM_READ_WRITE, (size_t)(c * scratch_bytes), "scratch", err, errlen); if (!g->scratch) goto fail; + g->S = mkbuf(g, CL_MEM_READ_WRITE, (size_t)(c * S_WORDS * sizeof(cl_ulong)), "S", err, errlen); if (!g->S) goto fail; + g->yctx = mkbuf(g, CL_MEM_READ_WRITE, (size_t)(c * ctx_bytes), "context", err, errlen); if (!g->yctx) goto fail; + + for (int seg = 0; seg < 4; seg++) { + const uint64_t segcount = (c + (uint64_t)(3 - seg)) / 4ULL; + char label[16]; + snprintf(label, sizeof(label), "V%d", seg); + g->V[seg] = mkbuf(g, CL_MEM_READ_WRITE, (size_t)(segcount * vbytes), label, err, errlen); + if (!g->V[seg]) goto fail; + } + + g->out = mkbuf(g, CL_MEM_WRITE_ONLY, (size_t)(c * 32ULL), "output", err, errlen); if (!g->out) goto fail; + if (capacity) *capacity = g->capacity; + return 0; + +fail: + release_buffers(g); + return -1; +} + +#define CLCHK(call, what) do { \ + cl_int _rc = (call); \ + if (_rc != CL_SUCCESS) return seterr(err, errlen, "%s failed: %d", what, _rc); \ +} while (0) + +#define SET_MEM(kernel, argno, memobj, label) do { \ + cl_mem _m = (memobj); \ + CLCHK(p_clSetKernelArg((kernel), (argno)++, sizeof(_m), &_m), label); \ +} while (0) + +#define SET_U32(kernel, argno, value, label) do { \ + cl_uint _v = (cl_uint)(value); \ + CLCHK(p_clSetKernelArg((kernel), (argno)++, sizeof(_v), &_v), label); \ +} while (0) + +int ycl_gpu_hash(ycl_gpu_ctx *g, + const unsigned char *passwords, + const uint32_t *pw_lens, + uint32_t count, + const unsigned char *salt, + uint32_t salt_len, + unsigned char *out, + char *err, size_t errlen) { + if (!g || !g->capacity || !g->program || count == 0 || count > g->capacity) + return seterr(err, errlen, "bad GPU batch count %u (capacity %u)", count, g ? g->capacity : 0); + if (salt_len > 64) return seterr(err, errlen, "salt too long for GPU (%u)", salt_len); + if (atomic_load_explicit(&g->cancel_requested, memory_order_relaxed)) + return YCL_GPU_CANCELLED; + + /* Queue transfers without host-side synchronization. The command queue is + * in-order, and the final blocking read keeps the Go-owned host buffers + * alive until all three GPU stages have consumed them. */ + CLCHK(p_clEnqueueWriteBuffer(g->queue, g->passwords, CL_FALSE, 0, + (size_t)count * MAX_PW, passwords, 0, NULL, NULL), + "write passwords"); + CLCHK(p_clEnqueueWriteBuffer(g->queue, g->pw_lens, CL_FALSE, 0, + (size_t)count * sizeof(cl_uint), pw_lens, 0, NULL, NULL), + "write password lengths"); + unsigned char saltbuf[64] = {0}; + if (salt_len) memcpy(saltbuf, salt, salt_len); + CLCHK(p_clEnqueueWriteBuffer(g->queue, g->salt, CL_FALSE, 0, 64, + saltbuf, 0, NULL, NULL), + "write salt"); + + /* Init: one work-item per candidate. */ + cl_uint a = 0; + SET_MEM(g->init_kernel, a, g->passwords, "set init passwords"); + SET_MEM(g->init_kernel, a, g->pw_lens, "set init lengths"); + SET_MEM(g->init_kernel, a, g->salt, "set init salt"); + SET_U32(g->init_kernel, a, salt_len, "set init salt_len"); + SET_U32(g->init_kernel, a, count, "set init count"); + SET_MEM(g->init_kernel, a, g->P, "set init P"); + SET_MEM(g->init_kernel, a, g->scratch, "set init scratch"); + SET_MEM(g->init_kernel, a, g->S, "set init S"); + SET_MEM(g->init_kernel, a, g->yctx, "set init context"); + SET_MEM(g->init_kernel, a, g->V[0], "set init V0"); + SET_MEM(g->init_kernel, a, g->V[1], "set init V1"); + SET_MEM(g->init_kernel, a, g->V[2], "set init V2"); + SET_MEM(g->init_kernel, a, g->V[3], "set init V3"); + size_t global_init = count; + CLCHK(p_clEnqueueNDRangeKernel(g->queue, g->init_kernel, 1, NULL, + &global_init, NULL, 0, NULL, NULL), + "launch yescrypt init"); + + /* Split the long SMix into resumable chunks. 2048 iterations gives Ctrl+C + * a chance to stop between launches without materially changing GPU throughput. */ + const uint64_t nloop = ((((uint64_t)g->N + 2ULL) / 3ULL) + 1ULL) & ~1ULL; + uint64_t remaining = (uint64_t)g->N + nloop; + const size_t local_loop = 32; + const size_t global_loop = (size_t)count * local_loop; + while (remaining) { + if (atomic_load_explicit(&g->cancel_requested, memory_order_relaxed)) { + (void)p_clFinish(g->queue); + return YCL_GPU_CANCELLED; + } + + const uint32_t chunk = (remaining > 2048ULL) ? 2048U : (uint32_t)remaining; + a = 0; + SET_U32(g->loop_kernel, a, count, "set loop count"); + SET_U32(g->loop_kernel, a, chunk, "set loop chunk"); + SET_MEM(g->loop_kernel, a, g->P, "set loop P"); + SET_MEM(g->loop_kernel, a, g->S, "set loop S"); + SET_MEM(g->loop_kernel, a, g->yctx, "set loop context"); + SET_MEM(g->loop_kernel, a, g->V[0], "set loop V0"); + SET_MEM(g->loop_kernel, a, g->V[1], "set loop V1"); + SET_MEM(g->loop_kernel, a, g->V[2], "set loop V2"); + SET_MEM(g->loop_kernel, a, g->V[3], "set loop V3"); + CLCHK(p_clEnqueueNDRangeKernel(g->queue, g->loop_kernel, 1, NULL, + &global_loop, &local_loop, 0, NULL, NULL), + "launch yescrypt cooperative loop"); + CLCHK(p_clFinish(g->queue), "wait for yescrypt cooperative loop"); + + if (atomic_load_explicit(&g->cancel_requested, memory_order_relaxed)) + return YCL_GPU_CANCELLED; + remaining -= chunk; + } + + if (atomic_load_explicit(&g->cancel_requested, memory_order_relaxed)) + return YCL_GPU_CANCELLED; + + /* Final: one work-item per candidate. */ + a = 0; + SET_U32(g->final_kernel, a, count, "set final count"); + SET_MEM(g->final_kernel, a, g->P, "set final P"); + SET_MEM(g->final_kernel, a, g->scratch, "set final scratch"); + SET_MEM(g->final_kernel, a, g->yctx, "set final context"); + SET_MEM(g->final_kernel, a, g->out, "set final output"); + size_t global_final = count; + CLCHK(p_clEnqueueNDRangeKernel(g->queue, g->final_kernel, 1, NULL, + &global_final, NULL, 0, NULL, NULL), + "launch yescrypt final"); + + CLCHK(p_clEnqueueReadBuffer(g->queue, g->out, CL_TRUE, 0, + (size_t)count * 32ULL, out, 0, NULL, NULL), + "read yescrypt output"); + CLCHK(p_clFinish(g->queue), "clFinish"); + return 0; +} + +#undef SET_U32 +#undef SET_MEM +#undef CLCHK diff --git a/gpu_opencl.go b/gpu_opencl.go new file mode 100644 index 0000000..ace39bf --- /dev/null +++ b/gpu_opencl.go @@ -0,0 +1,196 @@ +//go:build linux && cgo + +package main + +/* +#cgo LDFLAGS: -ldl +#include +#include "gpu_opencl.h" +*/ +import "C" + +import ( + "bytes" + _ "embed" + "errors" + "fmt" + "runtime" + "unsafe" +) + +//go:embed kernels/yescrypt.cl +var yescryptOpenCLSource string + +const gpuMaxPasswordLen = 256 + +var errGPUCancelled = errors.New("GPU cancelled") + +type GPUDeviceInfo struct { + Index int + Name string + GlobalMem uint64 + MaxAlloc uint64 +} + +func cErrorString(buf *[8192]C.char) string { + if buf[0] == 0 { + return "unknown OpenCL error" + } + return C.GoString(&buf[0]) +} + +func ListOpenCLGPUs() ([]GPUDeviceInfo, error) { + var errbuf [8192]C.char + n := int(C.ycl_opencl_device_count(&errbuf[0], C.size_t(len(errbuf)))) + if n < 0 { + return nil, errors.New(cErrorString(&errbuf)) + } + out := make([]GPUDeviceInfo, 0, n) + for i := 0; i < n; i++ { + var name [512]C.char + var gm, ma C.uint64_t + for j := range errbuf { + errbuf[j] = 0 + } + if C.ycl_opencl_device_info(C.int(i), &name[0], C.size_t(len(name)), &gm, &ma, &errbuf[0], C.size_t(len(errbuf))) != 0 { + return nil, errors.New(cErrorString(&errbuf)) + } + out = append(out, GPUDeviceInfo{Index: i, Name: C.GoString(&name[0]), GlobalMem: uint64(gm), MaxAlloc: uint64(ma)}) + } + return out, nil +} + +type OpenCLGPU struct { + ctx *C.ycl_gpu_ctx + Info GPUDeviceInfo + N uint32 + R uint32 + Capacity int +} + +func NewOpenCLGPU(info GPUDeviceInfo) (*OpenCLGPU, error) { + src := []byte(yescryptOpenCLSource) + if len(src) == 0 { + return nil, errors.New("embedded OpenCL kernel is empty") + } + p := C.CBytes(src) + defer C.free(p) + var errbuf [8192]C.char + ctx := C.ycl_gpu_create(C.int(info.Index), (*C.char)(p), C.size_t(len(src)), &errbuf[0], C.size_t(len(errbuf))) + if ctx == nil { + return nil, errors.New(cErrorString(&errbuf)) + } + return &OpenCLGPU{ctx: ctx, Info: info}, nil +} + +func (g *OpenCLGPU) Close() { + if g != nil && g.ctx != nil { + C.ycl_gpu_destroy(g.ctx) + g.ctx = nil + } +} + +func (g *OpenCLGPU) Cancel() { + if g != nil && g.ctx != nil { + C.ycl_gpu_cancel(g.ctx) + } +} + +func (g *OpenCLGPU) ResetCancel() { + if g != nil && g.ctx != nil { + C.ycl_gpu_reset_cancel(g.ctx) + } +} + +func (g *OpenCLGPU) SelfTest() error { + const vector = "$y$j9T$Gg3nKTjUa.Hrq3ZPL9S9J0$kJBFNrOZC2sCjYIOV3G4/NTOOFWUULJDthOnxtiTez9" + h, err := parseYescryptHash(vector) + if err != nil { + return fmt.Errorf("internal GPU self-test vector parse failed: %w", err) + } + if _, err := g.Configure(h.N, h.R, 1); err != nil { + return fmt.Errorf("GPU self-test configure failed: %w", err) + } + out, err := g.HashBatch([][]byte{[]byte("password")}, h.Salt) + if err != nil { + return fmt.Errorf("GPU self-test execution failed: %w", err) + } + if len(out) != 1 || !bytes.Equal(out[0][:], h.Expected[:]) { + if len(out) == 1 { + return fmt.Errorf("GPU self-test digest mismatch: got %x want %x", out[0], h.Expected) + } + return fmt.Errorf("GPU self-test returned %d digests", len(out)) + } + return nil +} + +func (g *OpenCLGPU) Configure(N, r uint32, batchHint int) (int, error) { + if g == nil || g.ctx == nil { + return 0, errors.New("GPU context is closed") + } + if batchHint < 0 { + batchHint = 0 + } + var cap C.uint32_t + var errbuf [8192]C.char + if C.ycl_gpu_configure(g.ctx, C.uint32_t(N), C.uint32_t(r), C.uint32_t(batchHint), &cap, &errbuf[0], C.size_t(len(errbuf))) != 0 { + return 0, errors.New(cErrorString(&errbuf)) + } + g.N, g.R, g.Capacity = N, r, int(cap) + return g.Capacity, nil +} + +func (g *OpenCLGPU) HashBatch(passwords [][]byte, salt []byte) ([][32]byte, error) { + if g == nil || g.ctx == nil || g.Capacity == 0 { + return nil, errors.New("GPU not configured") + } + if len(passwords) == 0 { + return nil, nil + } + if len(passwords) > g.Capacity { + return nil, fmt.Errorf("GPU batch %d exceeds configured capacity %d", len(passwords), g.Capacity) + } + if len(salt) > 64 { + return nil, fmt.Errorf("GPU salt length %d exceeds 64 bytes", len(salt)) + } + + packed := make([]byte, len(passwords)*gpuMaxPasswordLen) + lens := make([]uint32, len(passwords)) + for i, pw := range passwords { + if len(pw) > gpuMaxPasswordLen { + return nil, fmt.Errorf("password %d is %d bytes; GPU limit is %d", i, len(pw), gpuMaxPasswordLen) + } + copy(packed[i*gpuMaxPasswordLen:], pw) + lens[i] = uint32(len(pw)) + } + rawOut := make([]byte, len(passwords)*32) + var saltPtr *C.uchar + if len(salt) != 0 { + saltPtr = (*C.uchar)(unsafe.Pointer(&salt[0])) + } + var errbuf [8192]C.char + rc := C.ycl_gpu_hash( + g.ctx, + (*C.uchar)(unsafe.Pointer(&packed[0])), + (*C.uint32_t)(unsafe.Pointer(&lens[0])), + C.uint32_t(len(passwords)), + saltPtr, + C.uint32_t(len(salt)), + (*C.uchar)(unsafe.Pointer(&rawOut[0])), + &errbuf[0], C.size_t(len(errbuf)), + ) + runtime.KeepAlive(packed) + runtime.KeepAlive(lens) + runtime.KeepAlive(salt) + if rc == C.YCL_GPU_CANCELLED { + return nil, errGPUCancelled + } + if rc != 0 { + return nil, errors.New(cErrorString(&errbuf)) + } + out := make([][32]byte, len(passwords)) + for i := range out { + copy(out[i][:], rawOut[i*32:(i+1)*32]) + } + return out, nil +} diff --git a/gpu_opencl.h b/gpu_opencl.h new file mode 100644 index 0000000..fe693a9 --- /dev/null +++ b/gpu_opencl.h @@ -0,0 +1,33 @@ +#ifndef YESCRYPT_CRACK_GPU_OPENCL_H +#define YESCRYPT_CRACK_GPU_OPENCL_H + +#include +#include + +#define YCL_GPU_CANCELLED 1 + +typedef struct ycl_gpu_ctx ycl_gpu_ctx; + +int ycl_opencl_device_count(char *err, size_t errlen); +int ycl_opencl_device_info(int index, char *name, size_t namelen, + uint64_t *global_mem, uint64_t *max_alloc, + char *err, size_t errlen); + +ycl_gpu_ctx *ycl_gpu_create(int index, const char *source, size_t source_len, + char *err, size_t errlen); +int ycl_gpu_configure(ycl_gpu_ctx *g, uint32_t N, uint32_t r, + uint32_t batch_hint, uint32_t *capacity, + char *err, size_t errlen); +int ycl_gpu_hash(ycl_gpu_ctx *g, + const unsigned char *passwords, + const uint32_t *pw_lens, + uint32_t count, + const unsigned char *salt, + uint32_t salt_len, + unsigned char *out, + char *err, size_t errlen); +void ycl_gpu_cancel(ycl_gpu_ctx *g); +void ycl_gpu_reset_cancel(ycl_gpu_ctx *g); +void ycl_gpu_destroy(ycl_gpu_ctx *g); + +#endif diff --git a/gpu_opencl_stub.go b/gpu_opencl_stub.go new file mode 100644 index 0000000..60782f7 --- /dev/null +++ b/gpu_opencl_stub.go @@ -0,0 +1,47 @@ +//go:build !windows && (!linux || !cgo) + +package main + +import "errors" + +const gpuMaxPasswordLen = 256 + +var ( + errGPUCancelled = errors.New("GPU cancelled") + errOpenCLUnsupported = errors.New("OpenCL GPU backend is unavailable in this build") +) + +type GPUDeviceInfo struct { + Index int + Name string + GlobalMem uint64 + MaxAlloc uint64 +} + +type OpenCLGPU struct { + Info GPUDeviceInfo + N uint32 + R uint32 + Capacity int +} + +func ListOpenCLGPUs() ([]GPUDeviceInfo, error) { + return nil, errOpenCLUnsupported +} + +func NewOpenCLGPU(info GPUDeviceInfo) (*OpenCLGPU, error) { + return nil, errOpenCLUnsupported +} + +func (g *OpenCLGPU) Close() {} +func (g *OpenCLGPU) Cancel() {} +func (g *OpenCLGPU) ResetCancel() {} +func (g *OpenCLGPU) SelfTest() error { return errOpenCLUnsupported } + +func (g *OpenCLGPU) Configure(N, r uint32, batchHint int) (int, error) { + return 0, errOpenCLUnsupported +} + +func (g *OpenCLGPU) HashBatch(passwords [][]byte, salt []byte) ([][32]byte, error) { + return nil, errOpenCLUnsupported +} diff --git a/gpu_opencl_windows.go b/gpu_opencl_windows.go new file mode 100644 index 0000000..4f1f33d --- /dev/null +++ b/gpu_opencl_windows.go @@ -0,0 +1,922 @@ +//go:build windows + +package main + +import ( + "bytes" + _ "embed" + "errors" + "fmt" + "runtime" + "sync" + "sync/atomic" + "syscall" + "unsafe" +) + +//go:embed kernels/yescrypt.cl +var yescryptOpenCLSource string + +const gpuMaxPasswordLen = 256 + +var errGPUCancelled = errors.New("GPU cancelled") + +type GPUDeviceInfo struct { + Index int + Name string + GlobalMem uint64 + MaxAlloc uint64 +} + +const ( + clSuccess int32 = 0 + clDeviceNotFound int32 = -1 + clPlatformNotFoundKHR int32 = -1001 + clDeviceTypeGPU uint64 = 1 << 2 + clDeviceName uint32 = 0x102B + clDeviceGlobalMemSize uint32 = 0x101F + clDeviceMaxMemAllocSize uint32 = 0x1010 + clProgramBuildLog uint32 = 0x1183 + clMemReadWrite uint64 = 1 << 0 + clMemWriteOnly uint64 = 1 << 1 + clMemReadOnly uint64 = 1 << 2 + clFalse uint32 = 0 + clTrue uint32 = 1 + windowsMaxGPUDevices = 32 + windowsMaxR = 32 + windowsSWords = 1536 + windowsOneMiB uint64 = 1024 * 1024 +) + +type openCLWinAPI struct { + dll *syscall.DLL + clGetPlatformIDs *syscall.Proc + clGetDeviceIDs *syscall.Proc + clGetDeviceInfo *syscall.Proc + clCreateContext *syscall.Proc + clReleaseContext *syscall.Proc + clCreateCommandQueue *syscall.Proc + clReleaseCommandQueue *syscall.Proc + clCreateBuffer *syscall.Proc + clReleaseMemObject *syscall.Proc + clCreateProgramWithSource *syscall.Proc + clBuildProgram *syscall.Proc + clGetProgramBuildInfo *syscall.Proc + clReleaseProgram *syscall.Proc + clCreateKernel *syscall.Proc + clReleaseKernel *syscall.Proc + clSetKernelArg *syscall.Proc + clEnqueueWriteBuffer *syscall.Proc + clEnqueueReadBuffer *syscall.Proc + clEnqueueNDRangeKernel *syscall.Proc + clFinish *syscall.Proc +} + +var ( + openCLWinOnce sync.Once + openCLWin *openCLWinAPI + openCLWinErr error +) + +func loadOpenCLWin() (*openCLWinAPI, error) { + openCLWinOnce.Do(func() { + dll, err := syscall.LoadDLL("OpenCL.dll") + if err != nil { + openCLWinErr = fmt.Errorf("OpenCL loader not found: %w", err) + return + } + + a := &openCLWinAPI{dll: dll} + lookups := []struct { + name string + dst **syscall.Proc + }{ + {"clGetPlatformIDs", &a.clGetPlatformIDs}, + {"clGetDeviceIDs", &a.clGetDeviceIDs}, + {"clGetDeviceInfo", &a.clGetDeviceInfo}, + {"clCreateContext", &a.clCreateContext}, + {"clReleaseContext", &a.clReleaseContext}, + {"clCreateCommandQueue", &a.clCreateCommandQueue}, + {"clReleaseCommandQueue", &a.clReleaseCommandQueue}, + {"clCreateBuffer", &a.clCreateBuffer}, + {"clReleaseMemObject", &a.clReleaseMemObject}, + {"clCreateProgramWithSource", &a.clCreateProgramWithSource}, + {"clBuildProgram", &a.clBuildProgram}, + {"clGetProgramBuildInfo", &a.clGetProgramBuildInfo}, + {"clReleaseProgram", &a.clReleaseProgram}, + {"clCreateKernel", &a.clCreateKernel}, + {"clReleaseKernel", &a.clReleaseKernel}, + {"clSetKernelArg", &a.clSetKernelArg}, + {"clEnqueueWriteBuffer", &a.clEnqueueWriteBuffer}, + {"clEnqueueReadBuffer", &a.clEnqueueReadBuffer}, + {"clEnqueueNDRangeKernel", &a.clEnqueueNDRangeKernel}, + {"clFinish", &a.clFinish}, + } + for _, lookup := range lookups { + proc, err := dll.FindProc(lookup.name) + if err != nil { + _ = dll.Release() + openCLWinErr = fmt.Errorf("OpenCL symbol %s missing: %w", lookup.name, err) + return + } + *lookup.dst = proc + } + openCLWin = a + }) + if openCLWinErr != nil { + return nil, openCLWinErr + } + return openCLWin, nil +} + +func clStatus(p *syscall.Proc, args ...uintptr) int32 { + r1, _, _ := p.Call(args...) + return int32(r1) +} + +func ptr(v unsafe.Pointer) uintptr { return uintptr(v) } + +func enumerateOpenCLGPUs() ([]uintptr, error) { + a, err := loadOpenCLWin() + if err != nil { + return nil, err + } + + var platformCount uint32 + rc := clStatus(a.clGetPlatformIDs, 0, 0, ptr(unsafe.Pointer(&platformCount))) + if rc == clPlatformNotFoundKHR || platformCount == 0 { + return nil, nil + } + if rc != clSuccess { + return nil, fmt.Errorf("clGetPlatformIDs failed: %d", rc) + } + + platforms := make([]uintptr, platformCount) + rc = clStatus(a.clGetPlatformIDs, uintptr(platformCount), ptr(unsafe.Pointer(&platforms[0])), 0) + if rc != clSuccess { + return nil, fmt.Errorf("clGetPlatformIDs(list) failed: %d", rc) + } + + devices := make([]uintptr, 0, windowsMaxGPUDevices) + for _, platform := range platforms { + if len(devices) >= windowsMaxGPUDevices { + break + } + var deviceCount uint32 + rc = clStatus(a.clGetDeviceIDs, platform, uintptr(clDeviceTypeGPU), 0, 0, ptr(unsafe.Pointer(&deviceCount))) + if rc == clDeviceNotFound || deviceCount == 0 { + continue + } + if rc != clSuccess { + continue + } + + platformDevices := make([]uintptr, deviceCount) + rc = clStatus(a.clGetDeviceIDs, platform, uintptr(clDeviceTypeGPU), uintptr(deviceCount), ptr(unsafe.Pointer(&platformDevices[0])), 0) + if rc != clSuccess { + continue + } + for _, device := range platformDevices { + if len(devices) == windowsMaxGPUDevices { + break + } + devices = append(devices, device) + } + } + runtime.KeepAlive(platforms) + return devices, nil +} + +func openCLDeviceInfo(index int, device uintptr) (GPUDeviceInfo, error) { + a, err := loadOpenCLWin() + if err != nil { + return GPUDeviceInfo{}, err + } + + nameBuf := make([]byte, 512) + rc := clStatus(a.clGetDeviceInfo, device, uintptr(clDeviceName), uintptr(len(nameBuf)), ptr(unsafe.Pointer(&nameBuf[0])), 0) + name := fmt.Sprintf("OpenCL GPU %d", index) + if rc == clSuccess { + if end := bytes.IndexByte(nameBuf, 0); end >= 0 { + nameBuf = nameBuf[:end] + } + if len(nameBuf) != 0 { + name = string(nameBuf) + } + } + + var globalMem uint64 + rc = clStatus(a.clGetDeviceInfo, device, uintptr(clDeviceGlobalMemSize), unsafe.Sizeof(globalMem), ptr(unsafe.Pointer(&globalMem)), 0) + if rc != clSuccess { + return GPUDeviceInfo{}, errors.New("cannot query GPU global memory") + } + var maxAlloc uint64 + rc = clStatus(a.clGetDeviceInfo, device, uintptr(clDeviceMaxMemAllocSize), unsafe.Sizeof(maxAlloc), ptr(unsafe.Pointer(&maxAlloc)), 0) + if rc != clSuccess { + return GPUDeviceInfo{}, errors.New("cannot query GPU max allocation") + } + + return GPUDeviceInfo{Index: index, Name: name, GlobalMem: globalMem, MaxAlloc: maxAlloc}, nil +} + +func ListOpenCLGPUs() ([]GPUDeviceInfo, error) { + devices, err := enumerateOpenCLGPUs() + if err != nil { + return nil, err + } + out := make([]GPUDeviceInfo, 0, len(devices)) + for i, device := range devices { + info, err := openCLDeviceInfo(i, device) + if err != nil { + return nil, err + } + out = append(out, info) + } + return out, nil +} + +type OpenCLGPU struct { + Info GPUDeviceInfo + N uint32 + R uint32 + Capacity int + + dev uintptr + ctx uintptr + queue uintptr + + source []byte + program uintptr + initKernel uintptr + loopKernel uintptr + finalKernel uintptr + + globalMem uint64 + maxAlloc uint64 + cancel atomic.Bool + + passwords uintptr + pwLens uintptr + salt uintptr + p uintptr + scratch uintptr + s uintptr + yctx uintptr + v [4]uintptr + out uintptr +} + +func NewOpenCLGPU(info GPUDeviceInfo) (*OpenCLGPU, error) { + if yescryptOpenCLSource == "" { + return nil, errors.New("embedded OpenCL kernel is empty") + } + a, err := loadOpenCLWin() + if err != nil { + return nil, err + } + devices, err := enumerateOpenCLGPUs() + if err != nil { + return nil, err + } + if info.Index < 0 || info.Index >= len(devices) { + return nil, fmt.Errorf("OpenCL GPU index %d out of range (found %d)", info.Index, len(devices)) + } + + g := &OpenCLGPU{ + Info: info, + dev: devices[info.Index], + source: []byte(yescryptOpenCLSource), + globalMem: info.GlobalMem, + maxAlloc: info.MaxAlloc, + } + if g.globalMem == 0 || g.maxAlloc == 0 { + queried, qerr := openCLDeviceInfo(info.Index, g.dev) + if qerr != nil { + return nil, qerr + } + g.Info = queried + g.globalMem = queried.GlobalMem + g.maxAlloc = queried.MaxAlloc + } + + var rc int32 + dev := g.dev + r1, _, _ := a.clCreateContext.Call(0, 1, ptr(unsafe.Pointer(&dev)), 0, 0, ptr(unsafe.Pointer(&rc))) + g.ctx = r1 + if g.ctx == 0 || rc != clSuccess { + g.Close() + return nil, fmt.Errorf("clCreateContext failed: %d", rc) + } + + r1, _, _ = a.clCreateCommandQueue.Call(g.ctx, g.dev, 0, ptr(unsafe.Pointer(&rc))) + g.queue = r1 + if g.queue == 0 || rc != clSuccess { + g.Close() + return nil, fmt.Errorf("clCreateCommandQueue failed: %d", rc) + } + return g, nil +} + +func (g *OpenCLGPU) releaseBuffers() { + if g == nil { + return + } + a, err := loadOpenCLWin() + if err != nil { + return + } + all := []*uintptr{ + &g.passwords, &g.pwLens, &g.salt, &g.p, &g.scratch, &g.s, &g.yctx, + &g.v[0], &g.v[1], &g.v[2], &g.v[3], &g.out, + } + for _, mem := range all { + if *mem != 0 { + _, _, _ = a.clReleaseMemObject.Call(*mem) + *mem = 0 + } + } + g.Capacity = 0 +} + +func (g *OpenCLGPU) releaseProgram() { + if g == nil { + return + } + a, err := loadOpenCLWin() + if err != nil { + return + } + if g.initKernel != 0 { + _, _, _ = a.clReleaseKernel.Call(g.initKernel) + g.initKernel = 0 + } + if g.loopKernel != 0 { + _, _, _ = a.clReleaseKernel.Call(g.loopKernel) + g.loopKernel = 0 + } + if g.finalKernel != 0 { + _, _, _ = a.clReleaseKernel.Call(g.finalKernel) + g.finalKernel = 0 + } + if g.program != 0 { + _, _, _ = a.clReleaseProgram.Call(g.program) + g.program = 0 + } + g.N = 0 + g.R = 0 +} + +func (g *OpenCLGPU) Close() { + if g == nil { + return + } + a, _ := loadOpenCLWin() + g.releaseBuffers() + g.releaseProgram() + if a != nil { + if g.queue != 0 { + _, _, _ = a.clReleaseCommandQueue.Call(g.queue) + g.queue = 0 + } + if g.ctx != 0 { + _, _, _ = a.clReleaseContext.Call(g.ctx) + g.ctx = 0 + } + } + g.source = nil +} + +func (g *OpenCLGPU) Cancel() { + if g != nil { + g.cancel.Store(true) + } +} + +func (g *OpenCLGPU) ResetCancel() { + if g != nil { + g.cancel.Store(false) + } +} + +func (g *OpenCLGPU) SelfTest() error { + const vector = "$y$j9T$Gg3nKTjUa.Hrq3ZPL9S9J0$kJBFNrOZC2sCjYIOV3G4/NTOOFWUULJDthOnxtiTez9" + h, err := parseYescryptHash(vector) + if err != nil { + return fmt.Errorf("internal GPU self-test vector parse failed: %w", err) + } + if _, err := g.Configure(h.N, h.R, 1); err != nil { + return fmt.Errorf("GPU self-test configure failed: %w", err) + } + out, err := g.HashBatch([][]byte{[]byte("password")}, h.Salt) + if err != nil { + return fmt.Errorf("GPU self-test execution failed: %w", err) + } + if len(out) != 1 || !bytes.Equal(out[0][:], h.Expected[:]) { + if len(out) == 1 { + return fmt.Errorf("GPU self-test digest mismatch: got %x want %x", out[0], h.Expected) + } + return fmt.Errorf("GPU self-test returned %d digests", len(out)) + } + return nil +} + +func nulTerminated(s string) []byte { + b := make([]byte, len(s)+1) + copy(b, s) + return b +} + +func (g *OpenCLGPU) buildSpecializedProgram(N, r uint32) error { + if g.program != 0 && g.N == N && g.R == r { + return nil + } + g.releaseProgram() + + a, err := loadOpenCLWin() + if err != nil { + return err + } + if len(g.source) == 0 { + return errors.New("empty OpenCL source") + } + + sourcePtr := ptr(unsafe.Pointer(&g.source[0])) + sources := [1]uintptr{sourcePtr} + lengths := [1]uintptr{uintptr(len(g.source))} + var rc int32 + r1, _, _ := a.clCreateProgramWithSource.Call( + g.ctx, + 1, + ptr(unsafe.Pointer(&sources[0])), + ptr(unsafe.Pointer(&lengths[0])), + ptr(unsafe.Pointer(&rc)), + ) + g.program = r1 + if g.program == 0 || rc != clSuccess { + g.releaseProgram() + return fmt.Errorf("clCreateProgramWithSource failed: %d", rc) + } + runtime.KeepAlive(g.source) + + optionsText := fmt.Sprintf("-cl-std=CL1.2 -DYC_N=%d -DYC_R=%d", N, r) + options := nulTerminated(optionsText) + dev := g.dev + rc = clStatus(a.clBuildProgram, + g.program, + 1, + ptr(unsafe.Pointer(&dev)), + ptr(unsafe.Pointer(&options[0])), + 0, + 0, + ) + if rc != clSuccess { + var logSize uintptr + _ = clStatus(a.clGetProgramBuildInfo, g.program, g.dev, uintptr(clProgramBuildLog), 0, 0, ptr(unsafe.Pointer(&logSize))) + logText := "no build log" + if logSize > 0 { + logBuf := make([]byte, int(logSize)+1) + if clStatus(a.clGetProgramBuildInfo, g.program, g.dev, uintptr(clProgramBuildLog), logSize, ptr(unsafe.Pointer(&logBuf[0])), 0) == clSuccess { + if end := bytes.IndexByte(logBuf, 0); end >= 0 { + logBuf = logBuf[:end] + } + if len(logBuf) != 0 { + logText = string(logBuf) + } + } + } + g.releaseProgram() + return fmt.Errorf("OpenCL kernel build failed (%d) [%s]: %s", rc, optionsText, logText) + } + + createKernel := func(name string) (uintptr, error) { + kernelName := nulTerminated(name) + var kernelRC int32 + kernel, _, _ := a.clCreateKernel.Call(g.program, ptr(unsafe.Pointer(&kernelName[0])), ptr(unsafe.Pointer(&kernelRC))) + if kernel == 0 || kernelRC != clSuccess { + return 0, fmt.Errorf("clCreateKernel(%s) failed: %d", name, kernelRC) + } + return kernel, nil + } + if g.initKernel, err = createKernel("yescrypt_init"); err != nil { + g.releaseProgram() + return err + } + if g.loopKernel, err = createKernel("yescrypt_loop"); err != nil { + g.releaseProgram() + return err + } + if g.finalKernel, err = createKernel("yescrypt_final"); err != nil { + g.releaseProgram() + return err + } + + g.N = N + g.R = r + return nil +} + +func (g *OpenCLGPU) makeBuffer(flags uint64, size uint64, what string) (uintptr, error) { + a, err := loadOpenCLWin() + if err != nil { + return 0, err + } + if size == 0 { + size = 8 + } + var rc int32 + mem, _, _ := a.clCreateBuffer.Call(g.ctx, uintptr(flags), uintptr(size), 0, ptr(unsafe.Pointer(&rc))) + if mem == 0 || rc != clSuccess { + return 0, fmt.Errorf("clCreateBuffer(%s, %.2f MiB) failed: %d", what, float64(size)/(1024.0*1024.0), rc) + } + return mem, nil +} + +func (g *OpenCLGPU) Configure(N, r uint32, batchHint int) (int, error) { + if g == nil || g.ctx == 0 { + return 0, errors.New("GPU context is closed") + } + if batchHint < 0 { + batchHint = 0 + } + if N == 0 || r == 0 || r > windowsMaxR || N&(N-1) != 0 { + return 0, fmt.Errorf("invalid yescrypt GPU parameters N=%d r=%d", N, r) + } + + geometryChanged := g.N != N || g.R != r || g.program == 0 + if !geometryChanged && g.Capacity != 0 && (batchHint == 0 || batchHint == g.Capacity) { + return g.Capacity, nil + } + + g.releaseBuffers() + if err := g.buildSpecializedProgram(N, r); err != nil { + return 0, err + } + + stateBytes := uint64(128) * uint64(r) + vbytes := stateBytes * uint64(N) + scratchBytes := stateBytes + if scratchBytes < 256 { + scratchBytes = 256 + } + const ctxBytes = uint64(48) + workspace := uint64(gpuMaxPasswordLen) + 4 + 64 + stateBytes + scratchBytes + windowsSWords*8 + ctxBytes + 32 + + reserve := g.globalMem / 10 + if reserve < 512*windowsOneMiB { + reserve = 512 * windowsOneMiB + } + if reserve >= g.globalMem { + reserve = g.globalMem / 5 + } + usable := g.globalMem - reserve + capMem := usable / (vbytes + workspace) + + allocLimit := g.maxAlloc * 95 / 100 + slotsPerSegment := uint64(0) + if vbytes != 0 { + slotsPerSegment = allocLimit / vbytes + } + capAlloc := slotsPerSegment * 4 + cap := capMem + if cap > capAlloc { + cap = capAlloc + } + if cap > 4096 { + cap = 4096 + } + if batchHint > 0 { + if cap > uint64(batchHint) { + cap = uint64(batchHint) + } + } else { + cap = cap * 95 / 100 + if cap >= 32 { + cap = cap / 32 * 32 + } + } + if cap == 0 { + return 0, fmt.Errorf("GPU cannot allocate one yescrypt V region: need %.2f MiB/candidate, max allocation %.2f MiB", float64(vbytes)/float64(windowsOneMiB), float64(g.maxAlloc)/float64(windowsOneMiB)) + } + if cap >= 4 { + cap = cap / 4 * 4 + } + if cap == 0 { + cap = 1 + } + g.Capacity = int(cap) + + var err error + if g.passwords, err = g.makeBuffer(clMemReadOnly, cap*uint64(gpuMaxPasswordLen), "passwords"); err != nil { + goto fail + } + if g.pwLens, err = g.makeBuffer(clMemReadOnly, cap*4, "pw_lens"); err != nil { + goto fail + } + if g.salt, err = g.makeBuffer(clMemReadOnly, 64, "salt"); err != nil { + goto fail + } + if g.p, err = g.makeBuffer(clMemReadWrite, cap*stateBytes, "P"); err != nil { + goto fail + } + if g.scratch, err = g.makeBuffer(clMemReadWrite, cap*scratchBytes, "scratch"); err != nil { + goto fail + } + if g.s, err = g.makeBuffer(clMemReadWrite, cap*windowsSWords*8, "S"); err != nil { + goto fail + } + if g.yctx, err = g.makeBuffer(clMemReadWrite, cap*ctxBytes, "context"); err != nil { + goto fail + } + for seg := 0; seg < 4; seg++ { + segCount := (cap + uint64(3-seg)) / 4 + if g.v[seg], err = g.makeBuffer(clMemReadWrite, segCount*vbytes, fmt.Sprintf("V%d", seg)); err != nil { + goto fail + } + } + if g.out, err = g.makeBuffer(clMemWriteOnly, cap*32, "output"); err != nil { + goto fail + } + return g.Capacity, nil + +fail: + g.releaseBuffers() + return 0, err +} + +func (g *OpenCLGPU) setKernelMem(kernel uintptr, arg *uint32, mem uintptr, label string) error { + a, err := loadOpenCLWin() + if err != nil { + return err + } + value := mem + rc := clStatus(a.clSetKernelArg, kernel, uintptr(*arg), unsafe.Sizeof(value), ptr(unsafe.Pointer(&value))) + if rc != clSuccess { + return fmt.Errorf("%s failed: %d", label, rc) + } + (*arg)++ + return nil +} + +func (g *OpenCLGPU) setKernelU32(kernel uintptr, arg *uint32, value uint32, label string) error { + a, err := loadOpenCLWin() + if err != nil { + return err + } + v := value + rc := clStatus(a.clSetKernelArg, kernel, uintptr(*arg), unsafe.Sizeof(v), ptr(unsafe.Pointer(&v))) + if rc != clSuccess { + return fmt.Errorf("%s failed: %d", label, rc) + } + (*arg)++ + return nil +} + +func (g *OpenCLGPU) enqueueWrite(mem uintptr, size uintptr, data unsafe.Pointer, what string) error { + a, err := loadOpenCLWin() + if err != nil { + return err + } + rc := clStatus(a.clEnqueueWriteBuffer, g.queue, mem, uintptr(clFalse), 0, size, ptr(data), 0, 0, 0) + if rc != clSuccess { + return fmt.Errorf("%s failed: %d", what, rc) + } + return nil +} + +func (g *OpenCLGPU) enqueueKernel(kernel uintptr, global, local uintptr, what string) error { + a, err := loadOpenCLWin() + if err != nil { + return err + } + globalSize := global + var localPtr uintptr + var localSize uintptr + if local != 0 { + localSize = local + localPtr = ptr(unsafe.Pointer(&localSize)) + } + rc := clStatus(a.clEnqueueNDRangeKernel, g.queue, kernel, 1, 0, ptr(unsafe.Pointer(&globalSize)), localPtr, 0, 0, 0) + if rc != clSuccess { + return fmt.Errorf("%s failed: %d", what, rc) + } + return nil +} + +func (g *OpenCLGPU) finish(what string) error { + a, err := loadOpenCLWin() + if err != nil { + return err + } + rc := clStatus(a.clFinish, g.queue) + if rc != clSuccess { + return fmt.Errorf("%s failed: %d", what, rc) + } + return nil +} + +func (g *OpenCLGPU) enqueueRead(mem uintptr, size uintptr, data unsafe.Pointer, what string) error { + a, err := loadOpenCLWin() + if err != nil { + return err + } + rc := clStatus(a.clEnqueueReadBuffer, g.queue, mem, uintptr(clTrue), 0, size, ptr(data), 0, 0, 0) + if rc != clSuccess { + return fmt.Errorf("%s failed: %d", what, rc) + } + return nil +} + +// HashBatch mirrors the tuned Linux host sequence. Windows uses the same +// embedded kernel and workgroup geometry, but calls OpenCL.dll directly from Go. +func (g *OpenCLGPU) HashBatch(passwords [][]byte, salt []byte) ([][32]byte, error) { + if g == nil || g.ctx == 0 || g.Capacity == 0 { + return nil, errors.New("GPU not configured") + } + if len(passwords) == 0 { + return nil, nil + } + if len(passwords) > g.Capacity { + return nil, fmt.Errorf("GPU batch %d exceeds configured capacity %d", len(passwords), g.Capacity) + } + if len(salt) > 64 { + return nil, fmt.Errorf("GPU salt length %d exceeds 64 bytes", len(salt)) + } + if g.cancel.Load() { + return nil, errGPUCancelled + } + + packed := make([]byte, len(passwords)*gpuMaxPasswordLen) + lens := make([]uint32, len(passwords)) + for i, pw := range passwords { + if len(pw) > gpuMaxPasswordLen { + return nil, fmt.Errorf("password %d is %d bytes; GPU limit is %d", i, len(pw), gpuMaxPasswordLen) + } + copy(packed[i*gpuMaxPasswordLen:], pw) + lens[i] = uint32(len(pw)) + } + rawOut := make([]byte, len(passwords)*32) + + var hostPin runtime.Pinner + hostPin.Pin(&packed[0]) + hostPin.Pin(&lens[0]) + hostPin.Pin(&rawOut[0]) + defer hostPin.Unpin() + + if err := g.hashPacked(packed, lens, salt, rawOut); err != nil { + return nil, err + } + runtime.KeepAlive(packed) + runtime.KeepAlive(lens) + runtime.KeepAlive(salt) + runtime.KeepAlive(rawOut) + + out := make([][32]byte, len(passwords)) + for i := range out { + copy(out[i][:], rawOut[i*32:(i+1)*32]) + } + return out, nil +} + +func (g *OpenCLGPU) hashPacked(packed []byte, lens []uint32, salt []byte, rawOut []byte) error { + count := uint32(len(lens)) + if count == 0 || int(count) > g.Capacity || g.program == 0 { + return fmt.Errorf("bad GPU batch count %d (capacity %d)", count, g.Capacity) + } + if g.cancel.Load() { + return errGPUCancelled + } + + if err := g.enqueueWrite(g.passwords, uintptr(count)*gpuMaxPasswordLen, unsafe.Pointer(&packed[0]), "write passwords"); err != nil { + return err + } + if err := g.enqueueWrite(g.pwLens, uintptr(count)*4, unsafe.Pointer(&lens[0]), "write password lengths"); err != nil { + return err + } + var saltBuf [64]byte + copy(saltBuf[:], salt) + var saltPin runtime.Pinner + saltPin.Pin(&saltBuf[0]) + defer saltPin.Unpin() + if err := g.enqueueWrite(g.salt, 64, unsafe.Pointer(&saltBuf[0]), "write salt"); err != nil { + return err + } + + var a uint32 + for _, item := range []struct { + mem uintptr + label string + }{ + {g.passwords, "set init passwords"}, + {g.pwLens, "set init lengths"}, + {g.salt, "set init salt"}, + } { + if err := g.setKernelMem(g.initKernel, &a, item.mem, item.label); err != nil { + return err + } + } + if err := g.setKernelU32(g.initKernel, &a, uint32(len(salt)), "set init salt_len"); err != nil { + return err + } + if err := g.setKernelU32(g.initKernel, &a, count, "set init count"); err != nil { + return err + } + for _, item := range []struct { + mem uintptr + label string + }{ + {g.p, "set init P"}, + {g.scratch, "set init scratch"}, + {g.s, "set init S"}, + {g.yctx, "set init context"}, + {g.v[0], "set init V0"}, + {g.v[1], "set init V1"}, + {g.v[2], "set init V2"}, + {g.v[3], "set init V3"}, + } { + if err := g.setKernelMem(g.initKernel, &a, item.mem, item.label); err != nil { + return err + } + } + if err := g.enqueueKernel(g.initKernel, uintptr(count), 0, "launch yescrypt init"); err != nil { + return err + } + + nloop := (((uint64(g.N) + 2) / 3) + 1) &^ uint64(1) + remaining := uint64(g.N) + nloop + const localLoop = uintptr(32) + globalLoop := uintptr(count) * localLoop + for remaining != 0 { + if g.cancel.Load() { + _ = g.finish("cancel wait") + return errGPUCancelled + } + chunk := uint32(remaining) + if remaining > 2048 { + chunk = 2048 + } + a = 0 + if err := g.setKernelU32(g.loopKernel, &a, count, "set loop count"); err != nil { + return err + } + if err := g.setKernelU32(g.loopKernel, &a, chunk, "set loop chunk"); err != nil { + return err + } + for _, item := range []struct { + mem uintptr + label string + }{ + {g.p, "set loop P"}, + {g.s, "set loop S"}, + {g.yctx, "set loop context"}, + {g.v[0], "set loop V0"}, + {g.v[1], "set loop V1"}, + {g.v[2], "set loop V2"}, + {g.v[3], "set loop V3"}, + } { + if err := g.setKernelMem(g.loopKernel, &a, item.mem, item.label); err != nil { + return err + } + } + if err := g.enqueueKernel(g.loopKernel, globalLoop, localLoop, "launch yescrypt cooperative loop"); err != nil { + return err + } + if err := g.finish("wait for yescrypt cooperative loop"); err != nil { + return err + } + if g.cancel.Load() { + return errGPUCancelled + } + remaining -= uint64(chunk) + } + + if g.cancel.Load() { + return errGPUCancelled + } + a = 0 + if err := g.setKernelU32(g.finalKernel, &a, count, "set final count"); err != nil { + return err + } + for _, item := range []struct { + mem uintptr + label string + }{ + {g.p, "set final P"}, + {g.scratch, "set final scratch"}, + {g.yctx, "set final context"}, + {g.out, "set final output"}, + } { + if err := g.setKernelMem(g.finalKernel, &a, item.mem, item.label); err != nil { + return err + } + } + if err := g.enqueueKernel(g.finalKernel, uintptr(count), 0, "launch yescrypt final"); err != nil { + return err + } + if err := g.enqueueRead(g.out, uintptr(count)*32, unsafe.Pointer(&rawOut[0]), "read yescrypt output"); err != nil { + return err + } + if err := g.finish("clFinish"); err != nil { + return err + } + return nil +} diff --git a/hash.go b/hash.go index 4529894..58c51af 100644 --- a/hash.go +++ b/hash.go @@ -2,17 +2,76 @@ package main import ( "bufio" - "bytes" + "crypto/subtle" + "fmt" "os" - "regexp" "strings" +) + +type HashKind uint8 - "github.com/openwall/yescrypt-go" +const ( + HashYescrypt HashKind = iota + 1 + HashGostYescrypt ) type YescryptHash struct { - Hash []byte - Cracked int32 + Hash []byte + Cracked int32 + Kind HashKind + Params string + Salt []byte + Expected [32]byte + N uint32 + R uint32 + GPUOK bool + Setting []byte +} + +func parseYescryptHash(line string) (YescryptHash, error) { + var out YescryptHash + parts := strings.Split(line, "$") + if len(parts) != 5 || parts[0] != "" { + return out, fmt.Errorf("bad field count") + } + switch parts[1] { + case "y": + out.Kind = HashYescrypt + case "gy": + out.Kind = HashGostYescrypt + default: + return out, fmt.Errorf("unsupported signature") + } + if parts[2] == "" || len(parts[3]) > 86 || len(parts[4]) != 43 { + return out, fmt.Errorf("bad yescrypt field length") + } + + salt := decodeCryptBase64([]byte(parts[3])) + if salt == nil || len(salt) > 64 { + return out, fmt.Errorf("bad salt encoding") + } + digest := decodeCryptBase64([]byte(parts[4])) + if len(digest) != 32 { + return out, fmt.Errorf("bad digest encoding") + } + + out.Hash = []byte(line) + out.Params = parts[2] + out.Salt = salt + copy(out.Expected[:], digest) + last := strings.LastIndexByte(line, '$') + out.Setting = []byte(line[:last]) + + if len(out.Params) == 3 && out.Params[0] == 'j' { + nlog := cryptBase64Value(out.Params[1]) + 1 + r := cryptBase64Value(out.Params[2]) + 1 + if nlog >= 10 && nlog <= 18 && r >= 1 && r <= 32 { + out.N = 1 << uint(nlog) + out.R = uint32(r) + out.GPUOK = true + } + } + return out, nil } func ReadYescryptHashes(filePath string) ([]YescryptHash, error) { @@ -22,9 +81,6 @@ func ReadYescryptHashes(filePath string) ([]YescryptHash, error) { } defer file.Close() - // accept native yescrypt ($y$) and gost-yescrypt ($gy$) hashes - yescryptRegex := regexp.MustCompile(`^\$(y|gy)\$[./A-Za-z0-9]+\$[./A-Za-z0-9]{1,86}\$[./A-Za-z0-9]{43}$`) - var hashes []YescryptHash scanner := bufio.NewScanner(file) for scanner.Scan() { @@ -32,13 +88,11 @@ func ReadYescryptHashes(filePath string) ([]YescryptHash, error) { if line == "" { continue } - - // yescrypt / gost-yescrypt sanity check - if !yescryptRegex.MatchString(line) { + h, err := parseYescryptHash(line) + if err != nil { continue } - - hashes = append(hashes, YescryptHash{Hash: []byte(line), Cracked: 0}) + hashes = append(hashes, h) } if err := scanner.Err(); err != nil { return nil, err @@ -47,20 +101,33 @@ func ReadYescryptHashes(filePath string) ([]YescryptHash, error) { } func crackHash(password, fullHash []byte) bool { - switch { - case bytes.HasPrefix(fullHash, []byte("$y$")): - return crackYescrypt(password, fullHash) - case bytes.HasPrefix(fullHash, []byte("$gy$")): - return crackGostYescrypt(password, fullHash) - default: + h, err := parseYescryptHash(string(fullHash)) + if err != nil { return false } + return crackParsedHashCPU(password, &h) } -func crackYescrypt(password, fullHash []byte) bool { - generatedHash, err := yescrypt.Hash(password, fullHash) +func crackParsedHashCPU(password []byte, h *YescryptHash) bool { + if !h.GPUOK { + if len(h.Hash) != 0 { + parsed, err := parseYescryptHash(string(h.Hash)) + if err == nil && parsed.GPUOK { + return crackParsedHashCPU(password, &parsed) + } + } + return false + } + digest, err := yescryptKeyCPU(password, h.Salt, int(h.N), int(h.R), 1, 32) + if err != nil { + return false + } + if h.Kind == HashYescrypt { + return subtle.ConstantTimeCompare(digest, h.Expected[:]) == 1 + } + generated, err := gostFinalize(password, h.Setting, digest) if err != nil { return false } - return bytes.Equal(fullHash, generatedHash) + return constantTimeDigestEqual(generated, h.Expected[:]) } diff --git a/hash_test.go b/hash_test.go index c2af449..1fd6299 100644 --- a/hash_test.go +++ b/hash_test.go @@ -9,13 +9,14 @@ import ( "sync" "sync/atomic" "testing" + "time" ) const ( testPassword = "password" - testYescrypt = "$y$j9T$ukgaTIHHgVLdJH9qAK9Nz/$D6rr9OXstjx/QksjwzcO2M2UYgO.98RNVWt9jw1aW/9" - testGostYescrypt = "$gy$j9T$ukgaTIHHgVLdJH9qAK9Nz/$bH5kn7UF0Sk8ZgVzI6HWILrRemSMLVyJTiZgWbASi83" - testGostYescrypt2 = "$gy$jAT$0123456789abcdef0123456789abcdef$nhEOJFlM.wasdffY0jRHd0ACl2sH.1CuV0kREkGTKM8" + testYescrypt = "$y$j9T$Gg3nKTjUa.Hrq3ZPL9S9J0$kJBFNrOZC2sCjYIOV3G4/NTOOFWUULJDthOnxtiTez9" + testGostYescrypt = "$gy$j9T$Gg3nKTjUa.Hrq3ZPL9S9J0$NoLVpQb1ewFIXSN.3PShWUybuC8xZJ8HER18hi7lU73" + testGostYescrypt2 = "$gy$jAT$XDfTHhSYsW86PT.3V3l.UVh.yTta/MEA$gTqcPPYO9SMD3p0K0e0jHic0nwVNNFEgM01pftAVfdB" ) func TestCrackHashDispatch(t *testing.T) { @@ -27,7 +28,7 @@ func TestCrackHashDispatch(t *testing.T) { if !crackHash(password, []byte(testGostYescrypt)) { t.Fatal("gost-yescrypt test vector did not match") } - if !crackHash([]byte("hunter2"), []byte(testGostYescrypt2)) { + if !crackHash(password, []byte(testGostYescrypt2)) { t.Fatal("second gost-yescrypt test vector did not match") } if crackHash([]byte("wrong-password"), []byte(testGostYescrypt)) { @@ -61,8 +62,8 @@ func TestProcessPasswordChecksAllUncrackedHashes(t *testing.T) { writer := bufio.NewWriter(&output) var writerMu sync.Mutex var crackedCount int32 - var linesProcessed int32 - var totalHashesGenerated int32 + var linesProcessed uint64 + var totalHashesGenerated uint64 stopChan := make(chan struct{}) processPassword( @@ -85,10 +86,10 @@ func TestProcessPasswordChecksAllUncrackedHashes(t *testing.T) { if atomic.LoadInt32(&hashes[0].Cracked) != 1 || atomic.LoadInt32(&hashes[1].Cracked) != 1 { t.Fatalf("expected both hashes to be marked cracked: %+v", hashes) } - if got := atomic.LoadInt32(&linesProcessed); got != 1 { + if got := atomic.LoadUint64(&linesProcessed); got != 1 { t.Fatalf("expected one wordlist candidate to be processed, got %d", got) } - if got := atomic.LoadInt32(&totalHashesGenerated); got != 2 { + if got := atomic.LoadUint64(&totalHashesGenerated); got != 2 { t.Fatalf("expected candidate to be tested against both hashes, got %d hash attempts", got) } if got := strings.Count(output.String(), ":password\n"); got != 2 { @@ -102,3 +103,113 @@ func TestProcessPasswordChecksAllUncrackedHashes(t *testing.T) { t.Fatal("expected stop channel to close after all hashes were cracked") } } + +func TestCPUKeyMatchesKnownYescryptDigest(t *testing.T) { + h, err := parseYescryptHash(testYescrypt) + if err != nil { + t.Fatal(err) + } + got, err := yescryptKeyCPU([]byte(testPassword), h.Salt, int(h.N), int(h.R), 1, 32) + if err != nil { + t.Fatal(err) + } + if !bytes.Equal(got, h.Expected[:]) { + t.Fatalf("raw yescrypt mismatch: got %x want %x", got, h.Expected) + } +} + +func TestGostFinalizeKnownVector(t *testing.T) { + gy, err := parseYescryptHash(testGostYescrypt) + if err != nil { + t.Fatal(err) + } + y, err := parseYescryptHash(testYescrypt) + if err != nil { + t.Fatal(err) + } + got, err := gostFinalize([]byte(testPassword), gy.Setting, y.Expected[:]) + if err != nil { + t.Fatal(err) + } + if !constantTimeDigestEqual(got, gy.Expected[:]) { + t.Fatalf("gost finalization mismatch: got %x want %x", got, gy.Expected) + } +} + +func TestGPUCapacityJ9T24GiB(t *testing.T) { + info := GPUDeviceInfo{GlobalMem: 24 << 30, MaxAlloc: 6 << 30} + cap := estimateGPUCapacity(info, 4096, 32, 0) + if cap != 1280 { + t.Fatalf("expected tuned j9T auto batch 1280 on a 24GiB/6GiB-allocation GPU, got %d", cap) + } + + manual := estimateGPUCapacity(info, 4096, 32, 1352) + if manual != 1352 { + t.Fatalf("expected explicit GPU batch 1352 to be preserved, got %d", manual) + } +} + +func TestTuneAutoGPUCapacity4090(t *testing.T) { + if got := tuneAutoGPUCapacity(1352); got != 1280 { + t.Fatalf("expected measured 4090 cap 1352 to tune to 1280, got %d", got) + } +} + +func TestNormalizeGPUArgs(t *testing.T) { + tests := []struct { + in []string + want []string + }{ + {[]string{"yescrypt_crack.bin", "-gpu", "-h", "hashes.txt"}, []string{"yescrypt_crack.bin", "-gpu=all", "-h", "hashes.txt"}}, + {[]string{"yescrypt_crack.bin", "-gpu", "0,1", "-h", "hashes.txt"}, []string{"yescrypt_crack.bin", "-gpu=0,1", "-h", "hashes.txt"}}, + {[]string{"yescrypt_crack.bin", "-gpu", "list"}, []string{"yescrypt_crack.bin", "-gpu=list"}}, + } + + for _, tc := range tests { + got := normalizeGPUArgs(tc.in) + if strings.Join(got, "|") != strings.Join(tc.want, "|") { + t.Fatalf("normalizeGPUArgs(%q) = %q, want %q", tc.in, got, tc.want) + } + } +} + +func TestCountWordlistLines(t *testing.T) { + path := filepath.Join(t.TempDir(), "words.txt") + if err := os.WriteFile(path, []byte("one\ntwo\nthree"), 0600); err != nil { + t.Fatal(err) + } + + lines, err := countWordlistLines(path) + if err != nil { + t.Fatal(err) + } + if lines != 3 { + t.Fatalf("expected 3 wordlist lines, got %d", lines) + } +} + +func TestEstimateRemainingTime(t *testing.T) { + if got := estimateRemainingTime(500, 1000, 100); got != "01m" { + t.Fatalf("expected sub-minute ETA to round up to 01m, got %q", got) + } + if got := estimateRemainingTime(1000, 1000, 100); got != "00m" { + t.Fatalf("expected completed ETA, got %q", got) + } +} + +func TestFormatDurationMinutes(t *testing.T) { + tests := []struct { + d time.Duration + want string + }{ + {10 * time.Second, "00m"}, + {time.Minute, "01m"}, + {3*time.Hour + 7*time.Minute, "03h:07m"}, + {156*24*time.Hour + 7*time.Hour + 6*time.Minute, "156d:07h:06m"}, + } + for _, tc := range tests { + if got := formatDuration(tc.d); got != tc.want { + t.Fatalf("formatDuration(%v) = %q, want %q", tc.d, got, tc.want) + } + } +} diff --git a/kernels/yescrypt.cl b/kernels/yescrypt.cl new file mode 100644 index 0000000..cc98dcb --- /dev/null +++ b/kernels/yescrypt.cl @@ -0,0 +1,721 @@ +/* + * yescrypt_crack OpenCL yescrypt RW backend. + * + * The yescrypt algorithm follows Solar Designer's public yescrypt design and + * reference implementation. + * + * The GPU implementation in yescrypt.cl is adapted from + * hashcat's MIT-licensed yescrypt OpenCL implementation: + * + * https://github.com/hashcat/hashcat/blob/master/OpenCL/inc_hash_yescrypt.cl + * + * The surrounding kernel implementation, OpenCL host integration, batching, + * device management, and yescrypt_crack-specific runtime are part of + * yescrypt_crack. + * + * See THIRD_PARTY_NOTICES.md for upstream attribution and license information. + */ + +#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable + +#ifndef YC_N +#define YC_N 4096 +#endif +#ifndef YC_R +#define YC_R 32 +#endif + +#define YC_WG_SIZE 32 +#define YC_LOOP_CHUNK 2048U +#define YC_STATE_WORDS (16 * YC_R) +#define YC_STATE_BYTES (128 * YC_R) +#define YC_SCRATCH_WORDS ((YC_STATE_WORDS < 32) ? 32 : YC_STATE_WORDS) +#define YC_NLOOP ((((YC_N + 2U) / 3U) + 1U) & ~1U) +#define YC_PREHASH_N (YC_N / 64U) +#define YC_PREHASH_NEEDED ((YC_N >= 256U) && ((YC_N * YC_R) >= 0x20000U)) + +#define MAX_PW 256 +#define MAX_R 32 +#define B_STRIDE (128 * MAX_R) +#define XY_STRIDE_WORDS (32 * MAX_R) +#define S_WORDS 1536 +#define S_THIRD_WORDS 512 +#define PWX_SIMPLE 2 +#define PWX_GATHER 4 +#define PWX_ROUNDS 6 +#define S_MASK 4080U + +__constant uint K256[64] = { + 0x428a2f98U,0x71374491U,0xb5c0fbcfU,0xe9b5dba5U,0x3956c25bU,0x59f111f1U,0x923f82a4U,0xab1c5ed5U, + 0xd807aa98U,0x12835b01U,0x243185beU,0x550c7dc3U,0x72be5d74U,0x80deb1feU,0x9bdc06a7U,0xc19bf174U, + 0xe49b69c1U,0xefbe4786U,0x0fc19dc6U,0x240ca1ccU,0x2de92c6fU,0x4a7484aaU,0x5cb0a9dcU,0x76f988daU, + 0x983e5152U,0xa831c66dU,0xb00327c8U,0xbf597fc7U,0xc6e00bf3U,0xd5a79147U,0x06ca6351U,0x14292967U, + 0x27b70a85U,0x2e1b2138U,0x4d2c6dfcU,0x53380d13U,0x650a7354U,0x766a0abbU,0x81c2c92eU,0x92722c85U, + 0xa2bfe8a1U,0xa81a664bU,0xc24b8b70U,0xc76c51a3U,0xd192e819U,0xd6990624U,0xf40e3585U,0x106aa070U, + 0x19a4c116U,0x1e376c08U,0x2748774cU,0x34b0bcb5U,0x391c0cb3U,0x4ed8aa4aU,0x5b9cca4fU,0x682e6ff3U, + 0x748f82eeU,0x78a5636fU,0x84c87814U,0x8cc70208U,0x90befffaU,0xa4506cebU,0xbef9a3f7U,0xc67178f2U +}; + +typedef struct { + uint h[8]; + ulong total; + uchar buf[64]; + uint used; +} sha256_ctx_t; + +typedef struct { + sha256_ctx_t inner; + sha256_ctx_t outer; +} hmac256_base_t; + +inline uint rotr32(const uint x, const uint n) { return (x >> n) | (x << (32U - n)); } +inline uint load_be32_p(__private const uchar *p) { + return ((uint)p[0] << 24) | ((uint)p[1] << 16) | ((uint)p[2] << 8) | (uint)p[3]; +} +inline uint load_le32_g(__global const uchar *p) { + return (uint)p[0] | ((uint)p[1] << 8) | ((uint)p[2] << 16) | ((uint)p[3] << 24); +} +inline void store_le32_g(__global uchar *p, const uint v) { + p[0]=(uchar)v; p[1]=(uchar)(v>>8); p[2]=(uchar)(v>>16); p[3]=(uchar)(v>>24); +} +inline void store_be32_p(__private uchar *p, const uint v) { + p[0]=(uchar)(v>>24); p[1]=(uchar)(v>>16); p[2]=(uchar)(v>>8); p[3]=(uchar)v; +} + +inline void sha256_init(__private sha256_ctx_t *c) { + c->h[0]=0x6a09e667U; c->h[1]=0xbb67ae85U; c->h[2]=0x3c6ef372U; c->h[3]=0xa54ff53aU; + c->h[4]=0x510e527fU; c->h[5]=0x9b05688cU; c->h[6]=0x1f83d9abU; c->h[7]=0x5be0cd19U; + c->total=0; c->used=0; +} + +inline void sha256_transform(__private sha256_ctx_t *c, __private const uchar block[64]) { + uint w[16]; + for (uint i=0;i<16;i++) w[i]=load_be32_p(block + 4*i); + uint a=c->h[0], b=c->h[1], cc=c->h[2], d=c->h[3], e=c->h[4], f=c->h[5], g=c->h[6], h=c->h[7]; + for (uint i=0;i<64;i++) { + uint wi; + if (i < 16) { + wi = w[i]; + } else { + uint x=w[(i-15)&15], y=w[(i-2)&15]; + uint s0=rotr32(x,7)^rotr32(x,18)^(x>>3); + uint s1=rotr32(y,17)^rotr32(y,19)^(y>>10); + wi = w[i&15] = w[(i-16)&15] + s0 + w[(i-7)&15] + s1; + } + uint S1=rotr32(e,6)^rotr32(e,11)^rotr32(e,25); + uint ch=(e&f)^((~e)&g); + uint t1=h+S1+ch+K256[i]+wi; + uint S0=rotr32(a,2)^rotr32(a,13)^rotr32(a,22); + uint maj=(a&b)^(a&cc)^(b&cc); + uint t2=S0+maj; + h=g; g=f; f=e; e=d+t1; d=cc; cc=b; b=a; a=t1+t2; + } + c->h[0]+=a; c->h[1]+=b; c->h[2]+=cc; c->h[3]+=d; + c->h[4]+=e; c->h[5]+=f; c->h[6]+=g; c->h[7]+=h; +} + +inline void sha256_update_byte(__private sha256_ctx_t *c, const uchar x) { + c->buf[c->used++] = x; + c->total++; + if (c->used == 64) { sha256_transform(c,c->buf); c->used=0; } +} +inline void sha256_update_private(__private sha256_ctx_t *c, __private const uchar *p, const uint n) { + for (uint i=0;itotal*8UL; + c->buf[c->used++]=0x80; + if (c->used > 56) { + while (c->used<64) c->buf[c->used++]=0; + sha256_transform(c,c->buf); c->used=0; + } + while (c->used<56) c->buf[c->used++]=0; + for (uint i=0;i<8;i++) c->buf[56+i]=(uchar)(bits>>(56-8*i)); + sha256_transform(c,c->buf); + for (uint i=0;i<8;i++) store_be32_p(out+4*i,c->h[i]); +} + +inline void sha256_private_msg(__private const uchar *msg, const uint len, __private uchar out[32]) { + sha256_ctx_t c; sha256_init(&c); sha256_update_private(&c,msg,len); sha256_final(&c,out); +} +inline void sha256_global_msg(__global const uchar *msg, const uint len, __private uchar out[32]) { + sha256_ctx_t c; sha256_init(&c); sha256_update_global(&c,msg,len); sha256_final(&c,out); +} + +inline void hmac_init_private_key(__private const uchar *key, uint keylen, __private hmac256_base_t *base) { + uchar k0[64]; for (uint i=0;i<64;i++) k0[i]=0; + if (keylen > 64) { + uchar kh[32]; sha256_private_msg(key,keylen,kh); for (uint i=0;i<32;i++) k0[i]=kh[i]; + } else { + for (uint i=0;iinner); sha256_update_private(&base->inner,pad,64); + for (uint i=0;i<64;i++) pad[i]=k0[i]^0x5c; + sha256_init(&base->outer); sha256_update_private(&base->outer,pad,64); +} + +inline void hmac_init_global_key(__global const uchar *key, uint keylen, __private hmac256_base_t *base) { + uchar k0[64]; for (uint i=0;i<64;i++) k0[i]=0; + if (keylen > 64) { + uchar kh[32]; sha256_global_msg(key,keylen,kh); for (uint i=0;i<32;i++) k0[i]=kh[i]; + } else { + for (uint i=0;iinner); sha256_update_private(&base->inner,pad,64); + for (uint i=0;i<64;i++) pad[i]=k0[i]^0x5c; + sha256_init(&base->outer); sha256_update_private(&base->outer,pad,64); +} + +inline void hmac_finish(__private const hmac256_base_t *base, __private sha256_ctx_t *inner, __private uchar out[32]) { + uchar ih[32]; sha256_final(inner,ih); + sha256_ctx_t outer=base->outer; sha256_update_private(&outer,ih,32); sha256_final(&outer,out); +} + +inline void hmac_private_private(__private const uchar *key, uint keylen, __private const uchar *msg, uint msglen, __private uchar out[32]) { + hmac256_base_t b; hmac_init_private_key(key,keylen,&b); + sha256_ctx_t in=b.inner; sha256_update_private(&in,msg,msglen); hmac_finish(&b,&in,out); +} +inline void hmac_private_global(__private const uchar *key, uint keylen, __global const uchar *msg, uint msglen, __private uchar out[32]) { + hmac256_base_t b; hmac_init_private_key(key,keylen,&b); + sha256_ctx_t in=b.inner; sha256_update_global(&in,msg,msglen); hmac_finish(&b,&in,out); +} +inline void hmac_global_private(__global const uchar *key, uint keylen, __private const uchar *msg, uint msglen, __private uchar out[32]) { + hmac256_base_t b; hmac_init_global_key(key,keylen,&b); + sha256_ctx_t in=b.inner; sha256_update_private(&in,msg,msglen); hmac_finish(&b,&in,out); +} + +inline void pbkdf2_fill_b(__private const uchar key[32], __global const uchar *salt, uint saltlen, __global uchar *B, uint outlen) { + hmac256_base_t base; hmac_init_private_key(key,32,&base); + uint pos=0; + for (uint blk=1; pos>32), x2=(uint)d5, x3=(uint)(d3>>32); + uint x4=(uint)d2, x5=(uint)(d0>>32), x6=(uint)d7, x7=(uint)(d5>>32); + uint x8=(uint)d4, x9=(uint)(d2>>32), x10=(uint)d1, x11=(uint)(d7>>32); + uint x12=(uint)d6, x13=(uint)(d4>>32), x14=(uint)d3, x15=(uint)(d1>>32); + for (uint i=0;i>32)+x5)<<32); + d1=(ulong)((uint)d1+x10)|((ulong)((uint)(d1>>32)+x15)<<32); + d2=(ulong)((uint)d2+x4)|((ulong)((uint)(d2>>32)+x9)<<32); + d3=(ulong)((uint)d3+x14)|((ulong)((uint)(d3>>32)+x3)<<32); + d4=(ulong)((uint)d4+x8)|((ulong)((uint)(d4>>32)+x13)<<32); + d5=(ulong)((uint)d5+x2)|((ulong)((uint)(d5>>32)+x7)<<32); + d6=(ulong)((uint)d6+x12)|((ulong)((uint)(d6>>32)+x1)<<32); + d7=(ulong)((uint)d7+x6)|((ulong)((uint)(d7>>32)+x11)<<32); + out[0]=tmp[0]=d0; out[1]=tmp[1]=d1; out[2]=tmp[2]=d2; out[3]=tmp[3]=d3; + out[4]=tmp[4]=d4; out[5]=tmp[5]=d5; out[6]=tmp[6]=d6; out[7]=tmp[7]=d7; +} + +inline void block_mix_classic(__global ulong *in, __global ulong *out, uint r) { + ulong tmp[8]; + for (uint j=0;j<8;j++) tmp[j]=in[(2*r-1)*8+j]; + for (uint i=0;i<2*r;i+=2) { + salsa_xor(tmp,in+i*8,out+i*4,8); + salsa_xor(tmp,in+i*8+8,out+i*4+r*8,8); + } +} + +inline uint integerify(__global const ulong *x, uint r) { return (uint)x[(2*r-1)*8]; } +inline uint p2floor_u(uint x) { while ((x&(x-1U))!=0U) x&=x-1U; return x; } +inline uint wrap_u(uint x,uint i) { uint n=p2floor_u(i); return (x&(n-1U))+(i-n); } + +inline void load_b_to_x(__global const uchar *B, uint r, __global ulong *X) { + uint R=16*r, j=0; + for (uint i=0;i>32)); j+=4; + } +} + +inline void sbox_init(__global uchar *B, __global ulong *S, __global ulong *XY) { + __global ulong *x=XY; + __global ulong *y=XY+16; + load_b_to_x(B,1,x); + for (uint i=0;i<96;i+=2) { + copy_words(S+i*16,x,16); + block_mix_classic(x,y,1); + copy_words(S+(i+1)*16,y,16); + block_mix_classic(y,x,1); + } + store_x_to_b(x,1,B); +} + +inline void pwxform(__private ulong X[8], __global ulong *S, __private uint *s0off, __private uint *s1off, __private uint *s2off, __private uint *wptr) { + uint s0=*s0off, s1=*s1off, s2=*s2off, w=*wptr; + for (uint i=0;i>32); + x=(ulong)xh*(ulong)xl; + xl=(xl&S_MASK)>>3; xh=(xh&S_MASK)>>3; + x=(x+S[s0+xl])^S[s1+xh]; X[2*j]=x; + ulong y=X[2*j+1]; + y=((y>>32)*(ulong)((uint)y)+S[s0+xl+1])^S[s1+xh+1]; X[2*j+1]=y; + if (i!=0 && i!=PWX_ROUNDS-1) { S[s2+w]=x; S[s2+w+1]=y; w+=2; } + } + } + *s0off=s2; *s1off=s0; *s2off=s1; *wptr=w&511U; +} + +inline void block_mix_pwx(__global ulong *B, uint r, __global ulong *S, __private uint *s0, __private uint *s1, __private uint *s2, __private uint *w) { + ulong X[8]; uint r1=2*r; + for (uint j=0;j<8;j++) X[j]=B[(r1-1)*8+j]; + for (uint i=0;i1) { + uint j=wrap_u(integerify(x,r),i); + xor_words(x,V+(ulong)j*R,R); + } + block_mix_pwx(x,r,S,&s0,&s1,&s2,&w); + } + uint nloop=((N+2U)/3U+1U)&~1U; + for (uint i=0;i>32), x2=(uint)d5, x3=(uint)(d3>>32); + uint x4=(uint)d2, x5=(uint)(d0>>32), x6=(uint)d7, x7=(uint)(d5>>32); + uint x8=(uint)d4, x9=(uint)(d2>>32), x10=(uint)d1, x11=(uint)(d7>>32); + uint x12=(uint)d6, x13=(uint)(d4>>32), x14=(uint)d3, x15=(uint)(d1>>32); + for (uint i=0;i>32)+x5)<<32); + v[1]=(ulong)((uint)d1+x10)|((ulong)((uint)(d1>>32)+x15)<<32); + v[2]=(ulong)((uint)d2+x4)|((ulong)((uint)(d2>>32)+x9)<<32); + v[3]=(ulong)((uint)d3+x14)|((ulong)((uint)(d3>>32)+x3)<<32); + v[4]=(ulong)((uint)d4+x8)|((ulong)((uint)(d4>>32)+x13)<<32); + v[5]=(ulong)((uint)d5+x2)|((ulong)((uint)(d5>>32)+x7)<<32); + v[6]=(ulong)((uint)d6+x12)|((ulong)((uint)(d6>>32)+x1)<<32); + v[7]=(ulong)((uint)d7+x6)|((ulong)((uint)(d7>>32)+x11)<<32); +} + +/* + * Cooperative PWX blockmix. Four lanes own the four PWX gather lanes, two + * 64-bit words apiece. The other lanes participate in state/V transfers. + * S2 writes are independent during a PWX call; the barrier after each 64-byte + * segment makes them visible before the S-box thirds rotate for the next one. + */ +inline void coop_block_mix_pwx( + __local ulong *X, + __local ulong *S, + __private uint *s_state, + __private uint *w_ptr, + const uint lid) +{ + const uint r1 = 2U * YC_R; + const int do_lane = (lid < PWX_GATHER); + + ulong px0=0, px1=0; + if (do_lane) { + const uint last=(r1-1U)*8U + lid*2U; + px0=X[last+0U]; + px1=X[last+1U]; + } + + uint ss=*s_state; + uint ww=*w_ptr; + + for (uint i=0;i>32); + const uint p0=(xl&S_MASK)>>3; + const uint p1=(xh&S_MASK)>>3; + + /* p0/p1 are even, so each vload2 reads one aligned 16-byte S entry pair. */ + const ulong2 s0v=vload2(0,S+s0off+p0); + const ulong2 s1v=vload2(0,S+s1off+p1); + + ulong x=((ulong)((uint)(px0>>32))*(ulong)((uint)px0) + s0v.s0) ^ s1v.s0; + ulong y=((ulong)((uint)(px1>>32))*(ulong)((uint)px1) + s0v.s1) ^ s1v.s1; + px0=x; + px1=y; + + if (round!=0U && round!=(PWX_ROUNDS-1U)) { + const uint ai=round-1U; + const uint wslot=ww + ai*8U + lid*2U; + S[s2off+wslot+0U]=px0; + S[s2off+wslot+1U]=px1; + } + } + + X[base+0U]=px0; + X[base+1U]=px1; + } + + ss=(ss+1U)%3U; + ww=(ww+32U)&511U; + barrier(CLK_LOCAL_MEM_FENCE); + } + + if (lid==0U) { + ulong salsa[8]; + const uint off=(r1-1U)*8U; + for (uint i=0;i<8U;i++) salsa[i]=X[off+i]; + salsa_self_private(salsa,2U); + for (uint i=0;i<8U;i++) X[off+i]=salsa[i]; + } + barrier(CLK_LOCAL_MEM_FENCE); + + *s_state=ss; + *w_ptr=ww; +} + +inline void coop_smix1_step( + __local ulong *X, + __global ulong *V, + __local ulong *S, + __private uint *s_state, + __private uint *w_ptr, + const uint iter, + const uint lid, + const uint lsz) +{ + __global ulong *Vi=V+(ulong)iter*YC_STATE_WORDS; + for (uint j=lid;j1U) { + const uint idx=wrap_u(integerify_local(X),iter); + __global ulong *Vj=V+(ulong)idx*YC_STATE_WORDS; + /* Every lane must derive idx from the same pre-XOR state. */ + barrier(CLK_LOCAL_MEM_FENCE); + for (uint j=lid;j=count) return; + const uint pwlen=pw_lens[gid]; + if (pwlen>MAX_PW) return; + + __global const uchar *pw=passwords+(ulong)gid*MAX_PW; + __global ulong *P=P_all+(ulong)gid*YC_STATE_WORDS; + __global uchar *B=(__global uchar *)P; + __global ulong *scratch=scratch_all+(ulong)gid*YC_SCRATCH_WORDS; + __global ulong *S=S_all+(ulong)gid*S_WORDS; + __global yescrypt_gpu_ctx_t *ctx=&ctx_all[gid]; + + const ulong vwords=(ulong)YC_STATE_WORDS*(ulong)YC_N; + const ulong slot=(ulong)(gid>>2); + __global ulong *V; + switch (gid&3U) { + case 0U: V=V0+slot*vwords; break; + case 1U: V=V1+slot*vwords; break; + case 2U: V=V2+slot*vwords; break; + default: V=V3+slot*vwords; break; + } + + uchar ppass[32]; + uchar key[32]; + +#if YC_PREHASH_NEEDED + { + uchar prekey[16]={'y','e','s','c','r','y','p','t','-','p','r','e','h','a','s','h'}; + hmac256_base_t hb; hmac_init_private_key(prekey,16,&hb); + sha256_ctx_t hin=hb.inner; sha256_update_global(&hin,pw,pwlen); hmac_finish(&hb,&hin,ppass); + + pbkdf2_fill_b(ppass,salt,salt_len,B,YC_STATE_BYTES); + for (uint i=0;i<32U;i++) ppass[i]=B[i]; + smix_yescrypt(B,YC_R,YC_PREHASH_N,V,scratch,S,ppass); + pbkdf2_b_to_32(ppass,B,YC_STATE_BYTES,key); + for (uint i=0;i<32U;i++) ppass[i]=key[i]; + } +#endif + + { + uchar yeskey[8]={'y','e','s','c','r','y','p','t'}; + hmac256_base_t hb; hmac_init_private_key(yeskey,8,&hb); + sha256_ctx_t hin=hb.inner; +#if YC_PREHASH_NEEDED + sha256_update_private(&hin,ppass,32); +#else + sha256_update_global(&hin,pw,pwlen); +#endif + hmac_finish(&hb,&hin,ppass); + } + + /* Main-pass KDF setup, but stop before the expensive SMix. */ + pbkdf2_fill_b(ppass,salt,salt_len,B,YC_STATE_BYTES); + for (uint i=0;i<32U;i++) ppass[i]=B[i]; + sbox_init(B,S,scratch); + { + uchar next[32]; + hmac_global_private(B+64U*(2U*YC_R-1U),64,ppass,32,next); + for (uint i=0;i<32U;i++) ppass[i]=next[i]; + } + load_b_to_x(B,YC_R,scratch); + for (uint i=0;ipasswd[i]=ppass[i]; + ctx->phase=0U; + ctx->iter=0U; + ctx->s_state=0U; + ctx->w=0U; +} + +/* + * Stage 2. One 32-work-item workgroup owns one password. X and the entire + * 12 KiB S-box stay in local/shared memory for the duration of the launch. + */ +__attribute__((reqd_work_group_size(YC_WG_SIZE,1,1))) +__kernel void yescrypt_loop( + const uint count, + const uint loop_count, + __global ulong *P_all, + __global ulong *S_all, + __global yescrypt_gpu_ctx_t *ctx_all, + __global ulong *V0, + __global ulong *V1, + __global ulong *V2, + __global ulong *V3) +{ + const uint bid=get_group_id(0); + const uint lid=get_local_id(0); + const uint lsz=get_local_size(0); + if (bid>=count) return; + + __global ulong *P=P_all+(ulong)bid*YC_STATE_WORDS; + __global ulong *Sg=S_all+(ulong)bid*S_WORDS; + __global yescrypt_gpu_ctx_t *ctx=&ctx_all[bid]; + + __local ulong X[YC_STATE_WORDS]; + __local ulong S[S_WORDS]; + + for (uint i=lid;iphase; + uint iter=ctx->iter; + uint s_state=ctx->s_state; + uint w=ctx->w; + + const ulong vwords=(ulong)YC_STATE_WORDS*(ulong)YC_N; + const ulong slot=(ulong)(bid>>2); + __global ulong *V; + switch (bid&3U) { + case 0U: V=V0+slot*vwords; break; + case 1U: V=V1+slot*vwords; break; + case 2U: V=V2+slot*vwords; break; + default: V=V3+slot*vwords; break; + } + + barrier(CLK_LOCAL_MEM_FENCE); + + for (uint loop=0;loop=YC_N) { phase=1U; iter=0U; } + } + } else { + if (iter>=YC_NLOOP) break; + coop_smix2_step(X,V,S,&s_state,&w,lid,lsz); + iter++; + } + } + + barrier(CLK_LOCAL_MEM_FENCE); + for (uint i=lid;iphase=phase; + ctx->iter=iter; + ctx->s_state=s_state; + ctx->w=w; + } +} + +/* Stage 3. One ordinary work-item derives and writes the final 32-byte key. */ +__kernel void yescrypt_final( + const uint count, + __global ulong *P_all, + __global ulong *scratch_all, + __global yescrypt_gpu_ctx_t *ctx_all, + __global uchar *out_all) +{ + const uint gid=get_global_id(0); + if (gid>=count) return; + + __global ulong *P=P_all+(ulong)gid*YC_STATE_WORDS; + __global ulong *scratch=scratch_all+(ulong)gid*YC_SCRATCH_WORDS; + __global uchar *B=(__global uchar *)P; + __global yescrypt_gpu_ctx_t *ctx=&ctx_all[gid]; + + for (uint i=0;ipasswd[i]; + pbkdf2_b_to_32(ppass,B,YC_STATE_BYTES,key); + + uchar client[10]={'C','l','i','e','n','t',' ','K','e','y'}; + uchar h1[32], final[32]; + hmac_private_private(key,32,client,10,h1); + sha256_private_msg(h1,32,final); + for (uint i=0;i<32U;i++) out_all[(ulong)gid*32U+i]=final[i]; +} diff --git a/main.go b/main.go index 2886397..d27e35c 100644 --- a/main.go +++ b/main.go @@ -2,41 +2,48 @@ package main import ( "encoding/base64" + "errors" "flag" "fmt" + "log" "os" "runtime" + "strings" "sync" "time" ) /* Cyclone's Yescrypt Cracker -POC tool to crack Yescrypt hashes +POC tool to crack yescrypt /gost-yescrypt hashes https://github.com/cyclone-github/yescrypt_crack -written by cyclone in pure Go - GNU General Public License v2.0 https://github.com/cyclone-github/yescrypt_crack/blob/main/LICENSE Credits: -The yescrypt algo was written by Solar Designer: https://www.openwall.com/yescrypt/ -This tool uses their yescrypt-go implementation: https://github.com/openwall/yescrypt-go +yescrypt_crack was written by cyclone in Go +The yescrypt algorithm was designed by Solar Designer: https://www.openwall.com/yescrypt/ +The CPU yescrypt implementation in yescrypt_cpu.go is adapted from openwall/yescrypt-go: https://github.com/openwall/yescrypt-go +The GPU implementation in yescrypt.cl is adapted from hashcat's yescrypt OpenCL implementation: https://github.com/hashcat/hashcat/blob/master/OpenCL/inc_hash_yescrypt.cl +Streebog-256 is provided by github.com/tarantool/go-gostcrypto/streebog +See THIRD_PARTY_NOTICES.md for upstream copyright notices and license terms version history -v0.1.0; 2024-04-16 - initial POC -... -v0.2.0; 2025-03-06 - refactored code - github version -v0.3.1; 2026-08-08 - add gost-yescrypt support +v0.4.1-dev; 2026-08-22 + complete rewrite of codebase + add OpenCL GPU acceleration for yescrypt / gost-yescrypt with CPU fallback + add ETA status for -w wordlists */ +type wordlistCountResult struct { + lines uint64 + err error +} + // main func func main() { + os.Args = normalizeGPUArgs(os.Args) wordlistFileFlag := flag.String("w", "", "Input file to process (omit -w to read from stdin)") hashFileFlag := flag.String("h", "", "Yescrypt / gost-yescrypt hash file") outputFileFlag := flag.String("o", "", "Output file to write cracked hashes to (omit -o to print to console)") @@ -45,9 +52,26 @@ func main() { helpFlag := flag.Bool("help", false, "Prints help:") threadFlag := flag.Int("t", runtime.NumCPU(), "CPU threads to use (optional)") statsIntervalFlag := flag.Int("s", 60, "Interval in seconds for printing stats. Defaults to 60.") + gpuFlag := flag.String("gpu", "all", "GPU selection: all, 0, 0,1, or list") + cpuFlag := flag.Bool("cpu", false, "Force CPU mode") + gpuBatchFlag := flag.Int("b", 0, "GPU batch size (0 = auto)") flag.Parse() - clearScreen() + var ( + gpuSet bool + threadSet bool + batchSet bool + ) + flag.Visit(func(f *flag.Flag) { + switch f.Name { + case "gpu": + gpuSet = true + case "t": + threadSet = true + case "b": + batchSet = true + } + }) // run sanity checks for special flags if *versionFlag { @@ -65,25 +89,49 @@ func main() { os.Exit(0) } + if *cpuFlag && gpuSet { + fmt.Fprintln(os.Stderr, "-cpu and -gpu cannot be used together") + os.Exit(1) + } + if gpuSet && threadSet { + fmt.Fprintln(os.Stderr, "-gpu and -t cannot be used together (-t is for CPU mode)") + os.Exit(1) + } + if (*cpuFlag || threadSet) && batchSet { + fmt.Fprintln(os.Stderr, "-b is only valid in GPU mode") + os.Exit(1) + } + if *gpuBatchFlag < 0 { + fmt.Fprintln(os.Stderr, "-b must be >= 0") + os.Exit(1) + } + if *statsIntervalFlag < 0 { + fmt.Fprintln(os.Stderr, "-s must be >= 0") + os.Exit(1) + } + + if gpuSet && strings.EqualFold(strings.TrimSpace(*gpuFlag), "list") { + printGPUList() + os.Exit(0) + } + if *hashFileFlag == "" { fmt.Fprintln(os.Stderr, "-h (hash file) flag is required") fmt.Fprintln(os.Stderr, "Try running with -help for usage instructions") os.Exit(1) } - startTime := time.Now() - + cpuMode := *cpuFlag || threadSet numThreads := setNumThreads(*threadFlag) var ( crackedCount int32 - linesProcessed int32 + linesProcessed uint64 wg sync.WaitGroup - totalHashesGenerated int32 + totalHashesGenerated uint64 ) stopChan := make(chan struct{}) - handleGracefulShutdown(stopChan) hashes, err := ReadYescryptHashes(*hashFileFlag) @@ -92,17 +140,128 @@ func main() { os.Exit(1) } totalHashCount := len(hashes) + if totalHashCount == 0 { + fmt.Fprintln(os.Stderr, "No supported yescrypt / gost-yescrypt hashes found") + os.Exit(1) + } + + groups := makeHashGroups(hashes) + var selectedGPUs []GPUDeviceInfo + var fallbackReason string + displayBatch := 0 + + if !cpuMode { + devices, derr := ListOpenCLGPUs() + if derr != nil { + cpuMode = true + fallbackReason = fmt.Sprintf("OpenCL unavailable (%v); falling back to CPU", derr) + } else { + selectedGPUs, err = parseGPUSelection(*gpuFlag, devices) + if err != nil { + fmt.Fprintln(os.Stderr, err) + os.Exit(1) + } + if len(selectedGPUs) == 0 { + cpuMode = true + fallbackReason = "No OpenCL GPU devices found; falling back to CPU" + } else if len(groups) == 0 { + cpuMode = true + fallbackReason = "No GPU-compatible yescrypt groups found; falling back to CPU" + } else { + // RTX 4090s auto-tuned measured value was 1280 + for _, info := range selectedGPUs { + cap := minGPUCapacity(info, groups, *gpuBatchFlag) + if cap > 0 && (displayBatch == 0 || cap < displayBatch) { + displayBatch = cap + } + } + } + } + } + + welcome := WelcomeOptions{Backend: "CPU"} + if !cpuMode { + welcome.Backend = "GPU" + welcome.GPUSelection = *gpuFlag + welcome.GPUBatch = displayBatch + } + printWelcomeScreen(hashFileFlag, wordlistFileFlag, totalHashCount, numThreads, welcome) + if fallbackReason != "" { + log.Println(fallbackReason) + } + + var countCh chan wordlistCountResult + if *wordlistFileFlag != "" { + log.Println("Counting wordlist lines...") + countCh = make(chan wordlistCountResult, 1) + go func() { + lines, err := countWordlistLinesStop(*wordlistFileFlag, stopChan) + countCh <- wordlistCountResult{lines: lines, err: err} + }() + } + + var gpuPrepCh chan *gpuRuntime + if !cpuMode { + log.Println("Tuning GPU...") + gpuPrepCh = make(chan *gpuRuntime, 1) + go func() { + gpuPrepCh <- prepareGPUProcessor(selectedGPUs, *gpuBatchFlag, groups, hashes, stopChan) + }() + } - printWelcomeScreen(hashFileFlag, wordlistFileFlag, totalHashCount, numThreads) + var totalWordlistLines uint64 + if countCh != nil { + result := <-countCh + if result.err == nil { + totalWordlistLines = result.lines + } else if !errors.Is(result.err, errWordlistCountCancelled) { + log.Printf("Unable to count wordlist lines for ETA: %v", result.err) + } + } + + var preparedGPU *gpuRuntime + if gpuPrepCh != nil { + preparedGPU = <-gpuPrepCh + if preparedGPU == nil { + cpuMode = true + log.Println("GPU tuning failed; falling back to CPU") + } + } + + select { + case <-stopChan: + if preparedGPU != nil { + preparedGPU.Close() + } + return + default: + } + + startTime := time.Now() + log.Println("Working...") wg.Add(1) - go monitorPrintStats(&crackedCount, &linesProcessed, &totalHashesGenerated, stopChan, startTime, totalHashCount, &wg, *statsIntervalFlag) + go monitorPrintStats(&crackedCount, &linesProcessed, &totalHashesGenerated, stopChan, startTime, totalHashCount, totalWordlistLines, &wg, *statsIntervalFlag) - startProc(*wordlistFileFlag, *outputFileFlag, numThreads, hashes, &crackedCount, &linesProcessed, &totalHashesGenerated, stopChan) + procErr := startProcWithOptions( + *wordlistFileFlag, + *outputFileFlag, + numThreads, + hashes, + &crackedCount, + &linesProcessed, + &totalHashesGenerated, + stopChan, + RuntimeOptions{UseGPU: !cpuMode, GPUList: *gpuFlag, GPUBatch: *gpuBatchFlag, PreparedGPU: preparedGPU}, + ) closeStopChannel(stopChan) - wg.Wait() + if procErr != nil { + log.Printf("%v", procErr) + log.Println("GPU processing aborted. Re-run with -cpu to use the CPU backend.") + os.Exit(1) + } } // end code diff --git a/print_welcome.go b/print_welcome.go index d668aa7..49dc037 100644 --- a/print_welcome.go +++ b/print_welcome.go @@ -2,13 +2,12 @@ package main import ( "fmt" - "log" "os" ) // version func func versionFunc() { - fmt.Fprintln(os.Stderr, "Cyclone's Yescrypt Cracker v0.3.1; 2026-08-08\nhttps://github.com/cyclone-github/yescrypt_crack\n ") + fmt.Fprintln(os.Stderr, "Cyclone's Yescrypt Cracker v0.4.1-dev; 2026-08-22\nhttps://github.com/cyclone-github/yescrypt_crack ") } // help func @@ -19,24 +18,37 @@ func helpFunc() { -w {wordlist} (omit -w to read from stdin) -h {yescrypt_hash_file} -o {output} (omit -o to write to stdout) --t {cpu threads} +-t {cpu threads} (selects CPU mode) -s {print status every nth sec) +-gpu [all|0,1|list] (default: all GPUs; omit value to use all GPUs) +-cpu (force CPU mode) +-b {gpu batch size} (optional; 0 = auto) -version (version info) -help (usage instructions) -./yescrypt_crack.bin -h {yescrypt_hash_file} -w {wordlist} -o {output} -t {cpu threads} -s {print status every nth sec} +./yescrypt_crack.bin -h {yescrypt_hash_file} -w {wordlist} -o {output} -gpu -s {print status every nth sec} -./yescrypt_crack.bin -h hashes.txt -w wordlist.txt -o cracked.txt -t 16 -s 10 +./yescrypt_crack.bin -h hashes.txt -w wordlist.txt -o cracked.txt -gpu 0,1 -s 10 + +./yescrypt_crack.bin -h hashes.txt -w wordlist.txt -gpu 0,1 -b 1280 -s 10 + +./yescrypt_crack.bin -h hashes.txt -w wordlist.txt -cpu -t 16 -s 10 cat wordlist | ./yescrypt_crack.bin -h hashes.txt -./yescrypt_crack.bin -h hashes.txt -w wordlist.txt -o output.txt` +./yescrypt_crack.bin -gpu list` fmt.Fprintln(os.Stderr, str) } +type WelcomeOptions struct { + Backend string + GPUSelection string + GPUBatch int +} + // print welcome screen -func printWelcomeScreen(hashFileFlag, wordlistFileFlag *string, totalHashCount, numThreads int) { +func printWelcomeScreen(hashFileFlag, wordlistFileFlag *string, totalHashCount, numThreads int, opts WelcomeOptions) { fmt.Fprintln(os.Stderr, " -------------------------------------------------- ") fmt.Fprintln(os.Stderr, "| Cyclone's Yescrypt Cracker |") fmt.Fprintln(os.Stderr, "| https://github.com/cyclone-github/yescrypt_crack |") @@ -52,5 +64,29 @@ func printWelcomeScreen(hashFileFlag, wordlistFileFlag *string, totalHashCount, fmt.Fprintf(os.Stderr, "Wordlist:\t%s\n", *wordlistFileFlag) } - log.Println("Working...") + fmt.Fprintf(os.Stderr, "Backend:\t%s\n", opts.Backend) + if opts.Backend == "GPU" { + fmt.Fprintf(os.Stderr, "GPU selection:\t%s\n", opts.GPUSelection) + if opts.GPUBatch > 0 { + fmt.Fprintf(os.Stderr, "GPU batch:\t%d\n", opts.GPUBatch) + } else { + fmt.Fprintf(os.Stderr, "GPU batch:\tauto\n") + } + } +} + +func printGPUList() { + devices, err := ListOpenCLGPUs() + if err != nil { + fmt.Fprintln(os.Stderr, "OpenCL:", err) + return + } + if len(devices) == 0 { + fmt.Fprintln(os.Stderr, "No OpenCL GPU devices found") + return + } + + for _, device := range devices { + fmt.Printf("GPU %d: %s VRAM %.1f GiB\n", device.Index, device.Name, float64(device.GlobalMem)/(1<<30)) + } } diff --git a/process.go b/process.go index e1b94ce..7e3531b 100644 --- a/process.go +++ b/process.go @@ -2,107 +2,746 @@ package main import ( "bufio" + "bytes" + "crypto/subtle" + "errors" "fmt" "log" "os" + "sort" + "strconv" + "strings" "sync" "sync/atomic" ) -// process logic -func startProc(wordlistFileFlag string, outputPath string, numGoroutines int, hashes []YescryptHash, crackedCount *int32, linesProcessed *int32, totalHashesGenerated *int32, stopChan chan struct{}) { - var file *os.File - var err error +type RuntimeOptions struct { + UseGPU bool + GPUList string // "all" or comma-separated OpenCL GPU indexes + GPUBatch int // 0 = automatic + PreparedGPU *gpuRuntime +} - if wordlistFileFlag == "" { - file = os.Stdin - } else { - file, err = os.Open(wordlistFileFlag) - if err != nil { - log.Fatalf("Error opening file: %v\n", err) +type hashGroup struct { + N uint32 + R uint32 + Salt []byte + Indices []int +} + +func makeHashGroups(hashes []YescryptHash) []hashGroup { + type key struct { + N, R uint32 + Salt string + } + m := make(map[key]int) + var groups []hashGroup + for i := range hashes { + if !hashes[i].GPUOK { + continue } - defer file.Close() + k := key{hashes[i].N, hashes[i].R, string(hashes[i].Salt)} + if pos, ok := m[k]; ok { + groups[pos].Indices = append(groups[pos].Indices, i) + continue + } + m[k] = len(groups) + groups = append(groups, hashGroup{N: k.N, R: k.R, Salt: append([]byte(nil), hashes[i].Salt...), Indices: []int{i}}) } + sort.SliceStable(groups, func(i, j int) bool { + if groups[i].N != groups[j].N { + return groups[i].N < groups[j].N + } + if groups[i].R != groups[j].R { + return groups[i].R < groups[j].R + } + return bytes.Compare(groups[i].Salt, groups[j].Salt) < 0 + }) + return groups +} - var outputFile *os.File - if outputPath != "" { - outputFile, err = os.OpenFile(outputPath, os.O_APPEND|os.O_CREATE|os.O_WRONLY, 0644) - if err != nil { - log.Fatalf("Error opening output file: %v", err) +func groupHasUncracked(hashes []YescryptHash, g *hashGroup) bool { + for _, idx := range g.Indices { + if atomic.LoadInt32(&hashes[idx].Cracked) == 0 { + return true } - defer outputFile.Close() } + return false +} - var writer *bufio.Writer - if outputPath != "" { - writer = bufio.NewWriter(outputFile) +func parseGPUSelection(spec string, devices []GPUDeviceInfo) ([]GPUDeviceInfo, error) { + if len(devices) == 0 { + return nil, nil + } + spec = strings.TrimSpace(strings.ToLower(spec)) + if spec == "" || spec == "all" { + return append([]GPUDeviceInfo(nil), devices...), nil + } + seen := make(map[int]bool) + var out []GPUDeviceInfo + for _, field := range strings.Split(spec, ",") { + i, err := strconv.Atoi(strings.TrimSpace(field)) + if err != nil || i < 0 || i >= len(devices) { + return nil, fmt.Errorf("invalid -gpu value %q; available indexes are 0..%d", field, len(devices)-1) + } + if !seen[i] { + seen[i] = true + out = append(out, devices[i]) + } + } + return out, nil +} + +func tuneAutoGPUCapacity(cap uint64) uint64 { + if cap == 0 { + return 0 + } + + tuned := cap * 95 / 100 + if tuned >= 32 { + tuned = tuned / 32 * 32 + } + if tuned == 0 { + tuned = 1 + } + return tuned +} + +func estimateGPUCapacity(info GPUDeviceInfo, N, r uint32, batchHint int) int { + if N == 0 || r == 0 { + return 0 + } + vbytes := uint64(128) * uint64(N) * uint64(r) + stateBytes := uint64(128) * uint64(r) + scratchBytes := stateBytes + if scratchBytes < 256 { + scratchBytes = 256 + } + workspace := uint64(gpuMaxPasswordLen+4+64+32+48+1536*8) + stateBytes + scratchBytes + reserve := info.GlobalMem / 10 + const mib = uint64(1024 * 1024) + if reserve < 512*mib { + reserve = 512 * mib + } + if reserve >= info.GlobalMem { + reserve = info.GlobalMem / 5 + } + usable := info.GlobalMem - reserve + capMem := usable / (vbytes + workspace) + allocLimit := info.MaxAlloc * 95 / 100 + capAlloc := (allocLimit / vbytes) * 4 + cap := capMem + if cap > capAlloc { + cap = capAlloc + } + if cap > 4096 { + cap = 4096 + } + + if batchHint > 0 { + if cap > uint64(batchHint) { + cap = uint64(batchHint) + } } else { - writer = bufio.NewWriter(os.Stdout) + cap = tuneAutoGPUCapacity(cap) + } + + if cap >= 4 { + cap = cap / 4 * 4 + } + return int(cap) +} + +func minGPUCapacity(info GPUDeviceInfo, groups []hashGroup, batchHint int) int { + minCap := 0 + seen := make(map[[2]uint32]bool) + for i := range groups { + cfg := [2]uint32{groups[i].N, groups[i].R} + if seen[cfg] { + continue + } + seen[cfg] = true + cap := estimateGPUCapacity(info, cfg[0], cfg[1], batchHint) + if cap <= 0 { + return 0 + } + if minCap == 0 || cap < minCap { + minCap = cap + } + } + return minCap +} + +func startProcWithOptions(wordlistFileFlag string, outputPath string, numGoroutines int, hashes []YescryptHash, crackedCount *int32, linesProcessed *uint64, totalHashesGenerated *uint64, stopChan chan struct{}, opts RuntimeOptions) error { + file, err := openWordlist(wordlistFileFlag) + if err != nil { + log.Fatalf("Error opening wordlist: %v", err) + } + if file != os.Stdin { + defer file.Close() + } + writer, outputFile, err := openOutput(outputPath) + if err != nil { + log.Fatalf("Error opening output: %v", err) + } + if outputFile != nil { + defer outputFile.Close() } defer writer.Flush() - var ( - writerMu sync.Mutex - wg sync.WaitGroup - ) + var writerMu sync.Mutex + if opts.UseGPU { + if opts.PreparedGPU != nil { + usedGPU, gpuErr := runPreparedGPUProcessor(file, opts.PreparedGPU, &writerMu, writer, crackedCount, linesProcessed, totalHashesGenerated, stopChan) + if gpuErr != nil { + return gpuErr + } + if usedGPU { + log.Println("Finished") + return nil + } + log.Printf("No usable OpenCL yescrypt groups; falling back to CPU") + } else { + devices, derr := ListOpenCLGPUs() + if derr != nil { + log.Printf("OpenCL unavailable (%v); falling back to CPU", derr) + } else { + selected, serr := parseGPUSelection(opts.GPUList, devices) + if serr != nil { + log.Printf("%v; falling back to CPU", serr) + } else if len(selected) != 0 { + groups := makeHashGroups(hashes) + if len(groups) != 0 { + usedGPU, gpuErr := runGPUProcessor(file, selected, opts.GPUBatch, groups, hashes, &writerMu, writer, crackedCount, linesProcessed, totalHashesGenerated, stopChan) + if gpuErr != nil { + return gpuErr + } + if usedGPU { + log.Println("Finished") + return nil + } + } + log.Printf("No usable OpenCL yescrypt groups; falling back to CPU") + } else { + log.Printf("No OpenCL GPU devices found; falling back to CPU") + } + } + } + } + + runCPUProcessor(file, numGoroutines, hashes, &writerMu, writer, crackedCount, linesProcessed, totalHashesGenerated, stopChan) + log.Println("Finished") + return nil +} - // start worker goroutines +func openWordlist(path string) (*os.File, error) { + if path == "" { + return os.Stdin, nil + } + return os.Open(path) +} + +func openOutput(path string) (*bufio.Writer, *os.File, error) { + if path == "" { + return bufio.NewWriter(os.Stdout), nil, nil + } + f, err := os.OpenFile(path, os.O_APPEND|os.O_CREATE|os.O_WRONLY, 0644) + if err != nil { + return nil, nil, err + } + return bufio.NewWriter(f), f, nil +} + +func runCPUProcessor(file *os.File, numGoroutines int, hashes []YescryptHash, writerMu *sync.Mutex, writer *bufio.Writer, crackedCount *int32, linesProcessed, totalHashesGenerated *uint64, stopChan chan struct{}) { + var wg sync.WaitGroup linesCh := make(chan []byte, 1000) for i := 0; i < numGoroutines; i++ { wg.Add(1) go func() { defer wg.Done() for password := range linesCh { - processPassword(password, hashes, &writerMu, writer, crackedCount, linesProcessed, totalHashesGenerated, stopChan) + select { + case <-stopChan: + return + default: + } + processPassword(password, hashes, writerMu, writer, crackedCount, linesProcessed, totalHashesGenerated, stopChan) } }() } - // read lines from file and send them to workers scanner := bufio.NewScanner(file) + scanner.Buffer(make([]byte, 64*1024), 1024*1024) +readLoop: for scanner.Scan() { - line := scanner.Bytes() - decodedPassword, _, _ := checkForHexBytes(line) - password := make([]byte, len(decodedPassword)) - copy(password, decodedPassword) - linesCh <- password + select { + case <-stopChan: + break readLoop + default: + } + decodedPassword, _, _ := checkForHexBytes(scanner.Bytes()) + password := append([]byte(nil), decodedPassword...) + select { + case linesCh <- password: + case <-stopChan: + break readLoop + } } close(linesCh) - if err := scanner.Err(); err != nil { - log.Fatalf("Error reading file: %v\n", err) + log.Printf("Error reading wordlist: %v", err) + } + wg.Wait() +} + +type gpuTask struct{ Passwords [][]byte } + +type gpuWorker struct { + gpu *OpenCLGPU + info GPUDeviceInfo + groups []hashGroup + hashes []YescryptHash + batchCap int + failed bool +} + +type gpuRuntime struct { + workers []*gpuWorker + feedCap int +} + +func (rt *gpuRuntime) Close() { + if rt == nil { + return + } + for _, w := range rt.workers { + if w != nil && w.gpu != nil { + w.gpu.Close() + } + } + rt.workers = nil +} + +func prepareGPUProcessor(devices []GPUDeviceInfo, batchHint int, groups []hashGroup, hashes []YescryptHash, stopChan <-chan struct{}) *gpuRuntime { + if len(devices) == 0 || len(groups) == 0 { + return nil + } + + type prepResult struct { + worker *gpuWorker + } + + results := make(chan prepResult, len(devices)) + var prepWG sync.WaitGroup + for _, info := range devices { + info := info + prepWG.Add(1) + go func() { + defer prepWG.Done() + + select { + case <-stopChan: + return + default: + } + + cap := minGPUCapacity(info, groups, batchHint) + if cap <= 0 { + log.Printf("OpenCL GPU %d (%s) does not have enough allocatable VRAM for these yescrypt parameters", info.Index, info.Name) + return + } + + gpu, err := NewOpenCLGPU(info) + if err != nil { + log.Printf("OpenCL GPU %d (%s) initialization failed: %v", info.Index, info.Name, err) + return + } + + cancelDone := make(chan struct{}) + cancelExit := make(chan struct{}) + go func() { + defer close(cancelExit) + select { + case <-stopChan: + gpu.Cancel() + case <-cancelDone: + } + }() + finishCancelWatch := func() { + close(cancelDone) + <-cancelExit + } + + if err := gpu.SelfTest(); err != nil { + finishCancelWatch() + select { + case <-stopChan: + gpu.Close() + return + default: + } + log.Printf("OpenCL GPU %d (%s) self-test failed: %v", info.Index, info.Name, err) + gpu.Close() + return + } + gpu.ResetCancel() + + w := &gpuWorker{gpu: gpu, info: info, groups: groups, hashes: hashes, batchCap: cap} + + actualCap, err := configureGPUWithBackoff(w, groups[0].N, groups[0].R) + finishCancelWatch() + if err != nil { + select { + case <-stopChan: + gpu.Close() + return + default: + } + log.Printf("OpenCL GPU %d (%s) tuning failed: %v", info.Index, info.Name, err) + gpu.Close() + return + } + + select { + case <-stopChan: + gpu.Close() + return + default: + } + + log.Printf("OpenCL GPU %d: %s, %.1f GiB VRAM, batch cap %d, self-test passed", + info.Index, info.Name, float64(info.GlobalMem)/(1<<30), actualCap) + results <- prepResult{worker: w} + }() + } + + prepWG.Wait() + close(results) + + var workers []*gpuWorker + for result := range results { + if result.worker != nil { + workers = append(workers, result.worker) + } + } + if len(workers) == 0 { + return nil + } + + sort.Slice(workers, func(i, j int) bool { return workers[i].info.Index < workers[j].info.Index }) + feedCap := workers[0].batchCap + for _, w := range workers[1:] { + if w.batchCap < feedCap { + feedCap = w.batchCap + } } + if feedCap < 1 { + for _, w := range workers { + w.gpu.Close() + } + return nil + } + + return &gpuRuntime{workers: workers, feedCap: feedCap} +} + +func runGPUProcessor(file *os.File, devices []GPUDeviceInfo, batchHint int, groups []hashGroup, hashes []YescryptHash, writerMu *sync.Mutex, writer *bufio.Writer, crackedCount *int32, linesProcessed, totalHashesGenerated *uint64, stopChan chan struct{}) (bool, error) { + rt := prepareGPUProcessor(devices, batchHint, groups, hashes, stopChan) + if rt == nil { + return false, nil + } + return runPreparedGPUProcessor(file, rt, writerMu, writer, crackedCount, linesProcessed, totalHashesGenerated, stopChan) +} + +func runPreparedGPUProcessor(file *os.File, rt *gpuRuntime, writerMu *sync.Mutex, writer *bufio.Writer, crackedCount *int32, linesProcessed, totalHashesGenerated *uint64, stopChan chan struct{}) (bool, error) { + if rt == nil || len(rt.workers) == 0 || rt.feedCap < 1 { + return false, nil + } + + workers := rt.workers + feedCap := rt.feedCap + defer rt.Close() + + cancelWatchDone := make(chan struct{}) + cancelWatchExit := make(chan struct{}) + go func() { + defer close(cancelWatchExit) + select { + case <-stopChan: + for _, w := range workers { + w.gpu.Cancel() + } + case <-cancelWatchDone: + } + }() + defer func() { + close(cancelWatchDone) + <-cancelWatchExit + }() + tasks := make(chan gpuTask, len(workers)*2) + fatalSignal := make(chan struct{}) + var fatalOnce sync.Once + var fatalErr error + reportFatal := func(err error) { + if err == nil { + return + } + fatalOnce.Do(func() { + fatalErr = err + close(fatalSignal) + for _, worker := range workers { + worker.gpu.Cancel() + } + }) + } + + var wg sync.WaitGroup + for _, w := range workers { + w := w + wg.Add(1) + go func() { + defer wg.Done() + for task := range tasks { + select { + case <-stopChan: + return + case <-fatalSignal: + return + default: + } + if err := processGPUBatch(w, task.Passwords, writerMu, writer, crackedCount, linesProcessed, totalHashesGenerated, stopChan); err != nil { + reportFatal(err) + return + } + } + }() + } + + scanner := bufio.NewScanner(file) + scanner.Buffer(make([]byte, 64*1024), 1024*1024) + batch := make([][]byte, 0, feedCap) + sendBatch := func() bool { + if len(batch) == 0 { + return true + } + owned := batch + batch = make([][]byte, 0, feedCap) + select { + case tasks <- gpuTask{Passwords: owned}: + return true + case <-stopChan: + return false + case <-fatalSignal: + return false + } + } + +readLoop: + for scanner.Scan() { + select { + case <-stopChan: + break readLoop + case <-fatalSignal: + break readLoop + default: + } + decoded, _, _ := checkForHexBytes(scanner.Bytes()) + pw := append([]byte(nil), decoded...) + batch = append(batch, pw) + if len(batch) == feedCap { + if !sendBatch() { + break readLoop + } + } + } + if err := scanner.Err(); err != nil { + log.Printf("Error reading wordlist: %v", err) + } + select { + case <-fatalSignal: + // do not queue a partial trailing batch after a GPU runtime failure + default: + if len(batch) != 0 { + _ = sendBatch() + } + } + close(tasks) wg.Wait() + if fatalErr != nil { + return true, fatalErr + } + return true, nil +} - log.Println("Finished") +func configureGPUWithBackoff(w *gpuWorker, N, r uint32) (int, error) { + hint := w.batchCap + if hint < 1 { + hint = 1 + } + var lastErr error + for hint >= 1 { + cap, err := w.gpu.Configure(N, r, hint) + if err == nil { + w.batchCap = cap + return cap, nil + } + lastErr = err + if hint == 1 { + break + } + hint /= 2 + if hint < 1 { + hint = 1 + } + } + return 0, lastErr } -func processPassword(password []byte, hashes []YescryptHash, writerMu *sync.Mutex, writer *bufio.Writer, crackedCount, linesProcessed, totalHashesGenerated *int32, stopChan chan struct{}) { - atomic.AddInt32(linesProcessed, 1) - // check for hex, ignore hexErrCount - decodedPassword, _, _ := checkForHexBytes(password) +func processGPUBatch(w *gpuWorker, passwords [][]byte, writerMu *sync.Mutex, writer *bufio.Writer, crackedCount *int32, linesProcessed, totalHashesGenerated *uint64, stopChan chan struct{}) error { + select { + case <-stopChan: + return nil + default: + } - for i := range hashes { - if atomic.LoadInt32(&hashes[i].Cracked) == 0 { - atomic.AddInt32(totalHashesGenerated, 1) - if crackHash(decodedPassword, []byte(hashes[i].Hash)) { - if atomic.CompareAndSwapInt32(&hashes[i].Cracked, 0, 1) { - output := fmt.Sprintf("%s:%s\n", hashes[i].Hash, string(decodedPassword)) - if writer != nil { - writerMu.Lock() - atomic.AddInt32(crackedCount, 1) - writer.WriteString(output) - writer.Flush() - writerMu.Unlock() - } + gpuPW := make([][]byte, 0, len(passwords)) + for _, pw := range passwords { + select { + case <-stopChan: + return nil + default: + } + if len(pw) <= gpuMaxPasswordLen && !w.failed { + gpuPW = append(gpuPW, pw) + } else { + processPasswordNoLineCount(pw, w.hashes, writerMu, writer, crackedCount, totalHashesGenerated, stopChan) + } + } + if len(gpuPW) == 0 { + atomic.AddUint64(linesProcessed, uint64(len(passwords))) + return nil + } - // exit if all hashes are cracked - if isAllHashesCracked(hashes) { - closeStopChannel(stopChan) - return - } + handled := make([]bool, len(w.hashes)) + for gi := range w.groups { + g := &w.groups[gi] + if !groupHasUncracked(w.hashes, g) { + for _, idx := range g.Indices { + handled[idx] = true + } + continue + } + for _, idx := range g.Indices { + handled[idx] = true + } + cap, err := configureGPUWithBackoff(w, g.N, g.R) + if err != nil { + w.failed = true + return fmt.Errorf("OpenCL GPU %d configure failed: %w", w.info.Index, err) + } + for off := 0; off < len(gpuPW); off += cap { + end := off + cap + if end > len(gpuPW) { + end = len(gpuPW) + } + digests, err := w.gpu.HashBatch(gpuPW[off:end], g.Salt) + if err != nil { + if errors.Is(err, errGPUCancelled) { + return nil } + w.failed = true + return fmt.Errorf("OpenCL GPU %d execution failed: %w", w.info.Index, err) + } + atomic.AddUint64(totalHashesGenerated, uint64(len(digests))) + for j := range digests { + pw := gpuPW[off+j] + applyGPUResult(pw, digests[j], g, w.hashes, writerMu, writer, crackedCount, stopChan) + } + select { + case <-stopChan: + return nil + default: + } + } + } + + for idx := range w.hashes { + if handled[idx] || atomic.LoadInt32(&w.hashes[idx].Cracked) != 0 { + continue + } + for _, pw := range gpuPW { + select { + case <-stopChan: + return nil + default: + } + atomic.AddUint64(totalHashesGenerated, 1) + if crackParsedHashCPU(pw, &w.hashes[idx]) { + markCracked(idx, pw, w.hashes, writerMu, writer, crackedCount, stopChan) + } + } + } + + atomic.AddUint64(linesProcessed, uint64(len(passwords))) + return nil +} + +func applyGPUResult(password []byte, digest [32]byte, g *hashGroup, hashes []YescryptHash, writerMu *sync.Mutex, writer *bufio.Writer, crackedCount *int32, stopChan chan struct{}) { + for _, idx := range g.Indices { + if atomic.LoadInt32(&hashes[idx].Cracked) != 0 { + continue + } + matched := false + switch hashes[idx].Kind { + case HashYescrypt: + matched = subtle.ConstantTimeCompare(digest[:], hashes[idx].Expected[:]) == 1 + case HashGostYescrypt: + got, err := gostFinalize(password, hashes[idx].Setting, digest[:]) + if err == nil { + matched = constantTimeDigestEqual(got, hashes[idx].Expected[:]) + } + } + if matched { + markCracked(idx, password, hashes, writerMu, writer, crackedCount, stopChan) + } + } +} + +func markCracked(idx int, password []byte, hashes []YescryptHash, writerMu *sync.Mutex, writer *bufio.Writer, crackedCount *int32, stopChan chan struct{}) { + if !atomic.CompareAndSwapInt32(&hashes[idx].Cracked, 0, 1) { + return + } + output := fmt.Sprintf("%s:%s\n", hashes[idx].Hash, string(password)) + writerMu.Lock() + atomic.AddInt32(crackedCount, 1) + if writer != nil { + _, _ = writer.WriteString(output) + _ = writer.Flush() + } + writerMu.Unlock() + if isAllHashesCracked(hashes) { + closeStopChannel(stopChan) + } +} + +func processPassword(password []byte, hashes []YescryptHash, writerMu *sync.Mutex, writer *bufio.Writer, crackedCount *int32, linesProcessed, totalHashesGenerated *uint64, stopChan chan struct{}) { + atomic.AddUint64(linesProcessed, 1) + processPasswordNoLineCount(password, hashes, writerMu, writer, crackedCount, totalHashesGenerated, stopChan) +} + +func processPasswordNoLineCount(password []byte, hashes []YescryptHash, writerMu *sync.Mutex, writer *bufio.Writer, crackedCount *int32, totalHashesGenerated *uint64, stopChan chan struct{}) { + for i := range hashes { + if atomic.LoadInt32(&hashes[i].Cracked) != 0 { + continue + } + atomic.AddUint64(totalHashesGenerated, 1) + if crackParsedHashCPU(password, &hashes[i]) { + markCracked(i, password, hashes, writerMu, writer, crackedCount, stopChan) + select { + case <-stopChan: + return + default: } } } diff --git a/stats.go b/stats.go index 9a28f4d..bfcbdc5 100644 --- a/stats.go +++ b/stats.go @@ -1,53 +1,124 @@ package main import ( - "fmt" "log" - "os" "sync" "sync/atomic" "time" ) +const ( + statsSampleInterval = time.Second + rateEMAAlpha = 0.15 + etaEMAAlpha = 0.10 +) + // monitor status -func monitorPrintStats(crackedCount, linesProcessed, totalHashesGenerated *int32, stopChan <-chan struct{}, startTime time.Time, totalHashCount int, wg *sync.WaitGroup, interval int) { - var ticker *time.Ticker - if interval > 0 { - ticker = time.NewTicker(time.Duration(interval) * time.Second) - defer ticker.Stop() +func monitorPrintStats(crackedCount *int32, linesProcessed, totalHashesGenerated *uint64, stopChan <-chan struct{}, startTime time.Time, totalHashCount int, totalWordlistLines uint64, wg *sync.WaitGroup, interval int) { + defer wg.Done() + + ticker := time.NewTicker(statsSampleInterval) + defer ticker.Stop() + + lastSampleTime := startTime + lastPrintTime := startTime + var lastLines uint64 + var lastHashes uint64 + var candidateRateEMA float64 + var hashRateEMA float64 + var etaEMA time.Duration + var candidateRateReady bool + var hashRateReady bool + var etaReady bool + + sample := func(now time.Time) { + lines := atomic.LoadUint64(linesProcessed) + hashes := atomic.LoadUint64(totalHashesGenerated) + sampleDuration := now.Sub(lastSampleTime) + + if sampleDuration > 0 { + candidateRate := calcSampleRate(lines, lastLines, sampleDuration) + hashRate := calcSampleRate(hashes, lastHashes, sampleDuration) + + candidateRateEMA, candidateRateReady = updateRateEMA(candidateRateEMA, candidateRate, candidateRateReady, rateEMAAlpha) + hashRateEMA, hashRateReady = updateRateEMA(hashRateEMA, hashRate, hashRateReady, rateEMAAlpha) + + if totalWordlistLines > 0 && candidateRateReady { + if rawETA, ok := estimateRemainingDuration(lines, totalWordlistLines, candidateRateEMA); ok { + etaEMA, etaReady = updateDurationEMA(etaEMA, rawETA, etaReady, etaEMAAlpha) + } + } + } + + lastLines = lines + lastHashes = hashes + lastSampleTime = now + } + + print := func(now time.Time) { + lines := atomic.LoadUint64(linesProcessed) + printStats(now.Sub(startTime), int(atomic.LoadInt32(crackedCount)), totalHashCount, lines, totalWordlistLines, hashRateEMA, hashRateReady, etaEMA, etaReady) } for { select { case <-stopChan: - // print final stats and exit - printStats(time.Since(startTime), int(atomic.LoadInt32(crackedCount)), totalHashCount, int(atomic.LoadInt32(linesProcessed)), true, atomic.LoadInt32(totalHashesGenerated)) - wg.Done() + now := time.Now() + sample(now) + print(now) return - case <-func() <-chan time.Time { - if ticker != nil { - return ticker.C - } - return nil - }(): - if interval > 0 { - printStats(time.Since(startTime), int(atomic.LoadInt32(crackedCount)), totalHashCount, int(atomic.LoadInt32(linesProcessed)), false, atomic.LoadInt32(totalHashesGenerated)) + case now := <-ticker.C: + sample(now) + if interval > 0 && now.Sub(lastPrintTime) >= time.Duration(interval)*time.Second { + print(now) + lastPrintTime = now } } } } +func calcSampleRate(total, previous uint64, interval time.Duration) float64 { + if interval <= 0 || total < previous { + return 0 + } + return float64(total-previous) / interval.Seconds() +} + +func updateRateEMA(current, sample float64, initialized bool, alpha float64) (float64, bool) { + if !initialized { + if sample <= 0 { + return current, false + } + return sample, true + } + return current + alpha*(sample-current), true +} + +func updateDurationEMA(current, sample time.Duration, initialized bool, alpha float64) (time.Duration, bool) { + if !initialized { + return sample, true + } + return time.Duration(float64(current) + alpha*(float64(sample)-float64(current))), true +} + // printStats -func printStats(elapsedTime time.Duration, crackedCount int, totalHashCount, linesProcessed int, exitProgram bool, totalHashesGenerated int32) { - hours := int(elapsedTime.Hours()) - minutes := int(elapsedTime.Minutes()) % 60 - seconds := int(elapsedTime.Seconds()) % 60 - hashesPerSecond := float64(atomic.LoadInt32(&totalHashesGenerated)) / elapsedTime.Seconds() - log.Printf("Cracked: %d/%d %.2f h/s %02dh:%02dm:%02ds", crackedCount, totalHashCount, hashesPerSecond, hours, minutes, seconds) - - if exitProgram { - fmt.Println("") - time.Sleep(100 * time.Millisecond) - os.Exit(0) // exit only if exitProgram bool +func printStats(elapsedTime time.Duration, crackedCount int, totalHashCount int, linesProcessed, totalWordlistLines uint64, hashesPerSecond float64, hashRateReady bool, eta time.Duration, etaReady bool) { + timeText := formatDuration(elapsedTime) + + // stdin has no known line count, so only elapsed time can be displayed + if totalWordlistLines > 0 { + switch { + case linesProcessed >= totalWordlistLines: + timeText += "/00m" + case etaReady: + timeText += "/" + formatETADuration(eta) + default: + timeText += "/--" + } + } + + if !hashRateReady { + hashesPerSecond = 0 } + log.Printf("Cracked: %d/%d %.2f h/s %s", crackedCount, totalHashCount, hashesPerSecond, timeText) } diff --git a/streebog.go b/streebog.go new file mode 100644 index 0000000..04a25a0 --- /dev/null +++ b/streebog.go @@ -0,0 +1,39 @@ +package main + +import ( + "crypto/hmac" + "crypto/subtle" + + "github.com/tarantool/go-gostcrypto/streebog" +) + +func streebog256(data []byte) ([32]byte, error) { + return streebog.Sum256(data), nil +} + +func hmacStreebog256(key, msg []byte) ([32]byte, error) { + var out [32]byte + mac := hmac.New(streebog.New256, key) + _, _ = mac.Write(msg) + copy(out[:], mac.Sum(nil)) + return out, nil +} + +func gostFinalize(password, setting, yescryptDigest []byte) ([32]byte, error) { + hk, err := streebog256(password) + if err != nil { + return [32]byte{}, err + } + interm, err := hmacStreebog256(hk[:], setting) + if err != nil { + return [32]byte{}, err + } + return hmacStreebog256(interm[:], yescryptDigest) +} + +func constantTimeDigestEqual(a [32]byte, b []byte) bool { + if len(b) != 32 { + return false + } + return subtle.ConstantTimeCompare(a[:], b) == 1 +} diff --git a/utils.go b/utils.go index fcefed9..92b9752 100644 --- a/utils.go +++ b/utils.go @@ -1,31 +1,23 @@ package main import ( + "bytes" + "errors" "fmt" + "io" "os" - "os/exec" "os/signal" "runtime" + "strings" "sync/atomic" "syscall" + "time" ) -// clear screen function -func clearScreen() { - var cmd *exec.Cmd - switch runtime.GOOS { - case "linux", "darwin": - cmd = exec.Command("clear") - case "windows": - cmd = exec.Command("cmd", "/c", "cls") - default: - return - } - cmd.Stdout = os.Stdout - if err := cmd.Run(); err != nil { - fmt.Fprintln(os.Stderr, "Failed to clear screen:", err) - } -} +var ( + newlineSeparator = []byte{'\n'} + errWordlistCountCancelled = errors.New("wordlist line count cancelled") +) func closeStopChannel(stopChan chan struct{}) { select { @@ -62,3 +54,124 @@ func isAllHashesCracked(hashes []YescryptHash) bool { } return true } + +func normalizeGPUArgs(args []string) []string { + if len(args) <= 1 { + return args + } + + out := make([]string, 0, len(args)) + out = append(out, args[0]) + + for i := 1; i < len(args); i++ { + arg := args[i] + if arg != "-gpu" && arg != "--gpu" { + out = append(out, arg) + continue + } + + value := "all" + if i+1 < len(args) && !strings.HasPrefix(args[i+1], "-") { + value = args[i+1] + i++ + } + out = append(out, "-gpu="+value) + } + + return out +} + +func countWordlistLines(path string) (uint64, error) { + return countWordlistLinesStop(path, nil) +} + +func countWordlistLinesStop(path string, stopChan <-chan struct{}) (uint64, error) { + file, err := os.Open(path) + if err != nil { + return 0, err + } + defer file.Close() + + // 16 MiB keeps syscall overhead small without consuming excessive RAM + buf := make([]byte, 16*1024*1024) + var lines uint64 + var last byte + var readAny bool + + for { + select { + case <-stopChan: + return 0, errWordlistCountCancelled + default: + } + + n, err := file.Read(buf) + if n > 0 { + readAny = true + lines += uint64(bytes.Count(buf[:n], newlineSeparator)) + last = buf[n-1] + } + if err == io.EOF { + break + } + if err != nil { + return 0, err + } + } + + if readAny && last != '\n' { + lines++ + } + return lines, nil +} + +func formatDuration(d time.Duration) string { + return formatDurationMinutes(d, false) +} + +func formatETADuration(d time.Duration) string { + return formatDurationMinutes(d, true) +} + +func formatDurationMinutes(d time.Duration, roundUp bool) string { + if d < 0 { + d = 0 + } + + totalMinutes := uint64(d / time.Minute) + if roundUp && d%time.Minute != 0 { + totalMinutes++ + } + + days := totalMinutes / (24 * 60) + hours := (totalMinutes / 60) % 24 + minutes := totalMinutes % 60 + + if days > 0 { + return fmt.Sprintf("%02dd:%02dh:%02dm", days, hours, minutes) + } + if hours > 0 { + return fmt.Sprintf("%02dh:%02dm", hours, minutes) + } + return fmt.Sprintf("%02dm", minutes) +} + +func estimateRemainingDuration(linesProcessed, totalLines uint64, candidateRate float64) (time.Duration, bool) { + if linesProcessed >= totalLines { + return 0, true + } + if candidateRate <= 0 { + return 0, false + } + + remaining := float64(totalLines-linesProcessed) / candidateRate + return time.Duration(remaining * float64(time.Second)), true +} + +func estimateRemainingTime(linesProcessed, totalLines uint64, candidateRate float64) string { + remaining, ok := estimateRemainingDuration(linesProcessed, totalLines, candidateRate) + if !ok { + return "--" + } + return formatETADuration(remaining) +} diff --git a/yescrypt_cpu.go b/yescrypt_cpu.go new file mode 100644 index 0000000..461b1ce --- /dev/null +++ b/yescrypt_cpu.go @@ -0,0 +1,347 @@ +package main + +// CPU fallback implementation adapted from Solar Designer's openwall/yescrypt-go +// native yescrypt path. See THIRD_PARTY_NOTICES.md for attribution and license details. + +import ( + "crypto/hmac" + "crypto/sha256" + "encoding/binary" + "errors" + "math/bits" +) + +const maxInt = int(^uint(0) >> 1) + +func pbkdf2SHA256One(password, salt []byte, keyLen int) []byte { + if keyLen <= 0 { + return nil + } + out := make([]byte, 0, keyLen) + var ctr [4]byte + for block := uint32(1); len(out) < keyLen; block++ { + mac := hmac.New(sha256.New, password) + _, _ = mac.Write(salt) + binary.BigEndian.PutUint32(ctr[:], block) + _, _ = mac.Write(ctr[:]) + u := mac.Sum(nil) + remain := keyLen - len(out) + if remain < len(u) { + u = u[:remain] + } + out = append(out, u...) + } + return out +} + +func ycBlockCopy(dst, src []uint64, n int) { copy(dst, src[:n]) } + +func ycBlockXOR(dst, src []uint64, n int) { + for i, v := range src[:n] { + dst[i] ^= v + } +} + +func ycSalsaXOR(tmp *[8]uint64, in, out []uint64, rounds int) { + d0 := tmp[0] ^ in[0] + d1 := tmp[1] ^ in[1] + d2 := tmp[2] ^ in[2] + d3 := tmp[3] ^ in[3] + d4 := tmp[4] ^ in[4] + d5 := tmp[5] ^ in[5] + d6 := tmp[6] ^ in[6] + d7 := tmp[7] ^ in[7] + + x0, x1 := uint32(d0), uint32(d6>>32) + x2, x3 := uint32(d5), uint32(d3>>32) + x4, x5 := uint32(d2), uint32(d0>>32) + x6, x7 := uint32(d7), uint32(d5>>32) + x8, x9 := uint32(d4), uint32(d2>>32) + x10, x11 := uint32(d1), uint32(d7>>32) + x12, x13 := uint32(d6), uint32(d4>>32) + x14, x15 := uint32(d3), uint32(d1>>32) + + for i := 0; i < rounds; i += 2 { + x4 ^= bits.RotateLeft32(x0+x12, 7) + x8 ^= bits.RotateLeft32(x4+x0, 9) + x12 ^= bits.RotateLeft32(x8+x4, 13) + x0 ^= bits.RotateLeft32(x12+x8, 18) + + x9 ^= bits.RotateLeft32(x5+x1, 7) + x13 ^= bits.RotateLeft32(x9+x5, 9) + x1 ^= bits.RotateLeft32(x13+x9, 13) + x5 ^= bits.RotateLeft32(x1+x13, 18) + + x14 ^= bits.RotateLeft32(x10+x6, 7) + x2 ^= bits.RotateLeft32(x14+x10, 9) + x6 ^= bits.RotateLeft32(x2+x14, 13) + x10 ^= bits.RotateLeft32(x6+x2, 18) + + x3 ^= bits.RotateLeft32(x15+x11, 7) + x7 ^= bits.RotateLeft32(x3+x15, 9) + x11 ^= bits.RotateLeft32(x7+x3, 13) + x15 ^= bits.RotateLeft32(x11+x7, 18) + + x1 ^= bits.RotateLeft32(x0+x3, 7) + x2 ^= bits.RotateLeft32(x1+x0, 9) + x3 ^= bits.RotateLeft32(x2+x1, 13) + x0 ^= bits.RotateLeft32(x3+x2, 18) + + x6 ^= bits.RotateLeft32(x5+x4, 7) + x7 ^= bits.RotateLeft32(x6+x5, 9) + x4 ^= bits.RotateLeft32(x7+x6, 13) + x5 ^= bits.RotateLeft32(x4+x7, 18) + + x11 ^= bits.RotateLeft32(x10+x9, 7) + x8 ^= bits.RotateLeft32(x11+x10, 9) + x9 ^= bits.RotateLeft32(x8+x11, 13) + x10 ^= bits.RotateLeft32(x9+x8, 18) + + x12 ^= bits.RotateLeft32(x15+x14, 7) + x13 ^= bits.RotateLeft32(x12+x15, 9) + x14 ^= bits.RotateLeft32(x13+x12, 13) + x15 ^= bits.RotateLeft32(x14+x13, 18) + } + + d0 = uint64(uint32(d0)+x0) | uint64(uint32(d0>>32)+x5)<<32 + d1 = uint64(uint32(d1)+x10) | uint64(uint32(d1>>32)+x15)<<32 + d2 = uint64(uint32(d2)+x4) | uint64(uint32(d2>>32)+x9)<<32 + d3 = uint64(uint32(d3)+x14) | uint64(uint32(d3>>32)+x3)<<32 + d4 = uint64(uint32(d4)+x8) | uint64(uint32(d4>>32)+x13)<<32 + d5 = uint64(uint32(d5)+x2) | uint64(uint32(d5>>32)+x7)<<32 + d6 = uint64(uint32(d6)+x12) | uint64(uint32(d6>>32)+x1)<<32 + d7 = uint64(uint32(d7)+x6) | uint64(uint32(d7>>32)+x11)<<32 + + out[0], tmp[0] = d0, d0 + out[1], tmp[1] = d1, d1 + out[2], tmp[2] = d2, d2 + out[3], tmp[3] = d3, d3 + out[4], tmp[4] = d4, d4 + out[5], tmp[5] = d5, d5 + out[6], tmp[6] = d6, d6 + out[7], tmp[7] = d7, d7 +} + +func ycBlockMix(tmp *[8]uint64, in, out []uint64, r int) { + ycBlockCopy(tmp[:], in[(2*r-1)*8:], 8) + for i := 0; i < 2*r; i += 2 { + ycSalsaXOR(tmp, in[i*8:], out[i*4:], 8) + ycSalsaXOR(tmp, in[i*8+8:], out[i*4+r*8:], 8) + } +} + +const ( + ycPWXsimple = 2 + ycPWXgather = 4 + ycPWXrounds = 6 + ycSwidth = 8 + ycPWXbytes = ycPWXgather * ycPWXsimple * 8 + ycPWXwords = ycPWXbytes / 8 + ycSbytes = 3 * (1 << ycSwidth) * ycPWXsimple * 8 + ycSwords = ycSbytes / 8 + ycSmask = (((1 << ycSwidth) - 1) * ycPWXsimple * 8) +) + +type ycPwxformCtx struct { + S0, S1, S2 []uint64 + w uint32 +} + +func ycPwxform(X *[ycPWXwords]uint64, ctx *ycPwxformCtx) { + S0, S1, S2, w := ctx.S0, ctx.S1, ctx.S2, ctx.w + for i := 0; i < ycPWXrounds; i++ { + for j := 0; j < ycPWXgather; j++ { + x := X[j*ycPWXsimple] + xl := uint32(x) + xh := uint32(x >> 32) + x = uint64(xh) * uint64(xl) + xl = (xl & ycSmask) / 8 + xh = (xh & ycSmask) / 8 + x = (x + S0[xl]) ^ S1[xh] + X[j*ycPWXsimple] = x + + y := X[j*ycPWXsimple+1] + y = ((y>>32)*uint64(uint32(y)) + S0[xl+1]) ^ S1[xh+1] + X[j*ycPWXsimple+1] = y + + if i != 0 && i != ycPWXrounds-1 { + S2[w] = x + S2[w+1] = y + w += 2 + } + } + } + ctx.S0, ctx.S1, ctx.S2 = S2, S0, S1 + ctx.w = w & ((1< 1 { + jj := int(ycWrap(ycInteger(x, r), uint32(i))) + ycBlockXOR(x, v[jj*R:], R) + } + ycBlockMixPwxform(&tmp, x, r, ctx) + } + for i := 0; i < Nloop; i++ { + jj := int(ycInteger(x, r) & uint32(N-1)) + ycBlockXOR(x, v[jj*R:], R) + ycBlockCopy(v[jj*R:], x, R) + ycBlockMixPwxform(&tmp, x, r, ctx) + } + } else { + for i := 0; i < N; i += 2 { + ycBlockCopy(v[i*R:], x, R) + ycBlockMix(&tmp, x, y, r) + ycBlockCopy(v[(i+1)*R:], y, R) + ycBlockMix(&tmp, y, x, r) + } + for i := 0; i < Nloop; i += 2 { + jj := int(ycInteger(x, r) & uint32(N-1)) + ycBlockXOR(x, v[jj*R:], R) + ycBlockMix(&tmp, x, y, r) + jj = int(ycInteger(y, r) & uint32(N-1)) + ycBlockXOR(y, v[jj*R:], R) + ycBlockMix(&tmp, y, x, r) + } + } + + j = 0 + for _, vv := range x[:R] { + idx := (j & ^63) | ((j * 5) & 63) + binary.LittleEndian.PutUint32(b[idx:], uint32(vv)) + j += 4 + idx = (j & ^63) | ((j * 5) & 63) + binary.LittleEndian.PutUint32(b[idx:], uint32(vv>>32)) + j += 4 + } +} + +func ycSmixYescrypt(b []byte, r, N int, v, xy []uint64, passwordSha256 []byte) { + var ctx ycPwxformCtx + var S [ycSwords]uint64 + ycSmix(b, 1, ycSbytes/128, 0, S[:], xy, nil) + ctx.S2 = S[:] + ctx.S1 = S[(1< 1 and a power of 2") + } + if r <= 0 { + return nil, errors.New("yescrypt: r must be > 0") + } + if p != 1 { + return nil, errors.New("yescrypt: p must be 1") + } + if uint64(r)*uint64(p) >= 1<<30 || r > maxInt/128/p || r > maxInt/256 || N > maxInt/128/r { + return nil, errors.New("yescrypt: parameters are too large") + } + + original := password + ppassword := original + pass := 1 + prehash := []byte("yescrypt-prehash") + + v := make([]uint64, 16*N*r) + xy := make([]uint64, 16*max(r, 2)) + workN := N + + if workN/p >= 0x100 && workN/p*r >= 0x20000 { + pass = 0 + workN >>= 6 + } + + var key []byte + var pbuf [32]byte + + for pass <= 1 { + if pass == 1 { + prehash = prehash[:8] + } + h := hmac.New(sha256.New, prehash) + _, _ = h.Write(ppassword) + psum := h.Sum(nil) + copy(pbuf[:], psum) + ppassword = pbuf[:] + + b := pbkdf2SHA256One(ppassword, salt, p*128*r) + copy(ppassword, b[:32]) + ycSmixYescrypt(b, r, workN, v, xy, ppassword) + + key = pbkdf2SHA256One(ppassword, b, max(keyLen, 32)) + if pass == 0 { + copy(ppassword, key[:32]) + // Preserve the pass-0 result before pbuf is overwritten by the + // next prehash operation. + prev := make([]byte, 32) + copy(prev, ppassword) + ppassword = prev + workN <<= 6 + } else { + h1 := hmac.New(sha256.New, key[:32]) + _, _ = h1.Write([]byte("Client Key")) + h2 := sha256.Sum256(h1.Sum(nil)) + copy(key, h2[:]) + } + pass++ + } + + return key[:keyLen], nil +}