--- /dev/null
+// This file is part of BOINC.
+// http://boinc.berkeley.edu
+// Copyright (C) 2008 University of California
+//
+// BOINC is free software; you can redistribute it and/or modify it
+// under the terms of the GNU Lesser General Public License
+// as published by the Free Software Foundation,
+// either version 3 of the License, or (at your option) any later version.
+//
+// BOINC is distributed in the hope that it will be useful,
+// but WITHOUT ANY WARRANTY; without even the implied warranty of
+// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+// See the GNU Lesser General Public License for more details.
+//
+// You should have received a copy of the GNU Lesser General Public License
+// along with BOINC. If not, see <http://www.gnu.org/licenses/>.
+
+// This program serves as both
+// - An example BOINC application, illustrating the use of the BOINC API
+// - A program for testing various features of BOINC
+//
+// NOTE: this file exists as both
+// boinc/apps/upper_case.C
+// and
+// boinc_samples/example_app/uc2.C
+// If you update one, please update the other!
+
+// The program converts a mixed-case file to upper case:
+// read "in", convert to upper case, write to "out"
+//
+// command line options
+// -run_slow: sleep 1 second after each character
+// -cpu_time N: use about N CPU seconds after copying files
+// -early_exit: exit(10) after 30 chars
+// -early_crash: crash after 30 chars
+//
+
+#ifdef _WIN32
+#include "boinc_win.h"
+#else
+#include "config.h"
+#include <cstdio>
+#include <cctype>
+#include <ctime>
+#include <cstring>
+#include <cstdlib>
+#include <csignal>
+#include <unistd.h>
+#endif
+
+#include <string>
+#include <fstream>
+#include <iostream>
+#include "str_util.h"
+#include "util.h"
+#include "filesys.h"
+#include "boinc_api.h"
+#include "Public.h"
+// Rainbowcrack code
+#include "ChainWalkContext.h"
+//typedef unsigned int uint32;
+//typedef unsigned __int64 uint64;
+#include "rcuda.h"
+#include "rcuda_ext.h"
+
+
+using std::string;
+
+/*
+bool early_exit = false;
+bool early_crash = false;
+bool early_sleep = false;
+double cpu_time = 20, comp_result;
+*/
+int QuickSortPartition(RainbowChainCP* pChain, int nLow, int nHigh)
+{
+ int nRandomIndex = nLow + ((unsigned int)rand() * (RAND_MAX + 1) + (unsigned int)rand()) % (nHigh - nLow + 1);
+ RainbowChainCP TempChain;
+ TempChain = pChain[nLow];
+ pChain[nLow] = pChain[nRandomIndex];
+ pChain[nRandomIndex] = TempChain;
+
+ TempChain = pChain[nLow];
+ uint64 nPivotKey = pChain[nLow].nIndexE;
+ while (nLow < nHigh)
+ {
+ while (nLow < nHigh && pChain[nHigh].nIndexE >= nPivotKey)
+ nHigh--;
+ pChain[nLow] = pChain[nHigh];
+ while (nLow < nHigh && pChain[nLow].nIndexE <= nPivotKey)
+ nLow++;
+ pChain[nHigh] = pChain[nLow];
+ }
+ pChain[nLow] = TempChain;
+ return nLow;
+}
+
+void QuickSort(RainbowChainCP* pChain, int nLow, int nHigh)
+{
+ if (nLow < nHigh)
+ {
+ int nPivotLoc = QuickSortPartition(pChain, nLow, nHigh);
+ QuickSort(pChain, nLow, nPivotLoc - 1);
+ QuickSort(pChain, nPivotLoc + 1, nHigh);
+ }
+}
+
+int main(int argc, char **argv) {
+ int retval;
+ double fd;
+ char output_path[512], chkpt_path[512];
+ FILE* state;
+ retval = boinc_init();
+ if (retval) {
+ fprintf(stderr, "boinc_init returned %d\n", retval);
+ exit(retval);
+ }
+
+
+ // get size of input file (used to compute fraction done)
+ //
+ //file_size(input_path, fsize);
+
+ // See if there's a valid checkpoint file.
+ // If so seek input file and truncate output file
+ //
+
+
+ if(argc < 10)
+ {
+ fprintf(stderr, "Not enough parameters");
+ return -1;
+ }
+ string sHashRoutineName, sCharsetName, sSalt, sCheckPoints;
+ uint32 nRainbowChainCount, nPlainLenMin, nPlainLenMax, nRainbowTableIndex, nRainbowChainLen;
+ uint64 nChainStart;
+ sHashRoutineName = argv[1];
+ sCharsetName = argv[2];
+ nPlainLenMin = atoi(argv[3]);
+ nPlainLenMax = atoi(argv[4]);
+ nRainbowTableIndex = atoi(argv[5]);
+ nRainbowChainLen = atoi(argv[6]);
+ nRainbowChainCount = atoi(argv[7]);
+#ifdef _WIN32
+
+ nChainStart = _atoi64(argv[8]);
+
+#else
+ nChainStart = atoll(argv[8]);
+#endif
+ sCheckPoints = argv[9];
+ vector<int> vCPPositions;
+ char *cp = strtok((char *)sCheckPoints.c_str(), ",");
+ while(cp != NULL)
+ {
+ vCPPositions.push_back(atoi(cp));
+ cp = strtok(NULL, ",");
+ }
+ if(argc == 11)
+ {
+ sSalt = argv[10];
+ }
+ //std::cout << "Starting ChainGenerator" << std::endl;
+ // Setup CChainWalkContext
+ //std::cout << "ChainGenerator started." << std::endl;
+
+ if (!CChainWalkContext::SetHashRoutine(sHashRoutineName))
+ {
+ fprintf(stderr, "hash routine %s not supported\n", sHashRoutineName.c_str());
+ return 1;
+ }
+ //std::cout << "Hash routine validated" << std::endl;
+
+ if (!CChainWalkContext::SetPlainCharset(sCharsetName, nPlainLenMin, nPlainLenMax))
+ {
+ std::cerr << "charset " << sCharsetName << " not supported" << std::endl;
+ return 2;
+ }
+ //std::cout << "Plain charset validated" << std::endl;
+
+ if (!CChainWalkContext::SetRainbowTableIndex(nRainbowTableIndex))
+ {
+ std::cerr << "invalid rainbow table index " << nRainbowTableIndex << std::endl;
+ return 3;
+ }
+ //std::cout << "Rainbowtable index validated" << std::endl;
+
+ if(sHashRoutineName == "mscache")// || sHashRoutineName == "lmchall" || sHashRoutineName == "halflmchall")
+ {
+ // Convert username to unicode
+ const char *szSalt = sSalt.c_str();
+ int salt_length = strlen(szSalt);
+ unsigned char cur_salt[256];
+ for (int i=0; i<salt_length; i++)
+ {
+ cur_salt[i*2] = szSalt[i];
+ cur_salt[i*2+1] = 0x00;
+ }
+ CChainWalkContext::SetSalt(cur_salt, salt_length*2);
+ }
+ else if(sHashRoutineName == "halflmchall")
+ { // The salt is hardcoded into the hash routine
+ // CChainWalkContext::SetSalt((unsigned char*)&salt, 8);
+ }
+ else if(sHashRoutineName == "oracle")
+ {
+ CChainWalkContext::SetSalt((unsigned char *)sSalt.c_str(), sSalt.length());
+ }
+ //std::cout << "Opening chain file" << std::endl;
+
+
+ // Open file
+// fclose(fopen(sFilename.c_str(), "a"));
+// FILE* file = fopen(sFilename.c_str(), "r+b");
+ boinc_resolve_filename("result", output_path, sizeof(output_path));
+ fclose(boinc_fopen(output_path, "a"));
+ FILE *outfile = boinc_fopen(output_path, "r+b");
+
+ if (outfile == NULL)
+ {
+ std::cerr << "failed to create " << output_path << std::endl;
+ return 4;
+ }
+
+
+ // Check existing chains
+ unsigned int nDataLen = (unsigned int)GetFileLen(outfile);
+ unsigned int nFileLen;
+
+ // Round to boundary
+ nDataLen = nDataLen / 18 * 18;
+ if ((int)nDataLen == nRainbowChainCount * 18)
+ {
+ std::cerr << "precomputation of this rainbow table already finished" << std::endl;
+ fclose(outfile);
+ return 0;
+ }
+
+ fseek(outfile, nDataLen, SEEK_SET);
+ size_t nReturn;
+ CChainWalkContext cwc;
+ uint64 nIndex[2];
+
+// std::cout << "Starting to generate chains" << std::endl;
+ int maxCalcBuffSize = rcuda::GetChainsBufferSize(5000);
+ uint64 *calcBuff = new uint64[2*maxCalcBuffSize];
+ int ii;
+
+ CudaCWCExtender ex(&cwc);
+ rcuda::RCudaTask cuTask;
+ std::vector<unsigned char> stPlain;
+ ex.Init();
+time_t tStart, tStartFinal, tEndFinal;
+time_t tEnd;
+ tStartFinal = time(NULL);
+ for(int nCurrentCalculatedChains = nDataLen / 18, calcSize; nCurrentCalculatedChains < nRainbowChainCount; )
+ {
+ fd = (double)nCurrentCalculatedChains / (double)nRainbowChainCount;
+ boinc_fraction_done(fd);
+
+ cuTask.hash = ex.GetHash();
+ cuTask.startIdx = nChainStart + nCurrentCalculatedChains;
+ cuTask.idxCount = std::min<int>(nRainbowChainCount - nCurrentCalculatedChains, maxCalcBuffSize);
+ cuTask.stPlainSize = ex.IndexToStartPlain(0, stPlain);
+ cuTask.stPlain = &stPlain[0];
+ cuTask.dimVec = ex.GetPlainDimVec();
+ cuTask.dimVecSize = ex.GetPlainDimVecSize()/3;
+ cuTask.charSet = ex.GetCharSet();
+ cuTask.charSetSize = ex.GetCharSetSize();
+ cuTask.cpPositions = &vCPPositions[0];
+ cuTask.cpPosSize = vCPPositions.size();
+ cuTask.reduceOffset = ex.GetReduceOffset();
+ cuTask.plainSpaceTotal = ex.GetPlainSpaceTotal();
+ cuTask.rainbowChainLen = nRainbowChainLen;
+ for(int ii = 0; ii < cuTask.idxCount; ii++) {
+ calcBuff[2*ii] = cuTask.startIdx + ii;
+ calcBuff[2*ii+1] = 0;
+ }
+
+ tStart = time(NULL);
+
+ calcSize = rcuda::CalcChainsOnCUDA(&cuTask, calcBuff);
+ tEnd = time(NULL);
+ std::cerr << "CUDA time taken: " << tEnd - tStart << std::endl;
+ tStart = time(NULL);
+ if(calcSize > 0) {
+ nCurrentCalculatedChains += calcSize;
+ for(ii = 0; ii < cuTask.idxCount; ii++) {
+ nIndex[0] = cuTask.startIdx + ii;
+ nReturn = fwrite(nIndex, 1, 8, outfile);
+ nReturn += fwrite(calcBuff+(2*ii), 1, 8, outfile);
+ nReturn += fwrite(calcBuff+(2*ii+1), 1, 2, outfile);
+ if(nReturn != 18) {
+ std::cerr << "disk write fail" << std::endl;
+ fclose(outfile);
+ return 9;
+ }
+ }
+ } else {
+ std::cerr << "Calculations on CUDA failed!" << std::endl;
+ fclose(outfile);
+ return 0;
+ }
+ }
+ tEndFinal = time(NULL);
+ std::cerr << "Time taken: " << tEndFinal - tStartFinal << " secs" << std::endl;
+ delete [] calcBuff;
+#ifdef _DEBUG
+ std::cout << "Generation completed" << std::endl;
+#endif
+ fseek(outfile, 0, SEEK_SET);
+ nFileLen = GetFileLen(outfile);
+ nRainbowChainCount = nFileLen / 18;
+
+ RainbowChainCP* pChain = (RainbowChainCP*)new unsigned char[sizeof(RainbowChainCP) * nRainbowChainCount];
+
+ if (pChain != NULL)
+ {
+ // Load file
+#ifdef _DEBUG
+ std::cout << "Sorting file" << std::endl;
+#endif
+ fseek(outfile, 0, SEEK_SET);
+ for(int i = 0; i < nRainbowChainCount; i++)
+ {
+ if(fread(&pChain[i], 1, 16, outfile) != 16)
+ {
+ printf("disk read fail\n");
+ return 9;
+ }
+ if(fread(&pChain[i].nCheckPoint, 1, sizeof(pChain[i].nCheckPoint), outfile) != 2)
+ {
+ printf("disk read fail\n");
+ return 9;
+ }
+ }
+
+ // Sort file
+ QuickSort(pChain, 0, nRainbowChainCount - 1);
+
+ // Write file
+ fseek(outfile, 0, SEEK_SET);
+ for(int i = 0; i < nRainbowChainCount; i++)
+ {
+ fwrite(&pChain[i], 1, 16, outfile);
+ fwrite(&pChain[i].nCheckPoint, 2, 1, outfile);
+ }
+ delete[] pChain;
+ }
+
+ fclose(outfile);
+
+ // main loop - read characters, convert to UC, write
+ //
+
+ boinc_fraction_done(1);
+ boinc_finish(0);
+}
+
+#ifdef _WIN32
+int WINAPI WinMain(HINSTANCE hInst, HINSTANCE hPrevInst, LPSTR Args, int WinMode) {
+ LPSTR command_line;
+ char* argv[100];
+ int argc;
+
+ command_line = GetCommandLine();
+ argc = parse_command_line( command_line, argv );
+ return main(argc, argv);
+}
+#endif
+
+const char *BOINC_RCSID_33ac47a071 = "$Id: upper_case.C 12135 2007-02-21 20:04:14Z davea $";
+
--- /dev/null
+
+Microsoft Visual Studio Solution File, Format Version 10.00
+# Visual Studio 2008
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "distrrtgen", "distrrtgen.vcproj", "{A3BDF5F8-4D0A-4B27-B1D9-7E77CBDA86C7}"
+EndProject
+Global
+ GlobalSection(SolutionConfigurationPlatforms) = preSolution
+ Debug|Win32 = Debug|Win32
+ Release|Win32 = Release|Win32
+ EndGlobalSection
+ GlobalSection(ProjectConfigurationPlatforms) = postSolution
+ {A3BDF5F8-4D0A-4B27-B1D9-7E77CBDA86C7}.Debug|Win32.ActiveCfg = Debug|Win32
+ {A3BDF5F8-4D0A-4B27-B1D9-7E77CBDA86C7}.Debug|Win32.Build.0 = Debug|Win32
+ {A3BDF5F8-4D0A-4B27-B1D9-7E77CBDA86C7}.Release|Win32.ActiveCfg = Release|Win32
+ {A3BDF5F8-4D0A-4B27-B1D9-7E77CBDA86C7}.Release|Win32.Build.0 = Release|Win32
+ EndGlobalSection
+ GlobalSection(SolutionProperties) = preSolution
+ HideSolutionNode = FALSE
+ EndGlobalSection
+EndGlobal
--- /dev/null
+<?xml version="1.0" encoding="Windows-1252"?>
+<VisualStudioProject
+ ProjectType="Visual C++"
+ Version="9,00"
+ Name="distrrtgen"
+ ProjectGUID="{A3BDF5F8-4D0A-4B27-B1D9-7E77CBDA86C7}"
+ RootNamespace="distrrtgen"
+ Keyword="Win32Proj"
+ TargetFrameworkVersion="196613"
+ >
+ <Platforms>
+ <Platform
+ Name="Win32"
+ />
+ </Platforms>
+ <ToolFiles>
+ </ToolFiles>
+ <Configurations>
+ <Configuration
+ Name="Debug|Win32"
+ OutputDirectory="$(SolutionDir)$(ConfigurationName)"
+ IntermediateDirectory="$(ConfigurationName)"
+ ConfigurationType="1"
+ CharacterSet="0"
+ >
+ <Tool
+ Name="VCPreBuildEventTool"
+ />
+ <Tool
+ Name="VCCustomBuildTool"
+ CommandLine=""
+ />
+ <Tool
+ Name="VCXMLDataGeneratorTool"
+ />
+ <Tool
+ Name="VCWebServiceProxyGeneratorTool"
+ />
+ <Tool
+ Name="VCMIDLTool"
+ />
+ <Tool
+ Name="VCCLCompilerTool"
+ Optimization="0"
+ AdditionalIncludeDirectories="/boinc/lib;/boinc/api;"/boinc/zip";"/boinc/client/win";"/boinc/";../../rt api"
+ PreprocessorDefinitions="WIN32;_DEBUG;_MT;_CONSOLE"
+ MinimalRebuild="true"
+ BasicRuntimeChecks="3"
+ RuntimeLibrary="1"
+ UsePrecompiledHeader="0"
+ WarningLevel="3"
+ DebugInformationFormat="4"
+ />
+ <Tool
+ Name="VCManagedResourceCompilerTool"
+ />
+ <Tool
+ Name="VCResourceCompilerTool"
+ />
+ <Tool
+ Name="VCPreLinkEventTool"
+ />
+ <Tool
+ Name="VCLinkerTool"
+ AdditionalDependencies="libcmtd.lib libcpmtd.lib kernel32.lib user32.lib ole32.lib delayimp.lib ..\lib_debug\libboinc_staticcrt.lib ..\lib_debug\libboincapi_staticcrt.lib c:/openssl/lib/libeay32.lib cuda.lib cudart.lib"
+ LinkIncremental="0"
+ AdditionalLibraryDirectories="$(CUDA_LIB_PATH)\..\Win32"
+ IgnoreAllDefaultLibraries="true"
+ GenerateDebugInformation="true"
+ SubSystem="1"
+ TargetMachine="1"
+ />
+ <Tool
+ Name="VCALinkTool"
+ />
+ <Tool
+ Name="VCManifestTool"
+ />
+ <Tool
+ Name="VCXDCMakeTool"
+ />
+ <Tool
+ Name="VCBscMakeTool"
+ />
+ <Tool
+ Name="VCFxCopTool"
+ />
+ <Tool
+ Name="VCAppVerifierTool"
+ />
+ <Tool
+ Name="VCPostBuildEventTool"
+ />
+ </Configuration>
+ <Configuration
+ Name="Release|Win32"
+ OutputDirectory="$(SolutionDir)$(ConfigurationName)"
+ IntermediateDirectory="$(ConfigurationName)"
+ ConfigurationType="1"
+ CharacterSet="0"
+ WholeProgramOptimization="1"
+ >
+ <Tool
+ Name="VCPreBuildEventTool"
+ />
+ <Tool
+ Name="VCCustomBuildTool"
+ />
+ <Tool
+ Name="VCXMLDataGeneratorTool"
+ />
+ <Tool
+ Name="VCWebServiceProxyGeneratorTool"
+ />
+ <Tool
+ Name="VCMIDLTool"
+ />
+ <Tool
+ Name="VCCLCompilerTool"
+ Optimization="2"
+ InlineFunctionExpansion="2"
+ EnableIntrinsicFunctions="true"
+ FavorSizeOrSpeed="1"
+ WholeProgramOptimization="true"
+ AdditionalIncludeDirectories="/boinc/lib;/boinc/api;/boinc/zip;/boinc/client/win;/boinc/;"C:\Users\mwj\git\Common\rt api""
+ PreprocessorDefinitions="WIN32;NDEBUG;_MT;_CONSOLE"
+ RuntimeLibrary="0"
+ EnableFunctionLevelLinking="true"
+ EnableEnhancedInstructionSet="2"
+ UsePrecompiledHeader="0"
+ WarningLevel="3"
+ Detect64BitPortabilityProblems="true"
+ DebugInformationFormat="3"
+ />
+ <Tool
+ Name="VCManagedResourceCompilerTool"
+ />
+ <Tool
+ Name="VCResourceCompilerTool"
+ />
+ <Tool
+ Name="VCPreLinkEventTool"
+ />
+ <Tool
+ Name="VCLinkerTool"
+ AdditionalDependencies="odbc32.lib odbccp32.lib libcmt.lib libcpmt.lib C:\lib\boinc\Win32\Release\libboinc_staticcrt.lib C:\lib\boinc\Win32\Release\libboincapi_staticcrt.lib cuda.lib cudart.lib"
+ LinkIncremental="0"
+ AdditionalLibraryDirectories="$(CUDA_LIB_PATH)\..\Win32"
+ IgnoreAllDefaultLibraries="true"
+ GenerateDebugInformation="true"
+ SubSystem="2"
+ OptimizeReferences="2"
+ EnableCOMDATFolding="2"
+ RandomizedBaseAddress="1"
+ DataExecutionPrevention="0"
+ TargetMachine="1"
+ />
+ <Tool
+ Name="VCALinkTool"
+ />
+ <Tool
+ Name="VCManifestTool"
+ />
+ <Tool
+ Name="VCXDCMakeTool"
+ />
+ <Tool
+ Name="VCBscMakeTool"
+ />
+ <Tool
+ Name="VCFxCopTool"
+ />
+ <Tool
+ Name="VCAppVerifierTool"
+ />
+ <Tool
+ Name="VCPostBuildEventTool"
+ />
+ </Configuration>
+ </Configurations>
+ <References>
+ </References>
+ <Files>
+ <Filter
+ Name="Source Files"
+ Filter="cpp;c;cc;cxx;def;odl;idl;hpj;bat;asm;asmx"
+ UniqueIdentifier="{4FC737F1-C7A5-4376-A066-2A32D752A2FF}"
+ >
+ <File
+ RelativePath="..\..\..\Common\rt api\ChainWalkContext.cpp"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\des_enc.c"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\des_setkey.c"
+ >
+ </File>
+ <File
+ RelativePath=".\distrrtgen.cpp"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\ecb_enc.c"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\HashAlgorithm.cpp"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\HashRoutine.cpp"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\md4.cpp"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\md5.cpp"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\Public.cpp"
+ >
+ </File>
+ <File
+ RelativePath=".\rcuda.cu"
+ >
+ <FileConfiguration
+ Name="Debug|Win32"
+ >
+ <Tool
+ Name="VCCustomBuildTool"
+ CommandLine="\93$(CUDA_BIN_PATH)\nvcc.exe\94 --compile -m32 --output-file "$(OutDir)\$(InputName).obj" "$(InputFileName)"
"
+ Outputs="$(OutDir)\$(InputName).obj"
+ />
+ </FileConfiguration>
+ <FileConfiguration
+ Name="Release|Win32"
+ >
+ <Tool
+ Name="VCCustomBuildTool"
+ CommandLine="\93$(CUDA_BIN_PATH)\nvcc.exe\94 --compile -m32 --output-file "$(OutDir)\$(InputName).obj" "$(InputFileName)"
"
+ Outputs="$(OutDir)\$(InputName).obj"
+ />
+ </FileConfiguration>
+ </File>
+ <File
+ RelativePath=".\rcuda_ext.cpp"
+ >
+ </File>
+ </Filter>
+ <Filter
+ Name="Header Files"
+ Filter="h;hpp;hxx;hm;inl;inc;xsd"
+ UniqueIdentifier="{93995380-89BD-4b04-88EB-625FBE52EBFB}"
+ >
+ <File
+ RelativePath="..\..\..\Common\rt api\ChainWalkContext.h"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\des.h"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\des_locl.h"
+ >
+ </File>
+ <File
+ RelativePath=".\distrrtgen.h"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\HashAlgorithm.h"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\HashRoutine.h"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\md4.h"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\md5.h"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\podd.h"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\Public.h"
+ >
+ </File>
+ <File
+ RelativePath=".\rcuda.h"
+ >
+ </File>
+ <File
+ RelativePath=".\rcuda_ext.h"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\sk.h"
+ >
+ </File>
+ <File
+ RelativePath="..\..\..\Common\rt api\spr.h"
+ >
+ </File>
+ </Filter>
+ <Filter
+ Name="Resource Files"
+ Filter="rc;ico;cur;bmp;dlg;rc2;rct;bin;rgs;gif;jpg;jpeg;jpe;resx;tiff;tif;png;wav"
+ UniqueIdentifier="{67DA6AB6-F800-4c08-8B7A-83BB121AAD01}"
+ >
+ </Filter>
+ </Files>
+ <Globals>
+ </Globals>
+</VisualStudioProject>
--- /dev/null
+#include <stdio.h>
+#include <cuda.h>
+#include "rcuda.h"
+
+#define GRID_X_L2 6
+#define GRID_Y_L2 6
+#define BLOCK_X_L2 7
+#define GRID_X_SIZE (1<<GRID_X_L2)
+#define GRID_Y_SIZE (1<<GRID_Y_L2)
+#define BLOCK_X_SIZE (1<<BLOCK_X_L2)
+#define KERN_CHAIN_SIZE 100
+#define SHIDX(x) ((x)<<4)
+
+
+__device__ __constant__ unsigned int h[4] = { 0x67452301, 0xEFCDAB89, 0x98BADCFE, 0x10325476 };
+__device__ __constant__ unsigned char r[64] = { \
+ 7, 12, 17, 22, 7, 12, 17, 22, 7, 12, 17, 22, 7, 12, 17, 22, \
+ 5, 9, 14, 20, 5, 9, 14, 20, 5, 9, 14, 20, 5, 9, 14, 20, \
+ 4, 11, 16, 23, 4, 11, 16, 23, 4, 11, 16, 23, 4, 11, 16, 23, \
+ 6, 10, 15, 21, 6, 10, 15, 21, 6, 10, 15, 21, 6, 10, 15, 21 };
+__device__ __constant__ unsigned char g[64] = { \
+ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, \
+ 1, 6, 11, 0, 5, 10, 15, 4, 9, 14, 3, 8, 13, 2, 7, 12, \
+ 5, 8, 11, 14, 1, 4, 7, 10, 13, 0, 3, 6, 9, 12, 15, 2, \
+ 0, 7, 14, 5, 12, 3, 10, 1, 8, 15, 6, 13, 4, 11, 2, 9 };
+__device__ __constant__ unsigned int ac[64] = { \
+ 0xd76aa478, 0xe8c7b756, 0x242070db, 0xc1bdceee, \
+ 0xf57c0faf, 0x4787c62a, 0xa8304613, 0xfd469501, \
+ 0x698098d8, 0x8b44f7af, 0xffff5bb1, 0x895cd7be, \
+ 0x6b901122, 0xfd987193, 0xa679438e, 0x49b40821, \
+ \
+ 0xf61e2562, 0xc040b340, 0x265e5a51, 0xe9b6c7aa, \
+ 0xd62f105d, 0x02441453, 0xd8a1e681, 0xe7d3fbc8, \
+ 0x21e1cde6, 0xc33707d6, 0xf4d50d87, 0x455a14ed, \
+ 0xa9e3e905, 0xfcefa3f8, 0x676f02d9, 0x8d2a4c8a, \
+ \
+ 0xfffa3942, 0x8771f681, 0x6d9d6122, 0xfde5380c, \
+ 0xa4beea44, 0x4bdecfa9, 0xf6bb4b60, 0xbebfbc70, \
+ 0x289b7ec6, 0xeaa127fa, 0xd4ef3085, 0x04881d05, \
+ 0xd9d4d039, 0xe6db99e5, 0x1fa27cf8, 0xc4ac5665, \
+ \
+ 0xf4292244, 0x432aff97, 0xab9423a7, 0xfc93a039, \
+ 0x655b59c3, 0x8f0ccc92, 0xffeff47d, 0x85845dd1, \
+ 0x6fa87e4f, 0xfe2ce6e0, 0xa3014314, 0x4e0811a1, \
+ 0xf7537e82, 0xbd3af235, 0x2ad7d2bb, 0xeb86d391 };
+
+__device__ __constant__ unsigned int testData[16] = { \
+ 0x79706d63, 0x6d627667, 0x00000080, 0x00000000, \
+ 0x00000000, 0x00000000, 0x00000000, 0x00000000, \
+ 0x00000000, 0x00000000, 0x00000000, 0x00000000, \
+ 0x00000000, 0x00000000, 0x00000040, 0x00000000 };
+
+__device__ unsigned int FF(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) {
+ unsigned int ret;
+ ret = a + ((b&c)|((~b)&d)) + ac[i] + data[SHIDX(g[i])];
+ ret = (ret<<r[i])|(ret>>(32-r[i]));
+ ret += b;
+ return ret;
+}
+
+__device__ unsigned int GG(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) {
+ unsigned int ret;
+ ret = a + ((b&d)|(c&(~d))) + ac[i] + data[SHIDX(g[i])];
+ ret = (ret<<r[i])|(ret>>(32-r[i]));
+ ret += b;
+ return ret;
+}
+
+__device__ unsigned int HH(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) {
+ unsigned int ret;
+ ret = a + (b^c^d) + ac[i] + data[SHIDX(g[i])];
+ ret = (ret<<r[i])|(ret>>(32-r[i]));
+ ret += b;
+ return ret;
+}
+
+__device__ unsigned int II(unsigned int a, unsigned int b, unsigned int c, unsigned int d, int i, const unsigned int* data) {
+ unsigned int ret;
+ ret = a + (c^(b|(~d))) + ac[i] + data[SHIDX(g[i])];
+ ret = (ret<<r[i])|(ret>>(32-r[i]));
+ ret += b;
+ return ret;
+}
+
+__device__ void MD5(unsigned int* dataHash) {
+ unsigned int a = h[0], b = h[1], c = h[2], d = h[3], x;
+ int ii;
+
+ // Round 1
+ for(ii = 0; ii < 16; ii++) {
+ x = b;
+ b = FF(a, b, c, d, ii, dataHash);
+ a = d; d = c; c = x;
+ }
+
+ // Round 2
+ for(; ii < 32; ii++) {
+ x = b;
+ b = GG(a, b, c, d, ii, dataHash);
+ a = d; d = c; c = x;
+ }
+
+ // Round 3
+ for(; ii < 48; ii++) {
+ x = b;
+ b = HH(a, b, c, d, ii, dataHash);
+ a = d; d = c; c = x;
+ }
+
+ // Round 4
+ for(; ii < 64; ii++) {
+ x = b;
+ b = II(a, b, c, d, ii, dataHash);
+ a = d; d = c; c = x;
+ }
+
+ dataHash[SHIDX(0)] = a + h[0];
+ dataHash[SHIDX(1)] = b + h[1];
+ dataHash[SHIDX(2)] = c + h[2];
+ dataHash[SHIDX(3)] = d + h[3];
+}
+
+
+#define PLAIN_MAX_SIZE 20
+
+__device__ unsigned __int64 *dataHeap;
+__device__ unsigned char *plStart;
+__device__ uint3 *plDimVec;
+__device__ unsigned char *plChrSet;
+__device__ int *plCpPos;
+__device__ int plCpPosSize;
+__device__ unsigned __int64 reduceOffset;
+__device__ unsigned __int64 plainSpaceTotal;
+__device__ unsigned __int64 rPlainSpaceTotal;
+
+
+__global__ void RTGenMD5Kernel(unsigned int chainStart, unsigned int chainStop) {
+ unsigned int *hData;
+ uint3 dimItem;
+ unsigned int uiVal, uiMul, uiDiv, idx;
+ unsigned __int64 uiVal64, uiMul64, uiDiv64, idx64;
+ unsigned int nPos, size, ii, jj, kk;
+ unsigned int cpcheck, checkpoint;
+ unsigned int plain;
+
+ __shared__ unsigned int shData[SHIDX(BLOCK_X_SIZE)];
+
+ if(threadIdx.x == 0) {
+ nPos = ((((blockIdx.y<<GRID_X_L2) + blockIdx.x)<<BLOCK_X_L2) + threadIdx.x) << 1;
+ for(ii = 0; ii < BLOCK_X_SIZE; ii++, nPos+=2) {
+ hData = shData + ((ii>>4)<<8)+(ii&15);
+ hData[SHIDX(0)] = dataHeap[nPos];
+ hData[SHIDX(1)] = dataHeap[nPos]>>32;
+ hData[SHIDX(2)] = dataHeap[nPos+1];
+ }
+ }
+ __syncthreads();
+
+ hData = shData + ((threadIdx.x>>4)<<8)+(threadIdx.x&15);
+
+ idx64 = hData[SHIDX(1)];
+ idx64 = (idx64<<32) | hData[SHIDX(0)];
+ cpcheck = hData[SHIDX(2)];
+ checkpoint = cpcheck&0x0000ffff;
+ cpcheck = cpcheck>>16;
+
+ for(nPos = chainStart; nPos < chainStop; nPos++) {
+ // transform to the plain text
+ plain = 0x80;
+ jj = (PLAIN_MAX_SIZE>>2)+1;
+ for(ii = 0; idx64 > 0xfffffff0ull && ii < PLAIN_MAX_SIZE; ii++) {
+ uiVal64 = idx64 + plStart[ii];
+ uiVal64--;
+ dimItem = plDimVec[ii];
+ uiMul64 = (unsigned __int64)dimItem.y<<32;
+ idx64 = __umul64hi(uiVal64, uiMul64);
+ uiDiv64 = uiVal64 - idx64*(unsigned __int64)dimItem.x;
+ uiVal = __umulhi((unsigned int)uiDiv64, dimItem.y);
+ uiDiv = (unsigned int)uiDiv64 - uiVal * dimItem.x;
+ idx64 += uiVal;
+ if(uiDiv >= dimItem.x) {
+ uiDiv -= dimItem.x;
+ idx64++;
+ }
+ plain = (plain<<8) | plChrSet[dimItem.z + uiDiv];
+ if((ii&3) == 2) {
+ hData[SHIDX(jj--)] = plain;
+ plain = 0;
+ }
+ }
+ for(idx = (unsigned int)idx64; idx != 0 && ii < PLAIN_MAX_SIZE; ii++) {
+ uiVal = idx + plStart[ii];
+ uiVal--;
+ dimItem = plDimVec[ii];
+ idx = __umulhi(uiVal, dimItem.y);
+ uiDiv = uiVal - idx*dimItem.x;
+ if(uiDiv >= dimItem.x) {
+ uiDiv -= dimItem.x;
+ idx++;
+ }
+ plain = (plain<<8) | plChrSet[dimItem.z + uiDiv];
+ if((ii&3) == 2) {
+ hData[SHIDX(jj--)] = plain;
+ plain = 0;
+ }
+ }
+
+ // prepare for MD5
+ size = ii;
+ ii = ((((3-(ii&3))<<3)-1)&0x1f)+1;
+ plain = plain<<ii;
+ for(jj++, kk = 0; jj <= (PLAIN_MAX_SIZE>>2)+1; plain = hData[SHIDX(jj++)], kk++)
+ hData[SHIDX(kk)] = (plain>>ii)|(hData[SHIDX(jj)]<<(32-ii));
+ hData[SHIDX(kk)] = plain>>ii;
+ for(kk++; kk < 14; kk++)
+ hData[SHIDX(kk)] = 0;
+ hData[SHIDX(kk++)] = size<<3;
+ hData[SHIDX(kk)] = 0;
+
+ // hash
+ MD5(hData);
+
+ idx64 = hData[SHIDX(1)];
+ idx64 = (idx64<<32) | hData[SHIDX(0)];
+ idx64 += reduceOffset + nPos;
+ uiDiv64 = __umul64hi(idx64, rPlainSpaceTotal);
+ idx64 -= uiDiv64*plainSpaceTotal;
+ if(idx64 >= plainSpaceTotal)
+ idx64 -= plainSpaceTotal;
+
+ if(cpcheck < plCpPosSize && nPos == plCpPos[cpcheck]) {
+ checkpoint |= ((unsigned int)idx64&1) << cpcheck;
+ cpcheck++;
+ }
+ }
+
+ hData[SHIDX(0)] = idx64;
+ hData[SHIDX(1)] = idx64>>32;
+ hData[SHIDX(2)] = (cpcheck<<16)|(checkpoint&0x0000ffff);
+ __syncthreads();
+
+ if(threadIdx.x == 0) {
+ nPos = ((((blockIdx.y<<GRID_X_L2) + blockIdx.x)<<BLOCK_X_L2) + threadIdx.x) << 1;
+ for(ii = 0; ii < BLOCK_X_SIZE; ii++, nPos+=2) {
+ hData = shData + ((ii>>4)<<8)+(ii&15);
+ dataHeap[nPos] = ((unsigned __int64)hData[SHIDX(1)]<<32)|(unsigned __int64)hData[SHIDX(0)];
+ dataHeap[nPos+1] = hData[SHIDX(2)];
+ }
+ }
+ __syncthreads();
+}
+
+
+extern "C" int CalcChainsOnCUDA(const rcuda::RCudaTask* task, unsigned __int64 *resultBuff) {
+ cudaError_t cuErr;
+ char buff[PLAIN_MAX_SIZE];
+ unsigned __int64 *data;
+ unsigned char *stPlain;
+ uint3 *dimVec;
+ unsigned char *charSet;
+ int *cpPos;
+ unsigned __int64 uiVal64;
+ time_t tStart, tEnd;
+ if(task->hash != rcuda::RHASH_MD5)
+ return 0;
+
+ memset(buff, 0, PLAIN_MAX_SIZE);
+ cudaMalloc((void**)&data, task->idxCount*2*sizeof(unsigned __int64));
+ cudaMalloc((void**)&stPlain, PLAIN_MAX_SIZE);
+ cudaMalloc((void**)&dimVec, task->dimVecSize*sizeof(uint3));
+ cudaMalloc((void**)&charSet, task->charSetSize);
+ cudaMalloc((void**)&cpPos, task->cpPosSize*sizeof(int));
+
+ cudaMemcpy(data, resultBuff, task->idxCount*2*sizeof(unsigned __int64), cudaMemcpyHostToDevice);
+ cudaMemcpy(stPlain, buff, PLAIN_MAX_SIZE, cudaMemcpyHostToDevice);
+ cudaMemcpy(stPlain, task->stPlain, min(task->stPlainSize, PLAIN_MAX_SIZE), cudaMemcpyHostToDevice);
+ cudaMemcpy(dimVec, task->dimVec, task->dimVecSize*sizeof(uint3), cudaMemcpyHostToDevice);
+ cudaMemcpy(charSet, task->charSet, task->charSetSize, cudaMemcpyHostToDevice);
+ cudaMemcpy(cpPos, task->cpPositions, task->cpPosSize*sizeof(int), cudaMemcpyHostToDevice);
+
+ cudaMemcpyToSymbol(dataHeap, &data, sizeof(data));
+ cudaMemcpyToSymbol(plStart, &stPlain, sizeof(stPlain));
+ cudaMemcpyToSymbol(plDimVec, &dimVec, sizeof(dimVec));
+ cudaMemcpyToSymbol(plChrSet, &charSet, sizeof(charSet));
+ cudaMemcpyToSymbol(plCpPos, &cpPos, sizeof(cpPos));
+ cudaMemcpyToSymbol(plCpPosSize, &task->cpPosSize, sizeof(task->cpPosSize));
+ cudaMemcpyToSymbol(reduceOffset, &task->reduceOffset, sizeof(task->reduceOffset));
+ cudaMemcpyToSymbol(plainSpaceTotal, &task->plainSpaceTotal, sizeof(task->plainSpaceTotal));
+ uiVal64 = (unsigned __int64)-1/task->plainSpaceTotal;
+ cudaMemcpyToSymbol(rPlainSpaceTotal, &uiVal64, sizeof(uiVal64));
+
+ int grSizeX = (task->idxCount-1)/BLOCK_X_SIZE + 1;
+ int grSizeY = (grSizeX-1)/GRID_X_SIZE + 1;
+ grSizeX = GRID_X_SIZE;
+ dim3 numBlocks(grSizeX, grSizeY);
+ cuErr = cudaSuccess;
+ tStart = time(NULL);
+ for(int idx = 0; idx < task->rainbowChainLen-1 && cuErr == cudaSuccess; idx+=KERN_CHAIN_SIZE) {
+ RTGenMD5Kernel<<<numBlocks, BLOCK_X_SIZE>>>(idx, min(idx+KERN_CHAIN_SIZE, task->rainbowChainLen-1));
+ cuErr = cudaGetLastError();
+ if(cuErr == cudaSuccess)
+ cuErr = cudaThreadSynchronize();
+
+ }
+ tEnd = time(NULL);
+ fprintf(stderr, "Kernel run time: %i\n", (tEnd - tStart));
+
+ if(cuErr == cudaSuccess)
+ cudaMemcpy(resultBuff, data, task->idxCount*2*sizeof(unsigned __int64), cudaMemcpyDeviceToHost);
+ else
+ fprintf(stderr, "Error happened: %d (%s)\n", cuErr, cudaGetErrorString(cuErr));
+
+ cudaFree(cpPos);
+ cudaFree(charSet);
+ cudaFree(dimVec);
+ cudaFree(stPlain);
+ cudaFree(data);
+ return cuErr==cudaSuccess? task->idxCount : -1;
+}
+
+extern "C" int GetChainsBufferSize(int minSize) {
+ int grSizeX = (minSize-1)/BLOCK_X_SIZE + 1;
+ int grSizeY = (grSizeX-1)/GRID_X_SIZE + 1;
+ grSizeX = GRID_X_SIZE;
+ return grSizeX*grSizeY*BLOCK_X_SIZE;
+}
--- /dev/null
+#ifndef RCUDA_H
+#define RCUDA_H
+
+namespace rcuda {
+
+enum RHash { RHASH_UNDEF = -1, RHASH_LM, RHASH_MD4, RHASH_MD5, RHASH_SHA1, RHASH_NTLM };
+
+struct RCudaTask {
+ RHash hash;
+ unsigned __int64 startIdx;
+ int idxCount;
+ unsigned char* stPlain;
+ int stPlainSize;
+ unsigned int* dimVec;
+ int dimVecSize;
+ unsigned char* charSet;
+ int charSetSize;
+ int *cpPositions;
+ int cpPosSize;
+ unsigned __int64 reduceOffset;
+ unsigned __int64 plainSpaceTotal;
+ unsigned int rainbowChainLen;
+};
+
+extern "C" int CalcChainsOnCUDA(const RCudaTask* task, unsigned __int64 *resultBuff);
+extern "C" int GetChainsBufferSize(int minSize);
+
+}
+
+#endif //RCUDA_H
--- /dev/null
+//============================================================================
+// Name : rcuda_ext.cpp
+// Author : Jan Kyska
+// Version : 0.9
+// Description : A code to access internals of the CChainWalkContext
+// for the CUDA generator of FreeRainbowTables
+//============================================================================
+
+#include "rcuda_ext.h"
+#include <algorithm>
+#include <string>
+#include <time.h>
+#include <fstream>
+#include <iostream>
+
+CudaCWCExtender::CudaCWCExtender(CChainWalkContext *cwc) {
+ this->cwc = cwc;
+ hash = rcuda::RHASH_UNDEF;
+}
+
+void CudaCWCExtender::Init(void) {
+ std::string hashName;
+ int ii, jj;
+
+ plainDimVec.clear();
+ plainCharSet.clear();
+
+ hashName = CChainWalkContext::m_sHashRoutineName;
+ std::transform(hashName.begin(), hashName.end(), hashName.begin(), tolower);
+ if(hashName.compare("lm") == 0)
+ hash = rcuda::RHASH_LM;
+ else if(hashName.compare("md4") == 0)
+ hash = rcuda::RHASH_MD4;
+ else if(hashName.compare("md5") == 0)
+ hash = rcuda::RHASH_MD5;
+ else if(hashName.compare("sha1") == 0)
+ hash = rcuda::RHASH_SHA1;
+ else if(hashName.compare("ntlm") == 0)
+ hash = rcuda::RHASH_NTLM;
+ else
+ hash = rcuda::RHASH_UNDEF;
+
+ for(ii = 0; ii < (int)CChainWalkContext::m_vCharset.size(); ii++) {
+ stCharset &chs = CChainWalkContext::m_vCharset[ii];
+ int chSetOffset = plainCharSet.size();
+ plainCharSet.append((char*)chs.m_PlainCharset, chs.m_nPlainCharsetLen);
+ for(jj = 0; jj < chs.m_nPlainLenMax; jj++) {
+ plainDimVec.push_back((unsigned int)chs.m_nPlainCharsetLen);
+ plainDimVec.push_back((unsigned int)-1/(unsigned int)chs.m_nPlainCharsetLen);
+ plainDimVec.push_back((unsigned int)chSetOffset);
+ }
+ }
+}
+
+int CudaCWCExtender::IndexToStartPlain(const uint64 nIndex, std::vector<unsigned char>& stPlain) {
+ int nPlainLen, nCharsetLen;
+ int ii, jj;
+
+ stPlain.clear();
+ stPlain.reserve(0x20);
+ nPlainLen = 0;
+ for(ii = CChainWalkContext::m_nPlainLenMaxTotal - 1; ii >= CChainWalkContext::m_nPlainLenMinTotal - 1; ii--) {
+ if(nIndex >= CChainWalkContext::m_nPlainSpaceUpToX[ii]) {
+ nPlainLen = ii + 1;
+ break;
+ }
+ }
+ if(nPlainLen == 0)
+ nPlainLen = CChainWalkContext::m_nPlainLenMinTotal;
+ uint64 nIndexOfX = nIndex - CChainWalkContext::m_nPlainSpaceUpToX[nPlainLen - 1];
+
+ // Slow version, but never mind
+ for(ii = nPlainLen - 1; ii >= 0; ii--) {
+ nCharsetLen = 0;
+ for(jj = 0; jj < (int)CChainWalkContext::m_vCharset.size(); jj++) {
+ stCharset &chs = CChainWalkContext::m_vCharset[jj];
+ nCharsetLen += chs.m_nPlainLenMax;
+ if(ii < nCharsetLen) { // We found the correct charset
+ stPlain.push_back(nIndexOfX % chs.m_nPlainCharsetLen + 1);
+ nIndexOfX /= chs.m_nPlainCharsetLen;
+ }
+ }
+ }
+ return stPlain.size();
+}
--- /dev/null
+#ifndef RCUDA_EXT_H
+#define RCUDA_EXT_H
+
+#include "rcuda.h"
+#include "ChainWalkContext.h"
+
+class CudaCWCExtender {
+public:
+ CudaCWCExtender(CChainWalkContext *cwc);
+ void Init(void);
+ int IndexToStartPlain(const uint64 nIndex, std::vector<unsigned char>& stPlain);
+
+ inline rcuda::RHash GetHash(void) { return hash; }
+ inline unsigned int* GetPlainDimVec(void) { return &plainDimVec[0]; }
+ inline int GetPlainDimVecSize(void) { return plainDimVec.size(); }
+ inline unsigned char* GetCharSet(void) { return (unsigned char*)plainCharSet.c_str(); }
+ inline int GetCharSetSize(void) { return plainCharSet.size(); }
+ inline uint64 GetPlainSpaceTotal(void) { return CChainWalkContext::m_nPlainSpaceTotal; }
+ inline uint64 GetRainbowTableIndex(void) { return CChainWalkContext::m_nRainbowTableIndex; }
+ inline uint64 GetReduceOffset(void) { return CChainWalkContext::m_nReduceOffset; }
+
+protected:
+ CChainWalkContext *cwc;
+ rcuda::RHash hash;
+ std::vector<unsigned int> plainDimVec;
+ std::string plainCharSet;
+};
+
+#endif //RCUDA_EXT_H
projects / freerainbowtables / commitdiff