Initial commit

Makefile

@@ -0,0 +1 @@
+include src/Makefile

README.md

@@ -0,0 +1,47 @@
+## *oclHashcat* ##
+
+**oclHashcat** is the world's fastest and most advanced GPGPU-based password recovery utility, supporting five unique modes of attack for over 170 highly-optimized hashing algorithms. oclHashcat currently supports AMD (OpenCL) and Nvidia (CUDA) graphics processors on GNU/Linux and Windows 7/8/10, and has facilities to help enable distributed password cracking.
+
+### License ###
+
+**oclHashcat** is licensed under the MIT license. Refer to [docs/license.txt](docs/license.txt) for more information. 
+
+### Installation ###
+
+Download the [latest release](https://hashcat.net/oclhashcat/) and unpack it in the desired location. Please remember to use `7z x` when unpacking the archive from the command line to ensure full file paths remain intact.
+
+### Usage/Help ###
+
+Please refer to the [Hashcat Wiki](https://hashcat.net/wiki/) and the output of `--help` for usage information and general help. A list of frequently asked questions may also be found [here](https://hashcat.net/wiki/doku.php?id=frequently_asked_questions). The [Hashcat Forums](https://hashcat.net/forum/) also contain a plethora of information. 
+
+### Building ###
+
+Refer to [docs/BUILD.md](docs/BUILD.md) for instructions on how to build **oclHashcat** from source.
+
+### Contributing ###
+
+Contributions are welcome and encouraged, provided your code is of sufficient quality. Before submitting a pull request, please ensure your code adheres to the following requirements:
+
+1. Licensed under MIT license, or dedicated to public domain (BSD, GPL, etc. code is incompatible)
+2. Adheres to either C89, C90, or C99 standards
+2. Compiles cleanly with no warnings when compiled with `-W -Wall -std=c99`
+3. Uses [Allman-style](https://en.wikipedia.org/wiki/Indent_style#Allman_style) code blocks & indentation 
+4. Uses 2-character tabs and 2-character indentations
+5. Uses lower-case function and variable names 
+6. Avoids the use of `!` and uses positive conditionals wherever possible (e.g., `if (foo == 0)` instead of `if (!foo)`, and `if (foo)` instead of `if (foo !=0)`)
+
+You can use GNU Indent to help assist you with the style requirements:
+
+```
+indent -st -bad -bap -sc -bl -bli0 -ncdw -nce -cli0 -cbi0 -pcs -cs -npsl -bs -nbc -bls -blf -lp -i2 -ts2 -nut -l1024 -nbbo -fca -lc1024 -fc1
+```
+
+Your pull request should fully describe the functionality you are adding/removing or the problem you are solving. Regardless of whether your patch modifies one line or one thousand lines, you must describe what has prompted and/or motivated the change.
+
+Solve only one problem in each pull request. If you're fixing a bug and adding a new feature, you need to make two separate pull requests. If you're fixing three bugs, you need to make three separate pull requests. If you're adding four new features, you need to make four separate pull requests. So on, and so forth.
+
+If your patch fixes a bug, please be sure there is an [issue](https://github.com/hashcat/oclHashcat/issues) open for the bug before submitting a pull request. If your patch aims to improve performance or optimizes an algorithm, be sure to quantify your optimizations and document the trade-offs, and back up your claims with benchmarks and metrics.
+
+In order to maintain the quality and integrity of the **oclHashcat** source tree, all pull requests must be reviewed and signed off by at least two [board members](https://github.com/orgs/hashcat/people) before being merged. The [project lead](https://github.com/jsteube) has the ultimate authority in deciding whether to accept or reject a pull request. Do not be discouraged if your pull request is rejected!
+
+### Happy Cracking!

amd/amp_a0_v1.cl

@@ -0,0 +1,52 @@
+/**
+ * Author......: Jens Steube <[email protected]>
+ * License.....: MIT
+ */
+
+#define VECT_SIZE1
+
+#include "include/constants.h"
+#include "include/kernel_vendor.h"
+#include "types_amd.c"
+
+static u32 swap_workaround (const u32 v)
+{
+  return (as_uint (as_uchar4 (v).s3210));
+}
+
+#include "include/rp_gpu.h"
+#include "rp_amd.c"
+
+__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) amp (__global pw_t *pws, __global pw_t *pws_amp, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, const u32 combs_mode, const u32 gid_max)
+{
+  const u32 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  const u32 pw_len = pws[gid].pw_len;
+
+  u32x w0[4];
+  u32x w1[4];
+
+  w0[0] = pws[gid].i[ 0];
+  w0[1] = pws[gid].i[ 1];
+  w0[2] = pws[gid].i[ 2];
+  w0[3] = pws[gid].i[ 3];
+  w1[0] = pws[gid].i[ 4];
+  w1[1] = pws[gid].i[ 5];
+  w1[2] = pws[gid].i[ 6];
+  w1[3] = pws[gid].i[ 7];
+
+  const u32 out_len = apply_rules (rules_buf[0].cmds, w0, w1, pw_len);
+
+  pws_amp[gid].i[0] = w0[0];
+  pws_amp[gid].i[1] = w0[1];
+  pws_amp[gid].i[2] = w0[2];
+  pws_amp[gid].i[3] = w0[3];
+  pws_amp[gid].i[4] = w1[0];
+  pws_amp[gid].i[5] = w1[1];
+  pws_amp[gid].i[6] = w1[2];
+  pws_amp[gid].i[7] = w1[3];
+
+  pws_amp[gid].pw_len = out_len;
+}

amd/amp_a0_v2.cl

@@ -0,0 +1,52 @@
+/**
+ * Author......: Jens Steube <[email protected]>
+ * License.....: MIT
+ */
+
+#define VECT_SIZE2
+
+#include "include/constants.h"
+#include "include/kernel_vendor.h"
+#include "types_amd.c"
+
+static u32x swap_workaround (const u32x v)
+{
+  return (v << 24) + ((v & 0x0000FF00) << 8) + ((v & 0x00FF0000) >> 8) + (v >> 24);
+}
+
+#include "include/rp_gpu.h"
+#include "rp_amd.c"
+
+__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) amp (__global pw_t *pws, __global pw_t *pws_amp, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, const u32 combs_mode, const u32 gid_max)
+{
+  const u32 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  const u32 pw_len = pws[gid].pw_len;
+
+  u32x w0[4];
+  u32x w1[4];
+
+  w0[0] = pws[gid].i[ 0];
+  w0[1] = pws[gid].i[ 1];
+  w0[2] = pws[gid].i[ 2];
+  w0[3] = pws[gid].i[ 3];
+  w1[0] = pws[gid].i[ 4];
+  w1[1] = pws[gid].i[ 5];
+  w1[2] = pws[gid].i[ 6];
+  w1[3] = pws[gid].i[ 7];
+
+  const u32 out_len = apply_rules (rules_buf[0].cmds, w0, w1, pw_len);
+
+  pws_amp[gid].i[0] = w0[0];
+  pws_amp[gid].i[1] = w0[1];
+  pws_amp[gid].i[2] = w0[2];
+  pws_amp[gid].i[3] = w0[3];
+  pws_amp[gid].i[4] = w1[0];
+  pws_amp[gid].i[5] = w1[1];
+  pws_amp[gid].i[6] = w1[2];
+  pws_amp[gid].i[7] = w1[3];
+
+  pws_amp[gid].pw_len = out_len;
+}

amd/amp_a0_v4.cl

@@ -0,0 +1,52 @@
+/**
+ * Author......: Jens Steube <[email protected]>
+ * License.....: MIT
+ */
+
+#define VECT_SIZE4
+
+#include "include/constants.h"
+#include "include/kernel_vendor.h"
+#include "types_amd.c"
+
+static u32x swap_workaround (const u32x v)
+{
+  return (v << 24) + ((v & 0x0000FF00) << 8) + ((v & 0x00FF0000) >> 8) + (v >> 24);
+}
+
+#include "include/rp_gpu.h"
+#include "rp_amd.c"
+
+__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) amp (__global pw_t *pws, __global pw_t *pws_amp, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, const u32 combs_mode, const u32 gid_max)
+{
+  const u32 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  const u32 pw_len = pws[gid].pw_len;
+
+  u32x w0[4];
+  u32x w1[4];
+
+  w0[0] = pws[gid].i[ 0];
+  w0[1] = pws[gid].i[ 1];
+  w0[2] = pws[gid].i[ 2];
+  w0[3] = pws[gid].i[ 3];
+  w1[0] = pws[gid].i[ 4];
+  w1[1] = pws[gid].i[ 5];
+  w1[2] = pws[gid].i[ 6];
+  w1[3] = pws[gid].i[ 7];
+
+  const u32 out_len = apply_rules (rules_buf[0].cmds, w0, w1, pw_len);
+
+  pws_amp[gid].i[0] = w0[0];
+  pws_amp[gid].i[1] = w0[1];
+  pws_amp[gid].i[2] = w0[2];
+  pws_amp[gid].i[3] = w0[3];
+  pws_amp[gid].i[4] = w1[0];
+  pws_amp[gid].i[5] = w1[1];
+  pws_amp[gid].i[6] = w1[2];
+  pws_amp[gid].i[7] = w1[3];
+
+  pws_amp[gid].pw_len = out_len;
+}