d3c9999 · opencog-CompactHash

Changed files (7) +149 -221

neigh2d.c +0 -1

…	…	@@ -51,7 +51,6 @@
51	51	#include "gpu.h"
52	52	#include "HashFactory/HashFactory.h"
53	53	#include "simplehash/simplehash.h"
54		#include "simplehash/simplehashlib.h"
55	54	#include "timer/timer.h"
56	55
57	56	#ifdef HAVE_OPENCL

simplehash/Makefile.am +6 -6

…	…	@@ -1,20 +1,20 @@
1		default: simplehash.c simplehashlib.c libsimplehash.la
2		all: simplehash.c simplehashlib.c libsimplehash.la
	1	default: simplehash.c libsimplehash.la
	2	all: simplehash.c libsimplehash.la
3	3
4	4	AM_MAKEFLAGS = -j 4
5	5
6	6	DEFAULT_INCLUDES=-I. -I..
7	7
8	8	lib_LTLIBRARIES = libsimplehash.la
9		include_HEADERS = simplehash.h simplehashlib.h
	9	include_HEADERS = simplehash.h
10	10
11		H_SRCS = simplehash.h simplehashlib.h
	11	H_SRCS = simplehash.h
12	12
13		C_SRCS = simplehash.c simplehashlib.c
	13	C_SRCS = simplehash.c
14	14
15	15	libsimplehash_la_SOURCES = ${C_SRCS} ${H_SRCS}
16	16
17		simplehashlib.c: simplehashlib_kern.inc simplehash_kern.inc
	17	simplehash.c: simplehashlib_kern.inc simplehash_kern.inc
18	18
19	19	simplehashlib_kern.inc:
20	20	./embed_source.pl simplehashlib_kern.cl > simplehashlib_kern.inc

simplehash/Makefile.in +7 -8

…	…	@@ -122,7 +122,7 @@ am__uninstall_files_from_dir = { \
122	122	am__installdirs = "$(DESTDIR)$(libdir)" "$(DESTDIR)$(includedir)"
123	123	LTLIBRARIES = $(lib_LTLIBRARIES)
124	124	libsimplehash_la_LIBADD =
125		am__objects_1 = simplehash.lo simplehashlib.lo
	125	am__objects_1 = simplehash.lo
126	126	am__objects_2 =
127	127	am_libsimplehash_la_OBJECTS = $(am__objects_1) $(am__objects_2)
128	128	libsimplehash_la_OBJECTS = $(am_libsimplehash_la_OBJECTS)
…	…	@@ -318,9 +318,9 @@ top_srcdir = @top_srcdir@
318	318	AM_MAKEFLAGS = -j 4
319	319	DEFAULT_INCLUDES = -I. -I..
320	320	lib_LTLIBRARIES = libsimplehash.la
321		include_HEADERS = simplehash.h simplehashlib.h
322		H_SRCS = simplehash.h simplehashlib.h
323		C_SRCS = simplehash.c simplehashlib.c
	321	include_HEADERS = simplehash.h
	322	H_SRCS = simplehash.h
	323	C_SRCS = simplehash.c
324	324	libsimplehash_la_SOURCES = ${C_SRCS} ${H_SRCS}
325	325	EXTRA_DIST = simplehash_kern.cl simplehashlib_kern.cl embed_source.pl
326	326	all: all-am
…	…	@@ -403,7 +403,6 @@ distclean-compile:
403	403	-rm -f *.tab.c
404	404
405	405	@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/simplehash.Plo@am__quote@
406		@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/simplehashlib.Plo@am__quote@
407	406
408	407	.c.o:
409	408	@am__fastdepCC_TRUE@ $(AM_V_CC)$(COMPILE) -MT $@ -MD -MP -MF $(DEPDIR)/$*.Tpo -c -o $@ $<
…	…	@@ -661,10 +660,10 @@ uninstall-am: uninstall-includeHEADERS uninstall-libLTLIBRARIES
661	660	tags tags-am uninstall uninstall-am uninstall-includeHEADERS \
662	661	uninstall-libLTLIBRARIES
663	662
664		default: simplehash.c simplehashlib.c libsimplehash.la
665		all: simplehash.c simplehashlib.c libsimplehash.la
	663	default: simplehash.c libsimplehash.la
	664	all: simplehash.c libsimplehash.la
666	665
667		simplehashlib.c: simplehashlib_kern.inc simplehash_kern.inc
	666	simplehash.c: simplehashlib_kern.inc simplehash_kern.inc
668	667
669	668	simplehashlib_kern.inc:
670	669	./embed_source.pl simplehashlib_kern.cl > simplehashlib_kern.inc

simplehash/simplehash.c +122 -0

…	…	@@ -34,6 +34,12 @@
34	34	#include <stdlib.h>
35	35	#include "simplehash.h"
36	36
	37	#ifdef HAVE_OPENCL
	38	#include "simplehashlib_kern.inc"
	39	#include "simplehash_kern.inc"
	40	#endif
	41
	42
37	43	static ulong AA;
38	44	static ulong BB;
39	45	static ulong prime=4294967291;
…	…	@@ -1287,6 +1293,122 @@ void final_hash_collision_report(void){
1287	1293	printf("Final hash collision report -- write/read collisions per cell %lf/%lf\n",write_hash_collisions_runsum/(double)write_hash_collisions_count,read_hash_collisions_runsum/(double)read_hash_collisions_count);
1288	1294	}
1289	1295	}
	1296
	1297	#ifdef HAVE_OPENCL
	1298	char *get_hash_kernel_source_string(void){
	1299	return (char*) simplehash_kern_source;
	1300	}
	1301
	1302	cl_kernel init_kernel;
	1303	void hash_lib_init(cl_context context){
	1304
	1305	cl_int error;
	1306
	1307	cl_program program = clCreateProgramWithSource(context, 1, (const char **)&simplehashlib_kern_source, NULL, &error);
	1308	if (error != CL_SUCCESS){
	1309	printf("clCreateProgramWithSource returned an error %d at line %d in file %s\n", error,__LINE__,__FILE__);
	1310	}
	1311
	1312	size_t nReportSize;
	1313	char* BuildReport;
	1314	#ifdef HAVE_CL_DOUBLE
	1315	error = clBuildProgram(program, 0, NULL, "-DHAVE_CL_DOUBLE", NULL, NULL);
	1316	#else
	1317	error = clBuildProgram(program, 0, NULL, "-DNO_CL_DOUBLE -cl-single-precision-constant", NULL, NULL);
	1318	#endif
	1319	if (error != CL_SUCCESS){
	1320	printf("clBuildProgram returned an error %d at line %d in file %s\n", error,__LINE__,__FILE__);
	1321
	1322	cl_device_id device;
	1323
	1324	clGetContextInfo(context, CL_CONTEXT_DEVICES, 1, &device, NULL);
	1325	if (context == NULL){
	1326	printf("EZCL_DEVTYPE_INIT: Failed to find device and setup context in file %s at line %d\n", __FILE__, __LINE__);
	1327	exit(-1); // No device is available, something is wrong
	1328	}
	1329
	1330	error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &nReportSize);
	1331	if (error != CL_SUCCESS) {
	1332	switch (error){
	1333	case CL_INVALID_DEVICE:
	1334	printf("Invalid device in clProgramBuildInfo\n");
	1335	break;
	1336	case CL_INVALID_VALUE:
	1337	printf("Invalid value in clProgramBuildInfo\n");
	1338	break;
	1339	case CL_INVALID_PROGRAM:
	1340	printf("Invalid program in clProgramBuildInfo\n");
	1341	break;
	1342	}
	1343	}
	1344
	1345	BuildReport = (char *)malloc(nReportSize);
	1346
	1347	error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, nReportSize, BuildReport, NULL);
	1348	if (error != CL_SUCCESS) {
	1349	switch (error){
	1350	case CL_INVALID_DEVICE:
	1351	printf("Invalid device in clProgramBuildInfo\n");
	1352	break;
	1353	case CL_INVALID_VALUE:
	1354	printf("Invalid value in clProgramBuildInfo\n");
	1355	break;
	1356	case CL_INVALID_PROGRAM:
	1357	printf("Invalid program in clProgramBuildInfo\n");
	1358	break;
	1359	}
	1360	}
	1361	printf("%s\n", BuildReport);
	1362	}
	1363
	1364
	1365	init_kernel = clCreateKernel(program, "init_kern", &error);
	1366	if (error != CL_SUCCESS) printf("Error is %d at line %d\n",error,__LINE__);
	1367	}
	1368
	1369	cl_mem hash_init (int hash_size, int TILE_SIZE, cl_context context, cl_command_queue queue, long *gpu_time)
	1370	{
	1371	cl_int error;
	1372
	1373	long gpu_time_start, gpu_time_end;
	1374
	1375	cl_mem hash_buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, hash_size*sizeof(int), NULL, &error);
	1376	if (error != CL_SUCCESS) printf("Error is %d at line %d\n",error,__LINE__);
	1377
	1378	//init to -1
	1379
	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
	1385	size_t global_work_size;
	1386	size_t local_work_size;
	1387
	1388	local_work_size = TILE_SIZE;
	1389	global_work_size = ((hash_size+local_work_size-1)/local_work_size)*local_work_size;
	1390
	1391	cl_event hash_init_event;
	1392
	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__);
	1395
	1396	clWaitForEvents(1,&hash_init_event);
	1397
	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);
	1400
	1401
	1402	*gpu_time = gpu_time_end - gpu_time_start;
	1403
	1404	clReleaseEvent(hash_init_event);
	1405
	1406	//if (DETAILED_TIMING) printf("\n\tinit %.6lf,", (double)(gpu_time_end - gpu_time_start)*1.0e-9);
	1407
	1408	return(hash_buffer);
	1409	}
	1410	#endif
	1411
1290	1412	int read_dev_hash(int hash_method, ulong hashtablesize, ulong AA, ulong BB, ulong hashkey, int *hash){
1291	1413	//int hash_report_level = 3;
1292	1414	int max_collisions_allowed = 1000;

simplehash/simplehash.h +14 -0

…	…	@@ -37,6 +37,13 @@
37	37	#include <omp.h>
38	38	#endif
39	39
	40	#ifdef HAVE_OPENCL
	41	#ifdef __APPLE_CC__
	42	#include <OpenCL/OpenCL.h>
	43	#else
	44	#include <CL/cl.h>
	45	#endif
	46	#endif
40	47
41	48	enum choose_hash_method
42	49	{ METHOD_UNSET = 0, // use 0 for no method set
…	…	@@ -142,6 +149,13 @@ void compact_hash_delete(int *hash);
142	149	void read_hash_collision_report(void);
143	150	void final_hash_collision_report(void);
144	151
	152	char *get_hash_kernel_source_string(void);
	153	#ifdef HAVE_OPENCL
	154	void hash_lib_init(cl_context context);
	155	cl_mem hash_init (int hash_size, int TILE_SIZE, cl_context context, cl_command_queue queue, long *gpu_time);
	156	#endif
	157	void hash_lib_terminate(void);
	158
145	159	#ifdef __cplusplus
146	160	}
147	161	#endif

simplehash/simplehashlib.c +0 -156

…	…	@@ -1,156 +0,0 @@
1		/* Copyright 2013-14. Los Alamos National Security, LLC. This material was produced
2		* under U.S. Government contract DE-AC52-06NA25396 for Los Alamos National
3		* Laboratory (LANL), which is operated by Los Alamos National Security, LLC
4		* for the U.S. Department of Energy. The U.S. Government has rights to use,
5		* reproduce, and distribute this software. NEITHER THE GOVERNMENT NOR LOS
6		* ALAMOS NATIONAL SECURITY, LLC MAKES ANY WARRANTY, EXPRESS OR IMPLIED, OR
7		* ASSUMES ANY LIABILITY FOR THE USE OF THIS SOFTWARE. If software is modified
8		* to produce derivative works, such modified software should be clearly marked,
9		* so as not to confuse it with the version available from LANL.
10		*
11		* Licensed under the Apache License, Version 2.0 (the "License"); you may not
12		* use this file except in compliance with the License. You may obtain a copy
13		* of the License at
14		*
15		* http://www.apache.org/licenses/LICENSE-2.0
16		*
17		* Unless required by applicable law or agreed to in writing, software distributed
18		* under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
19		* CONDITIONS OF ANY KIND, either express or implied. See the License for the
20		* specific language governing permissions and limitations under the License.
21		*
22		* Under this license, it is required to include a reference to this work. We
23		* request that each derivative work contain a reference to LANL Copyright
24		* Disclosure C14043/LA-CC-14-003 so that this work's impact can be roughly
25		* measured. In addition, it is requested that a modifier is included as in
26		* the following example:
27		*
28		* //<Uses \| improves on \| modified from> LANL Copyright Disclosure C14043/LA-CC-14-003
29		*
30		* This is LANL Copyright Disclosure C14043/LA-CC-14-003
31		*/
32
33		#include <stdlib.h>
34		//#ifdef HAVE_CONFIG_H
35		//#include "config.h"
36		//#endif
37		#include "simplehashlib.h"
38		#include "simplehashlib_kern.inc"
39		#include "simplehash_kern.inc"
40
41		static int reportLevel = 0;
42
43		char *get_hash_kernel_source_string(void){
44		return (char*) simplehash_kern_source;
45		}
46
47		#ifdef HAVE_OPENCL
48		cl_kernel init_kernel;
49		void hash_lib_init(cl_context context){
50
51		cl_int error;
52
53		cl_program program = clCreateProgramWithSource(context, 1, (const char **)&simplehashlib_kern_source, NULL, &error);
54		if (error != CL_SUCCESS){
55		printf("clCreateProgramWithSource returned an error %d at line %d in file %s\n", error,__LINE__,__FILE__);
56		}
57
58		size_t nReportSize;
59		char* BuildReport;
60		#ifdef HAVE_CL_DOUBLE
61		error = clBuildProgram(program, 0, NULL, "-DHAVE_CL_DOUBLE", NULL, NULL);
62		#else
63		error = clBuildProgram(program, 0, NULL, "-DNO_CL_DOUBLE -cl-single-precision-constant", NULL, NULL);
64		#endif
65		if (error != CL_SUCCESS){
66		printf("clBuildProgram returned an error %d at line %d in file %s\n", error,__LINE__,__FILE__);
67
68		cl_device_id device;
69
70		clGetContextInfo(context, CL_CONTEXT_DEVICES, 1, &device, NULL);
71		if (context == NULL){
72		printf("EZCL_DEVTYPE_INIT: Failed to find device and setup context in file %s at line %d\n", __FILE__, __LINE__);
73		exit(-1); // No device is available, something is wrong
74		}
75
76		error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &nReportSize);
77		if (error != CL_SUCCESS) {
78		switch (error){
79		case CL_INVALID_DEVICE:
80		printf("Invalid device in clProgramBuildInfo\n");
81		break;
82		case CL_INVALID_VALUE:
83		printf("Invalid value in clProgramBuildInfo\n");
84		break;
85		case CL_INVALID_PROGRAM:
86		printf("Invalid program in clProgramBuildInfo\n");
87		break;
88		}
89		}
90
91		BuildReport = (char *)malloc(nReportSize);
92
93		error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, nReportSize, BuildReport, NULL);
94		if (error != CL_SUCCESS) {
95		switch (error){
96		case CL_INVALID_DEVICE:
97		printf("Invalid device in clProgramBuildInfo\n");
98		break;
99		case CL_INVALID_VALUE:
100		printf("Invalid value in clProgramBuildInfo\n");
101		break;
102		case CL_INVALID_PROGRAM:
103		printf("Invalid program in clProgramBuildInfo\n");
104		break;
105		}
106		}
107		printf("%s\n", BuildReport);
108		}
109
110
111		init_kernel = clCreateKernel(program, "init_kern", &error);
112		if (error != CL_SUCCESS) printf("Error is %d at line %d\n",error,__LINE__);
113		}
114
115		cl_mem hash_init (int hash_size, int TILE_SIZE, cl_context context, cl_command_queue queue, long *gpu_time)
116		{
117		cl_int error;
118
119		long gpu_time_start, gpu_time_end;
120
121		cl_mem hash_buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, hash_size*sizeof(int), NULL, &error);
122		if (error != CL_SUCCESS) printf("Error is %d at line %d\n",error,__LINE__);
123
124		//init to -1
125
126		error = clSetKernelArg(init_kernel, 0, sizeof(cl_uint), &hash_size);
127		if (error != CL_SUCCESS) printf("Error is %d at line %d\n",error,__LINE__);
128		error = clSetKernelArg(init_kernel, 1, sizeof(cl_mem), (void*)&hash_buffer);
129		if (error != CL_SUCCESS) printf("Error is %d at line %d\n",error,__LINE__);
130
131		size_t global_work_size;
132		size_t local_work_size;
133
134		local_work_size = TILE_SIZE;
135		global_work_size = ((hash_size+local_work_size-1)/local_work_size)*local_work_size;
136
137		cl_event hash_init_event;
138
139		error = clEnqueueNDRangeKernel(queue, init_kernel, 1, 0, &global_work_size, &local_work_size, 0, NULL, &hash_init_event);
140		if (error != CL_SUCCESS) printf("Error is %d at line %d\n",error,__LINE__);
141
142		clWaitForEvents(1,&hash_init_event);
143
144		clGetEventProfilingInfo(hash_init_event, CL_PROFILING_COMMAND_START, sizeof(gpu_time_start), &gpu_time_start, NULL);
145		clGetEventProfilingInfo(hash_init_event, CL_PROFILING_COMMAND_END, sizeof(gpu_time_end), &gpu_time_end, NULL);
146
147
148		*gpu_time = gpu_time_end - gpu_time_start;
149
150		clReleaseEvent(hash_init_event);
151
152		//if (DETAILED_TIMING) printf("\n\tinit %.6lf,", (double)(gpu_time_end - gpu_time_start)*1.0e-9);
153
154		return(hash_buffer);
155		}
156		#endif

simplehash/simplehashlib.h +0 -50

…	…	@@ -1,50 +0,0 @@
1		/* Copyright 2013-14. Los Alamos National Security, LLC. This material was produced
2		* under U.S. Government contract DE-AC52-06NA25396 for Los Alamos National
3		* Laboratory (LANL), which is operated by Los Alamos National Security, LLC
4		* for the U.S. Department of Energy. The U.S. Government has rights to use,
5		* reproduce, and distribute this software. NEITHER THE GOVERNMENT NOR LOS
6		* ALAMOS NATIONAL SECURITY, LLC MAKES ANY WARRANTY, EXPRESS OR IMPLIED, OR
7		* ASSUMES ANY LIABILITY FOR THE USE OF THIS SOFTWARE. If software is modified
8		* to produce derivative works, such modified software should be clearly marked,
9		* so as not to confuse it with the version available from LANL.
10		*
11		* Licensed under the Apache License, Version 2.0 (the "License"); you may not
12		* use this file except in compliance with the License. You may obtain a copy
13		* of the License at
14		*
15		* http://www.apache.org/licenses/LICENSE-2.0
16		*
17		* Unless required by applicable law or agreed to in writing, software distributed
18		* under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
19		* CONDITIONS OF ANY KIND, either express or implied. See the License for the
20		* specific language governing permissions and limitations under the License.
21		*
22		* Under this license, it is required to include a reference to this work. We
23		* request that each derivative work contain a reference to LANL Copyright
24		* Disclosure C14043/LA-CC-14-003 so that this work's impact can be roughly
25		* measured. In addition, it is requested that a modifier is included as in
26		* the following example:
27		*
28		* //<Uses \| improves on \| modified from> LANL Copyright Disclosure C14043/LA-CC-14-003
29		*
30		* This is LANL Copyright Disclosure C14043/LA-CC-14-003
31		*/
32
33		#include <stdio.h>
34		//#ifdef HAVE_CONFIG_H
35		//#include "config.h"
36		//#endif
37
38		#ifdef HAVE_OPENCL
39		#ifdef __APPLE_CC__
40		#include <OpenCL/OpenCL.h>
41		#else
42		#include <CL/cl.h>
43		#endif
44		#endif
45
46		char *get_hash_kernel_source_string();
47		#ifdef HAVE_OPENCL
48		void hash_lib_init(cl_context context);
49		cl_mem hash_init (int hash_size, int TILE_SIZE, cl_context context, cl_command_queue queue, long *gpu_time);
50		#endif

opencog-CompactHash commits / d3c9999

Changed files (7) +149 -221