--- nt_opencl_kernel.c	2010-12-21 19:38:14.983334057 -0800
+++ nt_opencl_kernel_fixed.c	2010-12-21 19:52:53.120000725 -0800
@@ -55,15 +55,18 @@ __kernel void nt_crypt(const __global ui
 	uint key_chars = keys[i];//Coalescing access to global memory
 	uint cache_key = GET_CHAR(key_chars,ELEM_0);
 	//Extract 4 chars by cycle
+	int jump = 0;
 	while(cache_key)
 	{
 		md4_size++;
 		uint temp = GET_CHAR(key_chars,ELEM_1);
 		nt_buffer[nt_index] = ((temp ? temp : 0x80) << 16) | cache_key;
 		
-		if(!temp)
-			goto key_cleaning;
-		
+		if(!temp) {
+			jump = 1;
+			break;
+		}
+			
 		md4_size++;
 		nt_index++;
 		cache_key = GET_CHAR(key_chars,ELEM_2);
@@ -76,8 +79,10 @@ __kernel void nt_crypt(const __global ui
 		temp = GET_CHAR(key_chars,ELEM_3);
 		nt_buffer[nt_index] = ((temp ? temp : 0x80) << 16) | cache_key;
 		
-		if(!temp)
-			goto key_cleaning;
+		if(!temp) {
+			jump = 1;
+			break;
+		}
 		
 		md4_size++;
 		nt_index++;
@@ -86,7 +91,8 @@ __kernel void nt_crypt(const __global ui
 		cache_key = GET_CHAR(key_chars,ELEM_0);
 	}
 	
-	nt_buffer[nt_index] = 0x80;
+	if(!jump)
+		nt_buffer[nt_index] = 0x80;
 	
 key_cleaning:
 	nt_index++;