tf_cuda.h

Go to the documentation of this file.

/*******************************************************************************
 * This file is part of Tissue Forge.
 * Copyright (c) 2022-2024 T.J. Sego
 * 
 * This program is free software: you can redistribute it and/or modify
 * it under the terms of the GNU Lesser General Public License as published
 * by the Free Software Foundation, either version 3 of the License, or
 * (at your option) any later version.
 * 
 * This program is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
 * GNU General Public License for more details.
 * 
 * You should have received a copy of the GNU Lesser General Public License
 * along with this program.  If not, see <http://www.gnu.org/licenses/>.
 * 
 ******************************************************************************/
 
// TODO: implement support for JIT-compiled programs and kernel usage in wrapped languages
 
#ifndef _SOURCE_TF_CUDA_H_
#define _SOURCE_TF_CUDA_H_
 
#include "tfError.h"
 
#include <cuda.h>
#include <cuda_runtime.h>
#include <nvrtc.h>
 
#include <vector>
#include <stdexcept>
#include <string>
 
 
namespace TissueForge::cuda {
 
 
    enum ErrorCode : int {
        TFCUDAERR_ok = 0,
        TFCUDAERR_setdevice,
        TFCUDAERR_setblocks,
        TFCUDAERR_setthreads,
        TFCUDAERR_ondevice,
        TFCUDAERR_notondevice,
        TFCUDAERR_cleardevices,
        TFCUDAERR_refresh,
        TFCUDAERR_send,
        TFCUDAERR_pull,
        TFCUDAERR_LAST
    };
 
    /* list of error messages. */
    static const char *tfcuda_err_msg[TFCUDAERR_LAST] = {
        "No CUDA errors.",
        "Failed to set device.",
        "Failed to set blocks.",
        "Failed to set threads.",
        "Already on device.",
        "Not on device.",
        "Failed to clear devices.",
        "Refresh failed.",
        "Attempting send to device failed.",
        "Attempting pull from device when not sent."
    };
 
    inline CUresult cuda_errorchk(CUresult retCode, const char *file, int line) {
        if(retCode != CUDA_SUCCESS) {
            std::string msg = "CUDA failed with error: ";
            const char *cmsg;
            cuGetErrorName(retCode, &cmsg);
            msg += std::string(cmsg);
            msg += ", " + std::string(file) + ", " + std::to_string(line);
            tf_exp(std::runtime_error(msg.c_str()));
        }
        return retCode;
    }
    #ifndef TF_CUDA_CALL
        #define TF_CUDA_CALL(res) cuda_errorchk(res, __FILE__, __LINE__)
    #endif
 
    inline nvrtcResult nvrtc_errorchk(nvrtcResult retCode, const char *file, int line) {
        if(retCode != NVRTC_SUCCESS) {
            std::string msg = "NVRTC failed with error: ";
            msg += std::string(nvrtcGetErrorString(retCode));
            msg += ", " + std::string(file) + ", " + std::to_string(line);
            tf_exp(std::runtime_error(msg.c_str()));
        }
        return retCode;
    }
    #ifndef TF_NVRTC_CALL
        #define TF_NVRTC_CALL(res) nvrtc_errorchk(res, __FILE__, __LINE__)
    #endif
 
    inline cudaError_t cudart_errorchk(cudaError_t retCode, const char *file, int line) {
        if(retCode != cudaSuccess) {
            std::string msg = "NVRTC failed with error: ";
            msg += std::string(cudaGetErrorString(retCode));
            msg += ", " + std::string(file) + ", " + std::to_string(line);
            tf_exp(std::runtime_error(msg.c_str()));
        }
        return retCode;
    }
    #ifndef TF_CUDART_CALL
        #define TF_CUDART_CALL(res) cudart_errorchk(res, __FILE__, __LINE__)
    #endif
 
 
    struct CUDARTSource {
        std::string source;
        const char *name;
 
        CUDARTSource(const char *filePath, const char *_name);
        const char *c_str() const;
    };
 
 
    struct CUDARTProgram {
 
        nvrtcProgram *prog;
        char *ptx;
        std::vector<std::string> opts;
        std::vector<std::string> namedExprs;
        std::vector<std::string> includePaths;
        int arch;
        bool is_compute;
 
        CUDARTProgram();
        ~CUDARTProgram();
 
        void addOpt(const std::string &opt);
 
        void addIncludePath(const std::string &ipath);
 
        void addNamedExpr(const std::string &namedExpr);
 
        void compile(const char *src, const char *name, int numHeaders=0, const char *const *headers=0, const char *const *includeNames=0);
 
        std::string loweredName(const std::string namedExpr);
 
    };
 
    struct CUDAContext;
 
 
    struct CUDAFunction {
        const std::string name;
        unsigned int gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes;
        CUstream hStream;
        void **extra;
 
        CUDAFunction(const std::string &name, CUDAContext *context);
        ~CUDAFunction();
 
        HRESULT autoConfig(const unsigned int &_nr_arrayElems, 
                        size_t dynamicSMemSize=0, 
                        size_t (*blockSizeToDynamicSMemSize)(int)=0, 
                        int blockSizeLimit=0);
 
        void operator()(void **args);
        void operator()(int nargs, ...);
 
    private:
        CUfunction *function;
        CUDAContext *context;
    };
 
 
    struct CUDAContext {
 
        CUcontext *context;
        CUdevice device;
        CUmodule *module;
        std::vector<CUjit_option> compileOpts;
        std::vector<void*> compileOptVals;
 
        /* Flag signifying whether this context is attached to a CPU thread. */
        bool attached;
 
        CUDAContext(CUdevice device=0);
        ~CUDAContext();
 
        void addOpt(CUjit_option opt, void *val);
 
        void loadProgram(const CUDARTProgram &prog);
 
        void loadPTX(const char *ptx);
 
        CUDAFunction *getFunction(const char *name);
 
        CUdeviceptr *getGlobal(const char *name);
 
        size_t getGlobalSize(const char *name);
 
        void pushCurrent();
 
        CUcontext *popCurrent();
 
        void destroy();
 
        int getAPIVersion();
 
        static void sync();
    };
 
 
    struct CUDADevice {
 
        CUdevice *device;
 
        CUDADevice();
        ~CUDADevice();
        
        void attachDevice(const int &deviceId=0);
 
        void detachDevice();
 
        std::string name();
 
        int arch();
 
        size_t totalMem();
 
        int getAttribute(const int &attrib);
 
        std::string PCIBusId();
 
        CUDAContext *createContext();
 
        CUDAContext *currentContext();
 
        static std::string getDeviceName(const int &deviceId);
 
        static size_t getDeviceTotalMem(const int &deviceId);
 
        static int getDeviceAttribute(const int &deviceId, const int &attrib);
 
        static int getNumDevices();
 
        static std::string getDevicePCIBusId(const int &deviceId);
 
        static int getCurrentDevice();
 
        int maxThreadsPerBlock();
 
        int maxBlockDimX();
 
        int maxBlockDimY();
 
        int maxBlockDimZ();
 
        int maxGridDimX();
 
        int maxGridDimY();
 
        int maxGridDimZ();
 
        int maxSharedMemPerBlock();
 
        int maxTotalMemConst();
 
        int warpSize();
 
        int maxRegsPerBlock();
 
        int clockRate();
 
        bool gpuOverlap();
 
        int numMultiprocessors();
 
        bool kernelExecTimeout();
 
        bool computeModeDefault();
 
        bool computeModeProhibited();
 
        bool computeModeExclusive();
 
        int PCIDeviceId();
 
        int PCIDomainId();
 
        int clockRateMem();
 
        int globalMemBusWidth();
 
        int L2CacheSize();
 
        int maxThreadsPerMultiprocessor();
 
        int computeCapabilityMajor();
 
        int computeCapabilityMinor();
 
        bool L1CacheSupportGlobal();
 
        bool L1CacheSupportLocal();
 
        int maxSharedMemPerMultiprocessor();
 
        int maxRegsPerMultiprocessor();
 
        bool managedMem();
 
        bool multiGPUBoard();
 
        int multiGPUBoardGroupId();
 
    private:
        
        void validateAttached();
        static void validateDeviceId(const int &deviceId);
    };
 
 
    // Tissue Forge CUDA interface
 
 
    CPPAPI_FUNC(void) init();
 
    CPPAPI_FUNC(void) setGLDevice(const int &deviceId);
 
    CPPAPI_FUNC(std::string) getDeviceName(const int &deviceId);
 
    CPPAPI_FUNC(size_t) getDeviceTotalMem(const int &deviceId);
 
    CPPAPI_FUNC(int) getDeviceAttribute(const int &deviceId, const int &attrib);
 
    CPPAPI_FUNC(int) getNumDevices();
 
    CPPAPI_FUNC(std::string) getDevicePCIBusId(const int &deviceId);
 
    CPPAPI_FUNC(int) getCurrentDevice();
 
    CPPAPI_FUNC(int) maxThreadsPerBlock(const int &deviceId);
 
    CPPAPI_FUNC(int) maxBlockDimX(const int &deviceId);
 
    CPPAPI_FUNC(int) maxBlockDimY(const int &deviceId);
 
    CPPAPI_FUNC(int) maxBlockDimZ(const int &deviceId);
 
    CPPAPI_FUNC(int) maxGridDimX(const int &deviceId);
 
    CPPAPI_FUNC(int) maxGridDimY(const int &deviceId);
 
    CPPAPI_FUNC(int) maxGridDimZ(const int &deviceId);
 
    CPPAPI_FUNC(int) maxSharedMemPerBlock(const int &deviceId);
 
    CPPAPI_FUNC(int) maxTotalMemConst(const int &deviceId);
 
    CPPAPI_FUNC(int) warpSize(const int &deviceId);
 
    CPPAPI_FUNC(int) maxRegsPerBlock(const int &deviceId);
 
    CPPAPI_FUNC(int) clockRate(const int &deviceId);
 
    CPPAPI_FUNC(bool) gpuOverlap(const int &deviceId);
 
    CPPAPI_FUNC(int) numMultiprocessors(const int &deviceId);
 
    CPPAPI_FUNC(bool) kernelExecTimeout(const int &deviceId);
 
    CPPAPI_FUNC(bool) computeModeDefault(const int &deviceId);
 
    CPPAPI_FUNC(bool) computeModeProhibited(const int &deviceId);
 
    CPPAPI_FUNC(bool) computeModeExclusive(const int &deviceId);
 
    CPPAPI_FUNC(int) PCIDeviceId(const int &deviceId);
 
    CPPAPI_FUNC(int) PCIDomainId(const int &deviceId);
 
    CPPAPI_FUNC(int) clockRateMem(const int &deviceId);
 
    CPPAPI_FUNC(int) globalMemBusWidth(const int &deviceId);
 
    CPPAPI_FUNC(int) L2CacheSize(const int &deviceId);
 
    CPPAPI_FUNC(int) maxThreadsPerMultiprocessor(const int &deviceId);
 
    CPPAPI_FUNC(int) computeCapabilityMajor(const int &deviceId);
 
    CPPAPI_FUNC(int) computeCapabilityMinor(const int &deviceId);
 
    CPPAPI_FUNC(bool) L1CacheSupportGlobal(const int &deviceId);
 
    CPPAPI_FUNC(bool) L1CacheSupportLocal(const int &deviceId);
 
    CPPAPI_FUNC(int) maxSharedMemPerMultiprocessor(const int &deviceId);
 
    CPPAPI_FUNC(int) maxRegsPerMultiprocessor(const int &deviceId);
 
    CPPAPI_FUNC(bool) managedMem(const int &deviceId);
 
    CPPAPI_FUNC(bool) multiGPUBoard(const int &deviceId);
 
    CPPAPI_FUNC(int) multiGPUBoardGroupId(const int &deviceId);
 
    CPPAPI_FUNC(void) test(const int &numBlocks, const int &numThreads, const int &numEls, const int &deviceId=0);
 
 
    CPPAPI_FUNC(std::string) tfIncludePath();
    
    CPPAPI_FUNC(HRESULT) setTfIncludePath(const std::string &_path);
 
    CPPAPI_FUNC(std::string) tfPrivateIncludePath();
 
    CPPAPI_FUNC(std::string) tfResourcePath();
 
    CPPAPI_FUNC(HRESULT) setTfResourcePath(const std::string &_path);
 
    CPPAPI_FUNC(std::string) CUDAPath();
 
    CPPAPI_FUNC(std::string) CUDAIncludePath();
 
    CPPAPI_FUNC(HRESULT) setCUDAIncludePath(const std::string &_path);
 
    CPPAPI_FUNC(std::string) CUDAResourcePath(const std::string &relativePath);
 
    CPPAPI_FUNC(std::string) CUDAPTXObjectRelPath();
 
    CPPAPI_FUNC(std::vector<std::string>) CUDAArchs();
 
};
 
#endif // _SOURCE_TF_CUDA_H_