continue removing constants (9182fde0) · Commits · Elphel / tile_processor_gpu

src/TileProcessor.cuh

+41 −18

Original line number	Original line	Diff line number	Diff line
	@@ -838,12 +838,15 @@ __global__ void mark_texture_tiles(
	__global__ void mark_texture_neighbor_tiles(		__global__ void mark_texture_neighbor_tiles(
	struct tp_task * gpu_tasks,		struct tp_task * gpu_tasks,
	int num_tiles, // number of tiles in task list		int num_tiles, // number of tiles in task list
			int width, // number of tiles in a row
	int * gpu_texture_indices, // packed tile + bits (now only (1 << 7)		int * gpu_texture_indices, // packed tile + bits (now only (1 << 7)
	int * woi); // x,y,width,height of the woi		int * woi); // x,y,width,height of the woi

	__global__ void gen_texture_list(		__global__ void gen_texture_list(
	struct tp_task * gpu_tasks,		struct tp_task * gpu_tasks,
	int num_tiles, // number of tiles in task list		int num_tiles, // number of tiles in task list
			int width, // number of tiles in a row
			int height, // number of tiles rows
	int * gpu_texture_indices, // packed tile + bits (now only (1 << 7)		int * gpu_texture_indices, // packed tile + bits (now only (1 << 7)
	int * num_texture_tiles, // number of texture tiles to process		int * num_texture_tiles, // number of texture tiles to process
	int * woi); // x,y,width,height of the woi		int * woi); // x,y,width,height of the woi
	@@ -1317,6 +1320,7 @@ extern "C" __global__ void generate_RBGA(
	mark_texture_neighbor_tiles <<<blocks,threads>>>(		mark_texture_neighbor_tiles <<<blocks,threads>>>(
	gpu_tasks,		gpu_tasks,
	num_tiles, // number of tiles in task list		num_tiles, // number of tiles in task list
			width, // number of tiles in a row
	gpu_texture_indices, // packed tile + bits (now only (1 << 7)		gpu_texture_indices, // packed tile + bits (now only (1 << 7)
	woi); // min_x, min_y, max_x, max_y		woi); // min_x, min_y, max_x, max_y

	@@ -1334,6 +1338,8 @@ extern "C" __global__ void generate_RBGA(
	gen_texture_list <<<blocks,threads>>>(		gen_texture_list <<<blocks,threads>>>(
	gpu_tasks,		gpu_tasks,
	num_tiles, // number of tiles in task list		num_tiles, // number of tiles in task list
			width, // number of tiles in a row
			height, // int height, // number of tiles rows
	gpu_texture_indices, // packed tile + bits (now only (1 << 7)		gpu_texture_indices, // packed tile + bits (now only (1 << 7)
	num_texture_tiles, // number of texture tiles to process		num_texture_tiles, // number of texture tiles to process
	woi); // x,y, here woi[2] = max_X, woi[3] - max-Y		woi); // x,y, here woi[2] = max_X, woi[3] - max-Y
	@@ -1495,6 +1501,7 @@ __global__ void prepare_texture_list(
	mark_texture_neighbor_tiles <<<blocks,threads>>>(		mark_texture_neighbor_tiles <<<blocks,threads>>>(
	gpu_tasks,		gpu_tasks,
	num_tiles, // number of tiles in task list		num_tiles, // number of tiles in task list
			width, // number of tiles in a row
	gpu_texture_indices, // packed tile + bits (now only (1 << 7)		gpu_texture_indices, // packed tile + bits (now only (1 << 7)
	woi); // min_x, min_y, max_x, max_y		woi); // min_x, min_y, max_x, max_y
	cudaDeviceSynchronize();		cudaDeviceSynchronize();
	@@ -1511,6 +1518,8 @@ __global__ void prepare_texture_list(
	gen_texture_list <<<blocks,threads>>>(		gen_texture_list <<<blocks,threads>>>(
	gpu_tasks,		gpu_tasks,
	num_tiles, // number of tiles in task list		num_tiles, // number of tiles in task list
			width, // number of tiles in a row
			height, // int height, // number of tiles rows
	gpu_texture_indices, // packed tile + bits (now only (1 << 7)		gpu_texture_indices, // packed tile + bits (now only (1 << 7)
	num_texture_tiles, // number of texture tiles to process		num_texture_tiles, // number of texture tiles to process
	woi); // x,y, here woi[2] = max_X, woi[3] - max-Y		woi); // x,y, here woi[2] = max_X, woi[3] - max-Y
	@@ -1586,6 +1595,7 @@ __global__ void mark_texture_tiles(
	__global__ void mark_texture_neighbor_tiles( // TODO: remove __global__?		__global__ void mark_texture_neighbor_tiles( // TODO: remove __global__?
	struct tp_task * gpu_tasks,		struct tp_task * gpu_tasks,
	int num_tiles, // number of tiles in task list		int num_tiles, // number of tiles in task list
			int width, // number of tiles in a row
	int * gpu_texture_indices, // packed tile + bits (now only (1 << 7)		int * gpu_texture_indices, // packed tile + bits (now only (1 << 7)
	int * woi) // x,y,width,height of the woi		int * woi) // x,y,width,height of the woi

	@@ -1606,10 +1616,15 @@ __global__ void mark_texture_neighbor_tiles( // TODO: remove __global__?
	atomicMax(woi+2, x);		atomicMax(woi+2, x);
	atomicMax(woi+3, y);		atomicMax(woi+3, y);
	int d = 0;		int d = 0;
	if ((y > 0) && (gpu_texture_indices + x + (y - 1) TILESX)) d \|= (1 << TASK_TEXTURE_N_BIT);		// if ((y > 0) && (gpu_texture_indices + x + (y - 1) TILESX)) d \|= (1 << TASK_TEXTURE_N_BIT);
	if ((x < (TILESX - 1)) && (gpu_texture_indices + (x + 1) + y TILESX)) d \|= (1 << TASK_TEXTURE_E_BIT);		// if ((x < (TILESX - 1)) && (gpu_texture_indices + (x + 1) + y TILESX)) d \|= (1 << TASK_TEXTURE_E_BIT);
	if ((y < (TILESY - 1)) && (gpu_texture_indices + x + (y + 1) TILESX)) d \|= (1 << TASK_TEXTURE_S_BIT);		// if ((y < (TILESY - 1)) && (gpu_texture_indices + x + (y + 1) TILESX)) d \|= (1 << TASK_TEXTURE_S_BIT);
	if ((x > 0) && (gpu_texture_indices + (x - 1) + y TILESX)) d \|= (1 << TASK_TEXTURE_W_BIT);		// if ((x > 0) && (gpu_texture_indices + (x - 1) + y TILESX)) d \|= (1 << TASK_TEXTURE_W_BIT);

			if ((y > 0) && (gpu_texture_indices + x + (y - 1) width)) d \|= (1 << TASK_TEXTURE_N_BIT);
			if ((x < (TILESX - 1)) && (gpu_texture_indices + (x + 1) + y width)) d \|= (1 << TASK_TEXTURE_E_BIT);
			if ((y < (TILESY - 1)) && (gpu_texture_indices + x + (y + 1) width)) d \|= (1 << TASK_TEXTURE_S_BIT);
			if ((x > 0) && (gpu_texture_indices + (x - 1) + y width)) d \|= (1 << TASK_TEXTURE_W_BIT);
	gpu_tasks[task_num].task = ((task ^ d) & TASK_TEXTURE_BITS) ^ task;		gpu_tasks[task_num].task = ((task ^ d) & TASK_TEXTURE_BITS) ^ task;
	}		}

	@@ -1629,11 +1644,14 @@ __global__ void mark_texture_neighbor_tiles( // TODO: remove __global__?
	__global__ void gen_texture_list(		__global__ void gen_texture_list(
	struct tp_task * gpu_tasks,		struct tp_task * gpu_tasks,
	int num_tiles, // number of tiles in task list		int num_tiles, // number of tiles in task list
			int width, // number of tiles in a row
			int height, // number of tiles rows
	int * gpu_texture_indices, // packed tile + bits (now only (1 << 7)		int * gpu_texture_indices, // packed tile + bits (now only (1 << 7)
	int * num_texture_tiles, // number of texture tiles to process		int * num_texture_tiles, // number of texture tiles to process
	int * woi) // min_x, min_y, max_x, max_y input		int * woi) // min_x, min_y, max_x, max_y input

	{		{
			int tilesya = ((height +3) & (~3)); //#define TILESYA ((TILESY +3) & (~3))
	int task_num = blockDim.x * blockIdx.x + threadIdx.x;		int task_num = blockDim.x * blockIdx.x + threadIdx.x;
	if (task_num >= num_tiles) {		if (task_num >= num_tiles) {
	return; // nothing to do		return; // nothing to do
	@@ -1659,19 +1677,23 @@ __global__ void gen_texture_list(


	// don't care if calculate extra pixels that still fit into memory		// don't care if calculate extra pixels that still fit into memory
	int is_border = (x == woi[0]) \|\| (y == woi[1]) \|\| (x == (TILESX - 1)) \|\| (y == woi[3]);		// int is_border = (x == woi[0]) \|\| (y == woi[1]) \|\| (x == (TILESX - 1)) \|\| (y == woi[3]);
			int is_border = (x == woi[0]) \|\| (y == woi[1]) \|\| (x == (width - 1)) \|\| (y == woi[3]);
	int buff_head = 0;		int buff_head = 0;
	int num_offset = 0;		int num_offset = 0;
	if (x & 1) {		if (x & 1) {
	buff_head += TILESX * (TILESYA >> 2); //TILESYA - 2 LSB == 00		// buff_head += TILESX * (TILESYA >> 2); //TILESYA - 2 LSB == 00
			buff_head += width * (tilesya >> 2); //TILESYA - 2 LSB == 00
	num_offset += 2; // int *		num_offset += 2; // int *
	}		}
	if (y & 1) {		if (y & 1) {
	buff_head += TILESX * (TILESYA >> 1);		// buff_head += TILESX * (TILESYA >> 1);
			buff_head += width * (tilesya >> 1);
	num_offset += 4; // int *		num_offset += 4; // int *
	}		}
	if (is_border){		if (is_border){
	buff_head += (TILESX * (TILESYA >> 2) - 1); // end of the buffer		// buff_head += (TILESX * (TILESYA >> 2) - 1); // end of the buffer
			buff_head += (width * (tilesya >> 2) - 1); // end of the buffer
	num_offset += 1; // int *		num_offset += 1; // int *
	}		}
	gpu_texture_indices += buff_head;		gpu_texture_indices += buff_head;
	@@ -1690,7 +1712,8 @@ __global__ void gen_texture_list(
	}		}
	__syncthreads();// __syncwarp();		__syncthreads();// __syncwarp();
	#endif // DEBUG12		#endif // DEBUG12
	(gpu_texture_indices + buf_offset) = task \| ((x + y TILESX) << CORR_NTILE_SHIFT) \| (1 << LIST_TEXTURE_BIT);		// (gpu_texture_indices + buf_offset) = task \| ((x + y TILESX) << CORR_NTILE_SHIFT) \| (1 << LIST_TEXTURE_BIT);
			(gpu_texture_indices + buf_offset) = task \| ((x + y width) << CORR_NTILE_SHIFT) \| (1 << LIST_TEXTURE_BIT);
	}		}

	/**		/**

Admin message