arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2025-07-23 06:58:31 +00:00
Code Issues Releases Wiki Activity

Merge branch 'master' into mega-password

Browse Source
This commit is contained in:
hashcat-bot 2025-07-16 20:22:00 +02:00 committed by GitHub
commit ad796638ff
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
2745 changed files with 288100 additions and 80641 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

76
.github/workflows/build.yml vendored
View File

@ -25,6 +25,7 @@ on:
- 'OpenCL/**.cl'
- 'include/**.h'
- 'src/**.c'
- 'src/**.mk'
- 'tools/**'
- '**/Makefile'
- '.github/workflows/build.yml'
@ -36,58 +37,36 @@ on:
- 'OpenCL/**.cl'
- 'include/**.h'
- 'src/**.c'
- 'src/**.mk'
- 'tools/**'
- '**/Makefile'
- '.github/workflows/build.yml'
jobs:
build-linux:
build:
strategy:
fail-fast: false
matrix:
os: [ubuntu-latest, macos-latest, windows-latest]
shared: [0, 1]
name: Build Linux (${{ matrix.shared == 0 && 'Static' || 'Shared' }})
runs-on: ubuntu-latest
steps:
- uses: actions/checkout@v3
- name: Build
env:
SHARED: ${{ matrix.shared }}
run: make
- name: Generate artifacts
uses: actions/upload-artifact@v3
with:
name: hashcat-linux-${{ matrix.shared == 0 && 'static' || 'shared' }}
path: ${{ env.include_paths }}
include:
- os: ubuntu-latest
os_name: Linux
os_name_lowercase: linux
- os: macos-latest
os_name: macOS
os_name_lowercase: macos
- os: windows-latest
os_name: Windows
os_name_lowercase: windows
build-macos:
strategy:
fail-fast: false
matrix:
shared: [0, 1]
name: Build macOS (${{ matrix.shared == 0 && 'Static' || 'Shared' }})
runs-on: macos-latest
name: Build ${{ matrix.os_name }} (${{ matrix.shared == 0 && 'Static' || 'Shared' }})
runs-on: ${{ matrix.os }}
steps:
- uses: actions/checkout@v3
- name: Build
env:
SHARED: ${{ matrix.shared }}
run: make
- name: Generate artifacts
uses: actions/upload-artifact@v3
with:
name: hashcat-macos-${{ matrix.shared == 0 && 'static' || 'shared' }}
path: ${{ env.include_paths }}
- uses: actions/checkout@v4
build-windows:
strategy:
fail-fast: false
matrix:
shared: [0, 1]
name: Build Windows (${{ matrix.shared == 0 && 'Static' || 'Shared' }})
runs-on: windows-latest
steps:
- name: Install libiconv
- name: Install dependencies (Windows only)
if: matrix.os_name_lowercase == 'windows'
uses: msys2/setup-msys2@v2
with:
update: true
@ -97,14 +76,23 @@ jobs:
libiconv
libiconv-devel
make
- uses: actions/checkout@v3
- name: Build
python
- name: Build (Windows)
if: matrix.os_name_lowercase == 'windows'
shell: msys2 {0}
env:
SHARED: ${{ matrix.shared }}
run: make
- name: Build (Linux/macOS)
if: matrix.os_name_lowercase != 'windows'
env:
SHARED: ${{ matrix.shared }}
run: make
- name: Generate artifacts
uses: actions/upload-artifact@v3
uses: actions/upload-artifact@v4
with:
name: hashcat-windows-${{ matrix.shared == 0 && 'static' || 'shared' }}
name: hashcat-${{ matrix.os_name_lowercase }}-${{ matrix.shared == 0 && 'static' || 'shared' }}
path: ${{ env.include_paths }}

11
.gitignore vendored
View File

@ -1,3 +1,5 @@
.DS_Store
*/.DS_Store
*.exe
*.bin
*.app
@ -17,6 +19,8 @@ hashcat.dll
*.dSYM
kernels/**
lib/*.a
bridges/*.dll
bridges/*.so
modules/*.dll
modules/*.so
obj/*/*/*.o
@ -24,3 +28,10 @@ obj/*.o
obj/*.a
include/CL
tools/luks_tests
.vscode
test_edge*
# Byte-compiled / optimized / DLL files
__pycache__/
*.py[cod]
*$py.class

165
BUILD.md
View File

@ -1,87 +1,138 @@
hashcat build documentation
=
### Revision ###
# Hashcat – Build Documentation
* 1.6
**Revision**: 1.7
**Author**: See `docs/credits.txt`
### Author ###
---
See docs/credits.txt
## ✅ Requirements
### Building hashcat for Linux and macOS ###
- **Python 3.12** or higher
Get a copy of the **hashcat** repository
Check your Python version:
```bash
$ python3 --version
# Expected output: Python 3.13.3
```
If you can't install Python ≥ 3.12 globally, you can use **pyenv**.
> If you're using `pyenv`, follow **all steps** below. Otherwise, follow only **steps 3 and 5**.
---
## 🛠️ Building Hashcat – Step-by-Step
### 🔹 Step 1: Install dependencies and pyenv
#### On Linux
Install required libraries to build Python:
```bash
$ sudo apt install libbz2-dev libssl-dev libncurses5-dev libffi-dev libreadline-dev libsqlite3-dev liblzma-dev
```
Install `pyenv`:
```bash
$ curl https://pyenv.run | bash
```
> Follow the instructions shown after installation to set up your shell correctly.
#### On macOS
Install `pyenv` via Homebrew:
```bash
$ brew install pyenv
```
---
### 🔹 Step 2: Install Python using pyenv
Install Python 3.12 (or newer):
```bash
$ pyenv install 3.12
```
Check installed versions:
```bash
$ pyenv versions
# Example:
# * system
# 3.12.11
```
---
### 🔹 Step 3: Clone the Hashcat repository
```bash
$ git clone https://github.com/hashcat/hashcat.git
$ cd hashcat
```
Run "make"
---
```
$ make
### 🔹 Step 4: Set the local Python version
```bash
$ pyenv local 3.12.11
```
### Install hashcat for Linux ###
---
The install target is linux FHS compatible and can be used like this:
### 🔹 Step 5: Build Hashcat
```bash
$ make clean && make
```
---
### 🔹 Step 6 (Optional): Install Hashcat (Linux only)
```bash
$ make install
```
If the $HOME/.hashcat folder exists, then:
Hashcat will use the following locations depending on your environment:
- Session related files go to: $HOME/.hashcat/sessions/
- Cached kernels go to: $HOME/.hashcat/kernels/
- Potfiles go to: $HOME/.hashcat/
| Condition | Session Files | Kernel Cache | Potfiles |
|--------------------------------------------|----------------------------------------|---------------------------------------|----------------------------------------|
| `$HOME/.hashcat` exists | `$HOME/.hashcat/sessions/` | `$HOME/.hashcat/kernels/` | `$HOME/.hashcat/` |
| `$XDG_DATA_HOME` and `$XDG_CACHE_HOME` set | `$XDG_DATA_HOME/hashcat/sessions/` | `$XDG_CACHE_HOME/hashcat/kernels/` | `$XDG_DATA_HOME/hashcat/` |
| Only `$XDG_DATA_HOME` set | `$XDG_DATA_HOME/hashcat/sessions/` | `$HOME/.cache/hashcat/` | `$XDG_DATA_HOME/hashcat/` |
| Only `$XDG_CACHE_HOME` set | `$HOME/.local/share/hashcat/sessions/` | `$XDG_CACHE_HOME/hashcat/kernels/` | `$HOME/.local/share/hashcat/` |
| None of the above | `$HOME/.local/share/hashcat/sessions/` | `$HOME/.cache/hashcat/` | `$HOME/.local/share/hashcat/` |
Otherwise, if environment variable XDG_DATA_HOME and XDG_CACHE_HOME exists, then:
---
- Session related files go to: $XDG_DATA_HOME/hashcat/sessions/
- Cached kernels go to: $XDG_CACHE_HOME/hashcat/kernels/
- Potfiles go to: $XDG_DATA_HOME/hashcat/
## 🐳 Building Hashcat with Docker
Otherwise, if environment variable XDG_DATA_HOME exists, then:
See: [BUILD_Docker.md](BUILD_Docker.md)
- Session related files go to: $XDG_DATA_HOME/hashcat/sessions/
- Cached kernels go to: $HOME/.cache/hashcat
- Potfiles go to: $XDG_DATA_HOME/hashcat/
---
Otherwise, if environment variable XDG_CACHE_HOME exists, then:
## 🪟 Building Hashcat for Windows
- Session related files go to: $HOME/.local/share/hashcat/sessions/
- Cached kernels go to: $XDG_CACHE_HOME/hashcat/kernels/
- Potfiles go to: $HOME/.local/share/hashcat/
| Method | Documentation |
|----------------------------------------|--------------------------------------|
| From macOS | [BUILD_macOS.md](BUILD_macOS.md) |
| Using Windows Subsystem for Linux (WSL)| [BUILD_WSL.md](BUILD_WSL.md) |
| Using Cygwin | [BUILD_CYGWIN.md](BUILD_CYGWIN.md) |
| Using MSYS2 | [BUILD_MSYS2.md](BUILD_MSYS2.md) |
| From Linux | Run: `make win` |
Otherwise:
---
- Session related files go to: $HOME/.local/share/hashcat/sessions/
- Cached kernels go to: $HOME/.cache/hashcat
- Potfiles go to: $HOME/.local/share/hashcat/
## 🎉 Done
### Building hashcat for Windows (using macOS) ###
Refer to [BUILD_macOS.md](BUILD_macOS.md)
### Building hashcat for Windows (using Windows Subsystem for Linux) ###
Refer to [BUILD_WSL.md](BUILD_WSL.md)
### Building hashcat for Windows (using Cygwin) ###
Refer to [BUILD_CYGWIN.md](BUILD_CYGWIN.md)
### Building hashcat for Windows (using MSYS2) ###
Refer to [BUILD_MSYS2.md](BUILD_MSYS2.md)
### Building hashcat for Windows from Linux ###
```
$ make win
```
=
Enjoy your fresh **hashcat** binaries ;)
Enjoy your fresh **Hashcat** binaries! 😎

2
BUILD_CYGWIN.md
View File

@ -14,6 +14,8 @@ gcc-core
gcc-g++
make
git
python312
python312-devel
```
### Building ###

40
BUILD_Docker.md Normal file
View File

@ -0,0 +1,40 @@
# Compiling hashcat binaries with Docker
To build both Linux and Windows binaries in a clean and reproducible environment a dockerfile is available.
It is not considered to be used as a runtime OS.
### Building ###
```bash
docker build -f docker/BinaryPackage.ubuntu20 -t hashcat-binaries .
```
This will create a Docker image with all required toolchains and dependencies.
Optionally you can place custom *.patch or *.diff files into `patches/` folder. They will be applied before compiling.
### Output ###
The resulting output package will be located in: `/root/xy/hashcat-<version>.7z`.
You can copy it to your host with this command:
```bash
docker run --rm \
-e HOST_UID=$(id -u) \
-e HOST_GID=$(id -g) \
-v $(pwd):/out \
hashcat-binaries \
bash -c "cp /root/xy/hashcat-*.7z /out && chown \$HOST_UID:\$HOST_GID /out/hashcat-*.7z"
```
The package will be available on your host machine in the `out` directory.
### Debug ###
In case you want to play around in the docker, run:
```bash
docker run --rm -it hashcat-binaries /bin/bash
```

1
BUILD_MSYS2.md
View File

@ -13,6 +13,7 @@ $ pacman -S git
$ pacman -S make
$ pacman -S gcc
$ pacman -S libiconv-devel
$ pacman -S python3
```
### Building ###

17
BUILD_WSL.md
View File

@ -1,24 +1,27 @@
# Compiling hashcat for Windows with Windows Subsystem for Linux.
# Compiling hashcat for Windows with Windows Subsystem for Linux 2.
Tested on Windows 10 x64, should also work to build hashcat for Windows on Linux.
Tested on Windows 11 x64, should also work to build hashcat for Windows on Linux.
I had it tested with WSL2 using Ubuntu_2004.2020.424.0_x64.appx.
I had it tested with WSL2 using "Ubuntu", which at the time of writing is Ubuntu 24.04
Make sure to have the system upgraded after install (otherwise it will fail to find the gcc-mingw-w64-x86-64 package).
### Installation ###
Enable WSL.
Enable WSL2.
Press the win + r key on your keyboard simultaneously and in the "Run" popup window type bash and make sure to install additional dependencies necessary for hashcat compilation
```
sudo apt install gcc-mingw-w64-x86-64 g++-mingw-w64-x86-64 make git
```bash
sudo apt install build-essential gcc-mingw-w64-x86-64 g++-mingw-w64-x86-64 make git zstd python3-dev cmake
git clone https://github.com/hashcat/hashcat
git clone https://github.com/win-iconv/win-iconv
cd win-iconv/
patch < ../hashcat/tools/win-iconv-64.diff
cmake -D WIN_ICONV_BUILD_EXECUTABLE=OFF -D CMAKE_INSTALL_PREFIX=/opt/win-iconv-64 -D CMAKE_CXX_COMPILER=$(which x86_64-w64-mingw32-g++) -D CMAKE_C_COMPILER=$(which x86_64-w64-mingw32-gcc) -D CMAKE_SYSTEM_NAME=Windows
sudo make install
cd ../
wget https://repo.msys2.org/mingw/mingw64/mingw-w64-x86_64-python-3.12.11-1-any.pkg.tar.zst
sudo mkdir /opt/win-python
sudo tar --zstd -xf mingw-w64-x86_64-python-3.12.11-1-any.pkg.tar.zst -C /opt/win-python
```
### Building ###

5
BUILD_macOS.md
View File

@ -11,9 +11,12 @@ brew install mingw-w64
git clone https://github.com/hashcat/hashcat
git clone https://github.com/win-iconv/win-iconv
cd win-iconv/
patch < ../hashcat/tools/win-iconv-64.diff
cmake -D WIN_ICONV_BUILD_EXECUTABLE=OFF -D CMAKE_INSTALL_PREFIX=/opt/win-iconv-64 -D CMAKE_CXX_COMPILER=$(which x86_64-w64-mingw32-g++) -D CMAKE_C_COMPILER=$(which x86_64-w64-mingw32-gcc) -D CMAKE_SYSTEM_NAME=Windows
sudo make install
cd ../
wget https://repo.msys2.org/mingw/mingw64/mingw-w64-x86_64-python-3.12.10-1-any.pkg.tar.zst
sudo mkdir /opt/win-python
sudo tar --zstd -xf mingw-w64-x86_64-python-3.12.10-1-any.pkg.tar.zst -C /opt/win-python
```
### Building ###

2
OpenCL/amp_a0.cl
View File

@ -13,7 +13,7 @@
#include M2S(INCLUDE_PATH/inc_amp.h)
#endif
KERNEL_FQ void amp (KERN_ATTR_AMP)
KERNEL_FQ KERNEL_FA void amp (KERN_ATTR_AMP)
{
const u64 gid = get_global_id (0);

2
OpenCL/amp_a1.cl
View File

@ -11,7 +11,7 @@
#include M2S(INCLUDE_PATH/inc_amp.h)
#endif
KERNEL_FQ void amp (KERN_ATTR_AMP)
KERNEL_FQ KERNEL_FA void amp (KERN_ATTR_AMP)
{
const u64 gid = get_global_id (0);

2
OpenCL/amp_a3.cl
View File

@ -10,7 +10,7 @@
#include M2S(INCLUDE_PATH/inc_amp.h)
#endif
KERNEL_FQ void amp (KERN_ATTR_AMP)
KERNEL_FQ KERNEL_FA void amp (KERN_ATTR_AMP)
{
const u64 gid = get_global_id (0);

2
OpenCL/inc_amp.h
View File

@ -16,7 +16,7 @@
GLOBAL_AS const bf_t *bfs_buf, \
CONSTANT_AS const u32 &combs_mode, \
CONSTANT_AS const u64 &gid_max, \
uint hc_gid [[ thread_position_in_grid ]]
uint3 hc_gid [[ thread_position_in_grid ]]
#else // CUDA, HIP, OpenCL

74
OpenCL/inc_cipher_cast.cl
View File

@ -569,22 +569,22 @@ DECLSPEC void Cast5Encrypt (PRIVATE_AS const u8 *inBlock, PRIVATE_AS u8 *outBloc
u32 t;
/* Do the work */
_CAST_F1(l, r, 0, 16);
_CAST_F2(r, l, 1, 16);
_CAST_F3(l, r, 2, 16);
_CAST_F1(r, l, 3, 16);
_CAST_F2(l, r, 4, 16);
_CAST_F3(r, l, 5, 16);
_CAST_F1(l, r, 6, 16);
_CAST_F2(r, l, 7, 16);
_CAST_F3(l, r, 8, 16);
_CAST_F1(r, l, 9, 16);
_CAST_F2(l, r, 10, 16);
_CAST_F3(r, l, 11, 16);
_CAST_F1(l, r, 12, 16);
_CAST_F2(r, l, 13, 16);
_CAST_F3(l, r, 14, 16);
_CAST_F1(r, l, 15, 16);
CAST_F1(l, r, 0, 16);
CAST_F2(r, l, 1, 16);
CAST_F3(l, r, 2, 16);
CAST_F1(r, l, 3, 16);
CAST_F2(l, r, 4, 16);
CAST_F3(r, l, 5, 16);
CAST_F1(l, r, 6, 16);
CAST_F2(r, l, 7, 16);
CAST_F3(l, r, 8, 16);
CAST_F1(r, l, 9, 16);
CAST_F2(l, r, 10, 16);
CAST_F3(r, l, 11, 16);
CAST_F1(l, r, 12, 16);
CAST_F2(r, l, 13, 16);
CAST_F3(l, r, 14, 16);
CAST_F1(r, l, 15, 16);
/* Put l,r into outblock */
PUT_UINT32BE(r, outBlock, 0);
@ -599,22 +599,22 @@ DECLSPEC void Cast5Decrypt (PRIVATE_AS const u8 *inBlock, PRIVATE_AS u8 *outBloc
u32 t;
/* Only do full 16 rounds if key length > 80 bits */
_CAST_F1(r, l, 15, 16);
_CAST_F3(l, r, 14, 16);
_CAST_F2(r, l, 13, 16);
_CAST_F1(l, r, 12, 16);
_CAST_F3(r, l, 11, 16);
_CAST_F2(l, r, 10, 16);
_CAST_F1(r, l, 9, 16);
_CAST_F3(l, r, 8, 16);
_CAST_F2(r, l, 7, 16);
_CAST_F1(l, r, 6, 16);
_CAST_F3(r, l, 5, 16);
_CAST_F2(l, r, 4, 16);
_CAST_F1(r, l, 3, 16);
_CAST_F3(l, r, 2, 16);
_CAST_F2(r, l, 1, 16);
_CAST_F1(l, r, 0, 16);
CAST_F1(r, l, 15, 16);
CAST_F3(l, r, 14, 16);
CAST_F2(r, l, 13, 16);
CAST_F1(l, r, 12, 16);
CAST_F3(r, l, 11, 16);
CAST_F2(l, r, 10, 16);
CAST_F1(r, l, 9, 16);
CAST_F3(l, r, 8, 16);
CAST_F2(r, l, 7, 16);
CAST_F1(l, r, 6, 16);
CAST_F3(r, l, 5, 16);
CAST_F2(l, r, 4, 16);
CAST_F1(r, l, 3, 16);
CAST_F3(l, r, 2, 16);
CAST_F2(r, l, 1, 16);
CAST_F1(l, r, 0, 16);
/* Put l,r into outblock */
PUT_UINT32BE(r, outBlock, 0);
PUT_UINT32BE(l, outBlock, 4);
@ -636,7 +636,8 @@ DECLSPEC void Cast5SetKey (PRIVATE_AS CAST_KEY *key, u32 keylength, PRIVATE_AS c
#define x(i) GETBYTE(X[i/4], 3-i%4)
#define z(i) GETBYTE(Z[i/4], 3-i%4)
for (i=0; i<=16; i+=16) {
for (i = 0; i <= 16; i += 16)
{
// this part is copied directly from RFC 2144 (with some search and replace) by Wei Dai
Z[0] = X[0] ^ s_S[4][x(0xD)] ^ s_S[5][x(0xF)] ^ s_S[6][x(0xC)] ^ s_S[7][x(0xE)] ^ s_S[6][x(0x8)];
Z[1] = X[2] ^ s_S[4][z(0x0)] ^ s_S[5][z(0x2)] ^ s_S[6][z(0x1)] ^ s_S[7][z(0x3)] ^ s_S[7][x(0xA)];
@ -673,11 +674,12 @@ DECLSPEC void Cast5SetKey (PRIVATE_AS CAST_KEY *key, u32 keylength, PRIVATE_AS c
}
u32 data[32];
for (i = 0; i < 16; i++) {
for (i = 0; i < 16; i++)
{
data[i * 2] = K[i];
data[i * 2 + 1] = ((K[i + 16]) + 16) & 0x1f; // here only the lowest 5 bits are set..
}
for (i=16; i<32; i++)
K[i] &= 0x1f;
for (i=16; i<32; i++) K[i] &= 0x1f;
}

22
OpenCL/inc_cipher_cast.h
View File

@ -1,7 +1,5 @@
#ifndef _OPENCL_CAST_H
#define _OPENCL_CAST_H
#ifndef INC_CIPHER_CAST_H
#define INC_CIPHER_CAST_H
// #include "opencl_misc.h"
#define GET_UINT32BE(n, b, i) \
@ -32,28 +30,26 @@ typedef struct {
#define U8d(x) GETBYTE(x,0)
/* CAST uses three different round functions */
#define _CAST_f1(l, r, km, kr) \
#define CAST_f1(l, r, km, kr) \
t = hc_rotl32_S(km + r, kr); \
l ^= ((s_S[0][U8a(t)] ^ s_S[1][U8b(t)]) - \
s_S[2][U8c(t)]) + s_S[3][U8d(t)];
#define _CAST_f2(l, r, km, kr) \
#define CAST_f2(l, r, km, kr) \
t = hc_rotl32_S(km ^ r, kr); \
l ^= ((s_S[0][U8a(t)] - s_S[1][U8b(t)]) + \
s_S[2][U8c(t)]) ^ s_S[3][U8d(t)];
#define _CAST_f3(l, r, km, kr) \
#define CAST_f3(l, r, km, kr) \
t = hc_rotl32_S(km - r, kr); \
l ^= ((s_S[0][U8a(t)] + s_S[1][U8b(t)]) ^ \
s_S[2][U8c(t)]) - s_S[3][U8d(t)];
#define _CAST_F1(l, r, i, j) _CAST_f1(l, r, K[i], K[i+j])
#define _CAST_F2(l, r, i, j) _CAST_f2(l, r, K[i], K[i+j])
#define _CAST_F3(l, r, i, j) _CAST_f3(l, r, K[i], K[i+j])
#define CAST_F1(l, r, i, j) CAST_f1(l, r, K[i], K[i+j])
#define CAST_F2(l, r, i, j) CAST_f2(l, r, K[i], K[i+j])
#define CAST_F3(l, r, i, j) CAST_f3(l, r, K[i], K[i+j])
/* OpenSSL API compatibility */
#define CAST_set_key(ckey, len, key) Cast5SetKey(ckey, len, key)
#define CAST_ecb_encrypt(in, out, ckey) Cast5Encrypt(in, out, ckey)
#define CAST_ecb_decrypt(in, out, ckey) Cast5Decrypt(in, out, ckey)
#endif /* _OPENCL_CAST_H */
#endif /* INC_CIPHER_CAST_H */

136
OpenCL/inc_cipher_rc4.cl
View File

@ -4,6 +4,12 @@
#include "inc_common.h"
#include "inc_cipher_rc4.h"
#ifdef IS_HIP
#define RC4_NOINLINE __attribute__ ((noinline))
#else
#define RC4_NOINLINE
#endif
#ifdef IS_CPU
// Pattern linear
@ -136,6 +142,113 @@ DECLSPEC void rc4_init_40 (LOCAL_AS u32 *S, PRIVATE_AS const u32 *key, const u64
j += GET_KEY8 (S, 255, lid) + d0; rc4_swap (S, 255, j, lid);
}
DECLSPEC void rc4_init_72 (LOCAL_AS u32 *S, PRIVATE_AS const u32 *key, const u64 lid)
{
u32 v = 0x03020100;
u32 a = 0x04040404;
#ifdef _unroll
#pragma unroll
#endif
for (u8 i = 0; i < 64; i++)
{
SET_KEY32 (S, i, v, lid); v += a;
}
const u8 d0 = v8a_from_v32_S (key[0]);
const u8 d1 = v8b_from_v32_S (key[0]);
const u8 d2 = v8c_from_v32_S (key[0]);
const u8 d3 = v8d_from_v32_S (key[0]);
const u8 d4 = v8a_from_v32_S (key[1]);
const u8 d5 = v8b_from_v32_S (key[1]);
const u8 d6 = v8c_from_v32_S (key[1]);
const u8 d7 = v8d_from_v32_S (key[1]);
const u8 d8 = v8a_from_v32_S (key[2]);
u8 j = 0;
#ifdef _unroll
#pragma unroll
#endif
for (u32 i = 0; i < 252; i += 9)
{
j += GET_KEY8 (S, i + 0, lid) + d0; rc4_swap (S, i + 0, j, lid);
j += GET_KEY8 (S, i + 1, lid) + d1; rc4_swap (S, i + 1, j, lid);
j += GET_KEY8 (S, i + 2, lid) + d2; rc4_swap (S, i + 2, j, lid);
j += GET_KEY8 (S, i + 3, lid) + d3; rc4_swap (S, i + 3, j, lid);
j += GET_KEY8 (S, i + 4, lid) + d4; rc4_swap (S, i + 4, j, lid);
j += GET_KEY8 (S, i + 5, lid) + d5; rc4_swap (S, i + 5, j, lid);
j += GET_KEY8 (S, i + 6, lid) + d6; rc4_swap (S, i + 6, j, lid);
j += GET_KEY8 (S, i + 7, lid) + d7; rc4_swap (S, i + 7, j, lid);
j += GET_KEY8 (S, i + 8, lid) + d8; rc4_swap (S, i + 8, j, lid);
}
j += GET_KEY8 (S, 252, lid) + d0; rc4_swap (S, 252, j, lid);
j += GET_KEY8 (S, 253, lid) + d1; rc4_swap (S, 253, j, lid);
j += GET_KEY8 (S, 254, lid) + d2; rc4_swap (S, 254, j, lid);
j += GET_KEY8 (S, 255, lid) + d3; rc4_swap (S, 255, j, lid);
}
DECLSPEC void rc4_init_104 (LOCAL_AS u32 *S, PRIVATE_AS const u32 *key, const u64 lid)
{
u32 v = 0x03020100;
u32 a = 0x04040404;
#ifdef _unroll
#pragma unroll
#endif
for (u8 i = 0; i < 64; i++)
{
SET_KEY32 (S, i, v, lid); v += a;
}
const u8 d0 = v8a_from_v32_S(key[0]);
const u8 d1 = v8b_from_v32_S(key[0]);
const u8 d2 = v8c_from_v32_S(key[0]);
const u8 d3 = v8d_from_v32_S(key[0]);
const u8 d4 = v8a_from_v32_S(key[1]);
const u8 d5 = v8b_from_v32_S(key[1]);
const u8 d6 = v8c_from_v32_S(key[1]);
const u8 d7 = v8d_from_v32_S(key[1]);
const u8 d8 = v8a_from_v32_S(key[2]);
const u8 d9 = v8b_from_v32_S(key[2]);
const u8 d10 = v8c_from_v32_S(key[2]);
const u8 d11 = v8d_from_v32_S(key[2]);
const u8 d12 = v8a_from_v32_S(key[3]);
u8 j = 0;
#ifdef _unroll
#pragma unroll
#endif
for (u32 i = 0; i < 247; i += 13)
{
j += GET_KEY8(S, i + 0, lid) + d0; rc4_swap(S, i + 0, j, lid);
j += GET_KEY8(S, i + 1, lid) + d1; rc4_swap(S, i + 1, j, lid);
j += GET_KEY8(S, i + 2, lid) + d2; rc4_swap(S, i + 2, j, lid);
j += GET_KEY8(S, i + 3, lid) + d3; rc4_swap(S, i + 3, j, lid);
j += GET_KEY8(S, i + 4, lid) + d4; rc4_swap(S, i + 4, j, lid);
j += GET_KEY8(S, i + 5, lid) + d5; rc4_swap(S, i + 5, j, lid);
j += GET_KEY8(S, i + 6, lid) + d6; rc4_swap(S, i + 6, j, lid);
j += GET_KEY8(S, i + 7, lid) + d7; rc4_swap(S, i + 7, j, lid);
j += GET_KEY8(S, i + 8, lid) + d8; rc4_swap(S, i + 8, j, lid);
j += GET_KEY8(S, i + 9, lid) + d9; rc4_swap(S, i + 9, j, lid);
j += GET_KEY8(S, i + 10, lid) + d10; rc4_swap(S, i + 10, j, lid);
j += GET_KEY8(S, i + 11, lid) + d11; rc4_swap(S, i + 11, j, lid);
j += GET_KEY8(S, i + 12, lid) + d12; rc4_swap(S, i + 12, j, lid);
}
j += GET_KEY8(S, 247, lid) + d0; rc4_swap(S, 247, j, lid);
j += GET_KEY8(S, 248, lid) + d1; rc4_swap(S, 248, j, lid);
j += GET_KEY8(S, 249, lid) + d2; rc4_swap(S, 249, j, lid);
j += GET_KEY8(S, 250, lid) + d3; rc4_swap(S, 250, j, lid);
j += GET_KEY8(S, 251, lid) + d4; rc4_swap(S, 251, j, lid);
j += GET_KEY8(S, 252, lid) + d5; rc4_swap(S, 252, j, lid);
j += GET_KEY8(S, 253, lid) + d6; rc4_swap(S, 253, j, lid);
j += GET_KEY8(S, 254, lid) + d7; rc4_swap(S, 254, j, lid);
j += GET_KEY8(S, 255, lid) + d8; rc4_swap(S, 255, j, lid);
}
DECLSPEC void rc4_init_128 (LOCAL_AS u32 *S, PRIVATE_AS const u32 *key, const u64 lid)
{
u32 v = 0x03020100;
@ -196,6 +309,27 @@ DECLSPEC void rc4_swap (LOCAL_AS u32 *S, const u8 i, const u8 j, const u64 lid)
SET_KEY8 (S, j, tmp, lid);
}
DECLSPEC void rc4_dropN (LOCAL_AS u32 *S, PRIVATE_AS u8 *i, PRIVATE_AS u8 *j, const u32 n, const u64 lid)
{
u8 a = *i;
u8 b = *j;
for (u32 z = 0; z < n; z++)
{
a += 1;
b += GET_KEY8 (S, a, lid);
rc4_swap (S, a, b, lid);
u8 idx = GET_KEY8 (S, a, lid) + GET_KEY8 (S, b, lid);
GET_KEY8 (S, idx, lid);
}
*i = a;
*j = b;
}
DECLSPEC u8 rc4_next_16 (LOCAL_AS u32 *S, const u8 i, const u8 j, PRIVATE_AS const u32 *in, PRIVATE_AS u32 *out, const u64 lid)
{
u8 a = i;
@ -262,7 +396,7 @@ DECLSPEC u8 rc4_next_16 (LOCAL_AS u32 *S, const u8 i, const u8 j, PRIVATE_AS con
return b;
}
DECLSPEC u8 rc4_next_16_global (LOCAL_AS u32 *S, const u8 i, const u8 j, GLOBAL_AS const u32 *in, PRIVATE_AS u32 *out, const u64 lid)
DECLSPEC RC4_NOINLINE u8 rc4_next_16_global (LOCAL_AS u32 *S, const u8 i, const u8 j, GLOBAL_AS const u32 *in, PRIVATE_AS u32 *out, const u64 lid)
{
u8 a = i;
u8 b = j;

3
OpenCL/inc_cipher_rc4.h
View File

@ -11,8 +11,11 @@ DECLSPEC void SET_KEY8 (LOCAL_AS u32 *S, const u8 k, const u8 v, const u64 lid)
DECLSPEC void SET_KEY32 (LOCAL_AS u32 *S, const u8 k, const u32 v, const u64 lid);
DECLSPEC void rc4_init_40 (LOCAL_AS u32 *S, PRIVATE_AS const u32 *key, const u64 lid);
DECLSPEC void rc4_init_72 (LOCAL_AS u32 *S, PRIVATE_AS const u32 *key, const u64 lid);
DECLSPEC void rc4_init_104 (LOCAL_AS u32 *S, PRIVATE_AS const u32 *key, const u64 lid);
DECLSPEC void rc4_init_128 (LOCAL_AS u32 *S, PRIVATE_AS const u32 *key, const u64 lid);
DECLSPEC void rc4_swap (LOCAL_AS u32 *S, const u8 i, const u8 j, const u64 lid);
DECLSPEC void rc4_dropN (LOCAL_AS u32 *S, PRIVATE_AS u8 *i, PRIVATE_AS u8 *j, const u32 n, const u64 lid);
DECLSPEC u8 rc4_next_16 (LOCAL_AS u32 *S, const u8 i, const u8 j, PRIVATE_AS const u32 *in, PRIVATE_AS u32 *out, const u64 lid);
DECLSPEC u8 rc4_next_16_global (LOCAL_AS u32 *S, const u8 i, const u8 j, GLOBAL_AS const u32 *in, PRIVATE_AS u32 *out, const u64 lid);

31407
OpenCL/inc_common.cl
View File

File diff suppressed because it is too large Load Diff

17
OpenCL/inc_common.h
View File

@ -124,9 +124,10 @@
#if defined IS_METAL
#define KERN_ATTR_MAIN_PARAMS \
uint hc_gid [[ thread_position_in_grid ]], \
uint hc_lid [[ thread_position_in_threadgroup ]], \
uint hc_lsz [[ threads_per_threadgroup ]]
uint3 hc_gid [[ thread_position_in_grid ]], \
uint3 hc_lid [[ thread_position_in_threadgroup ]], \
uint3 hc_lsz [[ threads_per_threadgroup ]], \
uint3 hc_bid [[ threadgroup_position_in_grid ]]
#endif // IS_METAL
/*
@ -283,6 +284,11 @@ DECLSPEC u32 hc_bfe_S (const u32 a, const u32 b, const u32 c);
DECLSPEC u32x hc_lop_0x96 (const u32x a, const u32x b, const u32x c);
DECLSPEC u32 hc_lop_0x96_S (const u32 a, const u32 b, const u32 c);
// arithmetic operations
DECLSPEC u32 hc_umulhi (const u32 x, const u32 y);
DECLSPEC u32 hc_umullo (const u32 x, const u32 y);
// legacy common code
DECLSPEC int ffz (const u32 v);
@ -298,10 +304,13 @@ DECLSPEC void hc_enc_init (PRIVATE_AS hc_enc_t *hc_enc);
DECLSPEC int hc_enc_has_next (PRIVATE_AS hc_enc_t *hc_enc, const int sz);
DECLSPEC int hc_enc_next (PRIVATE_AS hc_enc_t *hc_enc, PRIVATE_AS const u32 *src_buf, const int src_len, const int src_sz, PRIVATE_AS u32 *dst_buf, const int dst_sz);
DECLSPEC int hc_enc_next_global (PRIVATE_AS hc_enc_t *hc_enc, GLOBAL_AS const u32 *src_buf, const int src_len, const int src_sz, PRIVATE_AS u32 *dst_buf, const int dst_sz);
DECLSPEC int hc_enc_validate_utf8 (PRIVATE_AS const u32 *src_buf, const int src_pos, const int extraBytesToRead);
DECLSPEC int hc_enc_validate_utf8_global (GLOBAL_AS const u32 *src_buf, const int src_pos, const int extraBytesToRead);
DECLSPEC int pkcs_padding_bs8 (PRIVATE_AS const u32 *data_buf, const int data_len);
DECLSPEC int pkcs_padding_bs16 (PRIVATE_AS const u32 *data_buf, const int data_len);
DECLSPEC int asn1_detect (PRIVATE_AS const u32 *buf, const int len);
DECLSPEC int asn1_check_int_tag (PRIVATE_AS const u32 *buf, const int len);
DECLSPEC u32 check_bitmap (GLOBAL_AS const u32 *bitmap, const u32 bitmap_mask, const u32 bitmap_shift, const u32 digest);
DECLSPEC u32 check (PRIVATE_AS const u32 *digest, GLOBAL_AS const u32 *bitmap_s1_a, GLOBAL_AS const u32 *bitmap_s1_b, GLOBAL_AS const u32 *bitmap_s1_c, GLOBAL_AS const u32 *bitmap_s1_d, GLOBAL_AS const u32 *bitmap_s2_a, GLOBAL_AS const u32 *bitmap_s2_b, GLOBAL_AS const u32 *bitmap_s2_c, GLOBAL_AS const u32 *bitmap_s2_d, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2);
DECLSPEC void mark_hash (GLOBAL_AS plain_t *plains_buf, GLOBAL_AS u32 *d_result, const u32 salt_pos, const u32 digests_cnt, const u32 digest_pos, const u32 hash_pos, const u64 gid, const u32 il_pos, const u32 extra1, const u32 extra2);
@ -352,6 +361,7 @@ DECLSPEC void append_0x01_2x4_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, const u
DECLSPEC void append_0x06_2x4_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, const u32 offset);
DECLSPEC void append_0x01_4x4_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, const u32 offset);
DECLSPEC void append_0x2d_4x4_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, const u32 offset);
DECLSPEC void append_0x3a_4x4_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, const u32 offset);
DECLSPEC void append_0x80_1x4_S (PRIVATE_AS u32 *w0, const u32 offset);
DECLSPEC void append_0x80_2x4_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, const u32 offset);
DECLSPEC void append_0x80_3x4_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, const u32 offset);
@ -379,5 +389,6 @@ DECLSPEC void append_0x06_2x4_VV (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, cons
DECLSPEC void append_0x80_2x4_VV (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, const u32x offset);
DECLSPEC void append_0x80_4x4_VV (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, const u32x offset);
DECLSPEC void append_0x2d_4x4_VV (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, const u32x offset);
DECLSPEC void append_0x3a_4x4_VV (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, const u32x offset);
#endif // INC_COMMON_H

2
OpenCL/inc_comp_multi_iba.cl
View File

@ -3,7 +3,7 @@ for (int digest_pos = 0; digest_pos < DIGESTS_CNT; digest_pos++)
{
const u32 final_hash_pos = DIGESTS_OFFSET_HOST + digest_pos;
const digest_t *digest = digests_buf + final_hash_pos;
GLOBAL_AS const digest_t *digest = digests_buf + final_hash_pos;
const int invalid_bits = count_bits_32 (digest->digest_buf[0], r0)
+ count_bits_32 (digest->digest_buf[1], r1)

514
OpenCL/inc_hash_argon2.cl Normal file
View File

@ -0,0 +1,514 @@
/**
* Author......: Netherlands Forensic Institute
* License.....: MIT
*
* Warp code based on original work by Ondrej Mosnáček
*/
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_platform.h"
#include "inc_common.h"
#include "inc_hash_blake2b.h"
#include "inc_hash_argon2.h"
#define LBLOCKSIZE (128 / THREADS_PER_LANE)
DECLSPEC void argon2_initial_block (PRIVATE_AS const u32 *in, const u32 lane, const u32 blocknum, const u32 parallelism, GLOBAL_AS argon2_block_t *blocks)
{
blake2b_ctx_t ctx;
blake2b_init (&ctx);
ctx.m[0] = hl32_to_64 (in[ 0], sizeof(argon2_block_t));
ctx.m[1] = hl32_to_64 (in[ 2], in[ 1]);
ctx.m[2] = hl32_to_64 (in[ 4], in[ 3]);
ctx.m[3] = hl32_to_64 (in[ 6], in[ 5]);
ctx.m[4] = hl32_to_64 (in[ 8], in[ 7]);
ctx.m[5] = hl32_to_64 (in[10], in[ 9]);
ctx.m[6] = hl32_to_64 (in[12], in[11]);
ctx.m[7] = hl32_to_64 (in[14], in[13]);
ctx.m[8] = hl32_to_64 (blocknum, in[15]);
ctx.m[9] = hl32_to_64 (0, lane);
blake2b_transform (ctx.h, ctx.m, 76, (u64) BLAKE2B_FINAL);
GLOBAL_AS u64 *out = blocks[(blocknum * parallelism) + lane].values;
out[0] = ctx.h[0];
out[1] = ctx.h[1];
out[2] = ctx.h[2];
out[3] = ctx.h[3];
ctx.m[8] = 0;
ctx.m[9] = 0;
for (u32 off = 4; off < 124; off += 4)
{
for (u32 idx = 0; idx < 8; idx++) ctx.m[idx] = ctx.h[idx];
ctx.h[0] = BLAKE2B_IV_00 ^ 0x01010040; // default output length: 0x40 = 64 bytes
ctx.h[1] = BLAKE2B_IV_01;
ctx.h[2] = BLAKE2B_IV_02;
ctx.h[3] = BLAKE2B_IV_03;
ctx.h[4] = BLAKE2B_IV_04;
ctx.h[5] = BLAKE2B_IV_05;
ctx.h[6] = BLAKE2B_IV_06;
ctx.h[7] = BLAKE2B_IV_07;
blake2b_transform (ctx.h, ctx.m, 64, (u64) BLAKE2B_FINAL);
out[off + 0] = ctx.h[0];
out[off + 1] = ctx.h[1];
out[off + 2] = ctx.h[2];
out[off + 3] = ctx.h[3];
}
out[124] = ctx.h[4];
out[125] = ctx.h[5];
out[126] = ctx.h[6];
out[127] = ctx.h[7];
}
DECLSPEC void blake2b_update_8 (PRIVATE_AS blake2b_ctx_t *ctx, const u32 w0, const u32 w1, const int len)
{
const int pos = ctx->len & 127;
if (pos == 0)
{
if (ctx->len > 0) // if new block (pos == 0) AND the (old) len is not zero => transform
{
blake2b_transform (ctx->h, ctx->m, ctx->len, BLAKE2B_UPDATE);
}
}
const u64 m = hl32_to_64 (w1, w0);
const u32 s = (pos & 7) * 8;
const u64 m0 = (m << s);
const u64 m1 = (m >> 8) >> (56 - s);
const int idx = pos / 8;
ctx->m[ 0] |= (idx == 0) ? m0 : 0;
ctx->m[ 1] |= (idx == 1) ? m0 : (idx == 0) ? m1 : 0;
ctx->m[ 2] |= (idx == 2) ? m0 : (idx == 1) ? m1 : 0;
ctx->m[ 3] |= (idx == 3) ? m0 : (idx == 2) ? m1 : 0;
ctx->m[ 4] |= (idx == 4) ? m0 : (idx == 3) ? m1 : 0;
ctx->m[ 5] |= (idx == 5) ? m0 : (idx == 4) ? m1 : 0;
ctx->m[ 6] |= (idx == 6) ? m0 : (idx == 5) ? m1 : 0;
ctx->m[ 7] |= (idx == 7) ? m0 : (idx == 6) ? m1 : 0;
ctx->m[ 8] |= (idx == 8) ? m0 : (idx == 7) ? m1 : 0;
ctx->m[ 9] |= (idx == 9) ? m0 : (idx == 8) ? m1 : 0;
ctx->m[10] |= (idx == 10) ? m0 : (idx == 9) ? m1 : 0;
ctx->m[11] |= (idx == 11) ? m0 : (idx == 10) ? m1 : 0;
ctx->m[12] |= (idx == 12) ? m0 : (idx == 11) ? m1 : 0;
ctx->m[13] |= (idx == 13) ? m0 : (idx == 12) ? m1 : 0;
ctx->m[14] |= (idx == 14) ? m0 : (idx == 13) ? m1 : 0;
ctx->m[15] |= (idx == 15) ? m0 : (idx == 14) ? m1 : 0;
if ((pos + len) > 128)
{
const u32 cur_len = ((ctx->len + len) / 128) * 128;
blake2b_transform (ctx->h, ctx->m, cur_len, (u64) BLAKE2B_UPDATE);
for (u32 i = 1; i < 16; i++) ctx->m[i] = 0;
ctx->m[0] = m1;
}
ctx->len += len;
}
DECLSPEC void argon2_initial_hash (GLOBAL_AS const pw_t *pw, GLOBAL_AS const salt_t *salt, PRIVATE_AS const argon2_options_t *options, PRIVATE_AS u64 *blockhash)
{
blake2b_ctx_t ctx;
blake2b_init (&ctx);
ctx.m[0] = hl32_to_64 (options->digest_len, options->parallelism);
ctx.m[1] = hl32_to_64 (options->iterations, options->memory_usage_in_kib);
ctx.m[2] = hl32_to_64 (options->type, options->version);
ctx.len = 24;
const u32 pw_len = pw->pw_len;
blake2b_update_8 (&ctx, pw_len, 0, 4);
for (u32 i = 0, idx = 0; i < pw_len; i += 8, idx += 2)
{
blake2b_update_8 (&ctx, pw->i[idx + 0], pw->i[idx + 1], MIN((pw_len - i), 8));
}
const u32 salt_len = salt->salt_len;
blake2b_update_8 (&ctx, salt_len, 0, 4);
for (u32 i = 0, idx = 0; i < salt_len; i += 8, idx += 2)
{
blake2b_update_8 (&ctx, salt->salt_buf[idx + 0], salt->salt_buf[idx + 1], MIN((salt_len - i), 8));
}
blake2b_update_8 (&ctx, 0, 0, 8); // secret (K) and associated data (X)
blake2b_final (&ctx);
for (u32 idx = 0; idx < 8; idx++) blockhash[idx] = ctx.h[idx];
}
DECLSPEC void argon2_init (GLOBAL_AS const pw_t *pw, GLOBAL_AS const salt_t *salt,
PRIVATE_AS const argon2_options_t *options, GLOBAL_AS argon2_block_t *out)
{
u64 blockhash[16] = { 0 };
argon2_initial_hash (pw, salt, options, blockhash);
// Generate the first two blocks of each lane
for (u32 lane = 0; lane < options->parallelism; lane++)
{
argon2_initial_block ((PRIVATE_AS u32 *) blockhash, lane, 0, options->parallelism, out);
argon2_initial_block ((PRIVATE_AS u32 *) blockhash, lane, 1, options->parallelism, out);
}
}
DECLSPEC u64 trunc_mul (const u64 x, const u64 y)
{
const u32 xlo = l32_from_64_S (x);
const u32 ylo = l32_from_64_S (y);
const u32 xyhi = hc_umulhi (xlo, ylo);
const u32 xylo = hc_umullo (xlo, ylo);
return hl32_to_64_S (xyhi, xylo);
}
DECLSPEC inline u32 argon2_ref_address (PRIVATE_AS const argon2_options_t *options, PRIVATE_AS const argon2_pos_t *pos, u32 index, u64 pseudo_random)
{
u32 ref_lane = 0;
u32 ref_area = 0;
u32 ref_index = 0;
if ((pos->pass == 0) && (pos->slice == 0))
{
ref_lane = pos->lane;
}
else
{
ref_lane = h32_from_64_S (pseudo_random) % options->parallelism;
}
ref_area = (pos->pass == 0) ? pos->slice : (ARGON2_SYNC_POINTS - 1);
ref_area *= options->segment_length;
if ((ref_lane == pos->lane) || (index == 0))
{
ref_area += (index - 1);
}
const u32 j1 = l32_from_64_S (pseudo_random);
ref_index = (ref_area - 1 - hc_umulhi (ref_area, hc_umulhi (j1, j1)));
if (pos->pass > 0)
{
ref_index += (pos->slice + 1) * options->segment_length;
if (ref_index >= options->lane_length)
{
ref_index -= options->lane_length;
}
}
return (options->parallelism * ref_index) + ref_lane;
}
DECLSPEC void swap_u64 (PRIVATE_AS u64 *x, PRIVATE_AS u64 *y)
{
u64 tmp = *x;
*x = *y;
*y = tmp;
}
DECLSPEC void transpose_permute_block (u64 R[4], int argon2_thread)
{
if (argon2_thread & 0x08)
{
swap_u64 (&R[0], &R[2]);
swap_u64 (&R[1], &R[3]);
}
if (argon2_thread & 0x04)
{
swap_u64 (&R[0], &R[1]);
swap_u64 (&R[2], &R[3]);
}
}
DECLSPEC int argon2_shift (int idx, int argon2_thread)
{
const int delta = ((idx & 0x02) << 3) + (idx & 0x01);
return (argon2_thread & 0x0e) | (((argon2_thread & 0x11) + delta + 0x0e) & 0x11);
}
DECLSPEC void argon2_hash_block (u64 R[LBLOCKSIZE], int argon2_thread, LOCAL_AS u64 *shuffle_buf, int argon2_lsz)
{
#if THREADS_PER_LANE == 1
u64 v[16];
for (u32 i = 0, offset = 0; i < 8; i++, offset += 16)
{
for (u32 j = 0; j < 16; j++) v[j] = R[offset + j];
ARGON2_P();
for (u32 j = 0; j < 16; j++) R[offset + j] = v[j];
}
for (u32 i = 0, offset = 0; i < 8; i++, offset += 2)
{
for (u32 j = 0, k = offset; j < 16; j += 2, k += 16) {
v[j + 0] = R[k + 0];
v[j + 1] = R[k + 1];
}
ARGON2_P();
for (u32 j = 0, k = offset; j < 16; j += 2, k += 16)
{
R[k + 0] = v[j + 0];
R[k + 1] = v[j + 1];
}
}
#else
for (u32 idx = 1; idx < 4; idx++) R[idx] = hc__shfl_sync (shuffle_buf, FULL_MASK, R[idx], argon2_thread ^ (idx << 2), argon2_thread, argon2_lsz);
transpose_permute_block (R, argon2_thread);
for (u32 idx = 1; idx < 4; idx++) R[idx] = hc__shfl_sync (shuffle_buf, FULL_MASK, R[idx], argon2_thread ^ (idx << 2), argon2_thread, argon2_lsz);
ARGON2_G(R[0], R[1], R[2], R[3]);
for (u32 idx = 1; idx < 4; idx++) R[idx] = hc__shfl_sync (shuffle_buf, FULL_MASK, R[idx], (argon2_thread & 0x1c) | ((argon2_thread + idx) & 0x03), argon2_thread, argon2_lsz);
ARGON2_G(R[0], R[1], R[2], R[3]);
for (u32 idx = 1; idx < 4; idx++) R[idx] = hc__shfl_sync (shuffle_buf, FULL_MASK, R[idx], ((argon2_thread & 0x1c) | ((argon2_thread - idx) & 0x03)) ^ (idx << 2), argon2_thread, argon2_lsz);
transpose_permute_block (R, argon2_thread);
for (u32 idx = 1; idx < 4; idx++) R[idx] = hc__shfl_sync (shuffle_buf, FULL_MASK, R[idx], argon2_thread ^ (idx << 2), argon2_thread, argon2_lsz);
ARGON2_G(R[0], R[1], R[2], R[3]);
for (u32 idx = 1; idx < 4; idx++) R[idx] = hc__shfl_sync (shuffle_buf, FULL_MASK, R[idx], argon2_shift (idx, argon2_thread), argon2_thread, argon2_lsz);
ARGON2_G(R[0], R[1], R[2], R[3]);
for (u32 idx = 1; idx < 4; idx++) R[idx] = hc__shfl_sync (shuffle_buf, FULL_MASK, R[idx], argon2_shift ((4 - idx), argon2_thread), argon2_thread, argon2_lsz);
#endif
}
DECLSPEC void argon2_next_addresses (PRIVATE_AS const argon2_options_t *options, PRIVATE_AS const argon2_pos_t *pos, PRIVATE_AS u32 *addresses, u32 start_index, u32 argon2_thread, LOCAL_AS u64 *shuffle_buf, u32 argon2_lsz)
{
u64 Z[LBLOCKSIZE] = { 0 };
u64 tmp[LBLOCKSIZE] = { 0 };
for (u32 i = 0, index = argon2_thread; i < (LBLOCKSIZE / 4); i++, index += THREADS_PER_LANE)
{
switch (index)
{
case 0: Z[i] = pos->pass; break;
case 1: Z[i] = pos->lane; break;
case 2: Z[i] = pos->slice; break;
case 3: Z[i] = options->memory_block_count; break;
case 4: Z[i] = options->iterations; break;
case 5: Z[i] = options->type; break;
case 6: Z[i] = (start_index / 128) + 1; break;
default: Z[i] = 0; break;
}
tmp[i] = Z[i];
}
argon2_hash_block (Z, argon2_thread, shuffle_buf, argon2_lsz);
for (u32 idx = 0; idx < (LBLOCKSIZE / 4); idx++) Z[idx] ^= tmp[idx];
for (u32 idx = 0; idx < LBLOCKSIZE; idx++) tmp[idx] = Z[idx];
argon2_hash_block (Z, argon2_thread, shuffle_buf, argon2_lsz);
for (u32 idx = 0; idx < LBLOCKSIZE; idx++) Z[idx] ^= tmp[idx];
for (u32 i = 0, index = (start_index + argon2_thread); i < LBLOCKSIZE; i++, index += THREADS_PER_LANE)
{
addresses[i] = argon2_ref_address (options, pos, index, Z[i]);
}
}
DECLSPEC u32 index_u32x4 (const u32 array[4], u32 index)
{
switch (index)
{
case 0:
return array[0];
case 1:
return array[1];
case 2:
return array[2];
case 3:
return array[3];
}
return (u32) -1;
}
DECLSPEC GLOBAL_AS argon2_block_t *argon2_get_current_block (GLOBAL_AS argon2_block_t *blocks, PRIVATE_AS const argon2_options_t *options, u32 lane, u32 index_in_lane, u64 R[LBLOCKSIZE], u32 argon2_thread)
{
// Apply wrap-around to previous block index if the current block is the first block in the lane
const u32 prev_in_lane = (index_in_lane == 0) ? (options->lane_length - 1) : (index_in_lane - 1);
GLOBAL_AS argon2_block_t *prev_block = &blocks[(prev_in_lane * options->parallelism) + lane];
for (u32 idx = 0; idx < LBLOCKSIZE; idx++) R[idx] = prev_block->values[(idx * THREADS_PER_LANE) + argon2_thread];
return &blocks[(index_in_lane * options->parallelism) + lane];
}
DECLSPEC void argon2_fill_subsegment (GLOBAL_AS argon2_block_t *blocks, PRIVATE_AS const argon2_options_t *options, PRIVATE_AS const argon2_pos_t *pos, bool indep_addr, const u32 addresses[LBLOCKSIZE],
u32 start_index, u32 end_index, GLOBAL_AS argon2_block_t *cur_block, u64 R[LBLOCKSIZE], u32 argon2_thread, LOCAL_AS u64 *shuffle_buf, u32 argon2_lsz)
{
for (u32 index = start_index; index < end_index; index++, cur_block += options->parallelism)
{
u32 ref_address = 0;
if (indep_addr)
{
#if THREADS_PER_LANE == 1
ref_address = addresses[(index / THREADS_PER_LANE) % LBLOCKSIZE];
#else
ref_address = index_u32x4 (addresses, (index / THREADS_PER_LANE) % LBLOCKSIZE);
ref_address = hc__shfl_sync (shuffle_buf, FULL_MASK, ref_address, index, argon2_thread, argon2_lsz);
#endif
}
else
{
ref_address = argon2_ref_address (options, pos, index, R[0]);
#if THREADS_PER_LANE != 1
ref_address = hc__shfl_sync (shuffle_buf, FULL_MASK, ref_address, 0, argon2_thread, argon2_lsz);
#endif
}
GLOBAL_AS const argon2_block_t *ref_block = &blocks[ref_address];
u64 tmp[LBLOCKSIZE] = { 0 };
// First pass is overwrite, next passes are XOR with previous
if ((pos->pass > 0) && (options->version != ARGON2_VERSION_10))
{
for (u32 idx = 0; idx < LBLOCKSIZE; idx++) tmp[idx] = cur_block->values[(idx * THREADS_PER_LANE) + argon2_thread];
}
for (u32 idx = 0; idx < LBLOCKSIZE; idx++) R[idx] ^= ref_block->values[(idx * THREADS_PER_LANE) + argon2_thread];
for (u32 idx = 0; idx < LBLOCKSIZE; idx++) tmp[idx] ^= R[idx];
argon2_hash_block (R, argon2_thread, shuffle_buf, argon2_lsz);
for (u32 idx = 0; idx < LBLOCKSIZE; idx++) R[idx] ^= tmp[idx];
for (u32 idx = 0; idx < LBLOCKSIZE; idx++) cur_block->values[(idx * THREADS_PER_LANE) + argon2_thread] = R[idx];
}
}
DECLSPEC void argon2_fill_segment (GLOBAL_AS argon2_block_t *blocks, PRIVATE_AS const argon2_options_t *options, PRIVATE_AS const argon2_pos_t *pos, LOCAL_AS u64 *shuffle_buf, const u32 argon2_thread, const u32 argon2_lsz)
{
// We have already generated the first two blocks of each lane (for the first pass)
const u32 skip_blocks = (pos->pass == 0) && (pos->slice == 0) ? 2 : 0;
const u32 index_in_lane = (pos->slice * options->segment_length) + skip_blocks;
u64 R[LBLOCKSIZE] = { 0 };
GLOBAL_AS argon2_block_t *cur_block = argon2_get_current_block (blocks, options, pos->lane, index_in_lane, R, argon2_thread);
if ((options->type == TYPE_I) || ((options->type == TYPE_ID) && (pos->pass == 0) && (pos->slice <= 1)))
{
for (u32 block_index = 0; block_index < options->segment_length; block_index += 128)
{
const u32 start_index = (block_index == 0) ? skip_blocks : block_index;
const u32 end_index = MIN(((start_index | 127) + 1), options->segment_length);
u32 addresses[LBLOCKSIZE] = { 0 };
argon2_next_addresses (options, pos, addresses, block_index, argon2_thread, shuffle_buf, argon2_lsz);
argon2_fill_subsegment (blocks, options, pos, true, addresses, start_index, end_index, cur_block, R, argon2_thread, shuffle_buf, argon2_lsz);
cur_block += (end_index - start_index) * options->parallelism;
}
}
else
{
u32 addresses[LBLOCKSIZE] = { 0 };
argon2_fill_subsegment (blocks, options, pos, false, addresses, skip_blocks, options->segment_length, cur_block, R, argon2_thread, shuffle_buf, argon2_lsz);
}
}
DECLSPEC void argon2_final (GLOBAL_AS argon2_block_t *blocks, PRIVATE_AS const argon2_options_t *options, PRIVATE_AS u32 *out)
{
const u32 lane_length = options->lane_length;
const u32 lanes = options->parallelism;
blake2b_ctx_t ctx;
blake2b_init (&ctx);
// Override default (0x40) value in BLAKE2b
ctx.h[0] ^= 0x40 ^ options->digest_len;
u32 rem = options->digest_len;
for (u32 offset = 0; offset < 128; offset += 16)
{
for (u32 l = 0; l < lanes; l++)
{
for (u32 idx = 0; idx < 16; idx++)
{
ctx.m[idx] ^= blocks[((lane_length - 1) * lanes) + l].values[idx + offset];
}
}
for (u32 idx = 0; idx < 16; idx++)
{
const u64 value = ctx.m[idx];
ctx.m[idx] = hl32_to_64 (l32_from_64_S (value), rem);
rem = h32_from_64_S (value);
}
ctx.len += 128;
blake2b_transform (ctx.h, ctx.m, ctx.len, (u64) BLAKE2B_UPDATE);
for (u32 idx = 0; idx < 16; idx++) ctx.m[idx] = 0;
}
ctx.m[0] = hl32_to_64 (0, rem);
blake2b_transform (ctx.h, ctx.m, 1028, (u64) BLAKE2B_FINAL);
for (uint i = 0, idx = 0; i < (options->digest_len / 4); i += 2, idx += 1)
{
out [i + 0] = l32_from_64_S (ctx.h[idx]);
out [i + 1] = h32_from_64_S (ctx.h[idx]);
}
}
DECLSPEC GLOBAL_AS argon2_block_t *get_argon2_block (PRIVATE_AS const argon2_options_t *options, GLOBAL_AS void *buf, const int idx)
{
GLOBAL_AS u32 *buf32 = (GLOBAL_AS u32 *) buf;
#ifdef ARGON2_TMP_ELEM
return (GLOBAL_AS argon2_block_t *) buf32 + (ARGON2_TMP_ELEM * idx);
#else
return (GLOBAL_AS argon2_block_t *) buf32 + (options->memory_block_count * idx);
#endif
}

180
OpenCL/inc_hash_argon2.h Normal file
View File

@ -0,0 +1,180 @@
/**
* Author......: Netherlands Forensic Institute
* License.....: MIT
*/
#ifndef INC_HASH_ARGON2_H
#define INC_HASH_ARGON2_H
#define MIN(a,b) (((a) < (b)) ? (a) : (b))
#define ARGON2_VERSION_10 0x10
#define ARGON2_VERSION_13 0x13
#ifndef THREADS_PER_LANE
#define THREADS_PER_LANE 32
#endif
#define FULL_MASK 0xffffffff
#define BLAKE2B_OUTBYTES 64
#define ARGON2_SYNC_POINTS 4
#define ARGON2_ADDRESSES_IN_BLOCK 128
#define TYPE_D 0
#define TYPE_I 1
#define TYPE_ID 2
#if defined IS_CUDA
#define hc__shfl_sync(shfbuf,mask,var,srcLane,argon2_thread,argon2_lsz) __shfl_sync ((mask),(var),(srcLane))
#elif defined IS_HIP
// attention hard coded 32 warps for hip here
#define hc__shfl_sync(shfbuf,mask,var,srcLane,argon2_thread,argon2_lsz) __shfl ((var),(srcLane),32)
#elif defined IS_OPENCL
#define hc__shfl_sync(shfbuf,mask,var,srcLane,argon2_thread,argon2_lsz) hc__shfl ((shfbuf),(var),(srcLane),(argon2_thread),(argon2_lsz))
#if defined IS_AMD && defined IS_GPU
DECLSPEC u64 hc__shfl (MAYBE_UNUSED LOCAL_AS u64 *shuffle_buf, const u64 var, const int src_lane, const u32 argon2_thread, const u32 argon2_lsz)
{
const u32 idx = src_lane << 2;
const u32 l32 = l32_from_64_S (var);
const u32 h32 = h32_from_64_S (var);
const u32 l32r = __builtin_amdgcn_ds_bpermute (idx, l32);
const u32 h32r = __builtin_amdgcn_ds_bpermute (idx, h32);
const u64 out = hl32_to_64_S (h32r, l32r);
return out;
}
#elif defined IS_NV && defined IS_GPU
DECLSPEC u64 hc__shfl (MAYBE_UNUSED LOCAL_AS u64 *shuffle_buf, const u64 var, const int src_lane, const u32 argon2_thread, const u32 argon2_lsz)
{
const u32 l32 = l32_from_64_S (var);
const u32 h32 = h32_from_64_S (var);
u32 l32r;
u32 h32r;
asm("shfl.sync.idx.b32 %0, %1, %2, 0x1f, 0;"
: "=r"(l32r)
: "r"(l32), "r"(src_lane));
asm("shfl.sync.idx.b32 %0, %1, %2, 0x1f, 0;"
: "=r"(h32r)
: "r"(h32), "r"(src_lane));
const u64 out = hl32_to_64_S (h32r, l32r);
return out;
}
#else
DECLSPEC u64 hc__shfl (MAYBE_UNUSED LOCAL_AS u64 *shuffle_buf, const u64 var, const int src_lane, const u32 argon2_thread, const u32 argon2_lsz)
{
shuffle_buf[argon2_thread] = var;
barrier (CLK_LOCAL_MEM_FENCE);
const u64 out = shuffle_buf[src_lane & (argon2_lsz - 1)];
barrier (CLK_LOCAL_MEM_FENCE);
return out;
}
#endif
#elif defined IS_METAL
#define hc__shfl_sync(shfbuf,mask,var,srcLane,argon2_thread,argon2_lsz) simd_shuffle_64 ((var),(srcLane),(argon2_lsz))
DECLSPEC u64 simd_shuffle_64 (const u64 var, const int src_lane, const u32 argon2_lsz)
{
const u32 idx = src_lane & (argon2_lsz - 1);
const u32 l32 = l32_from_64_S (var);
const u32 h32 = h32_from_64_S (var);
u32 l32r = simd_shuffle (l32, idx);
u32 h32r = simd_shuffle (h32, idx);
const u64 out = hl32_to_64_S (h32r, l32r);
return out;
}
#endif
#ifdef IS_CPU
#define ARGON2_G(a,b,c,d) \
{ \
a = a + b + 2 * trunc_mul(a, b); \
d = hc_rotr64_S (d ^ a, 32); \
c = c + d + 2 * trunc_mul(c, d); \
b = hc_rotr64_S (b ^ c, 24); \
a = a + b + 2 * trunc_mul(a, b); \
d = hc_rotr64_S (d ^ a, 16); \
c = c + d + 2 * trunc_mul(c, d); \
b = hc_rotr64_S (b ^ c, 63); \
}
#else
#define ARGON2_G(a,b,c,d) \
{ \
a = a + b + 2 * trunc_mul(a, b); \
d = blake2b_rot32_S (d ^ a); \
c = c + d + 2 * trunc_mul(c, d); \
b = blake2b_rot24_S (b ^ c); \
a = a + b + 2 * trunc_mul(a, b); \
d = blake2b_rot16_S (d ^ a); \
c = c + d + 2 * trunc_mul(c, d); \
b = hc_rotr64_S (b ^ c, 63); \
}
#endif
#define ARGON2_P() \
{ \
ARGON2_G(v[0], v[4], v[8], v[12]); \
ARGON2_G(v[1], v[5], v[9], v[13]); \
ARGON2_G(v[2], v[6], v[10], v[14]); \
ARGON2_G(v[3], v[7], v[11], v[15]); \
\
ARGON2_G(v[0], v[5], v[10], v[15]); \
ARGON2_G(v[1], v[6], v[11], v[12]); \
ARGON2_G(v[2], v[7], v[8], v[13]); \
ARGON2_G(v[3], v[4], v[9], v[14]); \
}
typedef struct argon2_block
{
u64 values[128];
} argon2_block_t;
typedef struct argon2_options
{
u32 type;
u32 version;
u32 iterations;
u32 parallelism;
u32 memory_usage_in_kib;
u32 segment_length;
u32 lane_length;
u32 memory_block_count;
u32 digest_len;
} argon2_options_t;
typedef struct argon2_pos
{
u32 pass;
u32 slice;
u32 lane;
} argon2_pos_t;
DECLSPEC void argon2_init (GLOBAL_AS const pw_t *pw, GLOBAL_AS const salt_t *salt, PRIVATE_AS const argon2_options_t *options, GLOBAL_AS argon2_block_t *out);
DECLSPEC void argon2_fill_segment (GLOBAL_AS argon2_block_t *blocks, PRIVATE_AS const argon2_options_t *options, PRIVATE_AS const argon2_pos_t *pos, LOCAL_AS u64 *shuffle_buf, const u32 argon2_thread, const u32 argon2_lsz);
DECLSPEC void argon2_final (GLOBAL_AS argon2_block_t *blocks, PRIVATE_AS const argon2_options_t *options, PRIVATE_AS u32 *out);
DECLSPEC GLOBAL_AS argon2_block_t *get_argon2_block (PRIVATE_AS const argon2_options_t *options, GLOBAL_AS void *buf, const int idx);
#endif // INC_HASH_ARGON2_H

14
OpenCL/inc_hash_blake2b.cl
View File

@ -24,7 +24,7 @@ DECLSPEC u64 blake2b_rot16_S (const u64 a)
return out.v64;
#elif (defined IS_AMD || defined IS_HIP) && HAS_VPERM == 1
#elif (defined IS_AMD || defined IS_HIP)
vconv64_t in;
@ -98,7 +98,7 @@ DECLSPEC u64 blake2b_rot24_S (const u64 a)
return out.v64;
#elif (defined IS_AMD || defined IS_HIP) && HAS_VPERM == 1
#elif (defined IS_AMD || defined IS_HIP)
vconv64_t in;
@ -409,7 +409,7 @@ DECLSPEC void blake2b_update (PRIVATE_AS blake2b_ctx_t *ctx, PRIVATE_AS const u3
u32 w6[4];
u32 w7[4];
const int limit = (const int) len - 128; // int type needed, could be negative
const int limit = len - 128; // int type needed, could be negative
int pos1;
int pos4;
@ -499,7 +499,7 @@ DECLSPEC void blake2b_update_global (PRIVATE_AS blake2b_ctx_t *ctx, GLOBAL_AS co
u32 w6[4];
u32 w7[4];
const int limit = (const int) len - 128; // int type needed, could be negative
const int limit = len - 128; // int type needed, could be negative
int pos1;
int pos4;
@ -580,7 +580,7 @@ DECLSPEC void blake2b_update_global (PRIVATE_AS blake2b_ctx_t *ctx, GLOBAL_AS co
DECLSPEC void blake2b_final (PRIVATE_AS blake2b_ctx_t *ctx)
{
blake2b_transform (ctx->h, ctx->m, ctx->len, BLAKE2B_FINAL);
blake2b_transform (ctx->h, ctx->m, ctx->len, (u64) BLAKE2B_FINAL);
}
DECLSPEC void blake2b_transform_vector (PRIVATE_AS u64x *h, PRIVATE_AS const u64x *m, const u32x len, const u64 f0)
@ -813,7 +813,7 @@ DECLSPEC void blake2b_update_vector (PRIVATE_AS blake2b_ctx_vector_t *ctx, PRIVA
u32x w6[4];
u32x w7[4];
const int limit = (const int) len - 128; // int type needed, could be negative
const int limit = len - 128; // int type needed, could be negative
int pos1;
int pos4;
@ -894,5 +894,5 @@ DECLSPEC void blake2b_update_vector (PRIVATE_AS blake2b_ctx_vector_t *ctx, PRIVA
DECLSPEC void blake2b_final_vector (PRIVATE_AS blake2b_ctx_vector_t *ctx)
{
blake2b_transform_vector (ctx->h, ctx->m, (u32x) ctx->len, BLAKE2B_FINAL);
blake2b_transform_vector (ctx->h, ctx->m, (u32x) ctx->len, (u64) BLAKE2B_FINAL);
}

674
OpenCL/inc_hash_blake2s.cl
View File

@ -9,6 +9,7 @@
#include "inc_common.h"
#include "inc_hash_blake2s.h"
DECLSPEC u32 blake2s_rot16_S (const u32 a)
{
vconv32_t in;
@ -76,7 +77,7 @@ DECLSPEC u32 blake2s_rot08_S (const u32 a)
return out.v32;
#elif (defined IS_AMD || defined IS_HIP) && HAS_VPERM == 1
#elif (defined IS_AMD || defined IS_HIP)
vconv32_t in;
@ -217,7 +218,7 @@ DECLSPEC void blake2s_update_64 (PRIVATE_AS blake2s_ctx_t *ctx, PRIVATE_AS u32 *
if (pos == 0)
{
if (ctx->len > 0) // if new block (pos == 0) AND the (old) len is not zero => transform
if (ctx->len > 0)
{
blake2s_transform (ctx->h, ctx->m, ctx->len, BLAKE2S_UPDATE);
}
@ -288,8 +289,6 @@ DECLSPEC void blake2s_update_64 (PRIVATE_AS blake2s_ctx_t *ctx, PRIVATE_AS u32 *
ctx->m[14] |= w3[2];
ctx->m[15] |= w3[3];
// len must be a multiple of 64 (not ctx->len) for BLAKE2S_UPDATE:
const u32 cur_len = ((ctx->len + len) / 64) * 64;
blake2s_transform (ctx->h, ctx->m, cur_len, BLAKE2S_UPDATE);
@ -323,7 +322,7 @@ DECLSPEC void blake2s_update (PRIVATE_AS blake2s_ctx_t *ctx, PRIVATE_AS const u3
u32 w2[4];
u32 w3[4];
const int limit = (const int) len - 64; // int type needed, could be negative
const int limit = len - 64; // int type needed, could be negative
int pos1;
int pos4;
@ -377,7 +376,7 @@ DECLSPEC void blake2s_update_global (PRIVATE_AS blake2s_ctx_t *ctx, GLOBAL_AS co
u32 w2[4];
u32 w3[4];
const int limit = (const int) len - 64; // int type needed, could be negative
const int limit = len - 64; // int type needed, could be negative
int pos1;
int pos4;
@ -424,9 +423,516 @@ DECLSPEC void blake2s_update_global (PRIVATE_AS blake2s_ctx_t *ctx, GLOBAL_AS co
blake2s_update_64 (ctx, w0, w1, w2, w3, len - (u32) pos1);
}
DECLSPEC void blake2s_update_swap (PRIVATE_AS blake2s_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len)
{
u32 w0[4];
u32 w1[4];
u32 w2[4];
u32 w3[4];
int pos1;
int pos4;
for (pos1 = 0, pos4 = 0; pos1 < len - 64; pos1 += 64, pos4 += 16)
{
w0[0] = w[pos4 + 0];
w0[1] = w[pos4 + 1];
w0[2] = w[pos4 + 2];
w0[3] = w[pos4 + 3];
w1[0] = w[pos4 + 4];
w1[1] = w[pos4 + 5];
w1[2] = w[pos4 + 6];
w1[3] = w[pos4 + 7];
w2[0] = w[pos4 + 8];
w2[1] = w[pos4 + 9];
w2[2] = w[pos4 + 10];
w2[3] = w[pos4 + 11];
w3[0] = w[pos4 + 12];
w3[1] = w[pos4 + 13];
w3[2] = w[pos4 + 14];
w3[3] = w[pos4 + 15];
w0[0] = hc_swap32_S (w0[0]);
w0[1] = hc_swap32_S (w0[1]);
w0[2] = hc_swap32_S (w0[2]);
w0[3] = hc_swap32_S (w0[3]);
w1[0] = hc_swap32_S (w1[0]);
w1[1] = hc_swap32_S (w1[1]);
w1[2] = hc_swap32_S (w1[2]);
w1[3] = hc_swap32_S (w1[3]);
w2[0] = hc_swap32_S (w2[0]);
w2[1] = hc_swap32_S (w2[1]);
w2[2] = hc_swap32_S (w2[2]);
w2[3] = hc_swap32_S (w2[3]);
w3[0] = hc_swap32_S (w3[0]);
w3[1] = hc_swap32_S (w3[1]);
w3[2] = hc_swap32_S (w3[2]);
w3[3] = hc_swap32_S (w3[3]);
blake2s_update_64 (ctx, w0, w1, w2, w3, 64);
}
w0[0] = w[pos4 + 0];
w0[1] = w[pos4 + 1];
w0[2] = w[pos4 + 2];
w0[3] = w[pos4 + 3];
w1[0] = w[pos4 + 4];
w1[1] = w[pos4 + 5];
w1[2] = w[pos4 + 6];
w1[3] = w[pos4 + 7];
w2[0] = w[pos4 + 8];
w2[1] = w[pos4 + 9];
w2[2] = w[pos4 + 10];
w2[3] = w[pos4 + 11];
w3[0] = w[pos4 + 12];
w3[1] = w[pos4 + 13];
w3[2] = w[pos4 + 14];
w3[3] = w[pos4 + 15];
w0[0] = hc_swap32_S (w0[0]);
w0[1] = hc_swap32_S (w0[1]);
w0[2] = hc_swap32_S (w0[2]);
w0[3] = hc_swap32_S (w0[3]);
w1[0] = hc_swap32_S (w1[0]);
w1[1] = hc_swap32_S (w1[1]);
w1[2] = hc_swap32_S (w1[2]);
w1[3] = hc_swap32_S (w1[3]);
w2[0] = hc_swap32_S (w2[0]);
w2[1] = hc_swap32_S (w2[1]);
w2[2] = hc_swap32_S (w2[2]);
w2[3] = hc_swap32_S (w2[3]);
w3[0] = hc_swap32_S (w3[0]);
w3[1] = hc_swap32_S (w3[1]);
w3[2] = hc_swap32_S (w3[2]);
w3[3] = hc_swap32_S (w3[3]);
blake2s_update_64 (ctx, w0, w1, w2, w3, len - pos1);
}
DECLSPEC void blake2s_update_global_swap (PRIVATE_AS blake2s_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len)
{
u32 w0[4];
u32 w1[4];
u32 w2[4];
u32 w3[4];
const int limit = len - 64; // int type needed, could be negative
int pos1;
int pos4;
for (pos1 = 0, pos4 = 0; pos1 < limit; pos1 += 64, pos4 += 16)
{
w0[0] = w[pos4 + 0];
w0[1] = w[pos4 + 1];
w0[2] = w[pos4 + 2];
w0[3] = w[pos4 + 3];
w1[0] = w[pos4 + 4];
w1[1] = w[pos4 + 5];
w1[2] = w[pos4 + 6];
w1[3] = w[pos4 + 7];
w2[0] = w[pos4 + 8];
w2[1] = w[pos4 + 9];
w2[2] = w[pos4 + 10];
w2[3] = w[pos4 + 11];
w3[0] = w[pos4 + 12];
w3[1] = w[pos4 + 13];
w3[2] = w[pos4 + 14];
w3[3] = w[pos4 + 15];
w0[0] = hc_swap32_S (w0[0]);
w0[1] = hc_swap32_S (w0[1]);
w0[2] = hc_swap32_S (w0[2]);
w0[3] = hc_swap32_S (w0[3]);
w1[0] = hc_swap32_S (w1[0]);
w1[1] = hc_swap32_S (w1[1]);
w1[2] = hc_swap32_S (w1[2]);
w1[3] = hc_swap32_S (w1[3]);
w2[0] = hc_swap32_S (w2[0]);
w2[1] = hc_swap32_S (w2[1]);
w2[2] = hc_swap32_S (w2[2]);
w2[3] = hc_swap32_S (w2[3]);
w3[0] = hc_swap32_S (w3[0]);
w3[1] = hc_swap32_S (w3[1]);
w3[2] = hc_swap32_S (w3[2]);
w3[3] = hc_swap32_S (w3[3]);
blake2s_update_64 (ctx, w0, w1, w2, w3, 64);
}
w0[0] = w[pos4 + 0];
w0[1] = w[pos4 + 1];
w0[2] = w[pos4 + 2];
w0[3] = w[pos4 + 3];
w1[0] = w[pos4 + 4];
w1[1] = w[pos4 + 5];
w1[2] = w[pos4 + 6];
w1[3] = w[pos4 + 7];
w2[0] = w[pos4 + 8];
w2[1] = w[pos4 + 9];
w2[2] = w[pos4 + 10];
w2[3] = w[pos4 + 11];
w3[0] = w[pos4 + 12];
w3[1] = w[pos4 + 13];
w3[2] = w[pos4 + 14];
w3[3] = w[pos4 + 15];
w0[0] = hc_swap32_S (w0[0]);
w0[1] = hc_swap32_S (w0[1]);
w0[2] = hc_swap32_S (w0[2]);
w0[3] = hc_swap32_S (w0[3]);
w1[0] = hc_swap32_S (w1[0]);
w1[1] = hc_swap32_S (w1[1]);
w1[2] = hc_swap32_S (w1[2]);
w1[3] = hc_swap32_S (w1[3]);
w2[0] = hc_swap32_S (w2[0]);
w2[1] = hc_swap32_S (w2[1]);
w2[2] = hc_swap32_S (w2[2]);
w2[3] = hc_swap32_S (w2[3]);
w3[0] = hc_swap32_S (w3[0]);
w3[1] = hc_swap32_S (w3[1]);
w3[2] = hc_swap32_S (w3[2]);
w3[3] = hc_swap32_S (w3[3]);
blake2s_update_64 (ctx, w0, w1, w2, w3, len - (u32) pos1);
}
DECLSPEC void blake2s_final (PRIVATE_AS blake2s_ctx_t *ctx)
{
blake2s_transform (ctx->h, ctx->m, ctx->len, BLAKE2S_FINAL);
blake2s_transform (ctx->h, ctx->m, ctx->len, (u32) BLAKE2S_FINAL);
}
DECLSPEC void blake2s_hmac_init_64 (PRIVATE_AS blake2s_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w0, PRIVATE_AS const u32 *w1, PRIVATE_AS const u32 *w2, PRIVATE_AS const u32 *w3)
{
u32 a0[4];
u32 a1[4];
u32 a2[4];
u32 a3[4];
// ipad
a0[0] = w0[0] ^ 0x36363636;
a0[1] = w0[1] ^ 0x36363636;
a0[2] = w0[2] ^ 0x36363636;
a0[3] = w0[3] ^ 0x36363636;
a1[0] = w1[0] ^ 0x36363636;
a1[1] = w1[1] ^ 0x36363636;
a1[2] = w1[2] ^ 0x36363636;
a1[3] = w1[3] ^ 0x36363636;
a2[0] = w2[0] ^ 0x36363636;
a2[1] = w2[1] ^ 0x36363636;
a2[2] = w2[2] ^ 0x36363636;
a2[3] = w2[3] ^ 0x36363636;
a3[0] = w3[0] ^ 0x36363636;
a3[1] = w3[1] ^ 0x36363636;
a3[2] = w3[2] ^ 0x36363636;
a3[3] = w3[3] ^ 0x36363636;
blake2s_init (&ctx->ipad);
blake2s_update_64 (&ctx->ipad, a0, a1, a2, a3, 64);
// opad
u32 b0[4];
u32 b1[4];
u32 b2[4];
u32 b3[4];
b0[0] = w0[0] ^ 0x5c5c5c5c;
b0[1] = w0[1] ^ 0x5c5c5c5c;
b0[2] = w0[2] ^ 0x5c5c5c5c;
b0[3] = w0[3] ^ 0x5c5c5c5c;
b1[0] = w1[0] ^ 0x5c5c5c5c;
b1[1] = w1[1] ^ 0x5c5c5c5c;
b1[2] = w1[2] ^ 0x5c5c5c5c;
b1[3] = w1[3] ^ 0x5c5c5c5c;
b2[0] = w2[0] ^ 0x5c5c5c5c;
b2[1] = w2[1] ^ 0x5c5c5c5c;
b2[2] = w2[2] ^ 0x5c5c5c5c;
b2[3] = w2[3] ^ 0x5c5c5c5c;
b3[0] = w3[0] ^ 0x5c5c5c5c;
b3[1] = w3[1] ^ 0x5c5c5c5c;
b3[2] = w3[2] ^ 0x5c5c5c5c;
b3[3] = w3[3] ^ 0x5c5c5c5c;
blake2s_init (&ctx->opad);
blake2s_update_64 (&ctx->opad, b0, b1, b2, b3, 64);
}
DECLSPEC void blake2s_hmac_init (PRIVATE_AS blake2s_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len)
{
u32 w0[4];
u32 w1[4];
u32 w2[4];
u32 w3[4];
if (len > 64)
{
blake2s_ctx_t tmp;
blake2s_init (&tmp);
blake2s_update (&tmp, w, len);
blake2s_final (&tmp);
w0[0] = tmp.h[0];
w0[1] = tmp.h[1];
w0[2] = tmp.h[2];
w0[3] = tmp.h[3];
w1[0] = tmp.h[4];
w1[1] = tmp.h[5];
w1[2] = tmp.h[6];
w1[3] = tmp.h[7];
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
w3[3] = 0;
}
else
{
w0[0] = w[ 0];
w0[1] = w[ 1];
w0[2] = w[ 2];
w0[3] = w[ 3];
w1[0] = w[ 4];
w1[1] = w[ 5];
w1[2] = w[ 6];
w1[3] = w[ 7];
w2[0] = w[ 8];
w2[1] = w[ 9];
w2[2] = w[10];
w2[3] = w[11];
w3[0] = w[12];
w3[1] = w[13];
w3[2] = w[14];
w3[3] = w[15];
}
blake2s_hmac_init_64 (ctx, w0, w1, w2, w3);
}
DECLSPEC void blake2s_hmac_init_swap (PRIVATE_AS blake2s_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len)
{
u32 w0[4];
u32 w1[4];
u32 w2[4];
u32 w3[4];
if (len > 64)
{
blake2s_ctx_t tmp;
blake2s_init (&tmp);
blake2s_update_swap (&tmp, w, len);
blake2s_final (&tmp);
w0[0] = tmp.h[0];
w0[1] = tmp.h[1];
w0[2] = tmp.h[2];
w0[3] = tmp.h[3];
w1[0] = tmp.h[4];
w1[1] = tmp.h[5];
w1[2] = tmp.h[6];
w1[3] = tmp.h[7];
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
w3[3] = 0;
}
else
{
w0[0] = hc_swap32_S (w[ 0]);
w0[1] = hc_swap32_S (w[ 1]);
w0[2] = hc_swap32_S (w[ 2]);
w0[3] = hc_swap32_S (w[ 3]);
w1[0] = hc_swap32_S (w[ 4]);
w1[1] = hc_swap32_S (w[ 5]);
w1[2] = hc_swap32_S (w[ 6]);
w1[3] = hc_swap32_S (w[ 7]);
w2[0] = hc_swap32_S (w[ 8]);
w2[1] = hc_swap32_S (w[ 9]);
w2[2] = hc_swap32_S (w[10]);
w2[3] = hc_swap32_S (w[11]);
w3[0] = hc_swap32_S (w[12]);
w3[1] = hc_swap32_S (w[13]);
w3[2] = hc_swap32_S (w[14]);
w3[3] = hc_swap32_S (w[15]);
}
blake2s_hmac_init_64 (ctx, w0, w1, w2, w3);
}
DECLSPEC void blake2s_hmac_init_global (PRIVATE_AS blake2s_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len)
{
u32 w0[4];
u32 w1[4];
u32 w2[4];
u32 w3[4];
if (len > 64)
{
blake2s_ctx_t tmp;
blake2s_init (&tmp);
blake2s_update_global (&tmp, w, len);
blake2s_final (&tmp);
w0[0] = tmp.h[0];
w0[1] = tmp.h[1];
w0[2] = tmp.h[2];
w0[3] = tmp.h[3];
w1[0] = tmp.h[4];
w1[1] = tmp.h[5];
w1[2] = tmp.h[6];
w1[3] = tmp.h[7];
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
w3[3] = 0;
}
else
{
w0[0] = w[ 0];
w0[1] = w[ 1];
w0[2] = w[ 2];
w0[3] = w[ 3];
w1[0] = w[ 4];
w1[1] = w[ 5];
w1[2] = w[ 6];
w1[3] = w[ 7];
w2[0] = w[ 8];
w2[1] = w[ 9];
w2[2] = w[10];
w2[3] = w[11];
w3[0] = w[12];
w3[1] = w[13];
w3[2] = w[14];
w3[3] = w[15];
}
blake2s_hmac_init_64 (ctx, w0, w1, w2, w3);
}
DECLSPEC void blake2s_hmac_init_global_swap (PRIVATE_AS blake2s_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len)
{
u32 w0[4];
u32 w1[4];
u32 w2[4];
u32 w3[4];
if (len > 64)
{
blake2s_ctx_t tmp;
blake2s_init (&tmp);
blake2s_update_global_swap (&tmp, w, len);
blake2s_final (&tmp);
w0[0] = tmp.h[0];
w0[1] = tmp.h[1];
w0[2] = tmp.h[2];
w0[3] = tmp.h[3];
w1[0] = tmp.h[4];
w1[1] = tmp.h[5];
w1[2] = tmp.h[6];
w1[3] = tmp.h[7];
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
w3[3] = 0;
}
else
{
w0[0] = hc_swap32_S (w[ 0]);
w0[1] = hc_swap32_S (w[ 1]);
w0[2] = hc_swap32_S (w[ 2]);
w0[3] = hc_swap32_S (w[ 3]);
w1[0] = hc_swap32_S (w[ 4]);
w1[1] = hc_swap32_S (w[ 5]);
w1[2] = hc_swap32_S (w[ 6]);
w1[3] = hc_swap32_S (w[ 7]);
w2[0] = hc_swap32_S (w[ 8]);
w2[1] = hc_swap32_S (w[ 9]);
w2[2] = hc_swap32_S (w[10]);
w2[3] = hc_swap32_S (w[11]);
w3[0] = hc_swap32_S (w[12]);
w3[1] = hc_swap32_S (w[13]);
w3[2] = hc_swap32_S (w[14]);
w3[3] = hc_swap32_S (w[15]);
}
blake2s_hmac_init_64 (ctx, w0, w1, w2, w3);
}
DECLSPEC void blake2s_hmac_update_64 (PRIVATE_AS blake2s_hmac_ctx_t *ctx, PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, const int len)
{
blake2s_update_64 (&ctx->ipad, w0, w1, w2, w3, len);
}
DECLSPEC void blake2s_hmac_update (PRIVATE_AS blake2s_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len)
{
blake2s_update (&ctx->ipad, w, len);
}
DECLSPEC void blake2s_hmac_update_swap (PRIVATE_AS blake2s_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len)
{
blake2s_update_swap (&ctx->ipad, w, len);
}
DECLSPEC void blake2s_hmac_update_global (PRIVATE_AS blake2s_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len)
{
blake2s_update_global (&ctx->ipad, w, len);
}
DECLSPEC void blake2s_hmac_update_global_swap (PRIVATE_AS blake2s_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len)
{
blake2s_update_global_swap (&ctx->ipad, w, len);
}
DECLSPEC void blake2s_hmac_final (PRIVATE_AS blake2s_hmac_ctx_t *ctx)
{
blake2s_final (&ctx->ipad);
for (int n = 0; n < 8; n += 1)
{
blake2s_update(&ctx->opad, &ctx->ipad.h[n], 4);
}
ctx->opad.m[8] = 0;
ctx->opad.m[9] = 0;
ctx->opad.m[10]= 0;
ctx->opad.m[11]= 0;
ctx->opad.m[12]= 0;
ctx->opad.m[13]= 0;
ctx->opad.m[14]= 0;
ctx->opad.m[15]= 0;
blake2s_final (&ctx->opad);
}
DECLSPEC void blake2s_transform_vector (PRIVATE_AS u32x *h, PRIVATE_AS const u32x *m, const u32x len, const u32 f0)
@ -452,6 +958,7 @@ DECLSPEC void blake2s_transform_vector (PRIVATE_AS u32x *h, PRIVATE_AS const u32
v[14] = BLAKE2S_IV_06 ^ f0;
v[15] = BLAKE2S_IV_07; // ^ f1;
BLAKE2S_ROUND_VECTOR ( 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15);
BLAKE2S_ROUND_VECTOR (14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3);
BLAKE2S_ROUND_VECTOR (11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4);
@ -649,7 +1156,7 @@ DECLSPEC void blake2s_update_vector (PRIVATE_AS blake2s_ctx_vector_t *ctx, PRIVA
u32x w2[4];
u32x w3[4];
const int limit = (const int) len - 64; // int type needed, could be negative
const int limit = len - 64; // int type needed, could be negative
int pos1;
int pos4;
@ -698,5 +1205,154 @@ DECLSPEC void blake2s_update_vector (PRIVATE_AS blake2s_ctx_vector_t *ctx, PRIVA
DECLSPEC void blake2s_final_vector (PRIVATE_AS blake2s_ctx_vector_t *ctx)
{
blake2s_transform_vector (ctx->h, ctx->m, (u32x) ctx->len, BLAKE2S_FINAL);
blake2s_transform_vector (ctx->h, ctx->m, (u32x) ctx->len, (u32) BLAKE2S_FINAL);
}
DECLSPEC void blake2s_hmac_init_vector_64 (PRIVATE_AS blake2s_hmac_ctx_vector_t *ctx, PRIVATE_AS const u32x *w0, PRIVATE_AS const u32x *w1, PRIVATE_AS const u32x *w2, PRIVATE_AS const u32x *w3)
{
u32x a0[4];
u32x a1[4];
u32x a2[4];
u32x a3[4];
// ipad
a0[0] = w0[0] ^ 0x36363636;
a0[1] = w0[1] ^ 0x36363636;
a0[2] = w0[2] ^ 0x36363636;
a0[3] = w0[3] ^ 0x36363636;
a1[0] = w1[0] ^ 0x36363636;
a1[1] = w1[1] ^ 0x36363636;
a1[2] = w1[2] ^ 0x36363636;
a1[3] = w1[3] ^ 0x36363636;
a2[0] = w2[0] ^ 0x36363636;
a2[1] = w2[1] ^ 0x36363636;
a2[2] = w2[2] ^ 0x36363636;
a2[3] = w2[3] ^ 0x36363636;
a3[0] = w3[0] ^ 0x36363636;
a3[1] = w3[1] ^ 0x36363636;
a3[2] = w3[2] ^ 0x36363636;
a3[3] = w3[3] ^ 0x36363636;
blake2s_init_vector (&ctx->ipad);
blake2s_update_vector_64 (&ctx->ipad, a0, a1, a2, a3, 64);
// opad
u32x b0[4];
u32x b1[4];
u32x b2[4];
u32x b3[4];
b0[0] = w0[0] ^ 0x5c5c5c5c;
b0[1] = w0[1] ^ 0x5c5c5c5c;
b0[2] = w0[2] ^ 0x5c5c5c5c;
b0[3] = w0[3] ^ 0x5c5c5c5c;
b1[0] = w1[0] ^ 0x5c5c5c5c;
b1[1] = w1[1] ^ 0x5c5c5c5c;
b1[2] = w1[2] ^ 0x5c5c5c5c;
b1[3] = w1[3] ^ 0x5c5c5c5c;
b2[0] = w2[0] ^ 0x5c5c5c5c;
b2[1] = w2[1] ^ 0x5c5c5c5c;
b2[2] = w2[2] ^ 0x5c5c5c5c;
b2[3] = w2[3] ^ 0x5c5c5c5c;
b3[0] = w3[0] ^ 0x5c5c5c5c;
b3[1] = w3[1] ^ 0x5c5c5c5c;
b3[2] = w3[2] ^ 0x5c5c5c5c;
b3[3] = w3[3] ^ 0x5c5c5c5c;
blake2s_init_vector (&ctx->opad);
blake2s_update_vector_64 (&ctx->opad, b0, b1, b2, b3, 64);
}
DECLSPEC void blake2s_hmac_init_vector (PRIVATE_AS blake2s_hmac_ctx_vector_t *ctx, PRIVATE_AS const u32x *w, const int len)
{
u32x w0[4];
u32x w1[4];
u32x w2[4];
u32x w3[4];
if (len > 64)
{
blake2s_ctx_vector_t tmp;
blake2s_init_vector (&tmp);
blake2s_update_vector (&tmp, w, len);
blake2s_final_vector (&tmp);
w0[0] = tmp.h[0];
w0[1] = tmp.h[1];
w0[2] = tmp.h[2];
w0[3] = tmp.h[3];
w1[0] = tmp.h[4];
w1[1] = tmp.h[5];
w1[2] = tmp.h[6];
w1[3] = tmp.h[7];
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
w3[3] = 0;
}
else
{
w0[0] = w[ 0];
w0[1] = w[ 1];
w0[2] = w[ 2];
w0[3] = w[ 3];
w1[0] = w[ 4];
w1[1] = w[ 5];
w1[2] = w[ 6];
w1[3] = w[ 7];
w2[0] = w[ 8];
w2[1] = w[ 9];
w2[2] = w[10];
w2[3] = w[11];
w3[0] = w[12];
w3[1] = w[13];
w3[2] = w[14];
w3[3] = w[15];
}
blake2s_hmac_init_vector_64 (ctx, w0, w1, w2, w3);
}
DECLSPEC void blake2s_hmac_update_vector_64 (PRIVATE_AS blake2s_hmac_ctx_vector_t *ctx, PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, const int len)
{
blake2s_update_vector_64 (&ctx->ipad, w0, w1, w2, w3, len);
}
DECLSPEC void blake2s_hmac_update_vector (PRIVATE_AS blake2s_hmac_ctx_vector_t *ctx, PRIVATE_AS const u32x *w, const int len)
{
blake2s_update_vector (&ctx->ipad, w, len);
}
DECLSPEC void blake2s_hmac_final_vector (PRIVATE_AS blake2s_hmac_ctx_vector_t *ctx)
{
blake2s_final_vector (&ctx->ipad);
for (int n = 0; n < 8; n += 1)
{
blake2s_update_vector(&ctx->opad, &ctx->ipad.h[n], 4);
}
ctx->opad.m[8] = 0;
ctx->opad.m[9] = 0;
ctx->opad.m[10]= 0;
ctx->opad.m[11]= 0;
ctx->opad.m[12]= 0;
ctx->opad.m[13]= 0;
ctx->opad.m[14]= 0;
ctx->opad.m[15]= 0;
blake2s_final_vector (&ctx->opad);
}

37
OpenCL/inc_hash_blake2s.h
View File

@ -72,6 +72,14 @@ typedef struct blake2s_ctx
} blake2s_ctx_t;
typedef struct blake2s_hmac_ctx
{
blake2s_ctx_t ipad;
blake2s_ctx_t opad;
} blake2s_hmac_ctx_t;
typedef struct blake2s_ctx_vector
{
u32x m[16]; // buffer
@ -81,16 +89,45 @@ typedef struct blake2s_ctx_vector
} blake2s_ctx_vector_t;
typedef struct blake2s_hmac_ctx_vector
{
blake2s_ctx_vector_t ipad;
blake2s_ctx_vector_t opad;
} blake2s_hmac_ctx_vector_t;
DECLSPEC void blake2s_transform (PRIVATE_AS u32 *h, PRIVATE_AS const u32 *m, const int len, const u32 f0);
DECLSPEC void blake2s_init (PRIVATE_AS blake2s_ctx_t *ctx);
DECLSPEC void blake2s_update (PRIVATE_AS blake2s_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void blake2s_update_global (PRIVATE_AS blake2s_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void blake2s_update_global_swap (PRIVATE_AS blake2s_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void blake2s_final (PRIVATE_AS blake2s_ctx_t *ctx);
DECLSPEC void blake2s_hmac_init_64 (PRIVATE_AS blake2s_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w0, PRIVATE_AS const u32 *w1, PRIVATE_AS const u32 *w2, PRIVATE_AS const u32 *w3);
DECLSPEC void blake2s_hmac_init (PRIVATE_AS blake2s_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void blake2s_hmac_init_swap (PRIVATE_AS blake2s_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void blake2s_hmac_init_global (PRIVATE_AS blake2s_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void blake2s_hmac_init_global_swap (PRIVATE_AS blake2s_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void blake2s_hmac_update_64 (PRIVATE_AS blake2s_hmac_ctx_t *ctx, PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, const int len);
DECLSPEC void blake2s_hmac_update (PRIVATE_AS blake2s_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void blake2s_hmac_update_swap (PRIVATE_AS blake2s_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void blake2s_hmac_update_global (PRIVATE_AS blake2s_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void blake2s_hmac_update_global_swap (PRIVATE_AS blake2s_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void blake2s_hmac_final (PRIVATE_AS blake2s_hmac_ctx_t *ctx);
DECLSPEC void blake2s_transform_vector (PRIVATE_AS u32x *h, PRIVATE_AS const u32x *m, const u32x len, const u32 f0);
DECLSPEC void blake2s_init_vector (PRIVATE_AS blake2s_ctx_vector_t *ctx);
DECLSPEC void blake2s_init_vector_from_scalar (PRIVATE_AS blake2s_ctx_vector_t *ctx, PRIVATE_AS blake2s_ctx_t *ctx0);
DECLSPEC void blake2s_update_vector (PRIVATE_AS blake2s_ctx_vector_t *ctx, PRIVATE_AS const u32x *w, const int len);
DECLSPEC void blake2s_final_vector (PRIVATE_AS blake2s_ctx_vector_t *ctx);
DECLSPEC void blake2s_hmac_init_vector_64 (PRIVATE_AS blake2s_hmac_ctx_vector_t *ctx, PRIVATE_AS const u32x *w0, PRIVATE_AS const u32x *w1, PRIVATE_AS const u32x *w2, PRIVATE_AS const u32x *w3);
DECLSPEC void blake2s_hmac_init_vector (PRIVATE_AS blake2s_hmac_ctx_vector_t *ctx, PRIVATE_AS const u32x *w, const int len);
DECLSPEC void blake2s_hmac_update_vector_64 (PRIVATE_AS blake2s_hmac_ctx_vector_t *ctx, PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, const int len);
DECLSPEC void blake2s_hmac_update_vector (PRIVATE_AS blake2s_hmac_ctx_vector_t *ctx, PRIVATE_AS const u32x *w, const int len);
DECLSPEC void blake2s_hmac_final_vector (PRIVATE_AS blake2s_hmac_ctx_vector_t *ctx);
#endif // INC_HASH_BLAKE2S_H

2279
OpenCL/inc_hash_ripemd320.cl Normal file
View File

File diff suppressed because it is too large Load Diff

147
OpenCL/inc_hash_ripemd320.h Normal file
View File

@ -0,0 +1,147 @@
/**
* Author......: See docs/credits.txt
* License.....: MIT
*/
#ifndef INC_HASH_RIPEMD320_H
#define INC_HASH_RIPEMD320_H
#define RIPEMD320_F(x,y,z) ((x) ^ (y) ^ (z))
#define RIPEMD320_G(x,y,z) ((z) ^ ((x) & ((y) ^ (z)))) /* x ? y : z */
#define RIPEMD320_H(x,y,z) (((x) | ~(y)) ^ (z))
#define RIPEMD320_I(x,y,z) ((y) ^ ((z) & ((x) ^ (y)))) /* z ? x : y */
#define RIPEMD320_J(x,y,z) ((x) ^ ((y) | ~(z)))
#ifdef USE_BITSELECT
#define RIPEMD320_Go(x,y,z) (bitselect ((z), (y), (x)))
#define RIPEMD320_Io(x,y,z) (bitselect ((y), (x), (z)))
#else
#define RIPEMD320_Go(x,y,z) (RIPEMD320_G ((x), (y), (z)))
#define RIPEMD320_Io(x,y,z) (RIPEMD320_I ((x), (y), (z)))
#endif
#define RIPEMD320_STEP_S(f,a,b,c,d,e,x,K,s) \
{ \
a += K; \
a += x; \
a += f (b, c, d); \
a = hc_rotl32_S (a, s); \
a += e; \
c = hc_rotl32_S (c, 10u); \
}
#define RIPEMD320_STEP(f,a,b,c,d,e,x,K,s) \
{ \
a += make_u32x (K); \
a += x; \
a += f (b, c, d); \
a = hc_rotl32 (a, s); \
a += e; \
c = hc_rotl32 (c, 10u); \
}
#define ROTATE_LEFT_WORKAROUND_BUG(a,n) ((a << n) | (a >> (32 - n)))
#define RIPEMD320_STEP_S_WORKAROUND_BUG(f,a,b,c,d,e,x,K,s) \
{ \
a += K; \
a += x; \
a += f (b, c, d); \
a = ROTATE_LEFT_WORKAROUND_BUG (a, s); \
a += e; \
c = hc_rotl32_S (c, 10u); \
}
#define RIPEMD320_STEP_WORKAROUND_BUG(f,a,b,c,d,e,x,K,s) \
{ \
a += make_u32x (K); \
a += x; \
a += f (b, c, d); \
a = ROTATE_LEFT_WORKAROUND_BUG (a, s); \
a += e; \
c = hc_rotl32 (c, 10u); \
}
typedef struct ripemd320_ctx
{
u32 h[10];
u32 w0[4];
u32 w1[4];
u32 w2[4];
u32 w3[4];
int len;
} ripemd320_ctx_t;
typedef struct ripemd320_hmac_ctx
{
ripemd320_ctx_t ipad;
ripemd320_ctx_t opad;
} ripemd320_hmac_ctx_t;
typedef struct ripemd320_ctx_vector
{
u32x h[10];
u32x w0[4];
u32x w1[4];
u32x w2[4];
u32x w3[4];
int len;
} ripemd320_ctx_vector_t;
typedef struct ripemd320_hmac_ctx_vector
{
ripemd320_ctx_vector_t ipad;
ripemd320_ctx_vector_t opad;
} ripemd320_hmac_ctx_vector_t;
DECLSPEC void ripemd320_transform (PRIVATE_AS const u32 *w0, PRIVATE_AS const u32 *w1, PRIVATE_AS const u32 *w2, PRIVATE_AS const u32 *w3, PRIVATE_AS u32 *digest);
DECLSPEC void ripemd320_init (PRIVATE_AS ripemd320_ctx_t *ctx);
DECLSPEC void ripemd320_update_64 (PRIVATE_AS ripemd320_ctx_t *ctx, PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, const int len);
DECLSPEC void ripemd320_update (PRIVATE_AS ripemd320_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void ripemd320_update_swap (PRIVATE_AS ripemd320_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void ripemd320_update_utf16le (PRIVATE_AS ripemd320_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void ripemd320_update_utf16le_swap (PRIVATE_AS ripemd320_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void ripemd320_update_global (PRIVATE_AS ripemd320_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void ripemd320_update_global_swap (PRIVATE_AS ripemd320_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void ripemd320_update_global_utf16le (PRIVATE_AS ripemd320_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void ripemd320_update_global_utf16le_swap (PRIVATE_AS ripemd320_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void ripemd320_final (PRIVATE_AS ripemd320_ctx_t *ctx);
DECLSPEC void ripemd320_hmac_init_64 (PRIVATE_AS ripemd320_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w0, PRIVATE_AS const u32 *w1, PRIVATE_AS const u32 *w2, PRIVATE_AS const u32 *w3);
DECLSPEC void ripemd320_hmac_init (PRIVATE_AS ripemd320_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void ripemd320_hmac_init_swap (PRIVATE_AS ripemd320_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void ripemd320_hmac_init_global (PRIVATE_AS ripemd320_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void ripemd320_hmac_init_global_swap (PRIVATE_AS ripemd320_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void ripemd320_hmac_update_64 (PRIVATE_AS ripemd320_hmac_ctx_t *ctx, PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, const int len);
DECLSPEC void ripemd320_hmac_update (PRIVATE_AS ripemd320_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void ripemd320_hmac_update_swap (PRIVATE_AS ripemd320_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void ripemd320_hmac_update_utf16le (PRIVATE_AS ripemd320_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void ripemd320_hmac_update_utf16le_swap (PRIVATE_AS ripemd320_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void ripemd320_hmac_update_global (PRIVATE_AS ripemd320_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void ripemd320_hmac_update_global_swap (PRIVATE_AS ripemd320_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void ripemd320_hmac_update_global_utf16le (PRIVATE_AS ripemd320_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void ripemd320_hmac_update_global_utf16le_swap (PRIVATE_AS ripemd320_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void ripemd320_hmac_final (PRIVATE_AS ripemd320_hmac_ctx_t *ctx);
DECLSPEC void ripemd320_transform_vector (PRIVATE_AS const u32x *w0, PRIVATE_AS const u32x *w1, PRIVATE_AS const u32x *w2, PRIVATE_AS const u32x *w3, PRIVATE_AS u32x *digest);
DECLSPEC void ripemd320_init_vector (PRIVATE_AS ripemd320_ctx_vector_t *ctx);
DECLSPEC void ripemd320_init_vector_from_scalar (PRIVATE_AS ripemd320_ctx_vector_t *ctx, PRIVATE_AS ripemd320_ctx_t *ctx0);
DECLSPEC void ripemd320_update_vector_64 (PRIVATE_AS ripemd320_ctx_vector_t *ctx, PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, const int len);
DECLSPEC void ripemd320_update_vector (PRIVATE_AS ripemd320_ctx_vector_t *ctx, PRIVATE_AS const u32x *w, const int len);
DECLSPEC void ripemd320_update_vector_swap (PRIVATE_AS ripemd320_ctx_vector_t *ctx, PRIVATE_AS const u32x *w, const int len);
DECLSPEC void ripemd320_update_vector_utf16le (PRIVATE_AS ripemd320_ctx_vector_t *ctx, PRIVATE_AS const u32x *w, const int len);
DECLSPEC void ripemd320_update_vector_utf16le_swap (PRIVATE_AS ripemd320_ctx_vector_t *ctx, PRIVATE_AS const u32x *w, const int len);
DECLSPEC void ripemd320_final_vector (PRIVATE_AS ripemd320_ctx_vector_t *ctx);
DECLSPEC void ripemd320_hmac_init_vector_64 (PRIVATE_AS ripemd320_hmac_ctx_vector_t *ctx, PRIVATE_AS const u32x *w0, PRIVATE_AS const u32x *w1, PRIVATE_AS const u32x *w2, PRIVATE_AS const u32x *w3);
DECLSPEC void ripemd320_hmac_init_vector (PRIVATE_AS ripemd320_hmac_ctx_vector_t *ctx, PRIVATE_AS const u32x *w, const int len);
DECLSPEC void ripemd320_hmac_update_vector_64 (PRIVATE_AS ripemd320_hmac_ctx_vector_t *ctx, PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, const int len);
DECLSPEC void ripemd320_hmac_update_vector (PRIVATE_AS ripemd320_hmac_ctx_vector_t *ctx, PRIVATE_AS const u32x *w, const int len);
DECLSPEC void ripemd320_hmac_final_vector (PRIVATE_AS ripemd320_hmac_ctx_vector_t *ctx);
#endif // INC_HASH_RIPEMD320_H

523
OpenCL/inc_hash_scrypt.cl Normal file
View File

@ -0,0 +1,523 @@
/**
* Author......: See docs/credits.txt
* License.....: MIT
*/
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_platform.h"
#include "inc_common.h"
#include "inc_hash_scrypt.h"
DECLSPEC hc_uint4_t xor_uint4 (const hc_uint4_t a, const hc_uint4_t b)
{
hc_uint4_t r;
r.x = a.x ^ b.x;
r.y = a.y ^ b.y;
r.z = a.z ^ b.z;
r.w = a.w ^ b.w;
return r;
}
#if SCRYPT_R > 1
DECLSPEC void scrypt_shuffle (PRIVATE_AS u32 *TI)
{
u32 TT[STATE_CNT4 / 2];
for (int dst_off = 0, src_off = SALSA_CNT4; src_off < STATE_CNT4; dst_off += SALSA_CNT4, src_off += SALSA_CNT4 * 2)
{
for (int j = 0; j < SALSA_CNT4; j++) TT[dst_off + j] = TI[src_off + j];
}
for (int dst_off = SALSA_CNT4, src_off = SALSA_CNT4 * 2; src_off < STATE_CNT4; dst_off += SALSA_CNT4, src_off += SALSA_CNT4 * 2)
{
for (int j = 0; j < SALSA_CNT4; j++) TI[dst_off + j] = TI[src_off + j];
}
for (int dst_off = STATE_CNT4 / 2, src_off = 0; dst_off < STATE_CNT4; dst_off += SALSA_CNT4, src_off += SALSA_CNT4)
{
for (int j = 0; j < SALSA_CNT4; j++) TI[dst_off + j] = TT[src_off + j];
}
}
#endif
DECLSPEC void salsa_r (PRIVATE_AS u32 *TI)
{
u32 TT[16];
for (int j = 0; j < SALSA_CNT4; j++) TT[j] = TI[STATE_CNT4 - 16 + j];
for (int i = 0; i < STATE_CNT4; i += SALSA_CNT4)
{
for (int j = 0; j < SALSA_CNT4; j++) TT[j] ^= TI[i + j];
for (int j = 0; j < SALSA_CNT4; j++) TI[i + j] = TT[j];
for (int r = 0; r < 4; r++)
{
u32 t0, t1, t2, t3;
t0 = TT[ 0] + TT[12];
t1 = TT[ 1] + TT[13];
t2 = TT[ 2] + TT[14];
t3 = TT[ 3] + TT[15];
TT[ 4] ^= hc_rotl32_S (t0, 7);
TT[ 5] ^= hc_rotl32_S (t1, 7);
TT[ 6] ^= hc_rotl32_S (t2, 7);
TT[ 7] ^= hc_rotl32_S (t3, 7);
t0 = TT[ 4] + TT[ 0];
t1 = TT[ 5] + TT[ 1];
t2 = TT[ 6] + TT[ 2];
t3 = TT[ 7] + TT[ 3];
TT[ 8] ^= hc_rotl32_S (t0, 9);
TT[ 9] ^= hc_rotl32_S (t1, 9);
TT[10] ^= hc_rotl32_S (t2, 9);
TT[11] ^= hc_rotl32_S (t3, 9);
t0 = TT[ 8] + TT[ 4];
t1 = TT[ 9] + TT[ 5];
t2 = TT[10] + TT[ 6];
t3 = TT[11] + TT[ 7];
TT[12] ^= hc_rotl32_S (t0, 13);
TT[13] ^= hc_rotl32_S (t1, 13);
TT[14] ^= hc_rotl32_S (t2, 13);
TT[15] ^= hc_rotl32_S (t3, 13);
t0 = TT[12] + TT[ 8];
t1 = TT[13] + TT[ 9];
t2 = TT[14] + TT[10];
t3 = TT[15] + TT[11];
TT[ 0] ^= hc_rotl32_S (t0, 18);
TT[ 1] ^= hc_rotl32_S (t1, 18);
TT[ 2] ^= hc_rotl32_S (t2, 18);
TT[ 3] ^= hc_rotl32_S (t3, 18);
t0 = TT[ 4]; TT[ 4] = TT[ 7]; TT[ 7] = TT[ 6]; TT[ 6] = TT[ 5]; TT[ 5] = t0;
t0 = TT[ 8]; TT[ 8] = TT[10]; TT[10] = t0;
t0 = TT[ 9]; TT[ 9] = TT[11]; TT[11] = t0;
t0 = TT[12]; TT[12] = TT[13]; TT[13] = TT[14]; TT[14] = TT[15]; TT[15] = t0;
t0 = TT[ 0] + TT[ 4];
t1 = TT[ 1] + TT[ 5];
t2 = TT[ 2] + TT[ 6];
t3 = TT[ 3] + TT[ 7];
TT[12] ^= hc_rotl32_S (t0, 7);
TT[13] ^= hc_rotl32_S (t1, 7);
TT[14] ^= hc_rotl32_S (t2, 7);
TT[15] ^= hc_rotl32_S (t3, 7);
t0 = TT[12] + TT[ 0];
t1 = TT[13] + TT[ 1];
t2 = TT[14] + TT[ 2];
t3 = TT[15] + TT[ 3];
TT[ 8] ^= hc_rotl32_S (t0, 9);
TT[ 9] ^= hc_rotl32_S (t1, 9);
TT[10] ^= hc_rotl32_S (t2, 9);
TT[11] ^= hc_rotl32_S (t3, 9);
t0 = TT[ 8] + TT[12];
t1 = TT[ 9] + TT[13];
t2 = TT[10] + TT[14];
t3 = TT[11] + TT[15];
TT[ 4] ^= hc_rotl32_S (t0, 13);
TT[ 5] ^= hc_rotl32_S (t1, 13);
TT[ 6] ^= hc_rotl32_S (t2, 13);
TT[ 7] ^= hc_rotl32_S (t3, 13);
t0 = TT[ 4] + TT[ 8];
t1 = TT[ 5] + TT[ 9];
t2 = TT[ 6] + TT[10];
t3 = TT[ 7] + TT[11];
TT[ 0] ^= hc_rotl32_S (t0, 18);
TT[ 1] ^= hc_rotl32_S (t1, 18);
TT[ 2] ^= hc_rotl32_S (t2, 18);
TT[ 3] ^= hc_rotl32_S (t3, 18);
t0 = TT[ 4]; TT[ 4] = TT[ 5]; TT[ 5] = TT[ 6]; TT[ 6] = TT[ 7]; TT[ 7] = t0;
t0 = TT[ 8]; TT[ 8] = TT[10]; TT[10] = t0;
t0 = TT[ 9]; TT[ 9] = TT[11]; TT[11] = t0;
t0 = TT[15]; TT[15] = TT[14]; TT[14] = TT[13]; TT[13] = TT[12]; TT[12] = t0;
}
for (int j = 0; j < SALSA_CNT4; j++) TT[j] += TI[i + j];
for (int j = 0; j < SALSA_CNT4; j++) TI[i + j] = TT[j];
}
}
DECLSPEC void scrypt_smix_init (GLOBAL_AS u32 *P, PRIVATE_AS u32 *X, GLOBAL_AS void *V0, GLOBAL_AS void *V1, GLOBAL_AS void *V2, GLOBAL_AS void *V3, const u32 gid, const u32 lid, const u32 lsz, const u32 bid)
{
const u32 ySIZE = SCRYPT_N >> SCRYPT_TMTO;
const u32 zSIZE = STATE_CNT44;
const u32 xd4 = bid / 4;
const u32 xm4 = bid & 3;
PRIVATE_AS hc_uint4_t *X4 = (PRIVATE_AS hc_uint4_t *) X;
GLOBAL_AS hc_uint4_t *V;
switch (xm4)
{
case 0: V = (GLOBAL_AS hc_uint4_t *) ALIGN_PTR_1k (V0); break;
case 1: V = (GLOBAL_AS hc_uint4_t *) ALIGN_PTR_1k (V1); break;
case 2: V = (GLOBAL_AS hc_uint4_t *) ALIGN_PTR_1k (V2); break;
case 3: V = (GLOBAL_AS hc_uint4_t *) ALIGN_PTR_1k (V3); break;
}
GLOBAL_AS hc_uint4_t *Vx = V + (xd4 * lsz * ySIZE * zSIZE) + (lid * ySIZE * zSIZE);
for (u32 i = 0; i < STATE_CNT4; i++) X[i] = P[i];
for (u32 y = 0; y < ySIZE; y++)
{
GLOBAL_AS hc_uint4_t *Vxx = Vx + (y * zSIZE);
for (u32 z = 0; z < zSIZE; z++) Vxx[z] = X4[z];
for (u32 i = 0; i < (1 << SCRYPT_TMTO); i++)
{
salsa_r (X);
#if SCRYPT_R > 1
scrypt_shuffle (X);
#endif
}
}
for (u32 i = 0; i < STATE_CNT4; i++) P[i] = X[i];
}
DECLSPEC void scrypt_smix_loop (GLOBAL_AS u32 *P, PRIVATE_AS u32 *X, PRIVATE_AS u32 *T, GLOBAL_AS void *V0, GLOBAL_AS void *V1, GLOBAL_AS void *V2, GLOBAL_AS void *V3, const u32 gid, const u32 lid, const u32 lsz, const u32 bid)
{
const u32 ySIZE = SCRYPT_N >> SCRYPT_TMTO;
const u32 zSIZE = STATE_CNT44;
const u32 xd4 = bid / 4;
const u32 xm4 = bid & 3;
PRIVATE_AS hc_uint4_t *X4 = (PRIVATE_AS hc_uint4_t *) X;
PRIVATE_AS hc_uint4_t *T4 = (PRIVATE_AS hc_uint4_t *) T;
GLOBAL_AS hc_uint4_t *V;
switch (xm4)
{
case 0: V = (GLOBAL_AS hc_uint4_t *) ALIGN_PTR_1k (V0); break;
case 1: V = (GLOBAL_AS hc_uint4_t *) ALIGN_PTR_1k (V1); break;
case 2: V = (GLOBAL_AS hc_uint4_t *) ALIGN_PTR_1k (V2); break;
case 3: V = (GLOBAL_AS hc_uint4_t *) ALIGN_PTR_1k (V3); break;
}
GLOBAL_AS hc_uint4_t *Vx = V + (xd4 * lsz * ySIZE * zSIZE) + (lid * ySIZE * zSIZE);
for (u32 i = 0; i < STATE_CNT4; i++) X[i] = P[i];
// note: max 1024 iterations = forced -u 2048
const u32 N_max = (SCRYPT_N < 2048) ? SCRYPT_N : 2048;
for (u32 N_pos = 0; N_pos < N_max; N_pos++)
{
const u32 k = X4[zSIZE - 4].x & (SCRYPT_N - 1);
const u32 y = k >> SCRYPT_TMTO;
const u32 km = k - (y << SCRYPT_TMTO);
GLOBAL_AS hc_uint4_t *Vxx = Vx + (y * zSIZE);
for (u32 z = 0; z < zSIZE; z++) T4[z] = *Vxx++;
for (u32 i = 0; i < km; i++)
{
salsa_r (T);
#if SCRYPT_R > 1
scrypt_shuffle (T);
#endif
}
for (u32 z = 0; z < zSIZE; z++) X4[z] = xor_uint4 (X4[z], T4[z]);
salsa_r (X);
#if SCRYPT_R > 1
scrypt_shuffle (X);
#endif
}
for (u32 i = 0; i < STATE_CNT4; i++) P[i] = X[i];
}
DECLSPEC void scrypt_blockmix_in (GLOBAL_AS u32 *in_buf, GLOBAL_AS u32 *out_buf, const int out_len)
{
for (int i = 0, j = 0; i < out_len; i += SALSA_SZ, j += SALSA_CNT4)
{
u32 X[SALSA_CNT4];
X[ 0] = in_buf[j + 0];
X[ 1] = in_buf[j + 5];
X[ 2] = in_buf[j + 10];
X[ 3] = in_buf[j + 15];
X[ 4] = in_buf[j + 4];
X[ 5] = in_buf[j + 9];
X[ 6] = in_buf[j + 14];
X[ 7] = in_buf[j + 3];
X[ 8] = in_buf[j + 8];
X[ 9] = in_buf[j + 13];
X[10] = in_buf[j + 2];
X[11] = in_buf[j + 7];
X[12] = in_buf[j + 12];
X[13] = in_buf[j + 1];
X[14] = in_buf[j + 6];
X[15] = in_buf[j + 11];
out_buf[j + 0] = X[ 0];
out_buf[j + 1] = X[ 1];
out_buf[j + 2] = X[ 2];
out_buf[j + 3] = X[ 3];
out_buf[j + 4] = X[ 4];
out_buf[j + 5] = X[ 5];
out_buf[j + 6] = X[ 6];
out_buf[j + 7] = X[ 7];
out_buf[j + 8] = X[ 8];
out_buf[j + 9] = X[ 9];
out_buf[j + 10] = X[10];
out_buf[j + 11] = X[11];
out_buf[j + 12] = X[12];
out_buf[j + 13] = X[13];
out_buf[j + 14] = X[14];
out_buf[j + 15] = X[15];
}
}
DECLSPEC void scrypt_blockmix_out (GLOBAL_AS u32 *in_buf, GLOBAL_AS u32 *out_buf, const int out_len)
{
for (int i = 0, j = 0; i < out_len; i += SALSA_SZ, j += SALSA_CNT4)
{
u32 T[SALSA_CNT4];
T[ 0] = in_buf[j + 0];
T[ 1] = in_buf[j + 13];
T[ 2] = in_buf[j + 10];
T[ 3] = in_buf[j + 7];
T[ 4] = in_buf[j + 4];
T[ 5] = in_buf[j + 1];
T[ 6] = in_buf[j + 14];
T[ 7] = in_buf[j + 11];
T[ 8] = in_buf[j + 8];
T[ 9] = in_buf[j + 5];
T[10] = in_buf[j + 2];
T[11] = in_buf[j + 15];
T[12] = in_buf[j + 12];
T[13] = in_buf[j + 9];
T[14] = in_buf[j + 6];
T[15] = in_buf[j + 3];
out_buf[j + 0] = T[ 0];
out_buf[j + 1] = T[ 1];
out_buf[j + 2] = T[ 2];
out_buf[j + 3] = T[ 3];
out_buf[j + 4] = T[ 4];
out_buf[j + 5] = T[ 5];
out_buf[j + 6] = T[ 6];
out_buf[j + 7] = T[ 7];
out_buf[j + 8] = T[ 8];
out_buf[j + 9] = T[ 9];
out_buf[j + 10] = T[10];
out_buf[j + 11] = T[11];
out_buf[j + 12] = T[12];
out_buf[j + 13] = T[13];
out_buf[j + 14] = T[14];
out_buf[j + 15] = T[15];
}
}
DECLSPEC void scrypt_pbkdf2_body_pp (PRIVATE_AS sha256_hmac_ctx_t *sha256_hmac_ctx, PRIVATE_AS u32 *out_buf, const int out_len)
{
for (int i = 0, j = 1, k = 0; i < out_len; i += 32, j += 1, k += 8)
{
sha256_hmac_ctx_t sha256_hmac_ctx2 = *sha256_hmac_ctx;
u32 w0[4];
u32 w1[4];
u32 w2[4];
u32 w3[4];
w0[0] = j;
w0[1] = 0;
w0[2] = 0;
w0[3] = 0;
w1[0] = 0;
w1[1] = 0;
w1[2] = 0;
w1[3] = 0;
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
w3[3] = 0;
sha256_hmac_update_64 (&sha256_hmac_ctx2, w0, w1, w2, w3, 4);
sha256_hmac_final (&sha256_hmac_ctx2);
// this will not work if user specifies output length not a multiple of 4
// probably never happens...
// let's hope the compiler will auto optimize this since out_len is very likely
// a constant at caller level
if (out_len >= (i + 4)) out_buf[k + 0] = hc_swap32_S (sha256_hmac_ctx2.opad.h[0]);
if (out_len >= (i + 8)) out_buf[k + 1] = hc_swap32_S (sha256_hmac_ctx2.opad.h[1]);
if (out_len >= (i + 12)) out_buf[k + 2] = hc_swap32_S (sha256_hmac_ctx2.opad.h[2]);
if (out_len >= (i + 16)) out_buf[k + 3] = hc_swap32_S (sha256_hmac_ctx2.opad.h[3]);
if (out_len >= (i + 20)) out_buf[k + 4] = hc_swap32_S (sha256_hmac_ctx2.opad.h[4]);
if (out_len >= (i + 24)) out_buf[k + 5] = hc_swap32_S (sha256_hmac_ctx2.opad.h[5]);
if (out_len >= (i + 28)) out_buf[k + 6] = hc_swap32_S (sha256_hmac_ctx2.opad.h[6]);
if (out_len >= (i + 32)) out_buf[k + 7] = hc_swap32_S (sha256_hmac_ctx2.opad.h[7]);
}
}
DECLSPEC void scrypt_pbkdf2_body_pg (PRIVATE_AS sha256_hmac_ctx_t *sha256_hmac_ctx, GLOBAL_AS u32 *out_buf, const int out_len)
{
for (int i = 0, j = 1, k = 0; i < out_len; i += 32, j += 1, k += 8)
{
sha256_hmac_ctx_t sha256_hmac_ctx2 = *sha256_hmac_ctx;
u32 w0[4];
u32 w1[4];
u32 w2[4];
u32 w3[4];
w0[0] = j;
w0[1] = 0;
w0[2] = 0;
w0[3] = 0;
w1[0] = 0;
w1[1] = 0;
w1[2] = 0;
w1[3] = 0;
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
w3[3] = 0;
sha256_hmac_update_64 (&sha256_hmac_ctx2, w0, w1, w2, w3, 4);
sha256_hmac_final (&sha256_hmac_ctx2);
// this will not work if user specifies output length not a multiple of 4
// probably never happens...
// let's hope the compiler will auto optimize this since out_len is very likely
// a constant at caller level
if (out_len >= (i + 4)) out_buf[k + 0] = hc_swap32_S (sha256_hmac_ctx2.opad.h[0]);
if (out_len >= (i + 8)) out_buf[k + 1] = hc_swap32_S (sha256_hmac_ctx2.opad.h[1]);
if (out_len >= (i + 12)) out_buf[k + 2] = hc_swap32_S (sha256_hmac_ctx2.opad.h[2]);
if (out_len >= (i + 16)) out_buf[k + 3] = hc_swap32_S (sha256_hmac_ctx2.opad.h[3]);
if (out_len >= (i + 20)) out_buf[k + 4] = hc_swap32_S (sha256_hmac_ctx2.opad.h[4]);
if (out_len >= (i + 24)) out_buf[k + 5] = hc_swap32_S (sha256_hmac_ctx2.opad.h[5]);
if (out_len >= (i + 28)) out_buf[k + 6] = hc_swap32_S (sha256_hmac_ctx2.opad.h[6]);
if (out_len >= (i + 32)) out_buf[k + 7] = hc_swap32_S (sha256_hmac_ctx2.opad.h[7]);
}
}
DECLSPEC void scrypt_pbkdf2_ppp (PRIVATE_AS const u32 *pw_buf, const int pw_len, PRIVATE_AS const u32 *salt_buf, const int salt_len, PRIVATE_AS u32 *out_buf, const int out_len)
{
sha256_hmac_ctx_t sha256_hmac_ctx;
sha256_hmac_init_swap (&sha256_hmac_ctx, pw_buf, pw_len);
sha256_hmac_update_swap (&sha256_hmac_ctx, salt_buf, salt_len);
scrypt_pbkdf2_body_pp (&sha256_hmac_ctx, out_buf, out_len);
}
DECLSPEC void scrypt_pbkdf2_pgp (PRIVATE_AS const u32 *pw_buf, const int pw_len, GLOBAL_AS const u32 *salt_buf, const int salt_len, PRIVATE_AS u32 *out_buf, const int out_len)
{
sha256_hmac_ctx_t sha256_hmac_ctx;
sha256_hmac_init_swap (&sha256_hmac_ctx, pw_buf, pw_len);
sha256_hmac_update_global_swap (&sha256_hmac_ctx, salt_buf, salt_len);
scrypt_pbkdf2_body_pp (&sha256_hmac_ctx, out_buf, out_len);
}
DECLSPEC void scrypt_pbkdf2_gpp (GLOBAL_AS const u32 *pw_buf, const int pw_len, PRIVATE_AS const u32 *salt_buf, const int salt_len, PRIVATE_AS u32 *out_buf, const int out_len)
{
sha256_hmac_ctx_t sha256_hmac_ctx;
sha256_hmac_init_global_swap (&sha256_hmac_ctx, pw_buf, pw_len);
sha256_hmac_update_swap (&sha256_hmac_ctx, salt_buf, salt_len);
scrypt_pbkdf2_body_pp (&sha256_hmac_ctx, out_buf, out_len);
}
DECLSPEC void scrypt_pbkdf2_ggp (GLOBAL_AS const u32 *pw_buf, const int pw_len, GLOBAL_AS const u32 *salt_buf, const int salt_len, PRIVATE_AS u32 *out_buf, const int out_len)
{
sha256_hmac_ctx_t sha256_hmac_ctx;
sha256_hmac_init_global_swap (&sha256_hmac_ctx, pw_buf, pw_len);
sha256_hmac_update_global_swap (&sha256_hmac_ctx, salt_buf, salt_len);
scrypt_pbkdf2_body_pp (&sha256_hmac_ctx, out_buf, out_len);
}
DECLSPEC void scrypt_pbkdf2_ppg (PRIVATE_AS const u32 *pw_buf, const int pw_len, PRIVATE_AS const u32 *salt_buf, const int salt_len, GLOBAL_AS u32 *out_buf, const int out_len)
{
sha256_hmac_ctx_t sha256_hmac_ctx;
sha256_hmac_init_swap (&sha256_hmac_ctx, pw_buf, pw_len);
sha256_hmac_update_swap (&sha256_hmac_ctx, salt_buf, salt_len);
scrypt_pbkdf2_body_pg (&sha256_hmac_ctx, out_buf, out_len);
}
DECLSPEC void scrypt_pbkdf2_pgg (PRIVATE_AS const u32 *pw_buf, const int pw_len, GLOBAL_AS const u32 *salt_buf, const int salt_len, GLOBAL_AS u32 *out_buf, const int out_len)
{
sha256_hmac_ctx_t sha256_hmac_ctx;
sha256_hmac_init_swap (&sha256_hmac_ctx, pw_buf, pw_len);
sha256_hmac_update_global_swap (&sha256_hmac_ctx, salt_buf, salt_len);
scrypt_pbkdf2_body_pg (&sha256_hmac_ctx, out_buf, out_len);
}
DECLSPEC void scrypt_pbkdf2_gpg (GLOBAL_AS const u32 *pw_buf, const int pw_len, PRIVATE_AS const u32 *salt_buf, const int salt_len, GLOBAL_AS u32 *out_buf, const int out_len)
{
sha256_hmac_ctx_t sha256_hmac_ctx;
sha256_hmac_init_global_swap (&sha256_hmac_ctx, pw_buf, pw_len);
sha256_hmac_update_swap (&sha256_hmac_ctx, salt_buf, salt_len);
scrypt_pbkdf2_body_pg (&sha256_hmac_ctx, out_buf, out_len);
}
DECLSPEC void scrypt_pbkdf2_ggg (GLOBAL_AS const u32 *pw_buf, const int pw_len, GLOBAL_AS const u32 *salt_buf, const int salt_len, GLOBAL_AS u32 *out_buf, const int out_len)
{
sha256_hmac_ctx_t sha256_hmac_ctx;
sha256_hmac_init_global_swap (&sha256_hmac_ctx, pw_buf, pw_len);
sha256_hmac_update_global_swap (&sha256_hmac_ctx, salt_buf, salt_len);
scrypt_pbkdf2_body_pg (&sha256_hmac_ctx, out_buf, out_len);
}

60
OpenCL/inc_hash_scrypt.h Normal file
View File

@ -0,0 +1,60 @@
/**
* Author......: See docs/credits.txt
* License.....: MIT
*/
#ifndef INC_HASH_SCRYPT_H
#define INC_HASH_SCRYPT_H
#define GET_SCRYPT_SZ(r,p) (128 * (r) * (p))
#define GET_STATE_SZ(r) (128 * (r))
// _SZ is true sizes as bytes
#define SCRYPT_SZ GET_SCRYPT_SZ (SCRYPT_R, SCRYPT_P)
#define STATE_SZ GET_STATE_SZ (SCRYPT_R)
// _CNT is size as whatever /X datatype
#define SCRYPT_CNT4 (SCRYPT_SZ / 4)
#define STATE_CNT4 (STATE_SZ / 4)
// this would be uint4, feels more natural than 16
#define SCRYPT_CNT44 ((SCRYPT_SZ / 4) / 4)
#define STATE_CNT44 ((STATE_SZ / 4) / 4)
#define SALSA_SZ 64
#define SALSA_CNT4 (SALSA_SZ / 4)
#define SALSA_CNT44 ((SALSA_SZ / 4) / 4)
// should be safe, because in backend.c we use:
// u64 size_extra_buffer1 = 4096;
// size_extra_buffer1 += base_chunk_size;
#define ALIGN_PTR_1k(p) ((GLOBAL_AS hc_uint4_t *) (((u64) (p) + 1023) & ~1023UL))
#if defined IS_INTEL_SDK
typedef struct
{
u32 x, y, z, w;
} hc_uint4_t;
#else
typedef uint4 hc_uint4_t;
#endif
DECLSPEC hc_uint4_t xor_uint4 (const hc_uint4_t a, const hc_uint4_t b);
typedef struct
{
#ifndef SCRYPT_TMP_ELEM
#define SCRYPT_TMP_ELEM 1
#endif
u32 in[SCRYPT_TMP_ELEM / 2];
u32 out[SCRYPT_TMP_ELEM / 2];
} scrypt_tmp_t;
#endif

146
OpenCL/inc_hash_sha256.cl
View File

@ -1245,6 +1245,152 @@ DECLSPEC void sha256_hmac_init_global_swap (PRIVATE_AS sha256_hmac_ctx_t *ctx, G
sha256_hmac_init_64 (ctx, w0, w1, w2, w3);
}
DECLSPEC void sha256_hmac_init_global_utf16le_swap (PRIVATE_AS sha256_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len)
{
if (hc_enc_scan_global (w, len))
{
hc_enc_t hc_enc;
hc_enc_init (&hc_enc);
while (hc_enc_has_next (&hc_enc, len))
{
// forced full decode in one round
u32 enc_buf[256];
const int enc_len = hc_enc_next_global (&hc_enc, w, len, 256, enc_buf, sizeof (enc_buf));
if (enc_len == -1)
{
//hmac doesn't have password length
//ctx->len = -1;
return;
}
if (enc_len > 64)
{
sha256_ctx_t tmp;
sha256_init (&tmp);
sha256_update_utf16le_swap (&tmp, enc_buf, enc_len);
sha256_final (&tmp);
enc_buf[ 0] = tmp.h[0];
enc_buf[ 1] = tmp.h[1];
enc_buf[ 2] = tmp.h[2];
enc_buf[ 3] = tmp.h[3];
enc_buf[ 4] = tmp.h[4];
enc_buf[ 5] = tmp.h[5];
enc_buf[ 6] = tmp.h[6];
enc_buf[ 7] = tmp.h[7];
enc_buf[ 8] = 0;
enc_buf[ 9] = 0;
enc_buf[10] = 0;
enc_buf[11] = 0;
enc_buf[12] = 0;
enc_buf[13] = 0;
enc_buf[14] = 0;
enc_buf[15] = 0;
}
else
{
enc_buf[ 0] = hc_swap32_S (enc_buf[ 0]);
enc_buf[ 1] = hc_swap32_S (enc_buf[ 1]);
enc_buf[ 2] = hc_swap32_S (enc_buf[ 2]);
enc_buf[ 3] = hc_swap32_S (enc_buf[ 3]);
enc_buf[ 4] = hc_swap32_S (enc_buf[ 4]);
enc_buf[ 5] = hc_swap32_S (enc_buf[ 5]);
enc_buf[ 6] = hc_swap32_S (enc_buf[ 6]);
enc_buf[ 7] = hc_swap32_S (enc_buf[ 7]);
enc_buf[ 8] = hc_swap32_S (enc_buf[ 8]);
enc_buf[ 9] = hc_swap32_S (enc_buf[ 9]);
enc_buf[10] = hc_swap32_S (enc_buf[10]);
enc_buf[11] = hc_swap32_S (enc_buf[11]);
enc_buf[12] = hc_swap32_S (enc_buf[12]);
enc_buf[13] = hc_swap32_S (enc_buf[13]);
enc_buf[14] = hc_swap32_S (enc_buf[14]);
enc_buf[15] = hc_swap32_S (enc_buf[15]);
}
sha256_hmac_init_64 (ctx, enc_buf + 0, enc_buf + 4, enc_buf + 8, enc_buf + 12);
}
return;
}
u32 w0[4];
u32 w1[4];
u32 w2[4];
u32 w3[4];
const int len_new = len * 2;
if (len_new > 64)
{
sha256_ctx_t tmp;
sha256_init (&tmp);
sha256_update_global_utf16le_swap (&tmp, w, len);
sha256_final (&tmp);
w0[0] = tmp.h[0];
w0[1] = tmp.h[1];
w0[2] = tmp.h[2];
w0[3] = tmp.h[3];
w1[0] = tmp.h[4];
w1[1] = tmp.h[5];
w1[2] = tmp.h[6];
w1[3] = tmp.h[7];
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
w3[3] = 0;
}
else
{
w0[0] = w[0];
w0[1] = w[1];
w0[2] = w[2];
w0[3] = w[3];
w1[0] = w[4];
w1[1] = w[5];
w1[2] = w[6];
w1[3] = w[7];
make_utf16le_S (w1, w2, w3);
make_utf16le_S (w0, w0, w1);
w0[0] = hc_swap32_S (w0[0]);
w0[1] = hc_swap32_S (w0[1]);
w0[2] = hc_swap32_S (w0[2]);
w0[3] = hc_swap32_S (w0[3]);
w1[0] = hc_swap32_S (w1[0]);
w1[1] = hc_swap32_S (w1[1]);
w1[2] = hc_swap32_S (w1[2]);
w1[3] = hc_swap32_S (w1[3]);
w2[0] = hc_swap32_S (w2[0]);
w2[1] = hc_swap32_S (w2[1]);
w2[2] = hc_swap32_S (w2[2]);
w2[3] = hc_swap32_S (w2[3]);
w3[0] = hc_swap32_S (w3[0]);
w3[1] = hc_swap32_S (w3[1]);
w3[2] = hc_swap32_S (w3[2]);
w3[3] = hc_swap32_S (w3[3]);
}
sha256_hmac_init_64 (ctx, w0, w1, w2, w3);
}
DECLSPEC void sha256_hmac_update_64 (PRIVATE_AS sha256_hmac_ctx_t *ctx, PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, const int len)
{
sha256_update_64 (&ctx->ipad, w0, w1, w2, w3, len);

1
OpenCL/inc_hash_sha256.h
View File

@ -106,6 +106,7 @@ DECLSPEC void sha256_hmac_init (PRIVATE_AS sha256_hmac_ctx_t *ctx, PRIVATE_AS co
DECLSPEC void sha256_hmac_init_swap (PRIVATE_AS sha256_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void sha256_hmac_init_global (PRIVATE_AS sha256_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void sha256_hmac_init_global_swap (PRIVATE_AS sha256_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void sha256_hmac_init_global_utf16le_swap (PRIVATE_AS sha256_hmac_ctx_t *ctx, GLOBAL_AS const u32 *w, const int len);
DECLSPEC void sha256_hmac_update_64 (PRIVATE_AS sha256_hmac_ctx_t *ctx, PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, const int len);
DECLSPEC void sha256_hmac_update (PRIVATE_AS sha256_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);
DECLSPEC void sha256_hmac_update_swap (PRIVATE_AS sha256_hmac_ctx_t *ctx, PRIVATE_AS const u32 *w, const int len);

6
OpenCL/inc_markov.h
View File

@ -19,7 +19,7 @@
CONSTANT_AS const u32 &bits14, \
CONSTANT_AS const u32 &bits15, \
CONSTANT_AS const u64 &gid_max, \
uint hc_gid [[ thread_position_in_grid ]]
uint3 hc_gid [[ thread_position_in_grid ]]
#define KERN_ATTR_R_MARKOV \
GLOBAL_AS bf_t *pws_buf_r, \
@ -31,7 +31,7 @@
CONSTANT_AS const u32 &bits14, \
CONSTANT_AS const u32 &bits15, \
CONSTANT_AS const u64 &gid_max, \
uint hc_gid [[ thread_position_in_grid ]]
uint3 hc_gid [[ thread_position_in_grid ]]
#define KERN_ATTR_C_MARKOV \
GLOBAL_AS pw_t *pws_buf, \
@ -43,7 +43,7 @@
CONSTANT_AS const u32 &bits14, \
CONSTANT_AS const u32 &bits15, \
CONSTANT_AS const u64 &gid_max, \
uint hc_gid [[ thread_position_in_grid ]]
uint3 hc_gid [[ thread_position_in_grid ]]
#else // CUDA, HIP, OpenCL

309
OpenCL/inc_platform.cl
View File

@ -6,6 +6,7 @@
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_platform.h"
#include "inc_common.h"
#ifdef IS_NATIVE
#define FIXED_THREAD_COUNT(n)
@ -60,6 +61,33 @@ DECLSPEC u64 rotr64_S (const u64 a, const int n)
#endif // IS_AMD
// this applies to cuda and opencl
#if defined IS_NV
#ifdef USE_FUNNELSHIFT
DECLSPEC u32 hc_funnelshift_l (const u32 lo, const u32 hi, const int shift)
{
u32 result;
asm volatile ("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result) : "r"(lo), "r"(hi), "r"(shift));
return result;
}
DECLSPEC u32 hc_funnelshift_r (const u32 lo, const u32 hi, const int shift)
{
u32 result;
asm volatile ("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(result) : "r"(lo), "r"(hi), "r"(shift));
return result;
}
#endif
#endif // IS_NV
#if defined IS_CUDA
#if ATTACK_EXEC == 11
@ -104,55 +132,181 @@ DECLSPEC u32 hc_atomic_or (GLOBAL_AS u32 *p, volatile const u32 val)
return atomicOr (p, val);
}
DECLSPEC size_t get_group_id (const u32 dimindx)
{
switch (dimindx)
{
case 0:
return blockIdx.x;
case 1:
return blockIdx.y;
case 2:
return blockIdx.z;
}
return (size_t) -1;
}
DECLSPEC size_t get_global_id (const u32 dimindx __attribute__((unused)))
{
return (blockIdx.x * blockDim.x) + threadIdx.x;
}
DECLSPEC size_t get_local_id (const u32 dimindx __attribute__((unused)))
DECLSPEC size_t get_local_id (const u32 dimindx)
{
switch (dimindx)
{
case 0:
return threadIdx.x;
case 1:
return threadIdx.y;
case 2:
return threadIdx.z;
}
return (size_t) -1;
}
DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)))
DECLSPEC size_t get_local_size (const u32 dimindx)
{
// verify
switch (dimindx)
{
case 0:
return blockDim.x;
case 1:
return blockDim.y;
case 2:
return blockDim.z;
}
return (size_t) -1;
}
DECLSPEC u32x rotl32 (const u32x a, const int n)
{
return ((a << n) | ((a >> (32 - n))));
#if VECT_SIZE == 1
return rotl32_S (a, n);
#else
u32x t = 0;
#if VECT_SIZE >= 2
t.s0 = rotl32_S (a.s0, n);
t.s1 = rotl32_S (a.s1, n);
#endif
#if VECT_SIZE >= 4
t.s2 = rotl32_S (a.s2, n);
t.s3 = rotl32_S (a.s3, n);
#endif
#if VECT_SIZE >= 8
t.s4 = rotl32_S (a.s4, n);
t.s5 = rotl32_S (a.s5, n);
t.s6 = rotl32_S (a.s6, n);
t.s7 = rotl32_S (a.s7, n);
#endif
#if VECT_SIZE >= 16
t.s8 = rotl32_S (a.s8, n);
t.s9 = rotl32_S (a.s9, n);
t.sa = rotl32_S (a.sa, n);
t.sb = rotl32_S (a.sb, n);
t.sc = rotl32_S (a.sc, n);
t.sd = rotl32_S (a.sd, n);
t.se = rotl32_S (a.se, n);
t.sf = rotl32_S (a.sf, n);
#endif
return t;
#endif
}
DECLSPEC u32x rotr32 (const u32x a, const int n)
{
return ((a >> n) | ((a << (32 - n))));
#if VECT_SIZE == 1
return rotr32_S (a, n);
#else
u32x t = 0;
#if VECT_SIZE >= 2
t.s0 = rotr32_S (a.s0, n);
t.s1 = rotr32_S (a.s1, n);
#endif
#if VECT_SIZE >= 4
t.s2 = rotr32_S (a.s2, n);
t.s3 = rotr32_S (a.s3, n);
#endif
#if VECT_SIZE >= 8
t.s4 = rotr32_S (a.s4, n);
t.s5 = rotr32_S (a.s5, n);
t.s6 = rotr32_S (a.s6, n);
t.s7 = rotr32_S (a.s7, n);
#endif
#if VECT_SIZE >= 16
t.s8 = rotr32_S (a.s8, n);
t.s9 = rotr32_S (a.s9, n);
t.sa = rotr32_S (a.sa, n);
t.sb = rotr32_S (a.sb, n);
t.sc = rotr32_S (a.sc, n);
t.sd = rotr32_S (a.sd, n);
t.se = rotr32_S (a.se, n);
t.sf = rotr32_S (a.sf, n);
#endif
return t;
#endif
}
DECLSPEC u32 rotl32_S (const u32 a, const int n)
{
#ifdef USE_FUNNELSHIFT
return hc_funnelshift_l (a, a, n);
#else
return ((a << n) | ((a >> (32 - n))));
#endif
}
DECLSPEC u32 rotr32_S (const u32 a, const int n)
{
#ifdef USE_FUNNELSHIFT
return hc_funnelshift_r (a, a, n);
#else
return ((a >> n) | ((a << (32 - n))));
#endif
}
DECLSPEC u64x rotl64 (const u64x a, const int n)
{
#if VECT_SIZE == 1
return rotl64_S (a, n);
#else
return ((a << n) | ((a >> (64 - n))));
#endif
}
DECLSPEC u64x rotr64 (const u64x a, const int n)
{
#if VECT_SIZE == 1
return rotr64_S (a, n);
#else
return ((a >> n) | ((a << (64 - n))));
#endif
}
DECLSPEC u64 rotl64_S (const u64 a, const int n)
{
return ((a << n) | ((a >> (64 - n))));
return rotr64_S (a, 64 - n);
}
DECLSPEC u64 rotr64_S (const u64 a, const int n)
@ -208,40 +362,150 @@ DECLSPEC u32 hc_atomic_or (GLOBAL_AS u32 *p, volatile const u32 val)
return atomicOr (p, val);
}
DECLSPEC size_t get_group_id (const u32 dimindx)
{
switch (dimindx)
{
case 0:
return blockIdx.x;
case 1:
return blockIdx.y;
case 2:
return blockIdx.z;
}
return (size_t) -1;
}
DECLSPEC size_t get_global_id (const u32 dimindx __attribute__((unused)))
{
return (blockIdx.x * blockDim.x) + threadIdx.x;
}
DECLSPEC size_t get_local_id (const u32 dimindx __attribute__((unused)))
DECLSPEC size_t get_local_id (const u32 dimindx)
{
switch (dimindx)
{
case 0:
return threadIdx.x;
case 1:
return threadIdx.y;
case 2:
return threadIdx.z;
}
return (size_t) -1;
}
DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)))
DECLSPEC size_t get_local_size (const u32 dimindx)
{
// verify
switch (dimindx)
{
case 0:
return blockDim.x;
case 1:
return blockDim.y;
case 2:
return blockDim.z;
}
return (size_t) -1;
}
DECLSPEC u32x rotl32 (const u32x a, const int n)
{
return ((a << n) | ((a >> (32 - n))));
#if VECT_SIZE == 1
return rotl32_S (a, n);
#else
u32x t = 0;
#if VECT_SIZE >= 2
t.s0 = rotl32_S (a.s0, n);
t.s1 = rotl32_S (a.s1, n);
#endif
#if VECT_SIZE >= 4
t.s2 = rotl32_S (a.s2, n);
t.s3 = rotl32_S (a.s3, n);
#endif
#if VECT_SIZE >= 8
t.s4 = rotl32_S (a.s4, n);
t.s5 = rotl32_S (a.s5, n);
t.s6 = rotl32_S (a.s6, n);
t.s7 = rotl32_S (a.s7, n);
#endif
#if VECT_SIZE >= 16
t.s8 = rotl32_S (a.s8, n);
t.s9 = rotl32_S (a.s9, n);
t.sa = rotl32_S (a.sa, n);
t.sb = rotl32_S (a.sb, n);
t.sc = rotl32_S (a.sc, n);
t.sd = rotl32_S (a.sd, n);
t.se = rotl32_S (a.se, n);
t.sf = rotl32_S (a.sf, n);
#endif
return t;
#endif
}
DECLSPEC u32x rotr32 (const u32x a, const int n)
{
return ((a >> n) | ((a << (32 - n))));
#if VECT_SIZE == 1
return rotr32_S (a, n);
#else
u32x t = 0;
#if VECT_SIZE >= 2
t.s0 = rotr32_S (a.s0, n);
t.s1 = rotr32_S (a.s1, n);
#endif
#if VECT_SIZE >= 4
t.s2 = rotr32_S (a.s2, n);
t.s3 = rotr32_S (a.s3, n);
#endif
#if VECT_SIZE >= 8
t.s4 = rotr32_S (a.s4, n);
t.s5 = rotr32_S (a.s5, n);
t.s6 = rotr32_S (a.s6, n);
t.s7 = rotr32_S (a.s7, n);
#endif
#if VECT_SIZE >= 16
t.s8 = rotr32_S (a.s8, n);
t.s9 = rotr32_S (a.s9, n);
t.sa = rotr32_S (a.sa, n);
t.sb = rotr32_S (a.sb, n);
t.sc = rotr32_S (a.sc, n);
t.sd = rotr32_S (a.sd, n);
t.se = rotr32_S (a.se, n);
t.sf = rotr32_S (a.sf, n);
#endif
return t;
#endif
}
DECLSPEC u32 rotl32_S (const u32 a, const int n)
{
return ((a << n) | ((a >> (32 - n))));
return rotr32_S (a, 32 - n);
}
DECLSPEC u32 rotr32_S (const u32 a, const int n)
{
return ((a >> n) | ((a << (32 - n))));
return __builtin_amdgcn_alignbit (a, a, n);
}
DECLSPEC u64x rotl64 (const u64x a, const int n)
@ -249,15 +513,6 @@ DECLSPEC u64x rotl64 (const u64x a, const int n)
return rotr64 (a, 64 - n);
}
DECLSPEC u32 amd_bitalign_S (const u32 a, const u32 b, const int n)
{
u32 r = 0;
__asm__ ("V_ALIGNBIT_B32 %0, %1, %2, %3;" : "=v"(r): "v"(a), "v"(b), "I"(n));
return r;
}
DECLSPEC u64x rotr64 (const u64x a, const int n)
{
#if VECT_SIZE == 1
@ -283,15 +538,17 @@ DECLSPEC u64 rotr64_S (const u64 a, const int n)
vconv64_t out;
const int n31 = n & 31;
if (n < 32)
{
out.v32.a = amd_bitalign_S (a1, a0, n);
out.v32.b = amd_bitalign_S (a0, a1, n);
out.v32.a = __builtin_amdgcn_alignbit (a1, a0, n31);
out.v32.b = __builtin_amdgcn_alignbit (a0, a1, n31);
}
else
{
out.v32.a = amd_bitalign_S (a0, a1, n - 32);
out.v32.b = amd_bitalign_S (a1, a0, n - 32);
out.v32.a = __builtin_amdgcn_alignbit (a0, a1, n31);
out.v32.b = __builtin_amdgcn_alignbit (a1, a0, n31);
}
return out.v64;

35
OpenCL/inc_platform.h
View File

@ -21,14 +21,20 @@ DECLSPEC u64 rotl64_S (const u64 a, const int n);
DECLSPEC u64 rotr64_S (const u64 a, const int n);
#endif // IS_AMD
#ifdef IS_NV
DECLSPEC u32 hc_funnelshift_l (const u32 lo, const u32 hi, const int shift);
DECLSPEC u32 hc_funnelshift_r (const u32 lo, const u32 hi, const int shift);
#endif // IS_NV
#ifdef IS_CUDA
DECLSPEC u32 hc_atomic_dec (volatile GLOBAL_AS u32 *p);
DECLSPEC u32 hc_atomic_inc (volatile GLOBAL_AS u32 *p);
DECLSPEC u32 hc_atomic_or (volatile GLOBAL_AS u32 *p, volatile const u32 val);
DECLSPEC size_t get_global_id (const u32 dimindx __attribute__((unused)));
DECLSPEC size_t get_local_id (const u32 dimindx __attribute__((unused)));
DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)));
DECLSPEC size_t get_group_id (const u32 dimindx);
DECLSPEC size_t get_local_id (const u32 dimindx);
DECLSPEC size_t get_local_size (const u32 dimindx);
DECLSPEC u32x rotl32 (const u32x a, const int n);
DECLSPEC u32x rotr32 (const u32x a, const int n);
@ -48,7 +54,8 @@ DECLSPEC u32 hc_atomic_dec (volatile GLOBAL_AS u32 *p);
DECLSPEC u32 hc_atomic_inc (volatile GLOBAL_AS u32 *p);
DECLSPEC u32 hc_atomic_or (volatile GLOBAL_AS u32 *p, volatile const u32 val);
DECLSPEC size_t get_global_id (const u32 dimindx);
DECLSPEC size_t get_global_id (const u32 dimindx __attribute__((unused)));
DECLSPEC size_t get_group_id (const u32 dimindx);
DECLSPEC size_t get_local_id (const u32 dimindx);
DECLSPEC size_t get_local_size (const u32 dimindx);
@ -71,9 +78,25 @@ DECLSPEC u32 hc_atomic_dec (volatile GLOBAL_AS u32 *p);
DECLSPEC u32 hc_atomic_inc (volatile GLOBAL_AS u32 *p);
DECLSPEC u32 hc_atomic_or (volatile GLOBAL_AS u32 *p, volatile const u32 val);
#define get_global_id(param) hc_gid
#define get_local_id(param) hc_lid
#define get_local_size(param) hc_lsz
#define get_global_id(dimindx) \
((dimindx) == 0 ? hc_gid.x : \
(dimindx) == 1 ? hc_gid.y : \
(dimindx) == 2 ? hc_gid.z : -1)
#define get_group_id(dimindx) \
((dimindx) == 0 ? hc_bid.x : \
(dimindx) == 1 ? hc_bid.y : \
(dimindx) == 2 ? hc_bid.z : -1)
#define get_local_id(dimindx) \
((dimindx) == 0 ? hc_lid.x : \
(dimindx) == 1 ? hc_lid.y : \
(dimindx) == 2 ? hc_lid.z : -1)
#define get_local_size(dimindx) \
((dimindx) == 0 ? hc_lsz.x : \
(dimindx) == 1 ? hc_lsz.y : \
(dimindx) == 2 ? hc_lsz.z : -1)
DECLSPEC u32x rotl32 (const u32x a, const int n);
DECLSPEC u32x rotr32 (const u32x a, const int n);

134
OpenCL/inc_rp_optimized.cl
View File

@ -781,7 +781,6 @@ DECLSPEC void append_block8_optimized (const u32 offset, PRIVATE_AS u32 *buf0, P
const int offset_switch = offset / 4;
#if ((defined IS_AMD || defined IS_HIP) && HAS_VPERM == 0) || defined IS_GENERIC
const u32 src_r00 = src_r0[0];
const u32 src_r01 = src_r0[1];
const u32 src_r02 = src_r0[2];
@ -882,123 +881,6 @@ DECLSPEC void append_block8_optimized (const u32 offset, PRIVATE_AS u32 *buf0, P
s0 = 0;
break;
}
#endif
#if ((defined IS_AMD || defined IS_HIP) && HAS_VPERM == 1) || defined IS_NV
const int offset_mod_4 = offset & 3;
const int offset_minus_4 = 4 - offset_mod_4;
#if defined IS_NV
const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
#endif
#if (defined IS_AMD || defined IS_HIP)
const int selector = l32_from_64_S (0x0706050403020100UL >> (offset_minus_4 * 8));
#endif
const u32 src_r00 = src_r0[0];
const u32 src_r01 = src_r0[1];
const u32 src_r02 = src_r0[2];
const u32 src_r03 = src_r0[3];
const u32 src_r10 = src_r1[0];
const u32 src_r11 = src_r1[1];
const u32 src_r12 = src_r1[2];
const u32 src_r13 = src_r1[3];
switch (offset_switch)
{
case 0:
s7 = hc_byte_perm_S (src_r12, src_r13, selector);
s6 = hc_byte_perm_S (src_r11, src_r12, selector);
s5 = hc_byte_perm_S (src_r10, src_r11, selector);
s4 = hc_byte_perm_S (src_r03, src_r10, selector);
s3 = hc_byte_perm_S (src_r02, src_r03, selector);
s2 = hc_byte_perm_S (src_r01, src_r02, selector);
s1 = hc_byte_perm_S (src_r00, src_r01, selector);
s0 = hc_byte_perm_S ( 0, src_r00, selector);
break;
case 1:
s7 = hc_byte_perm_S (src_r11, src_r12, selector);
s6 = hc_byte_perm_S (src_r10, src_r11, selector);
s5 = hc_byte_perm_S (src_r03, src_r10, selector);
s4 = hc_byte_perm_S (src_r02, src_r03, selector);
s3 = hc_byte_perm_S (src_r01, src_r02, selector);
s2 = hc_byte_perm_S (src_r00, src_r01, selector);
s1 = hc_byte_perm_S ( 0, src_r00, selector);
s0 = 0;
break;
case 2:
s7 = hc_byte_perm_S (src_r10, src_r11, selector);
s6 = hc_byte_perm_S (src_r03, src_r10, selector);
s5 = hc_byte_perm_S (src_r02, src_r03, selector);
s4 = hc_byte_perm_S (src_r01, src_r02, selector);
s3 = hc_byte_perm_S (src_r00, src_r01, selector);
s2 = hc_byte_perm_S ( 0, src_r00, selector);
s1 = 0;
s0 = 0;
break;
case 3:
s7 = hc_byte_perm_S (src_r03, src_r10, selector);
s6 = hc_byte_perm_S (src_r02, src_r03, selector);
s5 = hc_byte_perm_S (src_r01, src_r02, selector);
s4 = hc_byte_perm_S (src_r00, src_r01, selector);
s3 = hc_byte_perm_S ( 0, src_r00, selector);
s2 = 0;
s1 = 0;
s0 = 0;
break;
case 4:
s7 = hc_byte_perm_S (src_r02, src_r03, selector);
s6 = hc_byte_perm_S (src_r01, src_r02, selector);
s5 = hc_byte_perm_S (src_r00, src_r01, selector);
s4 = hc_byte_perm_S ( 0, src_r00, selector);
s3 = 0;
s2 = 0;
s1 = 0;
s0 = 0;
break;
case 5:
s7 = hc_byte_perm_S (src_r01, src_r02, selector);
s6 = hc_byte_perm_S (src_r00, src_r01, selector);
s5 = hc_byte_perm_S ( 0, src_r00, selector);
s4 = 0;
s3 = 0;
s2 = 0;
s1 = 0;
s0 = 0;
break;
case 6:
s7 = hc_byte_perm_S (src_r00, src_r01, selector);
s6 = hc_byte_perm_S ( 0, src_r00, selector);
s5 = 0;
s4 = 0;
s3 = 0;
s2 = 0;
s1 = 0;
s0 = 0;
break;
case 7:
s7 = hc_byte_perm_S ( 0, src_r00, selector);
s6 = 0;
s5 = 0;
s4 = 0;
s3 = 0;
s2 = 0;
s1 = 0;
s0 = 0;
break;
}
#endif
buf0[0] = src_l0[0] | s0;
buf0[1] = src_l0[1] | s1;
@ -1202,9 +1084,18 @@ DECLSPEC HC_INLINE_RP u32 rule_op_mangle_toggle_at_sep (MAYBE_UNUSED const u32 p
{
ro = 1 << i;
break;
}
#ifdef IS_METAL
i = 32;
continue;
#else
break; // bug on Apple Intel/Silicon with Metal
#endif
}
occurence++;
}
}
@ -2353,6 +2244,8 @@ DECLSPEC u32 apply_rule_optimized (const u32 name, const u32 p0, const u32 p1, P
{
u32 out_len = in_len;
if (name == RULE_OP_MANGLE_NOOP) return out_len;
switch (name)
{
case RULE_OP_MANGLE_LREST: out_len = rule_op_mangle_lrest (p0, p1, buf0, buf1, out_len); break;
@ -2402,6 +2295,7 @@ DECLSPEC u32 apply_rule_optimized (const u32 name, const u32 p0, const u32 p1, P
return out_len;
}
//DECLSPEC u32 apply_rules_optimized (PRIVATE_AS const u32 *cmds, PRIVATE_AS u32 *buf0, PRIVATE_AS u32 *buf1, const u32 len)
DECLSPEC u32 apply_rules_optimized (CONSTANT_AS const u32 *cmds, PRIVATE_AS u32 *buf0, PRIVATE_AS u32 *buf1, const u32 len)
{
u32 out_len = len;

2
OpenCL/inc_rp_optimized.h
View File

@ -14,7 +14,7 @@
#define MAYBE_UNUSED
#endif
#ifdef IS_APPLE_SILICON
#ifdef IS_METAL
#define HC_INLINE_RP __attribute__ ((noinline))
#else
#define HC_INLINE_RP

12
OpenCL/inc_shared.h
View File

@ -13,28 +13,28 @@
GLOBAL_AS u32 *pws_comp, \
GLOBAL_AS pw_t *pws_buf, \
CONSTANT_AS const u64 &gid_max, \
uint hc_gid [[ thread_position_in_grid ]]
uint3 hc_gid [[ thread_position_in_grid ]]
#define KERN_ATTR_GPU_MEMSET \
GLOBAL_AS uint4 *buf, \
CONSTANT_AS const u32 &value, \
CONSTANT_AS const u64 &gid_max, \
uint hc_gid [[ thread_position_in_grid ]]
uint3 hc_gid [[ thread_position_in_grid ]]
#define KERN_ATTR_GPU_BZERO \
GLOBAL_AS uint4 *buf, \
CONSTANT_AS const u64 &gid_max, \
uint hc_gid [[ thread_position_in_grid ]]
uint3 hc_gid [[ thread_position_in_grid ]]
#define KERN_ATTR_GPU_ATINIT \
GLOBAL_AS pw_t *buf, \
CONSTANT_AS const u64 &gid_max, \
uint hc_gid [[ thread_position_in_grid ]]
uint3 hc_gid [[ thread_position_in_grid ]]
#define KERN_ATTR_GPU_UTF8_TO_UTF16 \
GLOBAL_AS pw_t *pws_buf, \
CONSTANT_AS const u64 &gid_max, \
uint hc_gid [[ thread_position_in_grid ]]
uint3 hc_gid [[ thread_position_in_grid ]]
#else // CUDA, HIP, OpenCL
@ -63,4 +63,6 @@
#endif // IS_METAL
DECLSPEC void gpu_decompress_entry (GLOBAL_AS pw_idx_t *pws_idx, GLOBAL_AS u32 *pws_comp, PRIVATE_AS pw_t *buf, const u64 gid);
#endif // INC_SHARED_H

201
OpenCL/inc_types.h
View File

@ -11,11 +11,13 @@
#define BITMAP_SHIFT1 kernel_param->bitmap_shift1
#define BITMAP_SHIFT2 kernel_param->bitmap_shift2
#define SALT_POS_HOST (kernel_param->pws_pos + gid)
#define SALT_POS_HOST_BID (kernel_param->pws_pos + bid)
#define LOOP_POS kernel_param->loop_pos
#define LOOP_CNT kernel_param->loop_cnt
#define IL_CNT kernel_param->il_cnt
#define DIGESTS_CNT 1
#define DIGESTS_OFFSET_HOST (kernel_param->pws_pos + gid)
#define DIGESTS_OFFSET_HOST_BID (kernel_param->pws_pos + bid)
#define COMBS_MODE kernel_param->combs_mode
#define SALT_REPEAT kernel_param->salt_repeat
#define PWS_POS kernel_param->pws_pos
@ -25,11 +27,13 @@
#define BITMAP_SHIFT1 kernel_param->bitmap_shift1
#define BITMAP_SHIFT2 kernel_param->bitmap_shift2
#define SALT_POS_HOST kernel_param->salt_pos_host
#define SALT_POS_HOST_BID SALT_POS_HOST
#define LOOP_POS kernel_param->loop_pos
#define LOOP_CNT kernel_param->loop_cnt
#define IL_CNT kernel_param->il_cnt
#define DIGESTS_CNT kernel_param->digests_cnt
#define DIGESTS_OFFSET_HOST kernel_param->digests_offset_host
#define DIGESTS_OFFSET_HOST_BID DIGESTS_OFFSET_HOST
#define COMBS_MODE kernel_param->combs_mode
#define SALT_REPEAT kernel_param->salt_repeat
#define PWS_POS kernel_param->pws_pos
@ -1565,6 +1569,202 @@ typedef enum ripemd160_constants
} ripemd160_constants_t;
typedef enum ripemd320_constants
{
RIPEMD320M_A=0x67452301U,
RIPEMD320M_B=0xefcdab89U,
RIPEMD320M_C=0x98badcfeU,
RIPEMD320M_D=0x10325476U,
RIPEMD320M_E=0xc3d2e1f0U,
RIPEMD320M_F=0x76543210U,
RIPEMD320M_G=0xfedcba98U,
RIPEMD320M_H=0x89abcdefU,
RIPEMD320M_I=0x01234567U,
RIPEMD320M_L=0x3c2d1e0fU,
RIPEMD320C00=0x00000000U,
RIPEMD320C10=0x5a827999U,
RIPEMD320C20=0x6ed9eba1U,
RIPEMD320C30=0x8f1bbcdcU,
RIPEMD320C40=0xa953fd4eU,
RIPEMD320C50=0x50a28be6U,
RIPEMD320C60=0x5c4dd124U,
RIPEMD320C70=0x6d703ef3U,
RIPEMD320C80=0x7a6d76e9U,
RIPEMD320C90=0x00000000U,
RIPEMD320S00=11,
RIPEMD320S01=14,
RIPEMD320S02=15,
RIPEMD320S03=12,
RIPEMD320S04=5,
RIPEMD320S05=8,
RIPEMD320S06=7,
RIPEMD320S07=9,
RIPEMD320S08=11,
RIPEMD320S09=13,
RIPEMD320S0A=14,
RIPEMD320S0B=15,
RIPEMD320S0C=6,
RIPEMD320S0D=7,
RIPEMD320S0E=9,
RIPEMD320S0F=8,
RIPEMD320S10=7,
RIPEMD320S11=6,
RIPEMD320S12=8,
RIPEMD320S13=13,
RIPEMD320S14=11,
RIPEMD320S15=9,
RIPEMD320S16=7,
RIPEMD320S17=15,
RIPEMD320S18=7,
RIPEMD320S19=12,
RIPEMD320S1A=15,
RIPEMD320S1B=9,
RIPEMD320S1C=11,
RIPEMD320S1D=7,
RIPEMD320S1E=13,
RIPEMD320S1F=12,
RIPEMD320S20=11,
RIPEMD320S21=13,
RIPEMD320S22=6,
RIPEMD320S23=7,
RIPEMD320S24=14,
RIPEMD320S25=9,
RIPEMD320S26=13,
RIPEMD320S27=15,
RIPEMD320S28=14,
RIPEMD320S29=8,
RIPEMD320S2A=13,
RIPEMD320S2B=6,
RIPEMD320S2C=5,
RIPEMD320S2D=12,
RIPEMD320S2E=7,
RIPEMD320S2F=5,
RIPEMD320S30=11,
RIPEMD320S31=12,
RIPEMD320S32=14,
RIPEMD320S33=15,
RIPEMD320S34=14,
RIPEMD320S35=15,
RIPEMD320S36=9,
RIPEMD320S37=8,
RIPEMD320S38=9,
RIPEMD320S39=14,
RIPEMD320S3A=5,
RIPEMD320S3B=6,
RIPEMD320S3C=8,
RIPEMD320S3D=6,
RIPEMD320S3E=5,
RIPEMD320S3F=12,
RIPEMD320S40=9,
RIPEMD320S41=15,
RIPEMD320S42=5,
RIPEMD320S43=11,
RIPEMD320S44=6,
RIPEMD320S45=8,
RIPEMD320S46=13,
RIPEMD320S47=12,
RIPEMD320S48=5,
RIPEMD320S49=12,
RIPEMD320S4A=13,
RIPEMD320S4B=14,
RIPEMD320S4C=11,
RIPEMD320S4D=8,
RIPEMD320S4E=5,
RIPEMD320S4F=6,
RIPEMD320S50=8,
RIPEMD320S51=9,
RIPEMD320S52=9,
RIPEMD320S53=11,
RIPEMD320S54=13,
RIPEMD320S55=15,
RIPEMD320S56=15,
RIPEMD320S57=5,
RIPEMD320S58=7,
RIPEMD320S59=7,
RIPEMD320S5A=8,
RIPEMD320S5B=11,
RIPEMD320S5C=14,
RIPEMD320S5D=14,
RIPEMD320S5E=12,
RIPEMD320S5F=6,
RIPEMD320S60=9,
RIPEMD320S61=13,
RIPEMD320S62=15,
RIPEMD320S63=7,
RIPEMD320S64=12,
RIPEMD320S65=8,
RIPEMD320S66=9,
RIPEMD320S67=11,
RIPEMD320S68=7,
RIPEMD320S69=7,
RIPEMD320S6A=12,
RIPEMD320S6B=7,
RIPEMD320S6C=6,
RIPEMD320S6D=15,
RIPEMD320S6E=13,
RIPEMD320S6F=11,
RIPEMD320S70=9,
RIPEMD320S71=7,
RIPEMD320S72=15,
RIPEMD320S73=11,
RIPEMD320S74=8,
RIPEMD320S75=6,
RIPEMD320S76=6,
RIPEMD320S77=14,
RIPEMD320S78=12,
RIPEMD320S79=13,
RIPEMD320S7A=5,
RIPEMD320S7B=14,
RIPEMD320S7C=13,
RIPEMD320S7D=13,
RIPEMD320S7E=7,
RIPEMD320S7F=5,
RIPEMD320S80=15,
RIPEMD320S81=5,
RIPEMD320S82=8,
RIPEMD320S83=11,
RIPEMD320S84=14,
RIPEMD320S85=14,
RIPEMD320S86=6,
RIPEMD320S87=14,
RIPEMD320S88=6,
RIPEMD320S89=9,
RIPEMD320S8A=12,
RIPEMD320S8B=9,
RIPEMD320S8C=12,
RIPEMD320S8D=5,
RIPEMD320S8E=15,
RIPEMD320S8F=8,
RIPEMD320S90=8,
RIPEMD320S91=5,
RIPEMD320S92=12,
RIPEMD320S93=9,
RIPEMD320S94=12,
RIPEMD320S95=5,
RIPEMD320S96=14,
RIPEMD320S97=6,
RIPEMD320S98=8,
RIPEMD320S99=13,
RIPEMD320S9A=6,
RIPEMD320S9B=5,
RIPEMD320S9C=15,
RIPEMD320S9D=13,
RIPEMD320S9E=11,
RIPEMD320S9F=11
} ripemd320_constants_t;
typedef enum keccak_constants
{
KECCAK_RNDC_00=0x0000000000000001UL,
@ -1812,6 +2012,7 @@ typedef struct salt
u32 salt_len_pc;
u32 salt_iter;
u32 salt_iter2;
u32 salt_dimy;
u32 salt_sign[2];
u32 salt_repeats;

41
OpenCL/inc_vendor.h
View File

@ -66,6 +66,12 @@ using namespace metal;
#define KERNEL_FQ __kernel
#endif
#if defined FIXED_LOCAL_SIZE
#define KERNEL_FA FIXED_THREAD_COUNT(FIXED_LOCAL_SIZE)
#else
#define KERNEL_FA
#endif
#ifndef MAYBE_UNUSED
#define MAYBE_UNUSED
#endif
@ -115,10 +121,6 @@ using namespace metal;
#define IS_GENERIC
#endif
#if defined IS_AMD && HAS_VPERM == 1
#define IS_ROCM
#endif
#define LOCAL_MEM_TYPE_LOCAL 1
#define LOCAL_MEM_TYPE_GLOBAL 2
@ -150,6 +152,8 @@ using namespace metal;
#if defined IS_AMD && defined IS_GPU
#define DECLSPEC HC_INLINE
#elif defined IS_CUDA
#define DECLSPEC __device__
#elif defined IS_HIP
#define DECLSPEC __device__ HC_INLINE
#else
@ -182,26 +186,25 @@ using namespace metal;
#define USE_ROTATE
#endif
#ifdef IS_ROCM
#define USE_BITSELECT
#define USE_ROTATE
#endif
#ifdef IS_INTEL_SDK
#ifdef IS_CPU
//#define USE_BITSELECT
//#define USE_ROTATE
#define USE_BITSELECT
#define USE_ROTATE
#endif
#endif
#ifdef IS_OPENCL
//#define USE_BITSELECT
//#define USE_ROTATE
//#define USE_SWIZZLE
#define USE_BITSELECT
#define USE_ROTATE
#define USE_SWIZZLE
#endif
#ifdef IS_METAL
#define USE_ROTATE
#ifndef IS_APPLE_SILICON
#define USE_BITSELECT
#define USE_SWIZZLE
#endif
// Metal support max VECT_SIZE = 4
#define s0 x
@ -210,4 +213,14 @@ using namespace metal;
#define s3 w
#endif
#if HAS_SHFW == 1
#define USE_FUNNELSHIFT
#endif
// some algorithms do not like this, eg 150, 1100, ...
#ifdef NO_FUNNELSHIFT
#undef USE_FUNNELSHIFT
#endif
#endif // INC_VENDOR_H

24
OpenCL/m00000_a0-optimized.cl
View File

@ -16,14 +16,8 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00000_m04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00000_m04 (KERN_ATTR_RULES ())
{
/**
* modifier
*/
const u64 lid = get_local_id (0);
/**
* base
*/
@ -143,22 +137,16 @@ KERNEL_FQ void m00000_m04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00000_m08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00000_m08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00000_m16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00000_m16 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00000_s04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00000_s04 (KERN_ATTR_RULES ())
{
/**
* modifier
*/
const u64 lid = get_local_id (0);
/**
* base
*/
@ -293,10 +281,10 @@ KERNEL_FQ void m00000_s04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00000_s08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00000_s08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00000_s16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00000_s16 (KERN_ATTR_RULES ())
{
}

4
OpenCL/m00000_a0-pure.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00000_mxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00000_mxx (KERN_ATTR_RULES ())
{
/**
* modifier
@ -60,7 +60,7 @@ KERNEL_FQ void m00000_mxx (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00000_sxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00000_sxx (KERN_ATTR_RULES ())
{
/**
* modifier

12
OpenCL/m00000_a1-optimized.cl
View File

@ -15,7 +15,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00000_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00000_m04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -201,15 +201,15 @@ KERNEL_FQ void m00000_m04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00000_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00000_m08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00000_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00000_m16 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00000_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00000_s04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -410,10 +410,10 @@ KERNEL_FQ void m00000_s04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00000_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00000_s08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00000_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00000_s16 (KERN_ATTR_BASIC ())
{
}

4
OpenCL/m00000_a1-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00000_mxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00000_mxx (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -56,7 +56,7 @@ KERNEL_FQ void m00000_mxx (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00000_sxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00000_sxx (KERN_ATTR_BASIC ())
{
/**
* modifier

12
OpenCL/m00000_a3-optimized.cl
View File

@ -423,7 +423,7 @@ DECLSPEC void m00000s (PRIVATE_AS u32 *w, const u32 pw_len, KERN_ATTR_FUNC_VECTO
}
}
KERNEL_FQ void m00000_m04 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00000_m04 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -463,7 +463,7 @@ KERNEL_FQ void m00000_m04 (KERN_ATTR_VECTOR ())
m00000m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00000_m08 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00000_m08 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -503,7 +503,7 @@ KERNEL_FQ void m00000_m08 (KERN_ATTR_VECTOR ())
m00000m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00000_m16 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00000_m16 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -543,7 +543,7 @@ KERNEL_FQ void m00000_m16 (KERN_ATTR_VECTOR ())
m00000m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00000_s04 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00000_s04 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -583,7 +583,7 @@ KERNEL_FQ void m00000_s04 (KERN_ATTR_VECTOR ())
m00000s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00000_s08 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00000_s08 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -623,7 +623,7 @@ KERNEL_FQ void m00000_s08 (KERN_ATTR_VECTOR ())
m00000s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00000_s16 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00000_s16 (KERN_ATTR_VECTOR ())
{
/**
* base

4
OpenCL/m00000_a3-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00000_mxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00000_mxx (KERN_ATTR_VECTOR ())
{
/**
* modifier
@ -69,7 +69,7 @@ KERNEL_FQ void m00000_mxx (KERN_ATTR_VECTOR ())
}
}
KERNEL_FQ void m00000_sxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00000_sxx (KERN_ATTR_VECTOR ())
{
/**
* modifier

12
OpenCL/m00010_a0-optimized.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00010_m04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00010_m04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -217,15 +217,15 @@ KERNEL_FQ void m00010_m04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00010_m08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00010_m08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00010_m16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00010_m16 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00010_s04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00010_s04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -441,10 +441,10 @@ KERNEL_FQ void m00010_s04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00010_s08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00010_s08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00010_s16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00010_s16 (KERN_ATTR_RULES ())
{
}

4
OpenCL/m00010_a0-pure.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00010_mxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00010_mxx (KERN_ATTR_RULES ())
{
/**
* modifier
@ -71,7 +71,7 @@ KERNEL_FQ void m00010_mxx (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00010_sxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00010_sxx (KERN_ATTR_RULES ())
{
/**
* modifier

12
OpenCL/m00010_a1-optimized.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00010_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00010_m04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -275,15 +275,15 @@ KERNEL_FQ void m00010_m04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00010_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00010_m08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00010_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00010_m16 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00010_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00010_s04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -559,10 +559,10 @@ KERNEL_FQ void m00010_s04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00010_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00010_s08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00010_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00010_s16 (KERN_ATTR_BASIC ())
{
}

4
OpenCL/m00010_a1-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00010_mxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00010_mxx (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -67,7 +67,7 @@ KERNEL_FQ void m00010_mxx (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00010_sxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00010_sxx (KERN_ATTR_BASIC ())
{
/**
* modifier

12
OpenCL/m00010_a3-optimized.cl
View File

@ -472,7 +472,7 @@ DECLSPEC void m00010s (PRIVATE_AS u32 *w, const u32 pw_len, KERN_ATTR_FUNC_VECTO
}
}
KERNEL_FQ void m00010_m04 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00010_m04 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -512,7 +512,7 @@ KERNEL_FQ void m00010_m04 (KERN_ATTR_VECTOR ())
m00010m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00010_m08 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00010_m08 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -552,7 +552,7 @@ KERNEL_FQ void m00010_m08 (KERN_ATTR_VECTOR ())
m00010m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00010_m16 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00010_m16 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -592,7 +592,7 @@ KERNEL_FQ void m00010_m16 (KERN_ATTR_VECTOR ())
m00010m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00010_s04 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00010_s04 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -632,7 +632,7 @@ KERNEL_FQ void m00010_s04 (KERN_ATTR_VECTOR ())
m00010s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00010_s08 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00010_s08 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -672,7 +672,7 @@ KERNEL_FQ void m00010_s08 (KERN_ATTR_VECTOR ())
m00010s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00010_s16 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00010_s16 (KERN_ATTR_VECTOR ())
{
/**
* base

4
OpenCL/m00010_a3-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00010_mxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00010_mxx (KERN_ATTR_VECTOR ())
{
/**
* modifier
@ -80,7 +80,7 @@ KERNEL_FQ void m00010_mxx (KERN_ATTR_VECTOR ())
}
}
KERNEL_FQ void m00010_sxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00010_sxx (KERN_ATTR_VECTOR ())
{
/**
* modifier

12
OpenCL/m00020_a0-optimized.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00020_m04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00020_m04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -197,15 +197,15 @@ KERNEL_FQ void m00020_m04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00020_m08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00020_m08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00020_m16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00020_m16 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00020_s04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00020_s04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -401,10 +401,10 @@ KERNEL_FQ void m00020_s04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00020_s08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00020_s08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00020_s16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00020_s16 (KERN_ATTR_RULES ())
{
}

4
OpenCL/m00020_a0-pure.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00020_mxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00020_mxx (KERN_ATTR_RULES ())
{
/**
* modifier
@ -64,7 +64,7 @@ KERNEL_FQ void m00020_mxx (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00020_sxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00020_sxx (KERN_ATTR_RULES ())
{
/**
* modifier

12
OpenCL/m00020_a1-optimized.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00020_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00020_m04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -253,15 +253,15 @@ KERNEL_FQ void m00020_m04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00020_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00020_m08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00020_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00020_m16 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00020_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00020_s04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -515,10 +515,10 @@ KERNEL_FQ void m00020_s04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00020_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00020_s08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00020_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00020_s16 (KERN_ATTR_BASIC ())
{
}

4
OpenCL/m00020_a1-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00020_mxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00020_mxx (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -58,7 +58,7 @@ KERNEL_FQ void m00020_mxx (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00020_sxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00020_sxx (KERN_ATTR_BASIC ())
{
/**
* modifier

12
OpenCL/m00020_a3-optimized.cl
View File

@ -373,7 +373,7 @@ DECLSPEC void m00020s (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w
}
}
KERNEL_FQ void m00020_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00020_m04 (KERN_ATTR_BASIC ())
{
/**
* base
@ -422,7 +422,7 @@ KERNEL_FQ void m00020_m04 (KERN_ATTR_BASIC ())
m00020m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00020_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00020_m08 (KERN_ATTR_BASIC ())
{
/**
* base
@ -471,7 +471,7 @@ KERNEL_FQ void m00020_m08 (KERN_ATTR_BASIC ())
m00020m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00020_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00020_m16 (KERN_ATTR_BASIC ())
{
/**
* base
@ -520,7 +520,7 @@ KERNEL_FQ void m00020_m16 (KERN_ATTR_BASIC ())
m00020m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00020_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00020_s04 (KERN_ATTR_BASIC ())
{
/**
* base
@ -569,7 +569,7 @@ KERNEL_FQ void m00020_s04 (KERN_ATTR_BASIC ())
m00020s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00020_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00020_s08 (KERN_ATTR_BASIC ())
{
/**
* base
@ -618,7 +618,7 @@ KERNEL_FQ void m00020_s08 (KERN_ATTR_BASIC ())
m00020s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00020_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00020_s16 (KERN_ATTR_BASIC ())
{
/**
* base

4
OpenCL/m00020_a3-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00020_mxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00020_mxx (KERN_ATTR_VECTOR ())
{
/**
* modifier
@ -75,7 +75,7 @@ KERNEL_FQ void m00020_mxx (KERN_ATTR_VECTOR ())
}
}
KERNEL_FQ void m00020_sxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00020_sxx (KERN_ATTR_VECTOR ())
{
/**
* modifier

12
OpenCL/m00030_a0-optimized.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00030_m04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00030_m04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -222,15 +222,15 @@ KERNEL_FQ void m00030_m04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00030_m08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00030_m08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00030_m16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00030_m16 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00030_s04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00030_s04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -451,10 +451,10 @@ KERNEL_FQ void m00030_s04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00030_s08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00030_s08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00030_s16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00030_s16 (KERN_ATTR_RULES ())
{
}

4
OpenCL/m00030_a0-pure.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00030_mxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00030_mxx (KERN_ATTR_RULES ())
{
/**
* modifier
@ -71,7 +71,7 @@ KERNEL_FQ void m00030_mxx (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00030_sxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00030_sxx (KERN_ATTR_RULES ())
{
/**
* modifier

12
OpenCL/m00030_a1-optimized.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00030_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00030_m04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -280,15 +280,15 @@ KERNEL_FQ void m00030_m04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00030_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00030_m08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00030_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00030_m16 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00030_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00030_s04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -569,10 +569,10 @@ KERNEL_FQ void m00030_s04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00030_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00030_s08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00030_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00030_s16 (KERN_ATTR_BASIC ())
{
}

4
OpenCL/m00030_a1-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00030_mxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00030_mxx (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -67,7 +67,7 @@ KERNEL_FQ void m00030_mxx (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00030_sxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00030_sxx (KERN_ATTR_BASIC ())
{
/**
* modifier

12
OpenCL/m00030_a3-optimized.cl
View File

@ -472,7 +472,7 @@ DECLSPEC void m00030s (PRIVATE_AS u32 *w, const u32 pw_len, KERN_ATTR_FUNC_VECTO
}
}
KERNEL_FQ void m00030_m04 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00030_m04 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -512,7 +512,7 @@ KERNEL_FQ void m00030_m04 (KERN_ATTR_VECTOR ())
m00030m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00030_m08 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00030_m08 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -552,7 +552,7 @@ KERNEL_FQ void m00030_m08 (KERN_ATTR_VECTOR ())
m00030m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00030_m16 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00030_m16 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -592,7 +592,7 @@ KERNEL_FQ void m00030_m16 (KERN_ATTR_VECTOR ())
m00030m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00030_s04 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00030_s04 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -632,7 +632,7 @@ KERNEL_FQ void m00030_s04 (KERN_ATTR_VECTOR ())
m00030s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00030_s08 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00030_s08 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -672,7 +672,7 @@ KERNEL_FQ void m00030_s08 (KERN_ATTR_VECTOR ())
m00030s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00030_s16 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00030_s16 (KERN_ATTR_VECTOR ())
{
/**
* base

4
OpenCL/m00030_a3-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00030_mxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00030_mxx (KERN_ATTR_VECTOR ())
{
/**
* modifier
@ -96,7 +96,7 @@ KERNEL_FQ void m00030_mxx (KERN_ATTR_VECTOR ())
}
}
KERNEL_FQ void m00030_sxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00030_sxx (KERN_ATTR_VECTOR ())
{
/**
* modifier

12
OpenCL/m00040_a0-optimized.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00040_m04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00040_m04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -202,15 +202,15 @@ KERNEL_FQ void m00040_m04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00040_m08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00040_m08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00040_m16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00040_m16 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00040_s04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00040_s04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -411,10 +411,10 @@ KERNEL_FQ void m00040_s04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00040_s08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00040_s08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00040_s16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00040_s16 (KERN_ATTR_RULES ())
{
}

4
OpenCL/m00040_a0-pure.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00040_mxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00040_mxx (KERN_ATTR_RULES ())
{
/**
* modifier
@ -64,7 +64,7 @@ KERNEL_FQ void m00040_mxx (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00040_sxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00040_sxx (KERN_ATTR_RULES ())
{
/**
* modifier

12
OpenCL/m00040_a1-optimized.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00040_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00040_m04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -258,15 +258,15 @@ KERNEL_FQ void m00040_m04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00040_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00040_m08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00040_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00040_m16 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00040_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00040_s04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -525,10 +525,10 @@ KERNEL_FQ void m00040_s04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00040_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00040_s08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00040_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00040_s16 (KERN_ATTR_BASIC ())
{
}

4
OpenCL/m00040_a1-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00040_mxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00040_mxx (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -58,7 +58,7 @@ KERNEL_FQ void m00040_mxx (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00040_sxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00040_sxx (KERN_ATTR_BASIC ())
{
/**
* modifier

12
OpenCL/m00040_a3-optimized.cl
View File

@ -373,7 +373,7 @@ DECLSPEC void m00040s (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w
}
}
KERNEL_FQ void m00040_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00040_m04 (KERN_ATTR_BASIC ())
{
/**
* base
@ -422,7 +422,7 @@ KERNEL_FQ void m00040_m04 (KERN_ATTR_BASIC ())
m00040m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00040_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00040_m08 (KERN_ATTR_BASIC ())
{
/**
* base
@ -471,7 +471,7 @@ KERNEL_FQ void m00040_m08 (KERN_ATTR_BASIC ())
m00040m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00040_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00040_m16 (KERN_ATTR_BASIC ())
{
/**
* base
@ -520,7 +520,7 @@ KERNEL_FQ void m00040_m16 (KERN_ATTR_BASIC ())
m00040m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00040_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00040_s04 (KERN_ATTR_BASIC ())
{
/**
* base
@ -569,7 +569,7 @@ KERNEL_FQ void m00040_s04 (KERN_ATTR_BASIC ())
m00040s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00040_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00040_s08 (KERN_ATTR_BASIC ())
{
/**
* base
@ -618,7 +618,7 @@ KERNEL_FQ void m00040_s08 (KERN_ATTR_BASIC ())
m00040s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00040_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00040_s16 (KERN_ATTR_BASIC ())
{
/**
* base

4
OpenCL/m00040_a3-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00040_mxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00040_mxx (KERN_ATTR_VECTOR ())
{
/**
* modifier
@ -87,7 +87,7 @@ KERNEL_FQ void m00040_mxx (KERN_ATTR_VECTOR ())
}
}
KERNEL_FQ void m00040_sxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00040_sxx (KERN_ATTR_VECTOR ())
{
/**
* modifier

12
OpenCL/m00050_a0-optimized.cl
View File

@ -101,7 +101,7 @@ DECLSPEC void hmac_md5_run (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS
md5_transform_vector (w0, w1, w2, w3, digest);
}
KERNEL_FQ void m00050_m04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00050_m04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -206,15 +206,15 @@ KERNEL_FQ void m00050_m04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00050_m08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00050_m08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00050_m16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00050_m16 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00050_s04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00050_s04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -331,10 +331,10 @@ KERNEL_FQ void m00050_s04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00050_s08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00050_s08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00050_s16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00050_s16 (KERN_ATTR_RULES ())
{
}

4
OpenCL/m00050_a0-pure.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00050_mxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00050_mxx (KERN_ATTR_RULES ())
{
/**
* modifier
@ -69,7 +69,7 @@ KERNEL_FQ void m00050_mxx (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00050_sxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00050_sxx (KERN_ATTR_RULES ())
{
/**
* modifier

12
OpenCL/m00050_a1-optimized.cl
View File

@ -99,7 +99,7 @@ DECLSPEC void hmac_md5_run (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS
md5_transform_vector (w0, w1, w2, w3, digest);
}
KERNEL_FQ void m00050_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00050_m04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -264,15 +264,15 @@ KERNEL_FQ void m00050_m04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00050_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00050_m08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00050_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00050_m16 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00050_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00050_s04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -449,10 +449,10 @@ KERNEL_FQ void m00050_s04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00050_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00050_s08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00050_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00050_s16 (KERN_ATTR_BASIC ())
{
}

4
OpenCL/m00050_a1-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00050_mxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00050_mxx (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -92,7 +92,7 @@ KERNEL_FQ void m00050_mxx (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00050_sxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00050_sxx (KERN_ATTR_BASIC ())
{
/**
* modifier

12
OpenCL/m00050_a3-optimized.cl
View File

@ -315,7 +315,7 @@ DECLSPEC void m00050s (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w
}
}
KERNEL_FQ void m00050_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00050_m04 (KERN_ATTR_BASIC ())
{
/**
* base
@ -364,7 +364,7 @@ KERNEL_FQ void m00050_m04 (KERN_ATTR_BASIC ())
m00050m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00050_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00050_m08 (KERN_ATTR_BASIC ())
{
/**
* base
@ -413,7 +413,7 @@ KERNEL_FQ void m00050_m08 (KERN_ATTR_BASIC ())
m00050m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00050_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00050_m16 (KERN_ATTR_BASIC ())
{
/**
* base
@ -462,7 +462,7 @@ KERNEL_FQ void m00050_m16 (KERN_ATTR_BASIC ())
m00050m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00050_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00050_s04 (KERN_ATTR_BASIC ())
{
/**
* base
@ -511,7 +511,7 @@ KERNEL_FQ void m00050_s04 (KERN_ATTR_BASIC ())
m00050s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00050_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00050_s08 (KERN_ATTR_BASIC ())
{
/**
* base
@ -560,7 +560,7 @@ KERNEL_FQ void m00050_s08 (KERN_ATTR_BASIC ())
m00050s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00050_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00050_s16 (KERN_ATTR_BASIC ())
{
/**
* base

4
OpenCL/m00050_a3-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00050_mxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00050_mxx (KERN_ATTR_VECTOR ())
{
/**
* modifier
@ -78,7 +78,7 @@ KERNEL_FQ void m00050_mxx (KERN_ATTR_VECTOR ())
}
}
KERNEL_FQ void m00050_sxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00050_sxx (KERN_ATTR_VECTOR ())
{
/**
* modifier

12
OpenCL/m00060_a0-optimized.cl
View File

@ -101,7 +101,7 @@ DECLSPEC void hmac_md5_run (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS
md5_transform_vector (w0, w1, w2, w3, digest);
}
KERNEL_FQ void m00060_m04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00060_m04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -228,15 +228,15 @@ KERNEL_FQ void m00060_m04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00060_m08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00060_m08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00060_m16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00060_m16 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00060_s04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00060_s04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -375,10 +375,10 @@ KERNEL_FQ void m00060_s04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00060_s08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00060_s08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00060_s16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00060_s16 (KERN_ATTR_RULES ())
{
}

4
OpenCL/m00060_a0-pure.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00060_mxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00060_mxx (KERN_ATTR_RULES ())
{
/**
* modifier
@ -71,7 +71,7 @@ KERNEL_FQ void m00060_mxx (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00060_sxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00060_sxx (KERN_ATTR_RULES ())
{
/**
* modifier

12
OpenCL/m00060_a1-optimized.cl
View File

@ -99,7 +99,7 @@ DECLSPEC void hmac_md5_run (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS
md5_transform_vector (w0, w1, w2, w3, digest);
}
KERNEL_FQ void m00060_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00060_m04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -267,15 +267,15 @@ KERNEL_FQ void m00060_m04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00060_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00060_m08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00060_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00060_m16 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00060_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00060_s04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -455,10 +455,10 @@ KERNEL_FQ void m00060_s04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00060_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00060_s08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00060_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00060_s16 (KERN_ATTR_BASIC ())
{
}

4
OpenCL/m00060_a1-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00060_mxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00060_mxx (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -94,7 +94,7 @@ KERNEL_FQ void m00060_mxx (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00060_sxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00060_sxx (KERN_ATTR_BASIC ())
{
/**
* modifier

12
OpenCL/m00060_a3-optimized.cl
View File

@ -311,7 +311,7 @@ DECLSPEC void m00060s (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w
}
}
KERNEL_FQ void m00060_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00060_m04 (KERN_ATTR_BASIC ())
{
/**
* base
@ -360,7 +360,7 @@ KERNEL_FQ void m00060_m04 (KERN_ATTR_BASIC ())
m00060m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00060_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00060_m08 (KERN_ATTR_BASIC ())
{
/**
* base
@ -409,7 +409,7 @@ KERNEL_FQ void m00060_m08 (KERN_ATTR_BASIC ())
m00060m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00060_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00060_m16 (KERN_ATTR_BASIC ())
{
/**
* base
@ -458,7 +458,7 @@ KERNEL_FQ void m00060_m16 (KERN_ATTR_BASIC ())
m00060m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00060_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00060_s04 (KERN_ATTR_BASIC ())
{
/**
* base
@ -507,7 +507,7 @@ KERNEL_FQ void m00060_s04 (KERN_ATTR_BASIC ())
m00060s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00060_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00060_s08 (KERN_ATTR_BASIC ())
{
/**
* base
@ -556,7 +556,7 @@ KERNEL_FQ void m00060_s08 (KERN_ATTR_BASIC ())
m00060s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00060_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00060_s16 (KERN_ATTR_BASIC ())
{
/**
* base

4
OpenCL/m00060_a3-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00060_mxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00060_mxx (KERN_ATTR_VECTOR ())
{
/**
* modifier
@ -80,7 +80,7 @@ KERNEL_FQ void m00060_mxx (KERN_ATTR_VECTOR ())
}
}
KERNEL_FQ void m00060_sxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00060_sxx (KERN_ATTR_VECTOR ())
{
/**
* modifier

12
OpenCL/m00070_a0-optimized.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00070_m04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00070_m04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -152,15 +152,15 @@ KERNEL_FQ void m00070_m04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00070_m08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00070_m08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00070_m16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00070_m16 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00070_s04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00070_s04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -311,10 +311,10 @@ KERNEL_FQ void m00070_s04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00070_s08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00070_s08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00070_s16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00070_s16 (KERN_ATTR_RULES ())
{
}

4
OpenCL/m00070_a0-pure.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00070_mxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00070_mxx (KERN_ATTR_RULES ())
{
/**
* modifier
@ -60,7 +60,7 @@ KERNEL_FQ void m00070_mxx (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00070_sxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00070_sxx (KERN_ATTR_RULES ())
{
/**
* modifier

12
OpenCL/m00070_a1-optimized.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00070_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00070_m04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -210,15 +210,15 @@ KERNEL_FQ void m00070_m04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00070_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00070_m08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00070_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00070_m16 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00070_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00070_s04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -429,10 +429,10 @@ KERNEL_FQ void m00070_s04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00070_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00070_s08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00070_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00070_s16 (KERN_ATTR_BASIC ())
{
}

4
OpenCL/m00070_a1-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00070_mxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00070_mxx (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -56,7 +56,7 @@ KERNEL_FQ void m00070_mxx (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00070_sxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00070_sxx (KERN_ATTR_BASIC ())
{
/**
* modifier

12
OpenCL/m00070_a3-optimized.cl
View File

@ -426,7 +426,7 @@ DECLSPEC void m00070s (PRIVATE_AS u32 *w, const u32 pw_len, KERN_ATTR_FUNC_VECTO
}
}
KERNEL_FQ void m00070_m04 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00070_m04 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -466,7 +466,7 @@ KERNEL_FQ void m00070_m04 (KERN_ATTR_VECTOR ())
m00070m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00070_m08 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00070_m08 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -506,7 +506,7 @@ KERNEL_FQ void m00070_m08 (KERN_ATTR_VECTOR ())
m00070m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00070_m16 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00070_m16 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -546,7 +546,7 @@ KERNEL_FQ void m00070_m16 (KERN_ATTR_VECTOR ())
m00070m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00070_s04 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00070_s04 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -586,7 +586,7 @@ KERNEL_FQ void m00070_s04 (KERN_ATTR_VECTOR ())
m00070s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00070_s08 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00070_s08 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -626,7 +626,7 @@ KERNEL_FQ void m00070_s08 (KERN_ATTR_VECTOR ())
m00070s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00070_s16 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00070_s16 (KERN_ATTR_VECTOR ())
{
/**
* base

4
OpenCL/m00070_a3-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_md5.cl)
#endif
KERNEL_FQ void m00070_mxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00070_mxx (KERN_ATTR_VECTOR ())
{
/**
* modifier
@ -83,7 +83,7 @@ KERNEL_FQ void m00070_mxx (KERN_ATTR_VECTOR ())
}
}
KERNEL_FQ void m00070_sxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00070_sxx (KERN_ATTR_VECTOR ())
{
/**
* modifier

12
OpenCL/m00100_a0-optimized.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00100_m04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00100_m04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -188,15 +188,15 @@ KERNEL_FQ void m00100_m04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00100_m08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00100_m08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00100_m16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00100_m16 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00100_s04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00100_s04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -389,10 +389,10 @@ KERNEL_FQ void m00100_s04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00100_s08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00100_s08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00100_s16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00100_s16 (KERN_ATTR_RULES ())
{
}

4
OpenCL/m00100_a0-pure.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00100_mxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00100_mxx (KERN_ATTR_RULES ())
{
/**
* modifier
@ -60,7 +60,7 @@ KERNEL_FQ void m00100_mxx (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00100_sxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00100_sxx (KERN_ATTR_RULES ())
{
/**
* modifier

12
OpenCL/m00100_a1-optimized.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00100_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00100_m04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -244,15 +244,15 @@ KERNEL_FQ void m00100_m04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00100_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00100_m08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00100_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00100_m16 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00100_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00100_s04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -503,10 +503,10 @@ KERNEL_FQ void m00100_s04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00100_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00100_s08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00100_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00100_s16 (KERN_ATTR_BASIC ())
{
}

4
OpenCL/m00100_a1-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00100_mxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00100_mxx (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -56,7 +56,7 @@ KERNEL_FQ void m00100_mxx (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00100_sxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00100_sxx (KERN_ATTR_BASIC ())
{
/**
* modifier

12
OpenCL/m00100_a3-optimized.cl
View File

@ -516,7 +516,7 @@ DECLSPEC void m00100s (PRIVATE_AS u32 *w, const u32 pw_len, KERN_ATTR_FUNC_VECTO
}
}
KERNEL_FQ void m00100_m04 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00100_m04 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -556,7 +556,7 @@ KERNEL_FQ void m00100_m04 (KERN_ATTR_VECTOR ())
m00100m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00100_m08 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00100_m08 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -596,7 +596,7 @@ KERNEL_FQ void m00100_m08 (KERN_ATTR_VECTOR ())
m00100m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00100_m16 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00100_m16 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -636,7 +636,7 @@ KERNEL_FQ void m00100_m16 (KERN_ATTR_VECTOR ())
m00100m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00100_s04 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00100_s04 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -676,7 +676,7 @@ KERNEL_FQ void m00100_s04 (KERN_ATTR_VECTOR ())
m00100s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00100_s08 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00100_s08 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -716,7 +716,7 @@ KERNEL_FQ void m00100_s08 (KERN_ATTR_VECTOR ())
m00100s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00100_s16 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00100_s16 (KERN_ATTR_VECTOR ())
{
/**
* base

4
OpenCL/m00100_a3-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00100_mxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00100_mxx (KERN_ATTR_VECTOR ())
{
/**
* modifier
@ -69,7 +69,7 @@ KERNEL_FQ void m00100_mxx (KERN_ATTR_VECTOR ())
}
}
KERNEL_FQ void m00100_sxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00100_sxx (KERN_ATTR_VECTOR ())
{
/**
* modifier

12
OpenCL/m00110_a0-optimized.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00110_m04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00110_m04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -261,15 +261,15 @@ KERNEL_FQ void m00110_m04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00110_m08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00110_m08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00110_m16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00110_m16 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00110_s04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00110_s04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -535,10 +535,10 @@ KERNEL_FQ void m00110_s04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00110_s08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00110_s08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00110_s16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00110_s16 (KERN_ATTR_RULES ())
{
}

4
OpenCL/m00110_a0-pure.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00110_mxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00110_mxx (KERN_ATTR_RULES ())
{
/**
* modifier
@ -71,7 +71,7 @@ KERNEL_FQ void m00110_mxx (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00110_sxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00110_sxx (KERN_ATTR_RULES ())
{
/**
* modifier

12
OpenCL/m00110_a1-optimized.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00110_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00110_m04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -319,15 +319,15 @@ KERNEL_FQ void m00110_m04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00110_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00110_m08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00110_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00110_m16 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00110_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00110_s04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -653,10 +653,10 @@ KERNEL_FQ void m00110_s04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00110_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00110_s08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00110_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00110_s16 (KERN_ATTR_BASIC ())
{
}

4
OpenCL/m00110_a1-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00110_mxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00110_mxx (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -67,7 +67,7 @@ KERNEL_FQ void m00110_mxx (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00110_sxx (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00110_sxx (KERN_ATTR_BASIC ())
{
/**
* modifier

12
OpenCL/m00110_a3-optimized.cl
View File

@ -564,7 +564,7 @@ DECLSPEC void m00110s (PRIVATE_AS u32 *w, const u32 pw_len, KERN_ATTR_FUNC_VECTO
}
}
KERNEL_FQ void m00110_m04 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00110_m04 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -604,7 +604,7 @@ KERNEL_FQ void m00110_m04 (KERN_ATTR_VECTOR ())
m00110m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00110_m08 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00110_m08 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -644,7 +644,7 @@ KERNEL_FQ void m00110_m08 (KERN_ATTR_VECTOR ())
m00110m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00110_m16 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00110_m16 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -684,7 +684,7 @@ KERNEL_FQ void m00110_m16 (KERN_ATTR_VECTOR ())
m00110m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00110_s04 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00110_s04 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -724,7 +724,7 @@ KERNEL_FQ void m00110_s04 (KERN_ATTR_VECTOR ())
m00110s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00110_s08 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00110_s08 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -764,7 +764,7 @@ KERNEL_FQ void m00110_s08 (KERN_ATTR_VECTOR ())
m00110s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, kernel_param, gid, lid, lsz);
}
KERNEL_FQ void m00110_s16 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00110_s16 (KERN_ATTR_VECTOR ())
{
/**
* base

4
OpenCL/m00110_a3-pure.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00110_mxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00110_mxx (KERN_ATTR_VECTOR ())
{
/**
* modifier
@ -80,7 +80,7 @@ KERNEL_FQ void m00110_mxx (KERN_ATTR_VECTOR ())
}
}
KERNEL_FQ void m00110_sxx (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00110_sxx (KERN_ATTR_VECTOR ())
{
/**
* modifier

12
OpenCL/m00120_a0-optimized.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00120_m04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00120_m04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -253,15 +253,15 @@ KERNEL_FQ void m00120_m04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00120_m08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00120_m08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00120_m16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00120_m16 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00120_s04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00120_s04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -507,10 +507,10 @@ KERNEL_FQ void m00120_s04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00120_s08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00120_s08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00120_s16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00120_s16 (KERN_ATTR_RULES ())
{
}

4
OpenCL/m00120_a0-pure.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00120_mxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00120_mxx (KERN_ATTR_RULES ())
{
/**
* modifier
@ -64,7 +64,7 @@ KERNEL_FQ void m00120_mxx (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00120_sxx (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00120_sxx (KERN_ATTR_RULES ())
{
/**
* modifier

Some files were not shown because too many files have changed in this diff Show More