From f8e84eefc92d7d63554472e0bb88fe691f40701d Mon Sep 17 00:00:00 2001
From: tgrogers <timrogers@gmail.com>
Date: Wed, 12 Jul 2017 17:03:30 -0400
Subject: Getting jenkins to work by changing the error on shell type to a
 warning. The way we detect the shell is not fool-proof and fails for jenkins,
 although we are using bash

---
 setup_environment | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

(limited to 'setup_environment')

diff --git a/setup_environment b/setup_environment
index 854a335..0895d44 100644
--- a/setup_environment
+++ b/setup_environment
@@ -1,6 +1,6 @@
 # see README before running this
 
-ps -p $$ | awk '/bash/ || / sh/ || /zsh/ {exit 1;}' && echo "ERROR ** source setup_environment must be run in a bash, zsh or sh shell; see README" && exit
+ps -p $$ | awk '/bash/ || / sh/ || /zsh/ {exit 1;}' && echo "WARNING ** source setup_environment must be run in a bash, zsh or sh shell; see README"
 
 export GPGPUSIM_SETUP_ENVIRONMENT_WAS_RUN=
 export GPGPUSIM_ROOT="$( cd "$( dirname "$BASH_SOURCE" )" && pwd )"
-- 
cgit v1.3


From 6347147dbd059410648aaf1045cfc73a7294b4e6 Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Tue, 12 Dec 2017 21:46:54 -0500
Subject: support for 9.1

---
 Makefile          | 1 +
 setup_environment | 2 +-
 2 files changed, 2 insertions(+), 1 deletion(-)

(limited to 'setup_environment')

diff --git a/Makefile b/Makefile
index e83de9a..95e4432 100644
--- a/Makefile
+++ b/Makefile
@@ -161,6 +161,7 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	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.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.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.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/setup_environment b/setup_environment
index 0895d44..8a529c3 100644
--- a/setup_environment
+++ b/setup_environment
@@ -43,7 +43,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/,//'`;
 CUDA_VERSION_NUMBER=`echo $CUDA_VERSION_STRING | sed 's/\./ /' | awk '{printf("%02u%02u", 10*int($1), 10*$2);}'`
-if [ $CUDA_VERSION_NUMBER -gt 8000 -o $CUDA_VERSION_NUMBER -lt 2030  ]; then
+if [ $CUDA_VERSION_NUMBER -gt 9100 -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
 elif [ $CUDA_VERSION_NUMBER -gt 4020 ]; then
-- 
cgit v1.3


From 25dcce6d91d088eb91d26fe0b6aa255090f24a4b Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Thu, 22 Feb 2018 15:59:37 -0500
Subject: run all the configs, do 9.1 on the sdk and stop the setup environment
 from complaining about cuda version

---
 Jenkinsfile       | 9 +++++++--
 setup_environment | 2 +-
 2 files changed, 8 insertions(+), 3 deletions(-)

(limited to 'setup_environment')

diff --git a/Jenkinsfile b/Jenkinsfile
index cf9d71e..0f6314f 100644
--- a/Jenkinsfile
+++ b/Jenkinsfile
@@ -46,12 +46,17 @@ pipeline {
                 }, "9.1-rodinia": {
                     sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
                     source `pwd`/setup_environment &&\
-                    ./gpgpu-sim_simulations/util/job_launching/run_simulations.py -B rodinia_2.0-ft -C PASCALTITANX -N regress-$$ && \
+                    ./gpgpu-sim_simulations/util/job_launching/run_simulations.py -B rodinia_2.0-ft -C PASCALTITANX,TITANX-P102 -N regress-$$ && \
                     ./gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -v -N regress-$$'
                 }, "4.2-sdk-4.2": {
                     sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/4.2_env_setup.sh &&\
                     source `pwd`/setup_environment &&\
-                    ./gpgpu-sim_simulations/util/job_launching/run_simulations.py -B sdk-4.2 -C GTX480,PASCALTITANX -N regress-$$ && \
+                    ./gpgpu-sim_simulations/util/job_launching/run_simulations.py -B sdk-4.2 -C GTX480,PASCALTITANX,TITANX-P102 -N regress-$$ && \
+                    ./gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -v -N regress-$$'
+                }, "9.1-sdk-4.2": {
+                    sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
+                    source `pwd`/setup_environment &&\
+                    ./gpgpu-sim_simulations/util/job_launching/run_simulations.py -B sdk-4.2 -C PASCALTITANX,TITANX-P102 -N regress-$$ && \
                     ./gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -v -N regress-$$'
                 }
             }
diff --git a/setup_environment b/setup_environment
index 8a529c3..cd6452e 100644
--- a/setup_environment
+++ b/setup_environment
@@ -46,7 +46,7 @@ CUDA_VERSION_NUMBER=`echo $CUDA_VERSION_STRING | sed 's/\./ /' | awk '{printf("%
 if [ $CUDA_VERSION_NUMBER -gt 9100 -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
-elif [ $CUDA_VERSION_NUMBER -gt 4020 ]; then
+elif [ $CUDA_VERSION_NUMBER -gt 9100 ]; then
 	echo "WARNING ** GPGPU-Sim version $GPGPUSIM_VERSION_STRING not fully tested with CUDA version $CUDA_VERSION_STRING (please see README)";
 fi
 
-- 
cgit v1.3


From 89094df5b4e861b104393ea0f41886494a7f26a4 Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Wed, 11 Apr 2018 15:46:01 -0400
Subject: This build status line was broken when there were cached files

---
 setup_environment | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

(limited to 'setup_environment')

diff --git a/setup_environment b/setup_environment
index 96cc362..3f2476e 100644
--- a/setup_environment
+++ b/setup_environment
@@ -8,7 +8,8 @@ export GPGPUSIM_ROOT="$( cd "$( dirname "$BASH_SOURCE" )" && pwd )"
 GPGPUSIM_VERSION_STRING=`cat $GPGPUSIM_ROOT/version | awk '/Version/ {print $8}'`
 #Detect Git branch and commit #
 GIT_COMMIT=`git log -n 1 | head -1 | sed -re 's/commit (.*)/\1/'`
-GIT_FILES_CHANGED=`git diff --numstat --cached && git diff --numstat | wc | sed -re 's/^\s+([0-9]+).*/\1/'`
+GIT_FILES_CHANGED=`git diff --numstat | wc | sed -re 's/^\s+([0-9]+).*/\1./'`
+GIT_FILES_CHANGED+=`git diff --numstat --cached | wc | sed -re 's/^\s+([0-9]+).*/\1/'`
 GPGPUSIM_BUILD_STRING="gpgpu-sim_git-commit-$GIT_COMMIT-modified_$GIT_FILES_CHANGED"
 
 echo -n "GPGPU-Sim version $GPGPUSIM_VERSION_STRING (build $GPGPUSIM_BUILD_STRING) ";
-- 
cgit v1.3


From 112a3a2e3c9d17af420d2389f68a35407c692744 Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Fri, 29 Jun 2018 17:34:22 -0400
Subject: Allowing modern GPU configrations to properly run PTXPLUS. There are
 several problems with this:

1) To get proper occupancy, based on register file usage, you must querry ptxas with
   the sm version that corresponds to the register usage in the config file. To enable this,
   a new config option has been added that determines what sm version you pass to ptxas to compute reg-usage.
   This configuration option is always required in the gpgpusim.config file
2) If you are running PTXPLUS with a modern card (i.e. volta/pascal), you need ptxas from CUDA 9.1.
   However since PTXPLUS only supports sm_13 - you need a version of CUDA where cudaobjectdump
   supports sm_1x. This ended at CUDA 5 - and PTXPLUS requires CUDA 4.2. Therefore, to run PTXPLUS
   on a modern card, you need CUDA 4.2 + modern CUDA installed. To fascilitate this, a new envronment
   varaible is added and the setup envrionment script prints an appropraite warning if you are using a newer CUDA.

We have tried to make this as fail-proof as possible - and die appropraitely when something is wrong.
---
 configs/GTX480/gpgpusim.config           |  1 +
 configs/GeForceGTX1080Ti/gpgpusim.config |  3 ++-
 configs/GeForceGTX750Ti/gpgpusim.config  |  1 +
 configs/QuadroFX5600/gpgpusim.config     |  3 +++
 configs/QuadroFX5800/gpgpusim.config     |  2 ++
 configs/TeslaC2050/gpgpusim.config       |  1 +
 libcuda/cuda_runtime_api.cc              | 43 +++++++++++++++++++++-----------
 libopencl/opencl_runtime_api.cc          |  2 +-
 setup_environment                        | 13 ++++++++++
 src/cuda-sim/ptx_loader.cc               | 18 ++++++++++---
 src/cuda-sim/ptx_loader.h                |  2 +-
 11 files changed, 67 insertions(+), 22 deletions(-)

(limited to 'setup_environment')

diff --git a/configs/GTX480/gpgpusim.config b/configs/GTX480/gpgpusim.config
index 436cb41..ee90c12 100644
--- a/configs/GTX480/gpgpusim.config
+++ b/configs/GTX480/gpgpusim.config
@@ -22,6 +22,7 @@
 
 # shader core pipeline config
 -gpgpu_shader_registers 32768
+-gpgpu_occupancy_sm_number 20
 
 # This implies a maximum of 48 warps/SM
 -gpgpu_shader_core_pipeline 1536:32 
diff --git a/configs/GeForceGTX1080Ti/gpgpusim.config b/configs/GeForceGTX1080Ti/gpgpusim.config
index 47c2b6a..fb044c6 100644
--- a/configs/GeForceGTX1080Ti/gpgpusim.config
+++ b/configs/GeForceGTX1080Ti/gpgpusim.config
@@ -3,7 +3,7 @@
 # functional simulator specification
 -gpgpu_ptx_instruction_classification 0
 -gpgpu_ptx_sim_mode 0
--gpgpu_ptx_force_max_capability 20 
+-gpgpu_ptx_force_max_capability 60
 
 # SASS execution (only supported with CUDA >= 4.0)
 -gpgpu_ptx_convert_to_ptxplus 0
@@ -23,6 +23,7 @@
 
 # shader core pipeline config
 -gpgpu_shader_registers 65536
+-gpgpu_occupancy_sm_number 60
 
 # This implies a maximum of 64 warps/SM
 -gpgpu_shader_core_pipeline 2048:32 
diff --git a/configs/GeForceGTX750Ti/gpgpusim.config b/configs/GeForceGTX750Ti/gpgpusim.config
index 8b030b6..c675aab 100644
--- a/configs/GeForceGTX750Ti/gpgpusim.config
+++ b/configs/GeForceGTX750Ti/gpgpusim.config
@@ -21,6 +21,7 @@
 
 # shader core pipeline config
 -gpgpu_shader_registers 65536
+-gpgpu_occupancy_sm_number 52
 
 # This implies a maximum of 64 warps/SM
 -gpgpu_shader_core_pipeline 2048:32 
diff --git a/configs/QuadroFX5600/gpgpusim.config b/configs/QuadroFX5600/gpgpusim.config
index cb87b65..6f836ee 100644
--- a/configs/QuadroFX5600/gpgpusim.config
+++ b/configs/QuadroFX5600/gpgpusim.config
@@ -11,6 +11,9 @@
 
 # shader core pipeline config
 -gpgpu_shader_registers 16384
+-gpgpu_occupancy_sm_number 12
+
+-gpgpu_occupancy_sm_number 12
 #8192 (registers per block as written by device Query and which used in this option in our other configurations but this break some benchmarks execution! it does not affect performance modeling though)
 -gpgpu_shader_core_pipeline 768:32 
 -gpgpu_shader_cta 8
diff --git a/configs/QuadroFX5800/gpgpusim.config b/configs/QuadroFX5800/gpgpusim.config
index 82243c2..fef1110 100644
--- a/configs/QuadroFX5800/gpgpusim.config
+++ b/configs/QuadroFX5800/gpgpusim.config
@@ -11,6 +11,8 @@
 
 # shader core pipeline config
 -gpgpu_shader_registers 16384
+-gpgpu_occupancy_sm_number 13
+
 -gpgpu_shader_core_pipeline 1024:32 
 -gpgpu_shader_cta 8
 -gpgpu_simd_model 1
diff --git a/configs/TeslaC2050/gpgpusim.config b/configs/TeslaC2050/gpgpusim.config
index 442ab8b..6ac2c12 100644
--- a/configs/TeslaC2050/gpgpusim.config
+++ b/configs/TeslaC2050/gpgpusim.config
@@ -25,6 +25,7 @@
 
 # shader core pipeline config
 -gpgpu_shader_registers 32768
+-gpgpu_occupancy_sm_number 20
 
 # This implies a maximum of 48 warps/SM
 -gpgpu_shader_core_pipeline 1536:32 
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 2b0f944..94705dc 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -1350,39 +1350,52 @@ std::string get_app_binary(){
  *	enabled
  * */
 void extract_code_using_cuobjdump(){
-	CUctx_st *context = GPGPUSim_Context();
-	char command[1000];
+    CUctx_st *context = GPGPUSim_Context();
+    std::string command;
 
-   std::string app_binary = get_app_binary(); 
+    std::string app_binary = get_app_binary(); 
 
 	char fname[1024];
 	snprintf(fname,1024,"_cuobjdump_complete_output_XXXXXX");
 	int fd=mkstemp(fname);
 	close(fd);
 	// Running cuobjdump using dynamic link to current process
-	snprintf(command,1000,"md5sum %s ", app_binary.c_str());
-	printf("Running md5sum using \"%s\"\n", command);
-	system(command);
+    command = "md5sum " + app_binary;
+	printf("Running md5sum using \"%s\"\n", command.c_str());
+	system(command.c_str());
 	// Running cuobjdump using dynamic link to current process
 	// Needs the option '-all' to extract PTX from CDP-enabled binary 
 	extern bool g_cdp_enabled;
+    const int current_cuda = atoi(getenv("CUDA_VERSION_NUMBER"));
+    if ( context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() && current_cuda > 42 ) {
+        const char* cuda_42_path = getenv("CUDA_42_INSTALL_PATH");
+        if ( NULL == cuda_42_path ) {
+            printf("You are using a new version of CUDA and PTXPLUS. You are required to have CUDA SDK 4.2 and explicitly "
+                   "point to it via the environment variable $CUDA_42_INSTALL_PATH.");
+            exit(1);
+        } else {
+            command = "$CUDA_42_INSTALL_PATH/bin/cuobjdump";
+        }
+    } else {
+        command = "$CUDA_INSTALL_PATH/bin/cuobjdump";
+    }
 	if(!g_cdp_enabled)
-	    snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass %s > %s", app_binary.c_str(), fname);
+        command += " -ptx -elf -sass " + app_binary + " > " + fname;
 	else
-	    snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass -all %s > %s", app_binary.c_str(), fname);
+        command += " -ptx -elf -sass -all " + app_binary + " > " + fname;
 	bool parse_output = true; 
-	int result = system(command);
+	int result = system(command.c_str());
 	if(result) {
 		if (context->get_device()->get_gpgpu()->get_config().experimental_lib_support() && (result == 65280)) {  
 			// Some CUDA application may exclusively use kernels provided by CUDA
 			// libraries (e.g. CUBLAS).  Skipping cuobjdump extraction from the
 			// executable for this case. 
 			// 65280 is the return code from cuobjdump denoting the specific error (tested on CUDA 4.0/4.1/4.2)
-			printf("WARNING: Failed to execute: %s\n", command); 
+			printf("WARNING: Failed to execute: %s\n", command.c_str()); 
 			printf("         Executable binary does not contain any GPU kernel.\n"); 
 			parse_output = false; 
 		} else {
-			printf("ERROR: Failed to execute: %s\n", command); 
+			printf("ERROR: Failed to execute: %s\n", command.c_str()); 
 			exit(1);
 		}
 	}
@@ -1435,7 +1448,7 @@ void extract_code_using_cuobjdump(){
 			std::cout << "Running cuobjdump on " << line << std::endl;
 			std::cout << "Using command: " << cmd.str() << std::endl;
 			result = system(cmd.str().c_str());
-			if(result) {printf("ERROR: Failed to execute: %s\n", command); exit(1);}
+			if(result) {printf("ERROR: Failed to execute: %s\n", command.c_str()); exit(1);}
 			std::cout << "Done" << std::endl;
 
 			std::cout << "Trying to parse " << libcodfn.str() << std::endl;
@@ -1737,13 +1750,13 @@ void cuobjdumpParseBinary(unsigned int handle){
         symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxplus_str, handle);
         printf("Adding %s with cubin handle %u\n", ptx->getPTXfilename().c_str(), handle);
         context->add_binary(symtab, handle);
-        gpgpu_ptxinfo_load_from_string( ptxcode, handle, max_capability );
+        gpgpu_ptxinfo_load_from_string( ptxcode, handle );
         delete[] ptxplus_str;
     } else {
         symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxcode, handle);
         printf("Adding %s with cubin handle %u\n", ptx->getPTXfilename().c_str(), handle);
         context->add_binary(symtab, handle);
-        gpgpu_ptxinfo_load_from_string( ptxcode, handle, max_capability );
+        gpgpu_ptxinfo_load_from_string( ptxcode, handle );
     }
     load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
     load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
@@ -1847,7 +1860,7 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
 			} else {
 				symtab=gpgpu_ptx_sim_load_ptx_from_string(ptx,source_num);
 				context->add_binary(symtab,fat_cubin_handle);
-				gpgpu_ptxinfo_load_from_string( ptx, source_num, max_capability );
+				gpgpu_ptxinfo_load_from_string( ptx, source_num );
 			}
 			source_num++;
 			load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index 63f7c81..97a54d8 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -577,7 +577,7 @@ void _cl_program::Build(const char *options)
       }
       info.m_asm = tmp;
       info.m_symtab = gpgpu_ptx_sim_load_ptx_from_string( tmp, source_num );
-      gpgpu_ptxinfo_load_from_string( tmp, source_num);
+      gpgpu_ptxinfo_load_from_string( tmp, source_num );
       free(tmp);
    }
    printf("GPGPU-Sim OpenCL API: finished compiling OpenCL kernels.\n"); 
diff --git a/setup_environment b/setup_environment
index 3f2476e..939fb85 100644
--- a/setup_environment
+++ b/setup_environment
@@ -53,6 +53,8 @@ if [ $CUDA_VERSION_NUMBER -gt 9100 -o $CUDA_VERSION_NUMBER -lt 2030  ]; then
 	return
 fi
 
+
+
 if [ $# = '1' ] ;
 then
     export GPGPUSIM_CONFIG=gcc-$CC_VERSION/cuda-$CUDA_VERSION_NUMBER/$1
@@ -134,4 +136,15 @@ fi
 
 echo "setup_environment succeeded";
 
+if [ $CUDA_VERSION_NUMBER -gt 4200 ]; then
+    echo ""
+    echo "----------------------------------------------------------------------------"
+    echo "WARNING - If you only care about PTX execution, ignore this warning.";
+    echo "If you are using PTXPLUS, only sm_13 is supported and it requires that the app binaries are compiled in CUDA 4.2 or less.";
+    echo "New versions of CUDA tools have dropped parsing support for sm_1x - therefore to use PTXPLUS with a modern card configuration, you must";
+    echo "have both CUDA 4.2 and a modern CUDA installed. When running PTXPLUS under a CUDA version >= 4.2, the simulator will fail unless \$CUDA_42_INSTALL_PATH";
+    echo "has been set properly. You must set this manually.";
+    echo "----------------------------------------------------------------------------"
+fi
+
 export GPGPUSIM_SETUP_ENVIRONMENT_WAS_RUN=1
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 6c1b595..99174cf 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -54,6 +54,7 @@ extern int ptxinfo_debug;
 extern FILE *ptxinfo_in;
 
 static bool g_save_embedded_ptx;
+static int g_occupancy_sm_number;
 bool g_keep_intermediate_files;
 bool m_ptx_save_converted_ptxplus;
 
@@ -71,6 +72,10 @@ void ptx_reg_options(option_parser_t opp)
                 &m_ptx_save_converted_ptxplus,
                 "Saved converted ptxplus to a file",
                 "0");
+   option_parser_register(opp, "-gpgpu_occupancy_sm_number", OPT_INT32, &g_occupancy_sm_number,
+                "The SM number to pass to ptxas when getting register usage for computing GPU occupancy."
+                "This parameter is required in the config.",
+                "0");
 }
 
 void print_ptx_file( const char *p, unsigned source_num, const char *filename )
@@ -287,7 +292,7 @@ void fix_duplicate_errors(char fname2[1024]) {
 	}
 }
 
-void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version )
+void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num )
 {
     char fname[1024];
     snprintf(fname,1024,"_ptx_XXXXXX");
@@ -321,12 +326,17 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
     extra_flags[0]=0;
 
 #if CUDART_VERSION >= 3000
-    if (sm_version == 0) sm_version = 20;
+    if ( g_occupancy_sm_number == 0 ) {
+        printf( "gpgpusim.config must specify the sm version for the GPU that you use to compute occupancy.\n"
+                "The register file size is specifically tied to the sm version used to querry ptxas for register usage."
+                "A register size/SM mismatch may result in occupancy differences." );
+        exit(1);
+    }
     extern bool g_cdp_enabled;
     if(!g_cdp_enabled)
-        snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version);
+        snprintf(extra_flags,1024,"--gpu-name=sm_%u", g_occupancy_sm_number);
     else
-        snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",sm_version);
+        snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",g_occupancy_sm_number);
 #endif
 
     snprintf(commandline,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file  /dev/null 2> %s",
diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h
index d3d0c92..cb02eda 100644
--- a/src/cuda-sim/ptx_loader.h
+++ b/src/cuda-sim/ptx_loader.h
@@ -32,7 +32,7 @@
 extern bool g_override_embedded_ptx;
  
 class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num );
-void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20 );
+void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num );
 char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str);
 bool keep_intermediate_files();
 
-- 
cgit v1.3


From d42cb5648f77cd2caed9c656f4d6ad75c23b6f5a Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Fri, 29 Jun 2018 23:43:31 -0400
Subject: Fixing the PTXPLUS + new CUDA execution case.

---
 libcuda/cuda_runtime_api.cc | 61 ++++++++++++++++++++++++++++++++-------------
 setup_environment           | 16 ++++++------
 src/cuda-sim/ptx_loader.cc  |  6 ++---
 3 files changed, 55 insertions(+), 28 deletions(-)

(limited to 'setup_environment')

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 94705dc..5964a4d 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -1367,17 +1367,17 @@ void extract_code_using_cuobjdump(){
 	// Needs the option '-all' to extract PTX from CDP-enabled binary 
 	extern bool g_cdp_enabled;
     const int current_cuda = atoi(getenv("CUDA_VERSION_NUMBER"));
-    if ( context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() && current_cuda > 42 ) {
+    if ( context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() && current_cuda > 4020 ) {
         const char* cuda_42_path = getenv("CUDA_42_INSTALL_PATH");
         if ( NULL == cuda_42_path ) {
             printf("You are using a new version of CUDA and PTXPLUS. You are required to have CUDA SDK 4.2 and explicitly "
-                   "point to it via the environment variable $CUDA_42_INSTALL_PATH.");
+                   "point to it via the environment variable $CUDA_42_INSTALL_PATH.\n\n");
             exit(1);
         } else {
             command = "$CUDA_42_INSTALL_PATH/bin/cuobjdump";
         }
     } else {
-        command = "$CUDA_INSTALL_PATH/bin/cuobjdump";
+        command = "$CUDA_42_INSTALL_PATH/bin/cuobjdump";
     }
 	if(!g_cdp_enabled)
         command += " -ptx -elf -sass " + app_binary + " > " + fname;
@@ -1778,21 +1778,46 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
 		if (sizeof(void*) == 4) 
 			printf("GPGPU-Sim PTX: FatBin file name extraction has not been tested on 32-bit system.\n"); 
 
-		#if (CUDART_VERSION <= 6000)
-		// FatBin handle from the .fatbin.c file (one of the intermediate files generated by NVCC)
-		typedef struct {int m; int v; const unsigned long long* d; char* f;} __fatDeviceText __attribute__ ((aligned (8))); 
-		__fatDeviceText * fatDeviceText = (__fatDeviceText *) fatCubin;
-
-		// Extract the source code file name that generate the given FatBin. 
-		// - Obtains the pointer to the actual fatbin structure from the FatBin handle (fatCubin).
-		// - An integer inside the fatbin structure contains the relative offset to the source code file name.
-		// - This offset differs among different CUDA and GCC versions. 
-		char * pfatbin = (char*) fatDeviceText->d; 
-		int offset = *((int*)(pfatbin+48)); 
-		char * filename = (pfatbin+16+offset); 
-		#else
-		const char * filename = "default";
-		#endif
+        // This code will get the CUDA version the app was compiled with.
+        // We need this to determine how to handle the parsing of the binary.
+        // Making this a runtime variable based on the app, enables GPGPU-Sim compiled
+        // with a newer version of CUDA to run apps compiled with older versions of
+        // CUDA. This is especially useful for PTXPLUS execution.
+        char fname[1024];
+        snprintf(fname,1024,"_app_cuda_version_XXXXXX");
+        int fd=mkstemp(fname);
+        close(fd);
+        int app_cuda_version = 0;
+        std::string app_cuda_version_command = "ldd " + get_app_binary() + " | grep libcudart.so | sed  's/.*libcudart.so.\\(.*\\) =>.*/\\1/' > " + fname;
+        system(app_cuda_version_command.c_str());
+        FILE * cmd = fopen(fname, "r");
+        char buf[256];
+        while (fgets(buf, sizeof(buf), cmd) != 0) {
+            std::cout << buf;
+            app_cuda_version = atoi(buf);
+        }
+        fclose(cmd);
+        if ( app_cuda_version == 0 ) {
+            printf( "Error - Cannot detect the app's CUDA version.\n" );
+            exit(1);
+        }
+        const char* filename;
+        if ( app_cuda_version < 6 ) {
+            // FatBin handle from the .fatbin.c file (one of the intermediate files generated by NVCC)
+            typedef struct {int m; int v; const unsigned long long* d; char* f;} __fatDeviceText __attribute__ ((aligned (8))); 
+            __fatDeviceText * fatDeviceText = (__fatDeviceText *) fatCubin;
+
+            // Extract the source code file name that generate the given FatBin. 
+            // - Obtains the pointer to the actual fatbin structure from the FatBin handle (fatCubin).
+            // - An integer inside the fatbin structure contains the relative offset to the source code file name.
+            // - This offset differs among different CUDA and GCC versions. 
+            char * pfatbin = (char*) fatDeviceText->d; 
+            int offset = *((int*)(pfatbin+48)); 
+            filename = (pfatbin+16+offset);
+        } else {
+            filename = "default";
+        }
+
 		// The extracted file name is associated with a fat_cubin_handle passed
 		// into cudaLaunch().  Inside cudaLaunch(), the associated file name is
 		// used to find the PTX/SASS section from cuobjdump, which contains the
diff --git a/setup_environment b/setup_environment
index 939fb85..8bcb9aa 100644
--- a/setup_environment
+++ b/setup_environment
@@ -47,7 +47,7 @@ fi
 CC_VERSION=`gcc --version | head -1 | awk '{for(i=1;i<=NF;i++){ if(match($i,/^[0-9]\.[0-9]\.[0-9]$/))  {print $i; exit 0}}}'`
 
 CUDA_VERSION_STRING=`$CUDA_INSTALL_PATH/bin/nvcc --version | awk '/release/ {print $5;}' | sed 's/,//'`;
-CUDA_VERSION_NUMBER=`echo $CUDA_VERSION_STRING | sed 's/\./ /' | awk '{printf("%02u%02u", 10*int($1), 10*$2);}'`
+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
 	echo "ERROR ** GPGPU-Sim version $GPGPUSIM_VERSION_STRING not tested with CUDA version $CUDA_VERSION_STRING (please see README)";
 	return
@@ -134,17 +134,19 @@ else
 	echo "configured without a power model.";
 fi
 
-echo "setup_environment succeeded";
-
 if [ $CUDA_VERSION_NUMBER -gt 4200 ]; then
-    echo ""
-    echo "----------------------------------------------------------------------------"
-    echo "WARNING - If you only care about PTX execution, ignore this warning.";
-    echo "If you are using PTXPLUS, only sm_13 is supported and it requires that the app binaries are compiled in CUDA 4.2 or less.";
+    echo "";
+    echo "----------------------------------------------------------------------------";
+    echo "WARNING - If you only care about PTX execution, ignore this warning. GPGPU-Sim supports PTX execution in modern CUDA."
+    echo "If you want to run PTXPLUS, download SDK 4.2 (in addition to your modern CUDA) and set \$CUDA_42_INSTALL_PATH."
+    echo "The following text describes why:";
+    echo "If you are using PTXPLUS, only sm_1x is supported and it requires that the app binaries are compiled in CUDA 4.2 or less.";
     echo "New versions of CUDA tools have dropped parsing support for sm_1x - therefore to use PTXPLUS with a modern card configuration, you must";
     echo "have both CUDA 4.2 and a modern CUDA installed. When running PTXPLUS under a CUDA version >= 4.2, the simulator will fail unless \$CUDA_42_INSTALL_PATH";
     echo "has been set properly. You must set this manually.";
     echo "----------------------------------------------------------------------------"
 fi
 
+echo "setup_environment succeeded";
+
 export GPGPUSIM_SETUP_ENVIRONMENT_WAS_RUN=1
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 99174cf..22f03b9 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -73,7 +73,7 @@ void ptx_reg_options(option_parser_t opp)
                 "Saved converted ptxplus to a file",
                 "0");
    option_parser_register(opp, "-gpgpu_occupancy_sm_number", OPT_INT32, &g_occupancy_sm_number,
-                "The SM number to pass to ptxas when getting register usage for computing GPU occupancy."
+                "The SM number to pass to ptxas when getting register usage for computing GPU occupancy. "
                 "This parameter is required in the config.",
                 "0");
 }
@@ -327,8 +327,8 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
 
 #if CUDART_VERSION >= 3000
     if ( g_occupancy_sm_number == 0 ) {
-        printf( "gpgpusim.config must specify the sm version for the GPU that you use to compute occupancy.\n"
-                "The register file size is specifically tied to the sm version used to querry ptxas for register usage."
+        printf( "gpgpusim.config must specify the sm version for the GPU that you use to compute occupancy \"-gpgpu_occupancy_sm_number XX\".\n"
+                "The register file size is specifically tied to the sm version used to querry ptxas for register usage.\n"
                 "A register size/SM mismatch may result in occupancy differences." );
         exit(1);
     }
-- 
cgit v1.3


From 691c57ea8596da35da3ce90a7b58796f3791036c Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Sat, 30 Jun 2018 10:12:20 -0400
Subject: changing the flow s.t. you cannot compile the simulator with a
 different version of CUDA from the app.

---
 diff.diff                   | 173 ++++++++++++++++++++++++++++++++++++++++++++
 libcuda/cuda_runtime_api.cc |  69 +++++++++---------
 setup_environment           |  28 ++++---
 src/cuda-sim/ptx_loader.cc  |   2 +-
 4 files changed, 225 insertions(+), 47 deletions(-)
 create mode 100644 diff.diff

(limited to 'setup_environment')

diff --git a/diff.diff b/diff.diff
new file mode 100644
index 0000000..8ab27fe
--- /dev/null
+++ b/diff.diff
@@ -0,0 +1,173 @@
+diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
+index 5964a4d..53436d3 100644
+--- a/libcuda/cuda_runtime_api.cc
++++ b/libcuda/cuda_runtime_api.cc
+@@ -182,6 +182,8 @@ cudaError_t g_last_cudaError = cudaSuccess;
+ 
+ extern stream_manager *g_stream_manager;
+ 
++
++
+ void register_ptx_function( const char *name, function_info *impl )
+ {
+ 	// no longer need this
+@@ -1341,6 +1343,29 @@ std::string get_app_binary(){
+    return self_exe_path; 
+ }
+ 
++static int get_app_cuda_version() {
++    int app_cuda_version = 0;
++    char fname[1024];
++    snprintf(fname,1024,"_app_cuda_version_XXXXXX");
++    int fd=mkstemp(fname);
++    close(fd);
++    std::string app_cuda_version_command = "ldd " + get_app_binary() + " | grep libcudart.so | sed  's/.*libcudart.so.\\(.*\\) =>.*/\\1/' > " + fname;
++    system(app_cuda_version_command.c_str());
++    FILE * cmd = fopen(fname, "r");
++    char buf[256];
++    while (fgets(buf, sizeof(buf), cmd) != 0) {
++        std::cout << buf;
++        app_cuda_version = atoi(buf);
++    }
++    fclose(cmd);
++    if ( app_cuda_version == 0 ) {
++        printf( "Error - Cannot detect the app's CUDA version.\n" );
++        exit(1);
++    }
++    return app_cuda_version;
++}
++
++
+ //! Call cuobjdump to extract everything (-elf -sass -ptx)
+ /*!
+  *	This Function extract the whole PTX (for all the files) using cuobjdump
+@@ -1366,23 +1391,10 @@ void extract_code_using_cuobjdump(){
+ 	// Running cuobjdump using dynamic link to current process
+ 	// Needs the option '-all' to extract PTX from CDP-enabled binary 
+ 	extern bool g_cdp_enabled;
+-    const int current_cuda = atoi(getenv("CUDA_VERSION_NUMBER"));
+-    if ( context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() && current_cuda > 4020 ) {
+-        const char* cuda_42_path = getenv("CUDA_42_INSTALL_PATH");
+-        if ( NULL == cuda_42_path ) {
+-            printf("You are using a new version of CUDA and PTXPLUS. You are required to have CUDA SDK 4.2 and explicitly "
+-                   "point to it via the environment variable $CUDA_42_INSTALL_PATH.\n\n");
+-            exit(1);
+-        } else {
+-            command = "$CUDA_42_INSTALL_PATH/bin/cuobjdump";
+-        }
+-    } else {
+-        command = "$CUDA_42_INSTALL_PATH/bin/cuobjdump";
+-    }
+ 	if(!g_cdp_enabled)
+-        command += " -ptx -elf -sass " + app_binary + " > " + fname;
++        command = "$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass " + app_binary + " > " + fname;
+ 	else
+-        command += " -ptx -elf -sass -all " + app_binary + " > " + fname;
++        command = "$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass -all " + app_binary + " > " + fname;
+ 	bool parse_output = true; 
+ 	int result = system(command.c_str());
+ 	if(result) {
+@@ -1765,6 +1777,7 @@ void cuobjdumpParseBinary(unsigned int handle){
+     //TODO: Remove temporarily files as per configurations
+ }
+ 
++
+ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
+ {
+ #if (CUDART_VERSION < 2010)
+@@ -1783,26 +1796,10 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
+         // Making this a runtime variable based on the app, enables GPGPU-Sim compiled
+         // with a newer version of CUDA to run apps compiled with older versions of
+         // CUDA. This is especially useful for PTXPLUS execution.
+-        char fname[1024];
+-        snprintf(fname,1024,"_app_cuda_version_XXXXXX");
+-        int fd=mkstemp(fname);
+-        close(fd);
+-        int app_cuda_version = 0;
+-        std::string app_cuda_version_command = "ldd " + get_app_binary() + " | grep libcudart.so | sed  's/.*libcudart.so.\\(.*\\) =>.*/\\1/' > " + fname;
+-        system(app_cuda_version_command.c_str());
+-        FILE * cmd = fopen(fname, "r");
+-        char buf[256];
+-        while (fgets(buf, sizeof(buf), cmd) != 0) {
+-            std::cout << buf;
+-            app_cuda_version = atoi(buf);
+-        }
+-        fclose(cmd);
+-        if ( app_cuda_version == 0 ) {
+-            printf( "Error - Cannot detect the app's CUDA version.\n" );
+-            exit(1);
+-        }
++        int app_cuda_version = get_app_cuda_version();
++        assert( app_cuda_version == CUDART_VERSION / 1000  && "The app must be compiled with same major version as the simulator." );
+         const char* filename;
+-        if ( app_cuda_version < 6 ) {
++#if CUDART_VERSION < 6000
+             // FatBin handle from the .fatbin.c file (one of the intermediate files generated by NVCC)
+             typedef struct {int m; int v; const unsigned long long* d; char* f;} __fatDeviceText __attribute__ ((aligned (8))); 
+             __fatDeviceText * fatDeviceText = (__fatDeviceText *) fatCubin;
+@@ -1814,9 +1811,9 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
+             char * pfatbin = (char*) fatDeviceText->d; 
+             int offset = *((int*)(pfatbin+48)); 
+             filename = (pfatbin+16+offset);
+-        } else {
++#else
+             filename = "default";
+-        }
++#endif
+ 
+ 		// The extracted file name is associated with a fat_cubin_handle passed
+ 		// into cudaLaunch().  Inside cudaLaunch(), the associated file name is
+diff --git a/setup_environment b/setup_environment
+index 8bcb9aa..e5cd4fa 100644
+--- a/setup_environment
++++ b/setup_environment
+@@ -134,18 +134,26 @@ else
+ 	echo "configured without a power model.";
+ fi
+ 
++export PTXAS_CUDA_INSTALL_PATH=$CUDA_INSTALL_PATH
++echo "";
++echo "----------------------------------------------------------------------------";
++echo "WARNING - If you only care about PTX execution, ignore this warning. GPGPU-Sim supports PTX execution in modern CUDA."
+ if [ $CUDA_VERSION_NUMBER -gt 4200 ]; then
+-    echo "";
+-    echo "----------------------------------------------------------------------------";
+-    echo "WARNING - If you only care about PTX execution, ignore this warning. GPGPU-Sim supports PTX execution in modern CUDA."
+-    echo "If you want to run PTXPLUS, download SDK 4.2 (in addition to your modern CUDA) and set \$CUDA_42_INSTALL_PATH."
+-    echo "The following text describes why:";
+-    echo "If you are using PTXPLUS, only sm_1x is supported and it requires that the app binaries are compiled in CUDA 4.2 or less.";
+-    echo "New versions of CUDA tools have dropped parsing support for sm_1x - therefore to use PTXPLUS with a modern card configuration, you must";
+-    echo "have both CUDA 4.2 and a modern CUDA installed. When running PTXPLUS under a CUDA version >= 4.2, the simulator will fail unless \$CUDA_42_INSTALL_PATH";
+-    echo "has been set properly. You must set this manually.";
+-    echo "----------------------------------------------------------------------------"
++    echo "If you want to run PTXPLUS (sm_1x SASS) with a modern card configuration, the apps and simulator must be compiled with CUDA 4.2."
++    echo "You can still run a PASCAL configuration when compiling with 4.2 by setting the \$PTXAS_CUDA_INSTALL_PATH directory environment variable."
++else
++    echo "If you want to run PTXPLUS (sm_1x SASS) with a modern card configuration - set the envronment variable"
++    echo "\$PTXAS_CUDA_INSTALL_PATH to point a CUDA version compabible with your card configurations (i.e. 8+ for PASCAL, 9+ for VOLTA etc..)"
++    echo "For example: \"export \$PTXAS_CUDA_INSTALL_PATH=/usr/local/cuda-9.1\""
+ fi
++echo "The following text describes why:";
++echo "If you are using PTXPLUS, only sm_1x is supported and it requires that the app and simulator binaries are compiled in CUDA 4.2 or less.";
++echo "The simulator requires it since CUDA headers desribe struct sizes in the exec which change from gen to gen.";
++echo "The apps require 4.2 because new versions of CUDA tools have dropped parsing support for generating sm_1x";
++echo "When running using modern config (i.e. pascal) and PTXPLUS with CUDA 4.2, the \$PTXAS_CUDA_INSTALL_PATH env variable is required to get proper register usage"
++echo "(and hence occupancy) using a version of CUDA that knows the register usage on the real card."
++echo "";
++echo "----------------------------------------------------------------------------";
+ 
+ echo "setup_environment succeeded";
+ 
+diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
+index 22f03b9..d8e04d4 100644
+--- a/src/cuda-sim/ptx_loader.cc
++++ b/src/cuda-sim/ptx_loader.cc
+@@ -339,7 +339,7 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
+         snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",g_occupancy_sm_number);
+ #endif
+ 
+-    snprintf(commandline,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file  /dev/null 2> %s",
++    snprintf(commandline,1024,"$PTXAS_CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file  /dev/null 2> %s",
+              extra_flags, fname2, tempfile_ptxinfo);
+     printf("GPGPU-Sim PTX: generating ptxinfo using \"%s\"\n", commandline);
+     result = system(commandline);
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 5964a4d..53436d3 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -182,6 +182,8 @@ cudaError_t g_last_cudaError = cudaSuccess;
 
 extern stream_manager *g_stream_manager;
 
+
+
 void register_ptx_function( const char *name, function_info *impl )
 {
 	// no longer need this
@@ -1341,6 +1343,29 @@ std::string get_app_binary(){
    return self_exe_path; 
 }
 
+static int get_app_cuda_version() {
+    int app_cuda_version = 0;
+    char fname[1024];
+    snprintf(fname,1024,"_app_cuda_version_XXXXXX");
+    int fd=mkstemp(fname);
+    close(fd);
+    std::string app_cuda_version_command = "ldd " + get_app_binary() + " | grep libcudart.so | sed  's/.*libcudart.so.\\(.*\\) =>.*/\\1/' > " + fname;
+    system(app_cuda_version_command.c_str());
+    FILE * cmd = fopen(fname, "r");
+    char buf[256];
+    while (fgets(buf, sizeof(buf), cmd) != 0) {
+        std::cout << buf;
+        app_cuda_version = atoi(buf);
+    }
+    fclose(cmd);
+    if ( app_cuda_version == 0 ) {
+        printf( "Error - Cannot detect the app's CUDA version.\n" );
+        exit(1);
+    }
+    return app_cuda_version;
+}
+
+
 //! Call cuobjdump to extract everything (-elf -sass -ptx)
 /*!
  *	This Function extract the whole PTX (for all the files) using cuobjdump
@@ -1366,23 +1391,10 @@ void extract_code_using_cuobjdump(){
 	// Running cuobjdump using dynamic link to current process
 	// Needs the option '-all' to extract PTX from CDP-enabled binary 
 	extern bool g_cdp_enabled;
-    const int current_cuda = atoi(getenv("CUDA_VERSION_NUMBER"));
-    if ( context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() && current_cuda > 4020 ) {
-        const char* cuda_42_path = getenv("CUDA_42_INSTALL_PATH");
-        if ( NULL == cuda_42_path ) {
-            printf("You are using a new version of CUDA and PTXPLUS. You are required to have CUDA SDK 4.2 and explicitly "
-                   "point to it via the environment variable $CUDA_42_INSTALL_PATH.\n\n");
-            exit(1);
-        } else {
-            command = "$CUDA_42_INSTALL_PATH/bin/cuobjdump";
-        }
-    } else {
-        command = "$CUDA_42_INSTALL_PATH/bin/cuobjdump";
-    }
 	if(!g_cdp_enabled)
-        command += " -ptx -elf -sass " + app_binary + " > " + fname;
+        command = "$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass " + app_binary + " > " + fname;
 	else
-        command += " -ptx -elf -sass -all " + app_binary + " > " + fname;
+        command = "$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass -all " + app_binary + " > " + fname;
 	bool parse_output = true; 
 	int result = system(command.c_str());
 	if(result) {
@@ -1765,6 +1777,7 @@ void cuobjdumpParseBinary(unsigned int handle){
     //TODO: Remove temporarily files as per configurations
 }
 
+
 void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
 {
 #if (CUDART_VERSION < 2010)
@@ -1783,26 +1796,10 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
         // Making this a runtime variable based on the app, enables GPGPU-Sim compiled
         // with a newer version of CUDA to run apps compiled with older versions of
         // CUDA. This is especially useful for PTXPLUS execution.
-        char fname[1024];
-        snprintf(fname,1024,"_app_cuda_version_XXXXXX");
-        int fd=mkstemp(fname);
-        close(fd);
-        int app_cuda_version = 0;
-        std::string app_cuda_version_command = "ldd " + get_app_binary() + " | grep libcudart.so | sed  's/.*libcudart.so.\\(.*\\) =>.*/\\1/' > " + fname;
-        system(app_cuda_version_command.c_str());
-        FILE * cmd = fopen(fname, "r");
-        char buf[256];
-        while (fgets(buf, sizeof(buf), cmd) != 0) {
-            std::cout << buf;
-            app_cuda_version = atoi(buf);
-        }
-        fclose(cmd);
-        if ( app_cuda_version == 0 ) {
-            printf( "Error - Cannot detect the app's CUDA version.\n" );
-            exit(1);
-        }
+        int app_cuda_version = get_app_cuda_version();
+        assert( app_cuda_version == CUDART_VERSION / 1000  && "The app must be compiled with same major version as the simulator." );
         const char* filename;
-        if ( app_cuda_version < 6 ) {
+#if CUDART_VERSION < 6000
             // FatBin handle from the .fatbin.c file (one of the intermediate files generated by NVCC)
             typedef struct {int m; int v; const unsigned long long* d; char* f;} __fatDeviceText __attribute__ ((aligned (8))); 
             __fatDeviceText * fatDeviceText = (__fatDeviceText *) fatCubin;
@@ -1814,9 +1811,9 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
             char * pfatbin = (char*) fatDeviceText->d; 
             int offset = *((int*)(pfatbin+48)); 
             filename = (pfatbin+16+offset);
-        } else {
+#else
             filename = "default";
-        }
+#endif
 
 		// The extracted file name is associated with a fat_cubin_handle passed
 		// into cudaLaunch().  Inside cudaLaunch(), the associated file name is
diff --git a/setup_environment b/setup_environment
index 8bcb9aa..e5cd4fa 100644
--- a/setup_environment
+++ b/setup_environment
@@ -134,18 +134,26 @@ else
 	echo "configured without a power model.";
 fi
 
+export PTXAS_CUDA_INSTALL_PATH=$CUDA_INSTALL_PATH
+echo "";
+echo "----------------------------------------------------------------------------";
+echo "WARNING - If you only care about PTX execution, ignore this warning. GPGPU-Sim supports PTX execution in modern CUDA."
 if [ $CUDA_VERSION_NUMBER -gt 4200 ]; then
-    echo "";
-    echo "----------------------------------------------------------------------------";
-    echo "WARNING - If you only care about PTX execution, ignore this warning. GPGPU-Sim supports PTX execution in modern CUDA."
-    echo "If you want to run PTXPLUS, download SDK 4.2 (in addition to your modern CUDA) and set \$CUDA_42_INSTALL_PATH."
-    echo "The following text describes why:";
-    echo "If you are using PTXPLUS, only sm_1x is supported and it requires that the app binaries are compiled in CUDA 4.2 or less.";
-    echo "New versions of CUDA tools have dropped parsing support for sm_1x - therefore to use PTXPLUS with a modern card configuration, you must";
-    echo "have both CUDA 4.2 and a modern CUDA installed. When running PTXPLUS under a CUDA version >= 4.2, the simulator will fail unless \$CUDA_42_INSTALL_PATH";
-    echo "has been set properly. You must set this manually.";
-    echo "----------------------------------------------------------------------------"
+    echo "If you want to run PTXPLUS (sm_1x SASS) with a modern card configuration, the apps and simulator must be compiled with CUDA 4.2."
+    echo "You can still run a PASCAL configuration when compiling with 4.2 by setting the \$PTXAS_CUDA_INSTALL_PATH directory environment variable."
+else
+    echo "If you want to run PTXPLUS (sm_1x SASS) with a modern card configuration - set the envronment variable"
+    echo "\$PTXAS_CUDA_INSTALL_PATH to point a CUDA version compabible with your card configurations (i.e. 8+ for PASCAL, 9+ for VOLTA etc..)"
+    echo "For example: \"export \$PTXAS_CUDA_INSTALL_PATH=/usr/local/cuda-9.1\""
 fi
+echo "The following text describes why:";
+echo "If you are using PTXPLUS, only sm_1x is supported and it requires that the app and simulator binaries are compiled in CUDA 4.2 or less.";
+echo "The simulator requires it since CUDA headers desribe struct sizes in the exec which change from gen to gen.";
+echo "The apps require 4.2 because new versions of CUDA tools have dropped parsing support for generating sm_1x";
+echo "When running using modern config (i.e. pascal) and PTXPLUS with CUDA 4.2, the \$PTXAS_CUDA_INSTALL_PATH env variable is required to get proper register usage"
+echo "(and hence occupancy) using a version of CUDA that knows the register usage on the real card."
+echo "";
+echo "----------------------------------------------------------------------------";
 
 echo "setup_environment succeeded";
 
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 22f03b9..d8e04d4 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -339,7 +339,7 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
         snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",g_occupancy_sm_number);
 #endif
 
-    snprintf(commandline,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file  /dev/null 2> %s",
+    snprintf(commandline,1024,"$PTXAS_CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file  /dev/null 2> %s",
              extra_flags, fname2, tempfile_ptxinfo);
     printf("GPGPU-Sim PTX: generating ptxinfo using \"%s\"\n", commandline);
     result = system(commandline);
-- 
cgit v1.3


From 263fd9540b90be452640b54c19c7403c2281f3d9 Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Sat, 30 Jun 2018 10:17:18 -0400
Subject: Preventing bad git commit numbers from printing when you run
 setup_env from a dir that is not the gpgpusim_distr dir

---
 setup_environment | 3 +++
 1 file changed, 3 insertions(+)

(limited to 'setup_environment')

diff --git a/setup_environment b/setup_environment
index e5cd4fa..65f8340 100644
--- a/setup_environment
+++ b/setup_environment
@@ -7,10 +7,13 @@ export GPGPUSIM_ROOT="$( cd "$( dirname "$BASH_SOURCE" )" && pwd )"
 
 GPGPUSIM_VERSION_STRING=`cat $GPGPUSIM_ROOT/version | awk '/Version/ {print $8}'`
 #Detect Git branch and commit #
+CURRENT_PWD=`pwd`
+cd $GPGPUSIM_ROOT
 GIT_COMMIT=`git log -n 1 | head -1 | sed -re 's/commit (.*)/\1/'`
 GIT_FILES_CHANGED=`git diff --numstat | wc | sed -re 's/^\s+([0-9]+).*/\1./'`
 GIT_FILES_CHANGED+=`git diff --numstat --cached | wc | sed -re 's/^\s+([0-9]+).*/\1/'`
 GPGPUSIM_BUILD_STRING="gpgpu-sim_git-commit-$GIT_COMMIT-modified_$GIT_FILES_CHANGED"
+cd $CURRENT_PWD
 
 echo -n "GPGPU-Sim version $GPGPUSIM_VERSION_STRING (build $GPGPUSIM_BUILD_STRING) ";
 
-- 
cgit v1.3


From a72f1d829205b9bfc4070cbb2ff46e7d7aba1fc6 Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Sat, 30 Jun 2018 12:54:21 -0400
Subject: Cleaning up the setup env message a bit and having the simulator
 properly dump to stderr on error

---
 setup_environment          |  2 +-
 src/cuda-sim/ptx_loader.cc | 26 +++++++++++++-------------
 2 files changed, 14 insertions(+), 14 deletions(-)

(limited to 'setup_environment')

diff --git a/setup_environment b/setup_environment
index 65f8340..55f3ea6 100644
--- a/setup_environment
+++ b/setup_environment
@@ -140,7 +140,7 @@ fi
 export PTXAS_CUDA_INSTALL_PATH=$CUDA_INSTALL_PATH
 echo "";
 echo "----------------------------------------------------------------------------";
-echo "WARNING - If you only care about PTX execution, ignore this warning. GPGPU-Sim supports PTX execution in modern CUDA."
+echo "INFO - If you only care about PTX execution, ignore this message. GPGPU-Sim supports PTX execution in modern CUDA."
 if [ $CUDA_VERSION_NUMBER -gt 4200 ]; then
     echo "If you want to run PTXPLUS (sm_1x SASS) with a modern card configuration, the apps and simulator must be compiled with CUDA 4.2."
     echo "You can still run a PASCAL configuration when compiling with 4.2 by setting the \$PTXAS_CUDA_INSTALL_PATH directory environment variable."
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index d8e04d4..11d9013 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -125,7 +125,7 @@ char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptxfilenam
 	fflush(stdout);
 	printf("GPGPU-Sim PTX: calling cuobjdump_to_ptxplus\ncommandline: %s\n", commandline);
 	result = system(commandline);
-	if(result){printf("GPGPU-Sim PTX: ERROR ** could not execute %s\n", commandline); exit(1);}
+	if(result){fprintf(stderr, "GPGPU-Sim PTX: ERROR ** could not execute %s\n", commandline); exit(1);}
 
 
 	// Get ptxplus from file
@@ -147,7 +147,7 @@ char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptxfilenam
 		printf("GPGPU-Sim PTX: removing temporary files using \"%s\"\n", rm_commandline);
 		int rm_result = system(rm_commandline);
 		if( rm_result != 0 ) {
-			printf("GPGPU-Sim PTX: ERROR ** while removing temporary files %d\n", rm_result);
+			fprintf(stderr, "GPGPU-Sim PTX: ERROR ** while removing temporary files %d\n", rm_result);
 			exit(1);
 		}
 	}
@@ -198,7 +198,7 @@ void fix_duplicate_errors(char fname2[1024]) {
 	printf("Running: %s\n", commandline);
 	int result = system(commandline);
 	if (result != 0) {
-		printf("GPGPU-Sim PTX: ERROR ** while changing filename from %s to %s", fname2, tempfile);
+		fprintf(stderr, "GPGPU-Sim PTX: ERROR ** while changing filename from %s to %s", fname2, tempfile);
 		exit(1);
 	}
 
@@ -287,7 +287,7 @@ void fix_duplicate_errors(char fname2[1024]) {
 	printf("Running: %s\n", commandline);
 	result = system(commandline);
 	if (result != 0) {
-		printf("GPGPU-Sim PTX: ERROR ** while deleting %s", tempfile);
+		fprintf(stderr, "GPGPU-Sim PTX: ERROR ** while deleting %s", tempfile);
 		exit(1);
 	}
 }
@@ -313,9 +313,9 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
     printf("Running: %s\n", commandline2);
     int result = system(commandline2);
     if( result != 0 ) {
-       printf("GPGPU-Sim PTX: ERROR ** while loading PTX (a) %d\n", result);
-       printf("               Ensure you have write access to simulation directory\n");
-       printf("               and have \'cat\' and \'sed\' in your path.\n");
+       fprintf(stderr, "GPGPU-Sim PTX: ERROR ** while loading PTX (a) %d\n", result);
+       fprintf(stderr, "               Ensure you have write access to simulation directory\n");
+       fprintf(stderr, "               and have \'cat\' and \'sed\' in your path.\n");
        exit(1);
     }
 
@@ -327,9 +327,9 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
 
 #if CUDART_VERSION >= 3000
     if ( g_occupancy_sm_number == 0 ) {
-        printf( "gpgpusim.config must specify the sm version for the GPU that you use to compute occupancy \"-gpgpu_occupancy_sm_number XX\".\n"
-                "The register file size is specifically tied to the sm version used to querry ptxas for register usage.\n"
-                "A register size/SM mismatch may result in occupancy differences." );
+        fprintf( stderr, "gpgpusim.config must specify the sm version for the GPU that you use to compute occupancy \"-gpgpu_occupancy_sm_number XX\".\n"
+                         "The register file size is specifically tied to the sm version used to querry ptxas for register usage.\n"
+                         "A register size/SM mismatch may result in occupancy differences." );
         exit(1);
     }
     extern bool g_cdp_enabled;
@@ -357,8 +357,8 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
     		result = system(commandline);
 	}
 	if (result != 0) {
-		printf("GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result);
-		printf("               Ensure ptxas is in your path.\n");
+		fprintf(stderr, "GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result);
+		fprintf(stderr, "               Ensure ptxas is in your path.\n");
 		exit(1);
 	}
     }
@@ -372,7 +372,7 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
         printf("GPGPU-Sim PTX: removing ptxinfo using \"%s\"\n", commandline);
         result = system(commandline);
         if( result != 0 ) {
-    	    printf("GPGPU-Sim PTX: ERROR ** while loading PTX (c) %d\n", result);
+    	    fprintf(stderr, "GPGPU-Sim PTX: ERROR ** while loading PTX (c) %d\n", result);
     	    exit(1);
         }
     }
-- 
cgit v1.3


From c37cd1e3fa9a3e4392ffc5453e9fd1696fd38396 Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Sat, 30 Jun 2018 13:53:16 -0400
Subject: Adding in a check to not blow away the PTXAS variable if you already
 had it set

---
 setup_environment | 4 +++-
 1 file changed, 3 insertions(+), 1 deletion(-)

(limited to 'setup_environment')

diff --git a/setup_environment b/setup_environment
index 55f3ea6..2dafba5 100644
--- a/setup_environment
+++ b/setup_environment
@@ -137,7 +137,9 @@ else
 	echo "configured without a power model.";
 fi
 
-export PTXAS_CUDA_INSTALL_PATH=$CUDA_INSTALL_PATH
+if [ ! -z "$PTXAS_CUDA_INSTALL_PATH" ]; then
+    export PTXAS_CUDA_INSTALL_PATH=$CUDA_INSTALL_PATH;
+fi
 echo "";
 echo "----------------------------------------------------------------------------";
 echo "INFO - If you only care about PTX execution, ignore this message. GPGPU-Sim supports PTX execution in modern CUDA."
-- 
cgit v1.3


From 5a6957fa2a45cd1856550e9426cb6bd743f19e98 Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Sat, 30 Jun 2018 23:25:20 -0400
Subject: oops - you should set this if it does not exist, not the other way
 around...

---
 setup_environment | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

(limited to 'setup_environment')

diff --git a/setup_environment b/setup_environment
index 2dafba5..f1af978 100644
--- a/setup_environment
+++ b/setup_environment
@@ -137,7 +137,7 @@ else
 	echo "configured without a power model.";
 fi
 
-if [ ! -z "$PTXAS_CUDA_INSTALL_PATH" ]; then
+if [ -z "$PTXAS_CUDA_INSTALL_PATH" ]; then
     export PTXAS_CUDA_INSTALL_PATH=$CUDA_INSTALL_PATH;
 fi
 echo "";
-- 
cgit v1.3


From 02f7f42c3b167350ecf85330f97ff991028ac49d Mon Sep 17 00:00:00 2001
From: Deval Shah <deval281shah@gmail.com>
Date: Fri, 9 Nov 2018 21:30:14 -0800
Subject: required environment variables

---
 setup_environment | 3 +++
 1 file changed, 3 insertions(+)

(limited to 'setup_environment')

diff --git a/setup_environment b/setup_environment
index 9578942..0fee641 100644
--- a/setup_environment
+++ b/setup_environment
@@ -4,6 +4,9 @@ ps -p $$ | awk '/bash/ || / sh/ || /zsh/ {exit 1;}' && echo "WARNING ** source s
 
 export GPGPUSIM_SETUP_ENVIRONMENT_WAS_RUN=
 export GPGPUSIM_ROOT="$( cd "$( dirname "$BASH_SOURCE" )" && pwd )"
+export PTX_SIM_USE_PTX_FILE=1.ptx
+export PTX_SIM_KERNELFILE=_1.ptx
+export CUOBJDUMP_SIM_FILE=jj
 
 GPGPUSIM_VERSION_STRING=`cat $GPGPUSIM_ROOT/version | awk '/Version/ {print $8}'`
 #Detect Git branch and commit #
-- 
cgit v1.3


From 1afe21bf29c7753e6e062a120c63f663c1b7c9d5 Mon Sep 17 00:00:00 2001
From: Deval Shah <deval281shah@gmail.com>
Date: Fri, 16 Nov 2018 11:31:31 -0800
Subject: changes for cuda 4.2  support

---
 setup_environment | 8 +++++---
 1 file changed, 5 insertions(+), 3 deletions(-)

(limited to 'setup_environment')

diff --git a/setup_environment b/setup_environment
index 0fee641..b420584 100644
--- a/setup_environment
+++ b/setup_environment
@@ -4,9 +4,6 @@ ps -p $$ | awk '/bash/ || / sh/ || /zsh/ {exit 1;}' && echo "WARNING ** source s
 
 export GPGPUSIM_SETUP_ENVIRONMENT_WAS_RUN=
 export GPGPUSIM_ROOT="$( cd "$( dirname "$BASH_SOURCE" )" && pwd )"
-export PTX_SIM_USE_PTX_FILE=1.ptx
-export PTX_SIM_KERNELFILE=_1.ptx
-export CUOBJDUMP_SIM_FILE=jj
 
 GPGPUSIM_VERSION_STRING=`cat $GPGPUSIM_ROOT/version | awk '/Version/ {print $8}'`
 #Detect Git branch and commit #
@@ -59,6 +56,11 @@ if [ $CUDA_VERSION_NUMBER -gt 9100 -o $CUDA_VERSION_NUMBER -lt 2030  ]; then
 	return
 fi
 
+if [ $CUDA_VERSION_NUMBER -ge 6000 ]; then
+	export PTX_SIM_USE_PTX_FILE=1.ptx
+	export PTX_SIM_KERNELFILE=_1.ptx
+	export CUOBJDUMP_SIM_FILE=jj
+fi
 
 
 if [ $# = '1' ] ;
-- 
cgit v1.3


From c7f515f6f5325c65f32dd64e1ad479660c751e99 Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
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 'setup_environment')

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 d26a6f5573c0cbeb2624cd4170729acf5d187679 Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Mon, 10 Jun 2019 10:08:33 -0400
Subject: Upgrading the tests to regress CUDA 10 and completely dropping GTX480
 - we should no longer care about this card. We will now regress TITANV with
 PTXPlus instead

---
 Jenkinsfile       | 48 +++++++++++++++++++++++-------------------------
 setup_environment | 14 +++++---------
 2 files changed, 28 insertions(+), 34 deletions(-)

(limited to 'setup_environment')

diff --git a/Jenkinsfile b/Jenkinsfile
index 2c06f49..a352d6a 100644
--- a/Jenkinsfile
+++ b/Jenkinsfile
@@ -14,8 +14,8 @@ pipeline {
                     sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/4.2_env_setup.sh &&\
                     source `pwd`/setup_environment &&\
                     make -j'
-                }, "9.1" : {
-                    sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
+                }, "10.1" : {
+                    sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/10.1_env_setup.sh &&\
                     source `pwd`/setup_environment &&\
                     make -j'
                 }
@@ -34,7 +34,7 @@ pipeline {
                     source ./benchmarks/src/setup_environment && \
                     make -j -C ./benchmarks/src rodinia_2.0-ft sdk-4.2 && \
                     make -C ./benchmarks/src data'
-                sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
+                sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/10.1_env_setup.sh &&\
                     source `pwd`/setup_environment &&\
                     cd gpgpu-sim_simulations && \
                     source ./benchmarks/src/setup_environment && \
@@ -47,17 +47,18 @@ pipeline {
                 parallel "4.2-rodinia": {
                     sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/4.2_env_setup.sh &&\
                     source `pwd`/setup_environment &&\
-                    ./gpgpu-sim_simulations/util/job_launching/run_simulations.py -B rodinia_2.0-ft -C GTX480,GTX480-PTXPLUS -N regress-$$ && \
+                    export PTXAS_CUDA_INSTALL_PATH=/home/tgrogers-raid/a/common/cuda-10.1 &&\
+                    ./gpgpu-sim_simulations/util/job_launching/run_simulations.py -B rodinia_2.0-ft -C TITANV-PTXPLUS -N regress-$$ && \
                     PLOTDIR="jenkins/${JOB_NAME}/${BUILD_NUMBER}/4.2-rodinia" && ssh tgrogers@dynamo.ecn.purdue.edu mkdir -p /home/dynamo/a/tgrogers/website/gpgpu-sim-plots/$PLOTDIR && \
                     ./gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -v -N regress-$$ -s stats-per-app-4.2.csv && \
                     ./gpgpu-sim_simulations/util/plotting/plot-get-stats.py -c stats-per-app-4.2.csv -p tgrogers@dynamo.ecn.purdue.edu:~/website/gpgpu-sim-plots/$PLOTDIR -w https://engineering.purdue.edu/tgrogers/gpgpu-sim-plots/$PLOTDIR -n $PLOTDIR '
-                }, "9.1-functest": {
-                    sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
+                }, "10.1-functest": {
+                    sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/10.1_env_setup.sh &&\
                     source `pwd`/setup_environment &&\
                     ./gpgpu-sim_simulations/util/job_launching/run_simulations.py -B rodinia_2.0-ft,sdk-4.2 -C TITANV -N regress-$$ && \
-                    PLOTDIR="jenkins/${JOB_NAME}/${BUILD_NUMBER}/9.1-rodinia" && ssh tgrogers@dynamo.ecn.purdue.edu mkdir -p /home/dynamo/a/tgrogers/website/gpgpu-sim-plots/$PLOTDIR && \
-                    ./gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -v  -s stats-per-app-9.1.csv -N regress-$$ && \
-                    ./gpgpu-sim_simulations/util/plotting/plot-get-stats.py -c stats-per-app-9.1.csv -p tgrogers@dynamo.ecn.purdue.edu:~/website/gpgpu-sim-plots/$PLOTDIR -w https://engineering.purdue.edu/tgrogers/gpgpu-sim-plots/$PLOTDIR -n $PLOTDIR'
+                    PLOTDIR="jenkins/${JOB_NAME}/${BUILD_NUMBER}/10.1-rodinia" && ssh tgrogers@dynamo.ecn.purdue.edu mkdir -p /home/dynamo/a/tgrogers/website/gpgpu-sim-plots/$PLOTDIR && \
+                    ./gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -v  -s stats-per-app-10.1.csv -N regress-$$ && \
+                    ./gpgpu-sim_simulations/util/plotting/plot-get-stats.py -c stats-per-app-10.1.csv -p tgrogers@dynamo.ecn.purdue.edu:~/website/gpgpu-sim-plots/$PLOTDIR -w https://engineering.purdue.edu/tgrogers/gpgpu-sim-plots/$PLOTDIR -n $PLOTDIR'
                 }
             }
         }
@@ -66,17 +67,17 @@ pipeline {
                 sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/4.2_env_setup.sh &&\
                     source `pwd`/setup_environment &&\
                     PLOTDIR="jenkins/${JOB_NAME}" &&\
-                    ./gpgpu-sim_simulations/util/job_launching/get_stats.py -R -K -k -B rodinia_2.0-ft -C GTX480,GTX480-PTXPLUS > stats-per-kernel-4.2.csv &&\
+                    ./gpgpu-sim_simulations/util/job_launching/get_stats.py -R -K -k -B rodinia_2.0-ft -C TITANV-PTXPLUS > stats-per-kernel-4.2.csv &&\
                     ./gpgpu-sim_simulations/util/plotting/correlate_and_publish.sh stats-per-kernel-4.2.csv $PLOTDIR ${BUILD_NUMBER} | grep "Correl=" | tee correl.4.2.txt'
             }
         }
-        stage('9.1-correlate'){
+        stage('10.1-correlate'){
             steps {
-                sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
+                sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/10.1_env_setup.sh &&\
                     source `pwd`/setup_environment &&\
                     PLOTDIR="jenkins/${JOB_NAME}" &&\
-                    ./gpgpu-sim_simulations/util/job_launching/get_stats.py -R -K -k -B rodinia_2.0-ft,sdk-4.2 -C TITANV > stats-per-kernel-9.1.csv &&\
-                    ./gpgpu-sim_simulations/util/plotting/correlate_and_publish.sh stats-per-kernel-9.1.csv $PLOTDIR ${BUILD_NUMBER} | grep "Correl=" | tee correl.9.1.txt'
+                    ./gpgpu-sim_simulations/util/job_launching/get_stats.py -R -K -k -B rodinia_2.0-ft,sdk-4.2 -C TITANV > stats-per-kernel-10.1.csv &&\
+                    ./gpgpu-sim_simulations/util/plotting/correlate_and_publish.sh stats-per-kernel-10.1.csv $PLOTDIR ${BUILD_NUMBER} | grep "Correl=" | tee correl.10.1.txt'
             }
         }
         stage('archive-and-delta') {
@@ -84,24 +85,21 @@ pipeline {
                     sh 'rm -rf gpgpu-sim-results-repo'
                     sh 'git clone git@github.com:purdue-aalp/gpgpu-sim-results-repo.git'
                     sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/4.2_env_setup.sh &&\
-                    ./gpgpu-sim_simulations/util/job_launching/get_stats.py -R -K -k -B rodinia_2.0-ft -C GTX480-PTXPLUS > stats-per-kernel-4.2-ptxplus.csv &&\
-                    ./gpgpu-sim_simulations/util/job_launching/get_stats.py -R -K -k -B rodinia_2.0-ft -C GTX480 > stats-per-kernel-4.2-ptx.csv'
-                    sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
-                    ./gpgpu-sim_simulations/util/job_launching/get_stats.py -R -K -k -B rodinia_2.0-ft,sdk-4.2 -C TITANV > stats-per-kernel-9.1.csv'
+                    ./gpgpu-sim_simulations/util/job_launching/get_stats.py -R -K -k -B rodinia_2.0-ft -C TITANV-PTXPLUS > stats-per-kernel-4.2-ptxplus.csv &&\
+                    sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/10.1_env_setup.sh &&\
+                    ./gpgpu-sim_simulations/util/job_launching/get_stats.py -R -K -k -B rodinia_2.0-ft,sdk-4.2 -C TITANV > stats-per-kernel-10.1.csv'
                     sh 'if [ ! -d ./gpgpu-sim-results-repo/${JOB_NAME} ]; then mkdir -p ./gpgpu-sim-results-repo/${JOB_NAME}/ ; cp ./gpgpu-sim-results-repo/purdue-aalp/gpgpu-sim_distribution/dev/* ./gpgpu-sim-results-repo/${JOB_NAME}/ ; fi'
                     sh './gpgpu-sim_simulations/util/plotting/merge-stats.py -c ./gpgpu-sim-results-repo/${JOB_NAME}/stats-per-app-4.2.csv,./stats-per-app-4.2.csv -R > per-app-merge-4.2.csv'
-                    sh './gpgpu-sim_simulations/util/plotting/merge-stats.py -c ./gpgpu-sim-results-repo/${JOB_NAME}/stats-per-app-9.1.csv,./stats-per-app-9.1.csv -R > per-app-merge-9.1.csv'
+                    sh './gpgpu-sim_simulations/util/plotting/merge-stats.py -c ./gpgpu-sim-results-repo/${JOB_NAME}/stats-per-app-10.1.csv,./stats-per-app-10.1.csv -R > per-app-merge-10.1.csv'
                     sh 'PLOTDIR="jenkins/${JOB_NAME}" &&\
                         ssh tgrogers@dynamo.ecn.purdue.edu mkdir -p /home/dynamo/a/tgrogers/website/gpgpu-sim-plots/$PLOTDIR/${BUILD_NUMBER}/deltas/4.2 && \
-                        ssh tgrogers@dynamo.ecn.purdue.edu mkdir -p /home/dynamo/a/tgrogers/website/gpgpu-sim-plots/$PLOTDIR/${BUILD_NUMBER}/deltas/9.1 && \
+                        ssh tgrogers@dynamo.ecn.purdue.edu mkdir -p /home/dynamo/a/tgrogers/website/gpgpu-sim-plots/$PLOTDIR/${BUILD_NUMBER}/deltas/10.1 && \
                         ./gpgpu-sim_simulations/util/plotting/plot-get-stats.py -c per-app-merge-4.2.csv -p tgrogers@dynamo.ecn.purdue.edu:~/website/gpgpu-sim-plots/$PLOTDIR/${BUILD_NUMBER}/deltas/4.2 -w https://engineering.purdue.edu/tgrogers/gpgpu-sim-plots/$PLOTDIR/${BUILD_NUMBER}/deltas/4.2 -n $PLOTDIR/${BUILD_NUMBER}/deltas/4.2 &&\
-                        ./gpgpu-sim_simulations/util/plotting/plot-get-stats.py -c per-app-merge-9.1.csv -p tgrogers@dynamo.ecn.purdue.edu:~/website/gpgpu-sim-plots/$PLOTDIR/${BUILD_NUMBER}/deltas/9.1 -w https://engineering.purdue.edu/tgrogers/gpgpu-sim-plots/$PLOTDIR/${BUILD_NUMBER}/deltas/9.1 -n $PLOTDIR/${BUILD_NUMBER}/deltas/9.1 &&\
-                        ./gpgpu-sim_simulations/util/plotting/merge-stats.py -c ./gpgpu-sim-results-repo/${JOB_NAME}/stats-per-kernel-4.2-ptx.csv,./stats-per-kernel-4.2-ptx.csv -R > per-kernel-merge-4.2-ptx.csv &&\
+                        ./gpgpu-sim_simulations/util/plotting/plot-get-stats.py -c per-app-merge-10.1.csv -p tgrogers@dynamo.ecn.purdue.edu:~/website/gpgpu-sim-plots/$PLOTDIR/${BUILD_NUMBER}/deltas/10.1 -w https://engineering.purdue.edu/tgrogers/gpgpu-sim-plots/$PLOTDIR/${BUILD_NUMBER}/deltas/10.1 -n $PLOTDIR/${BUILD_NUMBER}/deltas/10.1 &&\
                         ./gpgpu-sim_simulations/util/plotting/merge-stats.py -c ./gpgpu-sim-results-repo/${JOB_NAME}/stats-per-kernel-4.2-ptxplus.csv,./stats-per-kernel-4.2-ptxplus.csv -R > per-kernel-merge-4.2-ptxplus.csv &&\
-                        ./gpgpu-sim_simulations/util/plotting/merge-stats.py -c ./gpgpu-sim-results-repo/${JOB_NAME}/stats-per-kernel-9.1.csv,./stats-per-kernel-9.1.csv -R > per-kernel-merge-9.1.csv &&\
-                        ./gpgpu-sim_simulations/util/plotting/correlate_and_publish.sh per-kernel-merge-4.2-ptx.csv $PLOTDIR ${BUILD_NUMBER} &&\
+                        ./gpgpu-sim_simulations/util/plotting/merge-stats.py -c ./gpgpu-sim-results-repo/${JOB_NAME}/stats-per-kernel-10.1.csv,./stats-per-kernel-10.1.csv -R > per-kernel-merge-10.1.csv &&\
                         ./gpgpu-sim_simulations/util/plotting/correlate_and_publish.sh per-kernel-merge-4.2-ptxplus.csv $PLOTDIR ${BUILD_NUMBER} &&\
-                        ./gpgpu-sim_simulations/util/plotting/correlate_and_publish.sh per-kernel-merge-9.1.csv $PLOTDIR ${BUILD_NUMBER} &&\
+                        ./gpgpu-sim_simulations/util/plotting/correlate_and_publish.sh per-kernel-merge-10.1.csv $PLOTDIR ${BUILD_NUMBER} &&\
                         mkdir -p ./gpgpu-sim-results-repo/${JOB_NAME}/ && cp stats-per-*.csv ./gpgpu-sim-results-repo/${JOB_NAME}/ &&\
                         cd ./gpgpu-sim-results-repo &&\
                         git diff --quiet && git diff --staged --quiet || git commit -am "Jenkins automated checkin ${BUILD_NUMBER}" &&\
diff --git a/setup_environment b/setup_environment
index 17891ce..58f4713 100644
--- a/setup_environment
+++ b/setup_environment
@@ -148,19 +148,15 @@ fi
 echo "";
 echo "----------------------------------------------------------------------------";
 echo "INFO - If you only care about PTX execution, ignore this message. GPGPU-Sim supports PTX execution in modern CUDA."
-if [ $CUDA_VERSION_NUMBER -gt 4200 ]; then
-    echo "If you want to run PTXPLUS (sm_1x SASS) with a modern card configuration, the apps and simulator must be compiled with CUDA 4.2."
-    echo "You can still run a PASCAL configuration when compiling with 4.2 by setting the \$PTXAS_CUDA_INSTALL_PATH directory environment variable."
-else
-    echo "If you want to run PTXPLUS (sm_1x SASS) with a modern card configuration - set the envronment variable"
-    echo "\$PTXAS_CUDA_INSTALL_PATH to point a CUDA version compabible with your card configurations (i.e. 8+ for PASCAL, 9+ for VOLTA etc..)"
-    echo "For example: \"export \$PTXAS_CUDA_INSTALL_PATH=/usr/local/cuda-9.1\""
-fi
+echo "If you want to run PTXPLUS (sm_1x SASS) with a modern card configuration - set the envronment variable"
+echo "\$PTXAS_CUDA_INSTALL_PATH to point a CUDA version compabible with your card configurations (i.e. 8+ for PASCAL, 9+ for VOLTA etc..)"
+echo "For example: \"export \$PTXAS_CUDA_INSTALL_PATH=/usr/local/cuda-9.1\""
+echo ""
 echo "The following text describes why:";
 echo "If you are using PTXPLUS, only sm_1x is supported and it requires that the app and simulator binaries are compiled in CUDA 4.2 or less.";
 echo "The simulator requires it since CUDA headers desribe struct sizes in the exec which change from gen to gen.";
 echo "The apps require 4.2 because new versions of CUDA tools have dropped parsing support for generating sm_1x";
-echo "When running using modern config (i.e. pascal) and PTXPLUS with CUDA 4.2, the \$PTXAS_CUDA_INSTALL_PATH env variable is required to get proper register usage"
+echo "When running using modern config (i.e. volta) and PTXPLUS with CUDA 4.2, the \$PTXAS_CUDA_INSTALL_PATH env variable is required to get proper register usage"
 echo "(and hence occupancy) using a version of CUDA that knows the register usage on the real card."
 echo "";
 echo "----------------------------------------------------------------------------";
-- 
cgit v1.3


From 208654625cf668978d249359fcb85b1f7edd13dd Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Wed, 28 Aug 2019 13:02:33 -0400
Subject: Fixing the screwy way that setup env messes with the cd stack

---
 setup_environment | 9 +++------
 1 file changed, 3 insertions(+), 6 deletions(-)

(limited to 'setup_environment')

diff --git a/setup_environment b/setup_environment
index 58f4713..ca60d6b 100644
--- a/setup_environment
+++ b/setup_environment
@@ -7,13 +7,10 @@ export GPGPUSIM_ROOT="$( cd "$( dirname "$BASH_SOURCE" )" && pwd )"
 
 GPGPUSIM_VERSION_STRING=`cat $GPGPUSIM_ROOT/version | awk '/Version/ {print $8}'`
 #Detect Git branch and commit #
-CURRENT_PWD=`pwd`
-cd $GPGPUSIM_ROOT
-GIT_COMMIT=`git log -n 1 | head -1 | sed -re 's/commit (.*)/\1/'`
-GIT_FILES_CHANGED=`git diff --numstat | wc | sed -re 's/^\s+([0-9]+).*/\1./'`
-GIT_FILES_CHANGED+=`git diff --numstat --cached | wc | sed -re 's/^\s+([0-9]+).*/\1/'`
+GIT_COMMIT=`git --git-dir=$GPGPUSIM_ROOT/.git log -n 1 | head -1 | sed -re 's/commit (.*)/\1/'`
+GIT_FILES_CHANGED=`git --git-dir=$GPGPUSIM_ROOT/.git diff --numstat | wc | sed -re 's/^\s+([0-9]+).*/\1./'`
+GIT_FILES_CHANGED+=`git --git-dir=$GPGPUSIM_ROOT/.git diff --numstat --cached | wc | sed -re 's/^\s+([0-9]+).*/\1/'`
 GPGPUSIM_BUILD_STRING="gpgpu-sim_git-commit-$GIT_COMMIT-modified_$GIT_FILES_CHANGED"
-cd $CURRENT_PWD
 
 echo -n "GPGPU-Sim version $GPGPUSIM_VERSION_STRING (build $GPGPUSIM_BUILD_STRING) ";
 
-- 
cgit v1.3


From 555cce395bb8457a209b26f6f7ff62cc09839fdf Mon Sep 17 00:00:00 2001
From: bftf <f.demoullin@gmail.com>
Date: Mon, 11 Nov 2019 12:29:15 -0800
Subject: Enabled CUDA 10. Added implementations for
 __cudaPushCallConfiguration and __cudaPopCallConfiguration, neither of which
 are documented in the CUDA API but they preceed any cudaLaunchKernel() call
 and push the call configuration on a stack which is then popped when
 cudaLaunchKernel() is called. Currently the max stack size is 1, an assert
 will be triggered when this is violated.

---
 CHANGES                     |  1 +
 Makefile                    |  1 +
 libcuda/cuda_api.h          |  3 ++-
 libcuda/cuda_runtime_api.cc | 54 +++++++++++++++++++++++++++++++++++++++++++--
 linux-so-version.txt        |  2 ++
 setup_environment           |  5 +++--
 6 files changed, 61 insertions(+), 5 deletions(-)

(limited to 'setup_environment')

diff --git a/CHANGES b/CHANGES
index 2a4222b..31ae672 100644
--- a/CHANGES
+++ b/CHANGES
@@ -37,6 +37,7 @@ Version 4.0.0 (development branch) versus 3.2.3
 3- Addig new system stats: gpu occupancy, L2BW, etc
 -Library:
 1  Enabled CUTLASS Library on GPGPU-Sim
+2  Enabled CUDA 10
 -Regression:
 1- Added TensorCore Regression Kernel
 -Configs:
diff --git a/Makefile b/Makefile
index 3db8ce8..8df3b57 100644
--- a/Makefile
+++ b/Makefile
@@ -164,6 +164,7 @@ $(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
 
 $(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..b304a59 100644
--- a/libcuda/cuda_api.h
+++ b/libcuda/cuda_api.h
@@ -234,10 +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_VERSION < 1000
 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 44f0f4e..fd092e5 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -1448,6 +1448,50 @@ __host__ cudaError_t CUDARTAPI cudaConfigureCall(dim3 gridDim, dim3 blockDim, si
 	return g_last_cudaError = cudaSuccess;
 }
 
+
+#if CUDA_VERSION >= 1000
+/*
+* CUDA 10 requires a new CUDA kernel launch sequence
+* A call to __cudaPushCallConfiguration() preceeds any call to cudaLaunchKernel()
+* __cudaPushCallConfiguration is undocumented in the API but it simply sets up a buffer with the arguments which is accessed in cudaLaunchKernel()
+* __cudaPopCallConfiguration is undocumented in the API but it simply pops the configuration set in cudaLaunchKernel()
+*
+* pushing more than 1 configuration without popping is currently not implemented in GPGPU-Sim and will result in an assert error
+*/
+namespace g_cudaPushArgsBuffer
+{
+  bool g_is_initialized = false;
+  dim3 g_gridDim;
+  dim3 g_blockDim;
+  size_t g_sharedMem;
+  cudaStream_t g_stream;
+}
+
+__host__ cudaError_t CUDARTAPI __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim, size_t sharedMem, cudaStream_t stream)
+{
+  assert(g_cudaPushArgsBuffer::g_is_initialized == false);
+  printf("Pushing cuda call configuration \n");
+  g_cudaPushArgsBuffer::g_is_initialized = true;
+  g_cudaPushArgsBuffer::g_gridDim = gridDim;
+  g_cudaPushArgsBuffer::g_blockDim = blockDim;
+  g_cudaPushArgsBuffer::g_sharedMem = sharedMem;
+  g_cudaPushArgsBuffer::g_stream = stream;
+
+  return cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI __cudaPopCallConfiguration()
+{
+  printf("Inside __cudaPopCallConfiguration\n");
+  assert(g_cudaPushArgsBuffer::g_is_initialized == true);
+  printf("Poping cuda call configuration \n");
+  g_cudaPushArgsBuffer::g_is_initialized = false;
+  return cudaSuccess;
+}
+
+#endif // #if CUDA_VERSION >= 1000
+
+
 __host__ cudaError_t CUDARTAPI cudaSetupArgument(const void *arg, size_t size, size_t offset)
 {
 	if(g_debug_execution >= 3){
@@ -1551,8 +1595,14 @@ __host__ cudaError_t CUDARTAPI cudaLaunchKernel ( const char* hostFun, dim3 grid
     	}
         CUctx_st *context = GPGPUSim_Context();
         function_info *entry = context->get_kernel(hostFun);
-    
-	cudaConfigureCall(gridDim, blockDim, sharedMem, stream);
+
+#if CUDA_VERSION >= 1000
+  assert(g_cudaPushArgsBuffer::g_is_initialized == false);
+  cudaConfigureCall(g_cudaPushArgsBuffer::g_gridDim, g_cudaPushArgsBuffer::g_blockDim, g_cudaPushArgsBuffer::g_sharedMem, g_cudaPushArgsBuffer::g_stream);
+#else
+    cudaConfigureCall(gridDim, blockDim, sharedMem, stream);
+#endif // #if CUDA_VERSION >= 1000
+  
     	for(unsigned i = 0; i < entry->num_args(); i++){
         	std::pair<size_t, unsigned> p = entry->get_param_config(i);
         	cudaSetupArgument(args[i], p.first, p.second);
diff --git a/linux-so-version.txt b/linux-so-version.txt
index a7c2d3c..357c6cd 100644
--- a/linux-so-version.txt
+++ b/linux-so-version.txt
@@ -6,3 +6,5 @@ libcudart.so.9.2{
 };
 libcuda.so.1{
 };
+libcudart.so.10.0{
+};
\ No newline at end of file
diff --git a/setup_environment b/setup_environment
index b420584..154dea0 100644
--- a/setup_environment
+++ b/setup_environment
@@ -51,9 +51,10 @@ 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 10000 -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
+	echo $CUDA_VERSION_NUMBER
+  return
 fi
 
 if [ $CUDA_VERSION_NUMBER -ge 6000 ]; then
-- 
cgit v1.3