arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2025-07-21 22:18:44 +00:00
Code Issues Releases Wiki Activity

Merge remote-tracking branch 'upstream/master' into issue_4191

Browse Source
This commit is contained in:
Gabriele Gristina 2025-06-26 21:43:40 +02:00
commit a5e472c12e
No known key found for this signature in database
GPG Key ID: 9F68B59298F311F0
2609 changed files with 258480 additions and 44174 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 }}

8
.gitignore vendored
View File

@ -17,6 +17,8 @@ hashcat.dll
*.dSYM
kernels/**
lib/*.a
bridges/*.dll
bridges/*.so
modules/*.dll
modules/*.so
obj/*/*/*.o
@ -24,3 +26,9 @@ obj/*.o
obj/*.a
include/CL
tools/luks_tests
.vscode
# 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.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 ###

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);

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);

69
OpenCL/inc_common.cl
View File

@ -1317,7 +1317,7 @@ DECLSPEC u64x hc_swap64 (const u64x v)
asm volatile ("mov.b64 %0, {%1, %2};" : "=l"(r.sf) : "r"(tr.sf), "r"(tl.sf));
#endif
#elif defined IS_METAL
#elif defined IS_METAL || defined IS_APPLE_SILICON
const u32x a0 = h32_from_64 (v);
const u32x a1 = l32_from_64 (v);
@ -1391,7 +1391,7 @@ DECLSPEC u64 hc_swap64_S (const u64 v)
asm volatile ("mov.b64 %0, {%1, %2};" : "=l"(r) : "r"(tr), "r"(tl));
#elif defined IS_METAL
#elif defined IS_METAL || defined IS_APPLE_SILICON
const u32 v0 = h32_from_64_S (v);
const u32 v1 = l32_from_64_S (v);
@ -2701,6 +2701,71 @@ DECLSPEC int asn1_detect (PRIVATE_AS const u32 *buf, const int len)
return 1;
}
DECLSPEC int asn1_check_int_tag (PRIVATE_AS const u32 *buf, const int len)
{
PRIVATE_AS const u8 *bytes = (PRIVATE_AS const u8 *) buf;
int seq_len_offset = 0;
if (bytes[1] < 0x80)
{
seq_len_offset = 2;
}
else if (bytes[1] == 0x81)
{
seq_len_offset = 3;
}
else if (bytes[1] == 0x82)
{
seq_len_offset = 4;
}
else
{
return 0;
}
int pos = seq_len_offset;
if (pos >= len) return 0;
if (pos + 2 > len) return 0;
u8 tag = bytes[pos];
if (tag != 0x02) return 0;
u8 len_byte = bytes[pos + 1];
int val_len = 0;
int tmp_len = 1;
if (len_byte < 0x80)
{
val_len = len_byte;
}
else if (len_byte == 0x81)
{
if (pos + 2 >= len) return 0;
val_len = bytes[pos + 2];
tmp_len = 2;
}
else if (len_byte == 0x82)
{
if (pos + 3 >= len) return 0;
val_len = (bytes[pos + 2] << 8) | bytes[pos + 3];
tmp_len = 3;
}
else
{
return 0;
}
if (pos + 1 + tmp_len + val_len > len) return 0;
if (val_len != 1) return 0;
return 1;
}
DECLSPEC u32 check_bitmap (GLOBAL_AS const u32 *bitmap, const u32 bitmap_mask, const u32 bitmap_shift, const u32 digest)
{
return (bitmap[(digest >> bitmap_shift) & bitmap_mask] & (1 << (digest & 0x1f)));

4
OpenCL/inc_common.h
View File

@ -126,7 +126,8 @@
#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 ]]
uint hc_lsz [[ threads_per_threadgroup ]], \
uint hc_bid [[ threadgroup_position_in_grid ]]
#endif // IS_METAL
/*
@ -302,6 +303,7 @@ DECLSPEC int hc_enc_next_global (PRIVATE_AS hc_enc_t *hc_enc, GLOBAL_AS const u3
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);

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)

664
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;
@ -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);
@ -424,11 +423,520 @@ 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 = (const int) 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);
}
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)
{
const u32x t0 = len;
@ -452,6 +960,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);
@ -700,3 +1209,152 @@ DECLSPEC void blake2s_final_vector (PRIVATE_AS blake2s_ctx_vector_t *ctx)
{
blake2s_transform_vector (ctx->h, ctx->m, (u32x) ctx->len, 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

511
OpenCL/inc_hash_scrypt.cl Normal file
View File

@ -0,0 +1,511 @@
/**
* 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"
#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 uint4 *X4 = (PRIVATE_AS uint4 *) X;
GLOBAL_AS uint4 *V;
switch (xm4)
{
case 0: V = (GLOBAL_AS uint4 *) V0; break;
case 1: V = (GLOBAL_AS uint4 *) V1; break;
case 2: V = (GLOBAL_AS uint4 *) V2; break;
case 3: V = (GLOBAL_AS uint4 *) V3; break;
}
GLOBAL_AS uint4 *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 uint4 *Vxx = Vx + (y * zSIZE);
for (u32 z = 0; z < zSIZE; z++) *Vxx++ = 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 uint4 *X4 = (PRIVATE_AS uint4 *) X;
PRIVATE_AS uint4 *T4 = (PRIVATE_AS uint4 *) T;
GLOBAL_AS uint4 *V;
switch (xm4)
{
case 0: V = (GLOBAL_AS uint4 *) V0; break;
case 1: V = (GLOBAL_AS uint4 *) V1; break;
case 2: V = (GLOBAL_AS uint4 *) V2; break;
case 3: V = (GLOBAL_AS uint4 *) V3; break;
}
GLOBAL_AS uint4 *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 uint4 *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] = 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);
}

57
OpenCL/inc_hash_scrypt.h Normal file
View File

@ -0,0 +1,57 @@
/**
* 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)
//#define VIDX(bid4,lsz,lid,ySIZE,zSIZE,y,z) (((bid4) * (lsz) * (ySIZE) * (zSIZE)) + ((lid) * (ySIZE) * (zSIZE)) + ((y) * (zSIZE)) + (z))
#if defined IS_CUDA
DECLSPEC uint4 operator ^ (const uint4 a, const uint4 b)
{
uint4 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;
}
#endif
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);

223
OpenCL/inc_platform.cl
View File

@ -104,6 +104,11 @@ 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 __attribute__((unused)))
{
return blockIdx.x;
}
DECLSPEC size_t get_global_id (const u32 dimindx __attribute__((unused)))
{
return (blockIdx.x * blockDim.x) + threadIdx.x;
@ -122,37 +127,129 @@ DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)))
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 __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 __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,6 +305,11 @@ 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 __attribute__((unused)))
{
return blockIdx.x;
}
DECLSPEC size_t get_global_id (const u32 dimindx __attribute__((unused)))
{
return (blockIdx.x * blockDim.x) + threadIdx.x;
@ -226,36 +328,115 @@ DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)))
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 __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 __funnelshift_r (a, a, n);
#else
return ((a >> n) | ((a << (32 - n))));
#endif
}
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;
#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)
@ -274,6 +455,7 @@ DECLSPEC u64 rotl64_S (const u64 a, const int n)
DECLSPEC u64 rotr64_S (const u64 a, const int n)
{
#ifdef USE_FUNNELSHIFT
vconv64_t in;
in.v64 = a;
@ -285,16 +467,19 @@ DECLSPEC u64 rotr64_S (const u64 a, const int n)
if (n < 32)
{
out.v32.a = amd_bitalign_S (a1, a0, n);
out.v32.b = amd_bitalign_S (a0, a1, n);
out.v32.a = __funnelshift_r (a0, a1, n);
out.v32.b = __funnelshift_r (a1, a0, n);
}
else
{
out.v32.a = amd_bitalign_S (a0, a1, n - 32);
out.v32.b = amd_bitalign_S (a1, a0, n - 32);
out.v32.a = __funnelshift_r (a1, a0, n - 32);
out.v32.b = __funnelshift_r (a0, a1, n - 32);
}
return out.v64;
#else
return ((a >> n) | ((a << (64 - n))));
#endif
}
#define FIXED_THREAD_COUNT(n) __launch_bounds__((n), 0)

1
OpenCL/inc_platform.h
View File

@ -74,6 +74,7 @@ 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_group_id(param) hc_bid
DECLSPEC u32x rotl32 (const u32x a, const int n);
DECLSPEC u32x rotr32 (const u32x a, const int n);

17
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
@ -150,8 +156,10 @@ 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
#define DECLSPEC __device__
#else
#define DECLSPEC
#endif
@ -175,11 +183,13 @@ using namespace metal;
#ifdef IS_CUDA
#define USE_BITSELECT
#define USE_ROTATE
#define USE_FUNNELSHIFT
#endif
#ifdef IS_HIP
#define USE_BITSELECT
#define USE_ROTATE
#define USE_FUNNELSHIFT
#endif
#ifdef IS_ROCM
@ -210,4 +220,9 @@ using namespace metal;
#define s3 w
#endif
// some algorithms do not like this, eg 150, 1100, ...
#ifdef NO_FUNNELSHIFT
#undef USE_FUNNELSHIFT
#endif
#endif // INC_VENDOR_H

12
OpenCL/m00000_a0-optimized.cl
View File

@ -16,7 +16,7 @@
#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
@ -143,15 +143,15 @@ 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
@ -293,10 +293,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

12
OpenCL/m00120_a1-optimized.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00120_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00120_m04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -309,15 +309,15 @@ KERNEL_FQ void m00120_m04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00120_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00120_m08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00120_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00120_m16 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00120_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00120_s04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -621,10 +621,10 @@ KERNEL_FQ void m00120_s04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00120_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00120_s08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00120_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00120_s16 (KERN_ATTR_BASIC ())
{
}

4
OpenCL/m00120_a1-pure.cl
View File

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

12
OpenCL/m00120_a3-optimized.cl
View File

@ -467,7 +467,7 @@ DECLSPEC void m00120s (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w
}
}
KERNEL_FQ void m00120_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00120_m04 (KERN_ATTR_BASIC ())
{
/**
* base
@ -516,7 +516,7 @@ KERNEL_FQ void m00120_m04 (KERN_ATTR_BASIC ())
m00120m (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 m00120_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00120_m08 (KERN_ATTR_BASIC ())
{
/**
* base
@ -565,7 +565,7 @@ KERNEL_FQ void m00120_m08 (KERN_ATTR_BASIC ())
m00120m (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 m00120_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00120_m16 (KERN_ATTR_BASIC ())
{
/**
* base
@ -614,7 +614,7 @@ KERNEL_FQ void m00120_m16 (KERN_ATTR_BASIC ())
m00120m (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 m00120_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00120_s04 (KERN_ATTR_BASIC ())
{
/**
* base
@ -663,7 +663,7 @@ KERNEL_FQ void m00120_s04 (KERN_ATTR_BASIC ())
m00120s (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 m00120_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00120_s08 (KERN_ATTR_BASIC ())
{
/**
* base
@ -712,7 +712,7 @@ KERNEL_FQ void m00120_s08 (KERN_ATTR_BASIC ())
m00120s (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 m00120_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00120_s16 (KERN_ATTR_BASIC ())
{
/**
* base

4
OpenCL/m00120_a3-pure.cl
View File

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

12
OpenCL/m00130_a0-optimized.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00130_m04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00130_m04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -266,15 +266,15 @@ KERNEL_FQ void m00130_m04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00130_m08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00130_m08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00130_m16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00130_m16 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00130_s04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00130_s04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -545,10 +545,10 @@ KERNEL_FQ void m00130_s04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00130_s08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00130_s08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00130_s16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00130_s16 (KERN_ATTR_RULES ())
{
}

4
OpenCL/m00130_a0-pure.cl
View File

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

12
OpenCL/m00130_a1-optimized.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00130_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00130_m04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -324,15 +324,15 @@ KERNEL_FQ void m00130_m04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00130_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00130_m08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00130_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00130_m16 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00130_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00130_s04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -663,10 +663,10 @@ KERNEL_FQ void m00130_s04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00130_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00130_s08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00130_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00130_s16 (KERN_ATTR_BASIC ())
{
}

4
OpenCL/m00130_a1-pure.cl
View File

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

12
OpenCL/m00130_a3-optimized.cl
View File

@ -564,7 +564,7 @@ DECLSPEC void m00130s (PRIVATE_AS u32 *w, const u32 pw_len, KERN_ATTR_FUNC_VECTO
}
}
KERNEL_FQ void m00130_m04 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00130_m04 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -604,7 +604,7 @@ KERNEL_FQ void m00130_m04 (KERN_ATTR_VECTOR ())
m00130m (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 m00130_m08 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00130_m08 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -644,7 +644,7 @@ KERNEL_FQ void m00130_m08 (KERN_ATTR_VECTOR ())
m00130m (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 m00130_m16 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00130_m16 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -684,7 +684,7 @@ KERNEL_FQ void m00130_m16 (KERN_ATTR_VECTOR ())
m00130m (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 m00130_s04 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00130_s04 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -724,7 +724,7 @@ KERNEL_FQ void m00130_s04 (KERN_ATTR_VECTOR ())
m00130s (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 m00130_s08 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00130_s08 (KERN_ATTR_VECTOR ())
{
/**
* base
@ -764,7 +764,7 @@ KERNEL_FQ void m00130_s08 (KERN_ATTR_VECTOR ())
m00130s (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 m00130_s16 (KERN_ATTR_VECTOR ())
KERNEL_FQ KERNEL_FA void m00130_s16 (KERN_ATTR_VECTOR ())
{
/**
* base

4
OpenCL/m00130_a3-pure.cl
View File

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

12
OpenCL/m00140_a0-optimized.cl
View File

@ -16,7 +16,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00140_m04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00140_m04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -246,15 +246,15 @@ KERNEL_FQ void m00140_m04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00140_m08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00140_m08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00140_m16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00140_m16 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00140_s04 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00140_s04 (KERN_ATTR_RULES ())
{
/**
* modifier
@ -505,10 +505,10 @@ KERNEL_FQ void m00140_s04 (KERN_ATTR_RULES ())
}
}
KERNEL_FQ void m00140_s08 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00140_s08 (KERN_ATTR_RULES ())
{
}
KERNEL_FQ void m00140_s16 (KERN_ATTR_RULES ())
KERNEL_FQ KERNEL_FA void m00140_s16 (KERN_ATTR_RULES ())
{
}

4
OpenCL/m00140_a0-pure.cl
View File

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

12
OpenCL/m00140_a1-optimized.cl
View File

@ -14,7 +14,7 @@
#include M2S(INCLUDE_PATH/inc_hash_sha1.cl)
#endif
KERNEL_FQ void m00140_m04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00140_m04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -302,15 +302,15 @@ KERNEL_FQ void m00140_m04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00140_m08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00140_m08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00140_m16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00140_m16 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00140_s04 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00140_s04 (KERN_ATTR_BASIC ())
{
/**
* modifier
@ -619,10 +619,10 @@ KERNEL_FQ void m00140_s04 (KERN_ATTR_BASIC ())
}
}
KERNEL_FQ void m00140_s08 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00140_s08 (KERN_ATTR_BASIC ())
{
}
KERNEL_FQ void m00140_s16 (KERN_ATTR_BASIC ())
KERNEL_FQ KERNEL_FA void m00140_s16 (KERN_ATTR_BASIC ())
{
}

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