3c59d87 · atomspace-gpu

Changed files (28) +4756 -0

.gitignore +1 -0

…	…	@@ -0,0 +1 @@
	1	build/

CMakeLists.txt +204 -0

…	…	@@ -0,0 +1,204 @@
	1	#
	2	# Master AtomSpace GPU Storage Backend CMake file.
	3	#
	4	# General organization:
	5	# -- check for different compilers, OS'es
	6	# -- search for various required & optional libraries/tools
	7	# -- decide what to build based on above results.
	8	# -- configure various config files.
	9	# -- print pretty summary
	10	#
	11	# Supports CUDA (preferred for NVIDIA) and OpenCL (fallback for AMD/Intel).
	12	# At least one GPU backend must be available.
	13	#
	14
	15	# cogutils already requires 3.12, so may as well ask for that.
	16	CMAKE_MINIMUM_REQUIRED(VERSION 3.12)
	17	IF (COMMAND CMAKE_POLICY)
	18	# These must be explicitly set, as otherwise warnings are printed.
	19	CMAKE_POLICY(SET CMP0003 NEW)
	20	CMAKE_POLICY(SET CMP0037 NEW)
	21	CMAKE_POLICY(SET CMP0045 NEW)
	22	ENDIF (COMMAND CMAKE_POLICY)
	23
	24	IF(CMAKE_VERSION VERSION_GREATER 3.31)
	25	CMAKE_POLICY(SET CMP0177 NEW)
	26	ENDIF(CMAKE_VERSION VERSION_GREATER 3.31)
	27
	28	PROJECT(atomspace-gpu VERSION 1.0.0)
	29
	30	# ----------------------------------------------------------
	31	# User-modifiable options. Feel free to change these!
	32	#
	33	# Uncomment to be in Release mode [default].
	34	# SET(CMAKE_BUILD_TYPE Release)
	35
	36	# Uncomment to build in debug mode.
	37	# SET(CMAKE_BUILD_TYPE Debug)
	38
	39	# Uncomment to be in coverage testing mode.
	40	# SET(CMAKE_BUILD_TYPE Coverage)
	41
	42	# Uncomment to build in profile mode.
	43	# SET(CMAKE_BUILD_TYPE Profile)
	44
	45	# Uncomment to build in release mode with debug information.
	46	# SET(CMAKE_BUILD_TYPE RelWithDebInfo)
	47
	48	# default build type
	49	IF (CMAKE_BUILD_TYPE STREQUAL "")
	50	SET(CMAKE_BUILD_TYPE Release)
	51	ENDIF (CMAKE_BUILD_TYPE STREQUAL "")
	52
	53	MESSAGE(STATUS "Build type: ${CMAKE_BUILD_TYPE}")
	54
	55	ADD_DEFINITIONS(-DPROJECT_SOURCE_DIR="${CMAKE_SOURCE_DIR}"
	56	-DPROJECT_BINARY_DIR="${CMAKE_BINARY_DIR}")
	57
	58	# ===============================================================
	59	# Check for existence of various required, optional packages.
	60	# CogUtil must come first, because it supplies various FindXXX macros.
	61
	62	# Cogutil
	63	FIND_PACKAGE(CogUtil 2.0.1 CONFIG REQUIRED)
	64	IF (NOT COGUTIL_FOUND)
	65	MESSAGE(FATAL_ERROR "CogUtil missing: it is needed!")
	66	ENDIF ()
	67
	68	include(OpenCogGccOptions)
	69	include(OpenCogLibOptions)
	70	include(OpenCogInstallOptions)
	71	include(Summary)
	72
	73	# ===================================================================
	74	# Check for existence of various required, optional packages.
	75
	76	# AtomSpace
	77	FIND_PACKAGE(AtomSpace 5.0.3 CONFIG REQUIRED)
	78	IF (NOT ATOMSPACE_FOUND)
	79	MESSAGE(FATAL_ERROR "AtomSpace missing: it is needed!")
	80	ENDIF ()
	81
	82	FIND_PACKAGE(AtomSpaceStorage CONFIG REQUIRED)
	83	IF (NOT ATOMSPACE_STORAGE_FOUND)
	84	MESSAGE(FATAL_ERROR "AtomSpace Storage missing: it is needed!")
	85	ENDIF ()
	86
	87	# ----------------------------------------------------------
	88	# Find CUDA (optional, preferred for NVIDIA GPUs).
	89	INCLUDE(CheckLanguage)
	90	CHECK_LANGUAGE(CUDA)
	91	IF (CMAKE_CUDA_COMPILER)
	92	MESSAGE(STATUS "CUDA found.")
	93	ADD_DEFINITIONS(-DHAVE_CUDA)
	94	SET(HAVE_CUDA 1)
	95
	96	ENABLE_LANGUAGE(CUDA)
	97	CMAKE_POLICY(SET CMP0104 NEW)
	98
	99	ENDIF (CMAKE_CUDA_COMPILER)
	100
	101	# ----------------------------------------------------------
	102	# Find OpenCL (optional, fallback for AMD/Intel GPUs).
	103	FIND_PACKAGE(OpenCL)
	104	IF (OPENCL_FOUND)
	105	MESSAGE(STATUS "OpenCL found.")
	106	MESSAGE(STATUS "OpenCL includes: ${OpenCL_INCLUDE_DIRS}")
	107	MESSAGE(STATUS "OpenCL libraries: ${OpenCL_LIBRARIES}")
	108
	109	INCLUDE_DIRECTORIES(${OpenCL_INCLUDE_DIRS})
	110	ADD_DEFINITIONS(-DHAVE_OPENCL)
	111	SET(HAVE_OPENCL 1)
	112	ENDIF (OPENCL_FOUND)
	113
	114	# Require at least one GPU backend.
	115	IF (NOT HAVE_OPENCL AND NOT HAVE_CUDA)
	116	MESSAGE(FATAL_ERROR
	117	"No GPU backend found! Need at least one of CUDA or OpenCL.")
	118	ENDIF (NOT HAVE_OPENCL AND NOT HAVE_CUDA)
	119
	120	# ----------------------------------------------------------
	121	# Guile Python and Cython
	122
	123	include(OpenCogFindGuile)
	124	include(OpenCogFindPython)
	125
	126	# ===================================================================
	127	# Include configuration.
	128
	129	# Set default include paths.
	130	INCLUDE_DIRECTORIES(${PROJECT_SOURCE_DIR}
	131	${COGUTIL_INCLUDE_DIR} ${ATOMSPACE_INCLUDE_DIR})
	132
	133	INCLUDE(OpenCogMacros)
	134	INCLUDE(OpenCogGuile)
	135	INCLUDE(OpenCogCython)
	136
	137	# ==========================================================
	138	# Decide what to build, based on the packages found.
	139
	140	ADD_SUBDIRECTORY(cmake)
	141	ADD_SUBDIRECTORY(opencog)
	142
	143	INCLUDE(OpenCogTestOptions)
	144	OPENCOG_SETUP_TESTING()
	145	OPENCOG_ADD_TEST_TARGET(test_gpu
	146	tests/persist/gpu "Running GPU storage tests...")
	147
	148	ADD_CUSTOM_TARGET(cscope
	149	COMMAND find opencog tests -name '.cc' -o -name '.h' -o -name '.cxxtest' -o -name '.scm' -o -name '.cl' -o -name '.cu' > ${CMAKE_SOURCE_DIR}/cscope.files
	150	COMMAND cscope -b
	151	WORKING_DIRECTORY ${CMAKE_SOURCE_DIR}
	152	COMMENT "Generating CScope database"
	153	)
	154
	155	# ===================================================================
	156	# Packaging
	157	## Architecture the package is for.
	158	EXECUTE_PROCESS(COMMAND dpkg --print-architecture
	159	OUTPUT_VARIABLE PACKAGE_ARCHITECTURE
	160	OUTPUT_STRIP_TRAILING_WHITESPACE)
	161	STRING(TIMESTAMP UTC_DATE %Y%m%d UTC)
	162	FILE(WRITE "${PROJECT_BINARY_DIR}/install_manifest.txt")
	163
	164	## Cpack configuration
	165	SET(CPACK_GENERATOR "DEB")
	166	SET(CPACK_PACKAGE_CONTACT "opencog@googlegroups.com")
	167	SET(CPACK_PACKAGE_DESCRIPTION_FILE "${CMAKE_CURRENT_SOURCE_DIR}/README.md")
	168	SET(CPACK_PACKAGE_DIRECTORY "${CMAKE_BINARY_DIR}/packages")
	169	SET(CPACK_PACKAGE_DESCRIPTION_SUMMARY "The AtomSpace GPU Storage Backend (CUDA + OpenCL)")
	170	SET(CPACK_PACKAGE_NAME "atomspace-gpu-dev")
	171	SET(CPACK_PACKAGE_VENDOR "opencog.org")
	172	SET(CPACK_PACKAGE_VERSION "${PROJECT_VERSION}-${UTC_DATE}")
	173	SET(CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE")
	174	SET(CPACK_PACKAGE_FILE_NAME
	175	"${CPACK_PACKAGE_NAME}_${CPACK_PACKAGE_VERSION}_${PACKAGE_ARCHITECTURE}")
	176	SET(CPACK_PACKAGING_INSTALL_PREFIX "/usr/local")
	177
	178	## Debian specific configurations
	179	SET(DEPENDENCY_LIST
	180	"guile-3.0-dev (>= 3.0.0)"
	181	"libstdc++6 (>= 7.0)"
	182	"libcogutil-dev (>= 2.0.1)"
	183	"atomspace-dev (>= 5.0.3)"
	184	)
	185
	186	STRING(REPLACE ";" ", " MAIN_DEPENDENCIES "${DEPENDENCY_LIST}")
	187	SET(CPACK_DEBIAN_PACKAGE_DEPENDS "${MAIN_DEPENDENCIES}")
	188	SET(CPACK_DEBIAN_PACKAGE_SECTION "libdevel")
	189	SET(CPACK_DEBIAN_PACKAGE_HOMEPAGE "http://opencog.org")
	190	INCLUDE(CPack)
	191
	192	# ===================================================================
	193	# Documentation.
	194	FIND_PACKAGE(Doxygen)
	195
	196	# ===================================================================
	197	# Show a summary of what we found, what we will do.
	198
	199	SUMMARY_ADD("CUDA Backend" "CUDA GPU StorageNode" HAVE_CUDA)
	200	SUMMARY_ADD("OpenCL Backend" "OpenCL GPU StorageNode" HAVE_OPENCL)
	201	SUMMARY_ADD("Doxygen" "Code documentation" DOXYGEN_FOUND)
	202	SUMMARY_ADD("Unit tests" "Unit tests" CXXTEST_FOUND)
	203
	204	SUMMARY_SHOW()

LICENSE +693 -0

…	…	@@ -0,0 +1,693 @@
	1	Unless indicated otherwise (a license reference in file headers, or a
	2	LICENSE file in a directory), all OpenCog source files use the AGPL
	3	with additional permissions to link OpenCog libraries with non-AGPL
	4	works; the text of the OpenCog linking exception is located at the end
	5	of this file.
	6
	7
	8
	9	GNU AFFERO GENERAL PUBLIC LICENSE
	10	Version 3, 19 November 2007
	11
	12	Copyright (C) 2007 Free Software Foundation, Inc. <http://fsf.org/>
	13	Everyone is permitted to copy and distribute verbatim copies
	14	of this license document, but changing it is not allowed.
	15
	16	Preamble
	17
	18	The GNU Affero General Public License is a free, copyleft license for
	19	software and other kinds of works, specifically designed to ensure
	20	cooperation with the community in the case of network server software.
	21
	22	The licenses for most software and other practical works are designed
	23	to take away your freedom to share and change the works. By contrast,
	24	our General Public Licenses are intended to guarantee your freedom to
	25	share and change all versions of a program--to make sure it remains free
	26	software for all its users.
	27
	28	When we speak of free software, we are referring to freedom, not
	29	price. Our General Public Licenses are designed to make sure that you
	30	have the freedom to distribute copies of free software (and charge for
	31	them if you wish), that you receive source code or can get it if you
	32	want it, that you can change the software or use pieces of it in new
	33	free programs, and that you know you can do these things.
	34
	35	Developers that use our General Public Licenses protect your rights
	36	with two steps: (1) assert copyright on the software, and (2) offer
	37	you this License which gives you legal permission to copy, distribute
	38	and/or modify the software.
	39
	40	A secondary benefit of defending all users' freedom is that
	41	improvements made in alternate versions of the program, if they
	42	receive widespread use, become available for other developers to
	43	incorporate. Many developers of free software are heartened and
	44	encouraged by the resulting cooperation. However, in the case of
	45	software used on network servers, this result may fail to come about.
	46	The GNU General Public License permits making a modified version and
	47	letting the public access it on a server without ever releasing its
	48	source code to the public.
	49
	50	The GNU Affero General Public License is designed specifically to
	51	ensure that, in such cases, the modified source code becomes available
	52	to the community. It requires the operator of a network server to
	53	provide the source code of the modified version running there to the
	54	users of that server. Therefore, public use of a modified version, on
	55	a publicly accessible server, gives the public access to the source
	56	code of the modified version.
	57
	58	An older license, called the Affero General Public License and
	59	published by Affero, was designed to accomplish similar goals. This is
	60	a different license, not a version of the Affero GPL, but Affero has
	61	released a new version of the Affero GPL which permits relicensing under
	62	this license.
	63
	64	The precise terms and conditions for copying, distribution and
	65	modification follow.
	66
	67	TERMS AND CONDITIONS
	68
	69	0. Definitions.
	70
	71	"This License" refers to version 3 of the GNU Affero General Public License.
	72
	73	"Copyright" also means copyright-like laws that apply to other kinds of
	74	works, such as semiconductor masks.
	75
	76	"The Program" refers to any copyrightable work licensed under this
	77	License. Each licensee is addressed as "you". "Licensees" and
	78	"recipients" may be individuals or organizations.
	79
	80	To "modify" a work means to copy from or adapt all or part of the work
	81	in a fashion requiring copyright permission, other than the making of an
	82	exact copy. The resulting work is called a "modified version" of the
	83	earlier work or a work "based on" the earlier work.
	84
	85	A "covered work" means either the unmodified Program or a work based
	86	on the Program.
	87
	88	To "propagate" a work means to do anything with it that, without
	89	permission, would make you directly or secondarily liable for
	90	infringement under applicable copyright law, except executing it on a
	91	computer or modifying a private copy. Propagation includes copying,
	92	distribution (with or without modification), making available to the
	93	public, and in some countries other activities as well.
	94
	95	To "convey" a work means any kind of propagation that enables other
	96	parties to make or receive copies. Mere interaction with a user through
	97	a computer network, with no transfer of a copy, is not conveying.
	98
	99	An interactive user interface displays "Appropriate Legal Notices"
	100	to the extent that it includes a convenient and prominently visible
	101	feature that (1) displays an appropriate copyright notice, and (2)
	102	tells the user that there is no warranty for the work (except to the
	103	extent that warranties are provided), that licensees may convey the
	104	work under this License, and how to view a copy of this License. If
	105	the interface presents a list of user commands or options, such as a
	106	menu, a prominent item in the list meets this criterion.
	107
	108	1. Source Code.
	109
	110	The "source code" for a work means the preferred form of the work
	111	for making modifications to it. "Object code" means any non-source
	112	form of a work.
	113
	114	A "Standard Interface" means an interface that either is an official
	115	standard defined by a recognized standards body, or, in the case of
	116	interfaces specified for a particular programming language, one that
	117	is widely used among developers working in that language.
	118
	119	The "System Libraries" of an executable work include anything, other
	120	than the work as a whole, that (a) is included in the normal form of
	121	packaging a Major Component, but which is not part of that Major
	122	Component, and (b) serves only to enable use of the work with that
	123	Major Component, or to implement a Standard Interface for which an
	124	implementation is available to the public in source code form. A
	125	"Major Component", in this context, means a major essential component
	126	(kernel, window system, and so on) of the specific operating system
	127	(if any) on which the executable work runs, or a compiler used to
	128	produce the work, or an object code interpreter used to run it.
	129
	130	The "Corresponding Source" for a work in object code form means all
	131	the source code needed to generate, install, and (for an executable
	132	work) run the object code and to modify the work, including scripts to
	133	control those activities. However, it does not include the work's
	134	System Libraries, or general-purpose tools or generally available free
	135	programs which are used unmodified in performing those activities but
	136	which are not part of the work. For example, Corresponding Source
	137	includes interface definition files associated with source files for
	138	the work, and the source code for shared libraries and dynamically
	139	linked subprograms that the work is specifically designed to require,
	140	such as by intimate data communication or control flow between those
	141	subprograms and other parts of the work.
	142
	143	The Corresponding Source need not include anything that users
	144	can regenerate automatically from other parts of the Corresponding
	145	Source.
	146
	147	The Corresponding Source for a work in source code form is that
	148	same work.
	149
	150	2. Basic Permissions.
	151
	152	All rights granted under this License are granted for the term of
	153	copyright on the Program, and are irrevocable provided the stated
	154	conditions are met. This License explicitly affirms your unlimited
	155	permission to run the unmodified Program. The output from running a
	156	covered work is covered by this License only if the output, given its
	157	content, constitutes a covered work. This License acknowledges your
	158	rights of fair use or other equivalent, as provided by copyright law.
	159
	160	You may make, run and propagate covered works that you do not
	161	convey, without conditions so long as your license otherwise remains
	162	in force. You may convey covered works to others for the sole purpose
	163	of having them make modifications exclusively for you, or provide you
	164	with facilities for running those works, provided that you comply with
	165	the terms of this License in conveying all material for which you do
	166	not control copyright. Those thus making or running the covered works
	167	for you must do so exclusively on your behalf, under your direction
	168	and control, on terms that prohibit them from making any copies of
	169	your copyrighted material outside their relationship with you.
	170
	171	Conveying under any other circumstances is permitted solely under
	172	the conditions stated below. Sublicensing is not allowed; section 10
	173	makes it unnecessary.
	174
	175	3. Protecting Users' Legal Rights From Anti-Circumvention Law.
	176
	177	No covered work shall be deemed part of an effective technological
	178	measure under any applicable law fulfilling obligations under article
	179	11 of the WIPO copyright treaty adopted on 20 December 1996, or
	180	similar laws prohibiting or restricting circumvention of such
	181	measures.
	182
	183	When you convey a covered work, you waive any legal power to forbid
	184	circumvention of technological measures to the extent such circumvention
	185	is effected by exercising rights under this License with respect to
	186	the covered work, and you disclaim any intention to limit operation or
	187	modification of the work as a means of enforcing, against the work's
	188	users, your or third parties' legal rights to forbid circumvention of
	189	technological measures.
	190
	191	4. Conveying Verbatim Copies.
	192
	193	You may convey verbatim copies of the Program's source code as you
	194	receive it, in any medium, provided that you conspicuously and
	195	appropriately publish on each copy an appropriate copyright notice;
	196	keep intact all notices stating that this License and any
	197	non-permissive terms added in accord with section 7 apply to the code;
	198	keep intact all notices of the absence of any warranty; and give all
	199	recipients a copy of this License along with the Program.
	200
	201	You may charge any price or no price for each copy that you convey,
	202	and you may offer support or warranty protection for a fee.
	203
	204	5. Conveying Modified Source Versions.
	205
	206	You may convey a work based on the Program, or the modifications to
	207	produce it from the Program, in the form of source code under the
	208	terms of section 4, provided that you also meet all of these conditions:
	209
	210	a) The work must carry prominent notices stating that you modified
	211	it, and giving a relevant date.
	212
	213	b) The work must carry prominent notices stating that it is
	214	released under this License and any conditions added under section
	215	7. This requirement modifies the requirement in section 4 to
	216	"keep intact all notices".
	217
	218	c) You must license the entire work, as a whole, under this
	219	License to anyone who comes into possession of a copy. This
	220	License will therefore apply, along with any applicable section 7
	221	additional terms, to the whole of the work, and all its parts,
	222	regardless of how they are packaged. This License gives no
	223	permission to license the work in any other way, but it does not
	224	invalidate such permission if you have separately received it.
	225
	226	d) If the work has interactive user interfaces, each must display
	227	Appropriate Legal Notices; however, if the Program has interactive
	228	interfaces that do not display Appropriate Legal Notices, your
	229	work need not make them do so.
	230
	231	A compilation of a covered work with other separate and independent
	232	works, which are not by their nature extensions of the covered work,
	233	and which are not combined with it such as to form a larger program,
	234	in or on a volume of a storage or distribution medium, is called an
	235	"aggregate" if the compilation and its resulting copyright are not
	236	used to limit the access or legal rights of the compilation's users
	237	beyond what the individual works permit. Inclusion of a covered work
	238	in an aggregate does not cause this License to apply to the other
	239	parts of the aggregate.
	240
	241	6. Conveying Non-Source Forms.
	242
	243	You may convey a covered work in object code form under the terms
	244	of sections 4 and 5, provided that you also convey the
	245	machine-readable Corresponding Source under the terms of this License,
	246	in one of these ways:
	247
	248	a) Convey the object code in, or embodied in, a physical product
	249	(including a physical distribution medium), accompanied by the
	250	Corresponding Source fixed on a durable physical medium
	251	customarily used for software interchange.
	252
	253	b) Convey the object code in, or embodied in, a physical product
	254	(including a physical distribution medium), accompanied by a
	255	written offer, valid for at least three years and valid for as
	256	long as you offer spare parts or customer support for that product
	257	model, to give anyone who possesses the object code either (1) a
	258	copy of the Corresponding Source for all the software in the
	259	product that is covered by this License, on a durable physical
	260	medium customarily used for software interchange, for a price no
	261	more than your reasonable cost of physically performing this
	262	conveying of source, or (2) access to copy the
	263	Corresponding Source from a network server at no charge.
	264
	265	c) Convey individual copies of the object code with a copy of the
	266	written offer to provide the Corresponding Source. This
	267	alternative is allowed only occasionally and noncommercially, and
	268	only if you received the object code with such an offer, in accord
	269	with subsection 6b.
	270
	271	d) Convey the object code by offering access from a designated
	272	place (gratis or for a charge), and offer equivalent access to the
	273	Corresponding Source in the same way through the same place at no
	274	further charge. You need not require recipients to copy the
	275	Corresponding Source along with the object code. If the place to
	276	copy the object code is a network server, the Corresponding Source
	277	may be on a different server (operated by you or a third party)
	278	that supports equivalent copying facilities, provided you maintain
	279	clear directions next to the object code saying where to find the
	280	Corresponding Source. Regardless of what server hosts the
	281	Corresponding Source, you remain obligated to ensure that it is
	282	available for as long as needed to satisfy these requirements.
	283
	284	e) Convey the object code using peer-to-peer transmission, provided
	285	you inform other peers where the object code and Corresponding
	286	Source of the work are being offered to the general public at no
	287	charge under subsection 6d.
	288
	289	A separable portion of the object code, whose source code is excluded
	290	from the Corresponding Source as a System Library, need not be
	291	included in conveying the object code work.
	292
	293	A "User Product" is either (1) a "consumer product", which means any
	294	tangible personal property which is normally used for personal, family,
	295	or household purposes, or (2) anything designed or sold for incorporation
	296	into a dwelling. In determining whether a product is a consumer product,
	297	doubtful cases shall be resolved in favor of coverage. For a particular
	298	product received by a particular user, "normally used" refers to a
	299	typical or common use of that class of product, regardless of the status
	300	of the particular user or of the way in which the particular user
	301	actually uses, or expects or is expected to use, the product. A product
	302	is a consumer product regardless of whether the product has substantial
	303	commercial, industrial or non-consumer uses, unless such uses represent
	304	the only significant mode of use of the product.
	305
	306	"Installation Information" for a User Product means any methods,
	307	procedures, authorization keys, or other information required to install
	308	and execute modified versions of a covered work in that User Product from
	309	a modified version of its Corresponding Source. The information must
	310	suffice to ensure that the continued functioning of the modified object
	311	code is in no case prevented or interfered with solely because
	312	modification has been made.
	313
	314	If you convey an object code work under this section in, or with, or
	315	specifically for use in, a User Product, and the conveying occurs as
	316	part of a transaction in which the right of possession and use of the
	317	User Product is transferred to the recipient in perpetuity or for a
	318	fixed term (regardless of how the transaction is characterized), the
	319	Corresponding Source conveyed under this section must be accompanied
	320	by the Installation Information. But this requirement does not apply
	321	if neither you nor any third party retains the ability to install
	322	modified object code on the User Product (for example, the work has
	323	been installed in ROM).
	324
	325	The requirement to provide Installation Information does not include a
	326	requirement to continue to provide support service, warranty, or updates
	327	for a work that has been modified or installed by the recipient, or for
	328	the User Product in which it has been modified or installed. Access to a
	329	network may be denied when the modification itself materially and
	330	adversely affects the operation of the network or violates the rules and
	331	protocols for communication across the network.
	332
	333	Corresponding Source conveyed, and Installation Information provided,
	334	in accord with this section must be in a format that is publicly
	335	documented (and with an implementation available to the public in
	336	source code form), and must require no special password or key for
	337	unpacking, reading or copying.
	338
	339	7. Additional Terms.
	340
	341	"Additional permissions" are terms that supplement the terms of this
	342	License by making exceptions from one or more of its conditions.
	343	Additional permissions that are applicable to the entire Program shall
	344	be treated as though they were included in this License, to the extent
	345	that they are valid under applicable law. If additional permissions
	346	apply only to part of the Program, that part may be used separately
	347	under those permissions, but the entire Program remains governed by
	348	this License without regard to the additional permissions.
	349
	350	When you convey a copy of a covered work, you may at your option
	351	remove any additional permissions from that copy, or from any part of
	352	it. (Additional permissions may be written to require their own
	353	removal in certain cases when you modify the work.) You may place
	354	additional permissions on material, added by you to a covered work,
	355	for which you have or can give appropriate copyright permission.
	356
	357	Notwithstanding any other provision of this License, for material you
	358	add to a covered work, you may (if authorized by the copyright holders of
	359	that material) supplement the terms of this License with terms:
	360
	361	a) Disclaiming warranty or limiting liability differently from the
	362	terms of sections 15 and 16 of this License; or
	363
	364	b) Requiring preservation of specified reasonable legal notices or
	365	author attributions in that material or in the Appropriate Legal
	366	Notices displayed by works containing it; or
	367
	368	c) Prohibiting misrepresentation of the origin of that material, or
	369	requiring that modified versions of such material be marked in
	370	reasonable ways as different from the original version; or
	371
	372	d) Limiting the use for publicity purposes of names of licensors or
	373	authors of the material; or
	374
	375	e) Declining to grant rights under trademark law for use of some
	376	trade names, trademarks, or service marks; or
	377
	378	f) Requiring indemnification of licensors and authors of that
	379	material by anyone who conveys the material (or modified versions of
	380	it) with contractual assumptions of liability to the recipient, for
	381	any liability that these contractual assumptions directly impose on
	382	those licensors and authors.
	383
	384	All other non-permissive additional terms are considered "further
	385	restrictions" within the meaning of section 10. If the Program as you
	386	received it, or any part of it, contains a notice stating that it is
	387	governed by this License along with a term that is a further
	388	restriction, you may remove that term. If a license document contains
	389	a further restriction but permits relicensing or conveying under this
	390	License, you may add to a covered work material governed by the terms
	391	of that license document, provided that the further restriction does
	392	not survive such relicensing or conveying.
	393
	394	If you add terms to a covered work in accord with this section, you
	395	must place, in the relevant source files, a statement of the
	396	additional terms that apply to those files, or a notice indicating
	397	where to find the applicable terms.
	398
	399	Additional terms, permissive or non-permissive, may be stated in the
	400	form of a separately written license, or stated as exceptions;
	401	the above requirements apply either way.
	402
	403	8. Termination.
	404
	405	You may not propagate or modify a covered work except as expressly
	406	provided under this License. Any attempt otherwise to propagate or
	407	modify it is void, and will automatically terminate your rights under
	408	this License (including any patent licenses granted under the third
	409	paragraph of section 11).
	410
	411	However, if you cease all violation of this License, then your
	412	license from a particular copyright holder is reinstated (a)
	413	provisionally, unless and until the copyright holder explicitly and
	414	finally terminates your license, and (b) permanently, if the copyright
	415	holder fails to notify you of the violation by some reasonable means
	416	prior to 60 days after the cessation.
	417
	418	Moreover, your license from a particular copyright holder is
	419	reinstated permanently if the copyright holder notifies you of the
	420	violation by some reasonable means, this is the first time you have
	421	received notice of violation of this License (for any work) from that
	422	copyright holder, and you cure the violation prior to 30 days after
	423	your receipt of the notice.
	424
	425	Termination of your rights under this section does not terminate the
	426	licenses of parties who have received copies or rights from you under
	427	this License. If your rights have been terminated and not permanently
	428	reinstated, you do not qualify to receive new licenses for the same
	429	material under section 10.
	430
	431	9. Acceptance Not Required for Having Copies.
	432
	433	You are not required to accept this License in order to receive or
	434	run a copy of the Program. Ancillary propagation of a covered work
	435	occurring solely as a consequence of using peer-to-peer transmission
	436	to receive a copy likewise does not require acceptance. However,
	437	nothing other than this License grants you permission to propagate or
	438	modify any covered work. These actions infringe copyright if you do
	439	not accept this License. Therefore, by modifying or propagating a
	440	covered work, you indicate your acceptance of this License to do so.
	441
	442	10. Automatic Licensing of Downstream Recipients.
	443
	444	Each time you convey a covered work, the recipient automatically
	445	receives a license from the original licensors, to run, modify and
	446	propagate that work, subject to this License. You are not responsible
	447	for enforcing compliance by third parties with this License.
	448
	449	An "entity transaction" is a transaction transferring control of an
	450	organization, or substantially all assets of one, or subdividing an
	451	organization, or merging organizations. If propagation of a covered
	452	work results from an entity transaction, each party to that
	453	transaction who receives a copy of the work also receives whatever
	454	licenses to the work the party's predecessor in interest had or could
	455	give under the previous paragraph, plus a right to possession of the
	456	Corresponding Source of the work from the predecessor in interest, if
	457	the predecessor has it or can get it with reasonable efforts.
	458
	459	You may not impose any further restrictions on the exercise of the
	460	rights granted or affirmed under this License. For example, you may
	461	not impose a license fee, royalty, or other charge for exercise of
	462	rights granted under this License, and you may not initiate litigation
	463	(including a cross-claim or counterclaim in a lawsuit) alleging that
	464	any patent claim is infringed by making, using, selling, offering for
	465	sale, or importing the Program or any portion of it.
	466
	467	11. Patents.
	468
	469	A "contributor" is a copyright holder who authorizes use under this
	470	License of the Program or a work on which the Program is based. The
	471	work thus licensed is called the contributor's "contributor version".
	472
	473	A contributor's "essential patent claims" are all patent claims
	474	owned or controlled by the contributor, whether already acquired or
	475	hereafter acquired, that would be infringed by some manner, permitted
	476	by this License, of making, using, or selling its contributor version,
	477	but do not include claims that would be infringed only as a
	478	consequence of further modification of the contributor version. For
	479	purposes of this definition, "control" includes the right to grant
	480	patent sublicenses in a manner consistent with the requirements of
	481	this License.
	482
	483	Each contributor grants you a non-exclusive, worldwide, royalty-free
	484	patent license under the contributor's essential patent claims, to
	485	make, use, sell, offer for sale, import and otherwise run, modify and
	486	propagate the contents of its contributor version.
	487
	488	In the following three paragraphs, a "patent license" is any express
	489	agreement or commitment, however denominated, not to enforce a patent
	490	(such as an express permission to practice a patent or covenant not to
	491	sue for patent infringement). To "grant" such a patent license to a
	492	party means to make such an agreement or commitment not to enforce a
	493	patent against the party.
	494
	495	If you convey a covered work, knowingly relying on a patent license,
	496	and the Corresponding Source of the work is not available for anyone
	497	to copy, free of charge and under the terms of this License, through a
	498	publicly available network server or other readily accessible means,
	499	then you must either (1) cause the Corresponding Source to be so
	500	available, or (2) arrange to deprive yourself of the benefit of the
	501	patent license for this particular work, or (3) arrange, in a manner
	502	consistent with the requirements of this License, to extend the patent
	503	license to downstream recipients. "Knowingly relying" means you have
	504	actual knowledge that, but for the patent license, your conveying the
	505	covered work in a country, or your recipient's use of the covered work
	506	in a country, would infringe one or more identifiable patents in that
	507	country that you have reason to believe are valid.
	508
	509	If, pursuant to or in connection with a single transaction or
	510	arrangement, you convey, or propagate by procuring conveyance of, a
	511	covered work, and grant a patent license to some of the parties
	512	receiving the covered work authorizing them to use, propagate, modify
	513	or convey a specific copy of the covered work, then the patent license
	514	you grant is automatically extended to all recipients of the covered
	515	work and works based on it.
	516
	517	A patent license is "discriminatory" if it does not include within
	518	the scope of its coverage, prohibits the exercise of, or is
	519	conditioned on the non-exercise of one or more of the rights that are
	520	specifically granted under this License. You may not convey a covered
	521	work if you are a party to an arrangement with a third party that is
	522	in the business of distributing software, under which you make payment
	523	to the third party based on the extent of your activity of conveying
	524	the work, and under which the third party grants, to any of the
	525	parties who would receive the covered work from you, a discriminatory
	526	patent license (a) in connection with copies of the covered work
	527	conveyed by you (or copies made from those copies), or (b) primarily
	528	for and in connection with specific products or compilations that
	529	contain the covered work, unless you entered into that arrangement,
	530	or that patent license was granted, prior to 28 March 2007.
	531
	532	Nothing in this License shall be construed as excluding or limiting
	533	any implied license or other defenses to infringement that may
	534	otherwise be available to you under applicable patent law.
	535
	536	12. No Surrender of Others' Freedom.
	537
	538	If conditions are imposed on you (whether by court order, agreement or
	539	otherwise) that contradict the conditions of this License, they do not
	540	excuse you from the conditions of this License. If you cannot convey a
	541	covered work so as to satisfy simultaneously your obligations under this
	542	License and any other pertinent obligations, then as a consequence you may
	543	not convey it at all. For example, if you agree to terms that obligate you
	544	to collect a royalty for further conveying from those to whom you convey
	545	the Program, the only way you could satisfy both those terms and this
	546	License would be to refrain entirely from conveying the Program.
	547
	548	13. Remote Network Interaction; Use with the GNU General Public License.
	549
	550	Notwithstanding any other provision of this License, if you modify the
	551	Program, your modified version must prominently offer all users
	552	interacting with it remotely through a computer network (if your version
	553	supports such interaction) an opportunity to receive the Corresponding
	554	Source of your version by providing access to the Corresponding Source
	555	from a network server at no charge, through some standard or customary
	556	means of facilitating copying of software. This Corresponding Source
	557	shall include the Corresponding Source for any work covered by version 3
	558	of the GNU General Public License that is incorporated pursuant to the
	559	following paragraph.
	560
	561	Notwithstanding any other provision of this License, you have
	562	permission to link or combine any covered work with a work licensed
	563	under version 3 of the GNU General Public License into a single
	564	combined work, and to convey the resulting work. The terms of this
	565	License will continue to apply to the part which is the covered work,
	566	but the work with which it is combined will remain governed by version
	567	3 of the GNU General Public License.
	568
	569	14. Revised Versions of this License.
	570
	571	The Free Software Foundation may publish revised and/or new versions of
	572	the GNU Affero General Public License from time to time. Such new versions
	573	will be similar in spirit to the present version, but may differ in detail to
	574	address new problems or concerns.
	575
	576	Each version is given a distinguishing version number. If the
	577	Program specifies that a certain numbered version of the GNU Affero General
	578	Public License "or any later version" applies to it, you have the
	579	option of following the terms and conditions either of that numbered
	580	version or of any later version published by the Free Software
	581	Foundation. If the Program does not specify a version number of the
	582	GNU Affero General Public License, you may choose any version ever published
	583	by the Free Software Foundation.
	584
	585	If the Program specifies that a proxy can decide which future
	586	versions of the GNU Affero General Public License can be used, that proxy's
	587	public statement of acceptance of a version permanently authorizes you
	588	to choose that version for the Program.
	589
	590	Later license versions may give you additional or different
	591	permissions. However, no additional obligations are imposed on any
	592	author or copyright holder as a result of your choosing to follow a
	593	later version.
	594
	595	15. Disclaimer of Warranty.
	596
	597	THERE IS NO WARRANTY FOR THE PROGRAM, TO THE EXTENT PERMITTED BY
	598	APPLICABLE LAW. EXCEPT WHEN OTHERWISE STATED IN WRITING THE COPYRIGHT
	599	HOLDERS AND/OR OTHER PARTIES PROVIDE THE PROGRAM "AS IS" WITHOUT WARRANTY
	600	OF ANY KIND, EITHER EXPRESSED OR IMPLIED, INCLUDING, BUT NOT LIMITED TO,
	601	THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
	602	PURPOSE. THE ENTIRE RISK AS TO THE QUALITY AND PERFORMANCE OF THE PROGRAM
	603	IS WITH YOU. SHOULD THE PROGRAM PROVE DEFECTIVE, YOU ASSUME THE COST OF
	604	ALL NECESSARY SERVICING, REPAIR OR CORRECTION.
	605
	606	16. Limitation of Liability.
	607
	608	IN NO EVENT UNLESS REQUIRED BY APPLICABLE LAW OR AGREED TO IN WRITING
	609	WILL ANY COPYRIGHT HOLDER, OR ANY OTHER PARTY WHO MODIFIES AND/OR CONVEYS
	610	THE PROGRAM AS PERMITTED ABOVE, BE LIABLE TO YOU FOR DAMAGES, INCLUDING ANY
	611	GENERAL, SPECIAL, INCIDENTAL OR CONSEQUENTIAL DAMAGES ARISING OUT OF THE
	612	USE OR INABILITY TO USE THE PROGRAM (INCLUDING BUT NOT LIMITED TO LOSS OF
	613	DATA OR DATA BEING RENDERED INACCURATE OR LOSSES SUSTAINED BY YOU OR THIRD
	614	PARTIES OR A FAILURE OF THE PROGRAM TO OPERATE WITH ANY OTHER PROGRAMS),
	615	EVEN IF SUCH HOLDER OR OTHER PARTY HAS BEEN ADVISED OF THE POSSIBILITY OF
	616	SUCH DAMAGES.
	617
	618	17. Interpretation of Sections 15 and 16.
	619
	620	If the disclaimer of warranty and limitation of liability provided
	621	above cannot be given local legal effect according to their terms,
	622	reviewing courts shall apply local law that most closely approximates
	623	an absolute waiver of all civil liability in connection with the
	624	Program, unless a warranty or assumption of liability accompanies a
	625	copy of the Program in return for a fee.
	626
	627	END OF TERMS AND CONDITIONS
	628
	629	How to Apply These Terms to Your New Programs
	630
	631	If you develop a new program, and you want it to be of the greatest
	632	possible use to the public, the best way to achieve this is to make it
	633	free software which everyone can redistribute and change under these terms.
	634
	635	To do so, attach the following notices to the program. It is safest
	636	to attach them to the start of each source file to most effectively
	637	state the exclusion of warranty; and each file should have at least
	638	the "copyright" line and a pointer to where the full notice is found.
	639
	640	<one line to give the program's name and a brief idea of what it does.>
	641	Copyright (C) <year> <name of author>
	642
	643	This program is free software: you can redistribute it and/or modify
	644	it under the terms of the GNU Affero General Public License as published by
	645	the Free Software Foundation, either version 3 of the License, or
	646	(at your option) any later version.
	647
	648	This program is distributed in the hope that it will be useful,
	649	but WITHOUT ANY WARRANTY; without even the implied warranty of
	650	MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
	651	GNU Affero General Public License for more details.
	652
	653	You should have received a copy of the GNU Affero General Public License
	654	along with this program. If not, see <http://www.gnu.org/licenses/>.
	655
	656	Also add information on how to contact you by electronic and paper mail.
	657
	658	If your software can interact with users remotely through a computer
	659	network, you should also make sure that it provides a way for users to
	660	get its source. For example, if your program is a web application, its
	661	interface could display a "Source" link that leads users to an archive
	662	of the code. There are many ways you could offer source, and different
	663	solutions will be better for different programs; see section 13 for the
	664	specific requirements.
	665
	666	You should also get your employer (if you work as a programmer) or school,
	667	if any, to sign a "copyright disclaimer" for the program, if necessary.
	668	For more information on this, and how to apply and follow the GNU AGPL, see
	669	<http://www.gnu.org/licenses/>.
	670
	671
	672	Linking Exception
	673
	674	Linking this library statically or dynamically with other modules is making a
	675	combined work based on this library. Thus, the terms and conditions of the GNU
	676	Affero General Public License cover the whole combination.
	677
	678	As a special exception, the copyright holders of this library give you
	679	permission to link this library with independent modules to produce an
	680	executable, regardless of the license terms of these independent modules,
	681	and to copy and distribute the resulting executable under terms of your choice,
	682	provided that you also meet, for each linked independent module, the terms
	683	and conditions of the license of that module. An independent module is a
	684	module which is not derived from or based on this library. If you modify
	685	this library, you may extend this exception to your version of the library,
	686	but you are not obligated to do so. If you do not wish to do so, delete this
	687	exception statement from your version.
	688
	689	As such, it can be used to run, create and distribute a large class of
	690	MindAgents and other applications. When OpenCog is used unmodified as the core
	691	class library for an OpenCog instance, compiler for the combo language, or for
	692	a program written in the combo programming language it does not affect the
	693	licensing for distributing those programs directly.

README.md +146 -0

…	…	@@ -0,0 +1,146 @@
	1
	2	AtomSpace GPU Storage Backend
	3	=====================================
	4
	5	A GPU-backed [StorageNode](https://wiki.opencog.org/w/StorageNode)
	6	for the [OpenCog AtomSpace](https://github.com/opencog/atomspace).
	7	Allows atoms and values to be stored in GPU VRAM for high-throughput
	8	GPU-resident computation.
	9
	10	Same API as [RocksStorageNode](https://github.com/opencog/atomspace-rocks)
	11	(disk) and [CogStorageNode](https://github.com/opencog/atomspace-cog)
	12	(network), but backed by GPU SoA (Structure-of-Arrays) pools.
	13
	14	Supports CUDA (preferred for NVIDIA GPUs) and OpenCL (fallback
	15	for AMD/Intel). At least one must be available at build time. Backend
	16	is selected at runtime based on URI and hardware availability.
	17
	18	Architecture
	19	------------
	20	```
	21	CPU AtomSpace
	22	\|
	23	+-- RocksStorageNode -> disk (RocksDB)
	24	+-- CogStorageNode -> network (CogServer)
	25	+-- GpuStorageNode -> GPU VRAM (CUDA or OpenCL)
	26	\|
	27	+-- GpuBackend (abstract interface)
	28	\| \|
	29	\| +-- CudaBackend (preferred, NVIDIA)
	30	\| +-- OpenCLBackend (fallback, AMD/Intel/NVIDIA)
	31	\|
	32	+-- WordPool 128K nodes (name_hash, count, marginal, class_id)
	33	+-- PairPool 4M binary links (word_a, word_b, count, mi, flags)
	34	+-- SectionPool 1M sections (word, disjunct_hash, count)
	35	```
	36
	37	Each pool uses a lock-free GPU hash table for O(1) lookup and a bump
	38	allocator for slot allocation. CUDA backend uses device memory with
	39	managed staging buffers; OpenCL backend uses cl::Buffer with JIT
	40	kernel compilation.
	41
	42	URI Format
	43	----------
	44	```
	45	gpu://: first available device (CUDA preferred)
	46	gpu://NVIDIA:RTX NVIDIA device containing "RTX" (auto backend)
	47	gpu://cuda:: force CUDA backend, any device
	48	gpu://opencl:: force OpenCL backend, any device
	49	gpu://opencl:Intel: OpenCL on first Intel device
	50	```
	51
	52	Building
	53	--------
	54	Prerequisites: [CogUtil](https://github.com/opencog/cogutil),
	55	[AtomSpace](https://github.com/opencog/atomspace),
	56	[AtomSpace-Storage](https://github.com/opencog/atomspace-storage),
	57	and at least one of:
	58	- CUDA Toolkit (11.0+) for NVIDIA GPUs
	59	- OpenCL (1.2+ runtime and headers) for AMD/Intel/NVIDIA GPUs
	60
	61	```bash
	62	mkdir build && cd build
	63	cmake ..
	64	make -j$(nproc)
	65	make test
	66	sudo make install
	67	```
	68
	69	CMake will auto-detect available backends. To force a specific backend:
	70	```bash
	71	cmake .. -DHAVE_CUDA=OFF # OpenCL only
	72	cmake .. -DHAVE_OPENCL=OFF # CUDA only (not recommended)
	73	```
	74
	75	Usage (C++)
	76	-----------
	77	```cpp
	78	#include <opencog/persist/gpu/GpuStorageNode.h>
	79
	80	AtomSpacePtr as = createAtomSpace();
	81	Handle hsn = as->add_node(GPU_STORAGE_NODE, std::string("gpu://:"));
	82	GpuStorageNodePtr store = GpuStorageNodeCast(hsn);
	83
	84	store->open(); // Selects CUDA if available, else OpenCL
	85
	86	// Store a node with values
	87	Handle a(createNode(SCHEMA_NODE, "someWord"));
	88	a->setValue(truth_key(), createFloatValue({42.0, 3.7}));
	89	store->storeAtom(a, true);
	90	store->barrier();
	91
	92	// Fetch it back
	93	Handle b(createNode(SCHEMA_NODE, "someWord"));
	94	store->getAtom(b);
	95	// b now has the same FloatValue
	96
	97	store->close();
	98	```
	99
	100	Usage (Scheme)
	101	--------------
	102	```scheme
	103	(use-modules (opencog) (opencog persist) (opencog persist-gpu))
	104
	105	(define gpu (GpuStorageNode "gpu://:"))
	106	(cog-open gpu)
	107
	108	(define a (Concept "someWord"))
	109	(cog-set-value! a (Predicate "-TruthValueKey-")
	110	(FloatValue 42.0 3.7))
	111	(store-atom a)
	112	(barrier gpu)
	113
	114	; Fetch back
	115	(define b (Concept "someWord"))
	116	(fetch-atom b)
	117
	118	(cog-close gpu)
	119	```
	120
	121	GPU Memory Usage
	122	----------------
	123	\| Pool \| Capacity \| Hash Table \| GPU Memory \|
	124	\|----------\|----------\|------------\|------------\|
	125	\| Words \| 128K \| 256K \| ~4 MB \|
	126	\| Pairs \| 4M \| 8M \| ~160 MB \|
	127	\| Sections \| 1M \| 2M \| ~40 MB \|
	128	\| Total \| \| \| ~204 MB \|
	129
	130	Phases
	131	------
	132	- Phase 1 (complete): Store/fetch round-trip. 8/8 tests pass.
	133	- Phase 2 (planned): `runQuery()` -- execute Atomese expressions on GPU.
	134	- Phase 3 (planned): JIT compilation -- fused kernels from AST.
	135
	136	Known Limitations
	137	-----------------
	138	- GPU storage is volatile -- data is lost when the connection closes.
	139	Use RocksStorageNode for persistence.
	140	- Only `truth_key()` values are stored. Other value keys are silently
	141	ignored (GPU has fixed SoA layout).
	142	- Only binary links map to the pair pool. Higher-arity links store
	143	their outgoing nodes but not link-level values.
	144	- CUDA backend requires Compute Capability 5.0+ (Maxwell or newer).
	145	- OpenCL backend requires `enqueueWriteBuffer` with `CL_TRUE` (blocking)
	146	when the host buffer goes out of scope immediately.

cmake/AtomSpaceGPUConfig.cmake.in +15 -0

…	…	@@ -0,0 +1,15 @@
	1	@PACKAGE_INIT@
	2
	3	include("@CMAKE_INSTALL_PREFIX@/lib/cmake/AtomSpaceGPU/AtomSpaceGPUTargets.cmake")
	4
	5	link_directories(
	6	"@CMAKE_INSTALL_PREFIX@/lib/opencog/"
	7	)
	8	set(ATOMSPACE_GPU_LIBRARIES
	9	persist-gpu
	10	)
	11
	12	set(ATOMSPACE_GPU_DATA_DIR "@CMAKE_INSTALL_PREFIX@/share/opencog")
	13	set(ATOMSPACE_GPU_INCLUDE_DIR "@CMAKE_INSTALL_PREFIX@/include/")
	14	set(ATOMSPACE_GPU_VERSION "@PROJECT_VERSION@")
	15	set(ATOMSPACE_GPU_FOUND 1)

cmake/CMakeLists.txt +33 -0

…	…	@@ -0,0 +1,33 @@
	1
	2	# CMake boilerplate that allows users to do
	3	# find_package(AtomSpaceGPU REQUIRED 1.0.0)
	4	# and have it work.
	5
	6	include(CMakePackageConfigHelpers)
	7
	8	set(ConfigPackageLocation lib/cmake/AtomSpaceGPU)
	9
	10	install(EXPORT AtomSpaceGPUTargets
	11	FILE AtomSpaceGPUTargets.cmake
	12	DESTINATION ${ConfigPackageLocation}
	13	)
	14
	15	configure_package_config_file(AtomSpaceGPUConfig.cmake.in
	16	${CMAKE_CURRENT_BINARY_DIR}/AtomSpaceGPUConfig.cmake
	17	INSTALL_DESTINATION ${ConfigPackageLocation}
	18	PATH_VARS CMAKE_INSTALL_PREFIX
	19	)
	20
	21	write_basic_package_version_file(
	22	"${CMAKE_CURRENT_BINARY_DIR}/AtomSpaceGPUConfigVersion.cmake"
	23	VERSION ${PROJECT_VERSION}
	24	COMPATIBILITY SameMajorVersion
	25	)
	26
	27	INSTALL(FILES
	28	${CMAKE_CURRENT_BINARY_DIR}/AtomSpaceGPUConfigVersion.cmake
	29	${CMAKE_CURRENT_BINARY_DIR}/AtomSpaceGPUConfig.cmake
	30	DESTINATION ${ConfigPackageLocation}
	31	)
	32
	33	# -----------------------------------------------------------

opencog/CMakeLists.txt +16 -0

…	…	@@ -0,0 +1,16 @@
	1
	2	# This declaration must happen before subdirectory traversal.
	3	DECLARE_GUILE_CONFIG_TARGET(SCM_CONFIG "opencog gpu-config" "OPENCOG_TEST")
	4
	5	ADD_SUBDIRECTORY (persist)
	6
	7	# ------------------------------------------------------------------
	8	ADD_GUILE_MODULE (FILES
	9	scm/persist-gpu.scm
	10	MODULE_DESTINATION "${GUILE_SITE_DIR}/opencog"
	11	)
	12
	13	WRITE_GUILE_CONFIG(${GUILE_BIN_DIR}/opencog/gpu-config.scm SCM_CONFIG TRUE)
	14	WRITE_GUILE_CONFIG(${GUILE_BIN_DIR}/opencog/gpu-config-installable.scm SCM_CONFIG FALSE)
	15
	16	INSTALL(FILES ${GUILE_BIN_DIR}/opencog/gpu-config-installable.scm DESTINATION ${GUILE_SITE_DIR}/opencog RENAME gpu-config.scm)

opencog/persist/CMakeLists.txt +6 -0

…	…	@@ -0,0 +1,6 @@
	1
	2	# The atom_types.h file is written to the build directory
	3	INCLUDE_DIRECTORIES(${CMAKE_BINARY_DIR})
	4
	5	ADD_SUBDIRECTORY(gpu-types)
	6	ADD_SUBDIRECTORY(gpu)

opencog/persist/gpu-types/CMakeLists.txt +44 -0

…	…	@@ -0,0 +1,44 @@
	1	# Generate the atom types
	2	OPENCOG_GEN_CXX_ATOMTYPES(atom_types.script
	3	atom_types.h
	4	atom_types.definitions
	5	atom_types.inheritance)
	6
	7	OPENCOG_GEN_SCM_ATOMTYPES(atom_types.script persist_gpu_types.scm)
	8	OPENCOG_GEN_PYTHON_ATOMTYPES(atom_types.script persist_gpu_types.pyx)
	9
	10	ADD_CUSTOM_TARGET(persist_gpu_atom_types
	11	DEPENDS atom_types.h persist_gpu_types.scm persist_gpu_types.pyx)
	12
	13	# Build the shared library
	14	ADD_LIBRARY (persist-gpu-types SHARED
	15	atom_types_init.cc
	16	)
	17
	18	# The atom_types.h file is written to the build directory
	19	TARGET_INCLUDE_DIRECTORIES(persist-gpu-types PRIVATE ${CMAKE_BINARY_DIR})
	20
	21	TARGET_LINK_LIBRARIES(persist-gpu-types
	22	${ATOMSPACE_STORAGE_LIBRARIES}
	23	${ATOMSPACE_LIBRARIES}
	24	)
	25
	26	ADD_DEPENDENCIES(persist-gpu-types persist_gpu_atom_types)
	27
	28	# Install targets
	29	INSTALL (TARGETS persist-gpu-types EXPORT AtomSpaceGPUTargets
	30	DESTINATION "${CMAKE_INSTALL_LIBDIR}/opencog"
	31	)
	32
	33	INSTALL (FILES
	34	${CMAKE_CURRENT_BINARY_DIR}/atom_types.h
	35	DESTINATION "include/opencog/persist/gpu-types"
	36	)
	37
	38	ADD_GUILE_MODULE (FILES
	39	${CMAKE_CURRENT_BINARY_DIR}/persist_gpu_types.scm
	40	MODULE_DESTINATION "${GUILE_SITE_DIR}/opencog/persist/gpu-types"
	41	DEPENDS persist_gpu_atom_types
	42	)
	43
	44	ADD_GUILE_EXTENSION(SCM_CONFIG persist-gpu-types "opencog-ext-path-persist-gpu-types")

opencog/persist/gpu-types/atom_types.script +6 -0

…	…	@@ -0,0 +1,6 @@
	1	//
	2	// Script for generating the GpuStorageNode atom type.
	3	//
	4	GPU_STORAGE_NODE <- STORAGE_NODE
	5
	6	// That's all folks! Nothing more is needed here!

opencog/persist/gpu-types/atom_types_init.cc +29 -0

…	…	@@ -0,0 +1,29 @@
	1	/*
	2	* opencog/persist/gpu-types/atom_types_init.cc
	3	*
	4	* Copyright (C) 2025 OpenCog Foundation
	5	* All Rights Reserved
	6	*
	7	* This program is free software; you can redistribute it and/or modify
	8	* it under the terms of the GNU Affero General Public License v3 as
	9	* published by the Free Software Foundation and including the exceptions
	10	* at http://opencog.org/wiki/Licenses
	11	*
	12	* This program is distributed in the hope that it will be useful,
	13	* but WITHOUT ANY WARRANTY; without even the implied warranty of
	14	* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
	15	* GNU General Public License for more details.
	16	*
	17	* You should have received a copy of the GNU Affero General Public License
	18	* along with this program; if not, write to:
	19	* Free Software Foundation, Inc.,
	20	* 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
	21	*/
	22
	23	#include <opencog/persist/storage/storage_types.h>
	24	#include "opencog/persist/gpu-types/atom_types.definitions"
	25
	26	#define INHERITANCE_FILE "opencog/persist/gpu-types/atom_types.inheritance"
	27	#define INITNAME persist_gpu_types_init
	28
	29	#include <opencog/atoms/atom_types/atom_types.cc>

opencog/persist/gpu/CMakeLists.txt +71 -0

…	…	@@ -0,0 +1,71 @@
	1	#
	2	# Build the GPU driver for the AtomSpace.
	3	# Supports CUDA (preferred) and OpenCL (fallback).
	4	#
	5
	6	# -- Core StorageNode (backend-agnostic) --
	7	SET(GPU_SOURCES GpuStorageNode.cc)
	8	SET(GPU_LINK_LIBS
	9	persist-gpu-types
	10	persist
	11	${ATOMSPACE_STORAGE_LIBRARIES}
	12	${ATOMSPACE_LIBRARIES}
	13	)
	14	SET(GPU_INSTALL_HEADERS
	15	GpuStorageNode.h
	16	GpuBackend.h
	17	gpu-pool-defs.h
	18	)
	19
	20	# -- OpenCL backend (optional) --
	21	IF (HAVE_OPENCL)
	22	LIST(APPEND GPU_SOURCES OpenCLBackend.cc)
	23	LIST(APPEND GPU_LINK_LIBS OpenCL::OpenCL)
	24	LIST(APPEND GPU_INSTALL_HEADERS opencl-headers.h OpenCLBackend.h)
	25	ENDIF (HAVE_OPENCL)
	26
	27	# -- CUDA backend (optional, preferred) --
	28	IF (HAVE_CUDA)
	29	LIST(APPEND GPU_SOURCES CudaBackend.cu)
	30	LIST(APPEND GPU_INSTALL_HEADERS CudaBackend.h)
	31	ENDIF (HAVE_CUDA)
	32
	33	# -- Build the library --
	34	ADD_LIBRARY(persist-gpu SHARED ${GPU_SOURCES})
	35
	36	# The atom_types.h file is written to the build directory
	37	TARGET_INCLUDE_DIRECTORIES(persist-gpu PRIVATE ${CMAKE_BINARY_DIR})
	38
	39	TARGET_LINK_LIBRARIES(persist-gpu ${GPU_LINK_LIBS})
	40
	41	IF (HAVE_CUDA)
	42	SET_TARGET_PROPERTIES(persist-gpu PROPERTIES
	43	CUDA_SEPARABLE_COMPILATION ON
	44	CUDA_ARCHITECTURES "50;72;75;86;89"
	45	)
	46	ENDIF (HAVE_CUDA)
	47
	48	ADD_DEPENDENCIES(persist-gpu persist_gpu_atom_types)
	49
	50	ADD_GUILE_EXTENSION(SCM_CONFIG persist-gpu
	51	"opencog-ext-path-persist-gpu")
	52
	53	# The EXPORT is needed to autogenerate CMake boilerplate in the lib
	54	# directory that lets other packages FIND_PACKAGE(AtomSpaceGPU)
	55	INSTALL (TARGETS persist-gpu
	56	EXPORT AtomSpaceGPUTargets
	57	DESTINATION "${CMAKE_INSTALL_LIBDIR}/opencog"
	58	)
	59
	60	INSTALL (FILES ${GPU_INSTALL_HEADERS}
	61	DESTINATION "include/opencog/persist/gpu"
	62	)
	63
	64	# Install OpenCL kernel files (needed at runtime for JIT compilation)
	65	IF (HAVE_OPENCL)
	66	INSTALL (FILES
	67	gpu-hashtable.cl
	68	gpu-atomspace.cl
	69	DESTINATION "share/opencog/opencl/gpu"
	70	)
	71	ENDIF (HAVE_OPENCL)

opencog/persist/gpu/CudaBackend.cu +642 -0

…	…	@@ -0,0 +1,642 @@
	1	/*
	2	* opencog/persist/gpu/CudaBackend.cu
	3	*
	4	* CUDA implementation of GpuBackend.
	5	* Uses device memory with CUDA kernels for hash table operations,
	6	* and cudaMemcpy for direct value read/write.
	7	*
	8	* Managed staging buffers avoid per-operation cudaMalloc overhead.
	9	*
	10	* Copyright (C) 2025 OpenCog Foundation
	11	*
	12	* SPDX-License-Identifier: AGPL-3.0-or-later
	13	*/
	14
	15	#include <cstdio>
	16	#include <cstring>
	17
	18	#include <opencog/util/exceptions.h>
	19	#include <opencog/util/Logger.h>
	20
	21	#include "CudaBackend.h"
	22
	23	using namespace opencog;
	24
	25	// ==============================================================
	26	// CUDA Device Kernels
	27	// ==============================================================
	28
	29	// splitmix64 hash function (same as OpenCL gpu-hashtable.cl)
	30	__device__ __forceinline__
	31	uint64_t cuda_hash(uint64_t key)
	32	{
	33	key ^= key >> 30;
	34	key *= 0xBF58476D1CE4E5B9ULL;
	35	key ^= key >> 27;
	36	key *= 0x94D049BB133111EBULL;
	37	key ^= key >> 31;
	38	return key;
	39	}
	40
	41	#define CUDA_HT_EMPTY_KEY 0xFFFFFFFFFFFFFFFFULL
	42	#define CUDA_HT_EMPTY_VALUE 0xFFFFFFFFU
	43	#define CUDA_MAX_PROBES 4096
	44
	45	// ----------------------------------------------------------
	46	// Generic hash table lookup kernel
	47	__global__ void cuda_ht_lookup(
	48	const uint64_t* __restrict__ keys,
	49	const uint32_t* __restrict__ values,
	50	uint64_t capacity,
	51	const uint64_t* __restrict__ query_keys,
	52	uint32_t* __restrict__ results,
	53	uint32_t num_queries)
	54	{
	55	uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
	56	if (idx >= num_queries) return;
	57
	58	uint64_t key = query_keys[idx];
	59	uint64_t slot = cuda_hash(key) % capacity;
	60
	61	for (int probe = 0; probe < CUDA_MAX_PROBES; probe++)
	62	{
	63	uint64_t existing = keys[slot];
	64	if (existing == key)
	65	{
	66	results[idx] = values[slot];
	67	return;
	68	}
	69	if (existing == CUDA_HT_EMPTY_KEY)
	70	{
	71	results[idx] = CUDA_HT_EMPTY_VALUE;
	72	return;
	73	}
	74	slot = (slot + 1) % capacity;
	75	}
	76	results[idx] = CUDA_HT_EMPTY_VALUE;
	77	}
	78
	79	// ----------------------------------------------------------
	80	// Generic hash table delete kernel (tombstone)
	81	__global__ void cuda_ht_delete(
	82	uint64_t* __restrict__ keys,
	83	uint32_t* __restrict__ values,
	84	uint64_t capacity,
	85	const uint64_t* __restrict__ del_keys,
	86	uint32_t num_items)
	87	{
	88	uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
	89	if (idx >= num_items) return;
	90
	91	uint64_t key = del_keys[idx];
	92	uint64_t slot = cuda_hash(key) % capacity;
	93
	94	for (int probe = 0; probe < CUDA_MAX_PROBES; probe++)
	95	{
	96	uint64_t existing = keys[slot];
	97	if (existing == key)
	98	{
	99	// Tombstone: mark as deleted (use EMPTY-1 as tombstone)
	100	keys[slot] = 0xFFFFFFFFFFFFFFFEULL;
	101	values[slot] = CUDA_HT_EMPTY_VALUE;
	102	return;
	103	}
	104	if (existing == CUDA_HT_EMPTY_KEY) return; // not found
	105	slot = (slot + 1) % capacity;
	106	}
	107	}
	108
	109	// ----------------------------------------------------------
	110	// Word find-or-create: hash table insert + pool allocation
	111	__global__ void cuda_word_find_or_create(
	112	uint64_t* __restrict__ ht_keys,
	113	uint32_t* __restrict__ ht_values,
	114	uint64_t* __restrict__ name_hashes,
	115	double* __restrict__ word_counts,
	116	uint32_t* __restrict__ class_ids,
	117	uint32_t* __restrict__ next_free,
	118	const uint64_t* __restrict__ in_hashes,
	119	uint32_t* __restrict__ out_indices,
	120	uint32_t num_items)
	121	{
	122	uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
	123	if (idx >= num_items) return;
	124
	125	uint64_t key = in_hashes[idx];
	126	uint64_t slot = cuda_hash(key) % GPU_WORD_HT_CAPACITY;
	127
	128	for (int probe = 0; probe < CUDA_MAX_PROBES; probe++)
	129	{
	130	uint64_t old = atomicCAS(
	131	(unsigned long long*)&ht_keys[slot],
	132	(unsigned long long)CUDA_HT_EMPTY_KEY,
	133	(unsigned long long)key);
	134
	135	if (old == key)
	136	{
	137	// Already exists
	138	out_indices[idx] = ht_values[slot];
	139	return;
	140	}
	141	if (old == CUDA_HT_EMPTY_KEY)
	142	{
	143	// We won the slot; allocate pool entry
	144	uint32_t pool_idx = atomicAdd(next_free, 1);
	145	ht_values[slot] = pool_idx;
	146	name_hashes[pool_idx] = key;
	147	word_counts[pool_idx] = 0.0;
	148	class_ids[pool_idx] = 0;
	149	out_indices[idx] = pool_idx;
	150	return;
	151	}
	152	slot = (slot + 1) % GPU_WORD_HT_CAPACITY;
	153	}
	154	out_indices[idx] = CUDA_HT_EMPTY_VALUE;
	155	}
	156
	157	// ----------------------------------------------------------
	158	// Pair find-or-create: hash table insert + pool allocation
	159	__global__ void cuda_pair_find_or_create(
	160	uint64_t* __restrict__ ht_keys,
	161	uint32_t* __restrict__ ht_values,
	162	uint32_t* __restrict__ pair_wa,
	163	uint32_t* __restrict__ pair_wb,
	164	double* __restrict__ pair_counts,
	165	double* __restrict__ pair_mis,
	166	uint32_t* __restrict__ pair_flags,
	167	uint32_t* __restrict__ next_free,
	168	const uint32_t* __restrict__ in_wa,
	169	const uint32_t* __restrict__ in_wb,
	170	uint32_t* __restrict__ out_indices,
	171	uint32_t num_items)
	172	{
	173	uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
	174	if (idx >= num_items) return;
	175
	176	uint32_t wa = in_wa[idx];
	177	uint32_t wb = in_wb[idx];
	178	uint32_t lo = min(wa, wb);
	179	uint32_t hi = max(wa, wb);
	180	uint64_t key = ((uint64_t)lo << 32) \| (uint64_t)hi;
	181
	182	uint64_t slot = cuda_hash(key) % GPU_PAIR_HT_CAPACITY;
	183
	184	for (int probe = 0; probe < CUDA_MAX_PROBES; probe++)
	185	{
	186	uint64_t old = atomicCAS(
	187	(unsigned long long*)&ht_keys[slot],
	188	(unsigned long long)CUDA_HT_EMPTY_KEY,
	189	(unsigned long long)key);
	190
	191	if (old == key)
	192	{
	193	out_indices[idx] = ht_values[slot];
	194	return;
	195	}
	196	if (old == CUDA_HT_EMPTY_KEY)
	197	{
	198	uint32_t pool_idx = atomicAdd(next_free, 1);
	199	ht_values[slot] = pool_idx;
	200	pair_wa[pool_idx] = wa;
	201	pair_wb[pool_idx] = wb;
	202	pair_counts[pool_idx] = 0.0;
	203	pair_mis[pool_idx] = 0.0;
	204	pair_flags[pool_idx] = 0;
	205	out_indices[idx] = pool_idx;
	206	return;
	207	}
	208	slot = (slot + 1) % GPU_PAIR_HT_CAPACITY;
	209	}
	210	out_indices[idx] = CUDA_HT_EMPTY_VALUE;
	211	}
	212
	213	// ==============================================================
	214	// Host-side Implementation
	215	// ==============================================================
	216
	217	CudaBackend::CudaBackend()
	218	: _device_id(-1),
	219	_d_word_name_hash(nullptr), _d_word_count(nullptr),
	220	_d_word_mi_marginal(nullptr), _d_word_class_id(nullptr),
	221	_d_word_next_free(nullptr),
	222	_d_pair_word_a(nullptr), _d_pair_word_b(nullptr),
	223	_d_pair_count(nullptr), _d_pair_mi(nullptr),
	224	_d_pair_flags(nullptr), _d_pair_next_free(nullptr),
	225	_d_sec_word(nullptr), _d_sec_djh(nullptr),
	226	_d_sec_count(nullptr), _d_sec_next_free(nullptr),
	227	_d_word_ht_keys(nullptr), _d_word_ht_values(nullptr),
	228	_d_pair_ht_keys(nullptr), _d_pair_ht_values(nullptr),
	229	_d_sec_ht_keys(nullptr), _d_sec_ht_values(nullptr),
	230	_staging_key(nullptr), _staging_result(nullptr),
	231	_staging_wa(nullptr), _staging_wb(nullptr)
	232	{
	233	}
	234
	235	CudaBackend::~CudaBackend()
	236	{
	237	shutdown();
	238	}
	239
	240	void CudaBackend::check_error(cudaError_t err, const char* context)
	241	{
	242	if (err != cudaSuccess)
	243	throw RuntimeException(TRACE_INFO,
	244	"CudaBackend: %s failed: %s\n",
	245	context, cudaGetErrorString(err));
	246	}
	247
	248	// ==============================================================
	249	// Lifecycle
	250
	251	void CudaBackend::init(const std::string& platform_hint,
	252	const std::string& device_hint)
	253	{
	254	int device_count = 0;
	255	check_error(cudaGetDeviceCount(&device_count), "cudaGetDeviceCount");
	256
	257	if (device_count == 0)
	258	throw RuntimeException(TRACE_INFO,
	259	"CudaBackend: no CUDA devices found\n");
	260
	261	// Find a matching device (by name substring)
	262	_device_id = -1;
	263	for (int i = 0; i < device_count; i++)
	264	{
	265	cudaDeviceProp prop;
	266	check_error(cudaGetDeviceProperties(&prop, i),
	267	"cudaGetDeviceProperties");
	268
	269	std::string dname(prop.name);
	270
	271	// Check platform hint (vendor match)
	272	if (0 < platform_hint.size())
	273	{
	274	// NVIDIA is always the CUDA platform, so skip platform matching
	275	// unless the user explicitly asked for non-NVIDIA
	276	if (dname.find(platform_hint) == std::string::npos and
	277	platform_hint != "NVIDIA" and platform_hint != "nvidia")
	278	continue;
	279	}
	280
	281	// Check device hint (model match)
	282	if (0 < device_hint.size() and
	283	dname.find(device_hint) == std::string::npos)
	284	continue;
	285
	286	_device_id = i;
	287	_device_name = dname;
	288	break;
	289	}
	290
	291	if (_device_id < 0)
	292	throw RuntimeException(TRACE_INFO,
	293	"CudaBackend: no matching CUDA device for '%s:%s'\n",
	294	platform_hint.c_str(), device_hint.c_str());
	295
	296	check_error(cudaSetDevice(_device_id), "cudaSetDevice");
	297
	298	logger().info("CudaBackend: Using device %d '%s'\n",
	299	_device_id, _device_name.c_str());
	300	}
	301
	302	void CudaBackend::shutdown()
	303	{
	304	if (_device_id < 0) return;
	305
	306	cudaDeviceSynchronize();
	307
	308	// Free device memory
	309	if (_d_word_name_hash) cudaFree(_d_word_name_hash);
	310	if (_d_word_count) cudaFree(_d_word_count);
	311	if (_d_word_mi_marginal) cudaFree(_d_word_mi_marginal);
	312	if (_d_word_class_id) cudaFree(_d_word_class_id);
	313	if (_d_word_next_free) cudaFree(_d_word_next_free);
	314
	315	if (_d_pair_word_a) cudaFree(_d_pair_word_a);
	316	if (_d_pair_word_b) cudaFree(_d_pair_word_b);
	317	if (_d_pair_count) cudaFree(_d_pair_count);
	318	if (_d_pair_mi) cudaFree(_d_pair_mi);
	319	if (_d_pair_flags) cudaFree(_d_pair_flags);
	320	if (_d_pair_next_free) cudaFree(_d_pair_next_free);
	321
	322	if (_d_sec_word) cudaFree(_d_sec_word);
	323	if (_d_sec_djh) cudaFree(_d_sec_djh);
	324	if (_d_sec_count) cudaFree(_d_sec_count);
	325	if (_d_sec_next_free) cudaFree(_d_sec_next_free);
	326
	327	if (_d_word_ht_keys) cudaFree(_d_word_ht_keys);
	328	if (_d_word_ht_values) cudaFree(_d_word_ht_values);
	329	if (_d_pair_ht_keys) cudaFree(_d_pair_ht_keys);
	330	if (_d_pair_ht_values) cudaFree(_d_pair_ht_values);
	331	if (_d_sec_ht_keys) cudaFree(_d_sec_ht_keys);
	332	if (_d_sec_ht_values) cudaFree(_d_sec_ht_values);
	333
	334	// Free managed staging buffers
	335	if (_staging_key) cudaFree(_staging_key);
	336	if (_staging_result) cudaFree(_staging_result);
	337	if (_staging_wa) cudaFree(_staging_wa);
	338	if (_staging_wb) cudaFree(_staging_wb);
	339
	340	// Reset all pointers
	341	_d_word_name_hash = nullptr;
	342	_d_word_count = nullptr;
	343	_d_word_mi_marginal = nullptr;
	344	_d_word_class_id = nullptr;
	345	_d_word_next_free = nullptr;
	346	_d_pair_word_a = nullptr;
	347	_d_pair_word_b = nullptr;
	348	_d_pair_count = nullptr;
	349	_d_pair_mi = nullptr;
	350	_d_pair_flags = nullptr;
	351	_d_pair_next_free = nullptr;
	352	_d_sec_word = nullptr;
	353	_d_sec_djh = nullptr;
	354	_d_sec_count = nullptr;
	355	_d_sec_next_free = nullptr;
	356	_d_word_ht_keys = nullptr;
	357	_d_word_ht_values = nullptr;
	358	_d_pair_ht_keys = nullptr;
	359	_d_pair_ht_values = nullptr;
	360	_d_sec_ht_keys = nullptr;
	361	_d_sec_ht_values = nullptr;
	362	_staging_key = nullptr;
	363	_staging_result = nullptr;
	364	_staging_wa = nullptr;
	365	_staging_wb = nullptr;
	366
	367	_device_id = -1;
	368	}
	369
	370	std::string CudaBackend::device_info()
	371	{
	372	return "CUDA / " + _device_name;
	373	}
	374
	375	// ==============================================================
	376	// Pool allocation
	377
	378	void CudaBackend::alloc_pools()
	379	{
	380	// -- WordPool --
	381	check_error(cudaMalloc(&_d_word_name_hash,
	382	sizeof(uint64_t) * GPU_WORD_CAPACITY), "alloc word_name_hash");
	383	check_error(cudaMalloc(&_d_word_count,
	384	sizeof(double) * GPU_WORD_CAPACITY), "alloc word_count");
	385	check_error(cudaMalloc(&_d_word_mi_marginal,
	386	sizeof(double) * GPU_WORD_CAPACITY), "alloc word_mi_marginal");
	387	check_error(cudaMalloc(&_d_word_class_id,
	388	sizeof(uint32_t) * GPU_WORD_CAPACITY), "alloc word_class_id");
	389	check_error(cudaMalloc(&_d_word_next_free,
	390	sizeof(uint32_t)), "alloc word_next_free");
	391
	392	// -- PairPool --
	393	check_error(cudaMalloc(&_d_pair_word_a,
	394	sizeof(uint32_t) * GPU_PAIR_CAPACITY), "alloc pair_word_a");
	395	check_error(cudaMalloc(&_d_pair_word_b,
	396	sizeof(uint32_t) * GPU_PAIR_CAPACITY), "alloc pair_word_b");
	397	check_error(cudaMalloc(&_d_pair_count,
	398	sizeof(double) * GPU_PAIR_CAPACITY), "alloc pair_count");
	399	check_error(cudaMalloc(&_d_pair_mi,
	400	sizeof(double) * GPU_PAIR_CAPACITY), "alloc pair_mi");
	401	check_error(cudaMalloc(&_d_pair_flags,
	402	sizeof(uint32_t) * GPU_PAIR_CAPACITY), "alloc pair_flags");
	403	check_error(cudaMalloc(&_d_pair_next_free,
	404	sizeof(uint32_t)), "alloc pair_next_free");
	405
	406	// -- SectionPool --
	407	check_error(cudaMalloc(&_d_sec_word,
	408	sizeof(uint32_t) * GPU_SECTION_CAPACITY), "alloc sec_word");
	409	check_error(cudaMalloc(&_d_sec_djh,
	410	sizeof(uint64_t) * GPU_SECTION_CAPACITY), "alloc sec_djh");
	411	check_error(cudaMalloc(&_d_sec_count,
	412	sizeof(double) * GPU_SECTION_CAPACITY), "alloc sec_count");
	413	check_error(cudaMalloc(&_d_sec_next_free,
	414	sizeof(uint32_t)), "alloc sec_next_free");
	415
	416	// -- Hash tables --
	417	check_error(cudaMalloc(&_d_word_ht_keys,
	418	sizeof(uint64_t) * GPU_WORD_HT_CAPACITY), "alloc word_ht_keys");
	419	check_error(cudaMalloc(&_d_word_ht_values,
	420	sizeof(uint32_t) * GPU_WORD_HT_CAPACITY), "alloc word_ht_values");
	421	check_error(cudaMalloc(&_d_pair_ht_keys,
	422	sizeof(uint64_t) * GPU_PAIR_HT_CAPACITY), "alloc pair_ht_keys");
	423	check_error(cudaMalloc(&_d_pair_ht_values,
	424	sizeof(uint32_t) * GPU_PAIR_HT_CAPACITY), "alloc pair_ht_values");
	425	check_error(cudaMalloc(&_d_sec_ht_keys,
	426	sizeof(uint64_t) * GPU_SECTION_HT_CAPACITY), "alloc sec_ht_keys");
	427	check_error(cudaMalloc(&_d_sec_ht_values,
	428	sizeof(uint32_t) * GPU_SECTION_HT_CAPACITY), "alloc sec_ht_values");
	429
	430	// -- Managed staging buffers (CPU+GPU accessible) --
	431	check_error(cudaMallocManaged(&_staging_key, sizeof(uint64_t)),
	432	"alloc staging_key");
	433	check_error(cudaMallocManaged(&_staging_result, sizeof(uint32_t)),
	434	"alloc staging_result");
	435	check_error(cudaMallocManaged(&_staging_wa, sizeof(uint32_t)),
	436	"alloc staging_wa");
	437	check_error(cudaMallocManaged(&_staging_wb, sizeof(uint32_t)),
	438	"alloc staging_wb");
	439	}
	440
	441	void CudaBackend::init_pools()
	442	{
	443	// Zero bump allocators
	444	uint32_t zero = 0;
	445	cudaMemcpy(_d_word_next_free, &zero, sizeof(uint32_t), cudaMemcpyHostToDevice);
	446	cudaMemcpy(_d_pair_next_free, &zero, sizeof(uint32_t), cudaMemcpyHostToDevice);
	447	cudaMemcpy(_d_sec_next_free, &zero, sizeof(uint32_t), cudaMemcpyHostToDevice);
	448
	449	// Initialize hash table keys to EMPTY (0xFF bytes)
	450	cudaMemset(_d_word_ht_keys, 0xFF, sizeof(uint64_t) * GPU_WORD_HT_CAPACITY);
	451	cudaMemset(_d_word_ht_values, 0xFF, sizeof(uint32_t) * GPU_WORD_HT_CAPACITY);
	452	cudaMemset(_d_pair_ht_keys, 0xFF, sizeof(uint64_t) * GPU_PAIR_HT_CAPACITY);
	453	cudaMemset(_d_pair_ht_values, 0xFF, sizeof(uint32_t) * GPU_PAIR_HT_CAPACITY);
	454	cudaMemset(_d_sec_ht_keys, 0xFF, sizeof(uint64_t) * GPU_SECTION_HT_CAPACITY);
	455	cudaMemset(_d_sec_ht_values, 0xFF, sizeof(uint32_t) * GPU_SECTION_HT_CAPACITY);
	456
	457	cudaDeviceSynchronize();
	458
	459	logger().info("CudaBackend: GPU pools initialized\n");
	460	}
	461
	462	// ==============================================================
	463	// Word pool operations
	464
	465	uint32_t CudaBackend::word_find_or_create(uint64_t name_hash)
	466	{
	467	*_staging_key = name_hash;
	468	*_staging_result = GPU_NOT_FOUND;
	469
	470	cuda_word_find_or_create<<<1, 1>>>(
	471	_d_word_ht_keys, _d_word_ht_values,
	472	_d_word_name_hash, _d_word_count, _d_word_class_id,
	473	_d_word_next_free,
	474	_staging_key, _staging_result, 1);
	475	cudaDeviceSynchronize();
	476
	477	return *_staging_result;
	478	}
	479
	480	uint32_t CudaBackend::word_lookup(uint64_t name_hash)
	481	{
	482	*_staging_key = name_hash;
	483	*_staging_result = GPU_NOT_FOUND;
	484
	485	cuda_ht_lookup<<<1, 1>>>(
	486	_d_word_ht_keys, _d_word_ht_values,
	487	(uint64_t)GPU_WORD_HT_CAPACITY,
	488	_staging_key, _staging_result, 1);
	489	cudaDeviceSynchronize();
	490
	491	return *_staging_result;
	492	}
	493
	494	void CudaBackend::word_write_count(uint32_t idx, double count)
	495	{
	496	cudaMemcpy(_d_word_count + idx, &count,
	497	sizeof(double), cudaMemcpyHostToDevice);
	498	}
	499
	500	void CudaBackend::word_write_marginal(uint32_t idx, double marginal)
	501	{
	502	cudaMemcpy(_d_word_mi_marginal + idx, &marginal,
	503	sizeof(double), cudaMemcpyHostToDevice);
	504	}
	505
	506	void CudaBackend::word_read_values(uint32_t idx,
	507	double& count, double& marginal)
	508	{
	509	cudaMemcpy(&count, _d_word_count + idx,
	510	sizeof(double), cudaMemcpyDeviceToHost);
	511	cudaMemcpy(&marginal, _d_word_mi_marginal + idx,
	512	sizeof(double), cudaMemcpyDeviceToHost);
	513	}
	514
	515	void CudaBackend::word_delete(uint64_t name_hash)
	516	{
	517	*_staging_key = name_hash;
	518
	519	cuda_ht_delete<<<1, 1>>>(
	520	_d_word_ht_keys, _d_word_ht_values,
	521	(uint64_t)GPU_WORD_HT_CAPACITY,
	522	_staging_key, 1);
	523	cudaDeviceSynchronize();
	524	}
	525
	526	uint32_t CudaBackend::word_pool_count()
	527	{
	528	uint32_t cnt = 0;
	529	cudaMemcpy(&cnt, _d_word_next_free,
	530	sizeof(uint32_t), cudaMemcpyDeviceToHost);
	531	return cnt;
	532	}
	533
	534	void CudaBackend::word_read_bulk(uint32_t n, uint64_t* hashes,
	535	double* counts, double* marginals)
	536	{
	537	cudaMemcpy(hashes, _d_word_name_hash,
	538	sizeof(uint64_t) * n, cudaMemcpyDeviceToHost);
	539	cudaMemcpy(counts, _d_word_count,
	540	sizeof(double) * n, cudaMemcpyDeviceToHost);
	541	cudaMemcpy(marginals, _d_word_mi_marginal,
	542	sizeof(double) * n, cudaMemcpyDeviceToHost);
	543	}
	544
	545	// ==============================================================
	546	// Pair pool operations
	547
	548	uint32_t CudaBackend::pair_find_or_create(uint32_t word_a, uint32_t word_b)
	549	{
	550	*_staging_wa = word_a;
	551	*_staging_wb = word_b;
	552	*_staging_result = GPU_NOT_FOUND;
	553
	554	cuda_pair_find_or_create<<<1, 1>>>(
	555	_d_pair_ht_keys, _d_pair_ht_values,
	556	_d_pair_word_a, _d_pair_word_b,
	557	_d_pair_count, _d_pair_mi, _d_pair_flags,
	558	_d_pair_next_free,
	559	_staging_wa, _staging_wb, _staging_result, 1);
	560	cudaDeviceSynchronize();
	561
	562	return *_staging_result;
	563	}
	564
	565	uint32_t CudaBackend::pair_lookup(uint32_t word_a, uint32_t word_b)
	566	{
	567	uint32_t lo = std::min(word_a, word_b);
	568	uint32_t hi = std::max(word_a, word_b);
	569	uint64_t key = ((uint64_t)lo << 32) \| (uint64_t)hi;
	570
	571	*_staging_key = key;
	572	*_staging_result = GPU_NOT_FOUND;
	573
	574	cuda_ht_lookup<<<1, 1>>>(
	575	_d_pair_ht_keys, _d_pair_ht_values,
	576	(uint64_t)GPU_PAIR_HT_CAPACITY,
	577	_staging_key, _staging_result, 1);
	578	cudaDeviceSynchronize();
	579
	580	return *_staging_result;
	581	}
	582
	583	void CudaBackend::pair_write_count(uint32_t idx, double count)
	584	{
	585	cudaMemcpy(_d_pair_count + idx, &count,
	586	sizeof(double), cudaMemcpyHostToDevice);
	587	}
	588
	589	void CudaBackend::pair_write_mi(uint32_t idx, double mi)
	590	{
	591	cudaMemcpy(_d_pair_mi + idx, &mi,
	592	sizeof(double), cudaMemcpyHostToDevice);
	593	}
	594
	595	void CudaBackend::pair_read_values(uint32_t idx,
	596	double& count, double& mi)
	597	{
	598	cudaMemcpy(&count, _d_pair_count + idx,
	599	sizeof(double), cudaMemcpyDeviceToHost);
	600	cudaMemcpy(&mi, _d_pair_mi + idx,
	601	sizeof(double), cudaMemcpyDeviceToHost);
	602	}
	603
	604	uint32_t CudaBackend::pair_pool_count()
	605	{
	606	uint32_t cnt = 0;
	607	cudaMemcpy(&cnt, _d_pair_next_free,
	608	sizeof(uint32_t), cudaMemcpyDeviceToHost);
	609	return cnt;
	610	}
	611
	612	// ==============================================================
	613	// Section pool operations
	614
	615	uint32_t CudaBackend::section_pool_count()
	616	{
	617	uint32_t cnt = 0;
	618	cudaMemcpy(&cnt, _d_sec_next_free,
	619	sizeof(uint32_t), cudaMemcpyDeviceToHost);
	620	return cnt;
	621	}
	622
	623	// ==============================================================
	624
	625	void CudaBackend::barrier()
	626	{
	627	cudaDeviceSynchronize();
	628	}
	629
	630	// ==============================================================
	631	// Factory function
	632
	633	namespace opencog {
	634
	635	GpuBackend* create_cuda_backend()
	636	{
	637	return new CudaBackend();
	638	}
	639
	640	} // namespace opencog
	641
	642	// ==============================================================

opencog/persist/gpu/CudaBackend.h +109 -0

…	…	@@ -0,0 +1,109 @@
	1	/*
	2	* opencog/persist/gpu/CudaBackend.h
	3	*
	4	* CUDA implementation of the GpuBackend interface.
	5	* Primary backend for NVIDIA GPUs (more performant than OpenCL).
	6	*
	7	* NOTE: This header is only included by CudaBackend.cu, which is
	8	* compiled by nvcc. It must NOT be included by regular C++ files.
	9	*
	10	* Copyright (C) 2025 OpenCog Foundation
	11	*
	12	* SPDX-License-Identifier: AGPL-3.0-or-later
	13	*/
	14
	15	#ifndef _OPENCOG_CUDA_BACKEND_H
	16	#define _OPENCOG_CUDA_BACKEND_H
	17
	18	#include "GpuBackend.h"
	19	#include "gpu-pool-defs.h"
	20
	21	#include <cuda_runtime.h>
	22
	23	namespace opencog
	24	{
	25
	26	class CudaBackend : public GpuBackend
	27	{
	28	private:
	29	int _device_id;
	30	std::string _device_name;
	31
	32	// Device memory -- WordPool SoA
	33	uint64_t* _d_word_name_hash;
	34	double* _d_word_count;
	35	double* _d_word_mi_marginal;
	36	uint32_t* _d_word_class_id;
	37	uint32_t* _d_word_next_free;
	38
	39	// Device memory -- PairPool SoA
	40	uint32_t* _d_pair_word_a;
	41	uint32_t* _d_pair_word_b;
	42	double* _d_pair_count;
	43	double* _d_pair_mi;
	44	uint32_t* _d_pair_flags;
	45	uint32_t* _d_pair_next_free;
	46
	47	// Device memory -- SectionPool SoA
	48	uint32_t* _d_sec_word;
	49	uint64_t* _d_sec_djh;
	50	double* _d_sec_count;
	51	uint32_t* _d_sec_next_free;
	52
	53	// Device memory -- Hash tables
	54	uint64_t* _d_word_ht_keys;
	55	uint32_t* _d_word_ht_values;
	56	uint64_t* _d_pair_ht_keys;
	57	uint32_t* _d_pair_ht_values;
	58	uint64_t* _d_sec_ht_keys;
	59	uint32_t* _d_sec_ht_values;
	60
	61	// Managed staging buffers (CPU+GPU accessible, no cudaMemcpy)
	62	uint64_t* _staging_key;
	63	uint32_t* _staging_result;
	64	uint32_t* _staging_wa;
	65	uint32_t* _staging_wb;
	66
	67	void check_error(cudaError_t err, const char* context);
	68
	69	public:
	70	CudaBackend();
	71	~CudaBackend() override;
	72
	73	void init(const std::string& platform_hint,
	74	const std::string& device_hint) override;
	75	void shutdown() override;
	76	std::string device_info() override;
	77	std::string backend_name() override { return "CUDA"; }
	78
	79	void alloc_pools() override;
	80	void init_pools() override;
	81
	82	uint32_t word_find_or_create(uint64_t name_hash) override;
	83	uint32_t word_lookup(uint64_t name_hash) override;
	84	void word_write_count(uint32_t idx, double count) override;
	85	void word_write_marginal(uint32_t idx, double marginal) override;
	86	void word_read_values(uint32_t idx,
	87	double& count, double& marginal) override;
	88	void word_delete(uint64_t name_hash) override;
	89	uint32_t word_pool_count() override;
	90	void word_read_bulk(uint32_t n, uint64_t* hashes,
	91	double* counts, double* marginals) override;
	92
	93	uint32_t pair_find_or_create(uint32_t word_a,
	94	uint32_t word_b) override;
	95	uint32_t pair_lookup(uint32_t word_a, uint32_t word_b) override;
	96	void pair_write_count(uint32_t idx, double count) override;
	97	void pair_write_mi(uint32_t idx, double mi) override;
	98	void pair_read_values(uint32_t idx,
	99	double& count, double& mi) override;
	100	uint32_t pair_pool_count() override;
	101
	102	uint32_t section_pool_count() override;
	103
	104	void barrier() override;
	105	};
	106
	107	} // namespace opencog
	108
	109	#endif // _OPENCOG_CUDA_BACKEND_H

opencog/persist/gpu/GpuBackend.h +83 -0

…	…	@@ -0,0 +1,83 @@
	1	/*
	2	* opencog/persist/gpu/GpuBackend.h
	3	*
	4	* Abstract interface for GPU compute backends (CUDA, OpenCL).
	5	* GpuStorageNode delegates all GPU operations to a concrete
	6	* GpuBackend implementation, selected at open() time based on
	7	* hardware availability and URI hints.
	8	*
	9	* Priority: CUDA first (more performant on NVIDIA), OpenCL
	10	* as fallback (works on AMD/Intel/NVIDIA).
	11	*
	12	* Copyright (C) 2025 OpenCog Foundation
	13	*
	14	* SPDX-License-Identifier: AGPL-3.0-or-later
	15	*/
	16
	17	#ifndef _OPENCOG_GPU_BACKEND_H
	18	#define _OPENCOG_GPU_BACKEND_H
	19
	20	#include <cstdint>
	21	#include <string>
	22
	23	namespace opencog
	24	{
	25
	26	class GpuBackend
	27	{
	28	public:
	29	virtual ~GpuBackend() = default;
	30
	31	// -- Lifecycle --
	32	virtual void init(const std::string& platform_hint,
	33	const std::string& device_hint) = 0;
	34	virtual void shutdown() = 0;
	35	virtual std::string device_info() = 0;
	36	virtual std::string backend_name() = 0; // "CUDA" or "OpenCL"
	37
	38	// -- Pool allocation and initialization --
	39	virtual void alloc_pools() = 0;
	40	virtual void init_pools() = 0;
	41
	42	// -- Word pool operations --
	43	virtual uint32_t word_find_or_create(uint64_t name_hash) = 0;
	44	virtual uint32_t word_lookup(uint64_t name_hash) = 0;
	45	virtual void word_write_count(uint32_t idx, double count) = 0;
	46	virtual void word_write_marginal(uint32_t idx, double marginal) = 0;
	47	virtual void word_read_values(uint32_t idx,
	48	double& count, double& marginal) = 0;
	49	virtual void word_delete(uint64_t name_hash) = 0;
	50	virtual uint32_t word_pool_count() = 0;
	51	virtual void word_read_bulk(uint32_t n, uint64_t* hashes,
	52	double* counts, double* marginals) = 0;
	53
	54	// -- Pair pool operations --
	55	virtual uint32_t pair_find_or_create(uint32_t word_a,
	56	uint32_t word_b) = 0;
	57	virtual uint32_t pair_lookup(uint32_t word_a, uint32_t word_b) = 0;
	58	virtual void pair_write_count(uint32_t idx, double count) = 0;
	59	virtual void pair_write_mi(uint32_t idx, double mi) = 0;
	60	virtual void pair_read_values(uint32_t idx,
	61	double& count, double& mi) = 0;
	62	virtual uint32_t pair_pool_count() = 0;
	63
	64	// -- Section pool operations --
	65	virtual uint32_t section_pool_count() = 0;
	66
	67	// -- Synchronization --
	68	virtual void barrier() = 0;
	69	};
	70
	71	// Backend factory functions (defined in respective .cc/.cu files).
	72	// Availability depends on compile-time HAVE_CUDA / HAVE_OPENCL.
	73	#ifdef HAVE_CUDA
	74	GpuBackend* create_cuda_backend();
	75	#endif
	76
	77	#ifdef HAVE_OPENCL
	78	GpuBackend* create_opencl_backend();
	79	#endif
	80
	81	} // namespace opencog
	82
	83	#endif // _OPENCOG_GPU_BACKEND_H

opencog/persist/gpu/GpuStorageNode.cc +646 -0

…	…	@@ -0,0 +1,646 @@
	1	/*
	2	* opencog/persist/gpu/GpuStorageNode.cc
	3	*
	4	* StorageNode backed by GPU SoA pools.
	5	* Delegates all GPU operations to a GpuBackend (CUDA or OpenCL).
	6	*
	7	* Phase 1: Store/Fetch round-trip for atoms and values.
	8	*
	9	* Copyright (C) 2025 OpenCog Foundation
	10	*
	11	* SPDX-License-Identifier: AGPL-3.0-or-later
	12	*/
	13
	14	#include <sstream>
	15	#include <functional>
	16	#include <cstring>
	17
	18	#include <opencog/util/exceptions.h>
	19	#include <opencog/util/Logger.h>
	20	#include <opencog/atomspace/AtomSpace.h>
	21	#include <opencog/atoms/base/Node.h>
	22	#include <opencog/atoms/base/Link.h>
	23	#include <opencog/atoms/value/FloatValue.h>
	24	#include <opencog/atoms/atom_types/atom_types.h>
	25	#include <opencog/persist/gpu-types/atom_types.h>
	26
	27	#include "GpuStorageNode.h"
	28
	29	using namespace opencog;
	30
	31	// ==============================================================
	32	// Constructors / Destructor
	33
	34	GpuStorageNode::GpuStorageNode(Type t, const std::string&& uri) :
	35	StorageNode(t, std::move(uri)),
	36	_connected(false),
	37	_num_stores(0),
	38	_num_fetches(0)
	39	{
	40	_uri = get_name();
	41	parse_uri();
	42	}
	43
	44	GpuStorageNode::GpuStorageNode(const std::string&& uri) :
	45	StorageNode(GPU_STORAGE_NODE, std::move(uri)),
	46	_connected(false),
	47	_num_stores(0),
	48	_num_fetches(0)
	49	{
	50	_uri = get_name();
	51	parse_uri();
	52	}
	53
	54	GpuStorageNode::~GpuStorageNode()
	55	{
	56	if (_connected) close();
	57	}
	58
	59	// ==============================================================
	60	// URI parsing: "gpu://[backend:]platform:device"
	61	//
	62	// Extended URI format:
	63	// "gpu://:" → auto backend, any device
	64	// "gpu://NVIDIA:RTX" → auto backend, NVIDIA RTX
	65	// "gpu://cuda::" → force CUDA, any device
	66	// "gpu://opencl:Intel:" → force OpenCL, Intel device
	67	// "gpu://cuda:NVIDIA:RTX" → CUDA, NVIDIA RTX
	68
	69	void GpuStorageNode::parse_uri(void)
	70	{
	71	const std::string& url = _uri;
	72	if (0 != url.compare(0, 6, "gpu://"))
	73	throw RuntimeException(TRACE_INFO,
	74	"Unsupported URL \"%s\"\n"
	75	"\tExpecting 'gpu://[backend:]platform:device'",
	76	url.c_str());
	77
	78	std::string rest = url.substr(6);
	79
	80	// Check for explicit backend hint: "cuda:" or "opencl:"
	81	_backend_hint = "";
	82	if (rest.compare(0, 5, "cuda:") == 0)
	83	{
	84	_backend_hint = "cuda";
	85	rest = rest.substr(5);
	86	}
	87	else if (rest.compare(0, 7, "opencl:") == 0)
	88	{
	89	_backend_hint = "opencl";
	90	rest = rest.substr(7);
	91	}
	92
	93	// Parse platform:device from remainder
	94	size_t colon = rest.find(':');
	95	if (std::string::npos == colon)
	96	{
	97	_splat = rest;
	98	_sdev = "";
	99	}
	100	else
	101	{
	102	_splat = rest.substr(0, colon);
	103	_sdev = rest.substr(colon + 1);
	104	}
	105	}
	106
	107	// ==============================================================
	108	// Connection management
	109
	110	void GpuStorageNode::open(void)
	111	{
	112	if (_connected)
	113	throw RuntimeException(TRACE_INFO,
	114	"GpuStorageNode: already open! %s\n", _uri.c_str());
	115
	116	// Select backend: CUDA preferred, OpenCL fallback
	117	bool want_cuda = (_backend_hint == "cuda" or _backend_hint.empty());
	118	bool want_opencl = (_backend_hint == "opencl" or _backend_hint.empty());
	119
	120	GpuBackend* be = nullptr;
	121
	122	#ifdef HAVE_CUDA
	123	if (want_cuda and nullptr == be)
	124	{
	125	try
	126	{
	127	be = create_cuda_backend();
	128	be->init(_splat, _sdev);
	129	}
	130	catch (const RuntimeException& e)
	131	{
	132	delete be;
	133	be = nullptr;
	134	if (_backend_hint == "cuda")
	135	throw; // User explicitly asked for CUDA
	136	logger().info("GpuStorageNode: CUDA unavailable, "
	137	"trying OpenCL...\n");
	138	}
	139	}
	140	#endif
	141
	142	#ifdef HAVE_OPENCL
	143	if (want_opencl and nullptr == be)
	144	{
	145	try
	146	{
	147	be = create_opencl_backend();
	148	be->init(_splat, _sdev);
	149	}
	150	catch (const RuntimeException& e)
	151	{
	152	delete be;
	153	be = nullptr;
	154	if (_backend_hint == "opencl")
	155	throw; // User explicitly asked for OpenCL
	156	}
	157	}
	158	#endif
	159
	160	if (nullptr == be)
	161	throw RuntimeException(TRACE_INFO,
	162	"GpuStorageNode: no GPU backend available for '%s'\n"
	163	"\tCompiled with:"
	164	#ifdef HAVE_CUDA
	165	" CUDA"
	166	#endif
	167	#ifdef HAVE_OPENCL
	168	" OpenCL"
	169	#endif
	170	"\n",
	171	_uri.c_str());
	172
	173	_backend.reset(be);
	174	_backend->alloc_pools();
	175	_backend->init_pools();
	176
	177	_connected = true;
	178	_num_stores = 0;
	179	_num_fetches = 0;
	180
	181	logger().info("GpuStorageNode: opened %s (backend: %s, device: %s)\n",
	182	_uri.c_str(),
	183	_backend->backend_name().c_str(),
	184	_backend->device_info().c_str());
	185	}
	186
	187	void GpuStorageNode::close(void)
	188	{
	189	if (not _connected) return;
	190
	191	_backend->barrier();
	192	_backend->shutdown();
	193	_backend.reset();
	194
	195	// Clear CPU-side maps
	196	{
	197	std::lock_guard<std::mutex> lk(_name_mtx);
	198	_name_to_hash.clear();
	199	_hash_to_name.clear();
	200	}
	201	{
	202	std::lock_guard<std::mutex> lk(_atom_mtx);
	203	_atom_map.clear();
	204	}
	205
	206	_connected = false;
	207	logger().info("GpuStorageNode: closed %s\n", _uri.c_str());
	208	}
	209
	210	bool GpuStorageNode::connected(void)
	211	{
	212	return _connected;
	213	}
	214
	215	void GpuStorageNode::erase(void)
	216	{
	217	if (not _connected) return;
	218
	219	_backend->init_pools();
	220
	221	{
	222	std::lock_guard<std::mutex> lk(_name_mtx);
	223	_name_to_hash.clear();
	224	_hash_to_name.clear();
	225	}
	226	{
	227	std::lock_guard<std::mutex> lk(_atom_mtx);
	228	_atom_map.clear();
	229	}
	230	}
	231
	232	// ==============================================================
	233	// Name hashing
	234
	235	uint64_t GpuStorageNode::name_hash(const std::string& name)
	236	{
	237	std::lock_guard<std::mutex> lk(_name_mtx);
	238
	239	auto it = _name_to_hash.find(name);
	240	if (it != _name_to_hash.end())
	241	return it->second;
	242
	243	// splitmix64 on std::hash
	244	std::hash<std::string> hasher;
	245	uint64_t h = hasher(name);
	246
	247	if (h == GPU_HT_EMPTY_KEY) h = 0;
	248
	249	_name_to_hash[name] = h;
	250	_hash_to_name[h] = name;
	251	return h;
	252	}
	253
	254	// ==============================================================
	255	// GPU pool operations (delegated to backend)
	256
	257	uint32_t GpuStorageNode::store_word(const std::string& name)
	258	{
	259	uint64_t nhash = name_hash(name);
	260	uint32_t idx = _backend->word_find_or_create(nhash);
	261
	262	if (idx == GPU_NOT_FOUND)
	263	throw RuntimeException(TRACE_INFO,
	264	"GpuStorageNode: word pool full, cannot store '%s'\n",
	265	name.c_str());
	266
	267	return idx;
	268	}
	269
	270	uint32_t GpuStorageNode::lookup_word(const std::string& name)
	271	{
	272	uint64_t nhash = name_hash(name);
	273	return _backend->word_lookup(nhash);
	274	}
	275
	276	uint32_t GpuStorageNode::store_pair(uint32_t word_a, uint32_t word_b)
	277	{
	278	uint32_t idx = _backend->pair_find_or_create(word_a, word_b);
	279
	280	if (idx == GPU_NOT_FOUND)
	281	throw RuntimeException(TRACE_INFO,
	282	"GpuStorageNode: pair pool full\n");
	283
	284	return idx;
	285	}
	286
	287	uint32_t GpuStorageNode::lookup_pair(uint32_t word_a, uint32_t word_b)
	288	{
	289	return _backend->pair_lookup(word_a, word_b);
	290	}
	291
	292	// ==============================================================
	293	// Value storage helpers
	294
	295	void GpuStorageNode::store_node_values(const Handle& h, uint32_t pool_idx)
	296	{
	297	ValuePtr tv = h->getValue(truth_key());
	298	if (nullptr == tv) return;
	299
	300	FloatValuePtr fv = FloatValueCast(tv);
	301	if (nullptr == fv) return;
	302
	303	const std::vector<double>& vals = fv->value();
	304	if (vals.size() > 0)
	305	_backend->word_write_count(pool_idx, vals[0]);
	306	if (vals.size() > 1)
	307	_backend->word_write_marginal(pool_idx, vals[1]);
	308	}
	309
	310	void GpuStorageNode::load_node_values(const Handle& h, uint32_t pool_idx)
	311	{
	312	double count = 0.0, marginal = 0.0;
	313	_backend->word_read_values(pool_idx, count, marginal);
	314
	315	if (count != 0.0 or marginal != 0.0)
	316	{
	317	std::vector<double> vals = {count, marginal};
	318	h->setValue(truth_key(), createFloatValue(vals));
	319	}
	320	}
	321
	322	void GpuStorageNode::store_link_values(const Handle& h, uint32_t pool_idx)
	323	{
	324	ValuePtr tv = h->getValue(truth_key());
	325	if (nullptr == tv) return;
	326
	327	FloatValuePtr fv = FloatValueCast(tv);
	328	if (nullptr == fv) return;
	329
	330	const std::vector<double>& vals = fv->value();
	331	if (vals.size() > 0)
	332	_backend->pair_write_count(pool_idx, vals[0]);
	333	if (vals.size() > 1)
	334	_backend->pair_write_mi(pool_idx, vals[1]);
	335	}
	336
	337	void GpuStorageNode::load_link_values(const Handle& h, uint32_t pool_idx)
	338	{
	339	double count = 0.0, mi = 0.0;
	340	_backend->pair_read_values(pool_idx, count, mi);
	341
	342	if (count != 0.0 or mi != 0.0)
	343	{
	344	std::vector<double> vals = {count, mi};
	345	h->setValue(truth_key(), createFloatValue(vals));
	346	}
	347	}
	348
	349	// ==============================================================
	350	// StorageNode API: storeAtom
	351
	352	void GpuStorageNode::storeAtom(const Handle& h, bool synchronous)
	353	{
	354	if (not _connected)
	355	throw RuntimeException(TRACE_INFO,
	356	"GpuStorageNode: not connected!\n");
	357
	358	NodePtr np = NodeCast(h);
	359	if (np)
	360	{
	361	uint32_t idx = store_word(np->get_name());
	362	store_node_values(h, idx);
	363
	364	std::lock_guard<std::mutex> lk(_atom_mtx);
	365	_atom_map[h] = {idx, 0};
	366	_num_stores++;
	367
	368	if (synchronous) _backend->barrier();
	369	return;
	370	}
	371
	372	LinkPtr lp = LinkCast(h);
	373	if (lp)
	374	{
	375	for (const Handle& oh : lp->getOutgoingSet())
	376	storeAtom(oh, false);
	377
	378	if (lp->get_arity() == 2)
	379	{
	380	const Handle& ha = lp->getOutgoingAtom(0);
	381	const Handle& hb = lp->getOutgoingAtom(1);
	382
	383	NodePtr na = NodeCast(ha);
	384	NodePtr nb = NodeCast(hb);
	385	if (na and nb)
	386	{
	387	uint32_t ia = store_word(na->get_name());
	388	uint32_t ib = store_word(nb->get_name());
	389	uint32_t pidx = store_pair(ia, ib);
	390	store_link_values(h, pidx);
	391
	392	std::lock_guard<std::mutex> lk(_atom_mtx);
	393	_atom_map[h] = {pidx, 1};
	394	}
	395	}
	396
	397	_num_stores++;
	398	if (synchronous) _backend->barrier();
	399	return;
	400	}
	401	}
	402
	403	// ==============================================================
	404	// StorageNode API: getAtom
	405
	406	void GpuStorageNode::getAtom(const Handle& h)
	407	{
	408	if (not _connected)
	409	throw RuntimeException(TRACE_INFO,
	410	"GpuStorageNode: not connected!\n");
	411
	412	NodePtr np = NodeCast(h);
	413	if (np)
	414	{
	415	uint32_t idx = lookup_word(np->get_name());
	416	if (idx == GPU_NOT_FOUND) return;
	417
	418	load_node_values(h, idx);
	419	_num_fetches++;
	420	return;
	421	}
	422
	423	LinkPtr lp = LinkCast(h);
	424	if (lp and lp->get_arity() == 2)
	425	{
	426	const Handle& ha = lp->getOutgoingAtom(0);
	427	const Handle& hb = lp->getOutgoingAtom(1);
	428
	429	NodePtr na = NodeCast(ha);
	430	NodePtr nb = NodeCast(hb);
	431	if (na and nb)
	432	{
	433	uint32_t ia = lookup_word(na->get_name());
	434	uint32_t ib = lookup_word(nb->get_name());
	435	if (ia == GPU_NOT_FOUND or ib == GPU_NOT_FOUND) return;
	436
	437	uint32_t pidx = lookup_pair(ia, ib);
	438	if (pidx == GPU_NOT_FOUND) return;
	439
	440	load_link_values(h, pidx);
	441	_num_fetches++;
	442	}
	443	}
	444	}
	445
	446	// ==============================================================
	447	// StorageNode API: storeValue / loadValue
	448
	449	void GpuStorageNode::storeValue(const Handle& atom, const Handle& key)
	450	{
	451	if (not _connected)
	452	throw RuntimeException(TRACE_INFO,
	453	"GpuStorageNode: not connected!\n");
	454
	455	if (key != truth_key()) return;
	456
	457	NodePtr np = NodeCast(atom);
	458	if (np)
	459	{
	460	uint32_t idx = store_word(np->get_name());
	461	store_node_values(atom, idx);
	462	return;
	463	}
	464
	465	LinkPtr lp = LinkCast(atom);
	466	if (lp and lp->get_arity() == 2)
	467	{
	468	NodePtr na = NodeCast(lp->getOutgoingAtom(0));
	469	NodePtr nb = NodeCast(lp->getOutgoingAtom(1));
	470	if (na and nb)
	471	{
	472	uint32_t ia = store_word(na->get_name());
	473	uint32_t ib = store_word(nb->get_name());
	474	uint32_t pidx = store_pair(ia, ib);
	475	store_link_values(atom, pidx);
	476	}
	477	}
	478	}
	479
	480	void GpuStorageNode::loadValue(const Handle& atom, const Handle& key)
	481	{
	482	if (not _connected)
	483	throw RuntimeException(TRACE_INFO,
	484	"GpuStorageNode: not connected!\n");
	485
	486	if (key != truth_key()) return;
	487
	488	NodePtr np = NodeCast(atom);
	489	if (np)
	490	{
	491	uint32_t idx = lookup_word(np->get_name());
	492	if (idx != GPU_NOT_FOUND)
	493	load_node_values(atom, idx);
	494	return;
	495	}
	496
	497	LinkPtr lp = LinkCast(atom);
	498	if (lp and lp->get_arity() == 2)
	499	{
	500	NodePtr na = NodeCast(lp->getOutgoingAtom(0));
	501	NodePtr nb = NodeCast(lp->getOutgoingAtom(1));
	502	if (na and nb)
	503	{
	504	uint32_t ia = lookup_word(na->get_name());
	505	uint32_t ib = lookup_word(nb->get_name());
	506	if (ia != GPU_NOT_FOUND and ib != GPU_NOT_FOUND)
	507	{
	508	uint32_t pidx = lookup_pair(ia, ib);
	509	if (pidx != GPU_NOT_FOUND)
	510	load_link_values(atom, pidx);
	511	}
	512	}
	513	}
	514	}
	515
	516	// ==============================================================
	517	// StorageNode API: removeAtom
	518
	519	void GpuStorageNode::removeAtom(AtomSpace*, const Handle& h, bool recursive)
	520	{
	521	if (not _connected)
	522	throw RuntimeException(TRACE_INFO,
	523	"GpuStorageNode: not connected!\n");
	524
	525	NodePtr np = NodeCast(h);
	526	if (np)
	527	{
	528	uint64_t nhash = name_hash(np->get_name());
	529	_backend->word_delete(nhash);
	530
	531	std::lock_guard<std::mutex> lk(_atom_mtx);
	532	_atom_map.erase(h);
	533	}
	534	}
	535
	536	// ==============================================================
	537	// StorageNode API: bulk operations
	538
	539	void GpuStorageNode::fetchIncomingSet(AtomSpace*, const Handle&)
	540	{
	541	}
	542
	543	void GpuStorageNode::fetchIncomingByType(AtomSpace*, const Handle&, Type)
	544	{
	545	}
	546
	547	void GpuStorageNode::loadType(AtomSpace*, Type)
	548	{
	549	}
	550
	551	void GpuStorageNode::loadAtomSpace(AtomSpace* as)
	552	{
	553	if (not _connected)
	554	throw RuntimeException(TRACE_INFO,
	555	"GpuStorageNode: not connected!\n");
	556
	557	uint32_t nwords = _backend->word_pool_count();
	558	if (0 == nwords) return;
	559
	560	std::vector<uint64_t> hashes(nwords);
	561	std::vector<double> counts(nwords);
	562	std::vector<double> marginals(nwords);
	563	_backend->word_read_bulk(nwords, hashes.data(),
	564	counts.data(), marginals.data());
	565
	566	std::lock_guard<std::mutex> lk(_name_mtx);
	567	for (uint32_t i = 0; i < nwords; i++)
	568	{
	569	auto it = _hash_to_name.find(hashes[i]);
	570	if (it == _hash_to_name.end()) continue;
	571
	572	std::string name = it->second;
	573	Handle h = as->add_node(SCHEMA_NODE, std::move(name));
	574	if (counts[i] != 0.0 or marginals[i] != 0.0)
	575	{
	576	std::vector<double> vals = {counts[i], marginals[i]};
	577	h->setValue(truth_key(), createFloatValue(vals));
	578	}
	579	}
	580	}
	581
	582	void GpuStorageNode::storeAtomSpace(const AtomSpace* as)
	583	{
	584	if (not _connected)
	585	throw RuntimeException(TRACE_INFO,
	586	"GpuStorageNode: not connected!\n");
	587
	588	HandleSeq all_atoms;
	589	as->get_handles_by_type(all_atoms, ATOM, true);
	590	for (const Handle& h : all_atoms)
	591	storeAtom(h, false);
	592
	593	_backend->barrier();
	594	}
	595
	596	// ==============================================================
	597	// Barrier
	598
	599	void GpuStorageNode::barrier(AtomSpace*)
	600	{
	601	if (_connected)
	602	_backend->barrier();
	603	}
	604
	605	// ==============================================================
	606	// Monitor / Stats
	607
	608	std::string GpuStorageNode::monitor(void)
	609	{
	610	std::ostringstream ss;
	611	ss << "GpuStorageNode: " << _uri << "\n";
	612	ss << " Connected: " << (_connected ? "yes" : "no") << "\n";
	613
	614	if (_connected)
	615	{
	616	ss << " Backend: " << _backend->backend_name() << "\n";
	617	ss << " Device: " << _backend->device_info() << "\n";
	618	ss << " Words: " << _backend->word_pool_count()
	619	<< " / " << GPU_WORD_CAPACITY << "\n";
	620	ss << " Pairs: " << _backend->pair_pool_count()
	621	<< " / " << GPU_PAIR_CAPACITY << "\n";
	622	ss << " Sections: " << _backend->section_pool_count()
	623	<< " / " << GPU_SECTION_CAPACITY << "\n";
	624	ss << " Stores: " << _num_stores.load() << "\n";
	625	ss << " Fetches: " << _num_fetches.load() << "\n";
	626	}
	627
	628	return ss.str();
	629	}
	630
	631	// ==============================================================
	632	// Phase 2: runQuery (stub)
	633
	634	void GpuStorageNode::runQuery(const Handle& query, const Handle& key,
	635	const Handle& metadata_key, bool fresh)
	636	{
	637	throw RuntimeException(TRACE_INFO,
	638	"GpuStorageNode::runQuery() not yet implemented (Phase 2)\n");
	639	}
	640
	641	// ==============================================================
	642	// Factory
	643
	644	DEFINE_NODE_FACTORY(GpuStorageNode, GPU_STORAGE_NODE);
	645
	646	// ==============================================================

opencog/persist/gpu/GpuStorageNode.h +142 -0

…	…	@@ -0,0 +1,142 @@
	1	/*
	2	* opencog/persist/gpu/GpuStorageNode.h
	3	*
	4	* StorageNode implementation backed by GPU SoA pools.
	5	* Provides the standard StorageNode API (store/fetch/query)
	6	* using GPU-resident word, pair, and section pools.
	7	* Same pattern as RocksStorageNode (disk) and CogStorageNode
	8	* (network), but targeting GPU memory.
	9	*
	10	* Supports CUDA (primary, NVIDIA) and OpenCL (fallback, AMD/Intel).
	11	* Backend is selected at open() time based on hardware availability
	12	* and URI hints.
	13	*
	14	* URI format: "gpu://[backend:]platform:device"
	15	* "gpu://:" first available device (CUDA preferred)
	16	* "gpu://NVIDIA:RTX" first NVIDIA device containing "RTX"
	17	* "gpu://cuda::" force CUDA backend
	18	* "gpu://opencl::" force OpenCL backend
	19	* "gpu://opencl:Intel:" OpenCL on first Intel device
	20	*
	21	* Copyright (C) 2025 OpenCog Foundation
	22	*
	23	* SPDX-License-Identifier: AGPL-3.0-or-later
	24	*/
	25
	26	#ifndef _OPENCOG_GPU_STORAGE_NODE_H
	27	#define _OPENCOG_GPU_STORAGE_NODE_H
	28
	29	#include <memory>
	30	#include <mutex>
	31	#include <unordered_map>
	32	#include <functional>
	33
	34	#include <opencog/persist/api/StorageNode.h>
	35
	36	#include "gpu-pool-defs.h"
	37	#include "GpuBackend.h"
	38
	39	namespace opencog
	40	{
	41
	42	/** \addtogroup grp_persist
	43	* @{
	44	*/
	45
	46	class GpuStorageNode : public StorageNode
	47	{
	48	private:
	49	std::string _uri;
	50	std::string _backend_hint; // "cuda", "opencl", or "" (auto)
	51	std::string _splat; // platform substring
	52	std::string _sdev; // device substring
	53
	54	// GPU backend (CUDA or OpenCL, selected at open time)
	55	std::unique_ptr<GpuBackend> _backend;
	56	bool _connected;
	57
	58	// CPU-side name->hash mapping (GPU stores hashes, CPU knows strings)
	59	std::mutex _name_mtx;
	60	std::unordered_map<std::string, uint64_t> _name_to_hash;
	61	std::unordered_map<uint64_t, std::string> _hash_to_name;
	62
	63	// Atom->GPU-index mapping for Values
	64	struct ValueSlot {
	65	uint32_t pool_index; // index in word/pair/section pool
	66	int pool_type; // 0=word, 1=pair, 2=section
	67	};
	68	std::mutex _atom_mtx;
	69	std::unordered_map<Handle, ValueSlot> _atom_map;
	70
	71	// Internal helpers
	72	void parse_uri(void);
	73
	74	uint64_t name_hash(const std::string& name);
	75	uint32_t store_word(const std::string& name);
	76	uint32_t store_pair(uint32_t word_a, uint32_t word_b);
	77	uint32_t lookup_word(const std::string& name);
	78	uint32_t lookup_pair(uint32_t word_a, uint32_t word_b);
	79
	80	void store_node_values(const Handle&, uint32_t pool_idx);
	81	void load_node_values(const Handle&, uint32_t pool_idx);
	82	void store_link_values(const Handle&, uint32_t pool_idx);
	83	void load_link_values(const Handle&, uint32_t pool_idx);
	84
	85	// Stats tracking
	86	std::atomic<size_t> _num_stores;
	87	std::atomic<size_t> _num_fetches;
	88
	89	public:
	90	GpuStorageNode(Type t, const std::string&& uri);
	91	GpuStorageNode(const std::string&& uri);
	92	GpuStorageNode(const GpuStorageNode&) = delete;
	93	GpuStorageNode& operator=(const GpuStorageNode&) = delete;
	94	virtual ~GpuStorageNode();
	95
	96	// -- StorageNode interface --
	97	void open(void) override;
	98	void close(void) override;
	99	bool connected(void) override;
	100
	101	void create(void) override { erase(); }
	102	void destroy(void) override { erase(); }
	103	void erase(void) override;
	104
	105	// -- BackingStore interface --
	106	void getAtom(const Handle&) override;
	107	void storeAtom(const Handle&, bool synchronous = false) override;
	108	void removeAtom(AtomSpace*, const Handle&, bool recursive) override;
	109
	110	void storeValue(const Handle& atom, const Handle& key) override;
	111	void loadValue(const Handle& atom, const Handle& key) override;
	112
	113	void fetchIncomingSet(AtomSpace*, const Handle&) override;
	114	void fetchIncomingByType(AtomSpace*, const Handle&, Type) override;
	115	void loadType(AtomSpace*, Type) override;
	116	void loadAtomSpace(AtomSpace*) override;
	117	void storeAtomSpace(const AtomSpace*) override;
	118	void barrier(AtomSpace* = nullptr) override;
	119
	120	std::string monitor(void) override;
	121
	122	// -- Phase 2: Query execution on GPU --
	123	void runQuery(const Handle&, const Handle&,
	124	const Handle& = Handle::UNDEFINED,
	125	bool = false) override;
	126
	127	// Atom factory
	128	void setAtomSpace(AtomSpace* as)
	129	{
	130	if (nullptr == as) close();
	131	Atom::setAtomSpace(as);
	132	}
	133	static Handle factory(const Handle&);
	134	};
	135
	136	NODE_PTR_DECL(GpuStorageNode)
	137	#define createGpuStorageNode CREATE_DECL(GpuStorageNode)
	138
	139	/** @}*/
	140	} // namespace opencog
	141
	142	#endif // _OPENCOG_GPU_STORAGE_NODE_H

opencog/persist/gpu/OpenCLBackend.cc +499 -0

…	…	@@ -0,0 +1,499 @@
	1	/*
	2	* opencog/persist/gpu/OpenCLBackend.cc
	3	*
	4	* OpenCL implementation of GpuBackend.
	5	* Manages OpenCL context, command queue, kernel compilation,
	6	* and GPU buffer operations for SoA pools.
	7	*
	8	* Copyright (C) 2025 OpenCog Foundation
	9	*
	10	* SPDX-License-Identifier: AGPL-3.0-or-later
	11	*/
	12
	13	#include <fstream>
	14	#include <sstream>
	15
	16	#include <opencog/util/exceptions.h>
	17	#include <opencog/util/Logger.h>
	18
	19	#include "OpenCLBackend.h"
	20
	21	using namespace opencog;
	22
	23	// ==============================================================
	24
	25	OpenCLBackend::~OpenCLBackend()
	26	{
	27	// OpenCL resources are RAII via cl::* wrappers
	28	}
	29
	30	// ==============================================================
	31	// Device discovery
	32
	33	void OpenCLBackend::find_device(const std::string& splat,
	34	const std::string& sdev)
	35	{
	36	std::vector<cl::Platform> platforms;
	37	cl::Platform::get(&platforms);
	38
	39	for (const auto& plat : platforms)
	40	{
	41	std::string pname = plat.getInfo<CL_PLATFORM_NAME>();
	42	if (0 < splat.size() and pname.find(splat) == std::string::npos)
	43	continue;
	44
	45	std::vector<cl::Device> devices;
	46	plat.getDevices(CL_DEVICE_TYPE_ALL, &devices);
	47	for (const cl::Device& dev : devices)
	48	{
	49	std::string dname = dev.getInfo<CL_DEVICE_NAME>();
	50	if (0 < sdev.size() and dname.find(sdev) == std::string::npos)
	51	continue;
	52
	53	_platform = plat;
	54	_device = dev;
	55	_platform_name = pname;
	56	_device_name = dname;
	57
	58	logger().info("OpenCLBackend: Using platform '%s' device '%s'\n",
	59	pname.c_str(), dname.c_str());
	60	return;
	61	}
	62	}
	63
	64	throw RuntimeException(TRACE_INFO,
	65	"OpenCLBackend: unable to find platform '%s' device '%s'\n",
	66	splat.c_str(), sdev.c_str());
	67	}
	68
	69	// ==============================================================
	70	// Kernel compilation
	71
	72	void OpenCLBackend::compile_kernels()
	73	{
	74	std::ostringstream opts;
	75	opts << "-D WORD_CAPACITY=" << GPU_WORD_CAPACITY
	76	<< " -D PAIR_CAPACITY=" << GPU_PAIR_CAPACITY
	77	<< " -D SECTION_CAPACITY=" << GPU_SECTION_CAPACITY
	78	<< " -D WORD_HT_CAPACITY=" << GPU_WORD_HT_CAPACITY
	79	<< " -D PAIR_HT_CAPACITY=" << GPU_PAIR_HT_CAPACITY
	80	<< " -D SECTION_HT_CAPACITY=" << GPU_SECTION_HT_CAPACITY;
	81
	82	std::string ht_src, as_src;
	83	std::vector<std::string> search_paths = {
	84	"/usr/local/share/opencog/opencl/gpu/",
	85	"/usr/share/opencog/opencl/gpu/",
	86	};
	87
	88	search_paths.push_back("opencog/persist/gpu/");
	89	search_paths.push_back("../opencog/persist/gpu/");
	90	search_paths.push_back("../../opencog/persist/gpu/");
	91
	92	for (const auto& prefix : search_paths)
	93	{
	94	if (not ht_src.empty() and not as_src.empty()) break;
	95
	96	if (ht_src.empty())
	97	{
	98	std::ifstream f(prefix + "gpu-hashtable.cl");
	99	if (f.is_open())
	100	ht_src.assign(std::istreambuf_iterator<char>(f),
	101	std::istreambuf_iterator<char>());
	102	}
	103	if (as_src.empty())
	104	{
	105	std::ifstream f(prefix + "gpu-atomspace.cl");
	106	if (f.is_open())
	107	as_src.assign(std::istreambuf_iterator<char>(f),
	108	std::istreambuf_iterator<char>());
	109	}
	110	}
	111
	112	if (ht_src.empty() or as_src.empty())
	113	throw RuntimeException(TRACE_INFO,
	114	"OpenCLBackend: cannot find gpu-hashtable.cl "
	115	"or gpu-atomspace.cl kernel source files\n");
	116
	117	std::string combined = ht_src + "\n" + as_src;
	118
	119	cl::Program::Sources sources;
	120	sources.push_back(combined);
	121	_program = cl::Program(_context, sources);
	122
	123	try
	124	{
	125	_program.build(opts.str().c_str());
	126	}
	127	catch (const cl::Error& e)
	128	{
	129	std::string log = _program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(_device);
	130	logger().error("OpenCLBackend kernel build failed:\n%s\n", log.c_str());
	131	throw RuntimeException(TRACE_INFO,
	132	"OpenCLBackend: kernel compilation failed\n%s\n",
	133	log.c_str());
	134	}
	135
	136	logger().info("OpenCLBackend: kernels compiled successfully\n");
	137	}
	138
	139	// ==============================================================
	140	// Lifecycle
	141
	142	void OpenCLBackend::init(const std::string& platform_hint,
	143	const std::string& device_hint)
	144	{
	145	find_device(platform_hint, device_hint);
	146	_context = cl::Context(_device);
	147	_queue = cl::CommandQueue(_context, _device);
	148	compile_kernels();
	149	}
	150
	151	void OpenCLBackend::shutdown()
	152	{
	153	_queue.finish();
	154	}
	155
	156	std::string OpenCLBackend::device_info()
	157	{
	158	return _platform_name + " / " + _device_name;
	159	}
	160
	161	// ==============================================================
	162	// Pool allocation
	163
	164	void OpenCLBackend::alloc_pools()
	165	{
	166	// -- WordPool --
	167	_word_name_hash = cl::Buffer(_context, CL_MEM_READ_WRITE,
	168	sizeof(cl_ulong) * GPU_WORD_CAPACITY);
	169	_word_count = cl::Buffer(_context, CL_MEM_READ_WRITE,
	170	sizeof(cl_double) * GPU_WORD_CAPACITY);
	171	_word_mi_marginal = cl::Buffer(_context, CL_MEM_READ_WRITE,
	172	sizeof(cl_double) * GPU_WORD_CAPACITY);
	173	_word_class_id = cl::Buffer(_context, CL_MEM_READ_WRITE,
	174	sizeof(cl_uint) * GPU_WORD_CAPACITY);
	175	_word_next_free = cl::Buffer(_context, CL_MEM_READ_WRITE,
	176	sizeof(cl_uint));
	177
	178	// -- PairPool --
	179	_pair_word_a = cl::Buffer(_context, CL_MEM_READ_WRITE,
	180	sizeof(cl_uint) * GPU_PAIR_CAPACITY);
	181	_pair_word_b = cl::Buffer(_context, CL_MEM_READ_WRITE,
	182	sizeof(cl_uint) * GPU_PAIR_CAPACITY);
	183	_pair_count = cl::Buffer(_context, CL_MEM_READ_WRITE,
	184	sizeof(cl_double) * GPU_PAIR_CAPACITY);
	185	_pair_mi = cl::Buffer(_context, CL_MEM_READ_WRITE,
	186	sizeof(cl_double) * GPU_PAIR_CAPACITY);
	187	_pair_flags = cl::Buffer(_context, CL_MEM_READ_WRITE,
	188	sizeof(cl_uint) * GPU_PAIR_CAPACITY);
	189	_pair_next_free = cl::Buffer(_context, CL_MEM_READ_WRITE,
	190	sizeof(cl_uint));
	191
	192	// -- SectionPool --
	193	_sec_word = cl::Buffer(_context, CL_MEM_READ_WRITE,
	194	sizeof(cl_uint) * GPU_SECTION_CAPACITY);
	195	_sec_djh = cl::Buffer(_context, CL_MEM_READ_WRITE,
	196	sizeof(cl_ulong) * GPU_SECTION_CAPACITY);
	197	_sec_count = cl::Buffer(_context, CL_MEM_READ_WRITE,
	198	sizeof(cl_double) * GPU_SECTION_CAPACITY);
	199	_sec_next_free = cl::Buffer(_context, CL_MEM_READ_WRITE,
	200	sizeof(cl_uint));
	201
	202	// -- Hash tables --
	203	_word_ht_keys = cl::Buffer(_context, CL_MEM_READ_WRITE,
	204	sizeof(cl_ulong) * GPU_WORD_HT_CAPACITY);
	205	_word_ht_values = cl::Buffer(_context, CL_MEM_READ_WRITE,
	206	sizeof(cl_uint) * GPU_WORD_HT_CAPACITY);
	207	_pair_ht_keys = cl::Buffer(_context, CL_MEM_READ_WRITE,
	208	sizeof(cl_ulong) * GPU_PAIR_HT_CAPACITY);
	209	_pair_ht_values = cl::Buffer(_context, CL_MEM_READ_WRITE,
	210	sizeof(cl_uint) * GPU_PAIR_HT_CAPACITY);
	211	_sec_ht_keys = cl::Buffer(_context, CL_MEM_READ_WRITE,
	212	sizeof(cl_ulong) * GPU_SECTION_HT_CAPACITY);
	213	_sec_ht_values = cl::Buffer(_context, CL_MEM_READ_WRITE,
	214	sizeof(cl_uint) * GPU_SECTION_HT_CAPACITY);
	215	}
	216
	217	void OpenCLBackend::init_pools()
	218	{
	219	cl_uint zero = 0;
	220	_queue.enqueueWriteBuffer(_word_next_free, CL_TRUE, 0, sizeof(cl_uint), &zero);
	221	_queue.enqueueWriteBuffer(_pair_next_free, CL_TRUE, 0, sizeof(cl_uint), &zero);
	222	_queue.enqueueWriteBuffer(_sec_next_free, CL_TRUE, 0, sizeof(cl_uint), &zero);
	223
	224	// Initialize hash table keys to EMPTY, values to EMPTY.
	225	// Use CL_TRUE (blocking) because host vectors go out of scope.
	226	{
	227	std::vector<cl_ulong> ek(GPU_WORD_HT_CAPACITY, GPU_HT_EMPTY_KEY);
	228	_queue.enqueueWriteBuffer(_word_ht_keys, CL_TRUE, 0,
	229	sizeof(cl_ulong) * GPU_WORD_HT_CAPACITY, ek.data());
	230	}
	231	{
	232	std::vector<cl_uint> ev(GPU_WORD_HT_CAPACITY, GPU_HT_EMPTY_VALUE);
	233	_queue.enqueueWriteBuffer(_word_ht_values, CL_TRUE, 0,
	234	sizeof(cl_uint) * GPU_WORD_HT_CAPACITY, ev.data());
	235	}
	236	{
	237	std::vector<cl_ulong> ek(GPU_PAIR_HT_CAPACITY, GPU_HT_EMPTY_KEY);
	238	_queue.enqueueWriteBuffer(_pair_ht_keys, CL_TRUE, 0,
	239	sizeof(cl_ulong) * GPU_PAIR_HT_CAPACITY, ek.data());
	240	}
	241	{
	242	std::vector<cl_uint> ev(GPU_PAIR_HT_CAPACITY, GPU_HT_EMPTY_VALUE);
	243	_queue.enqueueWriteBuffer(_pair_ht_values, CL_TRUE, 0,
	244	sizeof(cl_uint) * GPU_PAIR_HT_CAPACITY, ev.data());
	245	}
	246	{
	247	std::vector<cl_ulong> ek(GPU_SECTION_HT_CAPACITY, GPU_HT_EMPTY_KEY);
	248	_queue.enqueueWriteBuffer(_sec_ht_keys, CL_TRUE, 0,
	249	sizeof(cl_ulong) * GPU_SECTION_HT_CAPACITY, ek.data());
	250	}
	251	{
	252	std::vector<cl_uint> ev(GPU_SECTION_HT_CAPACITY, GPU_HT_EMPTY_VALUE);
	253	_queue.enqueueWriteBuffer(_sec_ht_values, CL_TRUE, 0,
	254	sizeof(cl_uint) * GPU_SECTION_HT_CAPACITY, ev.data());
	255	}
	256
	257	logger().info("OpenCLBackend: GPU pools initialized\n");
	258	}
	259
	260	// ==============================================================
	261	// Word pool operations
	262
	263	uint32_t OpenCLBackend::word_find_or_create(uint64_t nhash)
	264	{
	265	cl::Kernel kern(_program, "word_find_or_create");
	266
	267	cl::Buffer in_hash(_context, CL_MEM_READ_ONLY \| CL_MEM_COPY_HOST_PTR,
	268	sizeof(cl_ulong), &nhash);
	269	cl::Buffer out_buf(_context, CL_MEM_WRITE_ONLY, sizeof(cl_uint));
	270	cl_uint num_items = 1;
	271
	272	kern.setArg(0, _word_ht_keys);
	273	kern.setArg(1, _word_ht_values);
	274	kern.setArg(2, _word_name_hash);
	275	kern.setArg(3, _word_count);
	276	kern.setArg(4, _word_class_id);
	277	kern.setArg(5, _word_next_free);
	278	kern.setArg(6, in_hash);
	279	kern.setArg(7, out_buf);
	280	kern.setArg(8, num_items);
	281
	282	_queue.enqueueNDRangeKernel(kern, cl::NullRange, cl::NDRange(1));
	283	cl_uint out_idx = GPU_NOT_FOUND;
	284	_queue.enqueueReadBuffer(out_buf, CL_TRUE, 0, sizeof(cl_uint), &out_idx);
	285
	286	return out_idx;
	287	}
	288
	289	uint32_t OpenCLBackend::word_lookup(uint64_t nhash)
	290	{
	291	cl::Kernel kern(_program, "ht_lookup");
	292
	293	cl::Buffer query_key(_context, CL_MEM_READ_ONLY \| CL_MEM_COPY_HOST_PTR,
	294	sizeof(cl_ulong), &nhash);
	295	cl::Buffer out_buf(_context, CL_MEM_WRITE_ONLY, sizeof(cl_uint));
	296	cl_ulong capacity = GPU_WORD_HT_CAPACITY;
	297	cl_uint num_queries = 1;
	298
	299	kern.setArg(0, _word_ht_keys);
	300	kern.setArg(1, _word_ht_values);
	301	kern.setArg(2, capacity);
	302	kern.setArg(3, query_key);
	303	kern.setArg(4, out_buf);
	304	kern.setArg(5, num_queries);
	305
	306	_queue.enqueueNDRangeKernel(kern, cl::NullRange, cl::NDRange(1));
	307	cl_uint result = GPU_NOT_FOUND;
	308	_queue.enqueueReadBuffer(out_buf, CL_TRUE, 0, sizeof(cl_uint), &result);
	309
	310	return result;
	311	}
	312
	313	void OpenCLBackend::word_write_count(uint32_t idx, double count)
	314	{
	315	cl_double v = count;
	316	_queue.enqueueWriteBuffer(_word_count, CL_FALSE,
	317	sizeof(cl_double) * idx, sizeof(cl_double), &v);
	318	}
	319
	320	void OpenCLBackend::word_write_marginal(uint32_t idx, double marginal)
	321	{
	322	cl_double v = marginal;
	323	_queue.enqueueWriteBuffer(_word_mi_marginal, CL_FALSE,
	324	sizeof(cl_double) * idx, sizeof(cl_double), &v);
	325	}
	326
	327	void OpenCLBackend::word_read_values(uint32_t idx,
	328	double& count, double& marginal)
	329	{
	330	cl_double c = 0.0, m = 0.0;
	331	_queue.enqueueReadBuffer(_word_count, CL_TRUE,
	332	sizeof(cl_double) * idx, sizeof(cl_double), &c);
	333	_queue.enqueueReadBuffer(_word_mi_marginal, CL_TRUE,
	334	sizeof(cl_double) * idx, sizeof(cl_double), &m);
	335	count = c;
	336	marginal = m;
	337	}
	338
	339	void OpenCLBackend::word_delete(uint64_t nhash)
	340	{
	341	cl::Kernel kern(_program, "ht_delete");
	342	cl::Buffer del_key(_context, CL_MEM_READ_ONLY \| CL_MEM_COPY_HOST_PTR,
	343	sizeof(cl_ulong), &nhash);
	344	cl_ulong capacity = GPU_WORD_HT_CAPACITY;
	345	cl_uint num = 1;
	346
	347	kern.setArg(0, _word_ht_keys);
	348	kern.setArg(1, _word_ht_values);
	349	kern.setArg(2, capacity);
	350	kern.setArg(3, del_key);
	351	kern.setArg(4, num);
	352
	353	_queue.enqueueNDRangeKernel(kern, cl::NullRange, cl::NDRange(1));
	354	_queue.finish();
	355	}
	356
	357	uint32_t OpenCLBackend::word_pool_count()
	358	{
	359	cl_uint cnt = 0;
	360	_queue.enqueueReadBuffer(_word_next_free, CL_TRUE, 0, sizeof(cl_uint), &cnt);
	361	return cnt;
	362	}
	363
	364	void OpenCLBackend::word_read_bulk(uint32_t n, uint64_t* hashes,
	365	double* counts, double* marginals)
	366	{
	367	_queue.enqueueReadBuffer(_word_name_hash, CL_TRUE, 0,
	368	sizeof(cl_ulong) * n, hashes);
	369	_queue.enqueueReadBuffer(_word_count, CL_TRUE, 0,
	370	sizeof(cl_double) * n, counts);
	371	_queue.enqueueReadBuffer(_word_mi_marginal, CL_TRUE, 0,
	372	sizeof(cl_double) * n, marginals);
	373	}
	374
	375	// ==============================================================
	376	// Pair pool operations
	377
	378	uint32_t OpenCLBackend::pair_find_or_create(uint32_t word_a, uint32_t word_b)
	379	{
	380	cl::Kernel kern(_program, "pair_find_or_create");
	381
	382	cl::Buffer in_wa(_context, CL_MEM_READ_ONLY \| CL_MEM_COPY_HOST_PTR,
	383	sizeof(cl_uint), &word_a);
	384	cl::Buffer in_wb(_context, CL_MEM_READ_ONLY \| CL_MEM_COPY_HOST_PTR,
	385	sizeof(cl_uint), &word_b);
	386	cl::Buffer out_buf(_context, CL_MEM_WRITE_ONLY, sizeof(cl_uint));
	387	cl_uint num_items = 1;
	388
	389	kern.setArg(0, _pair_ht_keys);
	390	kern.setArg(1, _pair_ht_values);
	391	kern.setArg(2, _pair_word_a);
	392	kern.setArg(3, _pair_word_b);
	393	kern.setArg(4, _pair_count);
	394	kern.setArg(5, _pair_mi);
	395	kern.setArg(6, _pair_flags);
	396	kern.setArg(7, _pair_next_free);
	397	kern.setArg(8, in_wa);
	398	kern.setArg(9, in_wb);
	399	kern.setArg(10, out_buf);
	400	kern.setArg(11, num_items);
	401
	402	_queue.enqueueNDRangeKernel(kern, cl::NullRange, cl::NDRange(1));
	403	cl_uint out_idx = GPU_NOT_FOUND;
	404	_queue.enqueueReadBuffer(out_buf, CL_TRUE, 0, sizeof(cl_uint), &out_idx);
	405
	406	return out_idx;
	407	}
	408
	409	uint32_t OpenCLBackend::pair_lookup(uint32_t word_a, uint32_t word_b)
	410	{
	411	uint32_t lo = std::min(word_a, word_b);
	412	uint32_t hi = std::max(word_a, word_b);
	413	uint64_t key = ((uint64_t)lo << 32) \| (uint64_t)hi;
	414
	415	cl::Kernel kern(_program, "ht_lookup");
	416
	417	cl::Buffer query_key(_context, CL_MEM_READ_ONLY \| CL_MEM_COPY_HOST_PTR,
	418	sizeof(cl_ulong), &key);
	419	cl::Buffer out_buf(_context, CL_MEM_WRITE_ONLY, sizeof(cl_uint));
	420	cl_ulong capacity = GPU_PAIR_HT_CAPACITY;
	421	cl_uint num_queries = 1;
	422
	423	kern.setArg(0, _pair_ht_keys);
	424	kern.setArg(1, _pair_ht_values);
	425	kern.setArg(2, capacity);
	426	kern.setArg(3, query_key);
	427	kern.setArg(4, out_buf);
	428	kern.setArg(5, num_queries);
	429
	430	_queue.enqueueNDRangeKernel(kern, cl::NullRange, cl::NDRange(1));
	431	cl_uint result = GPU_NOT_FOUND;
	432	_queue.enqueueReadBuffer(out_buf, CL_TRUE, 0, sizeof(cl_uint), &result);
	433
	434	return result;
	435	}
	436
	437	void OpenCLBackend::pair_write_count(uint32_t idx, double count)
	438	{
	439	cl_double v = count;
	440	_queue.enqueueWriteBuffer(_pair_count, CL_FALSE,
	441	sizeof(cl_double) * idx, sizeof(cl_double), &v);
	442	}
	443
	444	void OpenCLBackend::pair_write_mi(uint32_t idx, double mi)
	445	{
	446	cl_double v = mi;
	447	_queue.enqueueWriteBuffer(_pair_mi, CL_FALSE,
	448	sizeof(cl_double) * idx, sizeof(cl_double), &v);
	449	}
	450
	451	void OpenCLBackend::pair_read_values(uint32_t idx,
	452	double& count, double& mi)
	453	{
	454	cl_double c = 0.0, m = 0.0;
	455	_queue.enqueueReadBuffer(_pair_count, CL_TRUE,
	456	sizeof(cl_double) * idx, sizeof(cl_double), &c);
	457	_queue.enqueueReadBuffer(_pair_mi, CL_TRUE,
	458	sizeof(cl_double) * idx, sizeof(cl_double), &m);
	459	count = c;
	460	mi = m;
	461	}
	462
	463	uint32_t OpenCLBackend::pair_pool_count()
	464	{
	465	cl_uint cnt = 0;
	466	_queue.enqueueReadBuffer(_pair_next_free, CL_TRUE, 0, sizeof(cl_uint), &cnt);
	467	return cnt;
	468	}
	469
	470	// ==============================================================
	471	// Section pool operations
	472
	473	uint32_t OpenCLBackend::section_pool_count()
	474	{
	475	cl_uint cnt = 0;
	476	_queue.enqueueReadBuffer(_sec_next_free, CL_TRUE, 0, sizeof(cl_uint), &cnt);
	477	return cnt;
	478	}
	479
	480	// ==============================================================
	481
	482	void OpenCLBackend::barrier()
	483	{
	484	_queue.finish();
	485	}
	486
	487	// ==============================================================
	488	// Factory function
	489
	490	namespace opencog {
	491
	492	GpuBackend* create_opencl_backend()
	493	{
	494	return new OpenCLBackend();
	495	}
	496
	497	} // namespace opencog
	498
	499	// ==============================================================

opencog/persist/gpu/OpenCLBackend.h +106 -0

…	…	@@ -0,0 +1,106 @@
	1	/*
	2	* opencog/persist/gpu/OpenCLBackend.h
	3	*
	4	* OpenCL implementation of the GpuBackend interface.
	5	* Fallback backend for AMD/Intel GPUs, also works on NVIDIA.
	6	*
	7	* Copyright (C) 2025 OpenCog Foundation
	8	*
	9	* SPDX-License-Identifier: AGPL-3.0-or-later
	10	*/
	11
	12	#ifndef _OPENCOG_OPENCL_BACKEND_H
	13	#define _OPENCOG_OPENCL_BACKEND_H
	14
	15	#include "GpuBackend.h"
	16	#include "gpu-pool-defs.h"
	17	#include "opencl-headers.h"
	18
	19	namespace opencog
	20	{
	21
	22	class OpenCLBackend : public GpuBackend
	23	{
	24	private:
	25	cl::Platform _platform;
	26	cl::Device _device;
	27	cl::Context _context;
	28	cl::CommandQueue _queue;
	29	cl::Program _program;
	30
	31	// GPU SoA pool buffers -- WordPool
	32	cl::Buffer _word_name_hash; // ulong[WORD_CAPACITY]
	33	cl::Buffer _word_count; // double[WORD_CAPACITY]
	34	cl::Buffer _word_mi_marginal; // double[WORD_CAPACITY]
	35	cl::Buffer _word_class_id; // uint[WORD_CAPACITY]
	36	cl::Buffer _word_next_free; // uint[1]
	37
	38	// GPU SoA pool buffers -- PairPool
	39	cl::Buffer _pair_word_a; // uint[PAIR_CAPACITY]
	40	cl::Buffer _pair_word_b; // uint[PAIR_CAPACITY]
	41	cl::Buffer _pair_count; // double[PAIR_CAPACITY]
	42	cl::Buffer _pair_mi; // double[PAIR_CAPACITY]
	43	cl::Buffer _pair_flags; // uint[PAIR_CAPACITY]
	44	cl::Buffer _pair_next_free; // uint[1]
	45
	46	// GPU SoA pool buffers -- SectionPool
	47	cl::Buffer _sec_word; // uint[SECTION_CAPACITY]
	48	cl::Buffer _sec_djh; // ulong[SECTION_CAPACITY]
	49	cl::Buffer _sec_count; // double[SECTION_CAPACITY]
	50	cl::Buffer _sec_next_free; // uint[1]
	51
	52	// Hash tables (keys + values for each pool)
	53	cl::Buffer _word_ht_keys; // ulong[WORD_HT_CAPACITY]
	54	cl::Buffer _word_ht_values; // uint[WORD_HT_CAPACITY]
	55	cl::Buffer _pair_ht_keys; // ulong[PAIR_HT_CAPACITY]
	56	cl::Buffer _pair_ht_values; // uint[PAIR_HT_CAPACITY]
	57	cl::Buffer _sec_ht_keys; // ulong[SECTION_HT_CAPACITY]
	58	cl::Buffer _sec_ht_values; // uint[SECTION_HT_CAPACITY]
	59
	60	std::string _device_name;
	61	std::string _platform_name;
	62
	63	void find_device(const std::string& platform, const std::string& device);
	64	void compile_kernels();
	65
	66	public:
	67	OpenCLBackend() = default;
	68	~OpenCLBackend() override;
	69
	70	void init(const std::string& platform_hint,
	71	const std::string& device_hint) override;
	72	void shutdown() override;
	73	std::string device_info() override;
	74	std::string backend_name() override { return "OpenCL"; }
	75
	76	void alloc_pools() override;
	77	void init_pools() override;
	78
	79	uint32_t word_find_or_create(uint64_t name_hash) override;
	80	uint32_t word_lookup(uint64_t name_hash) override;
	81	void word_write_count(uint32_t idx, double count) override;
	82	void word_write_marginal(uint32_t idx, double marginal) override;
	83	void word_read_values(uint32_t idx,
	84	double& count, double& marginal) override;
	85	void word_delete(uint64_t name_hash) override;
	86	uint32_t word_pool_count() override;
	87	void word_read_bulk(uint32_t n, uint64_t* hashes,
	88	double* counts, double* marginals) override;
	89
	90	uint32_t pair_find_or_create(uint32_t word_a,
	91	uint32_t word_b) override;
	92	uint32_t pair_lookup(uint32_t word_a, uint32_t word_b) override;
	93	void pair_write_count(uint32_t idx, double count) override;
	94	void pair_write_mi(uint32_t idx, double mi) override;
	95	void pair_read_values(uint32_t idx,
	96	double& count, double& mi) override;
	97	uint32_t pair_pool_count() override;
	98
	99	uint32_t section_pool_count() override;
	100
	101	void barrier() override;
	102	};
	103
	104	} // namespace opencog
	105
	106	#endif // _OPENCOG_OPENCL_BACKEND_H

opencog/persist/gpu/gpu-atomspace.cl +407 -0

…	…	@@ -0,0 +1,407 @@
	1	/*
	2	* gpu-atomspace.cl -- GPU-resident AtomSpace pools
	3	*
	4	* Three flat SoA (struct-of-arrays) pools for uniform-size atoms:
	5	* - WordPool: words/concepts (nodes)
	6	* - PairPool: word pairs (binary links)
	7	* - SectionPool: word+disjunct (sections for connector vectors)
	8	*
	9	* Each pool has:
	10	* - A bump allocator (atomic counter for next free slot)
	11	* - A hash table mapping content hash → pool index
	12	* - SoA arrays for each field (coalesced GPU access)
	13	*
	14	* All counting uses atomic operations (CAS-loop for doubles).
	15	*
	16	* This file is concatenated after gpu-hashtable.cl at load time,
	17	* so ht_hash(), atom_cmpxchg, HT_EMPTY_KEY etc. are available.
	18	*/
	19
	20	/* ─── Double-precision atomic add via CAS loop ─── */
	21
	22	/*
	23	* OpenCL has no native atomicAdd for doubles. We implement it
	24	* via compare-and-swap on the 64-bit integer representation.
	25	* Uses atom_cmpxchg from cl_khr_int64_base_atomics.
	26	*/
	27	void atomic_add_double(__global volatile double* addr, double val)
	28	{
	29	union { ulong i; double f; } next, expected, current;
	30	current.f = *addr;
	31	do {
	32	expected = current;
	33	next.f = expected.f + val;
	34	current.i = atom_cmpxchg(
	35	(__global volatile ulong*)addr,
	36	expected.i, next.i);
	37	} while (current.i != expected.i);
	38	}
	39
	40	/* ─── Pool capacity constants (set by host at compile time) ─── */
	41
	42	/* These are defined as macros by the host via -D flags:
	43	* WORD_CAPACITY max words (e.g., 131072 = 128K)
	44	* PAIR_CAPACITY max pairs (e.g., 4194304 = 4M)
	45	* SECTION_CAPACITY max sections (e.g., 1048576 = 1M)
	46	* WORD_HT_CAPACITY word hash table size (power of 2)
	47	* PAIR_HT_CAPACITY pair hash table size (power of 2)
	48	* SECTION_HT_CAPACITY section hash table size (power of 2)
	49	*/
	50
	51	/* ═══════════════════════════════════════════════════════════════
	52	* WORD POOL — Nodes
	53	*
	54	* SoA layout:
	55	* word_name_hash[N] ulong — content hash (from string)
	56	* word_count[N] double — marginal count
	57	* word_mi_marginal[N] double — sum of MI for marginal
	58	* word_class_id[N] uint — class assignment (0 = none)
	59	*
	60	* Hash table: name_hash → word_index
	61	* ═══════════════════════════════════════════════════════════════ */
	62
	63	/*
	64	* Find or create words in the pool.
	65	*
	66	* For each input name_hash, looks up in the word hash table.
	67	* If found, returns the existing index. If not found, atomically
	68	* allocates a new slot and initializes it.
	69	*
	70	* Args:
	71	* wht_keys, wht_values - word hash table (keys + values)
	72	* word_name_hash - word pool: name hashes
	73	* word_count - word pool: counts (init to 0.0)
	74	* word_class_id - word pool: class IDs (init to 0)
	75	* word_next_free - atomic bump allocator (single uint)
	76	* in_name_hashes - input: name hashes to find/create
	77	* out_indices - output: pool index for each input
	78	* num_items - number of inputs
	79	*/
	80	__kernel void word_find_or_create(
	81	__global volatile ulong* wht_keys,
	82	__global volatile uint* wht_values,
	83	__global ulong* word_name_hash,
	84	__global double* word_count,
	85	__global uint* word_class_id,
	86	__global volatile uint* word_next_free,
	87	__global const ulong* in_name_hashes,
	88	__global uint* out_indices,
	89	const uint num_items)
	90	{
	91	uint tid = get_global_id(0);
	92	if (tid >= num_items) return;
	93
	94	ulong name_hash = in_name_hashes[tid];
	95	if (name_hash == HT_EMPTY_KEY) {
	96	out_indices[tid] = HT_EMPTY_VALUE;
	97	return;
	98	}
	99
	100	ulong ht_cap = WORD_HT_CAPACITY;
	101	ulong mask = ht_cap - 1;
	102	ulong slot = ht_hash(name_hash) & mask;
	103
	104	for (uint probe = 0; probe < HT_MAX_PROBES; probe++)
	105	{
	106	ulong prev = atom_cmpxchg(&wht_keys[slot], HT_EMPTY_KEY, name_hash);
	107
	108	if (prev == HT_EMPTY_KEY)
	109	{
	110	/* We won the slot — allocate from pool and publish value */
	111	uint idx = atomic_add(word_next_free, 1U);
	112	if (idx < WORD_CAPACITY) {
	113	word_name_hash[idx] = name_hash;
	114	word_count[idx] = 0.0;
	115	word_class_id[idx] = 0;
	116	/* Write value LAST so spin-waiters see valid index */
	117	mem_fence(CLK_GLOBAL_MEM_FENCE);
	118	wht_values[slot] = idx;
	119	out_indices[tid] = idx;
	120	} else {
	121	out_indices[tid] = HT_EMPTY_VALUE;
	122	}
	123	return;
	124	}
	125	if (prev == name_hash)
	126	{
	127	/* Already exists — spin until creator publishes value */
	128	uint val = wht_values[slot];
	129	while (val == HT_EMPTY_VALUE) {
	130	val = wht_values[slot];
	131	}
	132	out_indices[tid] = val;
	133	return;
	134	}
	135
	136	slot = (slot + 1) & mask;
	137	}
	138	out_indices[tid] = HT_EMPTY_VALUE;
	139	}
	140
	141	/* ═══════════════════════════════════════════════════════════════
	142	* PAIR POOL — Binary Links (word pairs)
	143	*
	144	* SoA layout:
	145	* pair_word_a[N] uint — index into WordPool
	146	* pair_word_b[N] uint — index into WordPool
	147	* pair_count[N] double — joint count
	148	* pair_mi[N] double — mutual information
	149	* pair_flags[N] uint — dirty bit, etc.
	150	*
	151	* Hash table key: (word_a << 32) \| word_b (packed 64-bit)
	152	* ═══════════════════════════════════════════════════════════════ */
	153
	154	/* Pack two 32-bit word indices into a single 64-bit hash key */
	155	inline ulong pair_key(uint word_a, uint word_b)
	156	{
	157	/* Always store with smaller index first for canonical ordering */
	158	uint lo = min(word_a, word_b);
	159	uint hi = max(word_a, word_b);
	160	return ((ulong)lo << 32) \| (ulong)hi;
	161	}
	162
	163	/*
	164	* Find or create word pairs in the pool.
	165	*
	166	* Args:
	167	* pht_keys, pht_values - pair hash table
	168	* pair_word_a, pair_word_b - pair pool: word indices
	169	* pair_count - pair pool: counts (init to 0.0)
	170	* pair_mi - pair pool: MI values (init to 0.0)
	171	* pair_flags - pair pool: flags (init to 0)
	172	* pair_next_free - atomic bump allocator
	173	* in_word_a, in_word_b - input: word index pairs
	174	* out_indices - output: pair pool index for each input
	175	* num_items - number of inputs
	176	*/
	177	__kernel void pair_find_or_create(
	178	__global volatile ulong* pht_keys,
	179	__global volatile uint* pht_values,
	180	__global uint* pair_word_a,
	181	__global uint* pair_word_b,
	182	__global double* pair_count,
	183	__global double* pair_mi,
	184	__global uint* pair_flags,
	185	__global volatile uint* pair_next_free,
	186	__global const uint* in_word_a,
	187	__global const uint* in_word_b,
	188	__global uint* out_indices,
	189	const uint num_items)
	190	{
	191	uint tid = get_global_id(0);
	192	if (tid >= num_items) return;
	193
	194	uint wa = in_word_a[tid];
	195	uint wb = in_word_b[tid];
	196	ulong key = pair_key(wa, wb);
	197
	198	ulong ht_cap = PAIR_HT_CAPACITY;
	199	ulong mask = ht_cap - 1;
	200	ulong slot = ht_hash(key) & mask;
	201
	202	for (uint probe = 0; probe < HT_MAX_PROBES; probe++)
	203	{
	204	ulong prev = atom_cmpxchg(&pht_keys[slot], HT_EMPTY_KEY, key);
	205
	206	if (prev == HT_EMPTY_KEY)
	207	{
	208	uint idx = atomic_add(pair_next_free, 1U);
	209	if (idx < PAIR_CAPACITY) {
	210	pair_word_a[idx] = min(wa, wb);
	211	pair_word_b[idx] = max(wa, wb);
	212	pair_count[idx] = 0.0;
	213	pair_mi[idx] = 0.0;
	214	pair_flags[idx] = 0;
	215	mem_fence(CLK_GLOBAL_MEM_FENCE);
	216	pht_values[slot] = idx;
	217	out_indices[tid] = idx;
	218	} else {
	219	out_indices[tid] = HT_EMPTY_VALUE;
	220	}
	221	return;
	222	}
	223	if (prev == key)
	224	{
	225	uint val = pht_values[slot];
	226	while (val == HT_EMPTY_VALUE) {
	227	val = pht_values[slot];
	228	}
	229	out_indices[tid] = val;
	230	return;
	231	}
	232
	233	slot = (slot + 1) & mask;
	234	}
	235	out_indices[tid] = HT_EMPTY_VALUE;
	236	}
	237
	238	/* ═══════════════════════════════════════════════════════════════
	239	* SECTION POOL — Word + Disjunct
	240	*
	241	* SoA layout:
	242	* sec_word[N] uint — index into WordPool
	243	* sec_disjunct_hash[N] ulong — hash of full disjunct string
	244	* sec_count[N] double — section count
	245	*
	246	* Hash table key: hash(word_idx, disjunct_hash)
	247	* ═══════════════════════════════════════════════════════════════ */
	248
	249	/* Combine word index and disjunct hash into a single 64-bit key */
	250	inline ulong section_key(uint word_idx, ulong disjunct_hash)
	251	{
	252	/* Mix the word index into the disjunct hash */
	253	ulong key = disjunct_hash ^ ((ulong)word_idx * 0x9E3779B97F4A7C15UL);
	254	/* Ensure it's never the sentinel */
	255	if (key == HT_EMPTY_KEY) key = 0;
	256	return key;
	257	}
	258
	259	/*
	260	* Find or create sections in the pool.
	261	*
	262	* Args:
	263	* sht_keys, sht_values - section hash table
	264	* sec_word - section pool: word indices
	265	* sec_disjunct_hash - section pool: disjunct hashes
	266	* sec_count - section pool: counts (init to 0.0)
	267	* sec_next_free - atomic bump allocator
	268	* in_word_indices - input: word indices
	269	* in_disjunct_hashes - input: disjunct hashes
	270	* out_indices - output: section pool index
	271	* num_items - number of inputs
	272	*/
	273	__kernel void section_find_or_create(
	274	__global volatile ulong* sht_keys,
	275	__global volatile uint* sht_values,
	276	__global uint* sec_word,
	277	__global ulong* sec_disjunct_hash,
	278	__global double* sec_count,
	279	__global volatile uint* sec_next_free,
	280	__global const uint* in_word_indices,
	281	__global const ulong* in_disjunct_hashes,
	282	__global uint* out_indices,
	283	const uint num_items)
	284	{
	285	uint tid = get_global_id(0);
	286	if (tid >= num_items) return;
	287
	288	uint widx = in_word_indices[tid];
	289	ulong dhash = in_disjunct_hashes[tid];
	290	ulong key = section_key(widx, dhash);
	291
	292	ulong ht_cap = SECTION_HT_CAPACITY;
	293	ulong mask = ht_cap - 1;
	294	ulong slot = ht_hash(key) & mask;
	295
	296	for (uint probe = 0; probe < HT_MAX_PROBES; probe++)
	297	{
	298	ulong prev = atom_cmpxchg(&sht_keys[slot], HT_EMPTY_KEY, key);
	299
	300	if (prev == HT_EMPTY_KEY)
	301	{
	302	uint idx = atomic_add(sec_next_free, 1U);
	303	if (idx < SECTION_CAPACITY) {
	304	sec_word[idx] = widx;
	305	sec_disjunct_hash[idx] = dhash;
	306	sec_count[idx] = 0.0;
	307	mem_fence(CLK_GLOBAL_MEM_FENCE);
	308	sht_values[slot] = idx;
	309	out_indices[tid] = idx;
	310	} else {
	311	out_indices[tid] = HT_EMPTY_VALUE;
	312	}
	313	return;
	314	}
	315	if (prev == key)
	316	{
	317	uint val = sht_values[slot];
	318	while (val == HT_EMPTY_VALUE) {
	319	val = sht_values[slot];
	320	}
	321	out_indices[tid] = val;
	322	return;
	323	}
	324
	325	slot = (slot + 1) & mask;
	326	}
	327	out_indices[tid] = HT_EMPTY_VALUE;
	328	}
	329
	330	/* ═══════════════════════════════════════════════════════════════
	331	* COUNTING OPERATIONS
	332	*
	333	* Atomic double-precision increments on pool arrays.
	334	* ═══════════════════════════════════════════════════════════════ */
	335
	336	/*
	337	* Increment pair counts by 1.0 for given pair indices.
	338	* Also increments word marginal counts for both words.
	339	*/
	340	__kernel void count_pairs(
	341	__global volatile double* pair_count,
	342	__global const uint* pair_word_a,
	343	__global const uint* pair_word_b,
	344	__global volatile double* word_count,
	345	__global volatile uint* pair_flags,
	346	__global const uint* pair_indices,
	347	const uint num_items)
	348	{
	349	uint tid = get_global_id(0);
	350	if (tid >= num_items) return;
	351
	352	uint pidx = pair_indices[tid];
	353	if (pidx == HT_EMPTY_VALUE) return;
	354
	355	/* Increment pair count */
	356	atomic_add_double(&pair_count[pidx], 1.0);
	357
	358	/* Increment word marginals */
	359	uint wa = pair_word_a[pidx];
	360	uint wb = pair_word_b[pidx];
	361	atomic_add_double(&word_count[wa], 1.0);
	362	atomic_add_double(&word_count[wb], 1.0);
	363
	364	/* Mark pair as dirty (needs MI recomputation) */
	365	pair_flags[pidx] = 1;
	366	}
	367
	368	/*
	369	* Increment section counts by 1.0 for given section indices.
	370	*/
	371	__kernel void count_sections(
	372	__global volatile double* sec_count,
	373	__global const uint* sec_indices,
	374	const uint num_items)
	375	{
	376	uint tid = get_global_id(0);
	377	if (tid >= num_items) return;
	378
	379	uint sidx = sec_indices[tid];
	380	if (sidx == HT_EMPTY_VALUE) return;
	381
	382	atomic_add_double(&sec_count[sidx], 1.0);
	383	}
	384
	385	/* ═══════════════════════════════════════════════════════════════
	386	* POOL STATISTICS
	387	*
	388	* Read back pool sizes (bump allocator values).
	389	* ═══════════════════════════════════════════════════════════════ */
	390
	391	/*
	392	* Read pool sizes into output array.
	393	* out_stats[0] = word count
	394	* out_stats[1] = pair count
	395	* out_stats[2] = section count
	396	*/
	397	__kernel void pool_stats(
	398	__global const uint* word_next_free,
	399	__global const uint* pair_next_free,
	400	__global const uint* sec_next_free,
	401	__global uint* out_stats)
	402	{
	403	if (get_global_id(0) != 0) return;
	404	out_stats[0] = *word_next_free;
	405	out_stats[1] = *pair_next_free;
	406	out_stats[2] = *sec_next_free;
	407	}

opencog/persist/gpu/gpu-hashtable.cl +303 -0

…	…	@@ -0,0 +1,303 @@
	1	/*
	2	* gpu-hashtable.cl -- Lock-free GPU hash table using linear probing
	3	*
	4	* Open-addressing hash table for mapping 64-bit keys to 32-bit values.
	5	* Designed for GPU-resident AtomSpace: maps atom content hashes to
	6	* pool indices (slots in flat SoA arrays).
	7	*
	8	* Based on: nosferalatu/SimpleGPUHashTable (CUDA) and
	9	* LANL/CompactHash (OpenCL)
	10	*
	11	* Key properties:
	12	* - Lock-free concurrent insert/lookup via atomic CAS
	13	* - Linear probing with power-of-two table size
	14	* - 64-bit keys (atom content hashes)
	15	* - 32-bit values (pool indices)
	16	* - Sentinel key 0xFFFFFFFFFFFFFFFF marks empty slots
	17	* - Sentinel value 0xFFFFFFFF marks deleted entries
	18	* - 50% load factor recommended (trade memory for speed)
	19	*
	20	* Requires: cl_khr_int64_base_atomics
	21	*/
	22
	23	#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
	24
	25	/* ─── Constants ─── */
	26
	27	/* Empty slot sentinel — no valid atom hash should equal this */
	28	#define HT_EMPTY_KEY 0xFFFFFFFFFFFFFFFFUL
	29	/* Deleted value sentinel — key stays, value marks as deleted */
	30	#define HT_EMPTY_VALUE 0xFFFFFFFFU
	31	/* Max probes before giving up (prevents infinite loops on full table) */
	32	#define HT_MAX_PROBES 4096
	33
	34	/* ─── Hash function ─── */
	35
	36	/*
	37	* 64-bit finalizer from splitmix64 / Murmur3.
	38	* Produces well-distributed hashes from sequential or clustered keys.
	39	*/
	40	inline ulong ht_hash(ulong key)
	41	{
	42	key ^= key >> 30;
	43	key *= 0xBF58476D1CE4E5B9UL;
	44	key ^= key >> 27;
	45	key *= 0x94D049BB133111EBUL;
	46	key ^= key >> 31;
	47	return key;
	48	}
	49
	50	/* ─── Insert ─── */
	51
	52	/*
	53	* Insert key-value pairs into the hash table.
	54	*
	55	* Each work-item handles one (key, value) pair.
	56	* Uses atomic CAS on the key slot: if the slot is empty or already
	57	* contains our key, we write the value. Otherwise linear-probe forward.
	58	*
	59	* Args:
	60	* table_keys - key array [capacity] (ulong)
	61	* table_values - value array [capacity] (uint)
	62	* capacity - table size (must be power of 2)
	63	* in_keys - keys to insert [num_items]
	64	* in_values - values to insert [num_items]
	65	* num_items - number of items to insert
	66	*/
	67	__kernel void ht_insert(
	68	__global volatile ulong* table_keys,
	69	__global uint* table_values,
	70	const ulong capacity,
	71	__global const ulong* in_keys,
	72	__global const uint* in_values,
	73	const uint num_items)
	74	{
	75	uint tid = get_global_id(0);
	76	if (tid >= num_items) return;
	77
	78	ulong key = in_keys[tid];
	79	uint value = in_values[tid];
	80
	81	/* Don't insert the sentinel key */
	82	if (key == HT_EMPTY_KEY) return;
	83
	84	ulong mask = capacity - 1; /* power-of-2 masking */
	85	ulong slot = ht_hash(key) & mask;
	86
	87	for (uint probe = 0; probe < HT_MAX_PROBES; probe++)
	88	{
	89	/* Try to claim this slot with our key */
	90	ulong prev = atom_cmpxchg(
	91	&table_keys[slot], HT_EMPTY_KEY, key);
	92
	93	if (prev == HT_EMPTY_KEY \|\| prev == key)
	94	{
	95	/* Slot was empty or already ours — write value */
	96	table_values[slot] = value;
	97	return;
	98	}
	99
	100	/* Slot taken by another key — linear probe */
	101	slot = (slot + 1) & mask;
	102	}
	103	/* Table full or too many collisions — item not inserted */
	104	}
	105
	106	/* ─── Lookup ─── */
	107
	108	/*
	109	* Look up keys in the hash table, returning their values.
	110	*
	111	* Each work-item looks up one key. If found, writes the value
	112	* to out_values[tid]. If not found, writes HT_EMPTY_VALUE.
	113	*
	114	* Args:
	115	* table_keys - key array [capacity] (ulong)
	116	* table_values - value array [capacity] (uint)
	117	* capacity - table size (must be power of 2)
	118	* query_keys - keys to look up [num_queries]
	119	* out_values - results written here [num_queries]
	120	* num_queries - number of lookups
	121	*/
	122	__kernel void ht_lookup(
	123	__global const ulong* table_keys,
	124	__global const uint* table_values,
	125	const ulong capacity,
	126	__global const ulong* query_keys,
	127	__global uint* out_values,
	128	const uint num_queries)
	129	{
	130	uint tid = get_global_id(0);
	131	if (tid >= num_queries) return;
	132
	133	ulong key = query_keys[tid];
	134	ulong mask = capacity - 1;
	135	ulong slot = ht_hash(key) & mask;
	136
	137	for (uint probe = 0; probe < HT_MAX_PROBES; probe++)
	138	{
	139	ulong slot_key = table_keys[slot];
	140
	141	if (slot_key == key)
	142	{
	143	/* Found — read value */
	144	out_values[tid] = table_values[slot];
	145	return;
	146	}
	147	if (slot_key == HT_EMPTY_KEY)
	148	{
	149	/* Empty slot — key not in table */
	150	out_values[tid] = HT_EMPTY_VALUE;
	151	return;
	152	}
	153
	154	/* Different key — linear probe */
	155	slot = (slot + 1) & mask;
	156	}
	157	/* Max probes exhausted — not found */
	158	out_values[tid] = HT_EMPTY_VALUE;
	159	}
	160
	161	/* ─── Delete ─── */
	162
	163	/*
	164	* Delete keys from the hash table.
	165	*
	166	* Tombstone deletion: key stays in table, value set to HT_EMPTY_VALUE.
	167	* Preserves probe chains for other keys. Deleted slots are NOT reused
	168	* for new inserts (keeps the algorithm simple).
	169	*
	170	* Args:
	171	* table_keys - key array [capacity] (ulong)
	172	* table_values - value array [capacity] (uint)
	173	* capacity - table size (must be power of 2)
	174	* del_keys - keys to delete [num_deletes]
	175	* num_deletes - number of deletions
	176	*/
	177	__kernel void ht_delete(
	178	__global const ulong* table_keys,
	179	__global uint* table_values,
	180	const ulong capacity,
	181	__global const ulong* del_keys,
	182	const uint num_deletes)
	183	{
	184	uint tid = get_global_id(0);
	185	if (tid >= num_deletes) return;
	186
	187	ulong key = del_keys[tid];
	188	ulong mask = capacity - 1;
	189	ulong slot = ht_hash(key) & mask;
	190
	191	for (uint probe = 0; probe < HT_MAX_PROBES; probe++)
	192	{
	193	ulong slot_key = table_keys[slot];
	194
	195	if (slot_key == key)
	196	{
	197	/* Found — tombstone the value */
	198	table_values[slot] = HT_EMPTY_VALUE;
	199	return;
	200	}
	201	if (slot_key == HT_EMPTY_KEY)
	202	{
	203	/* Not in table */
	204	return;
	205	}
	206
	207	slot = (slot + 1) & mask;
	208	}
	209	}
	210
	211	/* ─── Insert-or-increment ─── */
	212
	213	/*
	214	* Insert key with value, or atomically increment existing value.
	215	* This is the core operation for AtomSpace counting: if the pair
	216	* atom already exists, increment its count. Otherwise create it.
	217	*
	218	* IMPORTANT: table_values must be initialized to 0 (not 0xFF)
	219	* before using this kernel. Both the "new slot" and "existing slot"
	220	* paths use atomic_add(1), which avoids the race between CAS on
	221	* the key and write to the value.
	222	*
	223	* Args:
	224	* table_keys - key array [capacity] (ulong)
	225	* table_values - value array [capacity] (uint, init to 0)
	226	* capacity - table size (must be power of 2)
	227	* in_keys - keys to insert/increment [num_items]
	228	* num_items - number of items
	229	*/
	230	__kernel void ht_insert_or_increment(
	231	__global volatile ulong* table_keys,
	232	__global volatile uint* table_values,
	233	const ulong capacity,
	234	__global const ulong* in_keys,
	235	const uint num_items)
	236	{
	237	uint tid = get_global_id(0);
	238	if (tid >= num_items) return;
	239
	240	ulong key = in_keys[tid];
	241	if (key == HT_EMPTY_KEY) return;
	242
	243	ulong mask = capacity - 1;
	244	ulong slot = ht_hash(key) & mask;
	245
	246	for (uint probe = 0; probe < HT_MAX_PROBES; probe++)
	247	{
	248	ulong prev = atom_cmpxchg(
	249	&table_keys[slot], HT_EMPTY_KEY, key);
	250
	251	if (prev == HT_EMPTY_KEY \|\| prev == key)
	252	{
	253	/* Slot is ours (new or existing) — atomic increment.
	254	* Both paths do the same thing: no race between
	255	* key placement and value update. */
	256	atomic_add(&table_values[slot], 1U);
	257	return;
	258	}
	259
	260	slot = (slot + 1) & mask;
	261	}
	262	}
	263
	264	/* ─── Iterate (compact non-empty entries) ─── */
	265
	266	/*
	267	* Collect all non-empty, non-deleted entries into output arrays.
	268	* Uses atomic counter to pack results contiguously.
	269	*
	270	* Args:
	271	* table_keys - key array [capacity] (ulong)
	272	* table_values - value array [capacity] (uint)
	273	* capacity - table size
	274	* out_keys - output keys [max_output]
	275	* out_values - output values [max_output]
	276	* out_count - atomic counter (single uint, initialized to 0)
	277	* max_output - max entries to output
	278	*/
	279	__kernel void ht_iterate(
	280	__global const ulong* table_keys,
	281	__global const uint* table_values,
	282	const ulong capacity,
	283	__global ulong* out_keys,
	284	__global uint* out_values,
	285	__global uint* out_count,
	286	const uint max_output)
	287	{
	288	uint tid = get_global_id(0);
	289	if (tid >= capacity) return;
	290
	291	ulong key = table_keys[tid];
	292	uint val = table_values[tid];
	293
	294	if (key != HT_EMPTY_KEY && val != HT_EMPTY_VALUE)
	295	{
	296	uint idx = atomic_add(out_count, 1U);
	297	if (idx < max_output)
	298	{
	299	out_keys[idx] = key;
	300	out_values[idx] = val;
	301	}
	302	}
	303	}

opencog/persist/gpu/gpu-pool-defs.h +42 -0

…	…	@@ -0,0 +1,42 @@
	1	/*
	2	* opencog/persist/gpu/gpu-pool-defs.h
	3	*
	4	* Shared pool capacity constants for GPU SoA pools.
	5	* Used by both CUDA and OpenCL backends.
	6	*
	7	* Copyright (C) 2025 OpenCog Foundation
	8	*
	9	* SPDX-License-Identifier: AGPL-3.0-or-later
	10	*/
	11
	12	#ifndef _OPENCOG_GPU_POOL_DEFS_H
	13	#define _OPENCOG_GPU_POOL_DEFS_H
	14
	15	// Pool capacity defaults (can be overridden at compile time)
	16	#ifndef GPU_WORD_CAPACITY
	17	#define GPU_WORD_CAPACITY 131072 // 128K words
	18	#endif
	19	#ifndef GPU_PAIR_CAPACITY
	20	#define GPU_PAIR_CAPACITY 4194304 // 4M pairs
	21	#endif
	22	#ifndef GPU_SECTION_CAPACITY
	23	#define GPU_SECTION_CAPACITY 1048576 // 1M sections
	24	#endif
	25
	26	// Hash tables: 2x the pool capacity (50% load factor)
	27	#ifndef GPU_WORD_HT_CAPACITY
	28	#define GPU_WORD_HT_CAPACITY 262144
	29	#endif
	30	#ifndef GPU_PAIR_HT_CAPACITY
	31	#define GPU_PAIR_HT_CAPACITY 8388608
	32	#endif
	33	#ifndef GPU_SECTION_HT_CAPACITY
	34	#define GPU_SECTION_HT_CAPACITY 2097152
	35	#endif
	36
	37	// Sentinel values (must match gpu-hashtable.cl and CUDA kernels)
	38	#define GPU_HT_EMPTY_KEY 0xFFFFFFFFFFFFFFFFULL
	39	#define GPU_HT_EMPTY_VALUE 0xFFFFFFFFU
	40	#define GPU_NOT_FOUND 0xFFFFFFFFU
	41
	42	#endif // _OPENCOG_GPU_POOL_DEFS_H

opencog/persist/gpu/opencl-headers.h +27 -0

…	…	@@ -0,0 +1,27 @@
	1	/*
	2	* opencog/persist/gpu/opencl-headers.h
	3	*
	4	* Portable OpenCL C++ header include.
	5	*
	6	* Copyright (C) 2025 Linas Vepstas
	7	*
	8	* SPDX-License-Identifier: AGPL-3.0-or-later
	9	*/
	10
	11	#ifndef _OPENCOG_GPU_OPENCL_HEADERS_H
	12	#define _OPENCOG_GPU_OPENCL_HEADERS_H
	13
	14	#define CL_HPP_ENABLE_EXCEPTIONS
	15	#define CL_HPP_TARGET_OPENCL_VERSION 300
	16
	17	#if defined __has_include
	18	#if __has_include(<CL/opencl.hpp>)
	19	#include <CL/opencl.hpp>
	20	#else
	21	#include <CL/cl.hpp>
	22	#endif
	23	#else
	24	#include <CL/opencl.hpp>
	25	#endif
	26
	27	#endif // _OPENCOG_GPU_OPENCL_HEADERS_H

opencog/scm/persist-gpu.scm +15 -0

…	…	@@ -0,0 +1,15 @@
	1	;
	2	; persist-gpu.scm -- OpenCL GPU storage backend for the AtomSpace
	3	;
	4	; Load the GPU persistence module.
	5	;
	6
	7	(define-module (opencog persist-gpu))
	8
	9	(use-modules (opencog) (opencog exec) (opencog persist))
	10	(use-modules (opencog gpu-config))
	11
	12	; Load the C++ library that implements the GPU storage node.
	13	(load-extension (string-append opencog-ext-path-persist-gpu "libpersist-gpu") "opencog_persist_gpu_init")
	14
	15	(export GpuStorageNode)

tests/CMakeLists.txt +10 -0

…	…	@@ -0,0 +1,10 @@
	1
	2	INCLUDE(AddCxxtest)
	3
	4	# All tests should load the atomspace scm from the build dir
	5	SET(GUILE_LOAD_PATH "${PROJECT_BINARY_DIR}/opencog/scm")
	6
	7	# Perform tests in component-dependency order.
	8	IF (CXXTEST_FOUND)
	9	ADD_SUBDIRECTORY (persist)
	10	ENDIF (CXXTEST_FOUND)

tests/persist/CMakeLists.txt +5 -0

…	…	@@ -0,0 +1,5 @@
	1
	2	# The atom_types.h file is written to the build directory
	3	INCLUDE_DIRECTORIES(${CMAKE_BINARY_DIR})
	4
	5	ADD_SUBDIRECTORY(gpu)

tests/persist/gpu/CMakeLists.txt +12 -0

…	…	@@ -0,0 +1,12 @@
	1
	2	# Unit tests
	3
	4	LINK_LIBRARIES(
	5	persist-gpu
	6	persist
	7	${ATOMSPACE_STORAGE_LIBRARIES}
	8	${ATOMSPACE_LIBRARIES}
	9	OpenCL::OpenCL
	10	)
	11
	12	ADD_CXXTEST(GpuStorageNodeUTest)

tests/persist/gpu/GpuStorageNodeUTest.cxxtest +444 -0

…	…	@@ -0,0 +1,444 @@
	1	/*
	2	* tests/persist/gpu/GpuStorageNodeUTest.cxxtest
	3	*
	4	* Basic store/fetch round-trip tests for GpuStorageNode.
	5	* Modeled after atomspace-rocks BasicSaveUTest.cxxtest.
	6	*
	7	* Tests Phase 1: atoms and values survive a round-trip
	8	* through GPU SoA pools.
	9	*
	10	* Copyright (C) 2025 OpenCog Foundation
	11	*
	12	* SPDX-License-Identifier: AGPL-3.0-or-later
	13	*/
	14
	15	#include <cstdio>
	16	#include <cmath>
	17
	18	#include <opencog/atoms/base/Atom.h>
	19	#include <opencog/atoms/base/Link.h>
	20	#include <opencog/atoms/base/Node.h>
	21	#include <opencog/atoms/atom_types/atom_types.h>
	22	#include <opencog/atoms/value/FloatValue.h>
	23	#include <opencog/atoms/value/VoidValue.h>
	24	#include <opencog/atomspace/AtomSpace.h>
	25	#include <opencog/persist/gpu/GpuStorageNode.h>
	26	#include <opencog/persist/gpu-types/atom_types.h>
	27
	28	#include <opencog/util/Logger.h>
	29
	30	#include <cxxtest/TestSuite.h>
	31
	32	using namespace opencog;
	33
	34	// ============================================================
	35	// Compare two atoms for equality (type, name, truth value)
	36
	37	static void atomCompare(Handle a, Handle b, const char* where)
	38	{
	39	printf("Testing at %s\n", where);
	40	TSM_ASSERT("No atom found", b != nullptr);
	41	if (nullptr == b) return;
	42
	43	TSM_ASSERT_EQUALS("Type mismatch", a->get_type(), b->get_type());
	44
	45	NodePtr na(NodeCast(a));
	46	NodePtr nb(NodeCast(b));
	47	if (na and nb)
	48	{
	49	TSM_ASSERT_EQUALS("Name mismatch", na->get_name(), nb->get_name());
	50	}
	51
	52	LinkPtr la(LinkCast(a));
	53	LinkPtr lb(LinkCast(b));
	54	if (la and lb)
	55	{
	56	TSM_ASSERT_EQUALS("Arity mismatch", la->get_arity(), lb->get_arity());
	57	}
	58
	59	ValuePtr ta = a->getValue(truth_key());
	60	ValuePtr tb = b->getValue(truth_key());
	61	if (ta or tb)
	62	{
	63	TSM_ASSERT("Missing truth value on a", ta);
	64	TSM_ASSERT("Missing truth value on b", tb);
	65	if (ta and tb)
	66	{
	67	TSM_ASSERT("Truth value miscompare", (ta) == (tb));
	68
	69	if (not ((ta) == (tb)))
	70	{
	71	FloatValuePtr fva = FloatValueCast(ta);
	72	FloatValuePtr fvb = FloatValueCast(tb);
	73	if (fva and fvb)
	74	{
	75	fprintf(stderr,
	76	"Error at %s: val[0] a=%g b=%g, val[1] a=%g b=%g\n",
	77	where,
	78	fva->value()[0], fvb->value()[0],
	79	fva->value().size() > 1 ? fva->value()[1] : -999.0,
	80	fvb->value().size() > 1 ? fvb->value()[1] : -999.0);
	81	}
	82	}
	83	}
	84	}
	85	}
	86
	87	// ============================================================
	88
	89	class GpuStorageNodeUTest : public CxxTest::TestSuite
	90	{
	91	private:
	92	std::string uri;
	93
	94	public:
	95	GpuStorageNodeUTest(void)
	96	{
	97	logger().set_level(Logger::INFO);
	98	logger().set_print_to_stdout_flag(true);
	99	logger().set_sync_flag(true);
	100
	101	// Use first available GPU device
	102	uri = "gpu://:";
	103	}
	104
	105	~GpuStorageNodeUTest() {}
	106
	107	void setUp(void) {}
	108	void tearDown(void) {}
	109
	110	void test_open_close(void);
	111	void test_single_node(void);
	112	void test_node_with_value(void);
	113	void test_link_round_trip(void);
	114	void test_value_round_trip(void);
	115	void test_multiple_nodes(void);
	116	void test_erase(void);
	117	void test_monitor(void);
	118	};
	119
	120	// ============================================================
	121	// Test 1: Can we open and close a connection?
	122
	123	void GpuStorageNodeUTest::test_open_close(void)
	124	{
	125	logger().info("BEGIN TEST: %s", __FUNCTION__);
	126
	127	AtomSpacePtr as = createAtomSpace();
	128	Handle hsn = as->add_node(GPU_STORAGE_NODE, std::string(uri));
	129	GpuStorageNodePtr store = GpuStorageNodeCast(hsn);
	130
	131	TSM_ASSERT("Failed to create GpuStorageNode", store != nullptr);
	132	TSM_ASSERT("Should not be connected yet", !store->connected());
	133
	134	store->open();
	135	TSM_ASSERT("Should be connected after open", store->connected());
	136
	137	store->close();
	138	TSM_ASSERT("Should not be connected after close", !store->connected());
	139
	140	logger().info("END TEST: %s", __FUNCTION__);
	141	}
	142
	143	// ============================================================
	144	// Test 2: Store and fetch a single node (no values)
	145
	146	void GpuStorageNodeUTest::test_single_node(void)
	147	{
	148	logger().info("BEGIN TEST: %s", __FUNCTION__);
	149
	150	AtomSpacePtr as = createAtomSpace();
	151	Handle hsn = as->add_node(GPU_STORAGE_NODE, std::string(uri));
	152	GpuStorageNodePtr store = GpuStorageNodeCast(hsn);
	153	store->open();
	154
	155	// Create a node
	156	Handle a(createNode(SCHEMA_NODE, "someNode"));
	157
	158	// Store it
	159	store->storeAtom(a, true);
	160	store->barrier();
	161	printf("Stored single node\n");
	162
	163	// Create a fresh handle with the same name and fetch
	164	Handle b(createNode(SCHEMA_NODE, "someNode"));
	165	store->getAtom(b);
	166	printf("Fetched single node\n");
	167
	168	// Verify type and name match
	169	TSM_ASSERT_EQUALS("Type mismatch", a->get_type(), b->get_type());
	170	NodePtr na(NodeCast(a));
	171	NodePtr nb(NodeCast(b));
	172	TSM_ASSERT("Should be nodes", na and nb);
	173	if (na and nb)
	174	TSM_ASSERT_EQUALS("Name mismatch", na->get_name(), nb->get_name());
	175
	176	store->erase();
	177	store->close();
	178
	179	logger().info("END TEST: %s", __FUNCTION__);
	180	}
	181
	182	// ============================================================
	183	// Test 3: Store a node with a FloatValue, fetch it back
	184
	185	void GpuStorageNodeUTest::test_node_with_value(void)
	186	{
	187	logger().info("BEGIN TEST: %s", __FUNCTION__);
	188
	189	AtomSpacePtr as = createAtomSpace();
	190	Handle hsn = as->add_node(GPU_STORAGE_NODE, std::string(uri));
	191	GpuStorageNodePtr store = GpuStorageNodeCast(hsn);
	192	store->open();
	193
	194	// Create a node with a truth value
	195	Handle a(createNode(SCHEMA_NODE, "nodeWithValue"));
	196	ValuePtr stv(createFloatValue(std::vector<double>({0.55, 0.6})));
	197	a->setValue(truth_key(), stv);
	198
	199	// Store
	200	store->storeAtom(a, true);
	201	store->barrier();
	202	printf("Stored node with value\n");
	203
	204	// Fetch back
	205	Handle b(createNode(SCHEMA_NODE, "nodeWithValue"));
	206	store->getAtom(b);
	207	printf("Fetched node with value\n");
	208
	209	// Compare
	210	atomCompare(a, b, "Node with FloatValue round-trip");
	211
	212	store->erase();
	213	store->close();
	214
	215	logger().info("END TEST: %s", __FUNCTION__);
	216	}
	217
	218	// ============================================================
	219	// Test 4: Store a binary link (two nodes), fetch it back
	220
	221	void GpuStorageNodeUTest::test_link_round_trip(void)
	222	{
	223	logger().info("BEGIN TEST: %s", __FUNCTION__);
	224
	225	AtomSpacePtr as = createAtomSpace();
	226	Handle hsn = as->add_node(GPU_STORAGE_NODE, std::string(uri));
	227	GpuStorageNodePtr store = GpuStorageNodeCast(hsn);
	228	store->open();
	229
	230	// Create two nodes and a link with a value
	231	Handle n1(createNode(SCHEMA_NODE, "wordA"));
	232	Handle n2(createNode(SCHEMA_NODE, "wordB"));
	233
	234	HandleSeq hvec;
	235	hvec.push_back(n1);
	236	hvec.push_back(n2);
	237	Handle lnk(createLink(hvec, LIST_LINK));
	238	ValuePtr stv(createFloatValue(std::vector<double>({12.5, 3.7})));
	239	lnk->setValue(truth_key(), stv);
	240
	241	// Store
	242	store->storeAtom(lnk, true);
	243	store->barrier();
	244	printf("Stored link with value\n");
	245
	246	// Fetch back -- create fresh link with same structure
	247	Handle fn1(createNode(SCHEMA_NODE, "wordA"));
	248	Handle fn2(createNode(SCHEMA_NODE, "wordB"));
	249	Handle flnk(createLink(LIST_LINK, fn1, fn2));
	250	store->getAtom(flnk);
	251	printf("Fetched link with value\n");
	252
	253	// Compare the truth value
	254	atomCompare(lnk, flnk, "Link round-trip");
	255
	256	store->erase();
	257	store->close();
	258
	259	logger().info("END TEST: %s", __FUNCTION__);
	260	}
	261
	262	// ============================================================
	263	// Test 5: storeValue / loadValue specific key
	264
	265	void GpuStorageNodeUTest::test_value_round_trip(void)
	266	{
	267	logger().info("BEGIN TEST: %s", __FUNCTION__);
	268
	269	AtomSpacePtr as = createAtomSpace();
	270	Handle hsn = as->add_node(GPU_STORAGE_NODE, std::string(uri));
	271	GpuStorageNodePtr store = GpuStorageNodeCast(hsn);
	272	store->open();
	273
	274	// Create a node with truth value
	275	Handle a(createNode(SCHEMA_NODE, "valueTestNode"));
	276	double v1 = 1.0e33 / 3.0;
	277	double v2 = 1.0e-30 / 300.0;
	278	ValuePtr stv(createFloatValue(std::vector<double>({v1, v2})));
	279	a->setValue(truth_key(), stv);
	280
	281	// Store just the value
	282	store->storeValue(a, truth_key());
	283	store->barrier();
	284	printf("Stored value\n");
	285
	286	// Create a fresh atom and load just the value
	287	Handle b(createNode(SCHEMA_NODE, "valueTestNode"));
	288	store->loadValue(b, truth_key());
	289	printf("Loaded value\n");
	290
	291	// Compare
	292	atomCompare(a, b, "Value round-trip with extreme floats");
	293
	294	// Verify the actual doubles match
	295	FloatValuePtr fva = FloatValueCast(a->getValue(truth_key()));
	296	FloatValuePtr fvb = FloatValueCast(b->getValue(truth_key()));
	297	TSM_ASSERT("Missing value on a", fva);
	298	TSM_ASSERT("Missing value on b", fvb);
	299	if (fva and fvb)
	300	{
	301	TSM_ASSERT_EQUALS("Value size mismatch",
	302	fva->value().size(), fvb->value().size());
	303	for (size_t i = 0; i < fva->value().size(); i++)
	304	{
	305	TSM_ASSERT_DELTA("Float value mismatch",
	306	fva->value()[i], fvb->value()[i], 1e-10 * fabs(fva->value()[i]));
	307	}
	308	}
	309
	310	store->erase();
	311	store->close();
	312
	313	logger().info("END TEST: %s", __FUNCTION__);
	314	}
	315
	316	// ============================================================
	317	// Test 6: Store many nodes, verify pool counts
	318
	319	void GpuStorageNodeUTest::test_multiple_nodes(void)
	320	{
	321	logger().info("BEGIN TEST: %s", __FUNCTION__);
	322
	323	AtomSpacePtr as = createAtomSpace();
	324	Handle hsn = as->add_node(GPU_STORAGE_NODE, std::string(uri));
	325	GpuStorageNodePtr store = GpuStorageNodeCast(hsn);
	326	store->open();
	327
	328	const int N = 100;
	329
	330	// Store N nodes
	331	for (int i = 0; i < N; i++)
	332	{
	333	std::string name = "testWord_" + std::to_string(i);
	334	Handle h(createNode(SCHEMA_NODE, name));
	335	double count = (double)(i + 1);
	336	double mi = count * 0.1;
	337	h->setValue(truth_key(),
	338	createFloatValue(std::vector<double>({count, mi})));
	339	store->storeAtom(h, false);
	340	}
	341	store->barrier();
	342	printf("Stored %d nodes\n", N);
	343
	344	// Fetch them all back and verify
	345	int found = 0;
	346	for (int i = 0; i < N; i++)
	347	{
	348	std::string name = "testWord_" + std::to_string(i);
	349	Handle h(createNode(SCHEMA_NODE, name));
	350	store->getAtom(h);
	351
	352	ValuePtr tv = h->getValue(truth_key());
	353	if (tv)
	354	{
	355	FloatValuePtr fv = FloatValueCast(tv);
	356	if (fv and fv->value().size() >= 2)
	357	{
	358	double expected_count = (double)(i + 1);
	359	double expected_mi = expected_count * 0.1;
	360	TSM_ASSERT_DELTA("Count mismatch",
	361	fv->value()[0], expected_count, 1e-10);
	362	TSM_ASSERT_DELTA("MI mismatch",
	363	fv->value()[1], expected_mi, 1e-10);
	364	found++;
	365	}
	366	}
	367	}
	368
	369	printf("Verified %d / %d nodes\n", found, N);
	370	TSM_ASSERT_EQUALS("Not all nodes found", found, N);
	371
	372	store->erase();
	373	store->close();
	374
	375	logger().info("END TEST: %s", __FUNCTION__);
	376	}
	377
	378	// ============================================================
	379	// Test 7: Erase clears everything
	380
	381	void GpuStorageNodeUTest::test_erase(void)
	382	{
	383	logger().info("BEGIN TEST: %s", __FUNCTION__);
	384
	385	AtomSpacePtr as = createAtomSpace();
	386	Handle hsn = as->add_node(GPU_STORAGE_NODE, std::string(uri));
	387	GpuStorageNodePtr store = GpuStorageNodeCast(hsn);
	388	store->open();
	389
	390	// Store some data
	391	Handle a(createNode(SCHEMA_NODE, "eraseMe"));
	392	a->setValue(truth_key(),
	393	createFloatValue(std::vector<double>({42.0, 7.5})));
	394	store->storeAtom(a, true);
	395	store->barrier();
	396
	397	// Erase
	398	store->erase();
	399
	400	// Try to fetch -- should not find it
	401	Handle b(createNode(SCHEMA_NODE, "eraseMe"));
	402	store->getAtom(b);
	403
	404	ValuePtr tv = b->getValue(truth_key());
	405	TSM_ASSERT("Should not have value after erase", nullptr == tv);
	406
	407	store->close();
	408
	409	logger().info("END TEST: %s", __FUNCTION__);
	410	}
	411
	412	// ============================================================
	413	// Test 8: Monitor returns useful info
	414
	415	void GpuStorageNodeUTest::test_monitor(void)
	416	{
	417	logger().info("BEGIN TEST: %s", __FUNCTION__);
	418
	419	AtomSpacePtr as = createAtomSpace();
	420	Handle hsn = as->add_node(GPU_STORAGE_NODE, std::string(uri));
	421	GpuStorageNodePtr store = GpuStorageNodeCast(hsn);
	422	store->open();
	423
	424	// Store a node
	425	Handle a(createNode(SCHEMA_NODE, "monitorTest"));
	426	store->storeAtom(a, true);
	427	store->barrier();
	428
	429	std::string stats = store->monitor();
	430	printf("Monitor output:\n%s\n", stats.c_str());
	431
	432	TSM_ASSERT("Monitor should contain GPU", stats.find("gpu://") != std::string::npos);
	433	TSM_ASSERT("Monitor should show connected",
	434	stats.find("Connected: yes") != std::string::npos);
	435	TSM_ASSERT("Monitor should show word count",
	436	stats.find("Words:") != std::string::npos);
	437
	438	store->erase();
	439	store->close();
	440
	441	logger().info("END TEST: %s", __FUNCTION__);
	442	}
	443
	444	/* ============================= END OF FILE ================= */

atomspace-gpu commits / 3c59d87

Changed files (28) +4756 -0