From be5f1e4b948b484c405c0ca7e7692872bca7172e Mon Sep 17 00:00:00 2001 From: J Date: Mon, 13 Aug 2018 10:21:28 -0700 Subject: added cuda driver api --- linux-so-version.txt | 2 ++ 1 file changed, 2 insertions(+) (limited to 'linux-so-version.txt') diff --git a/linux-so-version.txt b/linux-so-version.txt index 40f775d..c6e03e9 100644 --- a/linux-so-version.txt +++ b/linux-so-version.txt @@ -1,2 +1,4 @@ libcudart.so.9.1{ }; +libcuda.so.1{ +}; -- cgit v1.3 From 37a709bb69fcd9f2b0fe53a189e92e548164cc4b Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Wed, 5 Jun 2019 19:18:29 -0400 Subject: adding new cuda 9 APIs to run the deepbench workloads --- Makefile | 2 ++ libcuda/cuda_runtime_api.cc | 70 +++++++++++++++++++++++++++++++++++++++++- linux-so-version.txt | 4 +++ src/abstract_hardware_model.cc | 22 +++++++++++++ src/abstract_hardware_model.h | 1 + src/gpgpu-sim/gpu-sim.cc | 10 ++++++ src/gpgpu-sim/gpu-sim.h | 2 ++ 7 files changed, 110 insertions(+), 1 deletion(-) (limited to 'linux-so-version.txt') diff --git a/Makefile b/Makefile index 2d0466e..e1e9aaa 100644 --- a/Makefile +++ b/Makefile @@ -159,10 +159,12 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib if [ ! -f $(SIM_LIB_DIR)/libcudart.so.5.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.5.5; fi if [ ! -f $(SIM_LIB_DIR)/libcudart.so.6.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.6.0; fi if [ ! -f $(SIM_LIB_DIR)/libcudart.so.6.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.6.5; fi + if [ ! -f $(SIM_LIB_DIR)/libcudart.so.7.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.7.0; fi if [ ! -f $(SIM_LIB_DIR)/libcudart.so.7.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.7.5; fi if [ ! -f $(SIM_LIB_DIR)/libcudart.so.8.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.8.0; fi if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.0; fi if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.1; fi + if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.2 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.2; fi $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\ diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 3a9d613..2cc84aa 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -974,6 +974,30 @@ __host__ cudaError_t CUDARTAPI cudaGetDeviceProperties(struct cudaDeviceProp *pr } } +#if (CUDART_VERSION >= 8000) +cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks, const char *hostFunc, int blockSize, size_t dynamicSMemSize, unsigned int flags) +{ + printf("GPGPU-Sim PTX: cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags %p\n", hostFunc); + CUctx_st *context = GPGPUSim_Context(); + function_info *entry = context->get_kernel(hostFunc); + printf("Calculate Maxium Active Block with function ptr=%p, blockSize=%d, SMemSize=%d\n", hostFunc, blockSize, dynamicSMemSize); + if (flags == cudaOccupancyDefault) { + //create kernel_info based on entry + dim3 gridDim(context->get_device()->get_gpgpu()->max_cta_per_core() + * context->get_device()->get_gpgpu()->get_config().num_shader()); + dim3 blockDim(blockSize); + kernel_info_t result(gridDim, blockDim, entry); + *numBlocks = context->get_device()->get_gpgpu()->get_max_cta(result); + printf("Maximum size is %d with gridDim %d and blockDim %d\n", *numBlocks, gridDim.x, blockDim.x); + return g_last_cudaError = cudaSuccess; + } else { + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; + } +} + +#endif + #if (CUDART_VERSION > 5000) __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device) { @@ -3169,6 +3193,23 @@ __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t v return g_last_cudaError = cudaSuccess; } +#if CUDART_VERSION >= 9000 +__host__ cudaError_t cudaFuncSetAttribute ( const void* func, enum cudaFuncAttribute attr, int value ) { + + //ignore this Attribute for now, and the default is that carveout = cudaSharedmemCarveoutDefault; // (-1) + return g_last_cudaError = cudaSuccess; +} + +#if CUDART_VERSION >= 9020 +__host__ __device__ unsigned __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim, + size_t sharedMem = 0, + void *stream = 0) { + + return 1; +} +#endif + +#endif #endif @@ -4071,6 +4112,33 @@ CUresult CUDAAPI cuMemHostRegister(void *p, size_t bytesize, unsigned int Flags) printf("WARNING: this function has not been implemented yet."); return CUDA_SUCCESS; } +__host__ cudaError_t cudaHostRegister(void* ptr, size_t size, unsigned int flags) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + printf("WARNING: this function has not been implemented yet."); + return g_last_cudaError = cudaSuccess; +} + +__host__ cudaError_t cudaProfilerStart ( ) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + printf("WARNING: this function has not been implemented yet."); + return g_last_cudaError = cudaSuccess; +} + +__host__ cudaError_t cudaProfilerStop ( ) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + printf("WARNING: this function has not been implemented yet."); + return g_last_cudaError = cudaSuccess; +} + #endif #if CUDART_VERSION >= 4000 @@ -4952,7 +5020,7 @@ CUresult CUDAAPI cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int *numBl printf("WARNING: this function has not been implemented yet."); return CUDA_SUCCESS; } - + CUresult CUDAAPI cuOccupancyMaxPotentialBlockSize(int *minGridSize, int *blockSize, CUfunction func, CUoccupancyB2DSize blockSizeToDynamicSMemSize, size_t dynamicSMemSize, int blockSizeLimit) { if(g_debug_execution >= 3){ diff --git a/linux-so-version.txt b/linux-so-version.txt index c6e03e9..a7c2d3c 100644 --- a/linux-so-version.txt +++ b/linux-so-version.txt @@ -1,4 +1,8 @@ +libcudart.so.9.0{ +}; libcudart.so.9.1{ }; +libcudart.so.9.2{ +}; libcuda.so.1{ }; diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index cebdb25..023f51b 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -691,6 +691,28 @@ unsigned g_kernel_launch_latency; unsigned kernel_info_t::m_next_uid = 1; +kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry) +{ + m_kernel_entry=entry; + m_grid_dim=gridDim; + m_block_dim=blockDim; + m_next_cta.x=0; + m_next_cta.y=0; + m_next_cta.z=0; + m_next_tid=m_next_cta; + m_num_cores_running=0; + m_uid = m_next_uid++; + m_param_mem = new memory_space_impl<8192>("param",64*1024); + + //Jin: parent and child kernel management for CDP + m_parent_kernel = NULL; + + //Jin: launch latency management + m_launch_latency = g_kernel_launch_latency; + + volta_cache_config_set=false; +} + /*A snapshot of the texture mappings needs to be stored in the kernel's info as kernels should use the texture bindings seen at the time of launch and textures can be bound/unbound asynchronously with respect to streams. */ diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 22ef509..64bbaa2 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -212,6 +212,7 @@ public: // m_num_cores_running=0; // m_param_mem=NULL; // } + kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry); kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTextureInfo); ~kernel_info_t(); diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index c1ba934..6de5845 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -806,6 +806,16 @@ int gpgpu_sim::shader_clock() const return m_config.core_freq/1000; } +int gpgpu_sim::max_cta_per_core() const +{ + return m_shader_config->max_cta_per_core; +} + +int gpgpu_sim::get_max_cta( const kernel_info_t &k ) const +{ + return m_shader_config->max_cta(k); +} + void gpgpu_sim::set_prop( cudaDeviceProp *prop ) { m_cuda_properties = prop; diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index c8dad89..7336cac 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -455,6 +455,8 @@ public: int num_registers_per_block() const; int wrp_size() const; int shader_clock() const; + int max_cta_per_core() const; + int get_max_cta( const kernel_info_t &k ) const; const struct cudaDeviceProp *get_prop() const; enum divergence_support_t simd_model() const; -- cgit v1.3 From ce1fabd955ee208a38e73d5a5acfba2b26369a3d Mon Sep 17 00:00:00 2001 From: tgrogers Date: Sun, 9 Jun 2019 22:20:15 -0400 Subject: A bunch of boilerbplate to get 10.1 to compile. Still does not yet run. The way CUDA calls kerenels (even on old code) has changed. --- Makefile | 2 ++ libcuda/cuda_api.h | 2 ++ libcuda/cuda_runtime_api.cc | 23 +++++++++++++++++++++++ linux-so-version.txt | 4 ++++ setup_environment | 2 +- 5 files changed, 32 insertions(+), 1 deletion(-) (limited to 'linux-so-version.txt') diff --git a/Makefile b/Makefile index e1e9aaa..7a3bb47 100644 --- a/Makefile +++ b/Makefile @@ -165,6 +165,8 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.0; fi if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.1; fi if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.2 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.2; fi + if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.0; fi + if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.1; fi $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\ diff --git a/libcuda/cuda_api.h b/libcuda/cuda_api.h index 3808e8a..7ee26dc 100644 --- a/libcuda/cuda_api.h +++ b/libcuda/cuda_api.h @@ -234,9 +234,11 @@ typedef struct CUgraphicsResource_st *CUgraphicsResource; /**< CUDA graphics int typedef unsigned long long CUtexObject; /**< An opaque value that represents a CUDA texture object */ typedef unsigned long long CUsurfObject; /**< An opaque value that represents a CUDA surface object */ +#if __CUDA_API_VERSION < 1010 typedef struct CUuuid_st { /**< CUDA definition of UUID */ char bytes[16]; } CUuuid; +#endif #if __CUDA_API_VERSION >= 4010 diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 2cc84aa..fb3e07a 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -2763,6 +2763,29 @@ cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ return g_last_cudaError = cudaSuccess; } +void CUDARTAPI __cudaRegisterFatBinaryEnd( void **fatCubinHandle ) +{ + +} + +unsigned CUDARTAPI __cudaPushCallConfiguration(dim3 gridDim, + dim3 blockDim, + size_t sharedMem = 0, + struct CUstream_st *stream = 0) +{ + +} + +cudaError_t CUDARTAPI __cudaPopCallConfiguration( + dim3 *gridDim, + dim3 *blockDim, + size_t *sharedMem, + void *stream +) +{ + return g_last_cudaError = cudaSuccess; +} + void CUDARTAPI __cudaRegisterFunction( void **fatCubinHandle, const char *hostFun, diff --git a/linux-so-version.txt b/linux-so-version.txt index a7c2d3c..45c40dd 100644 --- a/linux-so-version.txt +++ b/linux-so-version.txt @@ -4,5 +4,9 @@ libcudart.so.9.1{ }; libcudart.so.9.2{ }; +libcudart.so.10.0{ +}; +libcudart.so.10.1{ +}; libcuda.so.1{ }; diff --git a/setup_environment b/setup_environment index b420584..17891ce 100644 --- a/setup_environment +++ b/setup_environment @@ -51,7 +51,7 @@ CC_VERSION=`gcc --version | head -1 | awk '{for(i=1;i<=NF;i++){ if(match($i,/^[0 CUDA_VERSION_STRING=`$CUDA_INSTALL_PATH/bin/nvcc --version | awk '/release/ {print $5;}' | sed 's/,//'`; export CUDA_VERSION_NUMBER=`echo $CUDA_VERSION_STRING | sed 's/\./ /' | awk '{printf("%02u%02u", 10*int($1), 10*$2);}'` -if [ $CUDA_VERSION_NUMBER -gt 9100 -o $CUDA_VERSION_NUMBER -lt 2030 ]; then +if [ $CUDA_VERSION_NUMBER -gt 10100 -o $CUDA_VERSION_NUMBER -lt 2030 ]; then echo "ERROR ** GPGPU-Sim version $GPGPUSIM_VERSION_STRING not tested with CUDA version $CUDA_VERSION_STRING (please see README)"; return fi -- cgit v1.3