opencl/backend.h

Go to the documentation of this file.

#pragma once

// Standard C++ includes
#include <algorithm>
#include <array>
#include <functional>
#include <map>
#include <numeric>
#include <string>

// OpenCL includes
#include "../../../../share/genn/backends/opencl/cl2.hpp"

// GeNN includes
#include "backendExport.h"
#include "gennUtils.h"

// GeNN code generator includes
#include "code_generator/backendSIMT.h"
#include "code_generator/codeStream.h"
#include "code_generator/modelSpecMerged.h"
#include "code_generator/substitutions.h"

// Forward declarations
namespace filesystem
{
    class path;
}

//--------------------------------------------------------------------------
// CodeGenerator::OpenCL::DeviceSelectMethod
//--------------------------------------------------------------------------
namespace CodeGenerator
{
namespace OpenCL
{
enum class PlatformSelect
{
    MANUAL,         
};

enum class DeviceSelect
{
    MOST_MEMORY,    
    MANUAL,         
};

//--------------------------------------------------------------------------
// CodeGenerator::OpenCL::WorkGroupSizeSelect
//--------------------------------------------------------------------------
enum class WorkGroupSizeSelect
{
    MANUAL,     
};

//--------------------------------------------------------------------------
// CodeGenerator::OpenCL::Preferences
//--------------------------------------------------------------------------
struct Preferences : public PreferencesBase
{
    Preferences()
    {
        std::fill(manualWorkGroupSizes.begin(), manualWorkGroupSizes.end(), 32);
    }

    PlatformSelect platformSelectMethod = PlatformSelect::MANUAL;

    unsigned int manualPlatformID = 0;

    DeviceSelect deviceSelectMethod = DeviceSelect::MOST_MEMORY;

    unsigned int manualDeviceID = 0;

    WorkGroupSizeSelect workGroupSizeSelectMethod = WorkGroupSizeSelect::MANUAL;

    KernelBlockSize manualWorkGroupSizes;
};

//--------------------------------------------------------------------------
// CodeGenerator::OpenCL::Backend
//--------------------------------------------------------------------------
class BACKEND_EXPORT Backend : public BackendSIMT
{
public:
    Backend(const KernelBlockSize &kernelWorkGroupSizes, const Preferences &preferences,
            const std::string &scalarType, unsigned int platformIndex, unsigned int deviceIndex);

    //--------------------------------------------------------------------------
    // CodeGenerator::BackendSIMT virtuals
    //--------------------------------------------------------------------------
    virtual bool areSharedMemAtomicsSlow() const override;

    virtual std::string getSharedPrefix() const override { return "__local "; }

    virtual std::string getThreadID(unsigned int axis = 0) const override{ return "get_local_id(" + std::to_string(axis) + ")"; }

    virtual std::string getBlockID(unsigned int axis = 0) const override{ return "get_group_id(" + std::to_string(axis) + ")"; }

    virtual std::string getCLZ() const override { return "clz"; }

    virtual std::string getAtomic(const std::string &type, AtomicOperation op = AtomicOperation::ADD,
                                  AtomicMemSpace memSpace = AtomicMemSpace::GLOBAL) const override;

    virtual void genSharedMemBarrier(CodeStream &os) const override;

    virtual void genPopulationRNGInit(CodeStream &os, const std::string &globalRNG, const std::string &seed, const std::string &sequence) const override;

    virtual void genPopulationRNGPreamble(CodeStream &os, Substitutions &subs, const std::string &globalRNG, const std::string &name = "rng") const override;


    virtual void genPopulationRNGPostamble(CodeStream &os, const std::string &globalRNG) const override;

    virtual void genGlobalRNGSkipAhead(CodeStream &os, Substitutions &subs, const std::string &sequence, const std::string &name = "rng") const override;

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

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

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

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

    virtual void genDefinitionsPreamble(CodeStream &os, const ModelSpecMerged &modelMerged) const override;
    virtual void genDefinitionsInternalPreamble(CodeStream &os, const ModelSpecMerged &modelMerged) const override;
    virtual void genRunnerPreamble(CodeStream &os, const ModelSpecMerged &modelMerged, const MemAlloc &memAlloc) const override;
    virtual void genAllocateMemPreamble(CodeStream &os, const ModelSpecMerged &modelMerged, const MemAlloc &allocations) const override;
    virtual void genFreeMemPreamble(CodeStream &os, const ModelSpecMerged &modelMerged) const override;
    virtual void genStepTimeFinalisePreamble(CodeStream &os, const ModelSpecMerged &modelMerged) const override;

    virtual void genVariableDefinition(CodeStream &definitions, CodeStream &definitionsInternal, const std::string &type, const std::string &name, VarLocation loc) const override;
    virtual void genVariableImplementation(CodeStream &os, const std::string &type, const std::string &name, VarLocation loc) const override;
    virtual void genVariableAllocation(CodeStream &os, const std::string &type, const std::string &name, VarLocation loc, size_t count, MemAlloc &memAlloc) const override;
    virtual void genVariableFree(CodeStream &os, const std::string &name, VarLocation loc) const override;

    virtual void genExtraGlobalParamDefinition(CodeStream &definitions, CodeStream &definitionsInternal, const std::string &type, const std::string &name, VarLocation loc) const override;
    virtual void genExtraGlobalParamImplementation(CodeStream &os, const std::string &type, const std::string &name, VarLocation loc) const override;
    virtual void genExtraGlobalParamAllocation(CodeStream &os, const std::string &type, const std::string &name,
                                               VarLocation loc, const std::string &countVarName = "count", const std::string &prefix = "") const override;
    virtual void genExtraGlobalParamPush(CodeStream &os, const std::string &type, const std::string &name,
                                         VarLocation loc, const std::string &countVarName = "count", const std::string &prefix = "") const override;
    virtual void genExtraGlobalParamPull(CodeStream &os, const std::string &type, const std::string &name,
                                         VarLocation loc, const std::string &countVarName = "count", const std::string &prefix = "") const override;

    virtual void genMergedExtraGlobalParamPush(CodeStream &os, const std::string &suffix, size_t mergedGroupIdx,
                                               const std::string &groupIdx, const std::string &fieldName,
                                               const std::string &egpName) const override;

    virtual std::string getMergedGroupFieldHostType(const std::string &type) const override;

    virtual std::string getMergedGroupSimRNGType() const override { return "clrngLfsr113HostStream"; }

    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;

    virtual void genCurrentVariablePush(CodeStream &os, const NeuronGroupInternal &ng, const std::string &type,
                                        const std::string &name, VarLocation loc, unsigned int batchSize) const override;
    virtual void genCurrentVariablePull(CodeStream &os, const NeuronGroupInternal &ng, const std::string &type,
                                        const std::string &name, VarLocation loc, unsigned int batchSize) const override;

    virtual void genCurrentTrueSpikePush(CodeStream &os, const NeuronGroupInternal &ng, unsigned int batchSize) const override
    {
        genCurrentSpikePush(os, ng, batchSize, false);
    }
    virtual void genCurrentTrueSpikePull(CodeStream &os, const NeuronGroupInternal &ng, unsigned int batchSize) const override
    {
        genCurrentSpikePull(os, ng, batchSize, false);
    }
    virtual void genCurrentSpikeLikeEventPush(CodeStream &os, const NeuronGroupInternal &ng, unsigned int batchSize) const override
    {
        genCurrentSpikePush(os, ng, batchSize, true);
    }
    virtual void genCurrentSpikeLikeEventPull(CodeStream &os, const NeuronGroupInternal &ng, unsigned int batchSize) const override
    {
        genCurrentSpikePull(os, ng, batchSize, true);
    }

    virtual void genGlobalDeviceRNG(CodeStream &definitions, CodeStream &definitionsInternal, CodeStream &runner,
                                    CodeStream &allocations, CodeStream &free, MemAlloc &memAlloc) const override;
    virtual void genPopulationRNG(CodeStream &definitions, CodeStream &definitionsInternal, CodeStream &runner, CodeStream &allocations,
                                  CodeStream &free, const std::string &name, size_t count, MemAlloc &memAlloc) const override;
    virtual void genTimer(CodeStream &definitions, CodeStream &definitionsInternal, CodeStream &runner,
                          CodeStream &allocations, CodeStream &free, CodeStream &stepTimeFinalise,
                          const std::string &name, bool updateInStepTime) const override;

    virtual void genReturnFreeDeviceMemoryBytes(CodeStream &os) const override;

    virtual void genMakefilePreamble(std::ostream &os) const override;
    virtual void genMakefileLinkRule(std::ostream &os) const override;
    virtual void genMakefileCompileRule(std::ostream &os) const override;

    virtual void genMSBuildConfigProperties(std::ostream &os) const override;
    virtual void genMSBuildImportProps(std::ostream &os) const override;
    virtual void genMSBuildItemDefinitions(std::ostream &os) const override;
    virtual void genMSBuildCompileModule(const std::string &moduleName, std::ostream &os) const override;
    virtual void genMSBuildImportTarget(std::ostream &os) const override;


    virtual std::vector<filesystem::path> getFilesToCopy(const ModelSpecMerged &modelMerged) const override;

    virtual std::string getHostVarPrefix() const final { return "h_"; }

    virtual std::string getPointerPrefix() const override { return "__global "; };

    virtual bool isPopulationRNGInitialisedOnDevice() const override { return false; }

    virtual bool isHostReductionRequired() const override { return false; }

    virtual size_t getDeviceMemoryBytes() const override { return m_ChosenDevice.getInfo<CL_DEVICE_GLOBAL_MEM_SIZE>(); }
    
    virtual MemorySpaces getMergedGroupMemorySpaces(const ModelSpecMerged &modelMerged) const override;

    virtual bool supportsNamespace() const override { return false; };
    
    virtual boost::uuids::detail::sha1::digest_type getHashDigest() const override;
    
    //--------------------------------------------------------------------------
    // Public API
    //--------------------------------------------------------------------------
    const cl::Device &getChosenOpenCLDevice() const { return m_ChosenDevice; }
    
    std::string getFloatAtomicAdd(const std::string &ftype, const char* memoryType = "global") const;

private:
    //--------------------------------------------------------------------------
    // Private methods
    //--------------------------------------------------------------------------
    template<typename T>
    void genMergedStructPreamble(CodeStream &os, const ModelSpecMerged &modelMerged, const std::vector<T> &groups) const
    {
        // Loop through groups
        for(const auto &g : groups) {
            // Check there's no memory space assigned as OpenCL doesn't currently support them
            assert(g.getMemorySpace().empty());

            // Declare build kernel
            const std::string buildKernelName = "build" + T::name + std::to_string(g.getIndex()) + "Kernel";
            os << "cl::Kernel " << buildKernelName << ";" << std::endl;
            
            // Declare buffer
            os << "cl::Buffer d_merged" << T::name << "Group" << g.getIndex() << ";" << std::endl;

            // Write function to update
            os << "void pushMerged" << T::name << "Group" << g.getIndex() << "ToDevice(unsigned int idx, ";
            g.generateStructFieldArgumentDefinitions(os, *this);
            os << ")";
            {
                CodeStream::Scope b(os);

                // Add idx parameter
                os << "CHECK_OPENCL_ERRORS(" << buildKernelName << ".setArg(1, idx));" << std::endl;

                // Loop through sorted fields and add arguments
                const auto sortedFields = g.getSortedFields(*this);
                for(size_t fieldIndex = 0; fieldIndex < sortedFields.size(); fieldIndex++) {
                    const auto &f = sortedFields[fieldIndex];

                    os << "CHECK_OPENCL_ERRORS(" << buildKernelName << ".setArg(" << (2 + fieldIndex) << ", " << std::get<1>(f) << "));" << std::endl;
                }
                
                // Launch kernel
                os << "const cl::NDRange globalWorkSize(1, 1);" << std::endl;
                os << "const cl::NDRange localWorkSize(1, 1);" << std::endl;
                os << "CHECK_OPENCL_ERRORS(commandQueue.enqueueNDRangeKernel(" << buildKernelName << ", cl::NullRange, globalWorkSize, localWorkSize));" << std::endl;
            }
        }

        if(!groups.empty()) {
            // Get set of unique fields referenced in a merged group
            const auto mergedGroupFields = modelMerged.getMergedGroupFields<T>();

            // Loop through resultant fields and declare kernel for setting EGP
            for(auto f : mergedGroupFields) {
                os << "cl::Kernel setMerged" << T::name << f.mergedGroupIndex << f.fieldName << "Kernel;" << std::endl;
            }
        }
    }

    template<typename T>
    void genMergedStructBuild(CodeStream &os, const ModelSpecMerged &modelMerged, const std::vector<T> &groups, const std::string &programName) const
    {
        // Loop through groups 
        for(const auto &g : groups) {
            // Create kernel object
            const std::string kernelName = "build" + T::name + std::to_string(g.getIndex()) + "Kernel";
            os << "CHECK_OPENCL_ERRORS_POINTER(" << kernelName << " = cl::Kernel(" << programName << ", \"" << kernelName << "\", &error));" << std::endl;

            // Create group buffer
            os << "CHECK_OPENCL_ERRORS_POINTER(d_merged" << T::name << "Group" << g.getIndex() << " = cl::Buffer(clContext, CL_MEM_READ_WRITE, size_t{" << g.getStructArraySize(*this) << "}, nullptr, &error));" << std::endl;

            // Set group buffer as first kernel argument
            os << "CHECK_OPENCL_ERRORS(" << kernelName << ".setArg(0, d_merged" << T::name << "Group" << g.getIndex() << "));" << std::endl;
            os << std::endl;
        }

        if(!groups.empty()) {
            // Get set of unique fields referenced in a merged group
            const auto mergedGroupFields = modelMerged.getMergedGroupFields<T>();

            // Loop through resultant fields
            for(auto f : mergedGroupFields) {
                // Create kernel object
                const std::string kernelName = "setMerged" + T::name + std::to_string(f.mergedGroupIndex) + f.fieldName + "Kernel";
                os << "CHECK_OPENCL_ERRORS_POINTER(" << kernelName << " = cl::Kernel(" << programName << ", \"" << kernelName << "\", &error));" << std::endl;

                // Set group buffer as first kernel argument
                os << "CHECK_OPENCL_ERRORS(" << kernelName << ".setArg(0, d_merged" << T::name << "Group" << f.mergedGroupIndex << "));" << std::endl;
                os << std::endl;
            }
        }
    }

    template<typename T>
    void genMergedStructBuildKernels(CodeStream &os, const ModelSpecMerged &modelMerged, const std::vector<T> &groups) const
    {
        // Loop through groups
        for(const auto &g : groups) {
            // Generate kernel to build struct on device
            os << "__kernel void build" << T::name << g.getIndex() << "Kernel(";
            os << "__global struct Merged" << T::name << "Group" << g.getIndex() << " *group, unsigned int idx, ";

            // Loop through sorted struct fields
            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))) {
                    os << "__global ";
                }
                os << std::get<0>(f) << " " << std::get<1>(f);
                if(fieldIndex != (sortedFields.size() - 1)) {
                    os << ", ";
                }
            }
            os << ")";
            {
                CodeStream::Scope b(os);

                // Assign all structure fields to values passed through parameters
                for(const auto &f : sortedFields) {
                    os << "group[idx]." << std::get<1>(f) << " = " << std::get<1>(f) << ";" << std::endl;
                }
            }
            os << std::endl;
        }

        if(!groups.empty()) {
            // Get set of unique fields referenced in a merged group
            const auto mergedGroupFields = modelMerged.getMergedGroupFields<T>();

            // Loop through resultant fields and generate push function for pointer extra global parameters
            for(auto f : mergedGroupFields) {
                
                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)) {
                    os << "__global ";
                }
                os << f.type << " " << f.fieldName << ")";
                {
                    CodeStream::Scope b(os);
                    os << "group[idx]." << f.fieldName << " = " << f.fieldName << ";" << std::endl;
                }
                os << std::endl;
            }
        }
    }

    void genAtomicAddFloat(CodeStream &os, const std::string &memoryType) const;

    void genCurrentSpikePush(CodeStream &os, const NeuronGroupInternal &ng, unsigned int batchSize, bool spikeEvent) const
    {
        genCurrentSpikePushPull(os, ng, batchSize, spikeEvent, true);
    }

    void genCurrentSpikePull(CodeStream &os, const NeuronGroupInternal &ng, unsigned int batchSize, bool spikeEvent) const
    {
        genCurrentSpikePushPull(os, ng, batchSize, spikeEvent, false);
    }

    void genCurrentSpikePushPull(CodeStream &os, const NeuronGroupInternal &ng, unsigned int batchSize, bool spikeEvent, bool push) const;

    void genKernelDimensions(CodeStream &os, Kernel kernel, size_t numThreadsX, size_t batchSize, size_t numBlockThreadsY = 1) const;

    void genKernelPreamble(CodeStream &os, const ModelSpecMerged &modelMerged) const;

    void genBuildProgramFlagsString(CodeStream &os) const;

    void divideKernelStreamInParts(CodeStream &os, const std::stringstream &kernelCode, size_t partLength) const;

    bool isChosenDeviceAMD() const;

    bool isChosenDeviceNVIDIA() const;


    bool isChosenPlatformNVIDIA() const;


    bool shouldUseSubBufferAllocations() const;

    //--------------------------------------------------------------------------
    // Members
    //--------------------------------------------------------------------------
    const unsigned int m_ChosenPlatformIndex;
    const unsigned int m_ChosenDeviceIndex;
    unsigned int m_AllocationAlignementBytes;
    cl::Device m_ChosenDevice;
    cl::Platform m_ChosenPlatform;
};
}   // OpenCL
}   // CodeGenerator