From 35c741383e04b900729f68707aca9ce63985f8b0 Mon Sep 17 00:00:00 2001 From: Jiangqiu Shen Date: Thu, 23 May 2019 14:34:58 -0400 Subject: fix bug when launch a empty kernel --- libcuda/cuda_runtime_api.cc | 11 +++++++++++ 1 file changed, 11 insertions(+) (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 3a9d613..44f0f4e 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1473,6 +1473,17 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) sscanf(mode,"%u", &g_ptx_sim_mode); gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); kernel_config config = g_cuda_launch_stack.back(); + { + dim3 gridDim = config.grid_dim(); + dim3 blockDim = config.block_dim(); + if (gridDim.x * gridDim.y * gridDim.z == 0 || blockDim.x * blockDim.y * blockDim.z == 0) + { + //can't launch + printf("can't launch a empty kernel\n"); + g_cuda_launch_stack.pop_back(); + return g_last_cudaError = cudaErrorInvalidConfiguration; + } + } struct CUstream_st *stream = config.get_stream(); printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun, g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 ); -- cgit v1.3 From c7f515f6f5325c65f32dd64e1ad479660c751e99 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 | 25 +++++++++++++++++++++++++ linux-so-version.txt | 4 ++++ setup_environment | 2 +- 5 files changed, 34 insertions(+), 1 deletion(-) (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/Makefile b/Makefile index 3db8ce8..a69130c 100644 --- a/Makefile +++ b/Makefile @@ -164,6 +164,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 18a9abb..718db49 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -630,6 +630,8 @@ void gpgpu_context::cuobjdumpRegisterFatBinary(unsigned int handle, const char* fatbinmap[handle] = filename; } + + /******************************************************************************* * Add internal cuda runtime API call to accept gpgpu_context * *******************************************************************************/ @@ -2975,6 +2977,29 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) return cudaRegisterFatBinaryInternal(fatCubin); } +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 From 2ea18072618e7fe4e541f84de7d8575998299a1c Mon Sep 17 00:00:00 2001 From: tgrogers Date: Mon, 10 Jun 2019 09:19:30 -0400 Subject: Code to get CUDA 10 to work - looks like the gridDim/blockDim args given to cudaLaunchKernel are not complete garbage. We must rely on the PushConfig to get the proper sizing info --- libcuda/cuda_runtime_api.cc | 43 +++++++++++++++++++++++++++---------------- 1 file changed, 27 insertions(+), 16 deletions(-) (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 718db49..500b5f2 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -2108,20 +2108,21 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) __host__ cudaError_t CUDARTAPI cudaLaunchKernel ( const char* hostFun, dim3 gridDim, dim3 blockDim, const void** args, size_t sharedMem, cudaStream_t stream ) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUctx_st *context = GPGPUSim_Context(); - function_info *entry = context->get_kernel(hostFun); - - cudaConfigureCallInternal(gridDim, blockDim, sharedMem, stream); - for(unsigned i = 0; i < entry->num_args(); i++){ - std::pair p = entry->get_param_config(i); - cudaSetupArgumentInternal(args[i], p.first, p.second); - } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + CUctx_st *context = GPGPUSim_Context(); + function_info *entry = context->get_kernel(hostFun); +#if CUDART_VERSION < 10000 + cudaConfigureCallInternal(gridDim, blockDim, sharedMem, stream); +#endif + for(unsigned i = 0; i < entry->num_args(); i++){ + std::pair p = entry->get_param_config(i); + cudaSetupArgumentInternal(args[i], p.first, p.second); + } - cudaLaunchInternal(hostFun); - return g_last_cudaError = cudaSuccess; + cudaLaunchInternal(hostFun); + return g_last_cudaError = cudaSuccess; } @@ -2974,12 +2975,17 @@ extern "C" { void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) { + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } return cudaRegisterFatBinaryInternal(fatCubin); } void CUDARTAPI __cudaRegisterFatBinaryEnd( void **fatCubinHandle ) { - + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } } unsigned CUDARTAPI __cudaPushCallConfiguration(dim3 gridDim, @@ -2987,7 +2993,10 @@ unsigned CUDARTAPI __cudaPushCallConfiguration(dim3 gridDim, size_t sharedMem = 0, struct CUstream_st *stream = 0) { - + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + cudaConfigureCallInternal(gridDim, blockDim, sharedMem, stream); } cudaError_t CUDARTAPI __cudaPopCallConfiguration( @@ -2997,6 +3006,9 @@ cudaError_t CUDARTAPI __cudaPopCallConfiguration( void *stream ) { + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } return g_last_cudaError = cudaSuccess; } @@ -3056,7 +3068,6 @@ void __cudaUnregisterFatBinary(void **fatCubinHandle) if(g_debug_execution >= 3){ announce_call(__my_func__); } - ; } cudaError_t cudaDeviceReset ( void ) { -- cgit v1.3 From f63324eda157f742d06c889b1be73717771e60a5 Mon Sep 17 00:00:00 2001 From: tgrogers Date: Tue, 18 Jun 2019 22:34:07 -0400 Subject: Some stuff to get mnist to run with CUDA 9.1 --- libcuda/cuda_runtime_api.cc | 49 +++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 49 insertions(+) (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index f46d9f1..8ba2b0f 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1694,6 +1694,12 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic break; case 88: case 89: + case 90: + case 91: + case 92: + case 93: + case 94: + case 95: *value= 0; break; default: @@ -3246,6 +3252,49 @@ __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t v #endif +/** + * \brief Set attributes for a given function + * + * This function sets the attributes of a function specified via \p entry. + * The parameter \p entry must be a pointer to a function that executes + * on the device. The parameter specified by \p entry must be declared as a \p __global__ + * function. The enumeration defined by \p attr is set to the value defined by \p value + * If the specified function does not exist, then ::cudaErrorInvalidDeviceFunction is returned. + * If the specified attribute cannot be written, or if the value is incorrect, + * then ::cudaErrorInvalidValue is returned. + * + * Valid values for \p attr are: + * ::cuFuncAttrMaxDynamicSharedMem - Maximum size of dynamic shared memory per block + * ::cudaFuncAttributePreferredSharedMemoryCarveout - Preferred shared memory-L1 cache split ratio + * + * \param entry - Function to get attributes of + * \param attr - Attribute to set + * \param value - Value to set + * + * \return + * ::cudaSuccess, + * ::cudaErrorInitializationError, + * ::cudaErrorInvalidDeviceFunction, + * ::cudaErrorInvalidValue + * \notefnerr + * + * \ref ::cudaLaunchKernel(const T *func, dim3 gridDim, dim3 blockDim, void **args, size_t sharedMem, cudaStream_t stream) "cudaLaunchKernel (C++ API)", + * \ref ::cudaFuncSetCacheConfig(T*, enum cudaFuncCache) "cudaFuncSetCacheConfig (C++ API)", + * \ref ::cudaFuncGetAttributes(struct cudaFuncAttributes*, const void*) "cudaFuncGetAttributes (C API)", + * ::cudaSetDoubleForDevice, + * ::cudaSetDoubleForHost, + * \ref ::cudaSetupArgument(T, size_t) "cudaSetupArgument (C++ API)" + */ +cudaError_t CUDARTAPI cudaFuncSetAttribute(const void *func, enum cudaFuncAttribute attr, int value) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s ( func=%p, attr=%d, value=%d )\"\n", + __my_func__, func, attr, value ); + return g_last_cudaError = cudaSuccess; +} + cudaError_t CUDARTAPI cudaGLSetGLDevice(int device) { if(g_debug_execution >= 3){ -- cgit v1.3 From 48830687ede62b3acaebeba93633255b4d8ec9c8 Mon Sep 17 00:00:00 2001 From: tgrogers Date: Wed, 19 Jun 2019 10:22:41 -0400 Subject: Fixing 4.2 build --- libcuda/cuda_runtime_api.cc | 3 +++ 1 file changed, 3 insertions(+) (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 8ba2b0f..015fbc0 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -3252,6 +3252,8 @@ __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t v #endif + +#if CUDART_VERSION >= 9000 /** * \brief Set attributes for a given function * @@ -3294,6 +3296,7 @@ cudaError_t CUDARTAPI cudaFuncSetAttribute(const void *func, enum cudaFuncAttrib __my_func__, func, attr, value ); return g_last_cudaError = cudaSuccess; } +#endif cudaError_t CUDARTAPI cudaGLSetGLDevice(int device) { -- cgit v1.3