ce33dfa · opencog-CompactHash

Changed files (3) +41 -34

neigh2d.c +6 -6

…	…	@@ -1750,7 +1750,7 @@ struct neighbor2d *neighbors2d_hasholdlibcpu_opt_3( uint ncells, int mesh_size,
1750	1750	if(DETAILED_TIMING) cpu_timer_start(&tSection);
1751	1751
1752	1752	//Create Hash Table
1753		int *hash = compact_hash_init(smallestProthPrimeAbove(ncells / HASH_LOAD_FACTOR), imaxsize, jmaxsize, 1);
	1753	int *hash = compact_hash_init(smallestProthPrimeAbove(ncells / HASH_LOAD_FACTOR), imaxsize, jmaxsize, 0, 1);
1754	1754
1755	1755	if(DETAILED_TIMING) createCPUTime = cpu_timer_stop(tSection);
1756	1756	if(DETAILED_TIMING) cpu_timer_start(&tSection);
…	…	@@ -1886,7 +1886,7 @@ cl_mem neighbors2d_hashgpu( uint ncells, int mesh_size, int levmx, cl_mem i_buff
1886	1886	if(DETAILED_TIMING) GPUTime = 0;
1887	1887
1888	1888	//Create Hash Table
1889		hash_buffer = hash_init(hash_size, TILE_SIZE, context, queue, &GPUTime);
	1889	hash_buffer = hash_init(hash_size, TILE_SIZE, 0, context, queue, &GPUTime);
1890	1890
1891	1891	if(DETAILED_TIMING) createGPUTime = (double)(GPUTime)*1.0e-9;
1892	1892	if(DETAILED_TIMING) cpu_timer_start(&tSection);
…	…	@@ -2129,7 +2129,7 @@ cl_mem neighbors2d_hashgpu_opt_1( uint ncells, int mesh_size, int levmx, cl_mem
2129	2129	if(DETAILED_TIMING) GPUTime = 0;
2130	2130
2131	2131	//Create Hash Table
2132		hash_buffer = hash_init(hash_size, TILE_SIZE, context, queue, &GPUTime);
	2132	hash_buffer = hash_init(hash_size, TILE_SIZE, 0, context, queue, &GPUTime);
2133	2133
2134	2134	if(DETAILED_TIMING) createGPUTime = (double)(GPUTime)*1.0e-9;
2135	2135	if(DETAILED_TIMING) cpu_timer_start(&tSection);
…	…	@@ -2374,7 +2374,7 @@ cl_mem neighbors2d_hashgpu_opt_2( uint ncells, int mesh_size, int levmx, cl_mem
2374	2374	if(DETAILED_TIMING) GPUTime = 0;
2375	2375
2376	2376	//Create Hash Table
2377		hash_buffer = hash_init(hash_size, TILE_SIZE, context, queue, &GPUTime);
	2377	hash_buffer = hash_init(hash_size, TILE_SIZE, 0, context, queue, &GPUTime);
2378	2378
2379	2379	if(DETAILED_TIMING) createGPUTime = (double)(GPUTime)*1.0e-9;
2380	2380	if(DETAILED_TIMING) cpu_timer_start(&tSection);
…	…	@@ -2620,7 +2620,7 @@ cl_mem neighbors2d_hashgpu_opt_3( uint ncells, int mesh_size, int levmx, cl_mem
2620	2620	if(DETAILED_TIMING) GPUTime = 0;
2621	2621
2622	2622	//Create Hash Table
2623		hash_buffer = hash_init(hash_size, TILE_SIZE, context, queue, &GPUTime);
	2623	hash_buffer = hash_init(hash_size, TILE_SIZE, 0, context, queue, &GPUTime);
2624	2624
2625	2625	if(DETAILED_TIMING) createGPUTime = (double)(GPUTime)*1.0e-9;
2626	2626	if(DETAILED_TIMING) cpu_timer_start(&tSection);
…	…	@@ -2905,7 +2905,7 @@ cl_mem neighbors2d_hasholdlibgpu_opt_3( uint ncells, int mesh_size, int levmx, c
2905	2905	hash_size = gpu_perfect_hash_size;
2906	2906	}
2907	2907
2908		hash_buffer = hash_init(hash_size, TILE_SIZE, context, queue, &GPUTime);
	2908	hash_buffer = hash_init(hash_size, TILE_SIZE, 1, context, queue, &GPUTime);
2909	2909
2910	2910	if(DETAILED_TIMING) createGPUTime = (double)(GPUTime)*1.0e-9;
2911	2911	if(DETAILED_TIMING) cpu_timer_start(&tSection);

simplehash/simplehash.c +31 -24

…	…	@@ -81,7 +81,7 @@ long long get_hashtablesize(void) {
81	81	return(hashtablesize);
82	82	}
83	83
84		int *compact_hash_init(int ncells, uint isize, uint jsize, uint report_level){
	84	int *compact_hash_init(int ncells, uint isize, uint jsize, int do_init, uint report_level){
85	85	hash_ncells = 0;
86	86	write_hash_collisions = 0;
87	87	read_hash_collisions = 0;
…	…	@@ -159,8 +159,10 @@ int *compact_hash_init(int ncells, uint isize, uint jsize, uint report_level){
159	159	hashtablesize = perfect_hash_size;
160	160
161	161	hash = (int )malloc(hashtablesizesizeof(int));
162		for (uint ii = 0; ii<hashtablesize; ii++){
163		hash[ii] = -1;
	162	if (do_init) {
	163	for (uint ii = 0; ii<hashtablesize; ii++){
	164	hash[ii] = -1;
	165	}
164	166	}
165	167
166	168	read_hash = read_hash_perfect;
…	…	@@ -179,9 +181,9 @@ int *compact_hash_init(int ncells, uint isize, uint jsize, uint report_level){
179	181	#ifdef _OPENMP
180	182
181	183	#ifdef __GCC_HAVE_SYNC_COMPARE_AND_SWAP_4
182		int *compact_hash_init_openmp(int ncells, uint isize, uint jsize, uint report_level){
	184	int *compact_hash_init_openmp(int ncells, uint isize, uint jsize, int do_init, uint report_level){
183	185	#else
184		int compact_hash_init_openmp(int ncells, uint isize, uint jsize, uint report_level, omp_lock_t *lock){
	186	int compact_hash_init_openmp(int ncells, uint isize, uint jsize, int do_init, uint report_level, omp_lock_t *lock){
185	187	#endif
186	188
187	189	hash_ncells = 0;
…	…	@@ -287,8 +289,10 @@ int *compact_hash_init_openmp(int ncells, uint isize, uint jsize, uint report_le
287	289	#ifdef _OPENMP
288	290	#pragma omp parallel for
289	291	#endif
290		for (uint ii = 0; ii<hashtablesize; ii++){
291		hash[ii] = -1;
	292	if (do_init) {
	293	for (uint ii = 0; ii<hashtablesize; ii++){
	294	hash[ii] = -1;
	295	}
292	296	}
293	297
294	298	read_hash = read_hash_perfect;
…	…	@@ -1366,7 +1370,7 @@ void hash_lib_init(cl_context context){
1366	1370	if (error != CL_SUCCESS) printf("Error is %d at line %d\n",error,__LINE__);
1367	1371	}
1368	1372
1369		cl_mem hash_init (int hash_size, int TILE_SIZE, cl_context context, cl_command_queue queue, long *gpu_time)
	1373	cl_mem hash_init (int hash_size, int TILE_SIZE, int do_init, cl_context context, cl_command_queue queue, long *gpu_time)
1370	1374	{
1371	1375	cl_int error;
1372	1376
…	…	@@ -1377,31 +1381,34 @@ cl_mem hash_init (int hash_size, int TILE_SIZE, cl_context context, cl_command_q
1377	1381
1378	1382	//init to -1
1379	1383
1380		error = clSetKernelArg(init_kernel, 0, sizeof(cl_uint), &hash_size);
1381		if (error != CL_SUCCESS) printf("Error is %d at line %d\n",error,__LINE__);
1382		error = clSetKernelArg(init_kernel, 1, sizeof(cl_mem), (void*)&hash_buffer);
1383		if (error != CL_SUCCESS) printf("Error is %d at line %d\n",error,__LINE__);
	1384	if (do_init) {
1384	1385
1385		size_t global_work_size;
1386		size_t local_work_size;
	1386	error = clSetKernelArg(init_kernel, 0, sizeof(cl_uint), &hash_size);
	1387	if (error != CL_SUCCESS) printf("Error is %d at line %d\n",error,__LINE__);
	1388	error = clSetKernelArg(init_kernel, 1, sizeof(cl_mem), (void*)&hash_buffer);
	1389	if (error != CL_SUCCESS) printf("Error is %d at line %d\n",error,__LINE__);
1387	1390
1388		local_work_size = TILE_SIZE;
1389		global_work_size = ((hash_size+local_work_size-1)/local_work_size)*local_work_size;
	1391	size_t global_work_size;
	1392	size_t local_work_size;
1390	1393
1391		cl_event hash_init_event;
	1394	local_work_size = TILE_SIZE;
	1395	global_work_size = ((hash_size+local_work_size-1)/local_work_size)*local_work_size;
1392	1396
1393		error = clEnqueueNDRangeKernel(queue, init_kernel, 1, 0, &global_work_size, &local_work_size, 0, NULL, &hash_init_event);
1394		if (error != CL_SUCCESS) printf("Error is %d at line %d\n",error,__LINE__);
	1397	cl_event hash_init_event;
1395	1398
1396		clWaitForEvents(1,&hash_init_event);
	1399	error = clEnqueueNDRangeKernel(queue, init_kernel, 1, 0, &global_work_size, &local_work_size, 0, NULL, &hash_init_event);
	1400	if (error != CL_SUCCESS) printf("Error is %d at line %d\n",error,__LINE__);
1397	1401
1398		clGetEventProfilingInfo(hash_init_event, CL_PROFILING_COMMAND_START, sizeof(gpu_time_start), &gpu_time_start, NULL);
1399		clGetEventProfilingInfo(hash_init_event, CL_PROFILING_COMMAND_END, sizeof(gpu_time_end), &gpu_time_end, NULL);
	1402	clWaitForEvents(1,&hash_init_event);
1400	1403
	1404	clGetEventProfilingInfo(hash_init_event, CL_PROFILING_COMMAND_START, sizeof(gpu_time_start), &gpu_time_start, NULL);
	1405	clGetEventProfilingInfo(hash_init_event, CL_PROFILING_COMMAND_END, sizeof(gpu_time_end), &gpu_time_end, NULL);
1401	1406
1402		*gpu_time = gpu_time_end - gpu_time_start;
1403	1407
1404		clReleaseEvent(hash_init_event);
	1408	*gpu_time = gpu_time_end - gpu_time_start;
	1409
	1410	clReleaseEvent(hash_init_event);
	1411	}
1405	1412
1406	1413	//if (DETAILED_TIMING) printf("\n\tinit %.6lf,", (double)(gpu_time_end - gpu_time_start)*1.0e-9);
1407	1414

simplehash/simplehash.h +4 -4

…	…	@@ -63,13 +63,13 @@ typedef unsigned int uint;
63	63	#endif
64	64	typedef unsigned long ulong;
65	65
66		int *compact_hash_init(int ncells, uint isize, uint jsize, uint report_level);
	66	int *compact_hash_init(int ncells, uint isize, uint jsize, int do_init, uint report_level);
67	67
68	68	#ifdef _OPENMP
69	69	#ifdef __GCC_HAVE_SYNC_COMPARE_AND_SWAP_4
70		int *compact_hash_init_openmp(int ncells, uint isize, uint jsize, uint report_level);
	70	int *compact_hash_init_openmp(int ncells, uint isize, uint jsize, int do_init, uint report_level);
71	71	#else
72		int compact_hash_init_openmp(int ncells, uint isize, uint jsize, uint report_level, omp_lock_t *lock);
	72	int compact_hash_init_openmp(int ncells, uint isize, uint jsize, int do_init, uint report_level, omp_lock_t *lock);
73	73	#endif
74	74	#endif
75	75
…	…	@@ -155,7 +155,7 @@ void final_hash_collision_report(void);
155	155	char *get_hash_kernel_source_string(void);
156	156	#ifdef HAVE_OPENCL
157	157	void hash_lib_init(cl_context context);
158		cl_mem hash_init (int hash_size, int TILE_SIZE, cl_context context, cl_command_queue queue, long *gpu_time);
	158	cl_mem hash_init (int hash_size, int TILE_SIZE, int do_init, cl_context context, cl_command_queue queue, long *gpu_time);
159	159	#endif
160	160	void hash_lib_terminate(void);
161	161

opencog-CompactHash commits / ce33dfa

Changed files (3) +41 -34