__device__ unsigned char *plStart;
__device__ uint3 *plDimVec;
__device__ unsigned char *plChrSet;
__device__ int *plCpPos;
__device__ int plCpPosSize;
__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;
__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;
unsigned int nPos, size, ii, jj, kk;
unsigned int cpcheck, checkpoint;
unsigned int plain;
uiVal = __umulhi((unsigned int)uiDiv64, dimItem.y);
uiDiv = (unsigned int)uiDiv64 - uiVal * dimItem.x;
idx64 += uiVal;
uiVal = __umulhi((unsigned int)uiDiv64, dimItem.y);
uiDiv = (unsigned int)uiDiv64 - uiVal * dimItem.x;
idx64 += uiVal;
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);
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);
time_t tStart, tEnd;
if(task->hash != rcuda::RHASH_MD5)
return 0;
memset(buff, 0, PLAIN_MAX_SIZE);
time_t tStart, tEnd;
if(task->hash != rcuda::RHASH_MD5)
return 0;
memset(buff, 0, PLAIN_MAX_SIZE);
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));
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);
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);
cudaMemcpyToSymbol(plCpPosSize, &task->cpPosSize, sizeof(task->cpPosSize));
cudaMemcpyToSymbol(reduceOffset, &task->reduceOffset, sizeof(task->reduceOffset));
cudaMemcpyToSymbol(plainSpaceTotal, &task->plainSpaceTotal, sizeof(task->plainSpaceTotal));
cudaMemcpyToSymbol(plCpPosSize, &task->cpPosSize, sizeof(task->cpPosSize));
cudaMemcpyToSymbol(reduceOffset, &task->reduceOffset, sizeof(task->reduceOffset));
cudaMemcpyToSymbol(plainSpaceTotal, &task->plainSpaceTotal, sizeof(task->plainSpaceTotal));
cudaMemcpyToSymbol(rPlainSpaceTotal, &uiVal64, sizeof(uiVal64));
int grSizeX = (task->idxCount-1)/BLOCK_X_SIZE + 1;
cudaMemcpyToSymbol(rPlainSpaceTotal, &uiVal64, sizeof(uiVal64));
int grSizeX = (task->idxCount-1)/BLOCK_X_SIZE + 1;
- 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));
else
fprintf(stderr, "Error happened: %d (%s)\n", cuErr, cudaGetErrorString(cuErr));