diff -urpN john-1.7.6.orig//README_CUDA john-1.7.6/README_CUDA --- john-1.7.6.orig//README_CUDA 1970-01-01 00:00:00.000000000 +0000 +++ john-1.7.6/README_CUDA 2011-04-04 01:04:52.000000000 +0000 @@ -0,0 +1,12 @@ +README: +SHA256CUDA is supported/tested only on following builds: +linux-x86-mmx +linux-x86-64 +linux-x86-any +If you have problems while building patched JtR: + Please check that paths to yours CUDA library are proper in Makefile lines: + 23: CUDAPATH = /usr/local/cuda/lib + 24: CUDA64PATH = /usr/local/cuda/lib64 + Please check that yours PATH contains cuda lib,bin and include paths. + Please check that yours LD_LIBRARY_PATH contains cuda lib path + You can contact me at lukas[dot]odzioba[at]gmail[dot]com \ No newline at end of file diff -urpN john-1.7.6.orig//src/Makefile john-1.7.6/src/Makefile --- john-1.7.6.orig//src/Makefile 2010-06-13 21:12:37.000000000 +0000 +++ john-1.7.6/src/Makefile 2011-04-05 19:51:35.000000000 +0000 @@ -2,7 +2,7 @@ # This file is part of John the Ripper password cracker, # Copyright (c) 1996-2010 by Solar Designer # - +NVCC = nvcc CC = gcc AS = $(CC) LD = $(CC) @@ -19,7 +19,11 @@ OMPFLAGS = #OMPFLAGS = -fopenmp # Sun Studio with OpenMP (set the OMP_NUM_THREADS env var at runtime) #OMPFLAGS = -xopenmp -CFLAGS = -c -Wall -O2 -fomit-frame-pointer $(OMPFLAGS) +NVCCFLAGS = -c +CUDAPATH = /usr/local/cuda/lib +CUDA64PATH = /usr/local/cuda/lib64 +CUDALIB=$(CUDAPATH) +CFLAGS = -c -Wall -O2 -fomit-frame-pointer $(OMPFLAGS) ASFLAGS = -c $(OMPFLAGS) LDFLAGS = -s $(OMPFLAGS) OPT_NORMAL = -funroll-loops @@ -32,13 +36,14 @@ JOHN_OBJS = \ BF_fmt.o BF_std.o \ AFS_fmt.o \ LM_fmt.o \ + sha256cuda_fmt.o sha256.o \ batch.o bench.o charset.o common.o compiler.o config.o cracker.o \ crc32.o external.o formats.o getopt.o idle.o inc.o john.o list.o \ loader.o logger.o math.o memory.o misc.o options.o params.o path.o \ recovery.o rpp.o rules.o signals.o single.o status.o tty.o wordlist.o \ unshadow.o \ unafs.o \ - unique.o + unique.o BENCH_DES_OBJS_ORIG = \ DES_fmt.o DES_std.o @@ -148,7 +153,7 @@ linux-x86-64: $(MAKE) $(PROJ) \ JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86-64.o" \ CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \ - LDFLAGS="$(LDFLAGS) -lcrypt" + LDFLAGS="$(LDFLAGS) -L$(CUDA64PATH) -lcudart -lcrypt" linux-x86-64-32-sse2: $(LN) x86-sse.h arch.h @@ -178,14 +183,14 @@ linux-x86-mmx: $(MAKE) $(PROJ) \ JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86.o x86-mmx.o" \ CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \ - LDFLAGS="$(LDFLAGS) -lcrypt" + LDFLAGS="$(LDFLAGS) -L$(CUDAPATH) -lcudart -lcrypt" linux-x86-any: $(LN) x86-any.h arch.h $(MAKE) $(PROJ) \ JOHN_OBJS="$(JOHN_OBJS) c3_fmt.o x86.o" \ CFLAGS="$(CFLAGS) -DHAVE_CRYPT" \ - LDFLAGS="$(LDFLAGS) -lcrypt" + LDFLAGS="$(LDFLAGS) -L$(CUDAPATH) -lcudart -lcrypt" linux-alpha: $(LN) alpha.h arch.h @@ -776,6 +781,13 @@ generic.h: "$(BENCH_MD5_OBJS_DEPEND)" \ "$(BENCH_BF_OBJS_DEPEND)" +sha256.o: sha256.h cuda/sha256.cu + $(shell cd cuda; nvcc $(NVCCFLAGS) sha256.cu) + +sha256cuda_fmt.o: sha256.o sha256cuda_fmt.c sha256.h + cp cuda/sha256.o sha256.o + $(CC) $(CFLAGS) sha256cuda_fmt.c + bench: $(BENCH_OBJS) $(LD) $(BENCH_OBJS) $(LDFLAGS) -o bench @@ -855,8 +867,9 @@ depend: clean: $(RM) $(PROJ) $(PROJ_DOS) $(PROJ_WIN32) + $(RM) cuda/sha256.o $(RM) ../run/john.exe john-macosx-* *.o *.bak core $(RM) detect bench generic.h arch.h tmp.s $(CP) $(NULL) Makefile.dep -include Makefile.dep +include Makefile.dep \ No newline at end of file diff -urpN john-1.7.6.orig//src/cuda/sha256.cu john-1.7.6/src/cuda/sha256.cu --- john-1.7.6.orig//src/cuda/sha256.cu 1970-01-01 00:00:00.000000000 +0000 +++ john-1.7.6/src/cuda/sha256.cu 2011-05-16 21:19:12.293212458 +0000 @@ -0,0 +1,79 @@ +#include +#include +#include +#include +#include "../sha256.h" + +const uint DATA_IN_SIZE=KEYS_PER_CRYPT*sizeof(sha256_password); +const uint DATA_OUT_SIZE=KEYS_PER_CRYPT*sizeof(sha256_hash); + +sha256_password *cuda_data=NULL; ///candidates +sha256_hash *cuda_data_out=NULL; ///sha256(candidates) + +sha256_password *host_data=NULL; +sha256_hash *host_data_out=NULL; + +extern "C" void mem_init(sha256_password *,sha256_hash*); +extern "C" void mem_clear(void); +extern "C" void sha256(void); + +__global__ void kernel_sha256(sha256_password *data,sha256_hash* data_out); + +void sha256() +{ + dim3 dimGrid(BLOCKS); /// todo auto detection + dim3 dimBlock(THREADS); + kernel_sha256<<>>(cuda_data,cuda_data_out); + cudaThreadSynchronize(); +} + +void mem_init(sha256_password *p,sha256_hash *h) +{ + host_data=p; + host_data_out=h; + cudaMalloc(&cuda_data,DATA_IN_SIZE); + cudaMalloc(&cuda_data_out,DATA_OUT_SIZE); + cudaMemcpy(cuda_data,host_data,DATA_IN_SIZE,cudaMemcpyHostToDevice); +} + +void mem_clear() +{ + cudaMemcpy(host_data_out,cuda_data_out,DATA_OUT_SIZE,cudaMemcpyDeviceToHost); + cudaFree(cuda_data); + cudaFree(cuda_data_out); +} + +__global__ void kernel_sha256(sha256_password *data,sha256_hash* data_out){ /// todo - use shared memory + uint32_t idx = blockIdx.x*blockDim.x + threadIdx.x; + const uint32_t _h[]={ + 0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0xa54ff53a, 0x510e527f, 0x9b05688c, 0x1f83d9ab, 0x5be0cd19}; + const uint32_t k[]={ + 0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, 0x923f82a4, 0xab1c5ed5, + 0xd807aa98, 0x12835b01, 0x243185be, 0x550c7dc3, 0x72be5d74, 0x80deb1fe, 0x9bdc06a7, 0xc19bf174, + 0xe49b69c1, 0xefbe4786, 0x0fc19dc6, 0x240ca1cc, 0x2de92c6f, 0x4a7484aa, 0x5cb0a9dc, 0x76f988da, + 0x983e5152, 0xa831c66d, 0xb00327c8, 0xbf597fc7, 0xc6e00bf3, 0xd5a79147, 0x06ca6351, 0x14292967, + 0x27b70a85, 0x2e1b2138, 0x4d2c6dfc, 0x53380d13, 0x650a7354, 0x766a0abb, 0x81c2c92e, 0x92722c85, + 0xa2bfe8a1, 0xa81a664b, 0xc24b8b70, 0xc76c51a3, 0xd192e819, 0xd6990624, 0xf40e3585, 0x106aa070, + 0x19a4c116, 0x1e376c08, 0x2748774c, 0x34b0bcb5, 0x391c0cb3, 0x4ed8aa4a, 0x5b9cca4f, 0x682e6ff3, + 0x748f82ee, 0x78a5636f, 0x84c87814, 0x8cc70208, 0x90befffa, 0xa4506ceb, 0xbef9a3f7, 0xc67178f2}; + uint32_t w[64]={0}; + sha256_hash* out=&data_out[idx]; + + for(uint32_t j=0;j<64;j++){ ///todo unroll loop,nvcc option? + if(j<16) w[j]=data[idx].v[j]; + else w[j]=sigma1(w[j-2])+w[j-7]+sigma0(w[j-15])+w[j-16]; + } + + uint32_t a=_h[0];uint32_t b=_h[1];uint32_t c=_h[2];uint32_t d=_h[3]; + uint32_t e=_h[4];uint32_t f=_h[5];uint32_t g=_h[6];uint32_t h=_h[7]; + + for(uint32_t j=0;j<64;j++){ ///todo unroll loop, nvcc option? + uint32_t t1=h+Sigma1(e)+Ch(e,f,g)+k[j]+w[j]; + uint32_t t2=Sigma0(a)+Maj(a,b,c); + h=g;g=f;f=e;e=d+t1;d=c;c=b;b=a;a=t1+t2; + } + out->v[0]=a+_h[0];out->v[1]=b+_h[1]; + out->v[2]=c+_h[2];out->v[3]=d+_h[3]; + out->v[4]=e+_h[4];out->v[5]=f+_h[5]; + out->v[6]=g+_h[6];out->v[7]=h+_h[7]; +} \ No newline at end of file diff -urpN john-1.7.6.orig//src/formats.c john-1.7.6/src/formats.c --- john-1.7.6.orig//src/formats.c 2010-06-09 14:40:04.000000000 +0000 +++ john-1.7.6/src/formats.c 2011-05-16 21:17:42.679837535 +0000 @@ -53,7 +53,7 @@ char *fmt_self_test(struct fmt_main *for return "valid"; ciphertext = format->methods.split(current->ciphertext, 0); plaintext = current->plaintext; - + binary = format->methods.binary(ciphertext); salt = format->methods.salt(ciphertext); diff -urpN john-1.7.6.orig//src/john.c john-1.7.6/src/john.c --- john-1.7.6.orig//src/john.c 2010-06-14 01:34:04.000000000 +0000 +++ john-1.7.6/src/john.c 2011-04-05 19:51:35.000000000 +0000 @@ -38,6 +38,7 @@ extern int CPU_detect(void); extern struct fmt_main fmt_DES, fmt_BSDI, fmt_MD5, fmt_BF; extern struct fmt_main fmt_AFS, fmt_LM; +extern struct fmt_main fmt_SHA256CUDA; #ifdef HAVE_CRYPT extern struct fmt_main fmt_crypt; #endif @@ -69,10 +70,12 @@ static void john_register_all(void) john_register_one(&fmt_BF); john_register_one(&fmt_AFS); john_register_one(&fmt_LM); + john_register_one(&fmt_SHA256CUDA); + + #ifdef HAVE_CRYPT john_register_one(&fmt_crypt); #endif - if (!fmt_list) { fprintf(stderr, "Unknown ciphertext format name requested\n"); error(); diff -urpN john-1.7.6.orig//src/options.c john-1.7.6/src/options.c --- john-1.7.6.orig//src/options.c 2010-06-12 14:03:50.000000000 +0000 +++ john-1.7.6/src/options.c 2011-04-05 19:51:35.000000000 +0000 @@ -101,7 +101,7 @@ static struct opt_entry opt_list[] = { "--salts=[-]COUNT load salts with[out] at least COUNT passwords " \ "only\n" \ "--format=NAME force hash type NAME: " \ - "DES/BSDI/MD5/BF/AFS/LM/crypt\n" \ + "DES/BSDI/MD5/BF/AFS/LM/crypt/SHA256CUDA\n" \ "--save-memory=LEVEL enable memory saving, at LEVEL 1..3\n" void opt_init(char *name, int argc, char **argv) diff -urpN john-1.7.6.orig//src/sha256.h john-1.7.6/src/sha256.h --- john-1.7.6.orig//src/sha256.h 1970-01-01 00:00:00.000000000 +0000 +++ john-1.7.6/src/sha256.h 2011-05-16 21:19:47.347837074 +0000 @@ -0,0 +1,29 @@ +#ifndef _SHA256_H +#define _SHA256_H + +#ifndef uint32_t + #define uint32_t unsigned int +#endif + +#define rol(x,n) ((x << n) | (x >> (32-n))) /** todo - replace with hardware instructions **/ +#define ror(x,n) ((x >> n) | (x << (32-n))) /** todo - as above **/ +#define Ch(x,y,z) ((x & y) ^ ( (~x) & z)) +#define Maj(x,y,z) ((x & y) ^ (x & z) ^ (y & z)) +#define Sigma0(x) ((ror(x,2)) ^ (ror(x,13)) ^ (ror(x,22))) +#define Sigma1(x) ((ror(x,6)) ^ (ror(x,11)) ^ (ror(x,25))) +#define sigma0(x) ((ror(x,7)) ^ (ror(x,18)) ^(x>>3)) +#define sigma1(x) ((ror(x,17)) ^ (ror(x,19)) ^(x>>10)) + +#define THREADS 128 +#define BLOCKS 256 +#define KEYS_PER_CRYPT THREADS*BLOCKS ///256 blocks * 256 threads todo - auto detection + +typedef struct{ + uint32_t v[16]; ///512bit +}sha256_password; + +typedef struct{ + uint32_t v[8]; ///256bit +}sha256_hash; + +#endif \ No newline at end of file diff -urpN john-1.7.6.orig//src/sha256cuda_fmt.c john-1.7.6/src/sha256cuda_fmt.c --- john-1.7.6.orig//src/sha256cuda_fmt.c 1970-01-01 00:00:00.000000000 +0000 +++ john-1.7.6/src/sha256cuda_fmt.c 2011-05-16 21:17:18.541211784 +0000 @@ -0,0 +1,234 @@ +/* +* This software is Copyright (c) 2011 Lukas Odzioba +* +* and it is hereby released to the general public under the following terms: +* Redistribution and use in source and binary forms, with or without modification, are permitted. +*/ +#include +#include "arch.h" +#include "formats.h" +#include "common.h" ///for ARCH_WORD_32,atoi16 +#include "misc.h" ///for strnzcpy +#include "sha256.h" + +#define FORMAT_LABEL "sha256cuda" +#define FORMAT_NAME "SHA256CUDA" + +#define SHA_TYPE "SHA256" + +#define BENCHMARK_COMMENT "" +#define BENCHMARK_LENGTH -1 /// Raw benchmark + +#define PLAINTEXT_LENGTH 54 +#define CIPHERTEXT_LENGTH 64 ///256bit +#define BINARY_SIZE 32 + +#define SALT_SIZE 0 + +#define MIN_KEYS_PER_CRYPT KEYS_PER_CRYPT +#define MAX_KEYS_PER_CRYPT KEYS_PER_CRYPT + +extern void mem_init(sha256_password *,sha256_hash*); +extern void mem_clear(void); +extern void sha256(void); +static char saved_keys[MAX_KEYS_PER_CRYPT][PLAINTEXT_LENGTH+1]; /** plaintext ciphertexts **/ +static sha256_password inbuffer[MAX_KEYS_PER_CRYPT]; /** binary ciphertexts **/ +static sha256_hash outbuffer[MAX_KEYS_PER_CRYPT]; /** calculated hashes **/ + +static void preproc(char *key, int index){ /// todo - clean it + uint32_t dl=strlen(key),j; + uint32_t *blocks = inbuffer[index].v; + memset(inbuffer[index].v,0,sizeof(sha256_password)); + for(j=0;j='0' && ciphertext[i]<='9')|| + (ciphertext[i]>='a' && ciphertext[i]<='f')|| + (ciphertext[i]>='A' && ciphertext[i]<='Z'))) + return 0; + } + return 1; +}; + + +static void *binary(char *ciphertext){ + static char realcipher[BINARY_SIZE]; + memset(realcipher,0,BINARY_SIZE); + int i; + for(i=0;i