__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);
		//ÊÇ·ñÖ»ÓÐһλСÊý
	bool one_point_num = false;
	//ʱ¼ä
	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;
		//»ñÈ¡ÊÇ·ñÓе¥Î»
		int last_empty_index = -1;
		for (i = endx; 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) {
				last_empty_index = i;
				break;
			}
		}
		bool hasUnit = false;
		if (endx - last_empty_index > num_width) {
			hasUnit = true;
		}

		//´øÓе¥Î»
		if (hasUnit) {
			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) {
			//printf("ûÕÒµ½Ð¡Êýµã£º%d", index);
			if (hasUnit)
			{
				nps[0] = -1;
				nps[1] = -1;
				nps[2] = -1;
				nps[3] = -1;
			}
			else {
				nps[0] = startx;
				nps[1] = starty;
				nps[2] = endx;
				nps[3] = endy;
			}

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


		}
		else {
			//printf("ÕÒµ½ÁËСÊýµã£º%d", index);
			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;
			if (nps[6] - nps[4] + 1 <= num_width) {
				one_point_num = true;
			}
		}

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




	}
	//¼Û¸ñ
	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;
		}
		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;
	}
	//°Ù·Ö±È
	else if (type == 5) {

		//°Ù·Ö±ÈÊý×ÖÆðʼλÖÃ
		int percent_start = -1;
		int emptyIndex = -1;
		int i;
		for (i = startx; i <= endx; 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) {
				emptyIndex = i;
			}
			else if (emptyIndex > 0) {
				percent_start = i;
				break;
			}

		}


		//°Ù·Ö±ÈÊý×ֵĽáÊøλÖÃ
		int percent_end = -1;
		emptyIndex = -1;

		for (i = endx; 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) {
				emptyIndex = i;
			}
			else if (emptyIndex > 0) {
				percent_end = i;
				break;
			}
		}



		if (percent_start > -1) {
			startx = percent_start;
		}

		if (percent_end > -1) {
			endx = percent_end;
		}


		//°´ÕÕ¼Û¸ñµÄÐÎʽ×öʶ±ð
		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] = endx;
			nps[7] = endy;
		}

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

	}

	//ÊÖÊý£¬´øÓе¥Î»ÊÖ
	else if (type == 6) {
		//·Ö¸ôÊÖ
		int num_end = endx;
		int emptyIndex=-1;
		int i;
		for (i = endx; 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) {
				emptyIndex = i;
			}
			else if (emptyIndex > 0) {
				num_end = i;
				break;
			}
		}
		endx = num_end;
		//È¥³ýÉÏÏ¿հ×
		int startNy = -1;
		int endNy = -1;
		int y;
		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;
			}
		}
		if (endNy > 0)
		{
			endy = endNy;
		}
		if (startNy > 0)
		{
			starty = startNy;
		}
		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 - (one_point_num ? 1 : 0);
	//×î´óΪ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];
	}
	//һλСÊýÐèÒªÔÚĩβÌî³ä0
	if (type == 1 && one_point_num) {
		fresult[(numCount + zeroCount) * 2] = -1;
		fresult[(numCount + zeroCount) * 2 + 1] = -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);
	//ÊÇ·ñÖ»ÓÐһλСÊý
	bool one_point_num = false;
	//ʱ¼ä

	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;
			if (nps[6] - nps[4] <= num_width) {
				//Ö»ÓÐÒ»¸öÊý×Ö
				one_point_num = true;
			}

		}

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


	}
	//¼Û¸ñ
	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 - (one_point_num ? 1 : 0);
	//×î´óΪ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];
	}

	//TODO Èç¹ûÊǽð¶î£¬ÇÒºóÃæÖ»ÓÐ1λСÊýµÄ£¬ÐèÒªÔÚĩβÌî³ä0
	if (type == 1 && one_point_num) {
		fresult[(numCount + zeroCount) * 2] = -1;
		fresult[(numCount + zeroCount) * 2 + 1] = -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_trade_queue_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);
		//ÊÇ·ñÖ»ÓÐһλСÊý
	bool one_point_num = false;
	//ʱ¼ä
	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;
		//»ñÈ¡ÊÇ·ñÓе¥Î»
		int last_empty_index = -1;
		for (i = endx; 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) {
				last_empty_index = i;
				break;
			}
		}
		bool hasUnit = false;
		if (endx - last_empty_index > num_width) {
			hasUnit = true;
		}

		//´øÓе¥Î»
		if (hasUnit) {
			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) {
			//printf("ûÕÒµ½Ð¡Êýµã£º%d", index);
			if (hasUnit)
			{
				nps[0] = -1;
				nps[1] = -1;
				nps[2] = -1;
				nps[3] = -1;
			}
			else {
				nps[0] = startx;
				nps[1] = starty;
				nps[2] = endx;
				nps[3] = endy;
			}

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


		}
		else {
			//printf("ÕÒµ½ÁËСÊýµã£º%d", index);
			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;
			if (nps[6] - nps[4] + 1 <= num_width) {
				one_point_num = true;
			}
		}

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




	}
	//¼Û¸ñ
	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) {

				if (temp[0] + temp[1] + temp[4] + temp[5] + temp[6] + temp[7] == 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] = 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;
	}
	//°Ù·Ö±È
	else if (type == 5) {

		//°Ù·Ö±ÈÊý×ÖÆðʼλÖÃ
		int percent_start = -1;
		int emptyIndex = -1;
		int i;
		for (i = startx; i <= endx; 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) {
				emptyIndex = i;
			}
			else if (emptyIndex > 0) {
				percent_start = i;
				break;
			}

		}


		//°Ù·Ö±ÈÊý×ֵĽáÊøλÖÃ
		int percent_end = -1;
		emptyIndex = -1;

		for (i = endx; 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) {
				emptyIndex = i;
			}
			else if (emptyIndex > 0) {
				percent_end = i;
				break;
			}
		}



		if (percent_start > -1) {
			startx = percent_start;
		}

		if (percent_end > -1) {
			endx = percent_end;
		}


		//°´ÕÕ¼Û¸ñµÄÐÎʽ×öʶ±ð
		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] = endx;
			nps[7] = endy;
		}

		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 - (one_point_num ? 1 : 0);
	//×î´óΪ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];
	}
	//һλСÊýÐèÒªÔÚĩβÌî³ä0
	if (type == 1 && one_point_num) {
		fresult[(numCount + zeroCount) * 2] = -1;
		fresult[(numCount + zeroCount) * 2 + 1] = -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;
				//È¥³ýÉÏϵĿհ×,·ÀÖ¹hµ¥Î»³öÏÖÎÊÌâ
				//È¥³ýÉÏϵİױß
				int y = 0;
				for (y = nps[1]; y <= nps[3]; y++) {
					bool empty = 1;
					for (int x = nps[0]; x <= nps[2]; 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) {
						nps[1] = y;
						starty = y;
						break;
					}
				}

				for (y = nps[3]; y >= nps[1]; y--) {
					bool empty = 1;
					for (int x = nps[0]; x <= nps[2]; 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) {
						nps[3] = y;
						endy = y;
						break;
					}
				}
			}
			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");
}



__kernel void createTemplateNumbers(__global unsigned char* numbers, int num_width, int num_height, int countPerLine,
	__global unsigned char* result) {

	int gid = get_global_id(0);
	int LINE_NUMBER_COUNT = countPerLine;
	int NUMBER_COUNT = 10;
	//unsigned char* data = (unsigned char*)malloc(sizeof(unsigned char) * (ROWS * lines) *COLS * NUMBER_COUNT * LINE_NUMBER_COUNT);
	int outLineDataCount = num_height * num_width * NUMBER_COUNT * LINE_NUMBER_COUNT;
	int inLineDataCount = num_width * NUMBER_COUNT * LINE_NUMBER_COUNT;

	int outLineCount = outLineDataCount * gid;
	for (int re = 0; re < LINE_NUMBER_COUNT; re++) {
		for (int n = 0; n < NUMBER_COUNT; n++)
		{
			int index_num_0 = n * num_width * num_height;
			for (int r = 0; r < num_height; r++) {
				int intLineCount = inLineDataCount * r;

				for (int c = 0; c < num_width; c++) {
					int index_num = index_num_0 + r * num_width + c;
					unsigned char value = numbers[index_num];
					int index = outLineCount;
					index += intLineCount;
					int x = re * NUMBER_COUNT * num_width + n * num_width + c;
					index += x;
					result[index] = get_binary_value(value);
				}
			}
		}
	}

}


__kernel void splitL2RowData(__global unsigned char* imgs, __global int* rowIndexs, int width,
	__global int* result) {
	int gid = get_global_id(0);

	int startCol = rowIndexs[gid * 4];
	int startRow = rowIndexs[gid * 4 + 1];
	int endCol = rowIndexs[gid * 4 + 2];
	int endRow = rowIndexs[gid * 4 + 3];

	//printf(" index:%d - startCol:%d,startRow:%d,endCol:%d,endRow:%d\n", gid, startCol, startRow, endCol, endRow);


	int emptyStart = -1;
	int emptyEnd = -1;
	int dataStart = -1;
	int dataEnd = -1;

	int rowDataSize = 0;
	int rowDataIndexs[7 * 4];

	int startRowDataIndex = 7 * 4 * gid;

	for (int i = startCol; i <= endCol; i++) {

		bool empty = 1;
		for (int j = startRow; j <= endRow; j++) {
			unsigned char value = imgs[get_one_level_position(width, i, j)];
			if (value >= 64) {
				empty = 0;
				break;
			}
		}
		if (empty) {
			if (emptyStart < 0) {
				emptyStart = i;
				emptyEnd = i;
			}
			else {
				emptyEnd = i;
			}

			int SPACE = 5;
            //ʱ¼äµÄ¿í¶ÈΪ45-47 
			if (rowDataSize == 0 && (dataEnd - dataStart + 1 >= 45)) {
				SPACE = 1;
			}

			//3¸ö¼°ÒÔÉϵĿհ×Êý¾Ý²Å·ÖÁÐ
			if (emptyEnd - emptyStart  >= SPACE && dataEnd - dataStart > 0) {

				rowDataIndexs[rowDataSize * 4 + 0] = dataStart;
				rowDataIndexs[rowDataSize * 4 + 1] = startRow;
				rowDataIndexs[rowDataSize * 4 + 2] = dataEnd;
				rowDataIndexs[rowDataSize * 4 + 3] = endRow;
				rowDataSize++;

				dataEnd = -1;
				dataStart = -1;
			}
		}
		else {
			if (dataStart < 0) {
				dataStart = i;
				dataEnd = i;
			}
			else {
				dataEnd = i;
			}

			emptyStart = -1;
			emptyEnd = -1;
		}
	}
	//ºóÃæµÄÊý¾ÝûÓÐ×ã¹»µÄ¿Õ°×·Ö¸ô
	if (dataEnd - dataStart > 0) {
		rowDataIndexs[rowDataSize * 4 + 0] = dataStart;
		rowDataIndexs[rowDataSize * 4 + 1] = startRow;
		rowDataIndexs[rowDataSize * 4 + 2] = dataEnd;
		rowDataIndexs[rowDataSize * 4 + 3] = endRow;
		rowDataSize++;
	}

	//¿Õ°×Êý¾ÝµÄË÷ÒýλÖÃ
	int emptyColIndex1 = -1;
	int emptyColIndex2 = -1;
	//Èç¹ûûÓÐÂú7¸öÊý¾Ý£¬ÐèÒªÔÚÖм䲹¿ÕÊý¾Ý
	if (rowDataSize < 7) {
		//²¹Âò³· £¬²¹ÕÇÍ£
		if (rowDataSize == 5) {
			emptyColIndex1 = 1;
			emptyColIndex2 = 3;
		}
		else if (rowDataSize == 6) {
			//µÚ3¸öÔªËصĿí¶ÈСÓÚµÚ2¸öÔªËØ
			if ((rowDataIndexs[4 * 2 + 2] - rowDataIndexs[4 * 2 + 0]) < (rowDataIndexs[4 * 1 + 2] - rowDataIndexs[4 * 1 + 0])) {
				//ÓÐÕÇÍ££¬²¹Âò³·
				emptyColIndex1 = 1;
			}
			else {
				//ÎÞÂò³·
				emptyColIndex1 = 3;
			}
		}
		else {

		}
	}

	int emptyCount = 0;
	for (int i = 0; i < 7; i++) {
		if (i == emptyColIndex1 || i == emptyColIndex2) {
			result[startRowDataIndex + 4 * i + 0] = 0;
			result[startRowDataIndex + 4 * i + 1] = 0;
			result[startRowDataIndex + 4 * i + 2] = 0;
			result[startRowDataIndex + 4 * i + 3] = 0;
			emptyCount++;
			/*
			if (gid == 0)
				printf("%d-%d: %d %d %d %d \n", startRowDataIndex + 4 * i + 0, i, 0,0,0,0);
				*/
		}
		else {
			result[startRowDataIndex + 4 * i + 0] = rowDataIndexs[4 * (i - emptyCount) + 0];
			result[startRowDataIndex + 4 * i + 1] = rowDataIndexs[4 * (i - emptyCount) + 1];
			result[startRowDataIndex + 4 * i + 2] = rowDataIndexs[4 * (i - emptyCount) + 2];
			result[startRowDataIndex + 4 * i + 3] = rowDataIndexs[4 * (i - emptyCount) + 3];
			/*
			if(gid==0)

			printf("%d-%d: %d %d %d %d \n", startRowDataIndex + 4 * i + 0,i, rowDataIndexs[4 * (i - emptyCount) + 0], rowDataIndexs[4 * (i - emptyCount) + 1], rowDataIndexs[4 * (i - emptyCount) + 2], rowDataIndexs[4 * (i - emptyCount) + 3]);
		*/
		}

	}
}


int int_sort(int arr[], int len,int ele_count,int index)
{
	int j, k;    //¶¨ÒåÑ­»·Òò×Ó£¬Ç¶Ì×Ë«²ãÑ­»·
	for (j = 0; j < len - 1; j++)    //ÉèÖÃÑ­»·ºó½ç
	{
		for (k = 0; k < len - j - 1; k++)    //²»¶ÏÏòºó½øÐбȽÏ
		{
			if (arr[k] < arr[k + 1])    //±È½ÏÏàÁÚµÄÔªËØ
			{
				int temp = arr[k];    //Èý±­Ë®½»»»·¨
				arr[k] = arr[k + 1];
				arr[k + 1] = temp;
			}
		}
	}

	//if (index == 136) {
	//	for (j = 0; j < len; j++) {
	//		printf("%d,", arr[j]);
	//	}
	//}

	if (ele_count > len) {
	    return arr[len - 1];
	}
	else
	{
		return arr[ele_count - 2];
	}
}

//ÇзÖÐÐÊý¾Ý
//width£ºÍ¼Ïñ¿í¶È
//ele_count£ºÃ¿Ò»ÐеÄÔªËØÊýÁ¿
//min_empty_width£º×îС¿Õ°×£¬Ð¡ÓÚÕâ¸öÖµ²»Ëã¿Õ°×
//threshhold_color:СÓÚÕâ¸öÖµ±íʾÎÞÄÚÈÝ£¬·´Ö®±íʾÓÐÄÚÈÝ
__kernel void splitRowData(__global unsigned char* imgs, __global int* rowIndexs, int width,int ele_count,int min_empty_width,int threshhold_color,
	__global int* result) {
	int gid = get_global_id(0);

	int startCol = rowIndexs[gid * 4];
	int startRow = rowIndexs[gid * 4 + 1];
	int endCol = rowIndexs[gid * 4 + 2];
	int endRow = rowIndexs[gid * 4 + 3];

	//printf(" index:%d - startCol:%d,startRow:%d,endCol:%d,endRow:%d\n", gid, startCol, startRow, endCol, endRow);


	int emptyStart = -1;
	int emptyEnd = -1;
	int dataStart = -1;
	int dataEnd = -1;

	int rowDataSize = 0;
	// ×î´ó¼æÈÝ20¸öÊý¾Ý
	int rowDataIndexs[20 * 4];
	//×î´ó¼æÈÝ80¸öÊý¾Ý
	int startRowDataIndex = ele_count * 4 * gid;
	int empty_area_width[80];

	//ͳ¼Æ¿Õ°×µÄ¿í¶È
	int empty_index = 0;
	bool is_start = 0;
	for (int i = startCol; i <= endCol; i++) {
		bool empty = 1;
		for (int j = startRow; j <= endRow; j++) {
			unsigned char value = imgs[get_one_level_position(width, i, j)];
			if (value >= threshhold_color) {
				empty = 0;
				break;
			}
		}
		if (empty) {
			if (emptyStart < 0) {
				emptyStart = i;
				emptyEnd = i;
			}
			else {
				emptyEnd = i;
			}
		}
		else {
			if (emptyEnd - emptyStart >= min_empty_width && is_start) {
				//if (gid == 136) {
				//	printf("%d,%d,%d,%d\n", emptyStart, startRow, emptyEnd, endRow);
				//}
				empty_area_width[empty_index] = emptyEnd - emptyStart + 1;
				empty_index++;
			}
			emptyEnd = -1;
			emptyStart = -1;
			is_start = 1;
			
		}
	}
	//ÇóÇ°ele_count-1µÄ¿Õ°×¿í¶È
	int empty_width = int_sort(empty_area_width, empty_index,ele_count, gid);



	emptyStart = -1;
	emptyEnd = -1;
	dataStart = -1;
	dataEnd = -1;

	for (int i = startCol; i <= endCol; i++) {

		bool empty = 1;
		for (int j = startRow; j <= endRow; j++) {
			unsigned char value = imgs[get_one_level_position(width, i, j)];
			if (value >= threshhold_color) {
				empty = 0;
				break;
			}
		}
		if (empty) {
			if (emptyStart < 0) {
				emptyStart = i;
				emptyEnd = i;
			}
			else {
				emptyEnd = i;
			}

			//3¸ö¼°ÒÔÉϵĿհ×Êý¾Ý²Å·ÖÁÐ
			if (emptyEnd - emptyStart + 1 >= empty_width && dataEnd - dataStart > 0) {
				rowDataIndexs[rowDataSize * 4 + 0] = dataStart;
				rowDataIndexs[rowDataSize * 4 + 1] = startRow;
				rowDataIndexs[rowDataSize * 4 + 2] = dataEnd;
				rowDataIndexs[rowDataSize * 4 + 3] = endRow;
				rowDataSize++;
				dataEnd = -1;
				dataStart = -1;
			}
		}
		else {
			if (dataStart < 0) {
				dataStart = i;
				dataEnd = i;
			}
			else {
				dataEnd = i;
			}

			emptyStart = -1;
			emptyEnd = -1;
		}
	}
	//ºóÃæµÄÊý¾ÝûÓÐ×ã¹»µÄ¿Õ°×·Ö¸ô
	if (dataEnd - dataStart > 0) {
		rowDataIndexs[rowDataSize * 4 + 0] = dataStart;
		rowDataIndexs[rowDataSize * 4 + 1] = startRow;
		rowDataIndexs[rowDataSize * 4 + 2] = dataEnd;
		rowDataIndexs[rowDataSize * 4 + 3] = endRow;
		rowDataSize++;
	}
	//if (rowDataSize > 4)
	//{
	//	printf("data_size:%d  %d  %d\n", gid, rowDataSize, empty_width);
	//}
	//È¥³ýÉÏÏ¿հ×
	for (int i = 0; i < rowDataSize; i++) {
		int start_x = rowDataIndexs[i * 4 + 0];
		int start_y = rowDataIndexs[i * 4 + 1];
		int end_x = rowDataIndexs[i * 4 + 2];
		int end_y = rowDataIndexs[i * 4 + 3];
		//È¥³ýÉÏ¿Õ°×
		for (int r = start_y; r <= end_y; r++) {
			bool empty = 1;
			for (int c = start_x; c <= end_x; c++) {
				unsigned char value = imgs[get_one_level_position(width, c, r)];
				if (value >= threshhold_color) {
					empty = 0;
					break;
				}
			}
			if (!empty){
				start_y = r;
				break;
			}
		}
		rowDataIndexs[i * 4 + 1] = start_y;

		//È¥³ýÏ¿հ×
		for (int r = end_y; r >= start_y; r--) {
			bool empty = 1;
			for (int c = start_x; c <= end_x; c++) {
				unsigned char value = imgs[get_one_level_position(width, c, r)];
				if (value >= threshhold_color) {
					empty = 0;
					break;
				}
			}
			if (!empty) {
				end_y = r;
				break;
			}
		}
		rowDataIndexs[i * 4 + 3] = end_y;
	}



	int emptyColIndex1 = -1;
	int emptyColIndex2 = -1;
	int emptyCount = 0;
	for (int i = 0; i < ele_count; i++) {
		if (i == emptyColIndex1 || i == emptyColIndex2) {
			result[startRowDataIndex + 4 * i + 0] = 0;
			result[startRowDataIndex + 4 * i + 1] = 0;
			result[startRowDataIndex + 4 * i + 2] = 0;
			result[startRowDataIndex + 4 * i + 3] = 0;
			emptyCount++;
		}
		else {
			result[startRowDataIndex + 4 * i + 0] = rowDataIndexs[4 * (i - emptyCount) + 0];
			result[startRowDataIndex + 4 * i + 1] = rowDataIndexs[4 * (i - emptyCount) + 1];
			result[startRowDataIndex + 4 * i + 2] = rowDataIndexs[4 * (i - emptyCount) + 2];
			result[startRowDataIndex + 4 * i + 3] = rowDataIndexs[4 * (i - emptyCount) + 3];
			/*
			if(gid==0)

			printf("%d-%d: %d %d %d %d \n", startRowDataIndex + 4 * i + 0,i, rowDataIndexs[4 * (i - emptyCount) + 0], rowDataIndexs[4 * (i - emptyCount) + 1], rowDataIndexs[4 * (i - emptyCount) + 2], rowDataIndexs[4 * (i - emptyCount) + 3]);
		*/
		}

	}
}



__kernel void splitCommonRowData(__global unsigned char* imgs, __global int* rowIndexs, int width, int fromIndex, int toIndex,
	__global int* result) {
	int gid = get_global_id(0);

	int startCol = rowIndexs[gid * 4];
	int startRow = rowIndexs[gid * 4 + 1];
	int endCol = rowIndexs[gid * 4 + 2];
	int endRow = rowIndexs[gid * 4 + 3];

	//printf(" index:%d - startCol:%d,startRow:%d,endCol:%d,endRow:%d\n", gid, startCol, startRow, endCol, endRow);


	int emptyStart = -1;
	int emptyEnd = -1;
	int dataStart = -1;
	int dataEnd = -1;

	int rowDataSize = 0;
	int ELE_NUMBER_PER_LINE = toIndex - fromIndex + 1;
	//ÿÐÐ×î¶à50¸öÊý¾Ý
	int rowDataIndexs[50 * 4];

	//ÅųýµÚÒ»ÁÐÐòºÅÁÐ
	int startRowDataIndex = (ELE_NUMBER_PER_LINE) * 4 * gid;

	for (int i = startCol; i <= endCol; i++) {

		bool empty = 1;
		for (int j = startRow; j <= endRow; j++) {
			unsigned char value = imgs[get_one_level_position(width, i, j)];
			if (value >= 64) {
				empty = 0;
				break;
			}
		}
		if (empty) {
			if (emptyStart < 0) {
				emptyStart = i;
				emptyEnd = i;
			}
			else {
				emptyEnd = i;
			}

			//5¸ö¼°ÒÔÉϵĿհ×Êý¾Ý²Å·ÖÁÐ
			if (emptyEnd - emptyStart >= 5 && dataEnd - dataStart > 0) {
				if (rowDataSize >= fromIndex && rowDataSize <= toIndex)
				{
					rowDataIndexs[(rowDataSize - fromIndex) * 4 + 0] = dataStart;
					rowDataIndexs[(rowDataSize - fromIndex) * 4 + 1] = startRow;
					rowDataIndexs[(rowDataSize - fromIndex) * 4 + 2] = dataEnd;
					rowDataIndexs[(rowDataSize - fromIndex) * 4 + 3] = endRow;
				}
				rowDataSize++;


				dataEnd = -1;
				dataStart = -1;
			}
		}
		else {
			if (dataStart < 0) {
				dataStart = i;
				dataEnd = i;
			}
			else {
				dataEnd = i;
			}

			emptyStart = -1;
			emptyEnd = -1;
		}
	}
	//ºóÃæµÄÊý¾ÝûÓÐ×ã¹»µÄ¿Õ°×·Ö¸ô
	if (dataEnd - dataStart > 0) {
		if (rowDataSize >= fromIndex && rowDataSize <= toIndex)
		{
			rowDataIndexs[(rowDataSize - fromIndex) * 4 + 0] = dataStart;
			rowDataIndexs[(rowDataSize - fromIndex) * 4 + 1] = startRow;
			rowDataIndexs[(rowDataSize - fromIndex) * 4 + 2] = dataEnd;
			rowDataIndexs[(rowDataSize - fromIndex) * 4 + 3] = endRow;
		}
		rowDataSize++;
	}

	int emptyCount = 0;
	for (int i = 0; i < ELE_NUMBER_PER_LINE; i++) {
		result[startRowDataIndex + 4 * i + 0] = rowDataIndexs[4 * (i - emptyCount) + 0];
		result[startRowDataIndex + 4 * i + 1] = rowDataIndexs[4 * (i - emptyCount) + 1];
		result[startRowDataIndex + 4 * i + 2] = rowDataIndexs[4 * (i - emptyCount) + 2];
		result[startRowDataIndex + 4 * i + 3] = rowDataIndexs[4 * (i - emptyCount) + 3];
	}
}



//RGBͼתΪ»Ò¶Èͼ
__kernel void rgb2GrayImg(const __global unsigned char* imgs, int width,
	__global unsigned char* result) {

	int rowIndex = get_global_id(0);

	for (int c = 0; c < width; c++) {
		int index = rowIndex * width + c;
		int start = index * 3;
		unsigned char R = imgs[start];
		unsigned char G = imgs[start + 1];
		unsigned char B = imgs[start + 2];
		result[index] = (76 * R + 150 * G + 30 * B) >> 8;
	}
}

__kernel void rgb2ThresholdInvert(const __global unsigned char* imgs, int width,
	__global unsigned char* result, int threadhold) {

	int rowIndex = get_global_id(0);

	for (int c = 0; c < width; c++) {
		int index = rowIndex * width + c;
		int start = index * 3;
		unsigned char R = imgs[start];
		unsigned char G = imgs[start + 1];
		unsigned char B = imgs[start + 2];
		int v = (76 * R + 150 * G + 30 * B) >> 8;

		if (v > threadhold) {
			result[index] = 0;
		}
		else {
			result[index] = 255;
		}
	}
}

__kernel void rgba2GrayImg(const __global unsigned char* imgs, int width,
	__global unsigned char* result) {

	int rowIndex = get_global_id(0);

	for (int c = 0; c < width; c++) {
		int index = rowIndex * width + c;
		int start = index * 4;
		unsigned char R = imgs[start];
		unsigned char G = imgs[start + 1];
		unsigned char B = imgs[start + 2];
		result[index] = (76 * R + 150 * G + 30 * B) >> 8;
		//printf("%d\n", index);
	}
}


__kernel void rgba2ThresholdInvert(const __global unsigned char* imgs, int width,
	__global unsigned char* result, int threadhold) {

	int rowIndex = get_global_id(0);

	for (int c = 0; c < width; c++) {
		int index = rowIndex * width + c;
		int start = index * 4;
		unsigned char R = imgs[start];
		unsigned char G = imgs[start + 1];
		unsigned char B = imgs[start + 2];
		int v = (76 * R + 150 * G + 30 * B) >> 8;
		if (v > threadhold) {
			result[index] = 0;
		}
		else {
			result[index] = 255;
		}
	}
}


//L2·ÇÊý×Öʶ±ð
__kernel void recognition_not_num(__global unsigned char* imgs, __global int* rowIndexs, int width,
	__global int* result) {

	int index = get_global_id(0);
	int row = index / 3;
	int baseIndex = row * 4 * 7;
	if (index % 3 == 0) {
		//³·Ïúµ¥Î»Ê¶±ð
		int startx = baseIndex + 1 * 4;
		int starty = startx + 1;
		int endx = startx + 2;
		int endy = startx + 3;
		if (rowIndexs[startx] > 0) {
			int emptyX = -1;
			for (int x = rowIndexs[endx]; x >= rowIndexs[startx]; x--)
			{
				bool empty = true;
				for (int y = rowIndexs[starty]; y <= rowIndexs[endy]; y++) {
					unsigned char value = imgs[get_one_level_position(width, x, y)];
					if (get_binary_value(value)) {
						empty = 0;
						break;
					}
				}
				if (empty) {
					emptyX = x;
					break;
				}
			}
			if (emptyX > 0) {
				//¼ÆËãÖµ
				unsigned char count = 0;
				for (int x = emptyX + 1; x <= rowIndexs[endx]; x++)
				{
					for (int y = rowIndexs[starty]; y <= rowIndexs[endy]; y++) {
						unsigned char value = imgs[get_one_level_position(width, x, y)];
						count += get_binary_value(value) ? 1 : 0;
					}
				}
				//printf("row:%d count:%d\n", row, count);
				if (abs(count - 12) <= 1) {//S
					result[row * 3] = 0;
				}
				else if (abs(count - 16) <= 1) {//M
					result[row * 3] = 1;
				}
				else if (abs(count - 19) <= 1) {//H
					result[row * 3] = 2;
				}

			}
		}
		else {
			//ĬÈÏs
			result[row * 3] = 0;
			//printf("row:%d count:0\n", row);
		}
	}
	else if (index % 3 == 1) {
		//ÕǵøÍ£¼Û
		//»ñÈ¡Êý¾Ý×ø±ê
		int startx = baseIndex + 3 * 4;
		int starty = startx + 1;
		int endx = startx + 2;
		int endy = startx + 3;
		if (rowIndexs[startx] <= 0 && rowIndexs[starty] <= 0 && rowIndexs[endx] <= 0 && rowIndexs[endy] <= 0) {
			result[row * 3 + 1] = 0;
		}
		else {
			startx = rowIndexs[startx];
			starty = rowIndexs[starty];
			endx = rowIndexs[endx];
			endy = rowIndexs[endy];

			result[row * 2] = 1;
			//È¥³ýÉÏϵĿհ×
			for (int r = starty; r <= endy; r++)
			{
				bool empty = true;
				for (int c = startx; c <= endx; c++) {
					unsigned char value = imgs[get_one_level_position(width, c, r)];
					if (get_binary_value(value)) {
						empty = 0;
						break;
					}
				}
				if (!empty) {
					starty = r;
					break;
				}
			}
			for (int r = endy; r >= starty; r--)
			{
				bool empty = true;
				for (int c = startx; c <= endx; c++) {
					unsigned char value = imgs[get_one_level_position(width, c, r)];
					if (get_binary_value(value)) {
						empty = 0;
						break;
					}
				}
				if (!empty) {
					endy = r;
					break;
				}
			}
			int my = (starty + endy) / 2;
			//¼ÆËãÉϰ벿·ÖµÄÖµ
			int topValue = 0;
			for (int r = starty; r <= my; r++)
			{
				for (int c = startx; c <= endx; c++) {
					unsigned char value = imgs[get_one_level_position(width, c, r)];
					topValue += get_binary_value(value);

				}
			}
			//¼ÆËãÏ°벿·ÖµÄÖµ
			int bottomValue = 0;
			for (int r = my; r <= endy; r++)
			{
				for (int c = startx; c <= endx; c++) {
					unsigned char value = imgs[get_one_level_position(width, c, r)];
					bottomValue += get_binary_value(value);

				}
			}

			if (topValue > bottomValue) {
				//ÕÇÍ£
				result[row * 3 + 1] = 1;
			}
			else {
				//µøÍ£
				result[row * 3 + 1] = 2;
			}
		}
	}
	else {
		//²Ù×÷ÀàÐÍ
			//¼ÆËãÖµ
		int startx = baseIndex + 6 * 4;
		int starty = startx + 1;
		int endx = startx + 2;
		int endy = startx + 3;
		startx = rowIndexs[startx];
		starty = rowIndexs[starty];
		endx = rowIndexs[endx];
		endy = rowIndexs[endy];

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

		unsigned int pixelCount = 0;
		for (int r = starty; r <= endy; r++)
		{
			for (int c = startx; c <= endx; c++) {
				unsigned char val = imgs[get_one_level_position(width, c, r)];
				pixelCount += get_binary_value(val);
			}
		}

		//³õʼ»¯´íÎó¸´ÖÆ
		int value = 4;
		if (abs(pixelCount - 39) < 5) {
			//Âò
			value = 0;
		}
		else if (abs(pixelCount - 51) < 5) {
			//Âô
			value = 2;
		}
		else if (abs(pixelCount - 105) < 5) {
			//Âò³·
			value = 1;
		}
		else if (abs(pixelCount - 117) < 5) {
			//Âô³·
			value = 3;
		}
		result[row * 3 + 2] = value;
	}

}


//ʶ±ð°å¿éÖеĽð¶îµ¥Î»
__kernel void recognition_plate_money_unit(__global unsigned char* imgs, __global int* rowIndexs, int width,
	__global int* result) {
	int index = get_global_id(0);
	int start = index * 4;
	int startx = rowIndexs[start + 0];
	int starty = rowIndexs[start + 1];
	int endx = rowIndexs[start + 2];
	int endy = rowIndexs[start + 3];

	int unit_start = -1;
	int i;

	//²éѯÊÇ·ñÓе¥Î»
	  //»ñÈ¡ÊÇ·ñÓе¥Î»
	int last_empty_index = -1;
	for (i = endx; i >= startx; i--) {
		bool empty = 1;
		for (int y = starty; y < endy; y++) {
			if (get_binary_value(imgs[get_one_level_position(width, i, y)]) > 0) {
				empty = 0;
				break;
			}
		}
		if (empty) {
			last_empty_index = i;
			break;
		}
	}
	bool hasUnit = false;
	if (endx - last_empty_index > 6) {
		hasUnit = true;
	}
	else {
		//ûÓе¥Î»
		result[index] = 2;
		return;
	}

	for (i = endx - 11; i >= startx; i--) {
		bool empty = 1;
		for (int y = starty; y < endy; y++) {
			if (get_binary_value(imgs[get_one_level_position(width, i, y)]) > 0) {
				empty = 0;
				break;
			}
		}
		if (empty) {
			unit_start = i + 1;
			break;
		}
	}

	//printf("index:%d unit_start:%d\n", index, unit_start);

	int count = 0;
	if (unit_start > -1)
	{
		for (int c = unit_start; c <= endx; c++) {
			for (int y = starty; y <= endy; y++) {
				count += (get_binary_value(imgs[get_one_level_position(width, c, y)]) > 0 ? 1 : 0);
			}
		}
	}
	if (count >= 48) {
		result[index] = 0;
	}
	else {
		result[index] = 1;
	}
}