diff -urpN magnum-jumbo/src/Makefile magnum-jumbo_cuda_warnings_fix//src/Makefile --- magnum-jumbo/src/Makefile 2012-06-29 16:35:22.760939228 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/Makefile 2012-06-29 16:36:14.991939548 +0000 @@ -57,7 +57,7 @@ ifdef HAVE_NSS NSS_LDFLAGS = `pkg-config --libs nss` endif -CFLAGS = -c -Wall -O2 -fomit-frame-pointer -I/usr/local/include $(HAVE_NSS) $(OMPFLAGS) $(JOHN_CFLAGS) $(AMDAPP) +CFLAGS = -c -Wall -O2 -fomit-frame-pointer -Wdeclaration-after-statement -I/usr/local/include $(HAVE_NSS) $(OMPFLAGS) $(JOHN_CFLAGS) $(AMDAPP) # -DHAVE_SKEY # CFLAGS for use on the main john.c file only CFLAGS_MAIN = $(CFLAGS) diff -urpN magnum-jumbo/src/common-opencl.c magnum-jumbo_cuda_warnings_fix//src/common-opencl.c --- magnum-jumbo/src/common-opencl.c 2012-06-29 16:35:22.763939202 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/common-opencl.c 2012-06-29 18:11:44.360939478 +0000 @@ -54,7 +54,7 @@ static void dev_init(unsigned int dev_id { cl_platform_id platform[MAX_PLATFORMS]; cl_uint num_platforms, device_num; - + cl_context_properties properties[3]; assert(dev_id < MAXGPUS); ///Find CPU's HANDLE_CLERROR(clGetPlatformIDs(MAX_PLATFORMS, platform, @@ -67,11 +67,10 @@ static void dev_init(unsigned int dev_id "No OpenCL device of that type exist"); fprintf(stderr, "OpenCL platform %d: %s, %d device(s).\n", platform_id, opencl_log, device_num); - cl_context_properties properties[] = { - CL_CONTEXT_PLATFORM, - (cl_context_properties) platform[platform_id], - 0 - }; + + properties[0] = CL_CONTEXT_PLATFORM; + properties[1] = (cl_context_properties) platform[platform_id]; + properties[2] = 0; HANDLE_CLERROR(clGetDeviceInfo(devices[dev_id], CL_DEVICE_NAME, sizeof(opencl_log), opencl_log, NULL), "Error querying DEVICE_NAME"); @@ -107,14 +106,14 @@ static char *include_source(char *pathna static void build_kernel(int dev_id) { - assert(kernel_loaded); + cl_int build_code; const char *srcptr[] = { kernel_source }; + assert(kernel_loaded); program[dev_id] = clCreateProgramWithSource(context[dev_id], 1, srcptr, NULL, &ret_code); HANDLE_CLERROR(ret_code, "Error while creating program"); - cl_int build_code; build_code = clBuildProgram(program[dev_id], 0, NULL, include_source("$JOHN/", dev_id), NULL, NULL); @@ -162,11 +161,6 @@ void opencl_find_best_workgroup(struct f int i; size_t max_group_size, wg_multiple; - HANDLE_CLERROR(clGetKernelWorkGroupInfo(crypt_kernel, devices[gpu_id], - CL_KERNEL_WORK_GROUP_SIZE, sizeof(max_group_size), - &max_group_size, NULL), - "Error while getting CL_KERNEL_WORK_GROUP_SIZE"); - #if __OPENCL_VERSION__ < 110 cl_device_type device_type; clGetDeviceInfo(devices[gpu_id], CL_DEVICE_TYPE, @@ -181,6 +175,12 @@ void opencl_find_best_workgroup(struct f sizeof(wg_multiple), &wg_multiple, NULL), "Error while getting CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE"); #endif + + HANDLE_CLERROR(clGetKernelWorkGroupInfo(crypt_kernel, devices[gpu_id], + CL_KERNEL_WORK_GROUP_SIZE, sizeof(max_group_size), + &max_group_size, NULL), + "Error while getting CL_KERNEL_WORK_GROUP_SIZE"); + ///Command Queue changing: ///1) Delete old CQ clReleaseCommandQueue(queue[gpu_id]); diff -urpN magnum-jumbo/src/common_opencl_pbkdf2.c magnum-jumbo_cuda_warnings_fix//src/common_opencl_pbkdf2.c --- magnum-jumbo/src/common_opencl_pbkdf2.c 2012-06-29 16:35:22.763939202 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/common_opencl_pbkdf2.c 2012-06-29 17:16:51.737189157 +0000 @@ -74,17 +74,18 @@ static void find_best_workgroup(int pltf cl_device_type dTyp; - event_ctr=0; - - HANDLE_CLERROR(clGetDeviceInfo(devid[pltform_no][dev_no],CL_DEVICE_TYPE,sizeof(cl_device_type),&dTyp,NULL),"Failed Device Info"); - - ///Set Dummy DCC hash , unicode salt and ascii salt(username) length cl_uint *dcc_hash_host=(cl_uint*)malloc(4*sizeof(cl_uint)*64000); cl_uint *dcc2_hash_host=(cl_uint*)malloc(4*sizeof(cl_uint)*64000); cl_uint salt_api[9],length=10; + event_ctr=0; + + HANDLE_CLERROR(clGetDeviceInfo(devid[pltform_no][dev_no],CL_DEVICE_TYPE,sizeof(cl_device_type),&dTyp,NULL),"Failed Device Info"); + + ///Set Dummy DCC hash , unicode salt and ascii salt(username) length + memset(dcc_hash_host,0xb5,4*sizeof(cl_uint)*64000); memset(salt_api,0xfe,9*sizeof(cl_uint)); @@ -306,10 +307,10 @@ static gpu_mem_buffer exec_pbkdf2(cl_uin if(PROFILE){ - HANDLE_CLERROR(CL_SUCCESS!=clWaitForEvents(1,&evnt),"SYNC FAILED\n"); - cl_ulong startTime, endTime; + HANDLE_CLERROR(CL_SUCCESS!=clWaitForEvents(1,&evnt),"SYNC FAILED\n"); + HANDLE_CLERROR(clFinish(cmdq[platform_no][dev_no]), "clFinish error"); clGetEventProfilingInfo(evnt, CL_PROFILING_COMMAND_SUBMIT, sizeof(cl_ulong), &startTime, NULL); diff -urpN magnum-jumbo/src/cuda_cryptmd5_fmt.c magnum-jumbo_cuda_warnings_fix//src/cuda_cryptmd5_fmt.c --- magnum-jumbo/src/cuda_cryptmd5_fmt.c 2012-06-29 16:35:22.765939107 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/cuda_cryptmd5_fmt.c 2012-06-29 16:41:16.151189534 +0000 @@ -115,7 +115,8 @@ static void init(struct fmt_main *pFmt) static int valid(char *ciphertext, struct fmt_main *pFmt) { uint8_t i, len = strlen(ciphertext), prefix = 0; - + char *p; + if (strncmp(ciphertext, md5_salt_prefix, strlen(md5_salt_prefix)) == 0) prefix |= 1; if (strncmp(ciphertext, apr1_salt_prefix, @@ -123,7 +124,7 @@ static int valid(char *ciphertext, struc prefix |= 2; if (prefix == 0) return 0; - char *p = strrchr(ciphertext, '$'); + p = strrchr(ciphertext, '$'); if (p == NULL) return 0; for (i = p - ciphertext + 1; i < len; i++) { @@ -147,32 +148,36 @@ static void to_binary(char *crypt, char #define _24bit_from_b64(I,B2,B1,B0) \ {\ - unsigned char c1=findb64(crypt[I+0]);\ - unsigned char c2=findb64(crypt[I+1]);\ - unsigned char c3=findb64(crypt[I+2]);\ - unsigned char c4=findb64(crypt[I+3]);\ - unsigned int w=c4<<18|c3<<12|c2<<6|c1;\ - unsigned char b2=w&0xff;w>>=8;\ - unsigned char b1=w&0xff;w>>=8;\ - unsigned char b0=w&0xff;w>>=8;\ + uint8_t c1,c2,c3,c4,b0,b1,b2;\ + uint32_t w;\ + c1=findb64(crypt[I+0]);\ + c2=findb64(crypt[I+1]);\ + c3=findb64(crypt[I+2]);\ + c4=findb64(crypt[I+3]);\ + w=c4<<18|c3<<12|c2<<6|c1;\ + b2=w&0xff;w>>=8;\ + b1=w&0xff;w>>=8;\ + b0=w&0xff;w>>=8;\ alt[B2]=b0;\ alt[B1]=b1;\ alt[B0]=b2;\ } + uint32_t w; _24bit_from_b64(0, 0, 6, 12); _24bit_from_b64(4, 1, 7, 13); _24bit_from_b64(8, 2, 8, 14); _24bit_from_b64(12, 3, 9, 15); _24bit_from_b64(16, 4, 10, 5); - uint32_t w = findb64(crypt[21]) << 6 | findb64(crypt[20]) << 0; + w = findb64(crypt[21]) << 6 | findb64(crypt[20]) << 0; alt[11] = (w & 0xff); } static void *binary(char *ciphertext) { static char b[BINARY_SIZE]; + char *p; memset(b, 0, BINARY_SIZE); - char *p = strrchr(ciphertext, '$') + 1; + p = strrchr(ciphertext, '$') + 1; to_binary(p, b); return (void *) b; } @@ -185,7 +190,7 @@ static void *salt(char *ciphertext) #endif static crypt_md5_salt ret; uint8_t i, *pos = (uint8_t *) ciphertext, *end; - char *dest = ret.salt; + char *p,*dest = ret.salt; if (strncmp(ciphertext, md5_salt_prefix, strlen(md5_salt_prefix)) == 0) { pos += strlen(md5_salt_prefix); ret.prefix = '1'; @@ -200,7 +205,7 @@ static void *salt(char *ciphertext) while (pos != end) *dest++ = *pos++; ret.length = i; - char *p = strrchr(ciphertext, '$') + 1; + p = strrchr(ciphertext, '$') + 1; to_binary(p,(char*) ret.hash); #ifdef CUDA_DEBUG puts("salted:"); diff -urpN magnum-jumbo/src/cuda_cryptsha256_fmt.c magnum-jumbo_cuda_warnings_fix//src/cuda_cryptsha256_fmt.c --- magnum-jumbo/src/cuda_cryptsha256_fmt.c 2012-06-29 16:35:22.765939107 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/cuda_cryptsha256_fmt.c 2012-06-29 16:45:52.798064357 +0000 @@ -1,6 +1,5 @@ /* -* This software is Copyright (c) 2011 Lukas Odzioba -* +* 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. */ @@ -102,10 +101,10 @@ static int valid(char *ciphertext,struct { uint32_t i, j; int len = strlen(ciphertext); - + char *p; if (strncmp(ciphertext, "$5$", 3) != 0) return 0; - char *p = strrchr(ciphertext, '$'); + p = strrchr(ciphertext, '$'); if (p == NULL) return 0; for (i = p - ciphertext + 1; i < len; i++) { @@ -134,19 +133,21 @@ static void magic(char *crypt, char *alt #define _24bit_from_b64(I,B2,B1,B0) \ {\ - unsigned char c1=findb64(crypt[I+0]);\ - unsigned char c2=findb64(crypt[I+1]);\ - unsigned char c3=findb64(crypt[I+2]);\ - unsigned char c4=findb64(crypt[I+3]);\ - unsigned int w=c4<<18|c3<<12|c2<<6|c1;\ - unsigned char b2=w&0xff;w>>=8;\ - unsigned char b1=w&0xff;w>>=8;\ - unsigned char b0=w&0xff;w>>=8;\ + uint8_t c1,c2,c3,c4,b0,b1,b2;\ + uint32_t w;\ + c1=findb64(crypt[I+0]);\ + c2=findb64(crypt[I+1]);\ + c3=findb64(crypt[I+2]);\ + c4=findb64(crypt[I+3]);\ + w=c4<<18|c3<<12|c2<<6|c1;\ + b2=w&0xff;w>>=8;\ + b1=w&0xff;w>>=8;\ + b0=w&0xff;w>>=8;\ alt[B2]=b0;\ alt[B1]=b1;\ alt[B0]=b2;\ } - + uint32_t w; _24bit_from_b64(0, 0, 10, 20); _24bit_from_b64(4, 21, 1, 11); _24bit_from_b64(8, 12, 22, 2); @@ -157,7 +158,7 @@ static void magic(char *crypt, char *alt _24bit_from_b64(28, 27, 7, 17); _24bit_from_b64(32, 18, 28, 8); _24bit_from_b64(36, 9, 19, 29); - unsigned int w = + w = findb64(crypt[42]) << 12 | findb64(crypt[41]) << 6 | findb64(crypt[40]); alt[30] = w & 0xff; @@ -166,29 +167,27 @@ static void magic(char *crypt, char *alt w >>= 8; } - static void *binary(char *ciphertext) { static char b[BINARY_SIZE]; + char *p; memset(b, 0, BINARY_SIZE); - char *p = strrchr(ciphertext, '$'); + p = strrchr(ciphertext, '$'); if(p!=NULL) magic(p+1, b); return (void *) b; } - static void *salt(char *ciphertext) { int end = 0, i, len = strlen(ciphertext); + static unsigned char ret[64]; for (i = len - 1; i >= 0; i--) if (ciphertext[i] == '$') { end = i; break; } - - static unsigned char ret[64]; for (i = 0; i < end; i++) ret[i] = ciphertext[i]; ret[end] = 0; @@ -230,13 +229,12 @@ static int binary_hash_6(void *binary) return ((ARCH_WORD_32 *) binary)[0] & 0x7ffffff; } - static void set_salt(void *salt) { unsigned char *s = salt; int len = strlen(salt); - memcpy(currentsalt,s,len+1); unsigned char offset = 0; + memcpy(currentsalt,s,len+1); host_salt.rounds = ROUNDS_DEFAULT; if (strncmp((char *) "$5$", (char *) currentsalt, 3) == 0) diff -urpN magnum-jumbo/src/cuda_cryptsha512_fmt.c magnum-jumbo_cuda_warnings_fix//src/cuda_cryptsha512_fmt.c --- magnum-jumbo/src/cuda_cryptsha512_fmt.c 2012-06-29 16:35:22.765939107 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/cuda_cryptsha512_fmt.c 2012-06-29 16:58:38.627939797 +0000 @@ -1,5 +1,5 @@ /* -* This software is Copyright (c) 2011 Lukas Odzioba +* 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. */ @@ -82,10 +82,11 @@ static int valid(char *ciphertext,struct { uint32_t i, j; int len = strlen(ciphertext); - + char *p; + if (strncmp(ciphertext, "$6$", 3) != 0) return 0; - char *p = strrchr(ciphertext, '$'); + p = strrchr(ciphertext, '$'); if (p == NULL) return 0; for (i = p - ciphertext + 1; i < len; i++) { @@ -114,18 +115,21 @@ static void magic(char *crypt, unsigned #define _24bit_from_b64(I,B2,B1,B0) \ {\ - unsigned char c1=findb64(crypt[I+0]);\ - unsigned char c2=findb64(crypt[I+1]);\ - unsigned char c3=findb64(crypt[I+2]);\ - unsigned char c4=findb64(crypt[I+3]);\ - unsigned int w=c4<<18|c3<<12|c2<<6|c1;\ - unsigned char b2=w&0xff;w>>=8;\ - unsigned char b1=w&0xff;w>>=8;\ - unsigned char b0=w&0xff;w>>=8;\ + uint8_t c1,c2,c3,c4,b0,b1,b2;\ + uint32_t w;\ + c1=findb64(crypt[I+0]);\ + c2=findb64(crypt[I+1]);\ + c3=findb64(crypt[I+2]);\ + c4=findb64(crypt[I+3]);\ + w=c4<<18|c3<<12|c2<<6|c1;\ + b2=w&0xff;w>>=8;\ + b1=w&0xff;w>>=8;\ + b0=w&0xff;w>>=8;\ alt[B2]=b0;\ alt[B1]=b1;\ alt[B0]=b2;\ } + uint32_t w; _24bit_from_b64(0, 0, 21, 42); _24bit_from_b64(4, 22, 43, 1); _24bit_from_b64(8, 44, 2, 23); @@ -148,8 +152,7 @@ static void magic(char *crypt, unsigned _24bit_from_b64(76, 40, 61, 19); _24bit_from_b64(80, 62, 20, 41); - - uint32_t w = findb64(crypt[85]) << 6 | findb64(crypt[84]) << 0; + w = findb64(crypt[85]) << 6 | findb64(crypt[84]) << 0; alt[63] = (w & 0xff); } @@ -157,8 +160,8 @@ static void magic(char *crypt, unsigned static void *binary(char *ciphertext) { static unsigned char b[BINARY_SIZE]; - memset(b, 0, BINARY_SIZE); char *p = strrchr(ciphertext, '$'); + memset(b, 0, BINARY_SIZE); if(p!=NULL) magic(p+1, b); return (void *) b; @@ -168,13 +171,13 @@ static void *binary(char *ciphertext) static void *salt(char *ciphertext) { int end = 0, i, len = strlen(ciphertext); + static unsigned char ret[50]; for (i = len - 1; i >= 0; i--) if (ciphertext[i] == '$') { end = i; break; } - static unsigned char ret[50]; for (i = 0; i < end; i++) ret[i] = ciphertext[i]; ret[end] = 0; @@ -220,9 +223,9 @@ static void set_salt(void *salt) { unsigned char *s = salt; int len = strlen(salt); - memcpy(currentsalt,s,len+1); unsigned char offset = 0; _salt.rounds = ROUNDS_DEFAULT; + memcpy(currentsalt,s,len+1); if (strncmp((char *) "$6$", (char *) currentsalt, 3) == 0) offset += 3; diff -urpN magnum-jumbo/src/cuda_mscash2_fmt.c magnum-jumbo_cuda_warnings_fix//src/cuda_mscash2_fmt.c --- magnum-jumbo/src/cuda_mscash2_fmt.c 2012-06-29 16:35:22.765939107 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/cuda_mscash2_fmt.c 2012-06-29 17:04:38.173939633 +0000 @@ -1,5 +1,5 @@ /* -* This software is Copyright (c) 2011 Lukas Odzioba +* 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. * Based on S3nf implementation http://openwall.info/wiki/john/MSCash2 @@ -86,10 +86,12 @@ static void init(struct fmt_main *pFmt) static int valid(char *ciphertext, struct fmt_main *pFmt) { + char *pos,*hash = strrchr(ciphertext, '#') + 1; + int hashlength = 0; + int saltlength = 0; if (strncmp(ciphertext, mscash2_prefix, strlen(mscash2_prefix)) != 0) return 0; - char *hash = strrchr(ciphertext, '#') + 1; - int hashlength = 0; + if (hash == NULL) return 0; while (hash < ciphertext + strlen(ciphertext)) { @@ -100,8 +102,7 @@ static int valid(char *ciphertext, struc if (hashlength != 32) return 0; - int saltlength = 0; - char *pos = ciphertext + strlen(mscash2_prefix); + pos = ciphertext + strlen(mscash2_prefix); while (*pos++ != '#') { if (saltlength == 19) return 0; @@ -146,9 +147,9 @@ static void *binary(char *ciphertext) { static uint32_t binary[4]; char *hash = strrchr(ciphertext, '#') + 1; + int i; if (hash == NULL) return binary; - int i; for (i = 0; i < 4; i++) { sscanf(hash + (8 * i), "%08x", &binary[i]); binary[i] = SWAP(binary[i]); diff -urpN magnum-jumbo/src/cuda_phpass_fmt.c magnum-jumbo_cuda_warnings_fix//src/cuda_phpass_fmt.c --- magnum-jumbo/src/cuda_phpass_fmt.c 2012-06-29 16:35:22.766939174 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/cuda_phpass_fmt.c 2012-06-29 16:46:12.793064143 +0000 @@ -112,10 +112,10 @@ static int valid(char *ciphertext, struc ///code from historical JtR phpass patch static void pbinary(char *ciphertext, unsigned char *out) { - memset(out, 0, BINARY_SIZE); int i, bidx = 0; unsigned sixbits; char *pos = &ciphertext[3 + 1 + 8]; + memset(out, 0, BINARY_SIZE); for (i = 0; i < 5; i++) { sixbits = atoi64[ARCH_INDEX(*pos++)]; diff -urpN magnum-jumbo/src/cuda_pwsafe_fmt.c magnum-jumbo_cuda_warnings_fix//src/cuda_pwsafe_fmt.c --- magnum-jumbo/src/cuda_pwsafe_fmt.c 2012-06-29 16:35:22.766939174 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/cuda_pwsafe_fmt.c 2012-06-29 16:56:47.385939631 +0000 @@ -1,7 +1,7 @@ /* Password Safe cracker patch for JtR. Hacked together during May of * 2012 by Dhiru Kholia . * - * CUDA port by Lukas Odzioba + * CUDA port by Lukas Odzioba * * This software is Copyright © 2012, Dhiru Kholia , * and it is hereby released to the general public under the following terms: @@ -61,9 +61,9 @@ static void *get_salt(char *ciphertext) char *keeptr = ctcopy; char *p; int i; - ctcopy += 9; /* skip over "$pwsafe$*" */ pwsafe_salt *salt_struct = mem_alloc_tiny(sizeof(pwsafe_salt), MEM_ALIGN_WORD); + ctcopy += 9; /* skip over "$pwsafe$*" */ p = strtok(ctcopy, "*"); salt_struct->version = atoi(p); p = strtok(NULL, "*"); @@ -81,22 +81,18 @@ static void *get_salt(char *ciphertext) return (void *) salt_struct; } - static void set_salt(void *salt) { memcpy(host_salt, salt, SALT_SIZE); any_cracked = 0; } - - static void crypt_all(int count) { int i; - any_cracked = 0; - unsigned int *src = (unsigned int *) host_salt->hash; unsigned int *dst = (unsigned int *) host_salt->hash; + any_cracked = 0; for (i = 0; i < 8; i++) { dst[i] = SWAP32(src[i]); diff -urpN magnum-jumbo/src/cuda_rawsha256_fmt.c magnum-jumbo_cuda_warnings_fix//src/cuda_rawsha256_fmt.c --- magnum-jumbo/src/cuda_rawsha256_fmt.c 2012-06-29 16:35:22.766939174 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/cuda_rawsha256_fmt.c 2012-06-29 17:04:44.931939762 +0000 @@ -1,6 +1,5 @@ /* -* This software is Copyright (c) 2011 Lukas Odzioba -* +* 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. * This file is shared by cuda-rawsha224 and cuda-rawsha256 formats, @@ -101,8 +100,8 @@ static int valid(char * ciphertext,struc static void *binary(char *ciphertext){ static char realcipher[BINARY_SIZE]; - memset(realcipher,0,BINARY_SIZE); int i; + memset(realcipher,0,BINARY_SIZE); for(i=0;i= 'a')) pos++; return !*pos && pos - ciphertext == CIPHERTEXT_LENGTH; @@ -80,17 +78,16 @@ static int valid(char *ciphertext, struc static void *get_binary(char *ciphertext) { static unsigned char out[FULL_BINARY_SIZE]; - char *p; + char *p = ciphertext; int i; - - p = ciphertext; + uint64_t *b; for (i = 0; i < sizeof(out); i++) { out[i] = (atoi16[ARCH_INDEX(*p)] << 4) | atoi16[ARCH_INDEX(p[1])]; p += 2; } - uint64_t *b = (uint64_t*)out; + b = (uint64_t*)out; for (i = 0; i < 8; i++) { uint64_t t = SWAP64(b[i])-H[i]; b[i] = SWAP64(t); @@ -204,9 +201,9 @@ static int cmp_all(void *binary, int cou static int cmp_one(void *binary, int index) { - uint64_t *b = (uint64_t *) binary; + uint64_t *t,*b = (uint64_t *) binary; cuda_sha512_cpy_hash(ghash); - uint64_t *t = (uint64_t *)ghash; + t = (uint64_t *)ghash; if (b[3] != t[hash_addr(0, index)]) return 0; return 1; @@ -216,14 +213,15 @@ static int cmp_exact(char *source, int i { SHA512_CTX ctx; uint64_t crypt_out[8]; - + int i; + uint64_t *b,*c; + SHA512_Init(&ctx); SHA512_Update(&ctx, gkey[index].v, gkey[index].length); SHA512_Final((unsigned char *)(crypt_out), &ctx); - int i; - uint64_t *b = (uint64_t *)get_binary(source); - uint64_t *c = (uint64_t *)crypt_out; + b = (uint64_t *)get_binary(source); + c = (uint64_t *)crypt_out; for (i = 0; i < 8; i++) { uint64_t t = SWAP64(c[i])-H[i]; diff -urpN magnum-jumbo/src/cuda_xsha512_fmt.c magnum-jumbo_cuda_warnings_fix//src/cuda_xsha512_fmt.c --- magnum-jumbo/src/cuda_xsha512_fmt.c 2012-06-29 16:35:22.766939174 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/cuda_xsha512_fmt.c 2012-06-29 16:58:19.086064427 +0000 @@ -120,6 +120,7 @@ static void *get_binary(char *ciphertext { static unsigned char out[FULL_BINARY_SIZE]; char *p; + uint64_t *b; int i; ciphertext += 6; @@ -129,7 +130,7 @@ static void *get_binary(char *ciphertext (atoi16[ARCH_INDEX(*p)] << 4) | atoi16[ARCH_INDEX(p[1])]; p += 2; } - uint64_t *b = (uint64_t *) out; + b = (uint64_t *) out; for (i = 0; i < 8; i++) { uint64_t t = SWAP64(b[i]) - H[i]; b[i] = SWAP64(t); @@ -289,8 +290,8 @@ static int cmp_all(void *binary, int cou static int cmp_one(void *binary, int index) { uint64_t *b = (uint64_t *) binary; - cuda_xsha512_cpy_hash(ghash); uint64_t *t = (uint64_t *) ghash; + cuda_xsha512_cpy_hash(ghash); if (b[3] != t[hash_addr(0, index)]) return 0; return 1; @@ -301,7 +302,9 @@ static int cmp_exact(char *source, int i { SHA512_CTX ctx; uint64_t crypt_out[8]; - + int i; + uint64_t *b,*c; + SHA512_Init(&ctx); SHA512_Update(&ctx, gsalt.v, SALT_SIZE); if (gkey[index].length > PLAINTEXT_LENGTH) { @@ -312,9 +315,8 @@ static int cmp_exact(char *source, int i SHA512_Update(&ctx, gkey[index].v, gkey[index].length); SHA512_Final((unsigned char *) (crypt_out), &ctx); - int i; - uint64_t *b = (uint64_t *) get_binary(source); - uint64_t *c = (uint64_t *) crypt_out; + b = (uint64_t *) get_binary(source); + c = (uint64_t *) crypt_out; for (i = 0; i < 8; i++) { uint64_t t = SWAP64(c[i]) - H[i]; diff -urpN magnum-jumbo/src/opencl_bf_std.c magnum-jumbo_cuda_warnings_fix//src/opencl_bf_std.c --- magnum-jumbo/src/opencl_bf_std.c 2012-06-29 16:35:22.778939218 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/opencl_bf_std.c 2012-06-29 17:08:04.522939610 +0000 @@ -541,14 +541,13 @@ void opencl_BF_std_crypt(BF_salt *salt, int index=0,j; static unsigned int salt_api[4]; unsigned int rounds=salt->rounds; + static unsigned int BF_out[2*BF_N]; salt_api[0]=salt->salt[0]; salt_api[1]=salt->salt[1]; salt_api[2]=salt->salt[2]; salt_api[3]=salt->salt[3]; - static unsigned int BF_out[2*BF_N]; - exec_bf(salt_api,BF_out,rounds,pltfrmno,devno); for_each_index(){ j=2*index; diff -urpN magnum-jumbo/src/opencl_cryptmd5_fmt.c magnum-jumbo_cuda_warnings_fix//src/opencl_cryptmd5_fmt.c --- magnum-jumbo/src/opencl_cryptmd5_fmt.c 2012-06-29 16:35:22.778939218 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/opencl_cryptmd5_fmt.c 2012-06-29 17:10:47.047939275 +0000 @@ -1,5 +1,5 @@ /* -* This software is Copyright (c) 2011-2012 Lukas Odzioba +* This software is Copyright (c) 2011-2012 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. */ @@ -195,7 +195,7 @@ static void init(struct fmt_main *pFmt) static int valid(char *ciphertext, struct fmt_main *pFmt) { uint8_t i, len = strlen(ciphertext), prefix = 0; - + char *p; if (strncmp(ciphertext, md5_salt_prefix, strlen(md5_salt_prefix)) == 0) prefix |= 1; if (strncmp(ciphertext, apr1_salt_prefix, @@ -204,7 +204,7 @@ static int valid(char *ciphertext, struc if (prefix == 0) return 0; - char *p = strrchr(ciphertext, '$'); + p = strrchr(ciphertext, '$'); for (i = p - ciphertext + 1; i < len; i++) { uint8_t z = ARCH_INDEX(ciphertext[i]); if (ARCH_INDEX(atoi64[z]) == 0x7f) @@ -226,33 +226,35 @@ static void to_binary(char *crypt, char #define _24bit_from_b64(I,B2,B1,B0) \ {\ - unsigned char c1=findb64(crypt[I+0]);\ - unsigned char c2=findb64(crypt[I+1]);\ - unsigned char c3=findb64(crypt[I+2]);\ - unsigned char c4=findb64(crypt[I+3]);\ - unsigned int w=c4<<18|c3<<12|c2<<6|c1;\ - unsigned char b2=w&0xff;w>>=8;\ - unsigned char b1=w&0xff;w>>=8;\ - unsigned char b0=w&0xff;w>>=8;\ + uint8_t c1,c2,c3,c4,b0,b1,b2;\ + uint32_t w;\ + c1=findb64(crypt[I+0]);\ + c2=findb64(crypt[I+1]);\ + c3=findb64(crypt[I+2]);\ + c4=findb64(crypt[I+3]);\ + w=c4<<18|c3<<12|c2<<6|c1;\ + b2=w&0xff;w>>=8;\ + b1=w&0xff;w>>=8;\ + b0=w&0xff;w>>=8;\ alt[B2]=b0;\ alt[B1]=b1;\ alt[B0]=b2;\ } - + uint32_t w; _24bit_from_b64(0, 0, 6, 12); _24bit_from_b64(4, 1, 7, 13); _24bit_from_b64(8, 2, 8, 14); _24bit_from_b64(12, 3, 9, 15); _24bit_from_b64(16, 4, 10, 5); - uint32_t w = findb64(crypt[21]) << 6 | findb64(crypt[20]) << 0; + w = findb64(crypt[21]) << 6 | findb64(crypt[20]) << 0; alt[11] = (w & 0xff); } static void *binary(char *ciphertext) { static char b[BINARY_SIZE]; - memset(b, 0, BINARY_SIZE); char *p = strrchr(ciphertext, '$') + 1; + memset(b, 0, BINARY_SIZE); to_binary(p, b); return (void *) b; } @@ -261,8 +263,8 @@ static void *binary(char *ciphertext) static void *salt(char *ciphertext) { static uint8_t ret[SALT_SIZE]; - memset(ret, 0, SALT_SIZE); uint8_t i, *pos = (uint8_t *) ciphertext, *dest = ret, *end; + memset(ret, 0, SALT_SIZE); if (strncmp(ciphertext, md5_salt_prefix, strlen(md5_salt_prefix)) == 0) { pos += strlen(md5_salt_prefix); @@ -327,6 +329,8 @@ static void set_salt(void *salt) static void crypt_all(int count) { + size_t worksize = KEYS_PER_CRYPT; + size_t localworksize = local_work_size; ///Copy data to GPU memory HANDLE_CLERROR(clEnqueueWriteBuffer (queue[gpu_id], mem_in, CL_FALSE, 0, insize, inbuffer, 0, NULL, @@ -335,8 +339,6 @@ static void crypt_all(int count) 0, saltsize, &host_salt, 0, NULL, NULL), "Copy memsalt"); ///Run kernel - size_t worksize = KEYS_PER_CRYPT; - size_t localworksize = local_work_size; HANDLE_CLERROR(clEnqueueNDRangeKernel (queue[gpu_id], crypt_kernel, 1, NULL, &worksize, &localworksize, 0, NULL, &profilingEvent), "Set ND range"); diff -urpN magnum-jumbo/src/opencl_phpass_fmt.c magnum-jumbo_cuda_warnings_fix//src/opencl_phpass_fmt.c --- magnum-jumbo/src/opencl_phpass_fmt.c 2012-06-29 16:35:22.779939174 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/opencl_phpass_fmt.c 2012-06-29 17:14:33.138108923 +0000 @@ -124,6 +124,7 @@ static char *get_key(int index) static void init(struct fmt_main *pFmt) { + cl_int cl_error; atexit(release_all); opencl_init("$JOHN/phpass_kernel.cl", gpu_id,platform_id); @@ -136,8 +137,6 @@ static void init(struct fmt_main *pFmt) (phpass_hash *) calloc(MAX_KEYS_PER_CRYPT, sizeof(phpass_hash)); assert(inbuffer != NULL); - - cl_int cl_error; mem_in = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, insize, NULL, &cl_error); @@ -198,10 +197,10 @@ static int valid(char *ciphertext, struc static void *binary(char *ciphertext) { static unsigned char b[BINARY_SIZE]; - memset(b, 0, BINARY_SIZE); int i, bidx = 0; unsigned sixbits; char *pos = &ciphertext[3 + 1 + 8]; + memset(b, 0, BINARY_SIZE); for (i = 0; i < 5; i++) { sixbits = atoi64[ARCH_INDEX(*pos++)]; diff -urpN magnum-jumbo/src/opencl_pwsafe_fmt.c magnum-jumbo_cuda_warnings_fix//src/opencl_pwsafe_fmt.c --- magnum-jumbo/src/opencl_pwsafe_fmt.c 2012-06-29 16:35:22.779939174 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/opencl_pwsafe_fmt.c 2012-06-29 17:06:48.900064360 +0000 @@ -123,7 +123,6 @@ static void init(struct fmt_main *pFmt) opencl_find_best_workgroup(pFmt); //local_work_size=256; - atexit(release_all); } @@ -140,9 +139,9 @@ static void *get_salt(char *ciphertext) char *keeptr = ctcopy; char *p; int i; - ctcopy += 9; /* skip over "$pwsafe$*" */ pwsafe_salt *salt_struct = mem_alloc_tiny(sizeof(pwsafe_salt), MEM_ALIGN_WORD); + ctcopy += 9; /* skip over "$pwsafe$*" */ p = strtok(ctcopy, "*"); salt_struct->version = atoi(p); p = strtok(NULL, "*"); @@ -172,10 +171,11 @@ static void set_salt(void *salt) static void crypt_all(int count) { int i; - any_cracked = 0; - + size_t worksize = KEYS_PER_CRYPT; + size_t localworksize = local_work_size; unsigned int *src = (unsigned int *) host_salt->hash; unsigned int *dst = (unsigned int *) host_salt->hash; + any_cracked = 0; for (i = 0; i < 8; i++) { dst[i] = SWAP32(src[i]); @@ -189,8 +189,6 @@ static void crypt_all(int count) 0, saltsize, host_salt, 0, NULL, NULL), "Copy memsalt"); ///Run kernel - size_t worksize = KEYS_PER_CRYPT; - size_t localworksize = local_work_size; HANDLE_CLERROR(clEnqueueNDRangeKernel (queue[gpu_id], crypt_kernel, 1, NULL, &worksize, &localworksize, 0, NULL, &profilingEvent), "Set ND range"); diff -urpN magnum-jumbo/src/opencl_rawsha512_fmt.c magnum-jumbo_cuda_warnings_fix//src/opencl_rawsha512_fmt.c --- magnum-jumbo/src/opencl_rawsha512_fmt.c 2012-06-29 16:35:22.779939174 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/opencl_rawsha512_fmt.c 2012-06-29 17:39:26.331067099 +0000 @@ -147,19 +147,20 @@ static void find_best_workgroup() cl_int ret_code; int i; size_t max_group_size; + char *pass; size_t work_size = KEYS_PER_CRYPT; - HANDLE_CLERROR(clGetKernelWorkGroupInfo(crypt_kernel, devices[gpu_id], - CL_KERNEL_WORK_GROUP_SIZE,sizeof (max_group_size), &max_group_size, - NULL), "Error querying CL_DEVICE_MAX_WORK_GROUP_SIZE"); - cl_command_queue queue_prof = clCreateCommandQueue(context[gpu_id], devices[gpu_id], CL_QUEUE_PROFILING_ENABLE, &ret_code); HANDLE_CLERROR(ret_code, "Error while creating command queue"); + HANDLE_CLERROR(clGetKernelWorkGroupInfo(crypt_kernel, devices[gpu_id], + CL_KERNEL_WORK_GROUP_SIZE,sizeof (max_group_size), &max_group_size, + NULL), "Error querying CL_DEVICE_MAX_WORK_GROUP_SIZE"); + /// Set keys - char *pass = "password"; + pass = "password"; for (i = 0; i < MAX_KEYS_PER_CRYPT; i++) { set_key(pass, i); } @@ -260,7 +261,8 @@ static void *get_binary(char *ciphertext static unsigned char out[FULL_BINARY_SIZE]; char *p; int i; - + uint64_t *b; + p = ciphertext; for (i = 0; i < sizeof(out); i++) { out[i] = @@ -268,7 +270,7 @@ static void *get_binary(char *ciphertext atoi16[ARCH_INDEX(p[1])]; p += 2; } - uint64_t *b = (uint64_t*)out; + b = (uint64_t*)out; for (i = 0; i < 8; i++) { uint64_t t = SWAP64(b[i])-H[i]; b[i] = SWAP64(t); @@ -375,6 +377,8 @@ static int get_hash_6(int index) static void crypt_all(int count) { + size_t worksize = KEYS_PER_CRYPT; + size_t localworksize = local_work_size; ///Copy data to GPU memory if (sha512_key_changed) { HANDLE_CLERROR(clEnqueueWriteBuffer @@ -383,8 +387,6 @@ static void crypt_all(int count) } ///Run kernel - size_t worksize = KEYS_PER_CRYPT; - size_t localworksize = local_work_size; HANDLE_CLERROR(clEnqueueNDRangeKernel (queue[gpu_id], crypt_kernel, 1, NULL, &worksize, &localworksize, 0, NULL, NULL), "Set ND range"); @@ -398,18 +400,18 @@ static void crypt_all(int count) static int cmp_all(void *binary, int count) { + size_t worksize = KEYS_PER_CRYPT; + size_t localworksize = local_work_size; + uint32_t result; ///Copy binary to GPU memory HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_binary, CL_FALSE, 0, sizeof(uint64_t), ((uint64_t*)binary)+3, 0, NULL, NULL), "Copy mem_binary"); ///Run kernel - size_t worksize = KEYS_PER_CRYPT; - size_t localworksize = local_work_size; HANDLE_CLERROR(clEnqueueNDRangeKernel (queue[gpu_id], cmp_kernel, 1, NULL, &worksize, &localworksize, 0, NULL, NULL), "Set ND range"); - uint32_t result; /// Copy result out HANDLE_CLERROR(clEnqueueReadBuffer(queue[gpu_id], mem_cmp, CL_FALSE, 0, sizeof(uint32_t), &result, 0, NULL, NULL), "Copy data back"); @@ -422,12 +424,13 @@ static int cmp_all(void *binary, int cou static int cmp_one(void *binary, int index) { + uint64_t *b = (uint64_t *) binary; + uint64_t *t = (uint64_t *)ghash; + HANDLE_CLERROR(clEnqueueReadBuffer(queue[gpu_id], mem_out, CL_FALSE, 0, outsize, ghash, 0, NULL, NULL), "Copy data back"); HANDLE_CLERROR(clFinish(queue[gpu_id]), "clFinish error"); - uint64_t *b = (uint64_t *) binary; - uint64_t *t = (uint64_t *)ghash; if (b[3] != t[hash_addr(0, index)]) return 0; return 1; @@ -436,15 +439,14 @@ static int cmp_one(void *binary, int ind static int cmp_exact(char *source, int index) { SHA512_CTX ctx; - uint64_t crypt_out[8]; - + uint64_t *b,*c,crypt_out[8]; + int i; SHA512_Init(&ctx); SHA512_Update(&ctx, gkey[index].v, gkey[index].length); SHA512_Final((unsigned char *)(crypt_out), &ctx); - int i; - uint64_t *b = (uint64_t *)get_binary(source); - uint64_t *c = (uint64_t *)crypt_out; + b = (uint64_t *)get_binary(source); + c = (uint64_t *)crypt_out; for (i = 0; i < 8; i++) { uint64_t t = SWAP64(c[i])-H[i]; diff -urpN magnum-jumbo/src/opencl_wpapsk_fmt.c magnum-jumbo_cuda_warnings_fix//src/opencl_wpapsk_fmt.c --- magnum-jumbo/src/opencl_wpapsk_fmt.c 2012-06-29 16:35:22.779939174 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/opencl_wpapsk_fmt.c 2012-06-29 17:15:23.392938900 +0000 @@ -61,6 +61,7 @@ static void release_all(void) static void init(struct fmt_main *pFmt) { + cl_int cl_error; assert(sizeof(hccap_t) == HCCAP_SIZE); inbuffer = @@ -83,7 +84,6 @@ static void init(struct fmt_main *pFmt) //listOpenCLdevices(); opencl_init("$JOHN/wpapsk_kernel.cl", gpu_id, platform_id); /// Alocate memory - cl_int cl_error; mem_in = clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, insize, NULL, &cl_error); diff -urpN magnum-jumbo/src/opencl_xsha512_fmt.c magnum-jumbo_cuda_warnings_fix//src/opencl_xsha512_fmt.c --- magnum-jumbo/src/opencl_xsha512_fmt.c 2012-06-29 16:35:22.779939174 +0000 +++ magnum-jumbo_cuda_warnings_fix//src/opencl_xsha512_fmt.c 2012-06-29 17:43:15.977939462 +0000 @@ -151,21 +151,20 @@ static void find_best_workgroup() size_t my_work_group = 1; cl_int ret_code; int i; - size_t max_group_size; - size_t work_size = KEYS_PER_CRYPT; - HANDLE_CLERROR(clGetKernelWorkGroupInfo(crypt_kernel, devices[gpu_id], - CL_KERNEL_WORK_GROUP_SIZE, sizeof(max_group_size), - &max_group_size, NULL), - "Error querying CL_DEVICE_MAX_WORK_GROUP_SIZE"); - + char * pass; + size_t max_group_size, work_size = KEYS_PER_CRYPT; cl_command_queue queue_prof = clCreateCommandQueue(context[gpu_id], devices[gpu_id], CL_QUEUE_PROFILING_ENABLE, &ret_code); HANDLE_CLERROR(ret_code, "Error while creating command queue"); - + + HANDLE_CLERROR(clGetKernelWorkGroupInfo(crypt_kernel, devices[gpu_id], + CL_KERNEL_WORK_GROUP_SIZE, sizeof(max_group_size), + &max_group_size, NULL), + "Error querying CL_DEVICE_MAX_WORK_GROUP_SIZE"); /// Set keys - char *pass = "password"; + pass = "password"; for (i = 0; i < MAX_KEYS_PER_CRYPT; i++) { set_key(pass, i); } @@ -309,7 +308,7 @@ static void *get_binary(char *ciphertext static unsigned char out[FULL_BINARY_SIZE]; char *p; int i; - + uint64_t *b; ciphertext += 6; p = ciphertext + 8; for (i = 0; i < sizeof(out); i++) { @@ -317,7 +316,7 @@ static void *get_binary(char *ciphertext (atoi16[ARCH_INDEX(*p)] << 4) | atoi16[ARCH_INDEX(p[1])]; p += 2; } - uint64_t *b = (uint64_t *) out; + b = (uint64_t *) out; for (i = 0; i < 8; i++) { uint64_t t = SWAP64(b[i]) - H[i]; b[i] = SWAP64(t); @@ -454,6 +453,8 @@ static void set_salt(void *salt) static void crypt_all(int count) { + size_t worksize = KEYS_PER_CRYPT; + size_t localworksize = local_work_size; ///Copy data to GPU memory if (xsha512_key_changed) { HANDLE_CLERROR(clEnqueueWriteBuffer @@ -464,8 +465,6 @@ static void crypt_all(int count) 0, SALT_SIZE, &gsalt, 0, NULL, NULL), "Copy memsalt"); ///Run kernel - size_t worksize = KEYS_PER_CRYPT; - size_t localworksize = local_work_size; HANDLE_CLERROR(clEnqueueNDRangeKernel (queue[gpu_id], crypt_kernel, 1, NULL, &worksize, &localworksize, 0, NULL, NULL), "Set ND range"); @@ -479,19 +478,20 @@ static void crypt_all(int count) static int cmp_all(void *binary, int count) { + size_t worksize = KEYS_PER_CRYPT; + size_t localworksize = local_work_size; + uint32_t result; + ///Copy binary to GPU memory HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_binary, CL_FALSE, 0, sizeof(uint64_t), ((uint64_t *) binary) + 3, 0, NULL, NULL), "Copy mem_binary"); ///Run kernel - size_t worksize = KEYS_PER_CRYPT; - size_t localworksize = local_work_size; HANDLE_CLERROR(clEnqueueNDRangeKernel (queue[gpu_id], cmp_kernel, 1, NULL, &worksize, &localworksize, 0, NULL, NULL), "Set ND range"); - uint32_t result; /// Copy result out HANDLE_CLERROR(clEnqueueReadBuffer(queue[gpu_id], mem_cmp, CL_FALSE, 0, sizeof(uint32_t), &result, 0, NULL, NULL), "Copy data back"); @@ -504,38 +504,36 @@ static int cmp_all(void *binary, int cou static int cmp_one(void *binary, int index) { + uint64_t *b = (uint64_t *) binary; + uint64_t *t = (uint64_t *) ghash; HANDLE_CLERROR(clEnqueueReadBuffer(queue[gpu_id], mem_out, CL_FALSE, 0, outsize, ghash, 0, NULL, NULL), "Copy data back"); HANDLE_CLERROR(clFinish(queue[gpu_id]), "clFinish error"); - uint64_t *b = (uint64_t *) binary; - uint64_t *t = (uint64_t *) ghash; if (b[3] != t[hash_addr(0, index)]) return 0; return 1; - } static int cmp_exact(char *source, int index) { SHA512_CTX ctx; - uint64_t crypt_out[8]; + int i; + uint64_t *b,*c,crypt_out[8]; SHA512_Init(&ctx); SHA512_Update(&ctx, gsalt.v, SALT_SIZE); SHA512_Update(&ctx, gkey[index].v, gkey[index].length); SHA512_Final((unsigned char *) (crypt_out), &ctx); - int i; - uint64_t *b = (uint64_t *) get_binary(source); - uint64_t *c = (uint64_t *) crypt_out; + b = (uint64_t *) get_binary(source); + c = (uint64_t *) crypt_out; for (i = 0; i < 8; i++) { uint64_t t = SWAP64(c[i]) - H[i]; c[i] = SWAP64(t); } - for (i = 0; i < FULL_BINARY_SIZE / 8; i++) { //examin 512bits if (b[i] != c[i]) return 0;