OGR-27 OpenCL core write attempt

--nextPart1383547.jsE7uqTViX
Content-Transfer-Encoding: 7Bit
Content-Type: text/plain; charset="us-ascii"

   Hi

I try to write OpenCL OGR-27 core (see attachment), but I have encounter 
problem with choose() macro. This is in fact some big table generated in way I 
don't understood.

So I have question.
Can I replace it some way, or make some static pregenerated table for ALL 
working packets? Note that memory access is costly on GPU, so little more 
counting can be preferable over memory read.


PS. I'm new to OGR & OpenCL coding.
   
-- 
  Konrad Rzepecki (Hannibal)
--nextPart1383547.jsE7uqTViX
Content-Disposition: attachment; filename="ogr_tst.ocl.c"
Content-Transfer-Encoding: 7Bit
Content-Type: text/x-csrc; charset="UTF-8"; name="ogr_tst.ocl.c"

#define OGR_MAXDEPTH		27
#define OGR_MAXDEPTHM1		26
#define OGR_HALFDEPTH		13
#define OGR_HALFDEPTH2		14

#define COMP_LEFT_LIST_RIGHT() { \
	ulong4 tmp; \
	tmp = lin[lid].Levels[depth].comp >> (64 - s); \
	lin[lid].Levels[depth].comp = (lin[lid].Levels[depth].comp << s) | (ulong4)(tmp.s1, tmp.s2, tmp.s3, 0); \
	tmp = lin[lid].Levels[depth].list << (64 - s); \
	lin[lid].Levels[depth].list = (lin[lid].Levels[depth].list >> s) | (ulong4)((newbit >> s), tmp.s0, tmp.s1, tmp.s2); \
	newbit = 0; \
}

#define COMP_LEFT_LIST_RIGHT_WORD() { \
	lin[lid].Levels[depth].comp = (ulong4)(lin[lid].Levels[depth].comp.s1, lin[lid].Levels[depth].comp.s2, lin[lid].Levels[depth].comp.s3, 0); \
	lin[lid].Levels[depth].list = (ulong4)((newbit >> 63), lin[lid].Levels[depth].list.s0, lin[lid].Levels[depth].list.s1, lin[lid].Levels[depth].list.s2); \
	newbit = 0; \
}

#define PUSH_LEVEL_UPDATE_STATE() { \
	lin[lid].Levels[depth + 1].list = lin[lid].Levels[depth].list; \
	lin[lid].Levels[depth + 1].dist = lin[lid].Levels[depth].dist | lin[lid].Levels[depth + 1].list; \
	lin[lid].Levels[depth + 1].comp = lin[lid].Levels[depth].comp | lin[lid].Levels[depth + 1].dist; \
	newbit = 1; \
}

#define GOTO_EXIT() { \
	state_from_loc(out, lin, gid, lid); \
	ovar[gid].pnodes -= nodes;\
	ovar[gid].depth = depth; \
	return; \
}


struct OtherVars {
	int pnodes;
	int depth;
};

struct OgrLevel {
	ulong4 list;
	ulong4 dist;
	ulong4 comp;
	int mark;
	int limit;
};

struct OgrState {
	int max;
	int stopdepth;
	int depth;
	struct OgrLevel Levels[OGR_MAXDEPTH + 1];
};

void state_to_loc(global struct OgrState *in, local struct OgrState *lin, private int gid, private int lid) {
	int i;

	lin[lid].max = in[gid].max;
	lin[lid].stopdepth = in[gid].stopdepth;
	lin[lid].depth = in[gid].depth;

	for (i = 0; i <= OGR_MAXDEPTH; i++) {
		lin[lid].Levels[i].list = in[gid].Levels[i].list;
		lin[lid].Levels[i].dist = in[gid].Levels[i].dist;
		lin[lid].Levels[i].comp = in[gid].Levels[i].comp;
		lin[lid].Levels[i].mark = in[gid].Levels[i].mark;
		lin[lid].Levels[i].limit = in[gid].Levels[i].limit;
	}
}

void state_from_loc(global struct OgrState *out, local struct OgrState *lin, private int gid, private int lid) {
	int i;

	out[gid].max = lin[lid].max;
	out[gid].stopdepth = lin[lid].stopdepth;
	out[gid].depth = lin[lid].depth;

	for (i = 0; i <= OGR_MAXDEPTH; i++) {
		out[gid].Levels[i].list = lin[lid].Levels[i].list;
		out[gid].Levels[i].dist = lin[lid].Levels[i].dist;
		out[gid].Levels[i].comp = lin[lid].Levels[i].comp;
		out[gid].Levels[i].mark = lin[lid].Levels[i].mark;
		out[gid].Levels[i].limit = lin[lid].Levels[i].limit;
	}
}

kernel void ogr_cycle_256(global struct OgrState *in, global struct OgrState *out, global struct OtherVars *ivar, global struct OtherVars *ovar) {
	int gid = get_global_id(0);
	int lid = get_local_id(0);
	local struct OgrState *lin = 0; //???
	int depth = in[gid].depth;
	int nodes = ivar[gid].pnodes;
	ulong newbit = (depth < OGR_MAXDEPTHM1) ? 1 : 0;

	state_to_loc(in, lin, gid, lid);

	do {
		for(;;) {
			if (lin[lid].Levels[depth].comp.s0 > ~(ulong)(1)) {
				int s = clz(lin[lid].Levels[depth].comp.s0);

				if ((lin[lid].Levels[depth].mark += s) > lin[lid].Levels[depth].limit) {
					break;
				}
				COMP_LEFT_LIST_RIGHT();
			} else {
				if ((lin[lid].Levels[depth].mark += 64) > lin[lid].Levels[depth].limit) {
					break;
				}

				newbit = 0;
				if (lin[lid].Levels[depth].comp.s0 == ~(ulong)(0)) {
					COMP_LEFT_LIST_RIGHT_WORD();
					continue;
				}
				COMP_LEFT_LIST_RIGHT_WORD();
			}

			if (depth == OGR_MAXDEPTHM1) {
				GOTO_EXIT();
			}

			PUSH_LEVEL_UPDATE_STATE();
			++depth;

			#warning lin[lid].Levels[depth].limit = choose(lin[lid].Levels[depth].dist.s0, depth); !!!

			if (depth > OGR_HALFDEPTH && depth <= OGR_HALFDEPTH2) {
				int temp = lin[lid].max - lin[lid].Levels[OGR_HALFDEPTH].mark - 1;

				if (depth < OGR_HALFDEPTH2) {
					temp -= clz(lin[lid].Levels[depth].dist.s0);
				}

				if (lin[lid].Levels[depth].limit > temp) {
					lin[lid].Levels[depth].limit = temp;
				}
			}

			if (--nodes <= 0) {
				GOTO_EXIT();
			}

		}

		--depth;
		newbit = 0;
	} while (depth > lin[lid].stopdepth);

	GOTO_EXIT();
}

--nextPart1383547.jsE7uqTViX
Content-Type: text/plain; charset="us-ascii"
MIME-Version: 1.0
Content-Transfer-Encoding: 7bit
Content-Disposition: inline

_______________________________________________
Hardware mailing list
Hardware-Ra3b/QYEcJ3d140v2zMXi0fjHoOT/h/[email protected]
http://lists.distributed.net/mailman/listinfo/hardware

--nextPart1383547.jsE7uqTViX--