diff options
| -rw-r--r-- | Makefile | 1 | ||||
| -rw-r--r-- | README | 26 | ||||
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 11 | ||||
| -rw-r--r-- | linux-so-version.txt | 4 | ||||
| -rw-r--r-- | src/gpgpu-sim/gpu-cache.h | 6 |
5 files changed, 44 insertions, 4 deletions
@@ -163,6 +163,7 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib 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\ @@ -177,7 +177,7 @@ GPGPU-Sim dependencies: * flex * zlib * CUDA Toolkit - + GPGPU-Sim documentation dependencies: * doxygen * graphvi @@ -208,6 +208,9 @@ python-matplotlib" CUDA SDK dependencies: "sudo apt-get install libxi-dev libxmu-dev libglut3-dev" +If you are running applications which use NVIDIA libraries such as cuDNN and +cuBLAS, install them too. + Finally, ensure CUDA_INSTALL_PATH is set to the location where you installed the CUDA Toolkit (e.g., /usr/local/cuda) and that $CUDA_INSTALL_PATH/bin is in your PATH. You probably want to modify your .bashrc file to incude the @@ -216,6 +219,10 @@ following (this assumes the CUDA Toolkit was installed in /usr/local/cuda): export CUDA_INSTALL_PATH=/usr/local/cuda export PATH=$CUDA_INSTALL_PATH/bin +If running applications which use cuDNN or cuBLAS: + export CUDNN_PATH=<Path To cuDNN Directory> + export LD_LIBRARY_PATH=$CUDA_INSTALL_PATH/lib64:$CUDA_INSTALL_PATH/lib:$CUDNN_PATH/lib64 + Step 2: Build ============= @@ -261,6 +268,23 @@ ldd <your_application_name> You should see that your application is using libcudart.so file in GPGPUSim directory. +If running applications which use cuDNN or cuBLAS: + + * Modify the Makefile or the compilation command of the application to change + all the dynamic links to static ones, for example: + * -L$(CUDA_PATH)/lib64 -lcublas to + -L$(CUDA_PATH)/lib64 -lcublas_static + + * -L$(CUDNN_PATH)/lib64 -lcudnn to + -L$(CUDNN_PATH)/lib64 -lcudnn_static + + * Modify the Makefile or the compilation command such that the following + flags are used by the nvcc compiler: + -gencode arch=compute_61,code=compute_61 + + (the number 61 refers to the SM version. You would need to set it based + on the GPGPU-Sim config "-gpgpu ptx force max capability" you use) + Copy the contents of configs/QuadroFX5800/ or configs/GTX480/ to your application's working directory. These files configure the microarchitecture models to resemble the respective GPGPU architectures. 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 ); 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/gpgpu-sim/gpu-cache.h b/src/gpgpu-sim/gpu-cache.h index 19a1542..337f710 100644 --- a/src/gpgpu-sim/gpu-cache.h +++ b/src/gpgpu-sim/gpu-cache.h @@ -681,15 +681,15 @@ public: // tag + index is required to check for hit/miss. Tag is now identical to the block address. //return addr >> (m_line_sz_log2+m_nset_log2); - return addr & ~(m_line_sz-1); + return addr & ~(new_addr_type)(m_line_sz-1); } new_addr_type block_addr( new_addr_type addr ) const { - return addr & ~(m_line_sz-1); + return addr & ~(new_addr_type)(m_line_sz-1); } new_addr_type mshr_addr( new_addr_type addr ) const { - return addr & ~(m_atom_sz-1); + return addr & ~(new_addr_type)(m_atom_sz-1); } enum mshr_config_t get_mshr_type() const { |