Fixed NTLM9; added NTLM9 tests.

jtesta · Jun 28, 2019 · 865bbc0 · 865bbc0
1 parent 162639b
commit 865bbc0
Show file tree

Hide file tree

Showing 20 changed files with 681 additions and 28 deletions.
diff --git a/CL/crackalack_ntlm9.cl b/CL/crackalack_ntlm9.cl
@@ -8,15 +8,15 @@ __kernel void crackalack_ntlm9(
     __global unsigned int *unused3,
     __global unsigned int *unused4,
     __global unsigned int *unused5,
-    __global unsigned int *unused6,
+    __global unsigned int *g_chain_len,
     __global unsigned long *g_indices,
-    __global unsigned int *unused7) {
+    __global unsigned int *g_pos_start) {
   unsigned long index = g_indices[get_global_id(0)];
   unsigned char plaintext[9];
 
 
-  for (unsigned int pos = 0; pos < 1349999; pos++) {
-    index_to_plaintext_ntlm9(index, charset, plaintext);
+  for (unsigned int pos = *g_pos_start; pos < (*g_chain_len - 1); pos++) {
+    index_to_plaintext_ntlm9(index, plaintext);
     index = hash_to_index_ntlm9(hash_ntlm9(plaintext), pos);
   }
 

diff --git a/CL/ntlm9_functions.cl b/CL/ntlm9_functions.cl
@@ -1,16 +1,9 @@
 __constant char charset[] = " !\"#$%&'()*+,-./0123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[\\]^_`abcdefghijklmnopqrstuvwxyz{|}~";
 
 
-/* Since nobody else has made 9-character rainbow tables, we're free to take some of
- * our own artistic liberties...
- *
- * We have a 64-bit number that we need to map to a 9-character plaintext.  This
- * means if the character set is of length 128 or less, we can break the number into
- * nine 7-bit fragments, and use them to index into the character set.  This ends up
- * being 2.4x faster than the standard division method (below)! */
-inline void index_to_plaintext_ntlm9(unsigned long index, __constant char *charset, unsigned char *plaintext) {
-  for (int i = 8; i >=0; --i) {
-    plaintext[i] = charset[ index % 95 ];
+inline void index_to_plaintext_ntlm9(unsigned long index, unsigned char *plaintext) {
+  for (int i = 8; i >= 0; i--) {
+    plaintext[i] = charset[index % 95];
     index /= 95;
   }
 
@@ -156,7 +149,7 @@ inline unsigned long hash_to_index_ntlm9(unsigned long hash, unsigned int pos) {
   return (hash + pos) % 630249409724609375UL;
 }
 
-
+/*
 inline unsigned long hash_char_to_index_ntlm9(__global unsigned char *hash_value, unsigned int pos) {
   unsigned long ret = hash_value[7];
   ret <<= 8;
@@ -176,3 +169,4 @@ inline unsigned long hash_char_to_index_ntlm9(__global unsigned char *hash_value
 
   return (ret + pos) % 630249409724609375UL;
 }
+*/
diff --git a/CL/test_chain_ntlm9.cl b/CL/test_chain_ntlm9.cl
@@ -0,0 +1,15 @@
+#include "ntlm9_functions.cl"
+
+__kernel void test_chain_ntlm9(__global unsigned int *g_chain_len, __global unsigned long *g_start, __global unsigned long *g_end) {
+
+  unsigned long index = *g_start;
+  unsigned char plaintext[9];
+
+
+  for (unsigned int pos = 0; pos < (*g_chain_len - 1); pos++) {
+    index_to_plaintext_ntlm9(index, plaintext);
+    index = hash_to_index_ntlm9(hash_ntlm9(plaintext), pos);
+  }
+
+  *g_end = index;
+}
diff --git a/CL/test_hash_ntlm9.cl b/CL/test_hash_ntlm9.cl
@@ -0,0 +1,11 @@
+#include "ntlm9_functions.cl"
+
+__kernel void test_hash_ntlm9(__global char *g_input, __global unsigned long *g_output) {
+  unsigned char input[9];
+
+  for (int i = 0; i < 9; i++)
+    input[i] = g_input[i];
+
+  *g_output = hash_ntlm9(input);
+  return;
+}
diff --git a/CL/test_hash_to_index_ntlm9.cl b/CL/test_hash_to_index_ntlm9.cl
@@ -0,0 +1,7 @@
+#include "ntlm9_functions.cl"
+
+__kernel void test_hash_to_index_ntlm9(__global unsigned long *g_hash, __global unsigned int *g_pos, __global unsigned long *g_index) {
+
+  *g_index = hash_to_index_ntlm9((unsigned long)*g_hash, (unsigned int)*g_pos);
+
+}
diff --git a/CL/test_index_to_plaintext_ntlm9.cl b/CL/test_index_to_plaintext_ntlm9.cl
@@ -0,0 +1,13 @@
+#include "ntlm9_functions.cl"
+#include "shared.h"
+
+__kernel void test_index_to_plaintext_ntlm9(__global unsigned long *g_index, __global unsigned char *g_plaintext) {
+  unsigned char plaintext[9];
+
+  index_to_plaintext_ntlm9((unsigned long)*g_index, plaintext);
+
+  for (int i = 0; i < 9; i++)
+    g_plaintext[i] = plaintext[i];
+
+  return;
+}
diff --git a/Makefile b/Makefile
@@ -42,8 +42,8 @@ all:	$(GEN_PROG) $(UNITTEST_PROG) $(LOOKUP_PROG) $(RTC2RT_PROG) $(GETCHAIN_PROG)
 $(GEN_PROG):	charset.o clock.o cpu_rt_functions.o crackalack_gen.o file_lock.o gws.o hash_validate.o misc.o opencl_setup.o rtc_decompress.o verify.o
 	$(CC) $(COMPILE_OPTIONS) -o $(GEN_PROG) charset.o clock.o cpu_rt_functions.o crackalack_gen.o file_lock.o gws.o hash_validate.o misc.o opencl_setup.o rtc_decompress.o verify.o $(LINK_OPTIONS)
 
-$(UNITTEST_PROG):	charset.o cpu_rt_functions.o crackalack_unit_tests.o hash_validate.o misc.o opencl_setup.o  test_chain.o test_hash.o test_hash_to_index.o test_index_to_plaintext.o test_shared.o file_lock.o
-	$(CC) $(COMPILE_OPTIONS) -o $(UNITTEST_PROG) charset.o cpu_rt_functions.o crackalack_unit_tests.o hash_validate.o misc.o opencl_setup.o test_chain.o test_hash.o test_hash_to_index.o test_index_to_plaintext.o test_shared.o file_lock.o $(LINK_OPTIONS)
+$(UNITTEST_PROG):	charset.o cpu_rt_functions.o crackalack_unit_tests.o hash_validate.o misc.o opencl_setup.o  test_chain.o test_chain_ntlm9.o test_hash.o test_hash_ntlm9.o test_hash_to_index.o test_hash_to_index_ntlm9.o test_index_to_plaintext.o test_index_to_plaintext_ntlm9.o test_shared.o file_lock.o
+	$(CC) $(COMPILE_OPTIONS) -o $(UNITTEST_PROG) charset.o cpu_rt_functions.o crackalack_unit_tests.o hash_validate.o misc.o opencl_setup.o test_chain.o test_chain_ntlm9.o test_hash.o test_hash_ntlm9.o test_hash_to_index.o test_hash_to_index_ntlm9.o test_index_to_plaintext.o test_index_to_plaintext_ntlm9.o test_shared.o file_lock.o $(LINK_OPTIONS)
 
 $(GETCHAIN_PROG):	get_chain.o
 	$(CC) $(COMPILE_OPTIONS) -o $(GETCHAIN_PROG) get_chain.o $(LINK_OPTIONS)

diff --git a/crackalack_gen.c b/crackalack_gen.c
@@ -224,13 +224,13 @@ void *host_thread(void *ptr) {
     if (args->gpu.device_number == 0) { /* Only the first thread prints this. */
       printf("Note: optimized NTLM8 kernel will be used.\n"); fflush(stdout);
     }
-  } /*else if (is_ntlm9(args->hash_type, args->charset, args->plaintext_len_min, args->plaintext_len_max, args->reduction_offset, args->chain_len)) {
+  } else if (is_ntlm9(args->hash_type, args->charset, args->plaintext_len_min, args->plaintext_len_max, args->reduction_offset, args->chain_len)) {
     kernel_path = CRACKALACK_NTLM9_KERNEL_PATH;
     kernel_name = "crackalack_ntlm9";
-    if (args->gpu.device_number == 0) { * Only the first thread prints this. *
+    if (args->gpu.device_number == 0) { /* Only the first thread prints this. */
       printf("Note: optimized NTLM9 kernel will be used.\n"); fflush(stdout);
     }
-  }*/
+  }
 
   /* Get the number of compute units in this device. */
   get_device_uint(gpu->device, CL_DEVICE_MAX_COMPUTE_UNITS, &(gpu->num_work_units));

diff --git a/crackalack_unit_tests.c b/crackalack_unit_tests.c
@@ -25,10 +25,13 @@
 #include "opencl_setup.h"
 #include "shared.h"
 #include "test_chain.h"
-#include "test_des.h"
+#include "test_chain_ntlm9.h"
 #include "test_hash.h"
+#include "test_hash_ntlm9.h"
 #include "test_hash_to_index.h"
+#include "test_hash_to_index_ntlm9.h"
 #include "test_index_to_plaintext.h"
+#include "test_index_to_plaintext_ntlm9.h"
 #include "version.h"
 
 
@@ -96,9 +99,11 @@ int main(int ac, char **av) {
   CLRELEASEPROGRAM(program);
   */
 
+
   /* index_to_plaintext() tests. */
+  hash_type = HASH_NTLM;
   load_kernel(context, num_devices, devices, "test_index_to_plaintext.cl", "test_index_to_plaintext", &program, &kernel, hash_type);
-  printf("Running index_to_plaintext() tests... "); fflush(stdout);
+  printf("Running NTLM index_to_plaintext() tests... "); fflush(stdout);
   if (!test_index_to_plaintext(devices[0], context, kernel)) {
     ret = -1;
     all_tests_passed = 0;
@@ -110,7 +115,22 @@ int main(int ac, char **av) {
   CLRELEASEPROGRAM(program);
 
 
+  /* index_to_plaintext_ntlm9() tests. */
+  load_kernel(context, num_devices, devices, "test_index_to_plaintext_ntlm9.cl", "test_index_to_plaintext_ntlm9", &program, &kernel, hash_type);
+  printf("Running NTLM9 index_to_plaintext_ntlm9() tests... "); fflush(stdout);
+  if (!test_index_to_plaintext_ntlm9(devices[0], context, kernel)) {
+    ret = -1;
+    all_tests_passed = 0;
+    PRINT_FAILED();
+  } else
+    PRINT_PASSED();
+
+  CLRELEASEKERNEL(kernel);
+  CLRELEASEPROGRAM(program);
+
+
   /* Hash tests. */
+  /*
   printf("Running LM hash tests... "); fflush(stdout);
   hash_type = HASH_LM;
   load_kernel(context, num_devices, devices, "test_hash.cl", "test_hash", &program, &kernel, hash_type);
@@ -123,6 +143,7 @@ int main(int ac, char **av) {
 
   CLRELEASEKERNEL(kernel);
   CLRELEASEPROGRAM(program);
+  */
 
 
   printf("Running NTLM hash tests... "); fflush(stdout);
@@ -137,9 +158,23 @@ int main(int ac, char **av) {
 
   CLRELEASEKERNEL(kernel);
   CLRELEASEPROGRAM(program);
-
+
+
+  printf("Running NTLM9 hash tests... "); fflush(stdout);
+  load_kernel(context, num_devices, devices, "test_hash_ntlm9.cl", "test_hash_ntlm9", &program, &kernel, hash_type);
+  if (!test_hash_ntlm9(devices[0], context, kernel)) {
+    ret = -1;
+    all_tests_passed = 0;
+    PRINT_FAILED();
+  } else
+    PRINT_PASSED();
+
+  CLRELEASEKERNEL(kernel);
+  CLRELEASEPROGRAM(program);
+
 
   /* hash_to_index() tests. */
+  /*
   printf("Running LM hash_to_index() tests... "); fflush(stdout);
   hash_type = HASH_LM;
   load_kernel(context, num_devices, devices, "test_hash_to_index.cl", "test_hash_to_index", &program, &kernel, hash_type);
@@ -152,7 +187,7 @@ int main(int ac, char **av) {
 
   CLRELEASEKERNEL(kernel);
   CLRELEASEPROGRAM(program);
-
+  */
 
   printf("Running NTLM hash_to_index() tests... "); fflush(stdout);
   hash_type = HASH_NTLM;
@@ -168,7 +203,21 @@ int main(int ac, char **av) {
   CLRELEASEPROGRAM(program);
 
 
+  printf("Running NTLM9 hash_to_index() tests... "); fflush(stdout);
+  load_kernel(context, num_devices, devices, "test_hash_to_index_ntlm9.cl", "test_hash_to_index_ntlm9", &program, &kernel, hash_type);
+  if (!test_h2i_ntlm9(devices[0], context, kernel)) {
+    ret = -1;
+    all_tests_passed = 0;
+    PRINT_FAILED();
+  } else
+    PRINT_PASSED();
+
+  CLRELEASEKERNEL(kernel);
+  CLRELEASEPROGRAM(program);
+
+
   /* Chain tests. */
+  /*
   printf("Running LM chain tests... "); fflush(stdout);
   hash_type = HASH_LM;
   load_kernel(context, num_devices, devices, "test_chain.cl", "test_chain", &program, &kernel, hash_type);
@@ -181,6 +230,8 @@ int main(int ac, char **av) {
 
   CLRELEASEKERNEL(kernel);
   CLRELEASEPROGRAM(program);
+  */
+
 
   printf("Running NTLM chain tests... "); fflush(stdout);
   hash_type = HASH_NTLM;
@@ -196,6 +247,21 @@ int main(int ac, char **av) {
   CLRELEASEPROGRAM(program);
 
 
+  printf("Running NTLM9 chain tests... "); fflush(stdout);
+  hash_type = HASH_NTLM;
+  /*load_kernel(context, num_devices, devices, "test_chain_ntlm9.cl", "test_chain_ntlm9", &program, &kernel, hash_type);*/
+  load_kernel(context, num_devices, devices, "crackalack_ntlm9.cl", "crackalack_ntlm9", &program, &kernel, hash_type);
+  if (!test_chain_ntlm9(devices[0], context, kernel)) {
+    ret = -1;
+    all_tests_passed = 0;
+    PRINT_FAILED();
+  } else
+    PRINT_PASSED();
+
+  CLRELEASEKERNEL(kernel);
+  CLRELEASEPROGRAM(program);
+
+
   if (all_tests_passed)
     printf("\n\t%sALL UNIT TESTS PASS!%s\n\n", GREENB, CLR);
   else

diff --git a/misc.c b/misc.c
@@ -222,7 +222,7 @@ unsigned int is_ntlm9(unsigned int hash_type, char *charset, unsigned int plaint
       (plaintext_len_min == 9) && \
       (plaintext_len_max == 9) && \
       (reduction_offset == 0) && \
-      (chain_len == 1350000))
+      (chain_len == 803000))
     return 1;
   else
     return 0;