From 866b5a58b1bace57fa9da5a9d72bf4823304108f Mon Sep 17 00:00:00 2001 From: James Nobis Date: Sun, 7 Nov 2010 22:52:19 -0600 Subject: [PATCH] cleanup signed/unsigned warnings *nix compat unsigned __int64 to uint64 first pass at distrrtgen_cuda for *nix - compiles but results are incorrect --- .../BOINC client apps/distrrtgen/charset.txt | 27 +++++ .../distrrtgen/distrrtgen.cpp | 14 +-- .../distrrtgen_cuda/Makefile | 110 ++++++++++++++++++ .../distrrtgen_cuda/charset.txt | 61 ++++++++++ .../distrrtgen_cuda/distrrtgen.cpp | 12 +- .../distrrtgen_cuda/distrrtgen.h | 29 +++++ .../distrrtgen_cuda/rcuda.cu | 30 ++--- .../BOINC client apps/distrrtgen_cuda/rcuda.h | 10 +- .../distrrtgen_cuda/rcuda_ext.cpp | 2 +- 9 files changed, 262 insertions(+), 33 deletions(-) create mode 100644 BOINC software/BOINC client apps/distrrtgen_cuda/Makefile create mode 100644 BOINC software/BOINC client apps/distrrtgen_cuda/charset.txt create mode 100644 BOINC software/BOINC client apps/distrrtgen_cuda/distrrtgen.h diff --git a/BOINC software/BOINC client apps/distrrtgen/charset.txt b/BOINC software/BOINC client apps/distrrtgen/charset.txt index 6e749fa..d1e0179 100644 --- a/BOINC software/BOINC client apps/distrrtgen/charset.txt +++ b/BOINC software/BOINC client apps/distrrtgen/charset.txt @@ -32,3 +32,30 @@ mixalpha-numeric-symbol14 = [abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVW mixalpha-numeric-all = [abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789!@#$%^&*()-_+=~`[]{}|\:;"'<>,.?/] mixalpha-numeric-symbol32-space = [abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789!@#$%^&*()-_+=~`[]{}|\:;"'<>,.?/ ] mixalpha-numeric-all-space = [abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789!@#$%^&*()-_+=~`[]{}|\:;"'<>,.?/ ] + + + + + + + + + + + + + + + + + + + + + + + + + + + \ No newline at end of file diff --git a/BOINC software/BOINC client apps/distrrtgen/distrrtgen.cpp b/BOINC software/BOINC client apps/distrrtgen/distrrtgen.cpp index d936785..5ecd8cf 100644 --- a/BOINC software/BOINC client apps/distrrtgen/distrrtgen.cpp +++ b/BOINC software/BOINC client apps/distrrtgen/distrrtgen.cpp @@ -127,7 +127,7 @@ int main(int argc, char **argv) { return -1; } string sHashRoutineName, sCharsetName, sSalt, sCheckPoints; - int nRainbowChainCount, nPlainLenMin, nPlainLenMax, nRainbowTableIndex, nRainbowChainLen; + uint32 nRainbowChainCount, nPlainLenMin, nPlainLenMax, nRainbowTableIndex, nRainbowChainLen; uint64 nChainStart; sHashRoutineName = argv[1]; sCharsetName = argv[2]; @@ -144,7 +144,7 @@ int main(int argc, char **argv) { nChainStart = atoll(argv[8]); #endif sCheckPoints = argv[9]; - vector vCPPositions; + vector vCPPositions; char *cp = strtok((char *)sCheckPoints.c_str(), ","); while(cp != NULL) { @@ -224,7 +224,7 @@ int main(int argc, char **argv) { // Round to boundary nDataLen = nDataLen / 18 * 18; - if ((int)nDataLen == nRainbowChainCount * 18) + if (nDataLen == nRainbowChainCount * 18) { std::cerr << "precomputation of this rainbow table already finished" << std::endl; fclose(outfile); @@ -237,7 +237,7 @@ int main(int argc, char **argv) { uint64 nIndex[2]; time_t tStart = time(NULL); // std::cout << "Starting to generate chains" << std::endl; - for(int nCurrentCalculatedChains = nDataLen / 18; nCurrentCalculatedChains < nRainbowChainCount; nCurrentCalculatedChains++) + for(uint32 nCurrentCalculatedChains = nDataLen / 18; nCurrentCalculatedChains < nRainbowChainCount; nCurrentCalculatedChains++) { uint32 cpcheck = 0; unsigned short checkpoint = 0; @@ -246,7 +246,7 @@ int main(int argc, char **argv) { cwc.SetIndex(nChainStart++); // use a given index now! nIndex[0] = cwc.GetIndex(); - for (int nPos = 0; nPos < nRainbowChainLen - 1; nPos++) + for (uint32 nPos = 0; nPos < nRainbowChainLen - 1; nPos++) { // std::cout << "IndexToPlain()" << std::endl; cwc.IndexToPlain(); @@ -290,7 +290,7 @@ int main(int argc, char **argv) { { // Load file fseek(outfile, 0, SEEK_SET); - for(int i = 0; i < nRainbowChainCount; i++) + for(uint32 i = 0; i < nRainbowChainCount; i++) { if(fread(&pChain[i], 1, 16, outfile) != 16) { @@ -309,7 +309,7 @@ int main(int argc, char **argv) { // Write file fseek(outfile, 0, SEEK_SET); - for(int i = 0; i < nRainbowChainCount; i++) + for(uint32 i = 0; i < nRainbowChainCount; i++) { fwrite(&pChain[i], 1, 16, outfile); fwrite(&pChain[i].nCheckPoint, 2, 1, outfile); diff --git a/BOINC software/BOINC client apps/distrrtgen_cuda/Makefile b/BOINC software/BOINC client apps/distrrtgen_cuda/Makefile new file mode 100644 index 0000000..ccab6ed --- /dev/null +++ b/BOINC software/BOINC client apps/distrrtgen_cuda/Makefile @@ -0,0 +1,110 @@ +# freerainbowtables is a project for generating, distributing, and using +# perfect rainbow tables +# +# Copyright (C) Zhu Shuanglei +# Copyright 2010 Martin Westergaard Jørgensen +# Copyright 2009, 2010 Daniël Niggebrugge +# Copyright 2009, 2010 James Nobis +# +# This file is part of freerainbowtables. +# +# freerainbowtables is free software: you can redistribute it and/or modify +# it under the terms of the GNU General Public License as published by +# the Free Software Foundation, either version 2 of the License, or +# (at your option) any later version. +# +# freerainbowtables 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 General Public License for more details. +# +# You should have received a copy of the GNU General Public License +# along with freerainbowtables. If not, see . + +SHELL = /bin/sh + +#BOINC_DIR = /home/ubuntu/boinc +BOINC_DIR = /usr/share/boinc-dev +BOINC_API_DIR = $(BOINC_DIR)/api +BOINC_LIB_DIR = $(BOINC_DIR)/lib +CC = g++ +COMMON_LIB_DIR = /usr/lib +INCLUDES = -I../../../Common/rt\ api -I$(BOINC_DIR) -I$(BOINC_LIB_DIR) -I$(BOINC_API_DIR) -I/usr/local/cuda/include +COMMON_API_PATH = ../../../Common/rt\ api + +OPTIMIZATION = -O3 +WARNALL = -Wextra -Wunused-macros -Wunsafe-loop-optimizations -Wundef -Woverlength-strings -Wdisabled-optimization -Wformat-extra-args -Wformat-security -Winline + +# -static -static-libgcc +CXXFLAGS = -Wall -ansi -c -DBOINC $(INCLUDES) $(OPTIMIZATION) $(DEBUG) + +# -static-libgcc +LFLAGS = -Wall -ansi $(INCLUDES) $(OPTIMIZATION) -L$(BOINC_API_DIR) \ + -L$(BOINC_LIB_DIR) -L/usr/X11R6/lib -L/usr/lib -L/usr/local/cuda/lib64 -L/usr/local/cuda/lib -L. $(DEBUG) + +LIBS = -lboinc_api -lboinc -lpthread -lssl -lcudart libstdc++.a $(COMMON_LIB_DIR)/libssl.a $(COMMON_LIB_DIR)/libpthread.a +OBJS = ChainWalkContext.o des_enc.o des_setkey.o distrrtgen.o ecb_enc.o HashAlgorithm.o HashRoutine.o md4.o md5.o Public.o rcuda_ext.o rcuda.o + +OSNAME = $(shell uname -s) + +all: distrrtgen + +distrrtgen: libstdc++.a $(BOINC_API_DIR)/libboinc_api.a $(BOINC_LIB_DIR)/libboinc.a $(OBJS) + $(CXX) $(LFLAGS) $(OBJS) $(LIBS) -o distrrtgen + +libstdc++.a: + ln -s `$(CC) $(CXXFLAGS) -print-file-name=libstdc++.a` + +clean: + /bin/rm -f distrrtgen *.o + +debug: DEBUG += -DDEBUG -g +debug: all + +debugall: DEBUG += -DDEBUG -g $(WARNALL) +debugall: all + +distclean: + /bin/rm -f distrrtgen *.o libstdc++.a + +m32: DEBUG += -m32 +m32: all + +rebuild: clean all + +ChainWalkContext.o: $(COMMON_API_PATH)/ChainWalkContext.h $(COMMON_API_PATH)/ChainWalkContext.cpp $(COMMON_API_PATH)/HashRoutine.h $(COMMON_API_PATH)/Public.h + $(CXX) $(CXXFLAGS) $(COMMON_API_PATH)/ChainWalkContext.cpp + +des_enc.o: $(COMMON_API_PATH)/des_enc.c $(COMMON_API_PATH)/des_locl.h $(COMMON_API_PATH)/spr.h + $(CXX) $(CXXFLAGS) $(COMMON_API_PATH)/des_enc.c + +des_setkey.o: $(COMMON_API_PATH)/des_setkey.c $(COMMON_API_PATH)/des_locl.h $(COMMON_API_PATH)/podd.h $(COMMON_API_PATH)/sk.h + $(CXX) $(CXXFLAGS) $(COMMON_API_PATH)/des_setkey.c + +distrrtgen.o: distrrtgen.cpp distrrtgen.h $(COMMON_API_PATH)/ChainWalkContext.h $(COMMON_API_PATH)/Public.h + $(CXX) $(CXXFLAGS) distrrtgen.cpp + +ecb_enc.o: $(COMMON_API_PATH)/ecb_enc.c $(COMMON_API_PATH)/des_locl.h $(COMMON_API_PATH)/spr.h + $(CXX) $(CXXFLAGS) $(COMMON_API_PATH)/ecb_enc.c + +HashAlgorithm.o: $(COMMON_API_PATH)/HashAlgorithm.h $(COMMON_API_PATH)/HashAlgorithm.cpp $(COMMON_API_PATH)/Public.h $(COMMON_API_PATH)/md4.h + $(CXX) $(CXXFLAGS) $(COMMON_API_PATH)/HashAlgorithm.cpp + +HashRoutine.o: $(COMMON_API_PATH)/HashRoutine.h $(COMMON_API_PATH)/HashRoutine.cpp $(COMMON_API_PATH)/global.h $(COMMON_API_PATH)/HashAlgorithm.h + $(CXX) $(CXXFLAGS) $(COMMON_API_PATH)/HashRoutine.cpp + +md4.o: $(COMMON_API_PATH)/md4.h $(COMMON_API_PATH)/md4.cpp $(COMMON_API_PATH)/global.h + $(CXX) $(CXXFLAGS) $(COMMON_API_PATH)/md4.cpp + +md5.o: $(COMMON_API_PATH)/md5.h $(COMMON_API_PATH)/md5.cpp $(COMMON_API_PATH)/global.h + $(CXX) $(CXXFLAGS) $(COMMON_API_PATH)/md5.cpp + +Public.o: $(COMMON_API_PATH)/Public.h $(COMMON_API_PATH)/Public.cpp $(COMMON_API_PATH)/global.h + $(CXX) $(CXXFLAGS) $(COMMON_API_PATH)/Public.cpp + +rcuda_ext.o: rcuda_ext.h rcuda_ext.cpp $(COMMON_API_PATH)/ChainWalkContext.h + $(CXX) $(CXXFLAGS) rcuda_ext.cpp + +rcuda.o: rcuda.h rcuda.cu $(COMMON_API_PATH)/Public.h $(COMMON_API_PATH)/global.h + /usr/local/cuda/bin/nvcc -I../../../Common/rt\ api --compile rcuda.cu +# /usr/local/cuda/bin/nvcc --machine 32 -I../../../Common/rt\ api --compile rcuda.cu diff --git a/BOINC software/BOINC client apps/distrrtgen_cuda/charset.txt b/BOINC software/BOINC client apps/distrrtgen_cuda/charset.txt new file mode 100644 index 0000000..d1e0179 --- /dev/null +++ b/BOINC software/BOINC client apps/distrrtgen_cuda/charset.txt @@ -0,0 +1,61 @@ +# charset configuration file for DistrRTgen v3.2 by Martin Westergaard (martinwj2005@gmail.com) + +byte = [] +alpha = [ABCDEFGHIJKLMNOPQRSTUVWXYZ] +alpha-space = [ABCDEFGHIJKLMNOPQRSTUVWXYZ ] +alpha-numeric = [ABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789] +alpha-numeric-space = [ABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789 ] +alpha-numeric-symbol14 = [ABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789!@#$%^&*()-_+=] +alpha-numeric-symbol14-space= [ABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789!@#$%^&*()-_+= ] +all = [ABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789!@#$%^&*()-_+=~`[]{}|\:;"'<>,.?/] +all-space = [ABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789!@#$%^&*()-_+=~`[]{}|\:;"'<>,.?/ ] +alpha-numeric-symbol32-space = [ABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789!@#$%^&*()-_+=~`[]{}|\:;"'<>,.?/ ] +lm-frt-cp437 = [ !"#$%&'()*+,-./0123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[\]^_`{|}~€Ž’™š›œžŸ¥àáâãäæçèéêëî] +lm-frt-cp850 = [ !"#$%&'()*+,-./0123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[\]^_`{|}~€Ž’™šœŸ¥µ¶·½¾ÇÏÑÒÓÔÕÖ×ØÞàáâãåæèéêëíï] +lm-frt-cp437-850 = [ !"#$%&'()*+,-./0123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[\]^_`{|}~€Ž’™š›œžŸ¥µ¶·½¾ÇÏÑÒÓÔÕÖ×ØÞàáâãäåæçèéêëíîï] + +numeric = [0123456789] +numeric-space = [0123456789 ] +loweralpha = [abcdefghijklmnopqrstuvwxyz] +loweralpha-space = [abcdefghijklmnopqrstuvwxyz ] +loweralpha-numeric = [abcdefghijklmnopqrstuvwxyz0123456789] +loweralpha-numeric-space = [abcdefghijklmnopqrstuvwxyz0123456789 ] +loweralpha-numeric-symbol14 = [abcdefghijklmnopqrstuvwxyz0123456789!@#$%^&*()-_+=] +loweralpha-numeric-all = [abcdefghijklmnopqrstuvwxyz0123456789!@#$%^&*()-_+=~`[]{}|\:;"'<>,.?/] +loweralpha-numeric-symbol32-space= [abcdefghijklmnopqrstuvwxyz0123456789!@#$%^&*()-_+=~`[]{}|\:;"'<>,.?/ ] + +mixalpha = [abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ] +mixalpha-space = [abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ ] +mixalpha-numeric = [abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789] +mixalpha-numeric-space = [abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789 ] +mixalpha-numeric-symbol14 = [abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789!@#$%^&*()-_+=] +mixalpha-numeric-all = [abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789!@#$%^&*()-_+=~`[]{}|\:;"'<>,.?/] +mixalpha-numeric-symbol32-space = [abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789!@#$%^&*()-_+=~`[]{}|\:;"'<>,.?/ ] +mixalpha-numeric-all-space = [abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789!@#$%^&*()-_+=~`[]{}|\:;"'<>,.?/ ] + + + + + + + + + + + + + + + + + + + + + + + + + + + \ No newline at end of file diff --git a/BOINC software/BOINC client apps/distrrtgen_cuda/distrrtgen.cpp b/BOINC software/BOINC client apps/distrrtgen_cuda/distrrtgen.cpp index 001e54a..0d7145d 100644 --- a/BOINC software/BOINC client apps/distrrtgen_cuda/distrrtgen.cpp +++ b/BOINC software/BOINC client apps/distrrtgen_cuda/distrrtgen.cpp @@ -74,7 +74,7 @@ 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); + int nRandomIndex = nLow + ((uint32)rand() * ((uint32)RAND_MAX + 1) + (uint32)rand()) % (nHigh - nLow + 1); RainbowChainCP TempChain; TempChain = pChain[nLow]; pChain[nLow] = pChain[nRandomIndex]; @@ -229,7 +229,7 @@ int main(int argc, char **argv) { // Round to boundary nDataLen = nDataLen / 18 * 18; - if ((int)nDataLen == nRainbowChainCount * 18) + if (nDataLen == nRainbowChainCount * 18) { std::cerr << "precomputation of this rainbow table already finished" << std::endl; fclose(outfile); @@ -253,7 +253,7 @@ int main(int argc, char **argv) { time_t tStart, tStartFinal, tEndFinal; time_t tEnd; tStartFinal = time(NULL); - for(int nCurrentCalculatedChains = nDataLen / 18, calcSize; nCurrentCalculatedChains < nRainbowChainCount; ) + for(uint32 nCurrentCalculatedChains = nDataLen / 18, calcSize; nCurrentCalculatedChains < nRainbowChainCount; ) { fd = (double)nCurrentCalculatedChains / (double)nRainbowChainCount; boinc_fraction_done(fd); @@ -272,7 +272,7 @@ time_t tEnd; cuTask.reduceOffset = ex.GetReduceOffset(); cuTask.plainSpaceTotal = ex.GetPlainSpaceTotal(); cuTask.rainbowChainLen = nRainbowChainLen; - for(int ii = 0; ii < cuTask.idxCount; ii++) { + for(ii = 0; ii < cuTask.idxCount; ii++) { calcBuff[2*ii] = cuTask.startIdx + ii; calcBuff[2*ii+1] = 0; } @@ -321,7 +321,7 @@ time_t tEnd; std::cout << "Sorting file" << std::endl; #endif fseek(outfile, 0, SEEK_SET); - for(int i = 0; i < nRainbowChainCount; i++) + for(uint32 i = 0; i < nRainbowChainCount; i++) { if(fread(&pChain[i], 1, 16, outfile) != 16) { @@ -340,7 +340,7 @@ time_t tEnd; // Write file fseek(outfile, 0, SEEK_SET); - for(int i = 0; i < nRainbowChainCount; i++) + for(uint32 i = 0; i < nRainbowChainCount; i++) { fwrite(&pChain[i], 1, 16, outfile); fwrite(&pChain[i].nCheckPoint, 2, 1, outfile); diff --git a/BOINC software/BOINC client apps/distrrtgen_cuda/distrrtgen.h b/BOINC software/BOINC client apps/distrrtgen_cuda/distrrtgen.h new file mode 100644 index 0000000..d360617 --- /dev/null +++ b/BOINC software/BOINC client apps/distrrtgen_cuda/distrrtgen.h @@ -0,0 +1,29 @@ +// 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 . + +#include "boinc_api.h" + +struct UC_SHMEM { + double update_time; + double fraction_done; + double cpu_time; + BOINC_STATUS status; + int countdown; + // graphics app sets this to 5 repeatedly, + // main program decrements it once/sec. + // If it's zero, don't bother updating shmem +}; diff --git a/BOINC software/BOINC client apps/distrrtgen_cuda/rcuda.cu b/BOINC software/BOINC client apps/distrrtgen_cuda/rcuda.cu index e626bd4..187b105 100644 --- a/BOINC software/BOINC client apps/distrrtgen_cuda/rcuda.cu +++ b/BOINC software/BOINC client apps/distrrtgen_cuda/rcuda.cu @@ -123,22 +123,22 @@ __device__ void MD5(unsigned int* dataHash) { #define PLAIN_MAX_SIZE 20 -__device__ unsigned __int64 *dataHeap; +__device__ uint64 *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; +__device__ uint64 reduceOffset; +__device__ uint64 plainSpaceTotal; +__device__ uint64 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; + uint64 uiVal64, uiMul64, uiDiv64, idx64; unsigned int nPos, size, ii, jj, kk; unsigned int cpcheck, checkpoint; unsigned int plain; @@ -172,9 +172,9 @@ __global__ void RTGenMD5Kernel(unsigned int chainStart, unsigned int chainStop) uiVal64 = idx64 + plStart[ii]; uiVal64--; dimItem = plDimVec[ii]; - uiMul64 = (unsigned __int64)dimItem.y<<32; + uiMul64 = (uint64)dimItem.y<<32; idx64 = __umul64hi(uiVal64, uiMul64); - uiDiv64 = uiVal64 - idx64*(unsigned __int64)dimItem.x; + uiDiv64 = uiVal64 - idx64*(uint64)dimItem.x; uiVal = __umulhi((unsigned int)uiDiv64, dimItem.y); uiDiv = (unsigned int)uiDiv64 - uiVal * dimItem.x; idx64 += uiVal; @@ -243,7 +243,7 @@ __global__ void RTGenMD5Kernel(unsigned int chainStart, unsigned int chainStop) nPos = ((((blockIdx.y<>4)<<8)+(ii&15); - dataHeap[nPos] = ((unsigned __int64)hData[SHIDX(1)]<<32)|(unsigned __int64)hData[SHIDX(0)]; + dataHeap[nPos] = ((uint64)hData[SHIDX(1)]<<32)|(uint64)hData[SHIDX(0)]; dataHeap[nPos+1] = hData[SHIDX(2)]; } } @@ -251,27 +251,27 @@ __global__ void RTGenMD5Kernel(unsigned int chainStart, unsigned int chainStop) } -extern "C" int CalcChainsOnCUDA(const rcuda::RCudaTask* task, unsigned __int64 *resultBuff) { +extern "C" int CalcChainsOnCUDA(const rcuda::RCudaTask* task, uint64 *resultBuff) { cudaError_t cuErr; char buff[PLAIN_MAX_SIZE]; - unsigned __int64 *data; + uint64 *data; unsigned char *stPlain; uint3 *dimVec; unsigned char *charSet; int *cpPos; - unsigned __int64 uiVal64; + uint64 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**)&data, task->idxCount*2*sizeof(uint64)); 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(data, resultBuff, task->idxCount*2*sizeof(uint64), 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); @@ -286,7 +286,7 @@ extern "C" int CalcChainsOnCUDA(const rcuda::RCudaTask* task, unsigned __int64 * 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; + uiVal64 = (uint64)-1/task->plainSpaceTotal; cudaMemcpyToSymbol(rPlainSpaceTotal, &uiVal64, sizeof(uiVal64)); int grSizeX = (task->idxCount-1)/BLOCK_X_SIZE + 1; @@ -306,7 +306,7 @@ extern "C" int CalcChainsOnCUDA(const rcuda::RCudaTask* task, unsigned __int64 * fprintf(stderr, "Kernel run time: %i\n", (tEnd - tStart)); if(cuErr == cudaSuccess) - cudaMemcpy(resultBuff, data, task->idxCount*2*sizeof(unsigned __int64), cudaMemcpyDeviceToHost); + cudaMemcpy(resultBuff, data, task->idxCount*2*sizeof(uint64), cudaMemcpyDeviceToHost); else fprintf(stderr, "Error happened: %d (%s)\n", cuErr, cudaGetErrorString(cuErr)); diff --git a/BOINC software/BOINC client apps/distrrtgen_cuda/rcuda.h b/BOINC software/BOINC client apps/distrrtgen_cuda/rcuda.h index eeeac69..370b064 100644 --- a/BOINC software/BOINC client apps/distrrtgen_cuda/rcuda.h +++ b/BOINC software/BOINC client apps/distrrtgen_cuda/rcuda.h @@ -1,13 +1,15 @@ #ifndef RCUDA_H #define RCUDA_H +#include "Public.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; + uint64 startIdx; int idxCount; unsigned char* stPlain; int stPlainSize; @@ -17,12 +19,12 @@ struct RCudaTask { int charSetSize; int *cpPositions; int cpPosSize; - unsigned __int64 reduceOffset; - unsigned __int64 plainSpaceTotal; + uint64 reduceOffset; + uint64 plainSpaceTotal; unsigned int rainbowChainLen; }; -extern "C" int CalcChainsOnCUDA(const RCudaTask* task, unsigned __int64 *resultBuff); +extern "C" int CalcChainsOnCUDA(const RCudaTask* task, uint64 *resultBuff); extern "C" int GetChainsBufferSize(int minSize); } diff --git a/BOINC software/BOINC client apps/distrrtgen_cuda/rcuda_ext.cpp b/BOINC software/BOINC client apps/distrrtgen_cuda/rcuda_ext.cpp index 86f31ec..74b8d9e 100644 --- a/BOINC software/BOINC client apps/distrrtgen_cuda/rcuda_ext.cpp +++ b/BOINC software/BOINC client apps/distrrtgen_cuda/rcuda_ext.cpp @@ -26,7 +26,7 @@ void CudaCWCExtender::Init(void) { plainCharSet.clear(); hashName = CChainWalkContext::m_sHashRoutineName; - std::transform(hashName.begin(), hashName.end(), hashName.begin(), tolower); + std::transform(hashName.begin(), hashName.end(), hashName.begin(), ::tolower); if(hashName.compare("lm") == 0) hash = rcuda::RHASH_LM; else if(hashName.compare("md4") == 0) -- 2.39.2