be17f28 · opencog-CompactHash

Changed files (3) +1494 -519

HashFactory/HashFactory.c +1462 -519

…	…	@@ -3956,116 +3956,117 @@ int intintLCGLinearOpenCompactCLHash_BufferInsertNoOverwrite(intintHash_Table *
3956	3956	return (0);
3957	3957	}
3958	3958
3959		typedef struct intintLCGQuadraticOpenCompactHash_TableData {
	3959	typedef struct intintLCGLinearOpenCompactOpenMPHash_TableData {
3960	3960	int hashID;
3961	3961	unsigned int numBuckets;
3962	3962	intintHash_CompressLCGData compressFuncData;
3963		} intintLCGQuadraticOpenCompactHash_TableData;
3964		typedef struct intintLCGQuadraticOpenCompactHash_Bucket {
	3963	} intintLCGLinearOpenCompactOpenMPHash_TableData;
	3964	typedef struct intintLCGLinearOpenCompactOpenMPHash_Bucket {
3965	3965	int key;
3966	3966	int value;
3967		} intintLCGQuadraticOpenCompactHash_Bucket;
	3967	} intintLCGLinearOpenCompactOpenMPHash_Bucket;
3968	3968	intintHash_Table
3969		intintLCGQuadraticOpenCompactHash_CreateTable(intintHash_Factory factory,
3970		int hashIndex,
3971		size_t keyRange,
3972		size_t numEntries,
3973		float loadFactor) {
	3969	intintLCGLinearOpenCompactOpenMPHash_CreateTable(intintHash_Factory
	3970	factory, int hashIndex,
	3971	size_t keyRange,
	3972	size_t numEntries,
	3973	float loadFactor) {
3974	3974	intintHash_Table *table =
3975	3975	(intintHash_Table *) malloc(sizeof(intintHash_Table));
3976		table->destroyFunc = &intintLCGQuadraticOpenCompactHash_DestroyTable;
3977		table->emptyFunc = &intintLCGQuadraticOpenCompactHash_EmptyTable;
3978		table->queryFunc = &intintLCGQuadraticOpenCompactHash_Query;
3979		table->querySingleFunc = &intintLCGQuadraticOpenCompactHash_QuerySingle;
3980		table->insertFunc = &intintLCGQuadraticOpenCompactHash_Insert;
	3976	table->destroyFunc = &intintLCGLinearOpenCompactOpenMPHash_DestroyTable;
	3977	table->emptyFunc = &intintLCGLinearOpenCompactOpenMPHash_EmptyTable;
	3978	table->queryFunc = &intintLCGLinearOpenCompactOpenMPHash_Query;
	3979	table->querySingleFunc =
	3980	&intintLCGLinearOpenCompactOpenMPHash_QuerySingle;
	3981	table->insertFunc = &intintLCGLinearOpenCompactOpenMPHash_Insert;
3981	3982	table->insertSingleFunc =
3982		&intintLCGQuadraticOpenCompactHash_InsertSingle;
	3983	&intintLCGLinearOpenCompactOpenMPHash_InsertSingle;
3983	3984	table->insertNoOverwriteFunc =
3984		&intintLCGQuadraticOpenCompactHash_InsertNoOverwrite;
	3985	&intintLCGLinearOpenCompactOpenMPHash_InsertNoOverwrite;
3985	3986	table->insertSingleNoOverwriteFunc =
3986		&intintLCGQuadraticOpenCompactHash_InsertSingleNoOverwrite;
	3987	&intintLCGLinearOpenCompactOpenMPHash_InsertSingleNoOverwrite;
3987	3988	table->tableData =
3988		(char *)malloc(sizeof(intintLCGQuadraticOpenCompactHash_TableData));
3989		((intintLCGQuadraticOpenCompactHash_TableData *) table->tableData)->
3990		hashID = LCG_QUADRATIC_OPEN_COMPACT_HASH_ID;
3991		((intintLCGQuadraticOpenCompactHash_TableData *) table->tableData)->
	3989	(char *)
	3990	malloc(sizeof(intintLCGLinearOpenCompactOpenMPHash_TableData));
	3991	((intintLCGLinearOpenCompactOpenMPHash_TableData *) table->tableData)->
	3992	hashID = LCG_LINEAR_OPEN_COMPACT_OPENMP_HASH_ID;
	3993	((intintLCGLinearOpenCompactOpenMPHash_TableData *) table->tableData)->
3992	3994	numBuckets = (unsigned int)((double)numEntries / loadFactor);
3993		((intintLCGQuadraticOpenCompactHash_TableData *) table->tableData)->
	3995	((intintLCGLinearOpenCompactOpenMPHash_TableData *) table->tableData)->
3994	3996	compressFuncData.a = HASH_LCG_A;
3995		((intintLCGQuadraticOpenCompactHash_TableData *) table->tableData)->
	3997	((intintLCGLinearOpenCompactOpenMPHash_TableData *) table->tableData)->
3996	3998	compressFuncData.c = HASH_LCG_C;
3997		((intintLCGQuadraticOpenCompactHash_TableData *) table->tableData)->
	3999	((intintLCGLinearOpenCompactOpenMPHash_TableData *) table->tableData)->
3998	4000	compressFuncData.m = HASH_LCG_M;
3999		((intintLCGQuadraticOpenCompactHash_TableData *) table->tableData)->
	4001	((intintLCGLinearOpenCompactOpenMPHash_TableData *) table->tableData)->
4000	4002	compressFuncData.n =
4001		((intintLCGQuadraticOpenCompactHash_TableData *) table->tableData)->
4002		numBuckets;
4003		((intintLCGQuadraticOpenCompactHash_TableData *) table->tableData)->
4004		numBuckets =
4005		largestProthPrimeUnder(((intintLCGQuadraticOpenCompactHash_TableData
4006		*) table->tableData)->numBuckets);
	4003	((intintLCGLinearOpenCompactOpenMPHash_TableData *) table->
	4004	tableData)->numBuckets;
4007	4005	char *tempHashData =
4008		(char *)malloc(sizeof(intintLCGQuadraticOpenCompactHash_TableData) +
4009		((intintLCGQuadraticOpenCompactHash_TableData *)
4010		table->tableData)->numBuckets *
4011		sizeof(intintLCGQuadraticOpenCompactHash_Bucket));
	4006	(char *)
	4007	malloc(sizeof(intintLCGLinearOpenCompactOpenMPHash_TableData) +
	4008	((intintLCGLinearOpenCompactOpenMPHash_TableData *) table->
	4009	tableData)->numBuckets *
	4010	sizeof(intintLCGLinearOpenCompactOpenMPHash_Bucket));
4012	4011	memcpy(tempHashData, table->tableData,
4013		sizeof(intintLCGQuadraticOpenCompactHash_TableData));
	4012	sizeof(intintLCGLinearOpenCompactOpenMPHash_TableData));
4014	4013	free(table->tableData);
4015	4014	table->tableData = tempHashData;
4016	4015	return table;
4017	4016	}
4018		int intintLCGQuadraticOpenCompactHash_CreateFactory(intintHash_Factory *
4019		factory, int hashIndex) {
	4017	int intintLCGLinearOpenCompactOpenMPHash_CreateFactory(intintHash_Factory *
	4018	factory, int hashIndex) {
4020	4019	factory->createFunc[hashIndex] =
4021		&intintLCGQuadraticOpenCompactHash_CreateTable;
	4020	&intintLCGLinearOpenCompactOpenMPHash_CreateTable;
4022	4021	factory->destroyFunc[hashIndex] =
4023		&intintLCGQuadraticOpenCompactHash_DestroyFactory;;
	4022	&intintLCGLinearOpenCompactOpenMPHash_DestroyFactory;;
4024	4023	return HASH_EXIT_CODE_NORMAL;
4025	4024	}
4026		int intintLCGQuadraticOpenCompactHash_DestroyFactory(intintHash_Factory *
4027		factory, int hashIndex) {;
	4025	int intintLCGLinearOpenCompactOpenMPHash_DestroyFactory(intintHash_Factory *
	4026	factory,
	4027	int hashIndex) {;
4028	4028	return HASH_EXIT_CODE_NORMAL;
4029	4029	}
4030		int intintLCGQuadraticOpenCompactHash_DestroyTable(intintHash_Table * table) {
	4030	int intintLCGLinearOpenCompactOpenMPHash_DestroyTable(intintHash_Table * table) {
4031	4031	int exitCode = 0;
4032	4032	free(table->tableData);
4033	4033	free(table);
4034	4034	return exitCode;
4035	4035	}
4036		int intintLCGQuadraticOpenCompactHash_EmptyTable(intintHash_Table * table) {
	4036	int intintLCGLinearOpenCompactOpenMPHash_EmptyTable(intintHash_Table * table) {
4037	4037	int exitCode = 0;
4038		intintLCGQuadraticOpenCompactHash_Bucket *buckets =
4039		(intintLCGQuadraticOpenCompactHash_Bucket *) & table->
4040		tableData[sizeof(intintLCGQuadraticOpenCompactHash_TableData)];
	4038	intintLCGLinearOpenCompactOpenMPHash_Bucket *buckets =
	4039	(intintLCGLinearOpenCompactOpenMPHash_Bucket *) & table->
	4040	tableData[sizeof(intintLCGLinearOpenCompactOpenMPHash_TableData)];
4041	4041	int index;
4042	4042	for (index = 0;
4043	4043	index <
4044		((intintLCGQuadraticOpenCompactHash_TableData *) table->
	4044	((intintLCGLinearOpenCompactOpenMPHash_TableData *) table->
4045	4045	tableData)->numBuckets; index++) {
4046	4046	buckets[index].key = HASH_BUCKET_STATUS_EMPTY;
4047	4047	}
4048	4048	exitCode = HASH_EXIT_CODE_NORMAL;
4049	4049	return exitCode;
4050	4050	}
4051		int intintLCGQuadraticOpenCompactHash_InnerQuerySingle(char *tableData, int key,
4052		int *valueOutput) {
4053		intintLCGQuadraticOpenCompactHash_Bucket *buckets =
4054		(intintLCGQuadraticOpenCompactHash_Bucket *) &
4055		tableData[sizeof(intintLCGQuadraticOpenCompactHash_TableData)];
	4051	int intintLCGLinearOpenCompactOpenMPHash_InnerQuerySingle(char *tableData,
	4052	int key,
	4053	int *valueOutput) {
	4054	intintLCGLinearOpenCompactOpenMPHash_Bucket *buckets =
	4055	(intintLCGLinearOpenCompactOpenMPHash_Bucket *) &
	4056	tableData[sizeof(intintLCGLinearOpenCompactOpenMPHash_TableData)];
4056	4057	int index;
4057	4058	int exitCode;
4058		intintLCGQuadraticOpenCompactHash_TableData *mytableData =
4059		(intintLCGQuadraticOpenCompactHash_TableData *) tableData;
	4059	intintLCGLinearOpenCompactOpenMPHash_TableData *mytableData =
	4060	(intintLCGLinearOpenCompactOpenMPHash_TableData *) tableData;
4060	4061	intintHash_CompressLCGData compressFuncData =
4061	4062	mytableData->compressFuncData;
4062	4063	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
4063	4064	unsigned long int iteration = 0;
4064	4065	for (;;) {
4065	4066	index =
4066		((1 * iteration * iteration + 0 * iteration +
	4067	((1 * iteration +
4067	4068	c) %
4068		((intintLCGQuadraticOpenCompactHash_TableData *)
	4069	((intintLCGLinearOpenCompactOpenMPHash_TableData *)
4069	4070	tableData)->numBuckets);
4070	4071	if ((buckets[index].key) == HASH_BUCKET_STATUS_EMPTY) {
4071	4072	exitCode = HASH_SEARCH_CODE_EMPTY;
…	…	@@ -4073,10 +4074,7 @@ int intintLCGQuadraticOpenCompactHash_InnerQuerySingle(char *tableData, int key,
4073	4074	} else if (key == buckets[index].key) {
4074	4075	exitCode = HASH_SEARCH_CODE_MATCH;
4075	4076	break;
4076		} else
4077		if ((iteration >
4078		((intintLCGQuadraticOpenCompactHash_TableData *)
4079		tableData)->numBuckets)) {
	4077	} else if ((index == c && iteration > 0)) {
4080	4078	exitCode = HASH_EXIT_CODE_CYCLE;
4081	4079	break;
4082	4080	}
…	…	@@ -4093,12 +4091,13 @@ int intintLCGQuadraticOpenCompactHash_InnerQuerySingle(char *tableData, int key,
4093	4091	return exitCode;
4094	4092	}
4095	4093	}
4096		int intintLCGQuadraticOpenCompactHash_InnerQuery(char *tableData,
4097		unsigned int numKeys,
4098		int keys, int valuesOutput) {
4099		intintLCGQuadraticOpenCompactHash_Bucket *buckets =
4100		(intintLCGQuadraticOpenCompactHash_Bucket *) &
4101		tableData[sizeof(intintLCGQuadraticOpenCompactHash_TableData)];
	4094	int intintLCGLinearOpenCompactOpenMPHash_InnerQuery(char *tableData,
	4095	unsigned int numKeys,
	4096	int *keys,
	4097	int *valuesOutput) {
	4098	intintLCGLinearOpenCompactOpenMPHash_Bucket *buckets =
	4099	(intintLCGLinearOpenCompactOpenMPHash_Bucket *) &
	4100	tableData[sizeof(intintLCGLinearOpenCompactOpenMPHash_TableData)];
4102	4101	int key;
4103	4102	int *valueOutput;
4104	4103	int index;
…	…	@@ -4108,17 +4107,18 @@ int intintLCGQuadraticOpenCompactHash_InnerQuery(char *tableData,
4108	4107	for (i = 0; i < numKeys; i++) {
4109	4108	key = keys[i];
4110	4109	valueOutput = &valuesOutput[i];
4111		intintLCGQuadraticOpenCompactHash_TableData *mytableData =
4112		(intintLCGQuadraticOpenCompactHash_TableData *) tableData;
	4110	intintLCGLinearOpenCompactOpenMPHash_TableData *mytableData =
	4111	(intintLCGLinearOpenCompactOpenMPHash_TableData *)
	4112	tableData;
4113	4113	intintHash_CompressLCGData compressFuncData =
4114	4114	mytableData->compressFuncData;
4115	4115	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
4116	4116	unsigned long int iteration = 0;
4117	4117	for (;;) {
4118	4118	index =
4119		((1 * iteration * iteration + 0 * iteration +
	4119	((1 * iteration +
4120	4120	c) %
4121		((intintLCGQuadraticOpenCompactHash_TableData *)
	4121	((intintLCGLinearOpenCompactOpenMPHash_TableData *)
4122	4122	tableData)->numBuckets);
4123	4123	if ((buckets[index].key) == HASH_BUCKET_STATUS_EMPTY) {
4124	4124	exitCode = HASH_SEARCH_CODE_EMPTY;
…	…	@@ -4126,10 +4126,7 @@ int intintLCGQuadraticOpenCompactHash_InnerQuery(char *tableData,
4126	4126	} else if (key == buckets[index].key) {
4127	4127	exitCode = HASH_SEARCH_CODE_MATCH;
4128	4128	break;
4129		} else
4130		if ((iteration >
4131		((intintLCGQuadraticOpenCompactHash_TableData
4132		*) tableData)->numBuckets)) {
	4129	} else if ((index == c && iteration > 0)) {
4133	4130	exitCode = HASH_EXIT_CODE_CYCLE;
4134	4131	break;
4135	4132	}
…	…	@@ -4149,24 +4146,24 @@ int intintLCGQuadraticOpenCompactHash_InnerQuery(char *tableData,
4149	4146	}
4150	4147	return resultExitCode;
4151	4148	}
4152		int intintLCGQuadraticOpenCompactHash_InnerInsertSingle(char *tableData,
4153		int key, int value) {
4154		intintLCGQuadraticOpenCompactHash_Bucket *buckets =
4155		(intintLCGQuadraticOpenCompactHash_Bucket *) &
4156		tableData[sizeof(intintLCGQuadraticOpenCompactHash_TableData)];
	4149	int intintLCGLinearOpenCompactOpenMPHash_InnerInsertSingle(char *tableData,
	4150	int key, int value) {
	4151	intintLCGLinearOpenCompactOpenMPHash_Bucket *buckets =
	4152	(intintLCGLinearOpenCompactOpenMPHash_Bucket *) &
	4153	tableData[sizeof(intintLCGLinearOpenCompactOpenMPHash_TableData)];
4157	4154	int index;
4158	4155	int exitCode;
4159		intintLCGQuadraticOpenCompactHash_TableData *mytableData =
4160		(intintLCGQuadraticOpenCompactHash_TableData *) tableData;
	4156	intintLCGLinearOpenCompactOpenMPHash_TableData *mytableData =
	4157	(intintLCGLinearOpenCompactOpenMPHash_TableData *) tableData;
4161	4158	intintHash_CompressLCGData compressFuncData =
4162	4159	mytableData->compressFuncData;
4163	4160	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
4164	4161	unsigned long int iteration = 0;
4165	4162	for (;;) {
4166	4163	index =
4167		((1 * iteration * iteration + 0 * iteration +
	4164	((1 * iteration +
4168	4165	c) %
4169		((intintLCGQuadraticOpenCompactHash_TableData *)
	4166	((intintLCGLinearOpenCompactOpenMPHash_TableData *)
4170	4167	tableData)->numBuckets);
4171	4168	if (((buckets[index].key ==
4172	4169	HASH_BUCKET_STATUS_EMPTY) ? (buckets[index].key =
…	…	@@ -4178,10 +4175,7 @@ int intintLCGQuadraticOpenCompactHash_InnerInsertSingle(char *tableData,
4178	4175	} else if (key == buckets[index].key) {
4179	4176	exitCode = HASH_SEARCH_CODE_MATCH;
4180	4177	break;
4181		} else
4182		if ((iteration >
4183		((intintLCGQuadraticOpenCompactHash_TableData *)
4184		tableData)->numBuckets)) {
	4178	} else if ((index == c && iteration > 0)) {
4185	4179	exitCode = HASH_EXIT_CODE_CYCLE;
4186	4180	break;
4187	4181	}
…	…	@@ -4199,12 +4193,12 @@ int intintLCGQuadraticOpenCompactHash_InnerInsertSingle(char *tableData,
4199	4193	return exitCode;
4200	4194	}
4201	4195	}
4202		int intintLCGQuadraticOpenCompactHash_InnerInsert(char *tableData,
4203		unsigned int numEntries,
4204		int keys, int values) {
4205		intintLCGQuadraticOpenCompactHash_Bucket *buckets =
4206		(intintLCGQuadraticOpenCompactHash_Bucket *) &
4207		tableData[sizeof(intintLCGQuadraticOpenCompactHash_TableData)];
	4196	int intintLCGLinearOpenCompactOpenMPHash_InnerInsert(char *tableData,
	4197	unsigned int numEntries,
	4198	int keys, int values) {
	4199	intintLCGLinearOpenCompactOpenMPHash_Bucket *buckets =
	4200	(intintLCGLinearOpenCompactOpenMPHash_Bucket *) &
	4201	tableData[sizeof(intintLCGLinearOpenCompactOpenMPHash_TableData)];
4208	4202	int resultExitCode = HASH_EXIT_CODE_NORMAL;
4209	4203	int key;
4210	4204	int value;
…	…	@@ -4213,17 +4207,18 @@ int intintLCGQuadraticOpenCompactHash_InnerInsert(char *tableData,
4213	4207	uint i;
4214	4208	for (i = 0; i < numEntries; i++) {
4215	4209	key = keys[i];
4216		intintLCGQuadraticOpenCompactHash_TableData *mytableData =
4217		(intintLCGQuadraticOpenCompactHash_TableData *) tableData;
	4210	intintLCGLinearOpenCompactOpenMPHash_TableData *mytableData =
	4211	(intintLCGLinearOpenCompactOpenMPHash_TableData *)
	4212	tableData;
4218	4213	intintHash_CompressLCGData compressFuncData =
4219	4214	mytableData->compressFuncData;
4220	4215	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
4221	4216	unsigned long int iteration = 0;
4222	4217	for (;;) {
4223	4218	index =
4224		((1 * iteration * iteration + 0 * iteration +
	4219	((1 * iteration +
4225	4220	c) %
4226		((intintLCGQuadraticOpenCompactHash_TableData *)
	4221	((intintLCGLinearOpenCompactOpenMPHash_TableData *)
4227	4222	tableData)->numBuckets);
4228	4223	if (((buckets[index].key ==
4229	4224	HASH_BUCKET_STATUS_EMPTY) ? (buckets[index].key =
…	…	@@ -4236,10 +4231,7 @@ int intintLCGQuadraticOpenCompactHash_InnerInsert(char *tableData,
4236	4231	} else if (key == buckets[index].key) {
4237	4232	exitCode = HASH_SEARCH_CODE_MATCH;
4238	4233	break;
4239		} else
4240		if ((iteration >
4241		((intintLCGQuadraticOpenCompactHash_TableData
4242		*) tableData)->numBuckets)) {
	4234	} else if ((index == c && iteration > 0)) {
4243	4235	exitCode = HASH_EXIT_CODE_CYCLE;
4244	4236	break;
4245	4237	}
…	…	@@ -4258,26 +4250,27 @@ int intintLCGQuadraticOpenCompactHash_InnerInsert(char *tableData,
4258	4250	}
4259	4251	return resultExitCode;
4260	4252	}
4261		int intintLCGQuadraticOpenCompactHash_InnerInsertSingleNoOverwrite(char
4262		*tableData,
4263		int key,
4264		int value) {
4265		intintLCGQuadraticOpenCompactHash_Bucket *buckets =
4266		(intintLCGQuadraticOpenCompactHash_Bucket *) &
4267		tableData[sizeof(intintLCGQuadraticOpenCompactHash_TableData)];
	4253	int intintLCGLinearOpenCompactOpenMPHash_InnerInsertSingleNoOverwrite(char
	4254	*tableData,
	4255	int key,
	4256	int
	4257	value) {
	4258	intintLCGLinearOpenCompactOpenMPHash_Bucket *buckets =
	4259	(intintLCGLinearOpenCompactOpenMPHash_Bucket *) &
	4260	tableData[sizeof(intintLCGLinearOpenCompactOpenMPHash_TableData)];
4268	4261	int index;
4269	4262	int exitCode;
4270		intintLCGQuadraticOpenCompactHash_TableData *mytableData =
4271		(intintLCGQuadraticOpenCompactHash_TableData *) tableData;
	4263	intintLCGLinearOpenCompactOpenMPHash_TableData *mytableData =
	4264	(intintLCGLinearOpenCompactOpenMPHash_TableData *) tableData;
4272	4265	intintHash_CompressLCGData compressFuncData =
4273	4266	mytableData->compressFuncData;
4274	4267	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
4275	4268	unsigned long int iteration = 0;
4276	4269	for (;;) {
4277	4270	index =
4278		((1 * iteration * iteration + 0 * iteration +
	4271	((1 * iteration +
4279	4272	c) %
4280		((intintLCGQuadraticOpenCompactHash_TableData *)
	4273	((intintLCGLinearOpenCompactOpenMPHash_TableData *)
4281	4274	tableData)->numBuckets);
4282	4275	if (((buckets[index].key ==
4283	4276	HASH_BUCKET_STATUS_EMPTY) ? (buckets[index].key =
…	…	@@ -4289,10 +4282,7 @@ int intintLCGQuadraticOpenCompactHash_InnerInsertSingleNoOverwrite(char
4289	4282	} else if (key == buckets[index].key) {
4290	4283	exitCode = HASH_SEARCH_CODE_MATCH;
4291	4284	break;
4292		} else
4293		if ((iteration >
4294		((intintLCGQuadraticOpenCompactHash_TableData *)
4295		tableData)->numBuckets)) {
	4285	} else if ((index == c && iteration > 0)) {
4296	4286	exitCode = HASH_EXIT_CODE_CYCLE;
4297	4287	break;
4298	4288	}
…	…	@@ -4309,14 +4299,14 @@ int intintLCGQuadraticOpenCompactHash_InnerInsertSingleNoOverwrite(char
4309	4299	return exitCode;
4310	4300	}
4311	4301	}
4312		int intintLCGQuadraticOpenCompactHash_InnerInsertNoOverwrite(char *tableData,
4313		unsigned int
4314		numEntries,
4315		int *keys,
4316		int *values) {
4317		intintLCGQuadraticOpenCompactHash_Bucket *buckets =
4318		(intintLCGQuadraticOpenCompactHash_Bucket *) &
4319		tableData[sizeof(intintLCGQuadraticOpenCompactHash_TableData)];
	4302	int intintLCGLinearOpenCompactOpenMPHash_InnerInsertNoOverwrite(char *tableData,
	4303	unsigned int
	4304	numEntries,
	4305	int *keys,
	4306	int *values) {
	4307	intintLCGLinearOpenCompactOpenMPHash_Bucket *buckets =
	4308	(intintLCGLinearOpenCompactOpenMPHash_Bucket *) &
	4309	tableData[sizeof(intintLCGLinearOpenCompactOpenMPHash_TableData)];
4320	4310	int resultExitCode = HASH_EXIT_CODE_NORMAL;
4321	4311	int key;
4322	4312	int value;
…	…	@@ -4325,17 +4315,18 @@ int intintLCGQuadraticOpenCompactHash_InnerInsertNoOverwrite(char *tableData,
4325	4315	uint i;
4326	4316	for (i = 0; i < numEntries; i++) {
4327	4317	key = keys[i];
4328		intintLCGQuadraticOpenCompactHash_TableData *mytableData =
4329		(intintLCGQuadraticOpenCompactHash_TableData *) tableData;
	4318	intintLCGLinearOpenCompactOpenMPHash_TableData *mytableData =
	4319	(intintLCGLinearOpenCompactOpenMPHash_TableData *)
	4320	tableData;
4330	4321	intintHash_CompressLCGData compressFuncData =
4331	4322	mytableData->compressFuncData;
4332	4323	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
4333	4324	unsigned long int iteration = 0;
4334	4325	for (;;) {
4335	4326	index =
4336		((1 * iteration * iteration + 0 * iteration +
	4327	((1 * iteration +
4337	4328	c) %
4338		((intintLCGQuadraticOpenCompactHash_TableData *)
	4329	((intintLCGLinearOpenCompactOpenMPHash_TableData *)
4339	4330	tableData)->numBuckets);
4340	4331	if (((buckets[index].key ==
4341	4332	HASH_BUCKET_STATUS_EMPTY) ? (buckets[index].key =
…	…	@@ -4348,10 +4339,7 @@ int intintLCGQuadraticOpenCompactHash_InnerInsertNoOverwrite(char *tableData,
4348	4339	} else if (key == buckets[index].key) {
4349	4340	exitCode = HASH_SEARCH_CODE_MATCH;
4350	4341	break;
4351		} else
4352		if ((iteration >
4353		((intintLCGQuadraticOpenCompactHash_TableData
4354		*) tableData)->numBuckets)) {
	4342	} else if ((index == c && iteration > 0)) {
4355	4343	exitCode = HASH_EXIT_CODE_CYCLE;
4356	4344	break;
4357	4345	}
…	…	@@ -4371,302 +4359,768 @@ int intintLCGQuadraticOpenCompactHash_InnerInsertNoOverwrite(char *tableData,
4371	4359	}
4372	4360	return resultExitCode;
4373	4361	}
4374		int intintLCGQuadraticOpenCompactHash_QuerySingle(intintHash_Table * table,
4375		int key, int *valueOutput) {
4376		return intintLCGQuadraticOpenCompactHash_InnerQuerySingle(table->
4377		tableData,
4378		key,
4379		valueOutput);
	4362	int intintLCGLinearOpenCompactOpenMPHash_QuerySingle(intintHash_Table * table,
	4363	int key,
	4364	int *valueOutput) {
	4365	return intintLCGLinearOpenCompactOpenMPHash_InnerQuerySingle(table->
	4366	tableData,
	4367	key,
	4368	valueOutput);
4380	4369	}
4381		int intintLCGQuadraticOpenCompactHash_Query(intintHash_Table * table,
4382		size_t numKeys, int *keys,
4383		int *valuesOutput) {
4384		return intintLCGQuadraticOpenCompactHash_InnerQuery(table->tableData,
4385		numKeys, keys,
4386		valuesOutput);
	4370	int intintLCGLinearOpenCompactOpenMPHash_Query(intintHash_Table * table,
	4371	size_t numKeys, int *keys,
	4372	int *valuesOutput) {
	4373	return intintLCGLinearOpenCompactOpenMPHash_InnerQuery(table->tableData,
	4374	numKeys, keys,
	4375	valuesOutput);
4387	4376	}
4388		int intintLCGQuadraticOpenCompactHash_InsertSingle(intintHash_Table * table,
4389		int key, int value) {
4390		return intintLCGQuadraticOpenCompactHash_InnerInsertSingle(table->
4391		tableData,
4392		key, value);
	4377	int intintLCGLinearOpenCompactOpenMPHash_InsertSingle(intintHash_Table * table,
	4378	int key, int value) {
	4379	return intintLCGLinearOpenCompactOpenMPHash_InnerInsertSingle(table->
	4380	tableData,
	4381	key,
	4382	value);
4393	4383	}
4394		int intintLCGQuadraticOpenCompactHash_Insert(intintHash_Table * table,
4395		size_t numEntries, int *keys,
4396		int *values) {
4397		return intintLCGQuadraticOpenCompactHash_InnerInsert(table->tableData,
4398		numEntries, keys,
4399		values);
	4384	int intintLCGLinearOpenCompactOpenMPHash_Insert(intintHash_Table * table,
	4385	size_t numEntries, int *keys,
	4386	int *values) {
	4387	return intintLCGLinearOpenCompactOpenMPHash_InnerInsert(table->
	4388	tableData,
	4389	numEntries,
	4390	keys, values);
4400	4391	}
4401		int intintLCGQuadraticOpenCompactHash_InsertSingleNoOverwrite(intintHash_Table *
4402		table, int key,
4403		int value) {
	4392	int
	4393	intintLCGLinearOpenCompactOpenMPHash_InsertSingleNoOverwrite(intintHash_Table *
	4394	table, int key,
	4395	int value) {
4404	4396	return
4405		intintLCGQuadraticOpenCompactHash_InnerInsertSingleNoOverwrite
	4397	intintLCGLinearOpenCompactOpenMPHash_InnerInsertSingleNoOverwrite
4406	4398	(table->tableData, key, value);
4407	4399	}
4408		int intintLCGQuadraticOpenCompactHash_InsertNoOverwrite(intintHash_Table *
4409		table,
4410		size_t numEntries,
4411		int *keys,
4412		int *values) {
4413		return intintLCGQuadraticOpenCompactHash_InnerInsertNoOverwrite(table->
	4400	int intintLCGLinearOpenCompactOpenMPHash_InsertNoOverwrite(intintHash_Table *
	4401	table,
	4402	size_t numEntries,
	4403	int *keys,
	4404	int *values) {
	4405	return
	4406	intintLCGLinearOpenCompactOpenMPHash_InnerInsertNoOverwrite(table->
4414	4407	tableData,
4415	4408	numEntries,
4416	4409	keys,
4417	4410	values);
4418	4411	}
4419	4412
4420		typedef struct intintLCGQuadraticOpenCompactCLHash_TableData {
	4413	typedef struct intintLCGQuadraticOpenCompactHash_TableData {
4421	4414	int hashID;
4422	4415	unsigned int numBuckets;
4423	4416	intintHash_CompressLCGData compressFuncData;
4424		} intintLCGQuadraticOpenCompactCLHash_TableData;
4425		typedef struct intintLCGQuadraticOpenCompactCLHash_Bucket {
	4417	} intintLCGQuadraticOpenCompactHash_TableData;
	4418	typedef struct intintLCGQuadraticOpenCompactHash_Bucket {
4426	4419	int key;
4427	4420	int value;
4428		} intintLCGQuadraticOpenCompactCLHash_Bucket;
	4421	} intintLCGQuadraticOpenCompactHash_Bucket;
4429	4422	intintHash_Table
4430		intintLCGQuadraticOpenCompactCLHash_CreateTable(intintHash_Factory
4431		factory, int hashIndex,
4432		size_t keyRange,
4433		size_t numEntries,
4434		float loadFactor) {
	4423	intintLCGQuadraticOpenCompactHash_CreateTable(intintHash_Factory factory,
	4424	int hashIndex,
	4425	size_t keyRange,
	4426	size_t numEntries,
	4427	float loadFactor) {
4435	4428	intintHash_Table *table =
4436	4429	(intintHash_Table *) malloc(sizeof(intintHash_Table));
4437		table->destroyFunc = &intintLCGQuadraticOpenCompactCLHash_DestroyTable;
4438		table->emptyFunc = &intintLCGQuadraticOpenCompactCLHash_EmptyTable;
4439		table->queryFunc = &intintLCGQuadraticOpenCompactCLHash_Query;
4440		table->querySingleFunc =
4441		&intintLCGQuadraticOpenCompactCLHash_QuerySingle;
4442		table->insertFunc = &intintLCGQuadraticOpenCompactCLHash_Insert;
	4430	table->destroyFunc = &intintLCGQuadraticOpenCompactHash_DestroyTable;
	4431	table->emptyFunc = &intintLCGQuadraticOpenCompactHash_EmptyTable;
	4432	table->queryFunc = &intintLCGQuadraticOpenCompactHash_Query;
	4433	table->querySingleFunc = &intintLCGQuadraticOpenCompactHash_QuerySingle;
	4434	table->insertFunc = &intintLCGQuadraticOpenCompactHash_Insert;
4443	4435	table->insertSingleFunc =
4444		&intintLCGQuadraticOpenCompactCLHash_InsertSingle;
	4436	&intintLCGQuadraticOpenCompactHash_InsertSingle;
4445	4437	table->insertNoOverwriteFunc =
4446		&intintLCGQuadraticOpenCompactCLHash_InsertNoOverwrite;
	4438	&intintLCGQuadraticOpenCompactHash_InsertNoOverwrite;
4447	4439	table->insertSingleNoOverwriteFunc =
4448		&intintLCGQuadraticOpenCompactCLHash_InsertSingleNoOverwrite;
	4440	&intintLCGQuadraticOpenCompactHash_InsertSingleNoOverwrite;
4449	4441	table->tableData =
4450		(char *)
4451		malloc(sizeof(intintLCGQuadraticOpenCompactCLHash_TableData));
4452		((intintLCGQuadraticOpenCompactCLHash_TableData *) table->tableData)->
4453		hashID = LCG_QUADRATIC_OPEN_COMPACT_CL_HASH_ID;
4454		table->context = factory->context;
4455		table->queue = factory->queue;
4456		table->program = factory->program;
4457		table->localWorkSize = factory->localWorkSize;
4458		table->utilProgram = factory->utilProgram[hashIndex];
4459		table->emptyKernel = factory->emptyKernel[hashIndex];
4460		table->emptyKernelLocalWorkSize =
4461		factory->emptyKernelLocalWorkSize[hashIndex];
4462		table->querySingleKernel = factory->querySingleKernel[hashIndex];
4463		table->insertSingleKernel = factory->insertSingleKernel[hashIndex];
4464		table->insertSingleNoOverwriteKernel =
4465		factory->insertSingleNoOverwriteKernel[hashIndex];
4466		clRetainContext(table->context);
4467		clRetainCommandQueue(table->queue);
4468		clRetainProgram(table->program);
4469		clRetainProgram(table->utilProgram);
4470		clRetainKernel(table->emptyKernel);
4471		clRetainKernel(table->querySingleKernel);
4472		clRetainKernel(table->insertSingleKernel);
4473		clRetainKernel(table->insertSingleNoOverwriteKernel);;
4474		((intintLCGQuadraticOpenCompactCLHash_TableData *) table->tableData)->
	4442	(char *)malloc(sizeof(intintLCGQuadraticOpenCompactHash_TableData));
	4443	((intintLCGQuadraticOpenCompactHash_TableData *) table->tableData)->
	4444	hashID = LCG_QUADRATIC_OPEN_COMPACT_HASH_ID;
	4445	((intintLCGQuadraticOpenCompactHash_TableData *) table->tableData)->
4475	4446	numBuckets = (unsigned int)((double)numEntries / loadFactor);
4476		((intintLCGQuadraticOpenCompactCLHash_TableData *) table->tableData)->
	4447	((intintLCGQuadraticOpenCompactHash_TableData *) table->tableData)->
4477	4448	compressFuncData.a = HASH_LCG_A;
4478		((intintLCGQuadraticOpenCompactCLHash_TableData *) table->tableData)->
	4449	((intintLCGQuadraticOpenCompactHash_TableData *) table->tableData)->
4479	4450	compressFuncData.c = HASH_LCG_C;
4480		((intintLCGQuadraticOpenCompactCLHash_TableData *) table->tableData)->
	4451	((intintLCGQuadraticOpenCompactHash_TableData *) table->tableData)->
4481	4452	compressFuncData.m = HASH_LCG_M;
4482		((intintLCGQuadraticOpenCompactCLHash_TableData *) table->tableData)->
	4453	((intintLCGQuadraticOpenCompactHash_TableData *) table->tableData)->
4483	4454	compressFuncData.n =
4484		((intintLCGQuadraticOpenCompactCLHash_TableData *) table->
4485		tableData)->numBuckets;
4486		((intintLCGQuadraticOpenCompactCLHash_TableData *) table->tableData)->
	4455	((intintLCGQuadraticOpenCompactHash_TableData *) table->tableData)->
	4456	numBuckets;
	4457	((intintLCGQuadraticOpenCompactHash_TableData *) table->tableData)->
4487	4458	numBuckets =
4488		largestProthPrimeUnder(((intintLCGQuadraticOpenCompactCLHash_TableData *) table->tableData)->numBuckets);
	4459	largestProthPrimeUnder(((intintLCGQuadraticOpenCompactHash_TableData
	4460	*) table->tableData)->numBuckets);
4489	4461	char *tempHashData =
4490		(char *)malloc(sizeof(intintLCGQuadraticOpenCompactCLHash_TableData)
4491		+
4492		((intintLCGQuadraticOpenCompactCLHash_TableData *)
	4462	(char *)malloc(sizeof(intintLCGQuadraticOpenCompactHash_TableData) +
	4463	((intintLCGQuadraticOpenCompactHash_TableData *)
4493	4464	table->tableData)->numBuckets *
4494		sizeof(intintLCGQuadraticOpenCompactCLHash_Bucket));
	4465	sizeof(intintLCGQuadraticOpenCompactHash_Bucket));
4495	4466	memcpy(tempHashData, table->tableData,
4496		sizeof(intintLCGQuadraticOpenCompactCLHash_TableData));
	4467	sizeof(intintLCGQuadraticOpenCompactHash_TableData));
4497	4468	free(table->tableData);
4498	4469	table->tableData = tempHashData;
4499		cl_int err;
4500		table->tableDataBuffer =
4501		clCreateBuffer(table->context, CL_MEM_READ_WRITE,
4502		sizeof(intintLCGQuadraticOpenCompactHash_TableData) +
4503		((intintLCGQuadraticOpenCompactHash_TableData *)
4504		table->tableData)->numBuckets *
4505		sizeof(intintLCGQuadraticOpenCompactHash_Bucket),
4506		NULL, &err);
4507		CLHash_Utilities_HandleError(err,
4508		"intintLCGQuadraticOpenCompactCLHash_InitTable",
4509		"clCreateBuffer");
4510		err =
4511		clEnqueueWriteBuffer(table->queue, table->tableDataBuffer, CL_TRUE,
4512		0,
4513		sizeof
4514		(intintLCGQuadraticOpenCompactHash_TableData),
4515		table->tableData, 0, NULL, NULL);
4516		CLHash_Utilities_HandleError(err,
4517		"intintLCGQuadraticOpenCompactCLHash_InitTable",
4518		"clEnqueueWriteBuffer");
4519	4470	return table;
4520	4471	}
4521		int intintLCGQuadraticOpenCompactCLHash_CreateFactory(intintHash_Factory *
4522		factory, int hashIndex) {
	4472	int intintLCGQuadraticOpenCompactHash_CreateFactory(intintHash_Factory *
	4473	factory, int hashIndex) {
4523	4474	factory->createFunc[hashIndex] =
4524		&intintLCGQuadraticOpenCompactCLHash_CreateTable;
	4475	&intintLCGQuadraticOpenCompactHash_CreateTable;
4525	4476	factory->destroyFunc[hashIndex] =
4526		&intintLCGQuadraticOpenCompactCLHash_DestroyFactory;
4527		cl_int error;
4528		cl_device_id device;
4529		error =
4530		clGetContextInfo(factory->context, CL_CONTEXT_DEVICES,
4531		sizeof(device), &device, NULL);
4532		CLHash_Utilities_HandleError(error, "intintHash_CreateFactory",
4533		"clGetContextInfo");
4534		factory->querySingleKernel[hashIndex] =
4535		clCreateKernel(factory->program,
4536		"intintLCGQuadraticOpenCompactCLHash_RangeQuerySingle",
4537		&error);
4538		CLHash_Utilities_HandleError(error,
4539		"intintLCGQuadraticOpenCompactCLHash_CreateFactory",
4540		"clCreateKernel");
4541		factory->insertSingleKernel[hashIndex] =
4542		clCreateKernel(factory->program,
4543		"intintLCGQuadraticOpenCompactCLHash_RangeInsertSingle",
4544		&error);
4545		CLHash_Utilities_HandleError(error,
4546		"intintLCGQuadraticOpenCompactCLHash_CreateFactory",
4547		"clCreateKernel");
4548		factory->insertSingleNoOverwriteKernel[hashIndex] =
4549		clCreateKernel(factory->program,
4550		"intintLCGQuadraticOpenCompactCLHash_RangeInsertSingleNoOverwrite",
4551		&error);
4552		CLHash_Utilities_HandleError(error,
4553		"intintLCGQuadraticOpenCompactCLHash_CreateFactory",
4554		"clCreateKernel");
4555		factory->utilProgram[hashIndex] =
4556		CLHash_Utilities_BuildProgramString(factory->context, device,
4557		"static inline unsigned int intintHash_CompressIdentity(char data, int hashCode){ return hashCode; } typedef struct intintHash_CompressLCGData{ long unsigned int a; long unsigned int c; unsigned int m; unsigned int n; }intintHash_CompressLCGData; static inline unsigned int intintHash_CompressLCG(intintHash_CompressLCGData compressLCGData, int hashCode){ return ((compressLCGData.a * hashCode + compressLCGData.c) % compressLCGData.m) % compressLCGData.n; } typedef struct intintLCGQuadraticOpenCompactCLHash_TableData{ int hashID; unsigned int numBuckets; intintHash_CompressLCGData compressFuncData; }intintLCGQuadraticOpenCompactCLHash_TableData; typedef struct intintLCGQuadraticOpenCompactCLHash_Bucket{ int key; int value; }intintLCGQuadraticOpenCompactCLHash_Bucket; __kernel void intintLCGQuadraticOpenCompactCLHash_Empty(__global char tableData){ int index = get_global_id(0); if(index >= ((__global intintLCGQuadraticOpenCompactCLHash_TableData)tableData)->numBuckets){ return; } __global intintLCGQuadraticOpenCompactCLHash_Bucket buckets = (__global intintLCGQuadraticOpenCompactCLHash_Bucket)&tableData[sizeof(intintLCGQuadraticOpenCompactCLHash_TableData)]; buckets[index].key = -1;/HASH_BUCKET_STATUS_EMPTY/ }");
4558		factory->emptyKernel[hashIndex] =
4559		clCreateKernel(factory->utilProgram[hashIndex],
4560		"intintLCGQuadraticOpenCompactCLHash_Empty", &error);
4561		CLHash_Utilities_HandleError(error,
4562		"intintLCGQuadraticOpenCompactCLHash_CreateFactory",
4563		"clCreateKernel");
4564		error =
4565		clGetKernelWorkGroupInfo(factory->emptyKernel[hashIndex], device,
4566		CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t),
4567		&factory->
4568		emptyKernelLocalWorkSize[hashIndex], NULL);
4569		CLHash_Utilities_HandleError(error,
4570		"intintLCGQuadraticOpenCompactCLHash_CreateFactory",
4571		"clGetKernelWorkGroupInfo");;;
	4477	&intintLCGQuadraticOpenCompactHash_DestroyFactory;;
4572	4478	return HASH_EXIT_CODE_NORMAL;
4573	4479	}
4574		int intintLCGQuadraticOpenCompactCLHash_DestroyFactory(intintHash_Factory *
4575		factory,
4576		int hashIndex) {;
4577		clReleaseKernel(factory->emptyKernel[hashIndex]);
4578		clReleaseProgram(factory->utilProgram[hashIndex]);
4579		clReleaseKernel(factory->querySingleKernel[hashIndex]);
4580		clReleaseKernel(factory->insertSingleKernel[hashIndex]);
4581		clReleaseKernel(factory->insertSingleNoOverwriteKernel[hashIndex]);;
	4480	int intintLCGQuadraticOpenCompactHash_DestroyFactory(intintHash_Factory *
	4481	factory, int hashIndex) {;
4582	4482	return HASH_EXIT_CODE_NORMAL;
4583	4483	}
4584		int intintLCGQuadraticOpenCompactCLHash_DestroyTable(intintHash_Table * table) {
	4484	int intintLCGQuadraticOpenCompactHash_DestroyTable(intintHash_Table * table) {
4585	4485	int exitCode = 0;
4586		clReleaseMemObject(table->tableDataBuffer);
4587		clReleaseContext(table->context);
4588		clReleaseCommandQueue(table->queue);
4589		clReleaseProgram(table->utilProgram);
4590		clReleaseKernel(table->emptyKernel);
4591		clReleaseProgram(table->program);
4592		clReleaseKernel(table->querySingleKernel);
4593		clReleaseKernel(table->insertSingleKernel);
4594		clReleaseKernel(table->insertSingleNoOverwriteKernel);
4595	4486	free(table->tableData);
4596	4487	free(table);
4597	4488	return exitCode;
4598	4489	}
4599		int intintLCGQuadraticOpenCompactCLHash_EmptyTable(intintHash_Table * table) {
	4490	int intintLCGQuadraticOpenCompactHash_EmptyTable(intintHash_Table * table) {
4600	4491	int exitCode = 0;
4601		cl_int err;
4602		err =
4603		clSetKernelArg(table->emptyKernel, 0, sizeof(cl_mem),
4604		&table->tableDataBuffer);
4605		CLHash_Utilities_HandleError(err,
4606		"intintLCGQuadraticOpenCompactCLHash_EmptyTable",
4607		"clSetKernelArg");
4608		const size_t groupWorkSize =
4609		roundUpToNearest(((intintLCGQuadraticOpenCompactHash_TableData *)
4610		table->tableData)->numBuckets,
4611		table->emptyKernelLocalWorkSize);
4612		err =
4613		clEnqueueNDRangeKernel(table->queue, table->emptyKernel, 1, 0,
4614		&groupWorkSize,
4615		(const size_t *)&table->
4616		emptyKernelLocalWorkSize, 0, NULL, NULL);
4617		CLHash_Utilities_HandleError(err,
4618		"intintLCGQuadraticOpenCompactCLHash_EmptyTable",
4619		"clEnqueueNDRangeKernel");
4620		exitCode = HASH_EXIT_CODE_NORMAL;;
	4492	intintLCGQuadraticOpenCompactHash_Bucket *buckets =
	4493	(intintLCGQuadraticOpenCompactHash_Bucket *) & table->
	4494	tableData[sizeof(intintLCGQuadraticOpenCompactHash_TableData)];
	4495	int index;
	4496	for (index = 0;
	4497	index <
	4498	((intintLCGQuadraticOpenCompactHash_TableData *) table->
	4499	tableData)->numBuckets; index++) {
	4500	buckets[index].key = HASH_BUCKET_STATUS_EMPTY;
	4501	}
	4502	exitCode = HASH_EXIT_CODE_NORMAL;
4621	4503	return exitCode;
4622	4504	}
4623		int intintLCGQuadraticOpenCompactCLHash_QuerySingle(intintHash_Table * table,
4624		int key, int *valueOutput) {
4625		return intintLCGQuadraticOpenCompactCLHash_Query(table, 1, &key,
4626		valueOutput);
4627		}
4628		int intintLCGQuadraticOpenCompactCLHash_Query(intintHash_Table * table,
4629		size_t numKeys, int *keys,
4630		int *valuesOutput) {
4631		cl_int err;
4632		cl_mem keysBuffer =
4633		clCreateBuffer(table->context,
4634		CL_MEM_READ_ONLY \| CL_MEM_COPY_HOST_PTR,
4635		sizeof(int) * numKeys, keys, &err);
4636		CLHash_Utilities_HandleError(err,
4637		"intintLCGQuadraticOpenCompactCLHash_Query",
4638		"clCreateBuffer");
4639		cl_mem valuesOutputBuffer =
4640		clCreateBuffer(table->context, CL_MEM_WRITE_ONLY,
4641		sizeof(int) * numKeys, NULL, &err);
4642		CLHash_Utilities_HandleError(err,
4643		"intintLCGQuadraticOpenCompactCLHash_Query",
4644		"clCreateBuffer");
4645		intintLCGQuadraticOpenCompactCLHash_BufferQuery(table, numKeys,
4646		keysBuffer,
4647		valuesOutputBuffer);
4648		err =
4649		clEnqueueReadBuffer(table->queue, valuesOutputBuffer, CL_TRUE, 0,
4650		sizeof(int) * numKeys, valuesOutput, 0, NULL,
4651		NULL);
4652		CLHash_Utilities_HandleError(err,
4653		"intintLCGQuadraticOpenCompactCLHash_Query",
4654		"clEnqueueReadBuffer");
4655		clReleaseMemObject(keysBuffer);
4656		clReleaseMemObject(valuesOutputBuffer);
4657		return HASH_EXIT_CODE_NORMAL;
	4505	int intintLCGQuadraticOpenCompactHash_InnerQuerySingle(char *tableData, int key,
	4506	int *valueOutput) {
	4507	intintLCGQuadraticOpenCompactHash_Bucket *buckets =
	4508	(intintLCGQuadraticOpenCompactHash_Bucket *) &
	4509	tableData[sizeof(intintLCGQuadraticOpenCompactHash_TableData)];
	4510	int index;
	4511	int exitCode;
	4512	intintLCGQuadraticOpenCompactHash_TableData *mytableData =
	4513	(intintLCGQuadraticOpenCompactHash_TableData *) tableData;
	4514	intintHash_CompressLCGData compressFuncData =
	4515	mytableData->compressFuncData;
	4516	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
	4517	unsigned long int iteration = 0;
	4518	for (;;) {
	4519	index =
	4520	((1 * iteration * iteration + 0 * iteration +
	4521	c) %
	4522	((intintLCGQuadraticOpenCompactHash_TableData *)
	4523	tableData)->numBuckets);
	4524	if ((buckets[index].key) == HASH_BUCKET_STATUS_EMPTY) {
	4525	exitCode = HASH_SEARCH_CODE_EMPTY;
	4526	break;
	4527	} else if (key == buckets[index].key) {
	4528	exitCode = HASH_SEARCH_CODE_MATCH;
	4529	break;
	4530	} else
	4531	if ((iteration >
	4532	((intintLCGQuadraticOpenCompactHash_TableData *)
	4533	tableData)->numBuckets)) {
	4534	exitCode = HASH_EXIT_CODE_CYCLE;
	4535	break;
	4536	}
	4537	iteration++;
	4538	}
	4539	switch (exitCode) {
	4540	case HASH_SEARCH_CODE_MATCH:
	4541	*valueOutput = buckets[index].value;
	4542	return HASH_EXIT_CODE_NORMAL;
	4543	case HASH_SEARCH_CODE_MISMATCH:
	4544	case HASH_SEARCH_CODE_EMPTY:
	4545	return HASH_EXIT_CODE_KEY_DNE;
	4546	default:
	4547	return exitCode;
	4548	}
4658	4549	}
4659		int intintLCGQuadraticOpenCompactCLHash_BufferQuery(intintHash_Table * table,
4660		size_t numKeys,
4661		cl_mem keysBuffer,
4662		cl_mem valuesOutputBuffer) {
4663		cl_int err;
4664		err =
4665		clSetKernelArg(table->querySingleKernel, 0, sizeof(cl_mem),
4666		&table->tableDataBuffer);
4667		CLHash_Utilities_HandleError(err,
4668		"intintLCGQuadraticOpenCompactCLHash_BufferQuery",
4669		"clSetKernelArg");
	4550	int intintLCGQuadraticOpenCompactHash_InnerQuery(char *tableData,
	4551	unsigned int numKeys,
	4552	int keys, int valuesOutput) {
	4553	intintLCGQuadraticOpenCompactHash_Bucket *buckets =
	4554	(intintLCGQuadraticOpenCompactHash_Bucket *) &
	4555	tableData[sizeof(intintLCGQuadraticOpenCompactHash_TableData)];
	4556	int key;
	4557	int *valueOutput;
	4558	int index;
	4559	int exitCode;
	4560	uint i;
	4561	int resultExitCode = HASH_EXIT_CODE_NORMAL;
	4562	for (i = 0; i < numKeys; i++) {
	4563	key = keys[i];
	4564	valueOutput = &valuesOutput[i];
	4565	intintLCGQuadraticOpenCompactHash_TableData *mytableData =
	4566	(intintLCGQuadraticOpenCompactHash_TableData *) tableData;
	4567	intintHash_CompressLCGData compressFuncData =
	4568	mytableData->compressFuncData;
	4569	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
	4570	unsigned long int iteration = 0;
	4571	for (;;) {
	4572	index =
	4573	((1 * iteration * iteration + 0 * iteration +
	4574	c) %
	4575	((intintLCGQuadraticOpenCompactHash_TableData *)
	4576	tableData)->numBuckets);
	4577	if ((buckets[index].key) == HASH_BUCKET_STATUS_EMPTY) {
	4578	exitCode = HASH_SEARCH_CODE_EMPTY;
	4579	break;
	4580	} else if (key == buckets[index].key) {
	4581	exitCode = HASH_SEARCH_CODE_MATCH;
	4582	break;
	4583	} else
	4584	if ((iteration >
	4585	((intintLCGQuadraticOpenCompactHash_TableData
	4586	*) tableData)->numBuckets)) {
	4587	exitCode = HASH_EXIT_CODE_CYCLE;
	4588	break;
	4589	}
	4590	iteration++;
	4591	}
	4592	switch (exitCode) {
	4593	case HASH_SEARCH_CODE_MATCH:
	4594	*valueOutput = buckets[index].value;
	4595	break;
	4596	case HASH_SEARCH_CODE_MISMATCH:
	4597	case HASH_SEARCH_CODE_EMPTY:
	4598	resultExitCode = HASH_EXIT_CODE_KEY_DNE;
	4599	break;
	4600	default:
	4601	return exitCode;
	4602	}
	4603	}
	4604	return resultExitCode;
	4605	}
	4606	int intintLCGQuadraticOpenCompactHash_InnerInsertSingle(char *tableData,
	4607	int key, int value) {
	4608	intintLCGQuadraticOpenCompactHash_Bucket *buckets =
	4609	(intintLCGQuadraticOpenCompactHash_Bucket *) &
	4610	tableData[sizeof(intintLCGQuadraticOpenCompactHash_TableData)];
	4611	int index;
	4612	int exitCode;
	4613	intintLCGQuadraticOpenCompactHash_TableData *mytableData =
	4614	(intintLCGQuadraticOpenCompactHash_TableData *) tableData;
	4615	intintHash_CompressLCGData compressFuncData =
	4616	mytableData->compressFuncData;
	4617	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
	4618	unsigned long int iteration = 0;
	4619	for (;;) {
	4620	index =
	4621	((1 * iteration * iteration + 0 * iteration +
	4622	c) %
	4623	((intintLCGQuadraticOpenCompactHash_TableData *)
	4624	tableData)->numBuckets);
	4625	if (((buckets[index].key ==
	4626	HASH_BUCKET_STATUS_EMPTY) ? (buckets[index].key =
	4627	key,
	4628	HASH_BUCKET_STATUS_EMPTY) :
	4629	buckets[index].key) == HASH_BUCKET_STATUS_EMPTY) {
	4630	exitCode = HASH_SEARCH_CODE_EMPTY;
	4631	break;
	4632	} else if (key == buckets[index].key) {
	4633	exitCode = HASH_SEARCH_CODE_MATCH;
	4634	break;
	4635	} else
	4636	if ((iteration >
	4637	((intintLCGQuadraticOpenCompactHash_TableData *)
	4638	tableData)->numBuckets)) {
	4639	exitCode = HASH_EXIT_CODE_CYCLE;
	4640	break;
	4641	}
	4642	iteration++;
	4643	}
	4644	switch (exitCode) {
	4645	case HASH_SEARCH_CODE_MATCH:
	4646	case HASH_SEARCH_CODE_MISMATCH:
	4647	buckets[index].value = value;
	4648	return HASH_EXIT_CODE_OVERWRITE;
	4649	case HASH_SEARCH_CODE_EMPTY:
	4650	buckets[index].value = value;
	4651	return HASH_EXIT_CODE_NORMAL;
	4652	default:
	4653	return exitCode;
	4654	}
	4655	}
	4656	int intintLCGQuadraticOpenCompactHash_InnerInsert(char *tableData,
	4657	unsigned int numEntries,
	4658	int keys, int values) {
	4659	intintLCGQuadraticOpenCompactHash_Bucket *buckets =
	4660	(intintLCGQuadraticOpenCompactHash_Bucket *) &
	4661	tableData[sizeof(intintLCGQuadraticOpenCompactHash_TableData)];
	4662	int resultExitCode = HASH_EXIT_CODE_NORMAL;
	4663	int key;
	4664	int value;
	4665	int index;
	4666	int exitCode;
	4667	uint i;
	4668	for (i = 0; i < numEntries; i++) {
	4669	key = keys[i];
	4670	intintLCGQuadraticOpenCompactHash_TableData *mytableData =
	4671	(intintLCGQuadraticOpenCompactHash_TableData *) tableData;
	4672	intintHash_CompressLCGData compressFuncData =
	4673	mytableData->compressFuncData;
	4674	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
	4675	unsigned long int iteration = 0;
	4676	for (;;) {
	4677	index =
	4678	((1 * iteration * iteration + 0 * iteration +
	4679	c) %
	4680	((intintLCGQuadraticOpenCompactHash_TableData *)
	4681	tableData)->numBuckets);
	4682	if (((buckets[index].key ==
	4683	HASH_BUCKET_STATUS_EMPTY) ? (buckets[index].key =
	4684	key,
	4685	HASH_BUCKET_STATUS_EMPTY)
	4686	: buckets[index].key) ==
	4687	HASH_BUCKET_STATUS_EMPTY) {
	4688	exitCode = HASH_SEARCH_CODE_EMPTY;
	4689	break;
	4690	} else if (key == buckets[index].key) {
	4691	exitCode = HASH_SEARCH_CODE_MATCH;
	4692	break;
	4693	} else
	4694	if ((iteration >
	4695	((intintLCGQuadraticOpenCompactHash_TableData
	4696	*) tableData)->numBuckets)) {
	4697	exitCode = HASH_EXIT_CODE_CYCLE;
	4698	break;
	4699	}
	4700	iteration++;
	4701	}
	4702	switch (exitCode) {
	4703	case HASH_SEARCH_CODE_MATCH:
	4704	case HASH_SEARCH_CODE_MISMATCH:
	4705	resultExitCode = HASH_EXIT_CODE_OVERWRITE;
	4706	case HASH_SEARCH_CODE_EMPTY:
	4707	buckets[index].value = values[i];
	4708	break;
	4709	default:
	4710	return exitCode;
	4711	}
	4712	}
	4713	return resultExitCode;
	4714	}
	4715	int intintLCGQuadraticOpenCompactHash_InnerInsertSingleNoOverwrite(char
	4716	*tableData,
	4717	int key,
	4718	int value) {
	4719	intintLCGQuadraticOpenCompactHash_Bucket *buckets =
	4720	(intintLCGQuadraticOpenCompactHash_Bucket *) &
	4721	tableData[sizeof(intintLCGQuadraticOpenCompactHash_TableData)];
	4722	int index;
	4723	int exitCode;
	4724	intintLCGQuadraticOpenCompactHash_TableData *mytableData =
	4725	(intintLCGQuadraticOpenCompactHash_TableData *) tableData;
	4726	intintHash_CompressLCGData compressFuncData =
	4727	mytableData->compressFuncData;
	4728	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
	4729	unsigned long int iteration = 0;
	4730	for (;;) {
	4731	index =
	4732	((1 * iteration * iteration + 0 * iteration +
	4733	c) %
	4734	((intintLCGQuadraticOpenCompactHash_TableData *)
	4735	tableData)->numBuckets);
	4736	if (((buckets[index].key ==
	4737	HASH_BUCKET_STATUS_EMPTY) ? (buckets[index].key =
	4738	key,
	4739	HASH_BUCKET_STATUS_EMPTY) :
	4740	buckets[index].key) == HASH_BUCKET_STATUS_EMPTY) {
	4741	exitCode = HASH_SEARCH_CODE_EMPTY;
	4742	break;
	4743	} else if (key == buckets[index].key) {
	4744	exitCode = HASH_SEARCH_CODE_MATCH;
	4745	break;
	4746	} else
	4747	if ((iteration >
	4748	((intintLCGQuadraticOpenCompactHash_TableData *)
	4749	tableData)->numBuckets)) {
	4750	exitCode = HASH_EXIT_CODE_CYCLE;
	4751	break;
	4752	}
	4753	iteration++;
	4754	}
	4755	switch (exitCode) {
	4756	case HASH_SEARCH_CODE_MATCH:
	4757	case HASH_SEARCH_CODE_MISMATCH:
	4758	return HASH_EXIT_CODE_OVERWRITE;
	4759	case HASH_SEARCH_CODE_EMPTY:
	4760	buckets[index].value = value;
	4761	return HASH_EXIT_CODE_NORMAL;
	4762	default:
	4763	return exitCode;
	4764	}
	4765	}
	4766	int intintLCGQuadraticOpenCompactHash_InnerInsertNoOverwrite(char *tableData,
	4767	unsigned int
	4768	numEntries,
	4769	int *keys,
	4770	int *values) {
	4771	intintLCGQuadraticOpenCompactHash_Bucket *buckets =
	4772	(intintLCGQuadraticOpenCompactHash_Bucket *) &
	4773	tableData[sizeof(intintLCGQuadraticOpenCompactHash_TableData)];
	4774	int resultExitCode = HASH_EXIT_CODE_NORMAL;
	4775	int key;
	4776	int value;
	4777	int index;
	4778	int exitCode;
	4779	uint i;
	4780	for (i = 0; i < numEntries; i++) {
	4781	key = keys[i];
	4782	intintLCGQuadraticOpenCompactHash_TableData *mytableData =
	4783	(intintLCGQuadraticOpenCompactHash_TableData *) tableData;
	4784	intintHash_CompressLCGData compressFuncData =
	4785	mytableData->compressFuncData;
	4786	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
	4787	unsigned long int iteration = 0;
	4788	for (;;) {
	4789	index =
	4790	((1 * iteration * iteration + 0 * iteration +
	4791	c) %
	4792	((intintLCGQuadraticOpenCompactHash_TableData *)
	4793	tableData)->numBuckets);
	4794	if (((buckets[index].key ==
	4795	HASH_BUCKET_STATUS_EMPTY) ? (buckets[index].key =
	4796	key,
	4797	HASH_BUCKET_STATUS_EMPTY)
	4798	: buckets[index].key) ==
	4799	HASH_BUCKET_STATUS_EMPTY) {
	4800	exitCode = HASH_SEARCH_CODE_EMPTY;
	4801	break;
	4802	} else if (key == buckets[index].key) {
	4803	exitCode = HASH_SEARCH_CODE_MATCH;
	4804	break;
	4805	} else
	4806	if ((iteration >
	4807	((intintLCGQuadraticOpenCompactHash_TableData
	4808	*) tableData)->numBuckets)) {
	4809	exitCode = HASH_EXIT_CODE_CYCLE;
	4810	break;
	4811	}
	4812	iteration++;
	4813	}
	4814	switch (exitCode) {
	4815	case HASH_SEARCH_CODE_MATCH:
	4816	case HASH_SEARCH_CODE_MISMATCH:
	4817	resultExitCode = HASH_EXIT_CODE_OVERWRITE;
	4818	break;
	4819	case HASH_SEARCH_CODE_EMPTY:
	4820	buckets[index].value = values[i];
	4821	break;
	4822	default:
	4823	return exitCode;
	4824	}
	4825	}
	4826	return resultExitCode;
	4827	}
	4828	int intintLCGQuadraticOpenCompactHash_QuerySingle(intintHash_Table * table,
	4829	int key, int *valueOutput) {
	4830	return intintLCGQuadraticOpenCompactHash_InnerQuerySingle(table->
	4831	tableData,
	4832	key,
	4833	valueOutput);
	4834	}
	4835	int intintLCGQuadraticOpenCompactHash_Query(intintHash_Table * table,
	4836	size_t numKeys, int *keys,
	4837	int *valuesOutput) {
	4838	return intintLCGQuadraticOpenCompactHash_InnerQuery(table->tableData,
	4839	numKeys, keys,
	4840	valuesOutput);
	4841	}
	4842	int intintLCGQuadraticOpenCompactHash_InsertSingle(intintHash_Table * table,
	4843	int key, int value) {
	4844	return intintLCGQuadraticOpenCompactHash_InnerInsertSingle(table->
	4845	tableData,
	4846	key, value);
	4847	}
	4848	int intintLCGQuadraticOpenCompactHash_Insert(intintHash_Table * table,
	4849	size_t numEntries, int *keys,
	4850	int *values) {
	4851	return intintLCGQuadraticOpenCompactHash_InnerInsert(table->tableData,
	4852	numEntries, keys,
	4853	values);
	4854	}
	4855	int intintLCGQuadraticOpenCompactHash_InsertSingleNoOverwrite(intintHash_Table *
	4856	table, int key,
	4857	int value) {
	4858	return
	4859	intintLCGQuadraticOpenCompactHash_InnerInsertSingleNoOverwrite
	4860	(table->tableData, key, value);
	4861	}
	4862	int intintLCGQuadraticOpenCompactHash_InsertNoOverwrite(intintHash_Table *
	4863	table,
	4864	size_t numEntries,
	4865	int *keys,
	4866	int *values) {
	4867	return intintLCGQuadraticOpenCompactHash_InnerInsertNoOverwrite(table->
	4868	tableData,
	4869	numEntries,
	4870	keys,
	4871	values);
	4872	}
	4873
	4874	typedef struct intintLCGQuadraticOpenCompactCLHash_TableData {
	4875	int hashID;
	4876	unsigned int numBuckets;
	4877	intintHash_CompressLCGData compressFuncData;
	4878	} intintLCGQuadraticOpenCompactCLHash_TableData;
	4879	typedef struct intintLCGQuadraticOpenCompactCLHash_Bucket {
	4880	int key;
	4881	int value;
	4882	} intintLCGQuadraticOpenCompactCLHash_Bucket;
	4883	intintHash_Table
	4884	intintLCGQuadraticOpenCompactCLHash_CreateTable(intintHash_Factory
	4885	factory, int hashIndex,
	4886	size_t keyRange,
	4887	size_t numEntries,
	4888	float loadFactor) {
	4889	intintHash_Table *table =
	4890	(intintHash_Table *) malloc(sizeof(intintHash_Table));
	4891	table->destroyFunc = &intintLCGQuadraticOpenCompactCLHash_DestroyTable;
	4892	table->emptyFunc = &intintLCGQuadraticOpenCompactCLHash_EmptyTable;
	4893	table->queryFunc = &intintLCGQuadraticOpenCompactCLHash_Query;
	4894	table->querySingleFunc =
	4895	&intintLCGQuadraticOpenCompactCLHash_QuerySingle;
	4896	table->insertFunc = &intintLCGQuadraticOpenCompactCLHash_Insert;
	4897	table->insertSingleFunc =
	4898	&intintLCGQuadraticOpenCompactCLHash_InsertSingle;
	4899	table->insertNoOverwriteFunc =
	4900	&intintLCGQuadraticOpenCompactCLHash_InsertNoOverwrite;
	4901	table->insertSingleNoOverwriteFunc =
	4902	&intintLCGQuadraticOpenCompactCLHash_InsertSingleNoOverwrite;
	4903	table->tableData =
	4904	(char *)
	4905	malloc(sizeof(intintLCGQuadraticOpenCompactCLHash_TableData));
	4906	((intintLCGQuadraticOpenCompactCLHash_TableData *) table->tableData)->
	4907	hashID = LCG_QUADRATIC_OPEN_COMPACT_CL_HASH_ID;
	4908	table->context = factory->context;
	4909	table->queue = factory->queue;
	4910	table->program = factory->program;
	4911	table->localWorkSize = factory->localWorkSize;
	4912	table->utilProgram = factory->utilProgram[hashIndex];
	4913	table->emptyKernel = factory->emptyKernel[hashIndex];
	4914	table->emptyKernelLocalWorkSize =
	4915	factory->emptyKernelLocalWorkSize[hashIndex];
	4916	table->querySingleKernel = factory->querySingleKernel[hashIndex];
	4917	table->insertSingleKernel = factory->insertSingleKernel[hashIndex];
	4918	table->insertSingleNoOverwriteKernel =
	4919	factory->insertSingleNoOverwriteKernel[hashIndex];
	4920	clRetainContext(table->context);
	4921	clRetainCommandQueue(table->queue);
	4922	clRetainProgram(table->program);
	4923	clRetainProgram(table->utilProgram);
	4924	clRetainKernel(table->emptyKernel);
	4925	clRetainKernel(table->querySingleKernel);
	4926	clRetainKernel(table->insertSingleKernel);
	4927	clRetainKernel(table->insertSingleNoOverwriteKernel);;
	4928	((intintLCGQuadraticOpenCompactCLHash_TableData *) table->tableData)->
	4929	numBuckets = (unsigned int)((double)numEntries / loadFactor);
	4930	((intintLCGQuadraticOpenCompactCLHash_TableData *) table->tableData)->
	4931	compressFuncData.a = HASH_LCG_A;
	4932	((intintLCGQuadraticOpenCompactCLHash_TableData *) table->tableData)->
	4933	compressFuncData.c = HASH_LCG_C;
	4934	((intintLCGQuadraticOpenCompactCLHash_TableData *) table->tableData)->
	4935	compressFuncData.m = HASH_LCG_M;
	4936	((intintLCGQuadraticOpenCompactCLHash_TableData *) table->tableData)->
	4937	compressFuncData.n =
	4938	((intintLCGQuadraticOpenCompactCLHash_TableData *) table->
	4939	tableData)->numBuckets;
	4940	((intintLCGQuadraticOpenCompactCLHash_TableData *) table->tableData)->
	4941	numBuckets =
	4942	largestProthPrimeUnder(((intintLCGQuadraticOpenCompactCLHash_TableData *) table->tableData)->numBuckets);
	4943	char *tempHashData =
	4944	(char *)malloc(sizeof(intintLCGQuadraticOpenCompactCLHash_TableData)
	4945	+
	4946	((intintLCGQuadraticOpenCompactCLHash_TableData *)
	4947	table->tableData)->numBuckets *
	4948	sizeof(intintLCGQuadraticOpenCompactCLHash_Bucket));
	4949	memcpy(tempHashData, table->tableData,
	4950	sizeof(intintLCGQuadraticOpenCompactCLHash_TableData));
	4951	free(table->tableData);
	4952	table->tableData = tempHashData;
	4953	cl_int err;
	4954	table->tableDataBuffer =
	4955	clCreateBuffer(table->context, CL_MEM_READ_WRITE,
	4956	sizeof(intintLCGQuadraticOpenCompactHash_TableData) +
	4957	((intintLCGQuadraticOpenCompactHash_TableData *)
	4958	table->tableData)->numBuckets *
	4959	sizeof(intintLCGQuadraticOpenCompactHash_Bucket),
	4960	NULL, &err);
	4961	CLHash_Utilities_HandleError(err,
	4962	"intintLCGQuadraticOpenCompactCLHash_InitTable",
	4963	"clCreateBuffer");
	4964	err =
	4965	clEnqueueWriteBuffer(table->queue, table->tableDataBuffer, CL_TRUE,
	4966	0,
	4967	sizeof
	4968	(intintLCGQuadraticOpenCompactHash_TableData),
	4969	table->tableData, 0, NULL, NULL);
	4970	CLHash_Utilities_HandleError(err,
	4971	"intintLCGQuadraticOpenCompactCLHash_InitTable",
	4972	"clEnqueueWriteBuffer");
	4973	return table;
	4974	}
	4975	int intintLCGQuadraticOpenCompactCLHash_CreateFactory(intintHash_Factory *
	4976	factory, int hashIndex) {
	4977	factory->createFunc[hashIndex] =
	4978	&intintLCGQuadraticOpenCompactCLHash_CreateTable;
	4979	factory->destroyFunc[hashIndex] =
	4980	&intintLCGQuadraticOpenCompactCLHash_DestroyFactory;
	4981	cl_int error;
	4982	cl_device_id device;
	4983	error =
	4984	clGetContextInfo(factory->context, CL_CONTEXT_DEVICES,
	4985	sizeof(device), &device, NULL);
	4986	CLHash_Utilities_HandleError(error, "intintHash_CreateFactory",
	4987	"clGetContextInfo");
	4988	factory->querySingleKernel[hashIndex] =
	4989	clCreateKernel(factory->program,
	4990	"intintLCGQuadraticOpenCompactCLHash_RangeQuerySingle",
	4991	&error);
	4992	CLHash_Utilities_HandleError(error,
	4993	"intintLCGQuadraticOpenCompactCLHash_CreateFactory",
	4994	"clCreateKernel");
	4995	factory->insertSingleKernel[hashIndex] =
	4996	clCreateKernel(factory->program,
	4997	"intintLCGQuadraticOpenCompactCLHash_RangeInsertSingle",
	4998	&error);
	4999	CLHash_Utilities_HandleError(error,
	5000	"intintLCGQuadraticOpenCompactCLHash_CreateFactory",
	5001	"clCreateKernel");
	5002	factory->insertSingleNoOverwriteKernel[hashIndex] =
	5003	clCreateKernel(factory->program,
	5004	"intintLCGQuadraticOpenCompactCLHash_RangeInsertSingleNoOverwrite",
	5005	&error);
	5006	CLHash_Utilities_HandleError(error,
	5007	"intintLCGQuadraticOpenCompactCLHash_CreateFactory",
	5008	"clCreateKernel");
	5009	factory->utilProgram[hashIndex] =
	5010	CLHash_Utilities_BuildProgramString(factory->context, device,
	5011	"static inline unsigned int intintHash_CompressIdentity(char data, int hashCode){ return hashCode; } typedef struct intintHash_CompressLCGData{ long unsigned int a; long unsigned int c; unsigned int m; unsigned int n; }intintHash_CompressLCGData; static inline unsigned int intintHash_CompressLCG(intintHash_CompressLCGData compressLCGData, int hashCode){ return ((compressLCGData.a * hashCode + compressLCGData.c) % compressLCGData.m) % compressLCGData.n; } typedef struct intintLCGQuadraticOpenCompactCLHash_TableData{ int hashID; unsigned int numBuckets; intintHash_CompressLCGData compressFuncData; }intintLCGQuadraticOpenCompactCLHash_TableData; typedef struct intintLCGQuadraticOpenCompactCLHash_Bucket{ int key; int value; }intintLCGQuadraticOpenCompactCLHash_Bucket; __kernel void intintLCGQuadraticOpenCompactCLHash_Empty(__global char tableData){ int index = get_global_id(0); if(index >= ((__global intintLCGQuadraticOpenCompactCLHash_TableData)tableData)->numBuckets){ return; } __global intintLCGQuadraticOpenCompactCLHash_Bucket buckets = (__global intintLCGQuadraticOpenCompactCLHash_Bucket)&tableData[sizeof(intintLCGQuadraticOpenCompactCLHash_TableData)]; buckets[index].key = -1;/HASH_BUCKET_STATUS_EMPTY/ }");
	5012	factory->emptyKernel[hashIndex] =
	5013	clCreateKernel(factory->utilProgram[hashIndex],
	5014	"intintLCGQuadraticOpenCompactCLHash_Empty", &error);
	5015	CLHash_Utilities_HandleError(error,
	5016	"intintLCGQuadraticOpenCompactCLHash_CreateFactory",
	5017	"clCreateKernel");
	5018	error =
	5019	clGetKernelWorkGroupInfo(factory->emptyKernel[hashIndex], device,
	5020	CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t),
	5021	&factory->
	5022	emptyKernelLocalWorkSize[hashIndex], NULL);
	5023	CLHash_Utilities_HandleError(error,
	5024	"intintLCGQuadraticOpenCompactCLHash_CreateFactory",
	5025	"clGetKernelWorkGroupInfo");;;
	5026	return HASH_EXIT_CODE_NORMAL;
	5027	}
	5028	int intintLCGQuadraticOpenCompactCLHash_DestroyFactory(intintHash_Factory *
	5029	factory,
	5030	int hashIndex) {;
	5031	clReleaseKernel(factory->emptyKernel[hashIndex]);
	5032	clReleaseProgram(factory->utilProgram[hashIndex]);
	5033	clReleaseKernel(factory->querySingleKernel[hashIndex]);
	5034	clReleaseKernel(factory->insertSingleKernel[hashIndex]);
	5035	clReleaseKernel(factory->insertSingleNoOverwriteKernel[hashIndex]);;
	5036	return HASH_EXIT_CODE_NORMAL;
	5037	}
	5038	int intintLCGQuadraticOpenCompactCLHash_DestroyTable(intintHash_Table * table) {
	5039	int exitCode = 0;
	5040	clReleaseMemObject(table->tableDataBuffer);
	5041	clReleaseContext(table->context);
	5042	clReleaseCommandQueue(table->queue);
	5043	clReleaseProgram(table->utilProgram);
	5044	clReleaseKernel(table->emptyKernel);
	5045	clReleaseProgram(table->program);
	5046	clReleaseKernel(table->querySingleKernel);
	5047	clReleaseKernel(table->insertSingleKernel);
	5048	clReleaseKernel(table->insertSingleNoOverwriteKernel);
	5049	free(table->tableData);
	5050	free(table);
	5051	return exitCode;
	5052	}
	5053	int intintLCGQuadraticOpenCompactCLHash_EmptyTable(intintHash_Table * table) {
	5054	int exitCode = 0;
	5055	cl_int err;
	5056	err =
	5057	clSetKernelArg(table->emptyKernel, 0, sizeof(cl_mem),
	5058	&table->tableDataBuffer);
	5059	CLHash_Utilities_HandleError(err,
	5060	"intintLCGQuadraticOpenCompactCLHash_EmptyTable",
	5061	"clSetKernelArg");
	5062	const size_t groupWorkSize =
	5063	roundUpToNearest(((intintLCGQuadraticOpenCompactHash_TableData *)
	5064	table->tableData)->numBuckets,
	5065	table->emptyKernelLocalWorkSize);
	5066	err =
	5067	clEnqueueNDRangeKernel(table->queue, table->emptyKernel, 1, 0,
	5068	&groupWorkSize,
	5069	(const size_t *)&table->
	5070	emptyKernelLocalWorkSize, 0, NULL, NULL);
	5071	CLHash_Utilities_HandleError(err,
	5072	"intintLCGQuadraticOpenCompactCLHash_EmptyTable",
	5073	"clEnqueueNDRangeKernel");
	5074	exitCode = HASH_EXIT_CODE_NORMAL;;
	5075	return exitCode;
	5076	}
	5077	int intintLCGQuadraticOpenCompactCLHash_QuerySingle(intintHash_Table * table,
	5078	int key, int *valueOutput) {
	5079	return intintLCGQuadraticOpenCompactCLHash_Query(table, 1, &key,
	5080	valueOutput);
	5081	}
	5082	int intintLCGQuadraticOpenCompactCLHash_Query(intintHash_Table * table,
	5083	size_t numKeys, int *keys,
	5084	int *valuesOutput) {
	5085	cl_int err;
	5086	cl_mem keysBuffer =
	5087	clCreateBuffer(table->context,
	5088	CL_MEM_READ_ONLY \| CL_MEM_COPY_HOST_PTR,
	5089	sizeof(int) * numKeys, keys, &err);
	5090	CLHash_Utilities_HandleError(err,
	5091	"intintLCGQuadraticOpenCompactCLHash_Query",
	5092	"clCreateBuffer");
	5093	cl_mem valuesOutputBuffer =
	5094	clCreateBuffer(table->context, CL_MEM_WRITE_ONLY,
	5095	sizeof(int) * numKeys, NULL, &err);
	5096	CLHash_Utilities_HandleError(err,
	5097	"intintLCGQuadraticOpenCompactCLHash_Query",
	5098	"clCreateBuffer");
	5099	intintLCGQuadraticOpenCompactCLHash_BufferQuery(table, numKeys,
	5100	keysBuffer,
	5101	valuesOutputBuffer);
	5102	err =
	5103	clEnqueueReadBuffer(table->queue, valuesOutputBuffer, CL_TRUE, 0,
	5104	sizeof(int) * numKeys, valuesOutput, 0, NULL,
	5105	NULL);
	5106	CLHash_Utilities_HandleError(err,
	5107	"intintLCGQuadraticOpenCompactCLHash_Query",
	5108	"clEnqueueReadBuffer");
	5109	clReleaseMemObject(keysBuffer);
	5110	clReleaseMemObject(valuesOutputBuffer);
	5111	return HASH_EXIT_CODE_NORMAL;
	5112	}
	5113	int intintLCGQuadraticOpenCompactCLHash_BufferQuery(intintHash_Table * table,
	5114	size_t numKeys,
	5115	cl_mem keysBuffer,
	5116	cl_mem valuesOutputBuffer) {
	5117	cl_int err;
	5118	err =
	5119	clSetKernelArg(table->querySingleKernel, 0, sizeof(cl_mem),
	5120	&table->tableDataBuffer);
	5121	CLHash_Utilities_HandleError(err,
	5122	"intintLCGQuadraticOpenCompactCLHash_BufferQuery",
	5123	"clSetKernelArg");
4670	5124	err =
4671	5125	clSetKernelArg(table->querySingleKernel, 1, sizeof(unsigned int),
4672	5126	&numKeys);
…	…	@@ -4698,156 +5152,645 @@ int intintLCGQuadraticOpenCompactCLHash_BufferQuery(intintHash_Table * table,
4698	5152	clFinish(table->queue);
4699	5153	return HASH_EXIT_CODE_NORMAL;
4700	5154	}
4701		int intintLCGQuadraticOpenCompactCLHash_InsertSingle(intintHash_Table * table,
4702		int key, int value) {
4703		return intintLCGQuadraticOpenCompactCLHash_Insert(table, 1, &key,
4704		&value);
	5155	int intintLCGQuadraticOpenCompactCLHash_InsertSingle(intintHash_Table * table,
	5156	int key, int value) {
	5157	return intintLCGQuadraticOpenCompactCLHash_Insert(table, 1, &key,
	5158	&value);
	5159	}
	5160	int intintLCGQuadraticOpenCompactCLHash_Insert(intintHash_Table * table,
	5161	size_t numEntries, int *keys,
	5162	int *values) {
	5163	cl_int err;
	5164	cl_mem keysBuffer =
	5165	clCreateBuffer(table->context,
	5166	CL_MEM_READ_ONLY \| CL_MEM_COPY_HOST_PTR,
	5167	sizeof(int) * numEntries, keys, &err);
	5168	CLHash_Utilities_HandleError(err,
	5169	"intintLCGQuadraticOpenCompactCLHash_Insert",
	5170	"clCreateBuffer");
	5171	cl_mem valuesBuffer =
	5172	clCreateBuffer(table->context,
	5173	CL_MEM_READ_ONLY \| CL_MEM_COPY_HOST_PTR,
	5174	sizeof(int) * numEntries, values, &err);
	5175	CLHash_Utilities_HandleError(err,
	5176	"intintLCGQuadraticOpenCompactCLHash_Insert",
	5177	"clCreateBuffer");
	5178	intintLCGQuadraticOpenCompactCLHash_BufferInsert(table, numEntries,
	5179	keysBuffer,
	5180	valuesBuffer);
	5181	clReleaseMemObject(keysBuffer);
	5182	clReleaseMemObject(valuesBuffer);
	5183	return HASH_EXIT_CODE_NORMAL;
	5184	}
	5185	int intintLCGQuadraticOpenCompactCLHash_BufferInsert(intintHash_Table * table,
	5186	size_t numEntries,
	5187	cl_mem keysBuffer,
	5188	cl_mem valuesBuffer) {
	5189	cl_int err;
	5190	err =
	5191	clSetKernelArg(table->insertSingleKernel, 0, sizeof(cl_mem),
	5192	&table->tableDataBuffer);
	5193	CLHash_Utilities_HandleError(err,
	5194	"intintLCGQuadraticOpenCompactCLHash_BufferInsert",
	5195	"clSetKernelArg");
	5196	err =
	5197	clSetKernelArg(table->insertSingleKernel, 1, sizeof(unsigned int),
	5198	&numEntries);
	5199	CLHash_Utilities_HandleError(err,
	5200	"intintLCGQuadraticOpenCompactCLHash_BufferInsert",
	5201	"clSetKernelArg");
	5202	err =
	5203	clSetKernelArg(table->insertSingleKernel, 2, sizeof(cl_mem),
	5204	&keysBuffer);
	5205	CLHash_Utilities_HandleError(err,
	5206	"intintLCGQuadraticOpenCompactCLHash_BufferInsert",
	5207	"clSetKernelArg");
	5208	err =
	5209	clSetKernelArg(table->insertSingleKernel, 3, sizeof(cl_mem),
	5210	&valuesBuffer);
	5211	CLHash_Utilities_HandleError(err,
	5212	"intintLCGQuadraticOpenCompactCLHash_BufferInsert",
	5213	"clSetKernelArg");
	5214	const size_t groupWorkSize =
	5215	roundUpToNearest(numEntries, table->localWorkSize);
	5216	err =
	5217	clEnqueueNDRangeKernel(table->queue, table->insertSingleKernel, 1,
	5218	0, &groupWorkSize,
	5219	(const size_t *)&table->localWorkSize, 0,
	5220	NULL, NULL);
	5221	CLHash_Utilities_HandleError(err, NULL, "clEnqueueNDRangeKernel");
	5222	return (0);
	5223	}
	5224	int intintLCGQuadraticOpenCompactCLHash_InsertSingleNoOverwrite(intintHash_Table
	5225	* table,
	5226	int key,
	5227	int value) {
	5228	return intintLCGQuadraticOpenCompactCLHash_InsertNoOverwrite(table, 1,
	5229	&key,
	5230	&value);
	5231	}
	5232	int intintLCGQuadraticOpenCompactCLHash_InsertNoOverwrite(intintHash_Table *
	5233	table,
	5234	size_t numEntries,
	5235	int *keys,
	5236	int *values) {
	5237	cl_int err;
	5238	cl_mem keysBuffer =
	5239	clCreateBuffer(table->context,
	5240	CL_MEM_READ_ONLY \| CL_MEM_COPY_HOST_PTR,
	5241	sizeof(int) * numEntries, keys, &err);
	5242	CLHash_Utilities_HandleError(err,
	5243	"intintLCGQuadraticOpenCompactCLHash_InsertNoOverwrite",
	5244	"clCreateBuffer");
	5245	cl_mem valuesBuffer =
	5246	clCreateBuffer(table->context,
	5247	CL_MEM_READ_ONLY \| CL_MEM_COPY_HOST_PTR,
	5248	sizeof(int) * numEntries, values, &err);
	5249	CLHash_Utilities_HandleError(err,
	5250	"intintLCGQuadraticOpenCompactCLHash_InsertNoOverwrite",
	5251	"clCreateBuffer");
	5252	intintLCGQuadraticOpenCompactCLHash_BufferInsertNoOverwrite(table,
	5253	numEntries,
	5254	keysBuffer,
	5255	valuesBuffer);
	5256	clReleaseMemObject(keysBuffer);
	5257	clReleaseMemObject(valuesBuffer);
	5258	return HASH_EXIT_CODE_NORMAL;
	5259	}
	5260	int intintLCGQuadraticOpenCompactCLHash_BufferInsertNoOverwrite(intintHash_Table
	5261	* table,
	5262	size_t
	5263	numEntries,
	5264	cl_mem
	5265	keysBuffer,
	5266	cl_mem
	5267	valuesBuffer) {
	5268	cl_int err;
	5269	err =
	5270	clSetKernelArg(table->insertSingleNoOverwriteKernel, 0,
	5271	sizeof(cl_mem), &table->tableDataBuffer);
	5272	CLHash_Utilities_HandleError(err,
	5273	"intintLCGQuadraticOpenCompactCLHash_BufferInsertNoOverwrite",
	5274	"clSetKernelArg");
	5275	err =
	5276	clSetKernelArg(table->insertSingleNoOverwriteKernel, 1,
	5277	sizeof(unsigned int), &numEntries);
	5278	CLHash_Utilities_HandleError(err,
	5279	"intintLCGQuadraticOpenCompactCLHash_BufferInsertNoOverwrite",
	5280	"ClSetKernelArg");
	5281	err =
	5282	clSetKernelArg(table->insertSingleNoOverwriteKernel, 2,
	5283	sizeof(cl_mem), &keysBuffer);
	5284	CLHash_Utilities_HandleError(err,
	5285	"intintLCGQuadraticOpenCompactCLHash_BufferInsertNoOverwrite",
	5286	"clSetKernelArg");
	5287	err =
	5288	clSetKernelArg(table->insertSingleNoOverwriteKernel, 3,
	5289	sizeof(cl_mem), &valuesBuffer);
	5290	CLHash_Utilities_HandleError(err,
	5291	"intintLCGQuadraticOpenCompactCLHash_BufferInsertNoOverwrite",
	5292	"clSetKernelArg");
	5293	const size_t groupWorkSize =
	5294	roundUpToNearest(numEntries, table->localWorkSize);
	5295	err =
	5296	clEnqueueNDRangeKernel(table->queue,
	5297	table->insertSingleNoOverwriteKernel, 1, 0,
	5298	&groupWorkSize,
	5299	(const size_t *)&table->localWorkSize, 0,
	5300	NULL, NULL);
	5301	CLHash_Utilities_HandleError(err,
	5302	"intintLCGQuadraticOpenCompactCLHash_BufferInsertNoOverwrite",
	5303	"clEnqueueNDRangeKernel");
	5304	return (0);
	5305	}
	5306
	5307	typedef struct intintLCGQuadraticOpenCompactOpenMPHash_TableData {
	5308	int hashID;
	5309	unsigned int numBuckets;
	5310	intintHash_CompressLCGData compressFuncData;
	5311	} intintLCGQuadraticOpenCompactOpenMPHash_TableData;
	5312	typedef struct intintLCGQuadraticOpenCompactOpenMPHash_Bucket {
	5313	int key;
	5314	int value;
	5315	} intintLCGQuadraticOpenCompactOpenMPHash_Bucket;
	5316	intintHash_Table
	5317	intintLCGQuadraticOpenCompactOpenMPHash_CreateTable(intintHash_Factory
	5318	factory, int hashIndex,
	5319	size_t keyRange,
	5320	size_t numEntries,
	5321	float loadFactor) {
	5322	intintHash_Table *table =
	5323	(intintHash_Table *) malloc(sizeof(intintHash_Table));
	5324	table->destroyFunc =
	5325	&intintLCGQuadraticOpenCompactOpenMPHash_DestroyTable;
	5326	table->emptyFunc = &intintLCGQuadraticOpenCompactOpenMPHash_EmptyTable;
	5327	table->queryFunc = &intintLCGQuadraticOpenCompactOpenMPHash_Query;
	5328	table->querySingleFunc =
	5329	&intintLCGQuadraticOpenCompactOpenMPHash_QuerySingle;
	5330	table->insertFunc = &intintLCGQuadraticOpenCompactOpenMPHash_Insert;
	5331	table->insertSingleFunc =
	5332	&intintLCGQuadraticOpenCompactOpenMPHash_InsertSingle;
	5333	table->insertNoOverwriteFunc =
	5334	&intintLCGQuadraticOpenCompactOpenMPHash_InsertNoOverwrite;
	5335	table->insertSingleNoOverwriteFunc =
	5336	&intintLCGQuadraticOpenCompactOpenMPHash_InsertSingleNoOverwrite;
	5337	table->tableData =
	5338	(char *)
	5339	malloc(sizeof(intintLCGQuadraticOpenCompactOpenMPHash_TableData));
	5340	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *) table->
	5341	tableData)->hashID = LCG_QUADRATIC_OPEN_COMPACT_OPENMP_HASH_ID;
	5342	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *) table->
	5343	tableData)->numBuckets =
	5344	(unsigned int)((double)numEntries / loadFactor);
	5345	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *) table->
	5346	tableData)->compressFuncData.a = HASH_LCG_A;
	5347	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *) table->
	5348	tableData)->compressFuncData.c = HASH_LCG_C;
	5349	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *) table->
	5350	tableData)->compressFuncData.m = HASH_LCG_M;
	5351	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *) table->
	5352	tableData)->compressFuncData.n =
	5353	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *) table->tableData)->numBuckets;
	5354	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *) table->
	5355	tableData)->numBuckets =
	5356	largestProthPrimeUnder(((intintLCGQuadraticOpenCompactOpenMPHash_TableData *) table->tableData)->numBuckets);
	5357	char *tempHashData =
	5358	(char *)
	5359	malloc(sizeof(intintLCGQuadraticOpenCompactOpenMPHash_TableData) +
	5360	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *)
	5361	table->tableData)->numBuckets *
	5362	sizeof(intintLCGQuadraticOpenCompactOpenMPHash_Bucket));
	5363	memcpy(tempHashData, table->tableData,
	5364	sizeof(intintLCGQuadraticOpenCompactOpenMPHash_TableData));
	5365	free(table->tableData);
	5366	table->tableData = tempHashData;
	5367	return table;
	5368	}
	5369	int intintLCGQuadraticOpenCompactOpenMPHash_CreateFactory(intintHash_Factory *
	5370	factory,
	5371	int hashIndex) {
	5372	factory->createFunc[hashIndex] =
	5373	&intintLCGQuadraticOpenCompactOpenMPHash_CreateTable;
	5374	factory->destroyFunc[hashIndex] =
	5375	&intintLCGQuadraticOpenCompactOpenMPHash_DestroyFactory;;
	5376	return HASH_EXIT_CODE_NORMAL;
	5377	}
	5378	int intintLCGQuadraticOpenCompactOpenMPHash_DestroyFactory(intintHash_Factory *
	5379	factory,
	5380	int hashIndex) {;
	5381	return HASH_EXIT_CODE_NORMAL;
	5382	}
	5383	int intintLCGQuadraticOpenCompactOpenMPHash_DestroyTable(intintHash_Table *
	5384	table) {
	5385	int exitCode = 0;
	5386	free(table->tableData);
	5387	free(table);
	5388	return exitCode;
	5389	}
	5390	int intintLCGQuadraticOpenCompactOpenMPHash_EmptyTable(intintHash_Table * table) {
	5391	int exitCode = 0;
	5392	intintLCGQuadraticOpenCompactOpenMPHash_Bucket *buckets =
	5393	(intintLCGQuadraticOpenCompactOpenMPHash_Bucket *) & table->
	5394	tableData[sizeof
	5395	(intintLCGQuadraticOpenCompactOpenMPHash_TableData)];
	5396	int index;
	5397	for (index = 0;
	5398	index <
	5399	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *) table->
	5400	tableData)->numBuckets; index++) {
	5401	buckets[index].key = HASH_BUCKET_STATUS_EMPTY;
	5402	}
	5403	exitCode = HASH_EXIT_CODE_NORMAL;
	5404	return exitCode;
	5405	}
	5406	int intintLCGQuadraticOpenCompactOpenMPHash_InnerQuerySingle(char *tableData,
	5407	int key,
	5408	int *valueOutput) {
	5409	intintLCGQuadraticOpenCompactOpenMPHash_Bucket *buckets =
	5410	(intintLCGQuadraticOpenCompactOpenMPHash_Bucket *) &
	5411	tableData[sizeof
	5412	(intintLCGQuadraticOpenCompactOpenMPHash_TableData)];
	5413	int index;
	5414	int exitCode;
	5415	intintLCGQuadraticOpenCompactOpenMPHash_TableData *mytableData =
	5416	(intintLCGQuadraticOpenCompactOpenMPHash_TableData *) tableData;
	5417	intintHash_CompressLCGData compressFuncData =
	5418	mytableData->compressFuncData;
	5419	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
	5420	unsigned long int iteration = 0;
	5421	for (;;) {
	5422	index =
	5423	((1 * iteration * iteration + 0 * iteration +
	5424	c) %
	5425	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *)
	5426	tableData)->numBuckets);
	5427	if ((buckets[index].key) == HASH_BUCKET_STATUS_EMPTY) {
	5428	exitCode = HASH_SEARCH_CODE_EMPTY;
	5429	break;
	5430	} else if (key == buckets[index].key) {
	5431	exitCode = HASH_SEARCH_CODE_MATCH;
	5432	break;
	5433	} else
	5434	if ((iteration >
	5435	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *)
	5436	tableData)->numBuckets)) {
	5437	exitCode = HASH_EXIT_CODE_CYCLE;
	5438	break;
	5439	}
	5440	iteration++;
	5441	}
	5442	switch (exitCode) {
	5443	case HASH_SEARCH_CODE_MATCH:
	5444	*valueOutput = buckets[index].value;
	5445	return HASH_EXIT_CODE_NORMAL;
	5446	case HASH_SEARCH_CODE_MISMATCH:
	5447	case HASH_SEARCH_CODE_EMPTY:
	5448	return HASH_EXIT_CODE_KEY_DNE;
	5449	default:
	5450	return exitCode;
	5451	}
	5452	}
	5453	int intintLCGQuadraticOpenCompactOpenMPHash_InnerQuery(char *tableData,
	5454	unsigned int numKeys,
	5455	int *keys,
	5456	int *valuesOutput) {
	5457	intintLCGQuadraticOpenCompactOpenMPHash_Bucket *buckets =
	5458	(intintLCGQuadraticOpenCompactOpenMPHash_Bucket *) &
	5459	tableData[sizeof
	5460	(intintLCGQuadraticOpenCompactOpenMPHash_TableData)];
	5461	int key;
	5462	int *valueOutput;
	5463	int index;
	5464	int exitCode;
	5465	uint i;
	5466	int resultExitCode = HASH_EXIT_CODE_NORMAL;
	5467	for (i = 0; i < numKeys; i++) {
	5468	key = keys[i];
	5469	valueOutput = &valuesOutput[i];
	5470	intintLCGQuadraticOpenCompactOpenMPHash_TableData *mytableData =
	5471	(intintLCGQuadraticOpenCompactOpenMPHash_TableData *)
	5472	tableData;
	5473	intintHash_CompressLCGData compressFuncData =
	5474	mytableData->compressFuncData;
	5475	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
	5476	unsigned long int iteration = 0;
	5477	for (;;) {
	5478	index =
	5479	((1 * iteration * iteration + 0 * iteration +
	5480	c) %
	5481	((intintLCGQuadraticOpenCompactOpenMPHash_TableData
	5482	*) tableData)->numBuckets);
	5483	if ((buckets[index].key) == HASH_BUCKET_STATUS_EMPTY) {
	5484	exitCode = HASH_SEARCH_CODE_EMPTY;
	5485	break;
	5486	} else if (key == buckets[index].key) {
	5487	exitCode = HASH_SEARCH_CODE_MATCH;
	5488	break;
	5489	} else
	5490	if ((iteration >
	5491	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *) tableData)->numBuckets)) {
	5492	exitCode = HASH_EXIT_CODE_CYCLE;
	5493	break;
	5494	}
	5495	iteration++;
	5496	}
	5497	switch (exitCode) {
	5498	case HASH_SEARCH_CODE_MATCH:
	5499	*valueOutput = buckets[index].value;
	5500	break;
	5501	case HASH_SEARCH_CODE_MISMATCH:
	5502	case HASH_SEARCH_CODE_EMPTY:
	5503	resultExitCode = HASH_EXIT_CODE_KEY_DNE;
	5504	break;
	5505	default:
	5506	return exitCode;
	5507	}
	5508	}
	5509	return resultExitCode;
	5510	}
	5511	int intintLCGQuadraticOpenCompactOpenMPHash_InnerInsertSingle(char *tableData,
	5512	int key,
	5513	int value) {
	5514	intintLCGQuadraticOpenCompactOpenMPHash_Bucket *buckets =
	5515	(intintLCGQuadraticOpenCompactOpenMPHash_Bucket *) &
	5516	tableData[sizeof
	5517	(intintLCGQuadraticOpenCompactOpenMPHash_TableData)];
	5518	int index;
	5519	int exitCode;
	5520	intintLCGQuadraticOpenCompactOpenMPHash_TableData *mytableData =
	5521	(intintLCGQuadraticOpenCompactOpenMPHash_TableData *) tableData;
	5522	intintHash_CompressLCGData compressFuncData =
	5523	mytableData->compressFuncData;
	5524	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
	5525	unsigned long int iteration = 0;
	5526	for (;;) {
	5527	index =
	5528	((1 * iteration * iteration + 0 * iteration +
	5529	c) %
	5530	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *)
	5531	tableData)->numBuckets);
	5532	if (((buckets[index].key ==
	5533	HASH_BUCKET_STATUS_EMPTY) ? (buckets[index].key =
	5534	key,
	5535	HASH_BUCKET_STATUS_EMPTY) :
	5536	buckets[index].key) == HASH_BUCKET_STATUS_EMPTY) {
	5537	exitCode = HASH_SEARCH_CODE_EMPTY;
	5538	break;
	5539	} else if (key == buckets[index].key) {
	5540	exitCode = HASH_SEARCH_CODE_MATCH;
	5541	break;
	5542	} else
	5543	if ((iteration >
	5544	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *)
	5545	tableData)->numBuckets)) {
	5546	exitCode = HASH_EXIT_CODE_CYCLE;
	5547	break;
	5548	}
	5549	iteration++;
	5550	}
	5551	switch (exitCode) {
	5552	case HASH_SEARCH_CODE_MATCH:
	5553	case HASH_SEARCH_CODE_MISMATCH:
	5554	buckets[index].value = value;
	5555	return HASH_EXIT_CODE_OVERWRITE;
	5556	case HASH_SEARCH_CODE_EMPTY:
	5557	buckets[index].value = value;
	5558	return HASH_EXIT_CODE_NORMAL;
	5559	default:
	5560	return exitCode;
	5561	}
	5562	}
	5563	int intintLCGQuadraticOpenCompactOpenMPHash_InnerInsert(char *tableData,
	5564	unsigned int numEntries,
	5565	int *keys,
	5566	int *values) {
	5567	intintLCGQuadraticOpenCompactOpenMPHash_Bucket *buckets =
	5568	(intintLCGQuadraticOpenCompactOpenMPHash_Bucket *) &
	5569	tableData[sizeof
	5570	(intintLCGQuadraticOpenCompactOpenMPHash_TableData)];
	5571	int resultExitCode = HASH_EXIT_CODE_NORMAL;
	5572	int key;
	5573	int value;
	5574	int index;
	5575	int exitCode;
	5576	uint i;
	5577	for (i = 0; i < numEntries; i++) {
	5578	key = keys[i];
	5579	intintLCGQuadraticOpenCompactOpenMPHash_TableData *mytableData =
	5580	(intintLCGQuadraticOpenCompactOpenMPHash_TableData *)
	5581	tableData;
	5582	intintHash_CompressLCGData compressFuncData =
	5583	mytableData->compressFuncData;
	5584	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
	5585	unsigned long int iteration = 0;
	5586	for (;;) {
	5587	index =
	5588	((1 * iteration * iteration + 0 * iteration +
	5589	c) %
	5590	((intintLCGQuadraticOpenCompactOpenMPHash_TableData
	5591	*) tableData)->numBuckets);
	5592	if (((buckets[index].key ==
	5593	HASH_BUCKET_STATUS_EMPTY) ? (buckets[index].key =
	5594	key,
	5595	HASH_BUCKET_STATUS_EMPTY)
	5596	: buckets[index].key) ==
	5597	HASH_BUCKET_STATUS_EMPTY) {
	5598	exitCode = HASH_SEARCH_CODE_EMPTY;
	5599	break;
	5600	} else if (key == buckets[index].key) {
	5601	exitCode = HASH_SEARCH_CODE_MATCH;
	5602	break;
	5603	} else
	5604	if ((iteration >
	5605	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *) tableData)->numBuckets)) {
	5606	exitCode = HASH_EXIT_CODE_CYCLE;
	5607	break;
	5608	}
	5609	iteration++;
	5610	}
	5611	switch (exitCode) {
	5612	case HASH_SEARCH_CODE_MATCH:
	5613	case HASH_SEARCH_CODE_MISMATCH:
	5614	resultExitCode = HASH_EXIT_CODE_OVERWRITE;
	5615	case HASH_SEARCH_CODE_EMPTY:
	5616	buckets[index].value = values[i];
	5617	break;
	5618	default:
	5619	return exitCode;
	5620	}
	5621	}
	5622	return resultExitCode;
	5623	}
	5624	int intintLCGQuadraticOpenCompactOpenMPHash_InnerInsertSingleNoOverwrite(char
	5625	*tableData,
	5626	int
	5627	key,
	5628	int
	5629	value)
	5630	{
	5631	intintLCGQuadraticOpenCompactOpenMPHash_Bucket *buckets =
	5632	(intintLCGQuadraticOpenCompactOpenMPHash_Bucket *) &
	5633	tableData[sizeof
	5634	(intintLCGQuadraticOpenCompactOpenMPHash_TableData)];
	5635	int index;
	5636	int exitCode;
	5637	intintLCGQuadraticOpenCompactOpenMPHash_TableData *mytableData =
	5638	(intintLCGQuadraticOpenCompactOpenMPHash_TableData *) tableData;
	5639	intintHash_CompressLCGData compressFuncData =
	5640	mytableData->compressFuncData;
	5641	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
	5642	unsigned long int iteration = 0;
	5643	for (;;) {
	5644	index =
	5645	((1 * iteration * iteration + 0 * iteration +
	5646	c) %
	5647	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *)
	5648	tableData)->numBuckets);
	5649	if (((buckets[index].key ==
	5650	HASH_BUCKET_STATUS_EMPTY) ? (buckets[index].key =
	5651	key,
	5652	HASH_BUCKET_STATUS_EMPTY) :
	5653	buckets[index].key) == HASH_BUCKET_STATUS_EMPTY) {
	5654	exitCode = HASH_SEARCH_CODE_EMPTY;
	5655	break;
	5656	} else if (key == buckets[index].key) {
	5657	exitCode = HASH_SEARCH_CODE_MATCH;
	5658	break;
	5659	} else
	5660	if ((iteration >
	5661	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *)
	5662	tableData)->numBuckets)) {
	5663	exitCode = HASH_EXIT_CODE_CYCLE;
	5664	break;
	5665	}
	5666	iteration++;
	5667	}
	5668	switch (exitCode) {
	5669	case HASH_SEARCH_CODE_MATCH:
	5670	case HASH_SEARCH_CODE_MISMATCH:
	5671	return HASH_EXIT_CODE_OVERWRITE;
	5672	case HASH_SEARCH_CODE_EMPTY:
	5673	buckets[index].value = value;
	5674	return HASH_EXIT_CODE_NORMAL;
	5675	default:
	5676	return exitCode;
	5677	}
	5678	}
	5679	int intintLCGQuadraticOpenCompactOpenMPHash_InnerInsertNoOverwrite(char
	5680	*tableData,
	5681	unsigned int
	5682	numEntries,
	5683	int *keys,
	5684	int *values)
	5685	{
	5686	intintLCGQuadraticOpenCompactOpenMPHash_Bucket *buckets =
	5687	(intintLCGQuadraticOpenCompactOpenMPHash_Bucket *) &
	5688	tableData[sizeof
	5689	(intintLCGQuadraticOpenCompactOpenMPHash_TableData)];
	5690	int resultExitCode = HASH_EXIT_CODE_NORMAL;
	5691	int key;
	5692	int value;
	5693	int index;
	5694	int exitCode;
	5695	uint i;
	5696	for (i = 0; i < numEntries; i++) {
	5697	key = keys[i];
	5698	intintLCGQuadraticOpenCompactOpenMPHash_TableData *mytableData =
	5699	(intintLCGQuadraticOpenCompactOpenMPHash_TableData *)
	5700	tableData;
	5701	intintHash_CompressLCGData compressFuncData =
	5702	mytableData->compressFuncData;
	5703	unsigned int c = intintHash_CompressLCG(compressFuncData, key);
	5704	unsigned long int iteration = 0;
	5705	for (;;) {
	5706	index =
	5707	((1 * iteration * iteration + 0 * iteration +
	5708	c) %
	5709	((intintLCGQuadraticOpenCompactOpenMPHash_TableData
	5710	*) tableData)->numBuckets);
	5711	if (((buckets[index].key ==
	5712	HASH_BUCKET_STATUS_EMPTY) ? (buckets[index].key =
	5713	key,
	5714	HASH_BUCKET_STATUS_EMPTY)
	5715	: buckets[index].key) ==
	5716	HASH_BUCKET_STATUS_EMPTY) {
	5717	exitCode = HASH_SEARCH_CODE_EMPTY;
	5718	break;
	5719	} else if (key == buckets[index].key) {
	5720	exitCode = HASH_SEARCH_CODE_MATCH;
	5721	break;
	5722	} else
	5723	if ((iteration >
	5724	((intintLCGQuadraticOpenCompactOpenMPHash_TableData *) tableData)->numBuckets)) {
	5725	exitCode = HASH_EXIT_CODE_CYCLE;
	5726	break;
	5727	}
	5728	iteration++;
	5729	}
	5730	switch (exitCode) {
	5731	case HASH_SEARCH_CODE_MATCH:
	5732	case HASH_SEARCH_CODE_MISMATCH:
	5733	resultExitCode = HASH_EXIT_CODE_OVERWRITE;
	5734	break;
	5735	case HASH_SEARCH_CODE_EMPTY:
	5736	buckets[index].value = values[i];
	5737	break;
	5738	default:
	5739	return exitCode;
	5740	}
	5741	}
	5742	return resultExitCode;
4705	5743	}
4706		int intintLCGQuadraticOpenCompactCLHash_Insert(intintHash_Table * table,
4707		size_t numEntries, int *keys,
4708		int *values) {
4709		cl_int err;
4710		cl_mem keysBuffer =
4711		clCreateBuffer(table->context,
4712		CL_MEM_READ_ONLY \| CL_MEM_COPY_HOST_PTR,
4713		sizeof(int) * numEntries, keys, &err);
4714		CLHash_Utilities_HandleError(err,
4715		"intintLCGQuadraticOpenCompactCLHash_Insert",
4716		"clCreateBuffer");
4717		cl_mem valuesBuffer =
4718		clCreateBuffer(table->context,
4719		CL_MEM_READ_ONLY \| CL_MEM_COPY_HOST_PTR,
4720		sizeof(int) * numEntries, values, &err);
4721		CLHash_Utilities_HandleError(err,
4722		"intintLCGQuadraticOpenCompactCLHash_Insert",
4723		"clCreateBuffer");
4724		intintLCGQuadraticOpenCompactCLHash_BufferInsert(table, numEntries,
4725		keysBuffer,
4726		valuesBuffer);
4727		clReleaseMemObject(keysBuffer);
4728		clReleaseMemObject(valuesBuffer);
4729		return HASH_EXIT_CODE_NORMAL;
	5744	int intintLCGQuadraticOpenCompactOpenMPHash_QuerySingle(intintHash_Table *
	5745	table, int key,
	5746	int *valueOutput) {
	5747	return intintLCGQuadraticOpenCompactOpenMPHash_InnerQuerySingle(table->
	5748	tableData,
	5749	key,
	5750	valueOutput);
4730	5751	}
4731		int intintLCGQuadraticOpenCompactCLHash_BufferInsert(intintHash_Table * table,
4732		size_t numEntries,
4733		cl_mem keysBuffer,
4734		cl_mem valuesBuffer) {
4735		cl_int err;
4736		err =
4737		clSetKernelArg(table->insertSingleKernel, 0, sizeof(cl_mem),
4738		&table->tableDataBuffer);
4739		CLHash_Utilities_HandleError(err,
4740		"intintLCGQuadraticOpenCompactCLHash_BufferInsert",
4741		"clSetKernelArg");
4742		err =
4743		clSetKernelArg(table->insertSingleKernel, 1, sizeof(unsigned int),
4744		&numEntries);
4745		CLHash_Utilities_HandleError(err,
4746		"intintLCGQuadraticOpenCompactCLHash_BufferInsert",
4747		"clSetKernelArg");
4748		err =
4749		clSetKernelArg(table->insertSingleKernel, 2, sizeof(cl_mem),
4750		&keysBuffer);
4751		CLHash_Utilities_HandleError(err,
4752		"intintLCGQuadraticOpenCompactCLHash_BufferInsert",
4753		"clSetKernelArg");
4754		err =
4755		clSetKernelArg(table->insertSingleKernel, 3, sizeof(cl_mem),
4756		&valuesBuffer);
4757		CLHash_Utilities_HandleError(err,
4758		"intintLCGQuadraticOpenCompactCLHash_BufferInsert",
4759		"clSetKernelArg");
4760		const size_t groupWorkSize =
4761		roundUpToNearest(numEntries, table->localWorkSize);
4762		err =
4763		clEnqueueNDRangeKernel(table->queue, table->insertSingleKernel, 1,
4764		0, &groupWorkSize,
4765		(const size_t *)&table->localWorkSize, 0,
4766		NULL, NULL);
4767		CLHash_Utilities_HandleError(err, NULL, "clEnqueueNDRangeKernel");
4768		return (0);
	5752	int intintLCGQuadraticOpenCompactOpenMPHash_Query(intintHash_Table * table,
	5753	size_t numKeys, int *keys,
	5754	int *valuesOutput) {
	5755	return intintLCGQuadraticOpenCompactOpenMPHash_InnerQuery(table->
	5756	tableData,
	5757	numKeys, keys,
	5758	valuesOutput);
4769	5759	}
4770		int intintLCGQuadraticOpenCompactCLHash_InsertSingleNoOverwrite(intintHash_Table
	5760	int intintLCGQuadraticOpenCompactOpenMPHash_InsertSingle(intintHash_Table *
	5761	table, int key,
	5762	int value) {
	5763	return intintLCGQuadraticOpenCompactOpenMPHash_InnerInsertSingle(table->
	5764	tableData,
	5765	key,
	5766	value);
	5767	}
	5768	int intintLCGQuadraticOpenCompactOpenMPHash_Insert(intintHash_Table * table,
	5769	size_t numEntries, int *keys,
	5770	int *values) {
	5771	return intintLCGQuadraticOpenCompactOpenMPHash_InnerInsert(table->
	5772	tableData,
	5773	numEntries,
	5774	keys,
	5775	values);
	5776	}
	5777	int
	5778	intintLCGQuadraticOpenCompactOpenMPHash_InsertSingleNoOverwrite(intintHash_Table
4771	5779	* table,
4772	5780	int key,
4773	5781	int value) {
4774		return intintLCGQuadraticOpenCompactCLHash_InsertNoOverwrite(table, 1,
4775		&key,
4776		&value);
4777		}
4778		int intintLCGQuadraticOpenCompactCLHash_InsertNoOverwrite(intintHash_Table *
4779		table,
4780		size_t numEntries,
4781		int *keys,
4782		int *values) {
4783		cl_int err;
4784		cl_mem keysBuffer =
4785		clCreateBuffer(table->context,
4786		CL_MEM_READ_ONLY \| CL_MEM_COPY_HOST_PTR,
4787		sizeof(int) * numEntries, keys, &err);
4788		CLHash_Utilities_HandleError(err,
4789		"intintLCGQuadraticOpenCompactCLHash_InsertNoOverwrite",
4790		"clCreateBuffer");
4791		cl_mem valuesBuffer =
4792		clCreateBuffer(table->context,
4793		CL_MEM_READ_ONLY \| CL_MEM_COPY_HOST_PTR,
4794		sizeof(int) * numEntries, values, &err);
4795		CLHash_Utilities_HandleError(err,
4796		"intintLCGQuadraticOpenCompactCLHash_InsertNoOverwrite",
4797		"clCreateBuffer");
4798		intintLCGQuadraticOpenCompactCLHash_BufferInsertNoOverwrite(table,
4799		numEntries,
4800		keysBuffer,
4801		valuesBuffer);
4802		clReleaseMemObject(keysBuffer);
4803		clReleaseMemObject(valuesBuffer);
4804		return HASH_EXIT_CODE_NORMAL;
	5782	return
	5783	intintLCGQuadraticOpenCompactOpenMPHash_InnerInsertSingleNoOverwrite
	5784	(table->tableData, key, value);
4805	5785	}
4806		int intintLCGQuadraticOpenCompactCLHash_BufferInsertNoOverwrite(intintHash_Table
4807		* table,
4808		size_t
4809		numEntries,
4810		cl_mem
4811		keysBuffer,
4812		cl_mem
4813		valuesBuffer) {
4814		cl_int err;
4815		err =
4816		clSetKernelArg(table->insertSingleNoOverwriteKernel, 0,
4817		sizeof(cl_mem), &table->tableDataBuffer);
4818		CLHash_Utilities_HandleError(err,
4819		"intintLCGQuadraticOpenCompactCLHash_BufferInsertNoOverwrite",
4820		"clSetKernelArg");
4821		err =
4822		clSetKernelArg(table->insertSingleNoOverwriteKernel, 1,
4823		sizeof(unsigned int), &numEntries);
4824		CLHash_Utilities_HandleError(err,
4825		"intintLCGQuadraticOpenCompactCLHash_BufferInsertNoOverwrite",
4826		"ClSetKernelArg");
4827		err =
4828		clSetKernelArg(table->insertSingleNoOverwriteKernel, 2,
4829		sizeof(cl_mem), &keysBuffer);
4830		CLHash_Utilities_HandleError(err,
4831		"intintLCGQuadraticOpenCompactCLHash_BufferInsertNoOverwrite",
4832		"clSetKernelArg");
4833		err =
4834		clSetKernelArg(table->insertSingleNoOverwriteKernel, 3,
4835		sizeof(cl_mem), &valuesBuffer);
4836		CLHash_Utilities_HandleError(err,
4837		"intintLCGQuadraticOpenCompactCLHash_BufferInsertNoOverwrite",
4838		"clSetKernelArg");
4839		const size_t groupWorkSize =
4840		roundUpToNearest(numEntries, table->localWorkSize);
4841		err =
4842		clEnqueueNDRangeKernel(table->queue,
4843		table->insertSingleNoOverwriteKernel, 1, 0,
4844		&groupWorkSize,
4845		(const size_t *)&table->localWorkSize, 0,
4846		NULL, NULL);
4847		CLHash_Utilities_HandleError(err,
4848		"intintLCGQuadraticOpenCompactCLHash_BufferInsertNoOverwrite",
4849		"clEnqueueNDRangeKernel");
4850		return (0);
	5786	int intintLCGQuadraticOpenCompactOpenMPHash_InsertNoOverwrite(intintHash_Table *
	5787	table,
	5788	size_t numEntries,
	5789	int *keys,
	5790	int *values) {
	5791	return
	5792	intintLCGQuadraticOpenCompactOpenMPHash_InnerInsertNoOverwrite
	5793	(table->tableData, numEntries, keys, values);
4851	5794	}
4852	5795	const char *HashFactory_source =
4853	5796	"\n"

HashFactory/HashFactory.cm +4 -0

…	…	@@ -302,9 +302,13 @@
302	302	\
303	303	LCG_LINEAR_OPEN_COMPACT_CL_HASH_DEFINE(NAME, KEY_TYPE, VALUE_TYPE)\
304	304	\
	305	LCG_LINEAR_OPEN_COMPACT_OPENMP_HASH_DEFINE(NAME, KEY_TYPE, VALUE_TYPE, 1)\
	306	\
305	307	LCG_QUADRATIC_OPEN_COMPACT_HASH_DEFINE(NAME, KEY_TYPE, VALUE_TYPE, 1, 0)\
306	308	\
307	309	LCG_QUADRATIC_OPEN_COMPACT_CL_HASH_DEFINE(NAME, KEY_TYPE, VALUE_TYPE)\
	310	\
	311	LCG_QUADRATIC_OPEN_COMPACT_OPENMP_HASH_DEFINE(NAME, KEY_TYPE, VALUE_TYPE, 1, 0)\
308	312	\
309	313
310	314	#define KERN_SWITCH_LOOKUP_EXP(NAME, FUNCTION_NAME, ...)\

HashFactory/compactHash.cm +28 -0

…	…	@@ -55,6 +55,11 @@
55	55	CL_BASE_HASH_DEFINE(NAME, MODIFIER##Compact, HASH_ID, KEY_TYPE, VALUE_TYPE, KEY_VALUE_BASE_HASH_TYPES_DEFINE, COMPRESS_FUNC, COMPRESS_FUNC_DATA,,, INIT_EXP, KEY_VALUE_CL_HASH_INNER_EMPTY_EXP)\
56	56	\
57	57
	58	#define COMPACT_OPENMP_HASH_DEFINE(NAME, MODIFIER, HASH_ID, KEY_TYPE, VALUE_TYPE, COMPRESS_FUNC, COMPRESS_FUNC_DATA, INIT_EXP, SEARCH_EXP, ...)\
	59	\
	60	BASE_HASH_DEFINE(NAME, MODIFIER##CompactOpenMP, HASH_ID, KEY_TYPE, VALUE_TYPE, KEY_VALUE_BASE_HASH_TYPES_DEFINE, COMPRESS_FUNC, COMPRESS_FUNC_DATA,,, INIT_EXP, SEARCH_EXP, KEY_VALUE_HASH_INNER_EMPTY_EXP, ##__VA_ARGS__)\
	61	\
	62
58	63	#define OPEN_COMPACT_HASH_INIT_EXP(NAME, MODIFIER, INIT_EXP) \
59	64	((NAME##MODIFIER##Hash_TableData*)table->tableData)->numBuckets = (unsigned int)((double)numEntries / loadFactor);\
60	65	INIT_EXP\
…	…	@@ -100,6 +105,10 @@ unsigned int c = COMPRESS_FUNC(((CL_GLOBAL NAME##MODIFIER##Hash_TableData*)table
100	105	COMPACT_CL_HASH_DEFINE(NAME, MODIFIER##Open, HASH_ID, KEY_TYPE, VALUE_TYPE, COMPRESS_FUNC, COMPRESS_FUNC_DATA, OPEN_COMPACT_HASH_INIT_EXP(NAME, MODIFIER##OpenCompactCL, INIT_EXP))\
101	106	\
102	107
	108	#define OPEN_COMPACT_OPENMP_HASH_DEFINE(NAME, MODIFIER, HASH_ID, KEY_TYPE, VALUE_TYPE, COMPRESS_FUNC, COMPRESS_FUNC_DATA, INIT_EXP, ...)\
	109	\
	110	COMPACT_OPENMP_HASH_DEFINE(NAME, MODIFIER##Open, HASH_ID, KEY_TYPE, VALUE_TYPE, COMPRESS_FUNC, COMPRESS_FUNC_DATA, OPEN_COMPACT_HASH_INIT_EXP(NAME, MODIFIER##OpenCompactOpenMP, INIT_EXP), OPEN_COMPACT_HASH_SEARCH_EXP, ##__VA_ARGS__)\
	111
103	112	#define LINEAR_OPEN_COMPACT_HASH_JUMP_EXP(NAME, MODIFIER, B, CL_GLOBAL) ((B * iteration + c) % ((CL_GLOBAL NAME##MODIFIER##Hash_TableData*)tableData)->numBuckets)
104	113
105	114	#define LINEAR_OPEN_COMPACT_HASH_CYCLE_EXP (index == c && iteration > 0)
…	…	@@ -118,6 +127,10 @@ unsigned int c = COMPRESS_FUNC(((CL_GLOBAL NAME##MODIFIER##Hash_TableData*)table
118	127	OPEN_COMPACT_CL_HASH_DEFINE(NAME, MODIFIER##Linear, HASH_ID, KEY_TYPE, VALUE_TYPE, COMPRESS_FUNC, COMPRESS_FUNC_DATA, INIT_EXP)\
119	128	\
120	129
	130	#define LINEAR_OPEN_COMPACT_OPENMP_HASH_DEFINE(NAME, MODIFIER, HASH_ID, KEY_TYPE, VALUE_TYPE, COMPRESS_FUNC, COMPRESS_FUNC_DATA, INIT_EXP, B)\
	131	\
	132	OPEN_COMPACT_OPENMP_HASH_DEFINE(NAME, MODIFIER##Linear, HASH_ID, KEY_TYPE, VALUE_TYPE, COMPRESS_FUNC, COMPRESS_FUNC_DATA, INIT_EXP, LINEAR_OPEN_COMPACT_HASH_JUMP_EXP(NAME, MODIFIER##LinearOpenCompactOpenMP, B,), LINEAR_OPEN_COMPACT_HASH_CYCLE_EXP)\
	133
121	134	#define QUADRATIC_OPEN_COMPACT_HASH_JUMP_EXP(NAME, MODIFIER, A, B, CL_GLOBAL) ((A * iteration * iteration + B * iteration + c) % ((CL_GLOBAL NAME##MODIFIER##Hash_TableData*)tableData)->numBuckets)
122	135
123	136	#define QUADRATIC_OPEN_COMPACT_HASH_CYCLE_EXP(NAME, MODIFIER, CL_GLOBAL) (iteration > ((CL_GLOBAL NAME##MODIFIER##Hash_TableData*)tableData)->numBuckets)
…	…	@@ -142,6 +155,11 @@ unsigned int c = COMPRESS_FUNC(((CL_GLOBAL NAME##MODIFIER##Hash_TableData*)table
142	155	OPEN_COMPACT_CL_HASH_DEFINE(NAME, MODIFIER##Quadratic, HASH_ID, KEY_TYPE, VALUE_TYPE, COMPRESS_FUNC, COMPRESS_FUNC_DATA, QUADRATIC_OPEN_COMPACT_HASH_INIT_EXP(NAME, MODIFIER##QuadraticOpenCompactCL, INIT_EXP))
143	156	\
144	157
	158	#define QUADRATIC_OPEN_COMPACT_OPENMP_HASH_DEFINE(NAME, MODIFIER, HASH_ID, KEY_TYPE, VALUE_TYPE, COMPRESS_FUNC, COMPRESS_FUNC_DATA, INIT_EXP, A, B)\
	159	\
	160	OPEN_COMPACT_OPENMP_HASH_DEFINE(NAME, MODIFIER##Quadratic, HASH_ID, KEY_TYPE, VALUE_TYPE, COMPRESS_FUNC, COMPRESS_FUNC_DATA, QUADRATIC_OPEN_COMPACT_HASH_INIT_EXP(NAME, MODIFIER##QuadraticOpenCompactOpenMP, INIT_EXP), QUADRATIC_OPEN_COMPACT_HASH_JUMP_EXP(NAME, MODIFIER##QuadraticOpenCompactOpenMP, A, B,), QUADRATIC_OPEN_COMPACT_HASH_CYCLE_EXP(NAME, MODIFIER##QuadraticOpenCompactOpenMP,))\
	161	\
	162
145	163	#define LCG_OPEN_COMPACT_HASH_INIT_EXP(NAME, MODIFIER)\
146	164	((NAME##MODIFIER##Hash_TableData*)table->tableData)->compressFuncData.a = HASH_LCG_A;\
147	165	((NAME##MODIFIER##Hash_TableData*)table->tableData)->compressFuncData.c = HASH_LCG_C;\
…	…	@@ -164,6 +182,11 @@ unsigned int c = COMPRESS_FUNC(((CL_GLOBAL NAME##MODIFIER##Hash_TableData*)table
164	182	LINEAR_OPEN_COMPACT_CL_HASH_DEFINE(NAME, LCG, LCG_LINEAR_OPEN_COMPACT_CL_HASH_ID, KEY_TYPE, VALUE_TYPE, NAME##Hash_CompressLCG, NAME##Hash_CompressLCGData, LCG_OPEN_COMPACT_HASH_INIT_EXP(NAME, LCGLinearOpenCompactCL))\
165	183	\
166	184
	185	#define LCG_LINEAR_OPEN_COMPACT_OPENMP_HASH_DEFINE(NAME, KEY_TYPE, VALUE_TYPE, B)\
	186	\
	187	LINEAR_OPEN_COMPACT_OPENMP_HASH_DEFINE(NAME, LCG, LCG_LINEAR_OPEN_COMPACT_OPENMP_HASH_ID, KEY_TYPE, VALUE_TYPE, NAME##Hash_CompressLCG, NAME##Hash_CompressLCGData, LCG_OPEN_COMPACT_HASH_INIT_EXP(NAME, LCGLinearOpenCompactOpenMP), B)\
	188	\
	189
167	190	#define LCG_QUADRATIC_OPEN_COMPACT_HASH_DEFINE(NAME, KEY_TYPE, VALUE_TYPE, A, B)\
168	191	\
169	192	QUADRATIC_OPEN_COMPACT_HASH_DEFINE(NAME, LCG, LCG_QUADRATIC_OPEN_COMPACT_HASH_ID, KEY_TYPE, VALUE_TYPE, NAME##Hash_CompressLCG, NAME##Hash_CompressLCGData, LCG_OPEN_COMPACT_HASH_INIT_EXP(NAME, LCGQuadraticOpenCompact), A, B)\
…	…	@@ -179,3 +202,8 @@ unsigned int c = COMPRESS_FUNC(((CL_GLOBAL NAME##MODIFIER##Hash_TableData*)table
179	202	QUADRATIC_OPEN_COMPACT_CL_HASH_DEFINE(NAME, LCG, LCG_QUADRATIC_OPEN_COMPACT_CL_HASH_ID, KEY_TYPE, VALUE_TYPE, NAME##Hash_CompressLCG, NAME##Hash_CompressLCGData, LCG_OPEN_COMPACT_HASH_INIT_EXP(NAME, LCGQuadraticOpenCompactCL))\
180	203	\
181	204
	205	#define LCG_QUADRATIC_OPEN_COMPACT_OPENMP_HASH_DEFINE(NAME, KEY_TYPE, VALUE_TYPE, A, B)\
	206	\
	207	QUADRATIC_OPEN_COMPACT_OPENMP_HASH_DEFINE(NAME, LCG, LCG_QUADRATIC_OPEN_COMPACT_OPENMP_HASH_ID, KEY_TYPE, VALUE_TYPE, NAME##Hash_CompressLCG, NAME##Hash_CompressLCGData, LCG_OPEN_COMPACT_HASH_INIT_EXP(NAME, LCGQuadraticOpenCompactOpenMP), A, B)\
	208	\
	209

opencog-CompactHash commits / be17f28

Changed files (3) +1494 -519