__kernel void recognition_numbers_1(__global const unsigned char* a_in,
	__global const unsigned char* b_in, int width, int num_width, int num_height, int num_count,
	__global unsigned char* result) {

	int p = get_global_id(0);
	int startIndex = p / width * width * num_width * num_height + p % (num_count * 10) * num_width;
	unsigned char t = 0;
	for (int r = 0;r < num_height;r++)
		for (int c = 0;c < num_width;c++) {
			int index = startIndex + num_width * width * r + c;
			t += abs(a_in[index] - b_in[index]);
			if (p == 0) {
				//printf("Öµ:%d-%d \n", a_in[index], b_in[index]);
			}
		}
	result[p] = t;
}


__kernel void recognition_numbers_2(__global const unsigned char* a_in,
	__global unsigned char* result) {

	int index = get_global_id(0);
	int startIndex = index * 10;
	int endIndex = (index + 1) * 10;

	//»ñÈ¡×îСֵ
	int min = 255;
	int minIndex = 11;

	for (int i = startIndex;i < endIndex;i++)
	{
		if (a_in[i] < min) {
			min = a_in[i];
			minIndex = i;
		}
	}

	//printf("×îСֵ£º%d - %d - %d \n", index, min, minIndex - startIndex);

	result[index] = minIndex - startIndex;
}

int get_one_level_position(int width, int x, int y) {
	return width * y + x;
}

unsigned char get_binary_value(unsigned char v) {
	return v >= 64 ? 1 : 0;
}

//·Ö¸ô°å¿éÊý×Ö
__kernel void split_plate_num(__global const unsigned char* img_in, __global const int* pos_in, __global const unsigned char* zero, __global const unsigned char* types, int width, int num_width, int num_height, int ele_count_per_line, int length_per_num,
	__global unsigned char* result) {

	int index = get_global_id(0);
	//ÇøÓò×ø±ê
	int startP = index * 4;
	int startx = pos_in[startP];
	int starty = pos_in[startP + 1];
	int endx = pos_in[startP + 2];
	int endy = pos_in[startP + 3];

	if (startx > 0)
	{
		int startNy = -1;
		int endNy = -1;
		//È¥³ýÉÏϵİױß
		int y = 0;
		for (y = starty;y <= endy;y++) {
			bool empty = 1;
			for (int x = startx;x <= endx;x++)
			{
				int p = get_one_level_position(width, x, y);
				unsigned char value = get_binary_value(img_in[p]);
				//ÓÐÊý¾Ý
				if (value > 0) {
					empty = 0;
					break;
				}
			}
			if (!empty) {
				startNy = y;
				break;
			}
		}

		for (y = endy;y >= starty;y--) {
			bool empty = 1;
			for (int x = startx;x <= endx;x++)
			{
				int p = get_one_level_position(width, x, y);
				unsigned char value = get_binary_value(img_in[p]);
				//ÓÐÊý¾Ý
				if (value > 0) {
					empty = 0;
					break;
				}
			}
			if (!empty) {
				endNy = y;
				break;
			}
		}
		starty = startNy;
		endy = endNy;
	}
	int cols = endx - startx + 1;
	int rows = endy - starty + 1;

	//ÿ¸öÔªËØÕ¼ÓÃ6¸öÊý×ֵĿռ䣬²»¹»µÄÇ°ÃæÌî³ä0
	int num_count = ele_count_per_line * length_per_num;

	int nps[12];
	//Ö»ÄÜʶ±ð4ÖÖÊý×Ö
	unsigned char type = types[index % ele_count_per_line];
	//printf("ele_count_line:%d\n", ele_count_line);
	//ʱ¼ä
	if (type == 0) {
		//ʱ¼ä
		int s = cols / 2;
		int i;
		//ÍùÇ°²éÕÒðºÅ
		int m1_s = -1, m1_e = -1, m2_s = -1, m2_e = -1;
		uchar temp[10];
		for (i = s;i > -1;i--) {
			if (i < 5) {
				break;
			}
			temp[0] = get_binary_value(img_in[get_one_level_position(width, startx + i, endy)]);

			temp[1] = get_binary_value(img_in[get_one_level_position(width, startx + i - 1, endy)]);


			temp[2] = get_binary_value(img_in[get_one_level_position(width, startx + i - 2, endy)]);
			temp[3] = get_binary_value(img_in[get_one_level_position(width, startx + i - 3, endy)]);
			temp[4] = get_binary_value(img_in[get_one_level_position(width, startx + i - 4, endy)]);
			temp[5] = get_binary_value(img_in[get_one_level_position(width, startx + i - 5, endy)]);
			temp[6] = get_binary_value(img_in[get_one_level_position(width, startx + i - 2, endy - 1)]);
			temp[7] = get_binary_value(img_in[get_one_level_position(width, startx + i - 3, endy - 1)]);
			temp[8] = get_binary_value(img_in[get_one_level_position(width, startx + i - 2, endy - 2)]);
			temp[9] = get_binary_value(img_in[get_one_level_position(width, startx + i - 3, endy - 2)]);
			if (temp[3] > 0 && temp[2] > 0 && temp[6] > 0 && temp[7] > 0) {
				int t = temp[0] + temp[1] + temp[4] + temp[5] + temp[8] + temp[9];
				if (t == 0) {
					m1_s = i - 5;
					m1_e = i;
					break;
				}
			}

		}

		//Íùºó²éÕÒðºÅ

		for (i = s;i < cols;i++) {
			if (i > cols - 5) {
				break;
			}

			temp[0] = get_binary_value(img_in[get_one_level_position(width, startx + i, endy)]);
			temp[1] = get_binary_value(img_in[get_one_level_position(width, startx + i + 1, endy)]);
			temp[2] = get_binary_value(img_in[get_one_level_position(width, startx + i + 2, endy)]);
			temp[3] = get_binary_value(img_in[get_one_level_position(width, startx + i + 3, endy)]);
			temp[4] = get_binary_value(img_in[get_one_level_position(width, startx + i + 4, endy)]);
			temp[5] = get_binary_value(img_in[get_one_level_position(width, startx + i + 5, endy)]);
			temp[6] = get_binary_value(img_in[get_one_level_position(width, startx + i + 2, endy - 1)]);
			temp[7] = get_binary_value(img_in[get_one_level_position(width, startx + i + 3, endy - 1)]);
			temp[8] = get_binary_value(img_in[get_one_level_position(width, startx + i + 2, endy - 2)]);
			temp[9] = get_binary_value(img_in[get_one_level_position(width, startx + i + 3, endy - 2)]);
			if (temp[2] > 0 && temp[3] > 0 && temp[6] > 0 && temp[7] > 0) {
				int t = temp[0] + temp[1] + temp[4] + temp[5] + temp[8] + temp[9];
				if (t == 0) {
					m2_s = i;
					m2_e = i + 5;
					break;
				}
			}
		}
		if (m1_s < 0 || m1_e < 0 || m2_s < 0 || m2_e < 0) {
			//ûÓзָî³É¹¦µÄ»°Ìî³ä0
			nps[0] = -1;
			nps[1] = -1;
			nps[2] = -1;
			nps[3] = -1;

			nps[4] = -1;
			nps[5] = -1;
			nps[6] = -1;
			nps[7] = -1;

			nps[8] = -1;
			nps[9] = -1;
			nps[10] = -1;
			nps[11] = -1;
		}
		else {
			//3¸öʱ¼äµÄ×ø±ê
			nps[0] = startx;
			nps[1] = starty;
			nps[2] = startx + m1_s - 1;
			nps[3] = endy;

			nps[4] = startx + m1_e + 1;
			nps[5] = starty;
			nps[6] = startx + m2_s - 1;
			nps[7] = endy;

			nps[8] = startx + m2_e + 1;
			nps[9] = starty;
			nps[10] = endx;
			nps[11] = endy;
		}


	}
	//½ð¶î
	else if (type == 1) {
		//²éÕÒµ¥Î»Ç°µÄ¿Õ¸ñ
		int num_end = -1;
		int i;
		for (i = endx - 11;i >= startx;i--) {
			bool empty = 1;
			for (int y = starty;y < endy;y++) {
				if (get_binary_value(img_in[get_one_level_position(width, i, y)]) > 0) {
					empty = 0;
					break;
				}
			}
			if (empty) {
				num_end = i - 1;
				break;
			}
		}

		if (num_end >= 0) {
			endx = num_end;
		}

		//È¥³ýÉÏϵĿհ×
		int startNy = -1;
		int endNy = -1;
		//È¥³ýÉÏϵİױß
		int y = 0;
		for (y = starty;y <= endy;y++) {
			bool empty = 1;
			for (int x = startx;x <= endx;x++)
			{
				int p = get_one_level_position(width, x, y);
				unsigned char value = get_binary_value(img_in[p]);
				//ÓÐÊý¾Ý
				if (value > 0) {
					empty = 0;
					break;
				}
			}
			if (!empty) {
				startNy = y;
				break;
			}
		}

		for (y = endy;y >= starty;y--) {
			bool empty = 1;
			for (int x = startx;x <= endx;x++)
			{
				int p = get_one_level_position(width, x, y);
				unsigned char value = get_binary_value(img_in[p]);
				//ÓÐÊý¾Ý
				if (value > 0) {
					empty = 0;
					break;
				}
			}
			if (!empty) {
				endNy = y;
				break;
			}
		}
		starty = startNy;
		endy = endNy;



		cols = endx - startx + 1;




		int m_s = -1, m_e = -1;
		uchar temp[8];

		for (i = 0;i < cols;i++) {
			if (i > cols - 5) {
				break;
			}
			temp[0] = get_binary_value(img_in[get_one_level_position(width, startx + i, endy)]);
			temp[1] = get_binary_value(img_in[get_one_level_position(width, startx + i + 1, endy)]);
			temp[2] = get_binary_value(img_in[get_one_level_position(width, startx + i + 2, endy)]);
			temp[3] = get_binary_value(img_in[get_one_level_position(width, startx + i + 3, endy)]);
			temp[4] = get_binary_value(img_in[get_one_level_position(width, startx + i + 4, endy)]);
			temp[5] = get_binary_value(img_in[get_one_level_position(width, startx + i + 5, endy)]);
			temp[6] = get_binary_value(img_in[get_one_level_position(width, startx + i + 2, endy - 1)]);
			temp[7] = get_binary_value(img_in[get_one_level_position(width, startx + i + 3, endy - 1)]);
			if (temp[2] > 0 && temp[3] > 0 && temp[6] > 0 && temp[7] > 0) {

				if (temp[0] + temp[1] + temp[4] + temp[5] == 0)
				{

					bool empty = 1;
					for (int y = starty;y < endy - 1;y++) {
						if (get_binary_value(img_in[get_one_level_position(width, startx + i + 2, y)]) > 0) {
							empty = 0;
							break;
						}
					}
					//СÊýµãÉÏ·½Îª¿ÕÖµ
					if (empty) {
						m_s = i;
						m_e = i + 5;
						break;
					}
				}
			}
		}

		if (m_s == -1 || m_e == -1) {
			nps[0] = -1;
			nps[1] = -1;
			nps[2] = -1;
			nps[3] = -1;

			nps[4] = -1;
			nps[5] = -1;
			nps[6] = -1;
			nps[7] = -1;
		}
		else {
			nps[0] = startx;
			nps[1] = starty;
			nps[2] = startx + m_s - 1;
			nps[3] = endy;

			nps[4] = startx + m_e + 1;
			nps[5] = starty;
			nps[6] = num_end;
			nps[7] = endy;

		}




	}
	//¼Û¸ñ
	else if (type == 2) {
		//¹É¼Û

		//printf("startx:%d starty:%d endx:%d endy:%d\n",startx,starty,endx,endy);

			//ÍùÇ°²éÕÒСÊýµã
		int m_s = -1, m_e = -1;
		uchar temp[8];
		int i;
		for (i = 0;i < cols;i++) {
			if (i > cols - 5) {
				break;
			}
			temp[0] = get_binary_value(img_in[get_one_level_position(width, startx + i, endy)]);
			temp[1] = get_binary_value(img_in[get_one_level_position(width, startx + i + 1, endy)]);
			temp[2] = get_binary_value(img_in[get_one_level_position(width, startx + i + 2, endy)]);
			temp[3] = get_binary_value(img_in[get_one_level_position(width, startx + i + 3, endy)]);
			temp[4] = get_binary_value(img_in[get_one_level_position(width, startx + i + 4, endy)]);
			temp[5] = get_binary_value(img_in[get_one_level_position(width, startx + i + 5, endy)]);
			temp[6] = get_binary_value(img_in[get_one_level_position(width, startx + i + 2, endy - 1)]);
			temp[7] = get_binary_value(img_in[get_one_level_position(width, startx + i + 3, endy - 1)]);
			if (temp[2] > 0 && temp[3] > 0 && temp[6] > 0 && temp[7] > 0) {

				if (temp[0] + temp[1] + temp[4] + temp[5] == 0)
				{

					bool empty = 1;
					for (int y = starty;y < endy - 1;y++) {
						if (get_binary_value(img_in[get_one_level_position(width, startx + i + 2, y)]) > 0) {
							empty = 0;
							break;
						}
					}
					//СÊýµãÉÏ·½Îª¿ÕÖµ
					if (empty) {
						m_s = i;
						m_e = i + 5;
						break;
					}
				}
			}
		}

		if (m_s == -1 || m_e == -1) {
			nps[0] = -1;
			nps[1] = -1;
			nps[2] = -1;
			nps[3] = -1;

			nps[4] = -1;
			nps[5] = -1;
			nps[6] = -1;
			nps[7] = -1;

			nps[8] = -1;
			nps[9] = -1;
			nps[10] = -1;
			nps[11] = -1;
		}
		else {
			nps[0] = startx;
			nps[1] = starty;
			nps[2] = startx + m_s - 1;
			nps[3] = endy;

			nps[4] = startx + m_e + 1;
			nps[5] = starty;
			nps[6] = endx;
			nps[7] = endy;

			nps[8] = -1;
			nps[9] = -1;
			nps[10] = -1;
			nps[11] = -1;
		}

	}
	//´úÂëÓëÊý×Ö
	else if (type == 3 || type == 4) {
		//ÊÖÊý
		nps[0] = startx;
		nps[1] = starty;
		nps[2] = endx;
		nps[3] = endy;

		nps[4] = -1;
		nps[5] = -1;
		nps[6] = -1;
		nps[7] = -1;

		nps[8] = -1;
		nps[9] = -1;
		nps[10] = -1;
		nps[11] = -1;
	}
	//·Ö¸ôÊý¾Ý

	//×î´ó·Ö¸î20¸öÊý
	int numCount = 0;
	//¼Ç¼·Ö¸îµÃµ½µÄÊý¾Ý
	int  numberCols[20 * 2];

	int i;
	for (i = 0;i < 3;i++) {
		if (nps[i * 4] == -1) {
			break;
		}
		int _startx = nps[i * 4];
		int _starty = nps[i * 4 + 1];
		int _endx = nps[i * 4 + 2];
		int _endy = nps[i * 4 + 3];
		int _cols = _endx - _startx + 1;



		int c, startC = -1, endC = -1;
		for (c = 0;c < _cols;c++) {

			bool empty = 1;
			for (int y = _starty;y <= _endy;y++) {
				unsigned char value = get_binary_value(img_in[get_one_level_position(width, _startx + c, y)]);
				if (value > 0) {
					empty = 0;
					break;
				}
			}

			if (empty) {
				if (startC > -1 && endC > -1) {
					numberCols[numCount * 2] = _startx + startC;
					numberCols[numCount * 2 + 1] = _startx + endC;
					numCount++;
				}
				startC = -1;
				endC = -1;
			}
			else {
				if (startC == -1) {
					startC = c;
					endC = c;
				}
				else {
					endC = c;
				}
			}
		}

		if (startC > -1 && endC > -1) {
			numberCols[numCount * 2] = _startx + startC;
			numberCols[numCount * 2 + 1] = _startx + endC;
			numCount++;
		}
	}

	int zeroCount = length_per_num - numCount;
	//×î´óΪ20λÊý
	int fresult[20 * 2];
	for (i = 0;i < zeroCount;i++) {
		fresult[i * 2] = -1;
		fresult[i * 2 + 1] = -1;
	}

	for (i = 0;i < numCount;i++) {
		fresult[(i + zeroCount) * 2] = numberCols[i * 2];
		fresult[(i + zeroCount) * 2 + 1] = numberCols[i * 2 + 1];
	}

	//¿ªÊ¼Ìî³äÊý¾Ý
	for (i = 0;i < length_per_num;i++) {

		unsigned char numData[100 * 100];

		if (fresult[i * 2] == -1) {
			//Ìî³ä0
			for (int r = 0;r < num_height;r++) {
				for (int c = 0;c < num_width;c++) {
					unsigned char value = get_binary_value(zero[r * num_width + c]);
					//ÉèÖÃÊä³ö×ø±êµÄÖµ 
					numData[r * num_width + c] = value;
				}
			}

		}
		else {
			int _startx = fresult[i * 2];
			int _endx = fresult[i * 2 + 1];
			for (int r = starty;r <= endy;r++) {
				for (int c = 0;c < num_width;c++) {
					unsigned char value = 0;
					if (_startx + c <= _endx) {
						//Ìî³ä¿Õ°×0	
						value = get_binary_value(img_in[get_one_level_position(width, _startx + c, r)]);
						//ÉèÖÃÊä³ö×ø±êµÄÖµ 
						numData[(r - starty) * num_width + c] = value;
					}
					else {
						numData[(r - starty) * num_width + c] = 0;
					}
				}
			}
		}

		int rowDataCount = num_height * num_width * 10 * num_count;
		int rowIndex = index / ele_count_per_line;
		int index_0 = rowDataCount * rowIndex;
		//ÐÐÆ«ÒÆÁ¿
		int index_2 = (num_width * 10) * (length_per_num * (index % ele_count_per_line) + i);

		for (int re = 0;re < 10;re++) {
			int index_3 = re * num_width;
			for (int r = 0;r < num_height;r++) {
				int index_1 = r * (num_width * 10 * num_count);
				for (int c = 0;c < num_width;c++) {
					int findex = index_0 + index_1 + index_2 + index_3 + c;
					//printf("index:%d-findex:%d value:%d \n", index, findex, numData[r * num_width + c]);
					result[findex] = numData[r * num_width + c];
				}
			}
		}


	}

	//printf("\n");
}



__kernel void split_l2_num_new(__global const unsigned char* img_in, __global const int* pos_in, __global const unsigned char* zero, __global const unsigned char* types, int width, int num_width, int num_height, int ele_count_per_line, int length_per_num,
	__global unsigned char* result) {

	//printf("ele_count_per_line:%d  length_per_num:%d\n", ele_count_per_line, length_per_num);

	int index = get_global_id(0);
	//ÇøÓò×ø±ê
	int startP = index * 4;
	int startx = pos_in[startP];
	int starty = pos_in[startP + 1];
	int endx = pos_in[startP + 2];
	int endy = pos_in[startP + 3];

	if (startx > 0)
	{
		int startNy = -1;
		int endNy = -1;
		//È¥³ýÉÏϵİױß
		int y = 0;
		for (y = starty;y <= endy;y++) {
			bool empty = 1;
			for (int x = startx;x <= endx;x++)
			{
				int p = get_one_level_position(width, x, y);
				unsigned char value = get_binary_value(img_in[p]);
				//ÓÐÊý¾Ý
				if (value > 0) {
					empty = 0;
					break;
				}
			}
			if (!empty) {
				startNy = y;
				break;
			}
		}

		for (y = endy;y >= starty;y--) {
			bool empty = 1;
			for (int x = startx;x <= endx;x++)
			{
				int p = get_one_level_position(width, x, y);
				unsigned char value = get_binary_value(img_in[p]);
				//ÓÐÊý¾Ý
				if (value > 0) {
					empty = 0;
					break;
				}
			}
			if (!empty) {
				endNy = y;
				break;
			}
		}
		starty = startNy;
		endy = endNy;
	}
	int cols = endx - startx + 1;
	int rows = endy - starty + 1;

	//ÿ¸öÔªËØÕ¼ÓÃ6¸öÊý×ֵĿռ䣬²»¹»µÄÇ°ÃæÌî³ä0
	int num_count = ele_count_per_line * length_per_num;

	int nps[12];
	//Ö»ÄÜʶ±ð4ÖÖÊý×Ö
	unsigned char type = types[index % ele_count_per_line];
	//printf("ele_count_line:%d\n", ele_count_line);
	//ʱ¼ä
	if (type == 0) {
		//ʱ¼ä
		int s = cols / 2;
		int i;
		//ÍùÇ°²éÕÒðºÅ
		int m1_s = -1, m1_e = -1, m2_s = -1, m2_e = -1;
		uchar temp[7];
		for (i = s;i > -1;i--) {
			if (i < 6) {
				break;
			}
			temp[0] = get_binary_value(img_in[get_one_level_position(width, startx + i, endy)]);

			temp[1] = get_binary_value(img_in[get_one_level_position(width, startx + i - 1, endy)]);


			temp[2] = get_binary_value(img_in[get_one_level_position(width, startx + i - 2, endy)]);
			temp[3] = get_binary_value(img_in[get_one_level_position(width, startx + i - 3, endy)]);
			temp[4] = get_binary_value(img_in[get_one_level_position(width, startx + i - 4, endy)]);
			temp[5] = get_binary_value(img_in[get_one_level_position(width, startx + i - 5, endy)]);
			temp[6] = get_binary_value(img_in[get_one_level_position(width, startx + i - 6, endy)]);
			if (temp[3] > 0 && get_binary_value(img_in[get_one_level_position(width, startx + i - 3, endy - 1)]) == 0) {
				int t = temp[0] + temp[1] + temp[2] + temp[4] + temp[5] + temp[6];
				if (t == 0) {
					m1_s = i - 6;
					m1_e = i;
					break;
				}
			}

		}

		//Íùºó²éÕÒðºÅ

		for (i = s;i < cols;i++) {
			if (i > cols - 6) {
				break;
			}

			temp[0] = get_binary_value(img_in[get_one_level_position(width, startx + i, endy)]);
			temp[1] = get_binary_value(img_in[get_one_level_position(width, startx + i + 1, endy)]);
			temp[2] = get_binary_value(img_in[get_one_level_position(width, startx + i + 2, endy)]);
			temp[3] = get_binary_value(img_in[get_one_level_position(width, startx + i + 3, endy)]);
			temp[4] = get_binary_value(img_in[get_one_level_position(width, startx + i + 4, endy)]);
			temp[5] = get_binary_value(img_in[get_one_level_position(width, startx + i + 5, endy)]);
			temp[6] = get_binary_value(img_in[get_one_level_position(width, startx + i + 6, endy)]);
			if (temp[3] > 0 && get_binary_value(img_in[get_one_level_position(width, startx + i + 3, endy - 1)]) == 0) {
				int t = temp[0] + temp[1] + temp[2] + temp[4] + temp[5] + temp[6];
				if (t == 0) {
					m2_s = i;
					m2_e = i + 6;
					break;
				}
			}
		}

		if (m1_s < 0 || m1_e < 0 || m2_s < 0 || m2_e < 0) {
			//ûÓзָî³É¹¦µÄ»°Ìî³ä0
			nps[0] = -1;
			nps[1] = -1;
			nps[2] = -1;
			nps[3] = -1;

			nps[4] = -1;
			nps[5] = -1;
			nps[6] = -1;
			nps[7] = -1;

			nps[8] = -1;
			nps[9] = -1;
			nps[10] = -1;
			nps[11] = -1;
		}
		else {
			//3¸öʱ¼äµÄ×ø±ê
			nps[0] = startx;
			nps[1] = starty;
			nps[2] = startx + m1_s - 1;
			nps[3] = endy;

			nps[4] = startx + m1_e + 1;
			nps[5] = starty;
			nps[6] = startx + m2_s - 1;
			nps[7] = endy;

			nps[8] = startx + m2_e + 1;
			nps[9] = starty;
			nps[10] = endx;
			nps[11] = endy;
		}


	}
	//½ð¶î
	else if (type == 1) {
		//²éÕÒµ¥Î»Ç°µÄ¿Õ¸ñ
		int num_end = -1;
		int i;
		for (i = endx - 9;i >= startx;i--) {
			bool empty = 1;
			for (int y = starty;y < endy;y++) {
				if (get_binary_value(img_in[get_one_level_position(width, i, y)]) > 0) {
					empty = 0;
					break;
				}
			}
			if (empty) {
				num_end = i - 1;
				break;
			}
		}

		if (num_end >= 0) {
			endx = num_end;
		}

		//È¥³ýÉÏϵĿհ×
		int startNy = -1;
		int endNy = -1;
		//È¥³ýÉÏϵİױß
		int y = 0;
		for (y = starty;y <= endy;y++) {
			bool empty = 1;
			for (int x = startx;x <= endx;x++)
			{
				int p = get_one_level_position(width, x, y);
				unsigned char value = get_binary_value(img_in[p]);
				//ÓÐÊý¾Ý
				if (value > 0) {
					empty = 0;
					break;
				}
			}
			if (!empty) {
				startNy = y;
				break;
			}
		}

		for (y = endy;y >= starty;y--) {
			bool empty = 1;
			for (int x = startx;x <= endx;x++)
			{
				int p = get_one_level_position(width, x, y);
				unsigned char value = get_binary_value(img_in[p]);
				//ÓÐÊý¾Ý
				if (value > 0) {
					empty = 0;
					break;
				}
			}
			if (!empty) {
				endNy = y;
				break;
			}
		}
		starty = startNy;
		endy = endNy;



		cols = endx - startx + 1;



		//²éÕÒСÊýµã
		int m_s = -1, m_e = -1;
		uchar temp[6];
		for (i = 0;i < cols;i++) {
			if (i > cols - 5) {
				break;
			}
			temp[0] = get_binary_value(img_in[get_one_level_position(width, startx + i, endy)]);

			temp[1] = get_binary_value(img_in[get_one_level_position(width, startx + i + 1, endy)]);
			temp[2] = get_binary_value(img_in[get_one_level_position(width, startx + i + 2, endy)]);
			temp[3] = get_binary_value(img_in[get_one_level_position(width, startx + i + 3, endy)]);
			temp[4] = get_binary_value(img_in[get_one_level_position(width, startx + i + 4, endy)]);
			temp[5] = get_binary_value(img_in[get_one_level_position(width, startx + i + 5, endy)]);
			if (temp[2] > 0) {

				bool empty = 1;
				for (int y = starty;y < endy;y++) {
					if (get_binary_value(img_in[get_one_level_position(width, startx + i + 2, y)]) > 0) {
						empty = 0;
						break;
					}
				}
				//СÊýµãÉÏ·½Îª¿ÕÖµ
				if (empty) {
					unsigned char	t = temp[0] + temp[1] + temp[3] + temp[4] + temp[5];
					if (t == 0) {
						m_s = i;
						m_e = i + 5;
						break;
					}
				}
			}
		}

		if (m_s == -1 || m_e == -1) {
			nps[0] = -1;
			nps[1] = -1;
			nps[2] = -1;
			nps[3] = -1;

			nps[4] = -1;
			nps[5] = -1;
			nps[6] = -1;
			nps[7] = -1;
		}
		else {
			nps[0] = startx;
			nps[1] = starty;
			nps[2] = startx + m_s - 1;
			nps[3] = endy;

			nps[4] = startx + m_e + 1;
			nps[5] = starty;
			nps[6] = num_end;
			nps[7] = endy;

		}




	}
	//¼Û¸ñ
	else if (type == 2) {
		//¹É¼Û

		//printf("startx:%d starty:%d endx:%d endy:%d\n",startx,starty,endx,endy);

		int m_s = -1, m_e = -1;
		uchar temp[6];
		int i;
		for (i = 0;i < cols;i++) {
			if (i > cols - 5) {
				break;
			}
			temp[0] = get_binary_value(img_in[get_one_level_position(width, startx + i, endy)]);

			temp[1] = get_binary_value(img_in[get_one_level_position(width, startx + i + 1, endy)]);
			temp[2] = get_binary_value(img_in[get_one_level_position(width, startx + i + 2, endy)]);
			temp[3] = get_binary_value(img_in[get_one_level_position(width, startx + i + 3, endy)]);
			temp[4] = get_binary_value(img_in[get_one_level_position(width, startx + i + 4, endy)]);
			temp[5] = get_binary_value(img_in[get_one_level_position(width, startx + i + 5, endy)]);
			if (temp[2] > 0) {

				bool empty = 1;
				for (int y = starty;y < endy;y++) {
					if (get_binary_value(img_in[get_one_level_position(width, startx + i + 2, y)]) > 0) {
						empty = 0;
						break;
					}
				}
				//СÊýµãÉÏ·½Îª¿ÕÖµ
				if (empty) {
					unsigned char	t = temp[0] + temp[1] + temp[3] + temp[4] + temp[5];
					if (t == 0) {
						m_s = i;
						m_e = i + 5;
						break;
					}
				}
			}
		}

		if (m_s == -1 || m_e == -1) {
			nps[0] = -1;
			nps[1] = -1;
			nps[2] = -1;
			nps[3] = -1;

			nps[4] = -1;
			nps[5] = -1;
			nps[6] = -1;
			nps[7] = -1;

			nps[8] = -1;
			nps[9] = -1;
			nps[10] = -1;
			nps[11] = -1;
		}
		else {
			nps[0] = startx;
			nps[1] = starty;
			nps[2] = startx + m_s - 1;
			nps[3] = endy;

			nps[4] = startx + m_e + 1;
			nps[5] = starty;
			nps[6] = endx;
			nps[7] = endy;

			nps[8] = -1;
			nps[9] = -1;
			nps[10] = -1;
			nps[11] = -1;
		}

	}
	//´úÂëÓëÊý×Ö
	else if (type == 3 || type == 4) {
		//ÊÖÊý
		nps[0] = startx;
		nps[1] = starty;
		nps[2] = endx;
		nps[3] = endy;

		nps[4] = -1;
		nps[5] = -1;
		nps[6] = -1;
		nps[7] = -1;

		nps[8] = -1;
		nps[9] = -1;
		nps[10] = -1;
		nps[11] = -1;
	}
	//·Ö¸ôÊý¾Ý

	//×î´ó·Ö¸î20¸öÊý
	int numCount = 0;
	//¼Ç¼·Ö¸îµÃµ½µÄÊý¾Ý
	int  numberCols[20 * 2];

	int i;
	for (i = 0;i < 3;i++) {
		if (nps[i * 4] == -1) {
			break;
		}
		int _startx = nps[i * 4];
		int _starty = nps[i * 4 + 1];
		int _endx = nps[i * 4 + 2];
		int _endy = nps[i * 4 + 3];
		int _cols = _endx - _startx + 1;



		int c, startC = -1, endC = -1;
		for (c = 0;c < _cols;c++) {

			bool empty = 1;
			for (int y = _starty;y <= _endy;y++) {
				unsigned char value = get_binary_value(img_in[get_one_level_position(width, _startx + c, y)]);
				if (value > 0) {
					empty = 0;
					break;
				}
			}

			if (empty) {
				if (startC > -1 && endC > -1) {
					numberCols[numCount * 2] = _startx + startC;
					numberCols[numCount * 2 + 1] = _startx + endC;
					numCount++;
				}
				startC = -1;
				endC = -1;
			}
			else {
				if (startC == -1) {
					startC = c;
					endC = c;
				}
				else {
					endC = c;
				}
			}
		}

		if (startC > -1 && endC > -1) {
			numberCols[numCount * 2] = _startx + startC;
			numberCols[numCount * 2 + 1] = _startx + endC;
			numCount++;
		}
	}

	int zeroCount = length_per_num - numCount;
	//×î´óΪ20λÊý
	int fresult[20 * 2];
	for (i = 0;i < zeroCount;i++) {
		fresult[i * 2] = -1;
		fresult[i * 2 + 1] = -1;
	}

	for (i = 0;i < numCount;i++) {
		fresult[(i + zeroCount) * 2] = numberCols[i * 2];
		fresult[(i + zeroCount) * 2 + 1] = numberCols[i * 2 + 1];
	}

	//¿ªÊ¼Ìî³äÊý¾Ý
	for (i = 0;i < length_per_num;i++) {

		unsigned char numData[100 * 100];

		if (fresult[i * 2] == -1) {
			//Ìî³ä0
			for (int r = 0;r < num_height;r++) {
				for (int c = 0;c < num_width;c++) {
					unsigned char value = get_binary_value(zero[r * num_width + c]);
					//ÉèÖÃÊä³ö×ø±êµÄÖµ 
					numData[r * num_width + c] = value;
				}
			}

		}
		else {
			int _startx = fresult[i * 2];
			int _endx = fresult[i * 2 + 1];
			for (int r = starty;r <= endy;r++) {
				for (int c = 0;c < num_width;c++) {
					unsigned char value = 0;
					if (_startx + c <= _endx) {
						//Ìî³ä¿Õ°×0	
						value = get_binary_value(img_in[get_one_level_position(width, _startx + c, r)]);
						//ÉèÖÃÊä³ö×ø±êµÄÖµ 
						numData[(r - starty) * num_width + c] = value;
					}
					else {
						numData[(r - starty) * num_width + c] = 0;
					}
				}
			}
		}

		int rowDataCount = num_height * num_width * 10 * num_count;
		int rowIndex = index / ele_count_per_line;
		int index_0 = rowDataCount * rowIndex;
		//ÐÐÆ«ÒÆÁ¿
		int index_2 = (num_width * 10) * (length_per_num * (index % ele_count_per_line) + i);

		for (int re = 0;re < 10;re++) {
			int index_3 = re * num_width;
			for (int r = 0;r < num_height;r++) {
				int index_1 = r * (num_width * 10 * num_count);
				for (int c = 0;c < num_width;c++) {
					int findex = index_0 + index_1 + index_2 + index_3 + c;
					//printf("index:%d-findex:%d value:%d \n", index, findex, numData[r * num_width + c]);
					result[findex] = numData[r * num_width + c];
				}
			}
		}


	}

	//printf("\n");
}




//·Ö¸îL2Êý×Ö

__kernel void split_l2_num(__global const unsigned char* img_in, __global const int* pos_in, __global const unsigned char* zero, int width, int num_width, int num_height, int num_count,
	__global unsigned char* result) {

	int index = get_global_id(0);
	//ÇøÓò×ø±ê
	int startP = index * 4;
	int startx = pos_in[startP];
	int starty = pos_in[startP + 1];
	int endx = pos_in[startP + 2];
	int endy = pos_in[startP + 3];

	if (startx > 0)
	{
		int startNy = -1;
		int endNy = -1;
		//È¥³ýÉÏϵİױß
		int y = 0;
		for (y = starty;y <= endy;y++) {
			bool empty = 1;
			for (int x = startx;x <= endx;x++)
			{
				int p = get_one_level_position(width, x, y);
				unsigned char value = get_binary_value(img_in[p]);
				//ÓÐÊý¾Ý
				if (value > 0) {
					empty = 0;
					break;
				}
			}
			if (!empty) {
				startNy = y;
				break;
			}
		}

		for (y = endy;y >= starty;y--) {
			bool empty = 1;
			for (int x = startx;x <= endx;x++)
			{
				int p = get_one_level_position(width, x, y);
				unsigned char value = get_binary_value(img_in[p]);
				//ÓÐÊý¾Ý
				if (value > 0) {
					empty = 0;
					break;
				}
			}
			if (!empty) {
				endNy = y;
				break;
			}
		}
		starty = startNy;
		endy = endNy;
	}
	int cols = endx - startx + 1;
	int rows = endy - starty + 1;

	int nps[12];
	if (index % 4 == 0) {
		//ʱ¼ä
		int s = cols / 2;
		int i;
		//ÍùÇ°²éÕÒðºÅ
		int m1_s = -1, m1_e = -1, m2_s = -1, m2_e = -1;
		uchar temp[7];
		for (i = s;i > -1;i--) {
			if (i < 6) {
				break;
			}
			temp[0] = get_binary_value(img_in[get_one_level_position(width, startx + i, endy)]);

			temp[1] = get_binary_value(img_in[get_one_level_position(width, startx + i - 1, endy)]);


			temp[2] = get_binary_value(img_in[get_one_level_position(width, startx + i - 2, endy)]);
			temp[3] = get_binary_value(img_in[get_one_level_position(width, startx + i - 3, endy)]);
			temp[4] = get_binary_value(img_in[get_one_level_position(width, startx + i - 4, endy)]);
			temp[5] = get_binary_value(img_in[get_one_level_position(width, startx + i - 5, endy)]);
			temp[6] = get_binary_value(img_in[get_one_level_position(width, startx + i - 6, endy)]);
			if (temp[3] > 0 && get_binary_value(img_in[get_one_level_position(width, startx + i - 3, endy - 1)]) == 0) {
				int t = temp[0] + temp[1] + temp[2] + temp[4] + temp[5] + temp[6];
				if (t == 0) {
					m1_s = i - 6;
					m1_e = i;
					break;
				}
			}

		}

		//Íùºó²éÕÒðºÅ

		for (i = s;i < cols;i++) {
			if (i > cols - 6) {
				break;
			}

			temp[0] = get_binary_value(img_in[get_one_level_position(width, startx + i, endy)]);
			temp[1] = get_binary_value(img_in[get_one_level_position(width, startx + i + 1, endy)]);
			temp[2] = get_binary_value(img_in[get_one_level_position(width, startx + i + 2, endy)]);
			temp[3] = get_binary_value(img_in[get_one_level_position(width, startx + i + 3, endy)]);
			temp[4] = get_binary_value(img_in[get_one_level_position(width, startx + i + 4, endy)]);
			temp[5] = get_binary_value(img_in[get_one_level_position(width, startx + i + 5, endy)]);
			temp[6] = get_binary_value(img_in[get_one_level_position(width, startx + i + 6, endy)]);
			if (temp[3] > 0 && get_binary_value(img_in[get_one_level_position(width, startx + i + 3, endy - 1)]) == 0) {
				int t = temp[0] + temp[1] + temp[2] + temp[4] + temp[5] + temp[6];
				if (t == 0) {
					m2_s = i;
					m2_e = i + 6;
					break;
				}
			}
		}
		if (m1_s < 0 || m1_e < 0 || m2_s < 0 || m2_e < 0) {
			//printf("δ²éÕÒµ½Ã°ºÅ£ºm1_s:%d m1_end:%d m2_s:%d m2_e:%d\n", m1_s, m1_e, m2_s, m2_e);
			return;
		}

		//3¸öʱ¼äµÄ×ø±ê
		nps[0] = startx;
		nps[1] = starty;
		nps[2] = startx + m1_s - 1;
		nps[3] = endy;

		nps[4] = startx + m1_e + 1;
		nps[5] = starty;
		nps[6] = startx + m2_s - 1;
		nps[7] = endy;

		nps[8] = startx + m2_e + 1;
		nps[9] = starty;
		nps[10] = endx;
		nps[11] = endy;

	}
	else if (index % 4 == 1) {
		if (startx > 0 && endx > 0) {
			//·Ö¸ôs/m/h
			int emptyX = -1;
			for (int x = endx;x >= startx;x--) {

				bool empty = 1;
				for (int y = starty;y <= endy;y++)
				{
					int p = get_one_level_position(width, x, y);
					unsigned char value = get_binary_value(img_in[p]);
					//ÓÐÊý¾Ý
					if (value > 0) {
						empty = 0;
						break;
					}
				}
				if (empty) {
					emptyX = x;
					break;
				}
			}

			if (emptyX > 0) {
				nps[0] = startx;
				nps[1] = starty;
				nps[2] = emptyX - 1;
				nps[3] = endy;
			}
			else {
				//printf("³·µ¥Ê±¼äδ½âÎöµ½·Ö¸ô·û£º%d", index);
				nps[0] = -1;
				nps[1] = -1;
				nps[2] = -1;
				nps[3] = -1;
			}
		}
		else {
			nps[0] = -1;
			nps[1] = -1;
			nps[2] = -1;
			nps[3] = -1;
		}

		nps[4] = -1;
		nps[5] = -1;
		nps[6] = -1;
		nps[7] = -1;

		nps[8] = -1;
		nps[9] = -1;
		nps[10] = -1;
		nps[11] = -1;


	}

	else if (index % 4 == 2) {
		//¹É¼Û

		//printf("startx:%d starty:%d endx:%d endy:%d\n",startx,starty,endx,endy);

			//ÍùÇ°²éÕÒСÊýµã
		int m_s = -1, m_e = -1;
		uchar temp[6];
		int i;
		for (i = 0;i < cols;i++) {
			if (i > cols - 5) {
				break;
			}
			temp[0] = get_binary_value(img_in[get_one_level_position(width, startx + i, endy)]);

			temp[1] = get_binary_value(img_in[get_one_level_position(width, startx + i + 1, endy)]);
			temp[2] = get_binary_value(img_in[get_one_level_position(width, startx + i + 2, endy)]);
			temp[3] = get_binary_value(img_in[get_one_level_position(width, startx + i + 3, endy)]);
			temp[4] = get_binary_value(img_in[get_one_level_position(width, startx + i + 4, endy)]);
			temp[5] = get_binary_value(img_in[get_one_level_position(width, startx + i + 5, endy)]);
			if (temp[2] > 0) {

				bool empty = 1;
				for (int y = starty;y < endy;y++) {
					if (get_binary_value(img_in[get_one_level_position(width, startx + i + 2, y)]) > 0) {
						empty = 0;
						break;
					}
				}
				//СÊýµãÉÏ·½Îª¿ÕÖµ
				if (empty) {
					unsigned char	t = temp[0] + temp[1] + temp[3] + temp[4] + temp[5];
					if (t == 0) {
						m_s = i;
						m_e = i + 5;
						break;
					}
				}
			}
		}

		if (m_s == -1 || m_e == -1) {
			//throw("δ²éÕÒµ½Ð¡Êýµã");
			//printf("δ²éÕÒµ½Ð¡Êýµã m_s:%d m_e:%d\n", m_s,m_e);
			return;
		}


		nps[0] = startx;
		nps[1] = starty;
		nps[2] = startx + m_s - 1;
		nps[3] = endy;

		nps[4] = startx + m_e + 1;
		nps[5] = starty;
		nps[6] = endx;
		nps[7] = endy;

		nps[8] = -1;
		nps[9] = -1;
		nps[10] = -1;
		nps[11] = -1;

	}

	else if (index % 4 == 3) {
		//ÊÖÊý
		nps[0] = startx;
		nps[1] = starty;
		nps[2] = endx;
		nps[3] = endy;

		nps[4] = -1;
		nps[5] = -1;
		nps[6] = -1;
		nps[7] = -1;

		nps[8] = -1;
		nps[9] = -1;
		nps[10] = -1;
		nps[11] = -1;
	}
	//·Ö¸ôÊý¾Ý

	//×î´ó·Ö¸î6¸öÊý
	int numCount = 0;
	//¼Ç¼·Ö¸îµÃµ½µÄÊý¾Ý
	int  numberCols[12];

	int i;
	for (i = 0;i < 3;i++) {
		if (nps[i * 4] == -1) {
			break;
		}
		int _startx = nps[i * 4];
		int _starty = nps[i * 4 + 1];
		int _endx = nps[i * 4 + 2];
		int _endy = nps[i * 4 + 3];
		int _cols = _endx - _startx + 1;



		int c, startC = -1, endC = -1;
		for (c = 0;c < _cols;c++) {

			bool empty = 1;
			for (int y = _starty;y <= _endy;y++) {
				unsigned char value = get_binary_value(img_in[get_one_level_position(width, _startx + c, y)]);
				if (value > 0) {
					empty = 0;
					break;
				}
			}

			if (empty) {
				if (startC > -1 && endC > -1) {
					numberCols[numCount * 2] = _startx + startC;
					numberCols[numCount * 2 + 1] = _startx + endC;
					numCount++;
				}
				startC = -1;
				endC = -1;
			}
			else {
				if (startC == -1) {
					startC = c;
					endC = c;
				}
				else {
					endC = c;
				}
			}
		}

		if (startC > -1 && endC > -1) {
			numberCols[numCount * 2] = _startx + startC;
			numberCols[numCount * 2 + 1] = _startx + endC;
			numCount++;
		}



	}

	int maxNumberCount = 6;
	if (index % 4 == 3)
	{
		maxNumberCount = 5;
	}
	else if (index % 4 == 1) {
		maxNumberCount = 2;
	}

	int zeroCount = maxNumberCount - numCount;

	int fresult[6 * 2];
	for (i = 0;i < zeroCount;i++) {
		fresult[i * 2] = -1;
		fresult[i * 2 + 1] = -1;
	}

	for (i = 0;i < numCount;i++) {
		fresult[(i + zeroCount) * 2] = numberCols[i * 2];
		fresult[(i + zeroCount) * 2 + 1] = numberCols[i * 2 + 1];
	}

	//¿ªÊ¼Ìî³äÊý¾Ý
	for (i = 0;i < maxNumberCount;i++) {

		unsigned char numData[100 * 100];

		if (fresult[i * 2] == -1) {
			//Ìî³ä0
			for (int r = 0;r < num_height;r++) {
				for (int c = 0;c < num_width;c++) {
					unsigned char value = get_binary_value(zero[r * num_width + c]);
					//ÉèÖÃÊä³ö×ø±êµÄÖµ 
					numData[r * num_width + c] = value;
				}
			}

		}
		else {
			int _startx = fresult[i * 2];
			int _endx = fresult[i * 2 + 1];
			for (int r = starty;r <= endy;r++) {
				for (int c = 0;c < num_width;c++) {
					unsigned char value = 0;
					if (_startx + c <= _endx) {
						//Ìî³ä¿Õ°×0	
						value = get_binary_value(img_in[get_one_level_position(width, _startx + c, r)]);
						//ÉèÖÃÊä³ö×ø±êµÄÖµ 
						numData[(r - starty) * num_width + c] = value;
					}
					else {
						numData[(r - starty) * num_width + c] = 0;
					}
				}
			}
		}

		int rowData = num_height * num_width * 10 * num_count;
		int rowIndex = index / 4;
		int index_0 = rowData * rowIndex;
		//ÉèÖÃ×ø±êÖµ
		int index_2 = 0;
		if (index % 4 == 0) {
			index_2 += num_width * 10 * i;
		}
		else if (index % 4 == 1) {
			index_2 += (num_width * 10) * (6 + i);
		}
		else if (index % 4 == 2) {
			index_2 += (num_width * 10) * (6 + 2 + i);
		}
		else {
			index_2 += (num_width * 10) * (6 + 2 + 6 + i);
		}

		for (int re = 0;re < 10;re++) {
			int index_3 = re * num_width;
			for (int r = 0;r < num_height;r++) {
				int index_1 = r * (num_width * 10 * num_count);
				for (int c = 0;c < num_width;c++) {
					int findex = index_0 + index_1 + index_2 + index_3 + c;
					//printf("index:%d-findex:%d \n",index, findex);
					result[findex] = numData[r * num_width + c];
				}
			}
		}


	}

	//printf("\n");
}