// The kernel (this is the entrypoint of GPU code)
// Loads the 64-byte word to be hashed from g	lobal to shared memory and calls the calculation routine
__global__ void md5_calc(char *gwords, char *paddedWords3, uint *hash, int realthreads, int msg_size, int size_hash, uint launch){
	int linidx;
	uint a;
	uint b = 0;
	uint c;
	uint d;

	char msg[64] = {0};

	// assuming blockDim.y = 1 and threadIdx.y = 0, always
	int iblock= blockIdx.x + blockIdx.y * gridDim.x;
	linidx = threadIdx.x + iblock * blockDim.x;

	// this check slows down the code by ~0.4% (measured)
	if(linidx >= realthreads){
		return; 
	}

	// load the dictionary word for this thread
	//char *word = &memory[0] + threadIdx.x * msg_size;

	for(a = 0; a < msg_size; a++){
		//word[(a + threadIdx.x) % msg_size] = gwords[linidx * msg_size + a];
		//msg[a] = word[(a + threadIdx.x) % msg_size];
		msg[a] = gwords[linidx * msg_size + a];

		//end of the password
		if(msg[a] == 0){
			msg[a] = 0x80;
			b = a;

			break;
		}
	}

	if(!b){
		msg[msg_size] = 0x80;
	
		msg[56] = msg_size * 8;
	}
	else{
		msg[56] = b * 8;
	}


	for(a = 0; a < 64; a++){
		paddedWords3[linidx * 64 + a] = msg[a];
	}

	paddedWords3[linidx * 64 + 63] = linidx;
	paddedWords3[linidx * 64 + 62] = launch;
	paddedWords3[linidx * 64 + 61] = (launch * realthreads) + linidx + 1;

	// compute MD5 hash
	md5_gpu((uint *)msg, a, b, c, d);
	//md5_gpu((uint *)msg1, a1, b1, c1, d1);

	// return the hash converted to number
	// all this work around is because of different edian order
	if(size_hash == 1){
		hash[(linidx)] = ((a & 240) >> 4);
	}
	else{
		if(size_hash == 2){
			hash[(linidx)] = a & 255;
		}
		else{
			if(size_hash == 3){
				hash[(linidx)] = ((a & 61440) >> 12) + ((a & 240) << 4) + ((a & 15) << 4);
			}
			else{
				if(size_hash == 4){
					hash[(linidx)] = ((a & 61440) >> 8) + ((a & 3840) >> 8) + ((a & 240) << 8) + ((a & 15) << 8);
				}
			}
		}
	}
}