Skip to content

Commit

Permalink
Merge pull request #595 from genn-team/transpiler
Browse files Browse the repository at this point in the history
Step 1 - Transpiler
  • Loading branch information
neworderofjamie authored Jan 9, 2024
2 parents 5e5e6a5 + ca12c31 commit a1ab0d6
Show file tree
Hide file tree
Showing 603 changed files with 23,432 additions and 30,504 deletions.
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

0 comments on commit a1ab0d6

Please sign in to comment.