From 26b9853cc1f6a3b17de0e319b40f28a5703ad6bf Mon Sep 17 00:00:00 2001
From: Amruth <amruth@ece.ubc.ca>
Date: Sun, 25 Mar 2018 16:01:22 -0700
Subject: code for removing duplicates in embedded ptx

---
 src/cuda-sim/ptx_loader.cc | 155 ++++++++++++++++++++++++++++-----------------
 1 file changed, 98 insertions(+), 57 deletions(-)

(limited to 'src/cuda-sim/ptx_loader.cc')

diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 6c1b595..d98ca07 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -32,6 +32,7 @@
 #include <unistd.h>
 #include <dirent.h>
 #include <fstream>
+#include <sstream>
 
 /// globals
 
@@ -287,56 +288,87 @@ void fix_duplicate_errors(char fname2[1024]) {
 	}
 }
 
+//we need the application name here too.
+char* get_app_binary_name(){
+   char exe_path[1025];
+   char *self_exe_path;
+#ifdef __APPLE__
+   //AMRUTH:  get apple device and check the result.
+   printf("WARNING: not tested for Apple-mac devices \n");
+   abort();
+#else
+   std::stringstream exec_link;
+   exec_link << "/proc/self/exe";
+   ssize_t path_length = readlink(exec_link.str().c_str(), exe_path, 1024);
+   assert(path_length != -1);
+   exe_path[path_length] = '\0';
+
+   char *token = strtok(exe_path, "/");
+   while(token !=NULL){
+        self_exe_path = token;
+        token = strtok(NULL,"/");
+   }
+#endif
+   self_exe_path = strtok(self_exe_path, ".");
+   printf("self exe links to: %s\n", self_exe_path);
+   return self_exe_path;
+}
+
 void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version )
 {
-    char fname[1024];
-    snprintf(fname,1024,"_ptx_XXXXXX");
-    int fd=mkstemp(fname); 
-    close(fd);
-
-    printf("GPGPU-Sim PTX: extracting embedded .ptx to temporary file \"%s\"\n", fname);
-    FILE *ptxfile = fopen(fname,"w");
-    fprintf(ptxfile,"%s", p_for_info);
-    fclose(ptxfile);
-
-    char fname2[1024];
-    snprintf(fname2,1024,"_ptx2_XXXXXX");
-    fd=mkstemp(fname2); 
-    close(fd);
-    char commandline2[4096];
-    snprintf(commandline2,4096,"cat %s | sed 's/.version 1.5/.version 1.4/' | sed 's/, texmode_independent//' | sed 's/\\(\\.extern \\.const\\[1\\] .b8 \\w\\+\\)\\[\\]/\\1\\[1\\]/' | sed 's/const\\[.\\]/const\\[0\\]/g' > %s", fname, fname2);
-    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");
-       exit(1);
-    }
-
-    char tempfile_ptxinfo[1024];
-    snprintf(tempfile_ptxinfo,1024,"%sinfo",fname);
-    char commandline[1024];
-    char extra_flags[1024];
-    extra_flags[0]=0;
-
-#if CUDART_VERSION >= 3000
-    if (sm_version == 0) sm_version = 20;
-    extern bool g_cdp_enabled;
-    if(!g_cdp_enabled)
-        snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version);
-    else
-        snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",sm_version);
-#endif
+    //do ptxas for individual files instead of one big embedded ptx. This prevents the duplicate defs and declarations.
+    char ptx_file[1000];
+    char *name=get_app_binary_name();
+    char commandline[4096], fname[1024], fname2[1024];
+    for (int index=1; index <= no_of_ptx; index++){
+        snprintf(ptx_file, 1000, "%s.%d.sm_%u.ptx", name, index, sm_version);
+        snprintf(fname,1024,"_ptx_XXXXXX");
+        int fd=mkstemp(fname); 
+        close(fd);
 
-    snprintf(commandline,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file  /dev/null 2> %s",
+        printf("GPGPU-Sim PTX: extracting embedded .ptx to temporary file \"%s\"\n", fname);
+        snprintf(commandline,4096,"cat %s > %s",ptx_file, fname);
+        if (system(commandline) !=0) {
+            printf("ERROR: %s command failed\n", commandline);
+            exit(0);
+        }
+	
+        snprintf(fname2,1024,"_ptx2_XXXXXX");
+        fd=mkstemp(fname2); 
+        close(fd);
+        char commandline2[4096];
+        snprintf(commandline2,4096,"cat %s | sed 's/.version 1.5/.version 1.4/' | sed 's/, texmode_independent//' | sed 's/\\(\\.extern \\.const\\[1\\] .b8 \\w\\+\\)\\[\\]/\\1\\[1\\]/' | sed 's/const\\[.\\]/const\\[0\\]/g' > %s", fname, fname2);
+        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");
+       	    exit(1);
+    	}
+
+        char tempfile_ptxinfo[1024];
+        snprintf(tempfile_ptxinfo,1024,"%sinfo",fname);
+        char extra_flags[1024];
+        extra_flags[0]=0;
+
+        #if CUDART_VERSION >= 3000
+    	if (sm_version == 0) sm_version = 20;
+    	extern bool g_cdp_enabled;
+    	if(!g_cdp_enabled)
+        	snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version);
+    	else
+        	snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",sm_version);
+	#endif
+
+    	snprintf(commandline,1024,"$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);
-    if( result != 0 ) {
-    	// 65280 = duplicate errors
-    	if (result == 65280) {
-    		ptxinfo_in = fopen(tempfile_ptxinfo,"r");
+   	printf("GPGPU-Sim PTX: generating ptxinfo using \"%s\"\n", commandline);
+    	result = system(commandline);
+    	if( result != 0 ) {
+    	    // 65280 = duplicate errors
+    	    if (result == 65280) {
+    	 	ptxinfo_in = fopen(tempfile_ptxinfo,"r");
 		g_ptxinfo_filename = tempfile_ptxinfo;
 		ptxinfo_parse();
 
@@ -345,26 +377,35 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
     			 extra_flags, fname2, tempfile_ptxinfo);
     		printf("GPGPU-Sim PTX: regenerating ptxinfo using \"%s\"\n", commandline);
     		result = system(commandline);
-	}
-	if (result != 0) {
+	    }
+	    if (result != 0) {
 		printf("GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result);
 		printf("               Ensure ptxas is in your path.\n");
 		exit(1);
-	}
+	    }
+    	}
     }
-
-    ptxinfo_in = fopen(tempfile_ptxinfo,"r");
-    g_ptxinfo_filename = tempfile_ptxinfo;
+    //Now that we got resource usage per kernel in a ptx file, we dump all into one file and pass it to rest of the code as usual.
+    char commandline3[4096];
+    char final_tempfile_ptxinfo[1024];
+    snprintf(final_tempfile_ptxinfo,1024,"f_tempfile_ptx");
+    snprintf(commandline3,4096, "cat *info > %s", final_tempfile_ptxinfo);
+    if (system(commandline3)!=0) {
+	printf("ERROR: Either we dont have info files or cat is not working \n");
+        printf("ERROR: %s command failed\n",commandline3);
+	exit(1);
+    }    
+
+    ptxinfo_in = fopen(final_tempfile_ptxinfo,"r");
+    g_ptxinfo_filename = final_tempfile_ptxinfo;
     ptxinfo_parse();
 
     if( ! g_save_embedded_ptx ) {
-        snprintf(commandline,1024,"rm -f %s %s %s", fname, fname2, tempfile_ptxinfo);
+        snprintf(commandline,1024,"rm -f %s %s %s *info", fname, fname2, final_tempfile_ptxinfo);
         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);
+        if( system(commandline) != 0 ) {
+    	    printf("GPGPU-Sim PTX: ERROR ** while removing temporary files\n");
     	    exit(1);
         }
     }
 }
-
-- 
cgit v1.3


From 0d07b73b58224724cff0576d91878817ec3c01c1 Mon Sep 17 00:00:00 2001
From: Amruth <amruth@ece.ubc.ca>
Date: Tue, 27 Mar 2018 15:52:52 -0700
Subject: considered CDP scenario

---
 libcuda/cuda_runtime_api.cc | 24 +++++++-------
 src/cuda-sim/ptx_loader.cc  | 77 +++++++++++++++++++++++++++++++++++++--------
 2 files changed, 77 insertions(+), 24 deletions(-)

(limited to 'src/cuda-sim/ptx_loader.cc')

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 97d702c..57c5ea1 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -1083,8 +1083,8 @@ __host__ __device__ cudaError_t CUDARTAPI cudaDeviceGetStreamPriorityRange(int*
        	return cudaSuccess;	
 }
 
-__host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *stream, unsigned int flags) {
-	return cudaStreamCreate(stream);
+__host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *pStream, unsigned int flags) {
+	return cudaStreamCreate(pStream);
 }
 
 __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
@@ -1496,14 +1496,16 @@ void extract_code_using_cuobjdump(){
                     fclose(fp);
         }
     }
-    //based on the list above, dump ptx files individually. Format of dumped ptx file is prog_name.unique_no.sm_<>.ptx
-    for (int index=1; index<= no_of_ptx; index++){
-        snprintf(ptx_file, 1000, "%s.%d.sm_%u.ptx", get_app_binary_name(app_binary), index, forced_max_capability);
-        printf("Extracting specific PTX file named %s \n",ptx_file);
-        snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -arch=sm_%u -xptx %s %s", forced_max_capability, ptx_file, app_binary.c_str());
-        if (system(command)!=0) {
-            printf("ERROR: command: %s failed \n",command);
-            exit(0);
+    if(!g_cdp_enabled) {
+        //based on the list above, dump ptx files individually. Format of dumped ptx file is prog_name.unique_no.sm_<>.ptx
+        for (int index=1; index<= no_of_ptx; index++){
+            snprintf(ptx_file, 1000, "%s.%d.sm_%u.ptx", get_app_binary_name(app_binary), index, forced_max_capability);
+            printf("Extracting specific PTX file named %s \n",ptx_file);
+            snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -arch=sm_%u -xptx %s %s", forced_max_capability, ptx_file, app_binary.c_str());
+            if (system(command)!=0) {
+                printf("ERROR: command: %s failed \n",command);
+                exit(0);
+            }
         }
     }
     //TODO: redundant to dump twice. how can it be prevented?
@@ -1516,7 +1518,7 @@ void extract_code_using_cuobjdump(){
 	if(!g_cdp_enabled)
             snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass -arch=sm_%u %s > %s", forced_max_capability, app_binary.c_str(), fname);
 	else
-            snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass -arch=sm_%u -all %s > %s", forced_max_capability, app_binary.c_str(), fname);
+            snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass -all %s > %s", app_binary.c_str(), fname);
 	bool parse_output = true; 
 	result = system(command);
 	if(result) {
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index d98ca07..9ff0859 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -319,7 +319,7 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
     //do ptxas for individual files instead of one big embedded ptx. This prevents the duplicate defs and declarations.
     char ptx_file[1000];
     char *name=get_app_binary_name();
-    char commandline[4096], fname[1024], fname2[1024];
+    char commandline[4096], fname[1024], fname2[1024], final_tempfile_ptxinfo[1024], tempfile_ptxinfo[1024];
     for (int index=1; index <= no_of_ptx; index++){
         snprintf(ptx_file, 1000, "%s.%d.sm_%u.ptx", name, index, sm_version);
         snprintf(fname,1024,"_ptx_XXXXXX");
@@ -347,7 +347,6 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
        	    exit(1);
     	}
 
-        char tempfile_ptxinfo[1024];
         snprintf(tempfile_ptxinfo,1024,"%sinfo",fname);
         char extra_flags[1024];
         extra_flags[0]=0;
@@ -385,23 +384,75 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
 	    }
     	}
     }
+
+    //TODO: duplicate code! move it into a function so that it can be reused!
+    if(no_of_ptx==0) {
+        //For CDP, we dump everything. So no_of_ptx will be 0.
+	snprintf(fname,1024,"_ptx_XXXXXX");
+	int fd=mkstemp(fname);
+	close(fd);
+
+	printf("GPGPU-Sim PTX: extracting embedded .ptx to temporary file \"%s\"\n", fname);
+	FILE *ptxfile = fopen(fname,"w");
+	fprintf(ptxfile,"%s", p_for_info);
+	fclose(ptxfile);
+
+	snprintf(fname2,1024,"_ptx2_XXXXXX");
+	fd=mkstemp(fname2);
+	close(fd);
+	char commandline2[4096];
+	snprintf(commandline2,4096,"cat %s | sed 's/.version 1.5/.version 1.4/' | sed 's/, texmode_independent//' | sed 's/\\(\\.extern \\.const\\[1\\] .b8 \\w\\+\\)\\[\\]/\\1\\[1\\]/' | sed 's/const\\[.\\]/const\\[0\\]/g' > %s", fname, fname2);
+	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");
+		exit(1);
+	}
+	char tempfile_ptxinfo[1024];
+	snprintf(tempfile_ptxinfo,1024,"%sinfo",fname);
+	char extra_flags[1024];
+	extra_flags[0]=0;
+#if CUDART_VERSION >= 3000
+	snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version);
+#endif
+
+	snprintf(commandline,1024,"$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);
+	if( result != 0 ) {
+		printf("GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result);
+		printf("               Ensure ptxas is in your path.\n");
+		exit(1);
+	}
+    }
+
     //Now that we got resource usage per kernel in a ptx file, we dump all into one file and pass it to rest of the code as usual.
-    char commandline3[4096];
-    char final_tempfile_ptxinfo[1024];
-    snprintf(final_tempfile_ptxinfo,1024,"f_tempfile_ptx");
-    snprintf(commandline3,4096, "cat *info > %s", final_tempfile_ptxinfo);
-    if (system(commandline3)!=0) {
-	printf("ERROR: Either we dont have info files or cat is not working \n");
-        printf("ERROR: %s command failed\n",commandline3);
-	exit(1);
-    }    
+    if(no_of_ptx>0){
+        char commandline3[4096];
+        snprintf(final_tempfile_ptxinfo,1024,"f_tempfile_ptx");
+        snprintf(commandline3,4096, "cat *info > %s", final_tempfile_ptxinfo);
+        if (system(commandline3)!=0) {
+	    printf("ERROR: Either we dont have info files or cat is not working \n");
+            printf("ERROR: %s command failed\n",commandline3);
+	    exit(1);
+        }
+    }	
 
     ptxinfo_in = fopen(final_tempfile_ptxinfo,"r");
-    g_ptxinfo_filename = final_tempfile_ptxinfo;
+    if(no_of_ptx>0)
+        g_ptxinfo_filename = final_tempfile_ptxinfo;
+    else
+	g_ptxinfo_filename = tempfile_ptxinfo;
     ptxinfo_parse();
 
     if( ! g_save_embedded_ptx ) {
-        snprintf(commandline,1024,"rm -f %s %s %s *info", fname, fname2, final_tempfile_ptxinfo);
+	if(no_of_ptx>0)
+            snprintf(commandline,1024,"rm -f %s %s %s *info", fname, fname2, final_tempfile_ptxinfo);
+	else
+            snprintf(commandline,1024,"rm -f %s %s %s *info", fname, fname2, tempfile_ptxinfo);
         printf("GPGPU-Sim PTX: removing ptxinfo using \"%s\"\n", commandline);
         if( system(commandline) != 0 ) {
     	    printf("GPGPU-Sim PTX: ERROR ** while removing temporary files\n");
-- 
cgit v1.3


From 6d36a6a20f6e0c7082243c0dabe1a6b734c5002b Mon Sep 17 00:00:00 2001
From: Tor Aamodt <aamodt@ece.ubc.ca>
Date: Sat, 31 Mar 2018 08:48:47 -0700
Subject: masked declaration leading to using unitialized buffer as file name,
 leading to crash

this still fails regressions
---
 src/cuda-sim/ptx_loader.cc | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

(limited to 'src/cuda-sim/ptx_loader.cc')

diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 9ff0859..34870c4 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -410,7 +410,7 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
 		printf("               and have \'cat\' and \'sed\' in your path.\n");
 		exit(1);
 	}
-	char tempfile_ptxinfo[1024];
+	//char tempfile_ptxinfo[1024];
 	snprintf(tempfile_ptxinfo,1024,"%sinfo",fname);
 	char extra_flags[1024];
 	extra_flags[0]=0;
-- 
cgit v1.3


From 31ad7674de6dc4b25ba862bcd00b660fdb1a5cff Mon Sep 17 00:00:00 2001
From: Amruth <amruth@ece.ubc.ca>
Date: Wed, 4 Apr 2018 13:44:10 -0700
Subject: adding missing ptxas flags for cdp support

---
 README                      | 66 +++++++++++++++++++++++++++++++++++++++++++++
 libcuda/cuda_runtime_api.cc |  4 +--
 src/cuda-sim/ptx_loader.cc  | 27 ++++++++++++-------
 3 files changed, 85 insertions(+), 12 deletions(-)

(limited to 'src/cuda-sim/ptx_loader.cc')

diff --git a/README b/README
index 6e2d734..4426cdd 100644
--- a/README
+++ b/README
@@ -346,3 +346,69 @@ the applications you care about (implying these applications worked for you
 before you did the merge). You want to do this before making further changes to
 identify any compile time or runtime errors that occur due to the code merging
 process. 
+
+
+** Debugging failing GPGPU-Sim Regressions **
+
+To debug failing GPGPU-Sim regression tests you need to run them locally.  The fastest way to do this, assuming you are working with GPGPU-Sim versions more recent than the GPGPU-Sim dev branch circa March 28, 2018 (commit hash 2221d208a745a098a60b0d24c05007e92aaba092), is to install Docker.  The instructions below were tested with Docker CE version 18.03 on Ubuntu and Mac OS.  Docker will enable you to run the same set of regressions used by GPGPU-Sim when submitting a pull request to https://github.com/gpgpu-sim/gpgpu-sim_distribution and also allow you to log in and launch GPGPU-Sim in gdb so you can inspect failures.  
+
+1. Install Docker.  On Ubuntu 14.04 and 16.04 the following instructions work:  https://docs.docker.com/install/linux/docker-ce/ubuntu/#uninstall-old-versions 
+
+2. Clone GPGPU-Sim from your fork of GPGPU-Sim. For example:
+
+	git clone https://github.com/<YOUR GITHUB USERNAME>/gpgpu-sim_distribution.git
+
+
+3. Run the following command (this is all one line) to run the regressions in docker:
+
+	docker run --privileged -v `pwd`:/home/runner/gpgpu-sim_distribution:rw aamodt/gpgpu-sim_regress:latest /bin/bash -c "./start_torque.sh; chown -R runner /home/runner/gpgpu-sim_distribution; su - runner -c 'source /home/runner/gpgpu-sim_distribution/setup_environment && make -j -C /home/runner/gpgpu-sim_distribution && cd /home/runner/gpgpu-sim_simulations/ && git pull && /home/runner/gpgpu-sim_simulations/util/job_launching/run_simulations.py -c /home/runner/gpgpu-sim_simulations/util/job_launching/regression_recipies/rodinia_2.0-ft/configs.gtx1080ti.yml -N regress && /home/runner/gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -v -N regress’; tail -f /dev/null"
+
+Explanation: The last part of this command, "tail -f /dev/null" will keep the docker container running after the regressions finish.  This enables you to log into the container to run the same tests inside gdb so you can debug.   The "--privileged" part enables you to use breakpoints inside gdb in a container.  The "-v" part maps the current directory (with the GPGPU-Sim source code you want to test) into the container. The string "aamodt/gpgpu-sim_regress:latest" is a tag for a container setup to run regressions which will be downloaded from docker hub.  The portion starting with /bin/bash is a set of commands run inside a bash shell inside the container.  E.g., the command start_torque.sh starts up a queue manager inside the container.  
+
+If the above command stops with the message "fatal: unable to access 'https://github.com/tgrogers/gpgpu-sim_simulations.git/': Could not resolve host: github.com" this likely means your computer sits behind a firewall which is blocking access to Google's name servers (e.g., 8.8.8.8).  To get around this you will need to modify th above command to point to your local DNS server.  Lookup your DNS server IP address which we will call <DNS_IP_ADDRESS> below.  On Ubuntu run "ifconfig" to lookup the network interface connecting your computer to the network.  Then run "nmcli device show <interface name>" to find the IP address of your DNS server.  Modify the above command to include "--dns <DNS_IP_ADDRESS>" after "run", E.g.,
+
+	docker run --dns <DNS_IP_ADDRESS> --privileged -v `pwd`:/home/runner/gpgpu-sim_distribution:rw aamodt/gpgpu-sim_regress:latest /bin/bash -c "./start_torque.sh; chown -R runner /home/runner/gpgpu-sim_distribution; su - runner -c 'source /home/runner/gpgpu-sim_distribution/setup_environment && make -j -C /home/runner/gpgpu-sim_distribution && cd /home/runner/gpgpu-sim_simulations/ && git pull && /home/runner/gpgpu-sim_simulations/util/job_launching/run_simulations.py -c /home/runner/gpgpu-sim_simulations/util/job_launching/regression_recipies/rodinia_2.0-ft/configs.gtx1080ti.yml -N regress && /home/runner/gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -v -N regress’; tail -f /dev/null"
+
+4. Find the CONTAINER ID associated with your docker container by running "docker ps". 
+
+5. Log into the container by running the command:
+
+	docker exec -it <CONTAINER_ID> /bin/bash -c "su -l runner"
+
+The container is running Ubuntu 16.04 and has screen, cscope and vim installed (if you find a favorite Linux tool missing, it is fairly easy to create derived containers that have additional tools).
+
+6. Lookup the directory of the regression test you want to debug by going to the regression log file directory:
+
+	cd /home/runner/gpgpu-sim_simulations/util/job_launching/logfiles
+
+7.  The file "failed_job_log_sim_log.regress.<DATE>.txt" includes information about the failed test including its simulation directory.  For the following example, I'll assume the first failing test was "hotspot-rodinia-2.0-ft-30_6_40___data_result_30_6_40_txt--GTX1080Ti" for which the simulation directory is /home/runner/gpgpu-sim_simulations/util/job_launching/../../sim_run_4.2/hotspot-rodinia-2.0-ft/30_6_40___data_result_30_6_40_txt/GTX1080Ti/
+
+8.  Change to the simulation directory using:
+
+	cd <simulation_directory>
+
+E.g., "cd /home/runner/gpgpu-sim_simulations/util/job_launching/../../sim_run_4.2/hotspot-rodinia-2.0-ft/30_6_40___data_result_30_6_40_txt/GTX1080Ti/"
+
+This directory should contain a file called "torque.sim" that contains commands used to launch the simulation during regression tests.  We will modify this file to enable us to re-run the regression test in gdb.   This directory should also contain a file containing the standard output during the regression test.  This file will end in .o<number> where <number> is the torque queue manager job number.  For the running example for me this file is called "hotspot-rodinia-2.0-ft-30_6_40___data_result_30_6_40_txt.o2".  Open this file to determine the LD_LIBRARY_PATH settings used when launching the simulation.  Look for a line that starts "doing: export LD_LIBRARY_PATH" and copy the entire line starting with "export LD_LIBRARY_PATH ..."
+
+9. Paste the "export LD_LIBRARY_PATH ..." line into the bash shell to set LD_LIBRARY_PATH.  E.g.,
+
+	export LD_LIBRARY_PATH=/home/runner/gpgpu-sim_simulations/util/job_launching/../../sim_run_4.2/gpgpu-sim-builds/libcudart_gpgpu-sim_git-commit-177d02254ae38b6331b17dd6cd139b570a03c589_modified_0.so:/gpgpu-sim/usr/local/gcc-4.5.4/lib64:/gpgpu-sim/usr/local/gcc-4.5.4/lib:/gpgpu-sim/usr/local/gcc-4.5.4/lib/gcc/x86_64-unknown-linux-gnu/lib64/:/gpgpu-sim/usr/local/gcc-4.5.4/lib/gcc/x86_64-unknown-linux-gnu/4.5.4/:/usr/lib/x86_64-linux-gnu:/home/runner/gpgpu-sim_distribution/lib/gcc-4.5.4/cuda-4020/release:/gpgpu-sim/usr/local/gcc-4.5.4/lib64:/gpgpu-sim/usr/local/gcc-4.5.4/lib:/gpgpu-sim/usr/local/gcc-4.5.4/lib/gcc/x86_64-unknown-linux-gnu/lib64/:/gpgpu-sim/usr/local/gcc-4.5.4/lib/gcc/x86_64-unknown-linux-gnu/4.5.4/:/usr/lib/x86_64-linux-gnu:
+
+10. In the same shell, build the debug version of GPGPU-Sim then return to the directory above:
+
+	pushd ~/gpgpu-sim_distribution/
+	source setup_environment debug
+	make
+	popd
+
+11. Open and edit torque.sim and preface the very last line with "gdb --args ".  After editing the last line in torque.sim should look something like:
+
+	gdb --args /home/runner/gpgpu-sim_simulations/util/job_launching/../../benchmarks/bin/4.2/release/hotspot-rodinia-2.0-ft 30 6 40 ./data/result_30_6_40.txt
+
+12. Re-run the regression test in gdb by sourcing the torque.sim file:
+
+	. torque.sim
+
+This will put you in at the (gdb) prompt.  Setup any breakpoints needed and run.  
+
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index ded1aee..c103244 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -1085,8 +1085,8 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetStreamPriorityRange(int* leastPriori
        	return cudaSuccess;	
 }
 
-__host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *pStream, unsigned int flags) {
-	return cudaStreamCreate(pStream);
+__host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *stream, unsigned int flags) {
+	return cudaStreamCreate(stream);
 }
 
 __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 34870c4..7863be4 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -353,11 +353,7 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
 
         #if CUDART_VERSION >= 3000
     	if (sm_version == 0) sm_version = 20;
-    	extern bool g_cdp_enabled;
-    	if(!g_cdp_enabled)
-        	snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version);
-    	else
-        	snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",sm_version);
+        snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version);
 	#endif
 
     	snprintf(commandline,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file  /dev/null 2> %s",
@@ -414,9 +410,15 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
 	snprintf(tempfile_ptxinfo,1024,"%sinfo",fname);
 	char extra_flags[1024];
 	extra_flags[0]=0;
-#if CUDART_VERSION >= 3000
-	snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version);
-#endif
+
+	#if CUDART_VERSION >= 3000
+	if (sm_version == 0) sm_version = 20;
+    	extern bool g_cdp_enabled;
+	if(!g_cdp_enabled)
+	       	snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version);
+	else
+	       	snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",sm_version);
+	#endif
 
 	snprintf(commandline,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file  /dev/null 2> %s",
 			            extra_flags, fname2, tempfile_ptxinfo);
@@ -448,11 +450,16 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
 	g_ptxinfo_filename = tempfile_ptxinfo;
     ptxinfo_parse();
 
+    snprintf(commandline,1024,"rm -f *info");
+    if( system(commandline) != 0 ) {
+	    printf("GPGPU-Sim PTX: ERROR ** while removing temporary info files\n");
+	    exit(1);
+    }
     if( ! g_save_embedded_ptx ) {
 	if(no_of_ptx>0)
-            snprintf(commandline,1024,"rm -f %s %s %s *info", fname, fname2, final_tempfile_ptxinfo);
+            snprintf(commandline,1024,"rm -f %s %s %s", fname, fname2, final_tempfile_ptxinfo);
 	else
-            snprintf(commandline,1024,"rm -f %s %s %s *info", fname, fname2, tempfile_ptxinfo);
+            snprintf(commandline,1024,"rm -f %s %s %s", fname, fname2, tempfile_ptxinfo);
         printf("GPGPU-Sim PTX: removing ptxinfo using \"%s\"\n", commandline);
         if( system(commandline) != 0 ) {
     	    printf("GPGPU-Sim PTX: ERROR ** while removing temporary files\n");
-- 
cgit v1.3


From 335bb4b38974b1843e1786c5d86b0fcf2a7943c8 Mon Sep 17 00:00:00 2001
From: Amruth <amruth@ece.ubc.ca>
Date: Wed, 4 Apr 2018 16:01:32 -0700
Subject: Prog gets stuck because it doesnt recieve EOF

---
 src/cuda-sim/ptx_loader.cc | 3 +++
 1 file changed, 3 insertions(+)

(limited to 'src/cuda-sim/ptx_loader.cc')

diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 7863be4..03ea31a 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -448,6 +448,9 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
         g_ptxinfo_filename = final_tempfile_ptxinfo;
     else
 	g_ptxinfo_filename = tempfile_ptxinfo;
+
+    //The program might get stuck because the parser didnt receive a EOF.
+    printf("NOTE: If the program is stuck, please press ctrl+d for ubuntu and ctrl+z for windows users \n");
     ptxinfo_parse();
 
     snprintf(commandline,1024,"rm -f *info");
-- 
cgit v1.3


From 69c57b077d6799f46ad43b99cc0e0049b7c43775 Mon Sep 17 00:00:00 2001
From: Amruth <amruth@ece.ubc.ca>
Date: Wed, 4 Apr 2018 16:06:42 -0700
Subject: updating EOF condition for Mac users

---
 src/cuda-sim/ptx_loader.cc | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

(limited to 'src/cuda-sim/ptx_loader.cc')

diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 03ea31a..4ddc6bf 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -450,7 +450,7 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
 	g_ptxinfo_filename = tempfile_ptxinfo;
 
     //The program might get stuck because the parser didnt receive a EOF.
-    printf("NOTE: If the program is stuck, please press ctrl+d for ubuntu and ctrl+z for windows users \n");
+    printf("NOTE: If the program is stuck, please press ctrl+d for Ubuntu/Mac and ctrl+z for Windows users \n");
     ptxinfo_parse();
 
     snprintf(commandline,1024,"rm -f *info");
-- 
cgit v1.3


From e1dc8113aa2a51885541f96943bd8d90eaccd968 Mon Sep 17 00:00:00 2001
From: Amruth <amruth@ece.ubc.ca>
Date: Thu, 5 Apr 2018 16:31:44 -0700
Subject: fixing file pointer and attributes issues

---
 README                      |  1 +
 libcuda/cuda_runtime_api.cc | 34 ++++++++++++++++++++++++++++++++--
 src/cuda-sim/cuda-sim.cc    |  1 +
 src/cuda-sim/ptx_loader.cc  |  5 ++---
 4 files changed, 36 insertions(+), 5 deletions(-)

(limited to 'src/cuda-sim/ptx_loader.cc')

diff --git a/README b/README
index 4426cdd..543177c 100644
--- a/README
+++ b/README
@@ -349,6 +349,7 @@ process.
 
 
 ** Debugging failing GPGPU-Sim Regressions **
+Credits: Tor M Aamodt
 
 To debug failing GPGPU-Sim regression tests you need to run them locally.  The fastest way to do this, assuming you are working with GPGPU-Sim versions more recent than the GPGPU-Sim dev branch circa March 28, 2018 (commit hash 2221d208a745a098a60b0d24c05007e92aaba092), is to install Docker.  The instructions below were tested with Docker CE version 18.03 on Ubuntu and Mac OS.  Docker will enable you to run the same set of regressions used by GPGPU-Sim when submitting a pull request to https://github.com/gpgpu-sim/gpgpu-sim_distribution and also allow you to log in and launch GPGPU-Sim in gdb so you can inspect failures.  
 
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index c103244..3fd88dc 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -349,7 +349,15 @@ class _cuda_device_id *GPGPUSim_Init()
 		prop->maxGridSize[2] = 0x40000000;
 		prop->totalConstMem = 0x40000000;
 		prop->textureAlignment = 0;
-		prop->sharedMemPerBlock = the_gpu->shared_mem_size();
+		/*
+		 * TODO: Update the .config and xml files of all GPU config files with new value of sharedMemPerBlock.
+		 * Previously, this was thought as sharedMemPerMultiprocessor and is being used in many places.
+		 * Check whether all the instances of shared_mem_size(), gpgpu_shmem_size or sharedMemPerBlock are meant to use sharedMemPerBlock or sharedMemPerMultiprocessor.
+		*/
+		prop->sharedMemPerBlock = 49152;
+#if (CUDART_VERSION > 5000)
+		prop->sharedMemPerMultiprocessor = the_gpu->shared_mem_size();
+#endif
 		prop->regsPerBlock = the_gpu->num_registers_per_core();
 		prop->warpSize = the_gpu->wrp_size();
 		prop->clockRate = the_gpu->shader_clock();
@@ -840,6 +848,15 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic
         if (device <= dev->num_devices() )  {
                 prop = dev->get_prop();
                 switch (attr) {
+                case 2:
+                        *value= prop->maxThreadsDim[0];
+                        break;
+                case 3:
+                        *value= prop->maxThreadsDim[1];
+                        break;
+                case 4:
+                        *value= prop->maxThreadsDim[2];
+                        break;
                 case 5:
                         *value= prop->maxGridSize[0];
                         break;
@@ -849,6 +866,12 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic
                 case 7:
                         *value= prop->maxGridSize[2];
                         break;
+                case 8:
+                        *value= prop->sharedMemPerBlock;
+                        break;
+                case 9:
+                        *value= prop->totalConstMem;
+                        break;
                 case 10:
                         *value= prop->warpSize;
                         break;
@@ -861,11 +884,14 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic
                 case 16:
                         *value= prop->multiProcessorCount ;
                         break;
+                case 34:
+                        *value= 0;
+                        break;
                 case 39:
                         *value= dev->get_gpgpu()->threads_per_core();
                         break;
                 case 75:
-                        *value= 8 ;
+                        *value= 9 ;
                         break;
                 case 76:
                         *value= 3 ;
@@ -873,6 +899,9 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic
                 case 78:
                         *value= 0 ; //TODO: as of now, we dont support stream priorities.
                         break;
+                case 81:
+                        *value= prop->sharedMemPerMultiprocessor;
+                        break;
 		default:
 			printf("ERROR: implement the attribute numbered %d \n",attr);
 			abort();
@@ -1882,6 +1911,7 @@ void cuobjdumpParseBinary(unsigned int handle){
 		if (capability > max_capability) max_capability = capability;
 	}
 	if (max_capability > 20) printf("WARNING: No guarantee that PTX will be parsed for SM version %u\n", max_capability);
+	max_capability=context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
 
 	cuobjdumpPTXSection* ptx = NULL;
 	const char* pre_load = getenv("CUOBJDUMP_SIM_FILE");
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index dce35ca..b1eaf01 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -2155,3 +2155,4 @@ void functionalCoreSim::warp_exit( unsigned warp_id )
         }
     }
 }
+
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 4ddc6bf..33a4260 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -423,6 +423,7 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
 	snprintf(commandline,1024,"$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);
+	fflush(stdout);
 	result = system(commandline);
 	if( result != 0 ) {
 		printf("GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result);
@@ -443,14 +444,12 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
         }
     }	
 
-    ptxinfo_in = fopen(final_tempfile_ptxinfo,"r");
     if(no_of_ptx>0)
         g_ptxinfo_filename = final_tempfile_ptxinfo;
     else
 	g_ptxinfo_filename = tempfile_ptxinfo;
+    ptxinfo_in = fopen(g_ptxinfo_filename,"r");
 
-    //The program might get stuck because the parser didnt receive a EOF.
-    printf("NOTE: If the program is stuck, please press ctrl+d for Ubuntu/Mac and ctrl+z for Windows users \n");
     ptxinfo_parse();
 
     snprintf(commandline,1024,"rm -f *info");
-- 
cgit v1.3


From 2278609c8e279e7dfba49211256e818e3152318b Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Fri, 1 Jun 2018 13:04:30 -0700
Subject: WIP attempting to parse ptx files in order

---
 libcuda/cuda_runtime_api.cc | 257 +++++++++++++++++++++++++-------------------
 src/cuda-sim/ptx_loader.cc  |  20 ++++
 src/cuda-sim/ptx_loader.h   |   1 +
 3 files changed, 168 insertions(+), 110 deletions(-)

(limited to 'src/cuda-sim/ptx_loader.cc')

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index dc92522..ece958e 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -112,6 +112,7 @@
 #include <regex>
 #include <sstream>
 #include <fstream>
+#include <queue>
 #ifdef OPENGL_SUPPORT
 #define GL_GLEXT_PROTOTYPES
 #ifdef __APPLE__
@@ -148,6 +149,10 @@
 std::map<void *,void **> pinned_memory; //support for pinned memories added
 std::map<void *, size_t> pinned_memory_size;
 int no_of_ptx=0;
+char ptx_list_file_name[1024];
+std::map<int, std::string> fatbinmap;
+std::map<int, bool>fatbin_registered;
+std::map<std::string, symbol_table*> name_symtab;
 
 extern void synchronize();
 extern void exit_simulation();
@@ -1500,6 +1505,145 @@ char* get_app_binary_name(std::string abs_path){
    return self_exe_path;
 }
 
+//extracts all ptx files from binary and dumps into prog_name.unique_no.sm_<>.ptx files
+void extract_ptx_files_using_cuobjdump(bool g_cdp_enabled){
+    char command[1000];
+    std::string app_binary = get_app_binary(); 
+
+    snprintf(ptx_list_file_name,1024,"_cuobjdump_list_ptx_XXXXXX");
+    int fd2=mkstemp(ptx_list_file_name);
+    close(fd2);
+    //only want file names
+    snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -lptx %s  | cut -d \":\" -f 2 | awk '{$1=$1}1' > %s", app_binary.c_str(), ptx_list_file_name);
+    if( system(command) != 0 ) {
+        printf("WARNING: Failed to execute cuobjdump to get list of ptx files \n");
+        exit(0);
+    }   
+    if(!g_cdp_enabled) {
+        //based on the list above, dump ptx files individually. Format of dumped ptx file is prog_name.unique_no.sm_<>.ptx
+
+       std::ifstream infile(ptx_list_file_name);
+       std::string line;
+       while (std::getline(infile, line))
+       {
+            //int pos = line.find(std::string(get_app_binary_name(app_binary)));
+            const char *ptx_file = line.c_str();
+            printf("Extracting specific PTX file named %s \n",ptx_file);
+            snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -xptx %s %s", ptx_file, app_binary.c_str());
+            if (system(command)!=0) {
+                printf("ERROR: command: %s failed \n",command);
+                exit(0);
+            }
+            no_of_ptx++;
+       }
+    }
+
+	 if(!no_of_ptx){
+	 	printf("WARNING: Number of ptx in the executable file are 0. One of the reasons might be\n");
+	 	printf("\t1. CDP is enabled\n");
+	 }
+}
+
+void cuobjdumpParseBinary(unsigned int handle){
+
+
+	if(fatbin_registered[handle]) return;
+	fatbin_registered[handle] = true;
+	CUctx_st *context = GPGPUSim_Context();
+	std::string fname = fatbinmap[handle];
+
+	if (name_symtab.find(fname) != name_symtab.end()) {
+		symbol_table *symtab = name_symtab[fname];
+		context->add_binary(symtab, handle);
+		return;
+	}
+
+   std::map<int, std::set<std::string> > version_filename;
+
+   std::ifstream infile(ptx_list_file_name);
+   std::string line;
+   while (std::getline(infile, line))
+   {
+        //int pos = line.find(std::string(get_app_binary_name(app_binary)));
+        const char *ptx_file = line.c_str();
+        int pos1 = line.find("sm_");
+        int pos2 = line.find_last_of(".");
+        if (pos1==std::string::npos&&pos2==std::string::npos){
+            printf("ERROR: PTX list is not in correct format");
+            exit(0);
+        }
+        std::string vstr = line.substr(pos1+3,pos2-pos1-3);
+        int version = atoi(vstr.c_str());
+        if (version_filename.find(version)==version_filename.end()){
+           version_filename[version] = std::set<std::string>();
+        }
+        version_filename[version].insert(line);
+   }
+
+	symbol_table *symtab;
+   //loops through all ptx files from smallest sm version to largest
+   std::map<int,std::set<std::string> >::iterator itr_m;
+   for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){
+      std::set<std::string>::iterator itr_s;
+      for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){
+          std::string ptx_filename = *itr_s;
+          printf("GPGPU-Sim PTX: Parsing %s\n",ptx_filename.c_str());
+          symtab = gpgpu_ptx_sim_load_ptx_from_filename( ptx_filename.c_str() );
+      }
+   }
+	name_symtab[fname] = symtab;
+	context->add_binary(symtab, handle);
+
+
+//	unsigned max_capability = 0;
+//	for (	std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin();
+//			iter != cuobjdumpSectionList.end();
+//			iter++){
+//		unsigned capability = (*iter)->getArch();
+//		if (capability > max_capability) max_capability = capability;
+//	}
+//	if (max_capability > 20) printf("WARNING: No guarantee that PTX will be parsed for SM version %u\n", max_capability);
+//	if (max_capability == 0) max_capability=context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
+//
+//	cuobjdumpPTXSection* ptx = NULL;
+//	const char* pre_load = getenv("CUOBJDUMP_SIM_FILE");
+//	if(pre_load==NULL || strlen(pre_load)==0)
+//		ptx = findPTXSection(fname);
+//	symbol_table *symtab;
+//	char *ptxcode;
+//	const char *override_ptx_name = getenv("PTX_SIM_KERNELFILE"); 
+//	if (override_ptx_name == NULL or getenv("PTX_SIM_USE_PTX_FILE") == NULL or strlen(getenv("PTX_SIM_USE_PTX_FILE"))==0) {
+//		ptxcode = readfile(ptx->getPTXfilename());
+//	} else {
+//		printf("GPGPU-Sim PTX: overriding embedded ptx with '%s' (PTX_SIM_USE_PTX_FILE is set)\n", override_ptx_name);
+//		ptxcode = readfile(override_ptx_name);
+//	}
+//	if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) {
+//		cuobjdumpELFSection* elfsection = findELFSection(ptx->getIdentifier());
+//		assert (elfsection!= NULL);
+//		char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(
+//				ptx->getPTXfilename(),
+//				elfsection->getELFfilename(),
+//				elfsection->getSASSfilename());
+//		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 );
+//		delete[] ptxplus_str;
+//	} else {
+//		symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxcode, handle);
+//		//if CUOBJDUMP_SIM_FILE is not set, ptx is NULL. So comment below.
+//		//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 );
+//	}
+	load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
+	load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
+//	name_symtab[fname] = symtab;
+
+	//TODO: Remove temporarily files as per configurations
+}
+
 //! Call cuobjdump to extract everything (-elf -sass -ptx)
 /*!
  *	This Function extract the whole PTX (for all the files) using cuobjdump
@@ -1514,7 +1658,7 @@ void extract_code_using_cuobjdump(){
 
     //prevent the dumping by cuobjdump everytime we execute the code!
     const char *override_cuobjdump = getenv("CUOBJDUMP_SIM_FILE"); 
-    char command[1000], ptx_file[1000];
+    char command[1000];
     std::string app_binary = get_app_binary(); 
     //Running cuobjdump using dynamic link to current process
     snprintf(command,1000,"md5sum %s ", app_binary.c_str());
@@ -1527,50 +1671,8 @@ void extract_code_using_cuobjdump(){
     //dump ptx for all individial ptx files into sepearte files which is later used by ptxas.
     int result=0;
 #if (CUDART_VERSION >= 6000)
-    char fname2[1024];
-    snprintf(fname2,1024,"_cuobjdump_list_ptx_XXXXXX");
-    int fd2=mkstemp(fname2);
-    close(fd2);
-    snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -lptx -arch=sm_%u %s > %s", forced_max_capability, app_binary.c_str(), fname2);
-    result = system(command);
-    if( result != 0 ) {
-        printf("WARNING: Failed to execute cuobjdump to get list of ptx files \n");
-        exit(0);
-    } else {
-	/*
-          as we got list of ptx files, we need to extract one by one into seperate files so that ptxas can understand it. 
-          In this way, the duplicate definitions in a single embedded file can be prevented. 
-          No of lines in the file is equal to no of ptx fileis available.
-        */
-        FILE *fp = fopen(fname2,"r");
-        if (fp==NULL) {
-            printf("WARNING: cuobjdump file error! Could not open file %s \n", fname2);
-            exit(0);
-        } else {
-            for (char c = getc(fp); c != EOF; c = getc(fp))
-                if (c == '\n')
-                    no_of_ptx = no_of_ptx + 1;
-                    fclose(fp);
-        }
-	if(no_of_ptx==0){
-		printf("WARNING: Number of ptx in the executable file are 0. One of the reasons might be\n");
-		printf("\t1. CDP is enabled\n");
-		printf("\t2. cuobjdump -lptx doesnt recognize sm_%u\n",forced_max_capability);
-		printf("\t3. the application was not compiled with nvcc flag sm_%u\n",forced_max_capability);
-	}
-    }
-    if(!g_cdp_enabled) {
-        //based on the list above, dump ptx files individually. Format of dumped ptx file is prog_name.unique_no.sm_<>.ptx
-        for (int index=1; index<= no_of_ptx; index++){
-            snprintf(ptx_file, 1000, "%s.%d.sm_%u.ptx", get_app_binary_name(app_binary), index, forced_max_capability);
-            printf("Extracting specific PTX file named %s \n",ptx_file);
-            snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -arch=sm_%u -xptx %s %s", forced_max_capability, ptx_file, app_binary.c_str());
-            if (system(command)!=0) {
-                printf("ERROR: command: %s failed \n",command);
-                exit(0);
-            }
-        }
-    }
+    extract_ptx_files_using_cuobjdump(g_cdp_enabled);
+    cuobjdumpParseBinary(1);
 #endif
     //TODO: redundant to dump twice. how can it be prevented?
     //dump only for specific arch
@@ -1904,77 +2006,12 @@ void cuobjdumpInit(){
 	}
 }
 
-std::map<int, std::string> fatbinmap;
-std::map<int, bool>fatbin_registered;
-std::map<std::string, symbol_table*> name_symtab;
 
 //! Keep track of the association between filename and cubin handle
 void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename){
 	fatbinmap[handle] = filename;
 }
 
-//! Either submit PTX for simulation or convert SASS to PTXPlus and submit it
-void cuobjdumpParseBinary(unsigned int handle){
-
-	if(fatbin_registered[handle]) return;
-	fatbin_registered[handle] = true;
-	CUctx_st *context = GPGPUSim_Context();
-	std::string fname = fatbinmap[handle];
-
-	if (name_symtab.find(fname) != name_symtab.end()) {
-		symbol_table *symtab = name_symtab[fname];
-		context->add_binary(symtab, handle);
-		return;
-	}
-
-	unsigned max_capability = 0;
-	for (	std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin();
-			iter != cuobjdumpSectionList.end();
-			iter++){
-		unsigned capability = (*iter)->getArch();
-		if (capability > max_capability) max_capability = capability;
-	}
-	if (max_capability > 20) printf("WARNING: No guarantee that PTX will be parsed for SM version %u\n", max_capability);
-	if (max_capability == 0) max_capability=context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
-
-	cuobjdumpPTXSection* ptx = NULL;
-	const char* pre_load = getenv("CUOBJDUMP_SIM_FILE");
-	if(pre_load==NULL || strlen(pre_load)==0)
-		ptx = findPTXSection(fname);
-	symbol_table *symtab;
-	char *ptxcode;
-	const char *override_ptx_name = getenv("PTX_SIM_KERNELFILE"); 
-	if (override_ptx_name == NULL or getenv("PTX_SIM_USE_PTX_FILE") == NULL or strlen(getenv("PTX_SIM_USE_PTX_FILE"))==0) {
-		ptxcode = readfile(ptx->getPTXfilename());
-	} else {
-		printf("GPGPU-Sim PTX: overriding embedded ptx with '%s' (PTX_SIM_USE_PTX_FILE is set)\n", override_ptx_name);
-		ptxcode = readfile(override_ptx_name);
-	}
-	if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) {
-		cuobjdumpELFSection* elfsection = findELFSection(ptx->getIdentifier());
-		assert (elfsection!= NULL);
-		char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(
-				ptx->getPTXfilename(),
-				elfsection->getELFfilename(),
-				elfsection->getSASSfilename());
-		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 );
-		delete[] ptxplus_str;
-	} else {
-		symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxcode, handle);
-		//if CUOBJDUMP_SIM_FILE is not set, ptx is NULL. So comment below.
-		//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 );
-	}
-	load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
-	load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
-	name_symtab[fname] = symtab;
-
-	//TODO: Remove temporarily files as per configurations
-}
 
 void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
 {
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 33a4260..8deafc6 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -185,6 +185,26 @@ symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source
     return symtab;
 }
 
+symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename )
+{
+    symbol_table *symtab=init_parser(filename);
+    int errors = ptx_parse ();
+    if ( errors ) {
+        char fname[1024];
+        snprintf(fname,1024,"_ptx_errors_XXXXXX");
+        int fd=mkstemp(fname); 
+        close(fd);
+        printf("GPGPU-Sim PTX: parser error detected, exiting... but first extracting .ptx to \"%s\"\n", fname);
+        FILE *ptxfile = fopen(fname,"w");
+        fclose(ptxfile);
+        abort();
+        exit(40);
+    }
+
+    printf("GPGPU-Sim PTX: finished parsing EMBEDDED .ptx file %s\n",filename);
+    return symtab;
+}
+
 void fix_duplicate_errors(char fname2[1024]) {
 	char tempfile[1024] = "_temp_ptx";
 	char commandline[1024];
diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h
index a8ecda3..c5c9dd8 100644
--- a/src/cuda-sim/ptx_loader.h
+++ b/src/cuda-sim/ptx_loader.h
@@ -33,6 +33,7 @@ extern bool g_override_embedded_ptx;
 extern int no_of_ptx; //counter to track number of ptx files to be extracted in an application.
  
 class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num );
+class symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename );
 void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20 );
 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 73fea7152926dbc41cf008a4ed3402cc1feeefa7 Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Mon, 4 Jun 2018 16:53:23 -0700
Subject: parses through all ptx files, TODO: need to impl dp4a

---
 libcuda/cuda_runtime_api.cc  |  1 -
 src/cuda-sim/instructions.cc |  7 +++++++
 src/cuda-sim/opcodes.def     |  1 +
 src/cuda-sim/ptx.l           |  1 +
 src/cuda-sim/ptx_loader.cc   | 13 -------------
 src/cuda-sim/ptx_parser.cc   |  6 ++++++
 6 files changed, 15 insertions(+), 14 deletions(-)

(limited to 'src/cuda-sim/ptx_loader.cc')

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index ece958e..d03caf7 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -1546,7 +1546,6 @@ void extract_ptx_files_using_cuobjdump(bool g_cdp_enabled){
 
 void cuobjdumpParseBinary(unsigned int handle){
 
-
 	if(fatbin_registered[handle]) return;
 	fatbin_registered[handle] = true;
 	CUctx_st *context = GPGPUSim_Context();
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 35d1782..37438fa 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -2289,6 +2289,13 @@ void div_impl( const ptx_instruction *pI, ptx_thread_info *thread )
    thread->set_operand_value(dst,data, i_type, thread,pI);
 }
 
+void dp4a_impl( const ptx_instruction *pI, ptx_thread_info *thread ) 
+{ 
+   printf("instruction not implemented yet");
+   assert(0);
+
+}
+
 void ex2_impl( const ptx_instruction *pI, ptx_thread_info *thread ) 
 { 
    ptx_reg_t src1_data, src2_data, data;
diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def
index e1b1422..2dbfb78 100644
--- a/src/cuda-sim/opcodes.def
+++ b/src/cuda-sim/opcodes.def
@@ -60,6 +60,7 @@ OP_DEF(COS_OP,cos_impl,"cos",1,4)
 OP_DEF(CVT_OP,cvt_impl,"cvt",1,1)
 OP_DEF(CVTA_OP,cvta_impl,"cvta",1,1)
 OP_DEF(DIV_OP,div_impl,"div",1,1)
+OP_DEF(DP4A_OP,dp4a_impl,"dp4a",1,1)
 OP_DEF(EX2_OP,ex2_impl,"ex2",1,4)
 OP_DEF(EXIT_OP,exit_impl,"exit",1,3)
 OP_DEF(FMA_OP,fma_impl,"fma",1,2)
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index 908c5be..3bb0d17 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -77,6 +77,7 @@ cos	TC; ptx_lval.int_value = COS_OP; return OPCODE;
 cvt	TC; ptx_lval.int_value = CVT_OP; return OPCODE;
 cvta	TC; ptx_lval.int_value = CVTA_OP; return OPCODE;
 div	TC; ptx_lval.int_value = DIV_OP; return OPCODE;
+dp4a	TC; ptx_lval.int_value = DP4A_OP; return OPCODE;
 ex2	TC; ptx_lval.int_value = EX2_OP; return OPCODE;
 exit	TC; ptx_lval.int_value = EXIT_OP; return OPCODE;
 fma     TC; ptx_lval.int_value = FMA_OP; return OPCODE;
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 8deafc6..0348af0 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -188,19 +188,6 @@ symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source
 symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename )
 {
     symbol_table *symtab=init_parser(filename);
-    int errors = ptx_parse ();
-    if ( errors ) {
-        char fname[1024];
-        snprintf(fname,1024,"_ptx_errors_XXXXXX");
-        int fd=mkstemp(fname); 
-        close(fd);
-        printf("GPGPU-Sim PTX: parser error detected, exiting... but first extracting .ptx to \"%s\"\n", fname);
-        FILE *ptxfile = fopen(fname,"w");
-        fclose(ptxfile);
-        abort();
-        exit(40);
-    }
-
     printf("GPGPU-Sim PTX: finished parsing EMBEDDED .ptx file %s\n",filename);
     return symtab;
 }
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index a180da9..a51799a 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -32,6 +32,8 @@
 
 extern int ptx_error( const char *s );
 extern int ptx_lineno;
+extern int ptx_parse();
+extern FILE *ptx_in;
 
 static const struct core_config *g_shader_core_config;
 void set_ptx_warp_size(const struct core_config * warp_size)
@@ -135,6 +137,10 @@ symbol_table *init_parser( const char *ptx_filename )
    g_ptx_token_decode[generic_space] = "generic_space";
    g_ptx_token_decode[instruction_space] = "instruction_space";
 
+
+   ptx_in = fopen(ptx_filename, "r");
+   ptx_parse();
+   fclose(ptx_in);
    return g_global_symbol_table;
 }
 
-- 
cgit v1.3


From b79708980b80b6e50b1612f5436655c724377779 Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Tue, 5 Jun 2018 17:22:35 -0700
Subject: ptx parse only

---
 libcuda/cuda_runtime_api.cc | 48 ++++++++++++++++++++++-----------------------
 src/cuda-sim/ptx_loader.cc  | 37 ++++------------------------------
 src/cuda-sim/ptx_parser.cc  |  5 ++---
 3 files changed, 29 insertions(+), 61 deletions(-)

(limited to 'src/cuda-sim/ptx_loader.cc')

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index d03caf7..4efbd68 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -112,7 +112,6 @@
 #include <regex>
 #include <sstream>
 #include <fstream>
-#include <queue>
 #ifdef OPENGL_SUPPORT
 #define GL_GLEXT_PROTOTYPES
 #ifdef __APPLE__
@@ -149,7 +148,7 @@
 std::map<void *,void **> pinned_memory; //support for pinned memories added
 std::map<void *, size_t> pinned_memory_size;
 int no_of_ptx=0;
-char ptx_list_file_name[1024];
+std::map<int, std::set<std::string> > version_filename;
 std::map<int, std::string> fatbinmap;
 std::map<int, bool>fatbin_registered;
 std::map<std::string, symbol_table*> name_symtab;
@@ -1510,6 +1509,7 @@ void extract_ptx_files_using_cuobjdump(bool g_cdp_enabled){
     char command[1000];
     std::string app_binary = get_app_binary(); 
 
+    char ptx_list_file_name[1024];
     snprintf(ptx_list_file_name,1024,"_cuobjdump_list_ptx_XXXXXX");
     int fd2=mkstemp(ptx_list_file_name);
     close(fd2);
@@ -1542,6 +1542,27 @@ void extract_ptx_files_using_cuobjdump(bool g_cdp_enabled){
 	 	printf("WARNING: Number of ptx in the executable file are 0. One of the reasons might be\n");
 	 	printf("\t1. CDP is enabled\n");
 	 }
+
+    std::ifstream infile(ptx_list_file_name);
+    std::string line;
+    while (std::getline(infile, line))
+    {
+         //int pos = line.find(std::string(get_app_binary_name(app_binary)));
+         const char *ptx_file = line.c_str();
+         int pos1 = line.find("sm_");
+         int pos2 = line.find_last_of(".");
+         if (pos1==std::string::npos&&pos2==std::string::npos){
+             printf("ERROR: PTX list is not in correct format");
+             exit(0);
+         }
+         std::string vstr = line.substr(pos1+3,pos2-pos1-3);
+         int version = atoi(vstr.c_str());
+         if (version_filename.find(version)==version_filename.end()){
+            version_filename[version] = std::set<std::string>();
+         }
+         version_filename[version].insert(line);
+    }
+
 }
 
 void cuobjdumpParseBinary(unsigned int handle){
@@ -1557,28 +1578,6 @@ void cuobjdumpParseBinary(unsigned int handle){
 		return;
 	}
 
-   std::map<int, std::set<std::string> > version_filename;
-
-   std::ifstream infile(ptx_list_file_name);
-   std::string line;
-   while (std::getline(infile, line))
-   {
-        //int pos = line.find(std::string(get_app_binary_name(app_binary)));
-        const char *ptx_file = line.c_str();
-        int pos1 = line.find("sm_");
-        int pos2 = line.find_last_of(".");
-        if (pos1==std::string::npos&&pos2==std::string::npos){
-            printf("ERROR: PTX list is not in correct format");
-            exit(0);
-        }
-        std::string vstr = line.substr(pos1+3,pos2-pos1-3);
-        int version = atoi(vstr.c_str());
-        if (version_filename.find(version)==version_filename.end()){
-           version_filename[version] = std::set<std::string>();
-        }
-        version_filename[version].insert(line);
-   }
-
 	symbol_table *symtab;
    //loops through all ptx files from smallest sm version to largest
    std::map<int,std::set<std::string> >::iterator itr_m;
@@ -1671,7 +1670,6 @@ void extract_code_using_cuobjdump(){
     int result=0;
 #if (CUDART_VERSION >= 6000)
     extract_ptx_files_using_cuobjdump(g_cdp_enabled);
-    cuobjdumpParseBinary(1);
 #endif
     //TODO: redundant to dump twice. how can it be prevented?
     //dump only for specific arch
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 0348af0..97d5773 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -152,39 +152,6 @@ char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptxfilenam
 	return ptxplus_str;
 }
 
-
-symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num )
-{
-    char buf[1024];
-    snprintf(buf,1024,"_%u.ptx", source_num );
-    if( g_save_embedded_ptx ) {
-       FILE *fp = fopen(buf,"w");
-       fprintf(fp,"%s",p);
-       fclose(fp);
-    }
-    symbol_table *symtab=init_parser(buf);
-    ptx__scan_string(p);
-    int errors = ptx_parse ();
-    if ( errors ) {
-        char fname[1024];
-        snprintf(fname,1024,"_ptx_errors_XXXXXX");
-        int fd=mkstemp(fname); 
-        close(fd);
-        printf("GPGPU-Sim PTX: parser error detected, exiting... but first extracting .ptx to \"%s\"\n", fname);
-        FILE *ptxfile = fopen(fname,"w");
-        fprintf(ptxfile,"%s", p );
-        fclose(ptxfile);
-        abort();
-        exit(40);
-    }
-
-    if ( g_debug_execution >= 100 ) 
-       print_ptx_file(p,source_num,buf);
-
-    printf("GPGPU-Sim PTX: finished parsing EMBEDDED .ptx file %s\n",buf);
-    return symtab;
-}
-
 symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename )
 {
     symbol_table *symtab=init_parser(filename);
@@ -321,6 +288,10 @@ char* get_app_binary_name(){
    return self_exe_path;
 }
 
+void gpgpu_ptx_info_load_from_filename( const char *filename )
+{
+}
+
 void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version )
 {
     //do ptxas for individual files instead of one big embedded ptx. This prevents the duplicate defs and declarations.
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 06ca870..c418fac 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -142,9 +142,7 @@ symbol_table *init_parser( const char *ptx_filename )
        g_global_allfiles_symbol_table = new symbol_table("global_allfiles", 0, NULL);
        g_global_symbol_table = g_current_symbol_table = g_global_allfiles_symbol_table;
    }
-//   else {
-//       g_global_symbol_table = g_current_symbol_table = new symbol_table("global",0,g_global_allfiles_symbol_table);
-//   }
+
    ptx_lineno = 1;
 
 #define DEF(X,Y) g_ptx_token_decode[X] = Y;
@@ -164,6 +162,7 @@ symbol_table *init_parser( const char *ptx_filename )
    g_ptx_token_decode[global_space] = "global_space";
    g_ptx_token_decode[generic_space] = "generic_space";
    g_ptx_token_decode[instruction_space] = "instruction_space";
+
    init_directive_state();
    init_instruction_state();
 
-- 
cgit v1.3


From 40b275992791f34ecbb856e358af78cd3097da7e Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Wed, 6 Jun 2018 14:30:41 -0700
Subject: reverted some changed to pass regression

---
 libcuda/cuda_runtime_api.cc | 165 ++++++++++++++++++++++----------------------
 src/cuda-sim/ptx_loader.cc  |  32 +++++++++
 2 files changed, 116 insertions(+), 81 deletions(-)

(limited to 'src/cuda-sim/ptx_loader.cc')

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 4efbd68..08ee413 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -149,9 +149,6 @@ std::map<void *,void **> pinned_memory; //support for pinned memories added
 std::map<void *, size_t> pinned_memory_size;
 int no_of_ptx=0;
 std::map<int, std::set<std::string> > version_filename;
-std::map<int, std::string> fatbinmap;
-std::map<int, bool>fatbin_registered;
-std::map<std::string, symbol_table*> name_symtab;
 
 extern void synchronize();
 extern void exit_simulation();
@@ -1565,83 +1562,6 @@ void extract_ptx_files_using_cuobjdump(bool g_cdp_enabled){
 
 }
 
-void cuobjdumpParseBinary(unsigned int handle){
-
-	if(fatbin_registered[handle]) return;
-	fatbin_registered[handle] = true;
-	CUctx_st *context = GPGPUSim_Context();
-	std::string fname = fatbinmap[handle];
-
-	if (name_symtab.find(fname) != name_symtab.end()) {
-		symbol_table *symtab = name_symtab[fname];
-		context->add_binary(symtab, handle);
-		return;
-	}
-
-	symbol_table *symtab;
-   //loops through all ptx files from smallest sm version to largest
-   std::map<int,std::set<std::string> >::iterator itr_m;
-   for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){
-      std::set<std::string>::iterator itr_s;
-      for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){
-          std::string ptx_filename = *itr_s;
-          printf("GPGPU-Sim PTX: Parsing %s\n",ptx_filename.c_str());
-          symtab = gpgpu_ptx_sim_load_ptx_from_filename( ptx_filename.c_str() );
-      }
-   }
-	name_symtab[fname] = symtab;
-	context->add_binary(symtab, handle);
-
-
-//	unsigned max_capability = 0;
-//	for (	std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin();
-//			iter != cuobjdumpSectionList.end();
-//			iter++){
-//		unsigned capability = (*iter)->getArch();
-//		if (capability > max_capability) max_capability = capability;
-//	}
-//	if (max_capability > 20) printf("WARNING: No guarantee that PTX will be parsed for SM version %u\n", max_capability);
-//	if (max_capability == 0) max_capability=context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
-//
-//	cuobjdumpPTXSection* ptx = NULL;
-//	const char* pre_load = getenv("CUOBJDUMP_SIM_FILE");
-//	if(pre_load==NULL || strlen(pre_load)==0)
-//		ptx = findPTXSection(fname);
-//	symbol_table *symtab;
-//	char *ptxcode;
-//	const char *override_ptx_name = getenv("PTX_SIM_KERNELFILE"); 
-//	if (override_ptx_name == NULL or getenv("PTX_SIM_USE_PTX_FILE") == NULL or strlen(getenv("PTX_SIM_USE_PTX_FILE"))==0) {
-//		ptxcode = readfile(ptx->getPTXfilename());
-//	} else {
-//		printf("GPGPU-Sim PTX: overriding embedded ptx with '%s' (PTX_SIM_USE_PTX_FILE is set)\n", override_ptx_name);
-//		ptxcode = readfile(override_ptx_name);
-//	}
-//	if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) {
-//		cuobjdumpELFSection* elfsection = findELFSection(ptx->getIdentifier());
-//		assert (elfsection!= NULL);
-//		char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(
-//				ptx->getPTXfilename(),
-//				elfsection->getELFfilename(),
-//				elfsection->getSASSfilename());
-//		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 );
-//		delete[] ptxplus_str;
-//	} else {
-//		symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxcode, handle);
-//		//if CUOBJDUMP_SIM_FILE is not set, ptx is NULL. So comment below.
-//		//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 );
-//	}
-	load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
-	load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
-//	name_symtab[fname] = symtab;
-
-	//TODO: Remove temporarily files as per configurations
-}
-
 //! Call cuobjdump to extract everything (-elf -sass -ptx)
 /*!
  *	This Function extract the whole PTX (for all the files) using cuobjdump
@@ -1656,7 +1576,7 @@ void extract_code_using_cuobjdump(){
 
     //prevent the dumping by cuobjdump everytime we execute the code!
     const char *override_cuobjdump = getenv("CUOBJDUMP_SIM_FILE"); 
-    char command[1000];
+    char command[1000], ptx_file[1000];
     std::string app_binary = get_app_binary(); 
     //Running cuobjdump using dynamic link to current process
     snprintf(command,1000,"md5sum %s ", app_binary.c_str());
@@ -2003,12 +1923,95 @@ void cuobjdumpInit(){
 	}
 }
 
+std::map<int, std::string> fatbinmap;
+std::map<int, bool>fatbin_registered;
+std::map<std::string, symbol_table*> name_symtab;
 
 //! Keep track of the association between filename and cubin handle
 void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename){
 	fatbinmap[handle] = filename;
 }
 
+//! Either submit PTX for simulation or convert SASS to PTXPlus and submit it
+void cuobjdumpParseBinary(unsigned int handle){
+
+	if(fatbin_registered[handle]) return;
+	fatbin_registered[handle] = true;
+	CUctx_st *context = GPGPUSim_Context();
+	std::string fname = fatbinmap[handle];
+
+	if (name_symtab.find(fname) != name_symtab.end()) {
+		symbol_table *symtab = name_symtab[fname];
+		context->add_binary(symtab, handle);
+		return;
+	}
+	symbol_table *symtab;
+
+#if (CUDART_VERSION >= 6000)
+   //loops through all ptx files from smallest sm version to largest
+   std::map<int,std::set<std::string> >::iterator itr_m;
+   for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){
+      std::set<std::string>::iterator itr_s;
+      for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){
+          std::string ptx_filename = *itr_s;
+          printf("GPGPU-Sim PTX: Parsing %s\n",ptx_filename.c_str());
+          symtab = gpgpu_ptx_sim_load_ptx_from_filename( ptx_filename.c_str() );
+      }
+   }
+	name_symtab[fname] = symtab;
+	context->add_binary(symtab, 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());
+   return;
+#endif
+
+	unsigned max_capability = 0;
+	for (	std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin();
+			iter != cuobjdumpSectionList.end();
+			iter++){
+		unsigned capability = (*iter)->getArch();
+		if (capability > max_capability) max_capability = capability;
+	}
+	if (max_capability > 20) printf("WARNING: No guarantee that PTX will be parsed for SM version %u\n", max_capability);
+	if (max_capability == 0) max_capability=context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
+
+	cuobjdumpPTXSection* ptx = NULL;
+	const char* pre_load = getenv("CUOBJDUMP_SIM_FILE");
+	if(pre_load==NULL || strlen(pre_load)==0)
+		ptx = findPTXSection(fname);
+	char *ptxcode;
+	const char *override_ptx_name = getenv("PTX_SIM_KERNELFILE"); 
+	if (override_ptx_name == NULL or getenv("PTX_SIM_USE_PTX_FILE") == NULL or strlen(getenv("PTX_SIM_USE_PTX_FILE"))==0) {
+		ptxcode = readfile(ptx->getPTXfilename());
+	} else {
+		printf("GPGPU-Sim PTX: overriding embedded ptx with '%s' (PTX_SIM_USE_PTX_FILE is set)\n", override_ptx_name);
+		ptxcode = readfile(override_ptx_name);
+	}
+	if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) {
+		cuobjdumpELFSection* elfsection = findELFSection(ptx->getIdentifier());
+		assert (elfsection!= NULL);
+		char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(
+				ptx->getPTXfilename(),
+				elfsection->getELFfilename(),
+				elfsection->getSASSfilename());
+		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 );
+		delete[] ptxplus_str;
+	} else {
+		symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxcode, handle);
+		//if CUOBJDUMP_SIM_FILE is not set, ptx is NULL. So comment below.
+		//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 );
+	}
+	load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
+	load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
+	name_symtab[fname] = symtab;
+
+	//TODO: Remove temporarily files as per configurations
+}
 
 void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
 {
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 97d5773..2064b13 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -152,6 +152,38 @@ char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptxfilenam
 	return ptxplus_str;
 }
 
+symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num )
+{
+    char buf[1024];
+    snprintf(buf,1024,"_%u.ptx", source_num );
+    if( g_save_embedded_ptx ) {
+       FILE *fp = fopen(buf,"w");
+       fprintf(fp,"%s",p);
+       fclose(fp);
+    }
+    symbol_table *symtab=init_parser(buf);
+    ptx__scan_string(p);
+    int errors = ptx_parse ();
+    if ( errors ) {
+        char fname[1024];
+        snprintf(fname,1024,"_ptx_errors_XXXXXX");
+        int fd=mkstemp(fname); 
+        close(fd);
+        printf("GPGPU-Sim PTX: parser error detected, exiting... but first extracting .ptx to \"%s\"\n", fname);
+        FILE *ptxfile = fopen(fname,"w");
+        fprintf(ptxfile,"%s", p );
+        fclose(ptxfile);
+        abort();
+        exit(40);
+    }
+
+    if ( g_debug_execution >= 100 ) 
+       print_ptx_file(p,source_num,buf);
+
+    printf("GPGPU-Sim PTX: finished parsing EMBEDDED .ptx file %s\n",buf);
+    return symtab;
+}
+
 symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename )
 {
     symbol_table *symtab=init_parser(filename);
-- 
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 'src/cuda-sim/ptx_loader.cc')

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 'src/cuda-sim/ptx_loader.cc')

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 'src/cuda-sim/ptx_loader.cc')

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 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 'src/cuda-sim/ptx_loader.cc')

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 bd9a49ef576bfbbaf2be1d163eb99b7316f7bd20 Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Mon, 23 Jul 2018 14:20:14 -0700
Subject: brought back ptxinfo

---
 libcuda/cuda_runtime_api.cc | 19 +++++++++++++++----
 src/cuda-sim/ptx_loader.cc  | 23 ++++++++++++++++++++++-
 src/cuda-sim/ptx_loader.h   |  1 +
 3 files changed, 38 insertions(+), 5 deletions(-)

(limited to 'src/cuda-sim/ptx_loader.cc')

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index db4f58e..c87c6c3 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -152,7 +152,8 @@ std::map<void *,void **> pinned_memory; //support for pinned memories added
 std::map<void *, size_t> pinned_memory_size;
 std::map<unsigned long long, size_t> g_mallocPtr_Size;
 int no_of_ptx=0;
-std::map<int, std::set<std::string> > version_filename;
+//maps sm version number to set of filenames
+std::map<unsigned, std::set<std::string> > version_filename;
 
 extern void synchronize();
 extern void exit_simulation();
@@ -1718,7 +1719,8 @@ char* get_app_binary_name(std::string abs_path){
 }
 
 //extracts all ptx files from binary and dumps into prog_name.unique_no.sm_<>.ptx files
-void extract_ptx_files_using_cuobjdump(bool g_cdp_enabled){
+void extract_ptx_files_using_cuobjdump(){
+    extern bool g_cdp_enabled;
     char command[1000];
     std::string app_binary = get_app_binary(); 
 
@@ -1808,7 +1810,8 @@ void extract_code_using_cuobjdump(){
     //dump ptx for all individial ptx files into sepearte files which is later used by ptxas.
     int result=0;
 #if (CUDART_VERSION >= 6000)
-    extract_ptx_files_using_cuobjdump(g_cdp_enabled);
+    extract_ptx_files_using_cuobjdump();
+    return;
 #endif
     //TODO: redundant to dump twice. how can it be prevented?
     //dump only for specific arch
@@ -2168,7 +2171,7 @@ void cuobjdumpParseBinary(unsigned int handle){
 
 #if (CUDART_VERSION >= 6000)
    //loops through all ptx files from smallest sm version to largest
-   std::map<int,std::set<std::string> >::iterator itr_m;
+   std::map<unsigned,std::set<std::string> >::iterator itr_m;
    for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){
       std::set<std::string>::iterator itr_s;
       for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){
@@ -2181,6 +2184,14 @@ void cuobjdumpParseBinary(unsigned int handle){
    context->add_binary(symtab, 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());
+   for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){
+      std::set<std::string>::iterator itr_s;
+      for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){
+          std::string ptx_filename = *itr_s;
+          printf("GPGPU-Sim PTX: Loading PTXInfo from %s\n",ptx_filename.c_str());
+          gpgpu_ptx_info_load_from_filename( ptx_filename.c_str(), itr_m->first );
+      }
+   }
    return;
 #endif
 
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 2064b13..f5a1c77 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -320,8 +320,29 @@ char* get_app_binary_name(){
    return self_exe_path;
 }
 
-void gpgpu_ptx_info_load_from_filename( const char *filename )
+void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version)
 {
+    std::string ptxas_filename(std::string(filename) + "as");
+    char buff[1024], extra_flags[256];
+	extra_flags[0]=0;
+    extern bool g_cdp_enabled;
+	if(!g_cdp_enabled)
+	       	snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version);
+	else
+	       	snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",sm_version);
+   	snprintf(buff,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file  /dev/null 2> %s",
+         extra_flags, filename, ptxas_filename.c_str());
+	int result = system(buff);
+	if( result != 0 ) {
+		printf("GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result);
+		printf("               Ensure ptxas is in your path.\n");
+		exit(1);
+	}
+
+	g_ptxinfo_filename = strdup(ptxas_filename.c_str());
+    ptxinfo_in = fopen(g_ptxinfo_filename,"r");
+    ptxinfo_parse();
+    fclose(ptxinfo_in);
 }
 
 void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version )
diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h
index c5c9dd8..e5df6a9 100644
--- a/src/cuda-sim/ptx_loader.h
+++ b/src/cuda-sim/ptx_loader.h
@@ -35,6 +35,7 @@ extern int no_of_ptx; //counter to track number of ptx files to be extracted in
 class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num );
 class symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename );
 void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20 );
+void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version );
 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