Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Step 1 - Transpiler #595

Merged
merged 726 commits into from
Jan 9, 2024
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
726 commits
Select commit Hold shift + click to select a range
3373605
tidy whitespace
neworderofjamie Jul 27, 2023
f4aca43
re-enable fixed number total connectivity test
neworderofjamie Jul 27, 2023
af1453f
fixed some custom update code
neworderofjamie Jul 27, 2023
878a52f
start of custom update test
neworderofjamie Jul 27, 2023
c94ff77
fixed bug in ``EnvironmentLocalCacheBase`` which broke variable refer…
neworderofjamie Jul 27, 2023
585d6c3
custom update code consists of multiple statements not an expression!
neworderofjamie Jul 27, 2023
533ec36
fixed various small bugs relating to non-per neuron custom updates
neworderofjamie Jul 27, 2023
a3a90ee
more work on custom update test
neworderofjamie Jul 27, 2023
481d366
removed tests of deprecated features
neworderofjamie Jul 27, 2023
9a390bc
added current source to custom update test
neworderofjamie Jul 28, 2023
4b7e937
unit tests use experiments/filesystem on GCC older than 8
neworderofjamie Jul 28, 2023
7c90974
unit tests rename of Models::VarInit
neworderofjamie Jul 28, 2023
c5bb2bb
unit tests remove ModelSpecMerged backend constructor parameter
neworderofjamie Jul 28, 2023
8e9bee6
missing macro to add EGP refs to custom update model from C++
neworderofjamie Jul 28, 2023
61c24f1
unit tests fixed bad merge
neworderofjamie Jul 28, 2023
cc624b4
unit tests one more additional ModelSpecMerged constructor parameter
neworderofjamie Jul 28, 2023
3f8ed10
Correctly handle std::experimental::filesystem namespace and use stri…
neworderofjamie Jul 28, 2023
c46eac6
more bad mergery
neworderofjamie Jul 28, 2023
58b7023
small tidy
neworderofjamie Jul 31, 2023
9ba2f95
WIP test_spike_propagation/test_reverse feature test
neworderofjamie Jul 31, 2023
a0690a6
fixed bug in NeuronInitGroupMerged::OutSynPreOutput
neworderofjamie Jul 31, 2023
50d90a1
added 'reverse' spike propagation test for addToPre functionality
neworderofjamie Jul 31, 2023
32270f5
Fixed bugs and tidied up code generation for SIMT synapse groups with…
neworderofjamie Jul 31, 2023
e707a9e
removed re-implemented test
neworderofjamie Jul 31, 2023
24e0adf
removed re-implemented test
neworderofjamie Jul 31, 2023
da4efe0
feature test for presynaptic output from post-learn code
neworderofjamie Jul 31, 2023
ad3acf3
fixed a few typos in post-learn code generation
neworderofjamie Jul 31, 2023
ca02fec
fixed a few more bugs with post-learn code generation
neworderofjamie Jul 31, 2023
bc330ff
add some presynaptic paralellism tests
neworderofjamie Jul 31, 2023
5164f43
and test for presynaptic parallelism with dendritic delay
neworderofjamie Jul 31, 2023
1b6f65b
renamed test module
neworderofjamie Jul 31, 2023
8d7b286
and test it contains
neworderofjamie Jul 31, 2023
92bd3b3
more custom update tests
neworderofjamie Jul 31, 2023
80db5ee
more custom update tests and small bug fixes
neworderofjamie Jul 31, 2023
7a5f39b
extend custom update tests to custom update and kernel weights
neworderofjamie Jul 31, 2023
59bd45d
remove reimplemented test
neworderofjamie Jul 31, 2023
b3e327b
fixed couple of small bugs
neworderofjamie Jul 31, 2023
888df72
default empty parameters dictionary for init_var, init_connectivity etc
neworderofjamie Jul 31, 2023
fe64b99
expose binomial
neworderofjamie Jul 31, 2023
d64383e
started implementing var_init test
neworderofjamie Jul 31, 2023
b0b48a6
fixed typo in RNG naming
neworderofjamie Aug 1, 2023
056fbb2
sketch of var init test for discrete distributions - can't quite figu…
neworderofjamie Aug 1, 2023
30fb641
fixed several issues around transpose custom updates and added featur…
neworderofjamie Aug 1, 2023
2161469
remove re-implemented test
neworderofjamie Aug 1, 2023
529c6be
fixed duplicate code and missing ;
neworderofjamie Aug 1, 2023
3533148
WIP feature test for custom connectivity update
neworderofjamie Aug 1, 2023
f0d98b3
missing overrides for functions to create var references to custom co…
neworderofjamie Aug 1, 2023
cfc42d3
fixed a few small typos in test
neworderofjamie Aug 1, 2023
238995b
added some more fields to the standard custom connectivity update env…
neworderofjamie Aug 1, 2023
2a5642d
parser fix
neworderofjamie Aug 1, 2023
148d71b
inject standard custom connectivity update environment in appropriate…
neworderofjamie Aug 1, 2023
6f6d28d
fixed typo
neworderofjamie Aug 2, 2023
3763bf8
fixed GCC warning
neworderofjamie Aug 2, 2023
0027872
Fixed very nasty bug:
neworderofjamie Aug 2, 2023
d9d298c
fixed small bug in for_each_synapse loop generation
neworderofjamie Aug 2, 2023
311d6d6
fixed small typo in PyGeNN
neworderofjamie Aug 2, 2023
af81219
fixed small typo
neworderofjamie Aug 2, 2023
7baad0f
implemented first part of custom connectivity update test
neworderofjamie Aug 2, 2023
f92f80e
removed addition of duplicate fields in custom connectivity init group
neworderofjamie Aug 3, 2023
15fbd22
fixed unit tests
neworderofjamie Aug 3, 2023
6c5d821
check scanning and parsing errors separately in ``parseNumericType``
neworderofjamie Aug 3, 2023
4818df9
give error if there are extra unexpected tokens following expression …
neworderofjamie Aug 3, 2023
ca99523
unit tests for parser
neworderofjamie Aug 3, 2023
5490f28
more robust checking of error state in ``prettyPrintStatements`` and …
neworderofjamie Aug 3, 2023
3adc579
improved error message
neworderofjamie Aug 3, 2023
2c9ebe9
handling of legacy EGP pointer type strings
neworderofjamie Aug 3, 2023
16f31a2
expose Var, VarRef, EGP and EGPRef constructors which take resolved …
neworderofjamie Aug 3, 2023
fdb2c1e
added EGP host custom connectivity update
neworderofjamie Aug 3, 2023
5e24440
switched from & operator to +
neworderofjamie Aug 3, 2023
00f1d21
fixed a few typos
neworderofjamie Aug 3, 2023
828921c
almost complete custom connectivity update test
neworderofjamie Aug 3, 2023
0d471c7
* seperate CustomUpdateMixin and CustomUpdateWUMixin
neworderofjamie Aug 3, 2023
ff55d7b
finished custom connectivity update test
neworderofjamie Aug 3, 2023
e637a94
remove now reimplemented-test
neworderofjamie Aug 3, 2023
7cb3807
generalised callbacks implemented for transpose to also insert inner …
neworderofjamie Aug 3, 2023
8c16073
implemented neuron reduction batch size one test
neworderofjamie Aug 3, 2023
4509b47
removed reimpleemented test
neworderofjamie Aug 3, 2023
89e79c9
fixed issues resulting to the different parallelism strategies used w…
neworderofjamie Aug 3, 2023
9caba45
implemented batched neuron reduction test
neworderofjamie Aug 3, 2023
4077a38
removed reimpleemented test
neworderofjamie Aug 3, 2023
41daf90
renamed some tests
neworderofjamie Aug 3, 2023
78923d6
fixed similar grossness resulting from different parallelism for cust…
neworderofjamie Aug 3, 2023
119d2bd
WIP custom_update_transpose_batch test
neworderofjamie Aug 3, 2023
f1b01e0
fixed warnings
neworderofjamie Aug 4, 2023
efa88d4
for custom updates, split the building of standard environments into two
neworderofjamie Aug 4, 2023
921d90b
fixed crash in corner case where merged group has empty structure - s…
neworderofjamie Aug 4, 2023
848d699
fixed numerous small issues around batching
neworderofjamie Aug 4, 2023
2f50c0c
removed reimpleemented test
neworderofjamie Aug 4, 2023
5d0a4fc
fixed much more exciting bug in type system where qualified rather th…
neworderofjamie Aug 4, 2023
1af2362
made batching an additional test parameter not seperate test
neworderofjamie Aug 4, 2023
95c3ffa
fixed compilation warning
neworderofjamie Aug 4, 2023
b3f5b5d
missing batch size
neworderofjamie Aug 4, 2023
1574ce5
removed reimpleemented test
neworderofjamie Aug 4, 2023
6043cde
extended test_custom_update to cover batching
neworderofjamie Aug 4, 2023
05806f1
tidy up kernel size for initialisation
neworderofjamie Aug 7, 2023
fc7dc50
need kernel size even if batching isn't enabled
neworderofjamie Aug 7, 2023
f5f6cc4
add _kernel_size as well as _size for custom updates
neworderofjamie Aug 7, 2023
4e68f53
wrong variable
neworderofjamie Aug 7, 2023
8d9c084
fixed a couple of tiny custom update WU issues
neworderofjamie Aug 7, 2023
23f0fd1
exposed _is_batched property in CustomUpdate and CustomUpdateWU and u…
neworderofjamie Aug 7, 2023
3980e23
refactor of custom update to enable future plans and enable fix for c…
neworderofjamie Aug 7, 2023
43167bb
fixed nasty issue with custom update finalisation order meaning batch…
neworderofjamie Aug 7, 2023
bc32536
move custom models out of module scope into tests
neworderofjamie Aug 7, 2023
58abde6
fixed very small bug
neworderofjamie Aug 7, 2023
3b1db72
WIP test_custom_update_batch_reduction
neworderofjamie Aug 7, 2023
ce39327
GCC fix
neworderofjamie Aug 8, 2023
b298cfc
fixed a couple more small bugs in batch reductions
neworderofjamie Aug 8, 2023
c474ebe
small PyGeNN fixes
neworderofjamie Aug 8, 2023
4194ebf
added level of protection to variable views to prevent e.g. ``vars["x…
neworderofjamie Aug 8, 2023
e75412d
fixed ``SynapseGroupMixin._get_view_values``
neworderofjamie Aug 8, 2023
fc5502b
completed test implementation
neworderofjamie Aug 8, 2023
b36d347
removed reimplemented tests
neworderofjamie Aug 8, 2023
fb78563
fixed a couple of typos in tests
neworderofjamie Aug 8, 2023
bb3b7ef
fixed up single-threaded CPU
neworderofjamie Aug 8, 2023
8ac547c
moved models into test
neworderofjamie Aug 8, 2023
3e9dbf2
WIP WU VAR feature test
neworderofjamie Aug 8, 2023
f7c56ae
extended Scanner/DecimalFloat unit test to cover floating point numbe…
neworderofjamie Aug 8, 2023
b793dd1
fixed scanning logic for floating point numbers specified in scientif…
neworderofjamie Aug 8, 2023
844a018
at least partially fixed logic in ``InSynWUMPostCode::genCopyDelayedV…
neworderofjamie Aug 8, 2023
69c27e4
fixed PyGeNN access to spkQuePtr
neworderofjamie Aug 8, 2023
a591f12
use correct number of delay steps when initialising pre and postsynap…
neworderofjamie Aug 8, 2023
b42b78a
updated empty code and identifier referenced tests
neworderofjamie Aug 8, 2023
4b6ed32
remove re-implemented test
neworderofjamie Aug 8, 2023
184df7f
only use a subset of the WUM hash for determining whether to fuse WU …
neworderofjamie Aug 8, 2023
bf85ba0
extended test to post wu vars in sim and synapse dynamics code as wel…
neworderofjamie Aug 8, 2023
57315dd
added test_wu_vars_pre test
neworderofjamie Aug 8, 2023
bc5b0f4
remove re-implemented test
neworderofjamie Aug 8, 2023
4231747
fixed typo in single-threaded CPU backend
neworderofjamie Aug 8, 2023
0085736
merge pre and postsynaptic WU var tests together
neworderofjamie Aug 8, 2023
989fb00
converted convolution test data to python
neworderofjamie Aug 8, 2023
5425465
added, currently totally broken, toeplitz convolution test to spike p…
neworderofjamie Aug 8, 2023
60278da
renamed $(_src_spk_evnt_cnt) to $(_src_spk_cnt_evnt) for consistency …
neworderofjamie Aug 8, 2023
ff8e5f1
removed diagonal row build vars from toeplitz connectivity snippet an…
neworderofjamie Aug 8, 2023
5ee357e
missing PyGeNN create_toeplitz_connect_init_snippet implementation
neworderofjamie Aug 8, 2023
ae7be0c
removed diagonal build state vars from PyGeNN
neworderofjamie Aug 8, 2023
cd97c48
renamed addSynapse to add_synapse - not sure if this is a good idea o…
neworderofjamie Aug 8, 2023
208fd82
small tidy
neworderofjamie Aug 8, 2023
12219e6
WIP new single-threaded CPU implementation of Toeplitz connectivity
neworderofjamie Aug 8, 2023
8cde4c6
add extra functionality to environment to allow it to be constructed …
neworderofjamie Aug 8, 2023
b15e736
hooked extra functionality into toeplitz connectivity generation but …
neworderofjamie Aug 8, 2023
cb0afa2
finished test_forward_toeplitz test
neworderofjamie Aug 9, 2023
7502072
somewhat hacky fix for CPU toeplitz connectivity
neworderofjamie Aug 9, 2023
5e4980f
SIMT toeplitz connectivity
neworderofjamie Aug 9, 2023
62ea0f3
strip out old support code system
neworderofjamie Aug 9, 2023
b7d061d
* renamed add_synapse back to addSynapse
neworderofjamie Aug 9, 2023
233df81
moved stubbing out of PROCEDURAL weights in a bit :)
neworderofjamie Aug 9, 2023
29f86df
fixed a couple of small typos
neworderofjamie Aug 9, 2023
4ef5537
test_forward_procedural test with constant weights
neworderofjamie Aug 9, 2023
e072756
removed reimpleemented test
neworderofjamie Aug 9, 2023
c7b9d92
fixed up procedural weight generation
neworderofjamie Aug 9, 2023
2f43501
added test for DENSE_PROCEDURALG
neworderofjamie Aug 9, 2023
a851180
fixed a couple of bugs with procedural connectivity
neworderofjamie Aug 9, 2023
dd30447
removed reimpleemented test
neworderofjamie Aug 9, 2023
ccd34e4
fixed typos in ``InitSparseConnectivitySnippet::Conv2D``
neworderofjamie Aug 9, 2023
19cd987
added integer abs, min and max to standard library
neworderofjamie Aug 9, 2023
7d6fb26
reordered spike propagation tests and added new one for procedural ke…
neworderofjamie Aug 9, 2023
1c0f4d8
remove comment
neworderofjamie Aug 9, 2023
f5ad079
fix typos in ``InitVarSnippet::Kernel``
neworderofjamie Aug 9, 2023
7c80da0
purge of raw string based "is there any code?" checks
neworderofjamie Aug 9, 2023
63f8ba3
test for initialising sparse connectivity from kernel
neworderofjamie Aug 9, 2023
5ab25c2
fixed up single-threaded CPU generation of connectivity from kernels
neworderofjamie Aug 9, 2023
12ff99d
added cast to model
neworderofjamie Aug 9, 2023
ccba0f6
fixed up SIMT generation of connectivity from kernels
neworderofjamie Aug 9, 2023
9f08a1f
remove unimplemented tests
neworderofjamie Aug 9, 2023
c42828f
added missing fixture
neworderofjamie Aug 9, 2023
e13a52e
properly disable generation of dendritic delay update groups in singl…
neworderofjamie Aug 9, 2023
7601bfc
added hybrid parallelism test
neworderofjamie Aug 10, 2023
6b7b20d
removed unnecessary logic and renamed lrevInSyn to lOutPre for consis…
neworderofjamie Aug 10, 2023
431aa30
test for pre-output with presynaptic parallelism
neworderofjamie Aug 10, 2023
f81d50c
update of spike time access
neworderofjamie Aug 10, 2023
6276c4c
removed wrapper code for removed bits of ``WeightUpdateModelBase``
neworderofjamie Aug 10, 2023
312acd6
WIP spike time test
neworderofjamie Aug 10, 2023
c36582c
no need for this
neworderofjamie Aug 10, 2023
436041d
fixed nasty bug in printSubs where single character variable names li…
neworderofjamie Aug 10, 2023
5cb0502
made naming of spike times more consistent and fixed small bugs
neworderofjamie Aug 10, 2023
1e94416
neuron groups require spike times if neuron code references them too
neworderofjamie Aug 10, 2023
984d27a
completed test
neworderofjamie Aug 10, 2023
3c5c91d
pass dt through everywhere
neworderofjamie Aug 10, 2023
ecdfe2d
no need for convoluted checks of non-merged groups in CUDA backend no…
neworderofjamie Aug 10, 2023
5537672
Revert "pass dt through everywhere"
neworderofjamie Aug 10, 2023
4bb6136
need to download connectivity in test
neworderofjamie Aug 10, 2023
1161827
removed reimplemented tests
neworderofjamie Aug 10, 2023
7955a8b
also test post-synaptic spike times
neworderofjamie Aug 10, 2023
ed967b9
removed reimplemented tests
neworderofjamie Aug 10, 2023
e64c7e8
set confidence interval for sim RNG tests using same logic as var ini…
neworderofjamie Aug 10, 2023
9864618
rough and ready Jenkins CI
neworderofjamie Aug 10, 2023
5234c47
yey! now groups are merged during ModelSpecMerged constructor, no nee…
neworderofjamie Aug 10, 2023
c5a0339
temporarily disable OpenCL
neworderofjamie Aug 10, 2023
0c64c67
one pop too many!
neworderofjamie Aug 10, 2023
c38d773
try another way of disabling OPENCL
neworderofjamie Aug 10, 2023
c65b768
wrong output directory for GeNN building
neworderofjamie Aug 10, 2023
7de7bd7
remove very elderly MPI crud
neworderofjamie Aug 10, 2023
2bc6e87
add support for coverage builds of PyGeNN to setup.py
neworderofjamie Aug 10, 2023
b74fd23
two small fixes for custom connectivity updates relating to variables…
neworderofjamie Aug 11, 2023
a2199b1
extend test_sim to test batched simulation
neworderofjamie Aug 11, 2023
ab1f4d6
remove reimplemented test
neworderofjamie Aug 11, 2023
1724977
custom update variables are never batched and references should refle…
neworderofjamie Aug 11, 2023
b5d9251
extend custom connectivity update test to batching
neworderofjamie Aug 11, 2023
17cf6d3
remove reimplemented test
neworderofjamie Aug 11, 2023
2958cd6
WIP var init egp test
neworderofjamie Aug 11, 2023
13fbdc6
extend test - crashing
neworderofjamie Aug 11, 2023
28f4616
correct pop
neworderofjamie Aug 11, 2023
9a33b81
remove reimplemented test
neworderofjamie Aug 11, 2023
683b856
corrected test
neworderofjamie Aug 11, 2023
54bb6fa
implemented various missing bits of EGP reference wrapping
neworderofjamie Aug 11, 2023
b6e6fd4
todo in RNG test
neworderofjamie Aug 11, 2023
51d3e59
``Model::Base::EGPRef`` string types should also be processed in same…
neworderofjamie Aug 11, 2023
c6a3b7d
EGP reference test
neworderofjamie Aug 11, 2023
201690a
remove reimplemented test
neworderofjamie Aug 11, 2023
2fab2dc
remove unsupported test
neworderofjamie Aug 11, 2023
2cad78f
remove reimplemented test
neworderofjamie Aug 11, 2023
54547b0
spike recording test
neworderofjamie Aug 11, 2023
3d135e7
parameterised delay and added test_wu_var_cont test
neworderofjamie Aug 11, 2023
4c79759
remove reimplemented test
neworderofjamie Aug 11, 2023
0bfbc92
In neuron kernel, if variables are delayed, they should be copied bac…
neworderofjamie Aug 11, 2023
660fecf
fixed typo
neworderofjamie Aug 14, 2023
ffd547e
re-enabled caching
neworderofjamie Aug 14, 2023
31e3895
WIP add numeric value type
neworderofjamie Aug 14, 2023
df18460
replaced various literal printing routines with Type::writeNumeric
neworderofjamie Aug 14, 2023
8df4c6a
missing return - how did GCC not detect this!?
neworderofjamie Aug 14, 2023
44b9ace
use 64-bit indices for synapses where required
neworderofjamie Aug 14, 2023
3f76f31
fixed typos in CPU backend
neworderofjamie Aug 14, 2023
c8dc6b2
removed arbitrary code initialisers for additional input vars from un…
neworderofjamie Aug 14, 2023
d79ac4e
typo in wrapper
neworderofjamie Aug 14, 2023
c4d4fac
maybe fhopefully fixed feature test filename
neworderofjamie Aug 14, 2023
5da5393
SPACE
neworderofjamie Aug 14, 2023
a261ca6
maybe you need to install for pytest to work
neworderofjamie Aug 14, 2023
a92ef1b
tidy up escaping
neworderofjamie Aug 14, 2023
295d784
* no pushd and popd in sh
neworderofjamie Aug 14, 2023
0234dfa
more dumb paths
neworderofjamie Aug 14, 2023
6064f9b
missing quote
neworderofjamie Aug 14, 2023
22805e0
improved numeric printing logic
neworderofjamie Aug 15, 2023
c5c0b6a
unit test for numeric printing
neworderofjamie Aug 15, 2023
3df1efa
* option to force variables and variable referneces to be read-only i…
neworderofjamie Aug 15, 2023
b9e36c8
fixed typo in pytest
neworderofjamie Aug 15, 2023
c1ddd05
isTrueSpikeRequired needed in neuron hash
neworderofjamie Aug 15, 2023
2b1c0c0
Merge branch 'transpiler' of https://github.com/genn-team/genn into t…
neworderofjamie Aug 15, 2023
b4b2975
fixed test - time_min is different depending on precision
neworderofjamie Aug 15, 2023
2af369b
build pygenn with coverage support
neworderofjamie Aug 15, 2023
3ad62eb
bloody pip
neworderofjamie Aug 15, 2023
5ee79de
give up on pip entirely
neworderofjamie Aug 15, 2023
087d317
add coverage compiler flags if PyGeNN is built with coverage
neworderofjamie Aug 15, 2023
07aeeaf
missing colon
neworderofjamie Aug 15, 2023
a31da6e
coverage linker arguments
neworderofjamie Aug 15, 2023
b08fd1d
new bash script for gathering combined feature and unit test coverage
neworderofjamie Aug 15, 2023
cbf088b
run gather coverage from jenkins and upload both
neworderofjamie Aug 15, 2023
45c5b69
* tee compile output to seperate file for warnings
neworderofjamie Aug 15, 2023
46dfb23
Always copy logic in ``EnvironmentLocalCacheBase`` was overly pessimi…
neworderofjamie Aug 16, 2023
3870b19
fixed failing test
neworderofjamie Aug 17, 2023
ca12c31
fix bug in environment
neworderofjamie Sep 1, 2023
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
The table of contents is too big for display.
Diff view
Diff view
  •  
  •  
  •  
189 changes: 128 additions & 61 deletions Jenkinsfile

Large diffs are not rendered by default.

309 changes: 193 additions & 116 deletions include/genn/backends/cuda/backend.h

Large diffs are not rendered by default.

16 changes: 7 additions & 9 deletions include/genn/backends/cuda/optimiser.h
Original file line number Diff line number Diff line change
Expand Up @@ -7,25 +7,23 @@
#include "backend.h"

// Forward declarations
namespace GeNN
{
class ModelSpecInternal;
}

namespace plog
{
class IAppender;
}


//--------------------------------------------------------------------------
// CodeGenerator::CUDA::Optimiser
// GeNN::CodeGenerator::CUDA::Optimiser
//--------------------------------------------------------------------------
namespace CodeGenerator
{
namespace CUDA
{
namespace Optimiser
namespace GeNN::CodeGenerator::CUDA::Optimiser
{
BACKEND_EXPORT Backend createBackend(const ModelSpecInternal &model, const filesystem::path &outputPath,
plog::Severity backendLevel, plog::IAppender *backendAppender,
const Preferences &preferences);
} // namespace Optimiser
} // namespace CUDA
} // namespace CodeGenerator
} // namespace GeNN::CodeGenerator::CUDA::Optimiser
34 changes: 12 additions & 22 deletions include/genn/backends/opencl/backend.h
Original file line number Diff line number Diff line change
Expand Up @@ -28,11 +28,9 @@ namespace filesystem
}

//--------------------------------------------------------------------------
// CodeGenerator::OpenCL::DeviceSelectMethod
// GeNN::CodeGenerator::OpenCL::DeviceSelectMethod
//--------------------------------------------------------------------------
namespace CodeGenerator
{
namespace OpenCL
namespace GeNN::CodeGenerator::OpenCL
{
//! Methods for selecting OpenCL platform
enum class PlatformSelect
Expand Down Expand Up @@ -124,29 +122,25 @@ class BACKEND_EXPORT Backend : public BackendSIMT
virtual void genPopulationRNGInit(CodeStream &os, const std::string &globalRNG, const std::string &seed, const std::string &sequence) const override;

//! Generate a preamble to add substitution name for population RNG
virtual void genPopulationRNGPreamble(CodeStream &os, Substitutions &subs, const std::string &globalRNG, const std::string &name = "rng") const override;
virtual std::string genPopulationRNGPreamble(CodeStream &os, const std::string &globalRNG) const override;

//! If required, generate a postamble for population RNG
/*! For example, in OpenCL, this is used to write local RNG state back to global memory*/
virtual void genPopulationRNGPostamble(CodeStream &os, const std::string &globalRNG) const override;

//! Generate code to skip ahead local copy of global RNG
virtual void genGlobalRNGSkipAhead(CodeStream &os, Substitutions &subs, const std::string &sequence, const std::string &name = "rng") const override;
virtual std::string genGlobalRNGSkipAhead(CodeStream &os, const std::string &sequence) const override;

//--------------------------------------------------------------------------
// CodeGenerator::BackendBase:: virtuals
//--------------------------------------------------------------------------
virtual void genNeuronUpdate(CodeStream &os, const ModelSpecMerged &modelMerged,
HostHandler preambleHandler, HostHandler pushEGPHandler) const override;
virtual void genNeuronUpdate(CodeStream &os, const ModelSpecMerged &modelMerged, HostHandler preambleHandler) const override;

virtual void genSynapseUpdate(CodeStream &os, const ModelSpecMerged &modelMerged,
HostHandler preambleHandler, HostHandler pushEGPHandler) const override;
virtual void genSynapseUpdate(CodeStream &os, const ModelSpecMerged &modelMerged, HostHandler preambleHandler) const override;

virtual void genCustomUpdate(CodeStream &os, const ModelSpecMerged &modelMerged,
HostHandler preambleHandler, HostHandler pushEGPHandler) const override;
virtual void genCustomUpdate(CodeStream &os, const ModelSpecMerged &modelMerged, HostHandler preambleHandler) const override;

virtual void genInit(CodeStream &os, const ModelSpecMerged &modelMerged,
HostHandler preambleHandler, HostHandler initPushEGPHandler, HostHandler initSparsePushEGPHandler) const override;
virtual void genInit(CodeStream &os, const ModelSpecMerged &modelMerged, HostHandler preambleHandler) const override;

virtual void genDefinitionsPreamble(CodeStream &os, const ModelSpecMerged &modelMerged) const override;
virtual void genDefinitionsInternalPreamble(CodeStream &os, const ModelSpecMerged &modelMerged) const override;
Expand Down Expand Up @@ -175,10 +169,7 @@ class BACKEND_EXPORT Backend : public BackendSIMT
const std::string &egpName) const override;

//! When generating function calls to push to merged groups, backend without equivalent of Unified Virtual Addressing e.g. OpenCL 1.2 may use different types on host
virtual std::string getMergedGroupFieldHostType(const std::string &type) const override;

//! When generating merged structures what type to use for simulation RNGs
virtual std::string getMergedGroupSimRNGType() const override { return "clrngLfsr113HostStream"; }
virtual std::string getMergedGroupFieldHostTypeName(const Type::Base *type, const Type::TypeContext &context) const override;

virtual void genVariablePush(CodeStream &os, const std::string &type, const std::string &name, VarLocation loc, bool autoInitialized, size_t count) const override;
virtual void genVariablePull(CodeStream &os, const std::string &type, const std::string &name, VarLocation loc, size_t count) const override;
Expand Down Expand Up @@ -367,7 +358,7 @@ class BACKEND_EXPORT Backend : public BackendSIMT
const auto sortedFields = g.getSortedFields(*this);
for(size_t fieldIndex = 0; fieldIndex < sortedFields.size(); fieldIndex++) {
const auto &f = sortedFields[fieldIndex];
if(::Utils::isTypePointer(std::get<0>(f))) {
if(GeNN::Utils::isTypePointer(std::get<0>(f))) {
os << "__global ";
}
os << std::get<0>(f) << " " << std::get<1>(f);
Expand Down Expand Up @@ -396,7 +387,7 @@ class BACKEND_EXPORT Backend : public BackendSIMT

os << "__kernel void setMerged" << T::name << f.mergedGroupIndex << f.fieldName << "Kernel(";
os << "__global struct Merged" << T::name << "Group" << f.mergedGroupIndex << " *group, unsigned int idx, ";
if(::Utils::isTypePointer(f.type)) {
if(GeNN::Utils::isTypePointer(f.type)) {
os << "__global ";
}
os << f.type << " " << f.fieldName << ")";
Expand Down Expand Up @@ -456,5 +447,4 @@ class BACKEND_EXPORT Backend : public BackendSIMT
cl::Device m_ChosenDevice;
cl::Platform m_ChosenPlatform;
};
} // OpenCL
} // CodeGenerator
} // GeNN::CodeGenerator::OpenCL
16 changes: 7 additions & 9 deletions include/genn/backends/opencl/optimiser.h
Original file line number Diff line number Diff line change
Expand Up @@ -7,25 +7,23 @@
#include "backend.h"

// Forward declarations
namespace GeNN
{
class ModelSpecInternal;
}

namespace plog
{
class IAppender;
}


//--------------------------------------------------------------------------
// CodeGenerator::OpenCL::Optimiser
// GeNN::CodeGenerator::OpenCL::Optimiser
//--------------------------------------------------------------------------
namespace CodeGenerator
{
namespace OpenCL
{
namespace Optimiser
namespace GeNN::CodeGenerator::OpenCL::Optimiser
{
BACKEND_EXPORT Backend createBackend(const ModelSpecInternal &model, const filesystem::path &outputPath,
plog::Severity backendLevel, plog::IAppender *backendAppender,
const Preferences &preferences);
} // namespace Optimiser
} // namespace CUDA
} // namespace CodeGenerator
} // namespace GeNN::CodeGenerator::OpenCL::Optimiser
12 changes: 3 additions & 9 deletions include/genn/backends/opencl/utils.h
Original file line number Diff line number Diff line change
Expand Up @@ -20,13 +20,9 @@
}

//--------------------------------------------------------------------------
// CodeGenerator::OpenCL::Utils
// GeNN::CodeGenerator::OpenCL::Utils
//--------------------------------------------------------------------------
namespace CodeGenerator
{
namespace OpenCL
{
namespace Utils
namespace GeNN::CodeGenerator::OpenCL::Utils
{
// OpenCL error string
const char *clGetErrorString(cl_int error)
Expand Down Expand Up @@ -100,6 +96,4 @@ const char *clGetErrorString(cl_int error)
}
#undef GEN_CL_ERROR_CASE
}
} // namespace Utils
} // namespace OpenCL
} // namespace CodeGenerator
} // namespace GeNN::Utils::OpenCL::CodeGenerator
Loading