From 18e0b0614611edcb19ee0a3b315c7b45e50b5595 Mon Sep 17 00:00:00 2001
From: Ahmed ElTantawy <ahmede@ece.ubc.ca>
Date: Wed, 4 Mar 2015 12:02:52 -0800
Subject: initial support for CUDA 5.0, 5.5, 6.0 to get template from SDK
 running

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

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index 0334168..080ca2d 100644
--- a/Makefile
+++ b/Makefile
@@ -152,6 +152,9 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.2 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.2; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.3 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.3; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.4 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.4; fi
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.5.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.5.0; fi
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.5.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.5.5; fi
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.6.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.6.0; fi
 
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
-- 
cgit v1.3


From c0ed7da45b3836cefc08c6072fd183d2b33bbdeb Mon Sep 17 00:00:00 2001
From: Ahmed ElTantawy <ahmede@ece.ubc.ca>
Date: Thu, 12 May 2016 14:25:14 -0700
Subject: updating Makefile and setup environment for 7.5 support

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

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index 080ca2d..9e2c576 100644
--- a/Makefile
+++ b/Makefile
@@ -155,7 +155,7 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.5.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.5.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.5.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.5.5; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.6.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.6.0; fi
-
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.7.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.7.5; fi
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
 			$(SIM_OBJ_FILES_DIR)/libcuda/*.o \
diff --git a/setup_environment b/setup_environment
index c55d71f..2684514 100644
--- a/setup_environment
+++ b/setup_environment
@@ -43,7 +43,7 @@ CC_VERSION=`gcc --version | head -1 | awk '{for(i=1;i<=NF;i++){ if(match($i,/^[0
 
 CUDA_VERSION_STRING=`$CUDA_INSTALL_PATH/bin/nvcc --version | awk '/release/ {print $5;}' | sed 's/,//'`;
 CUDA_VERSION_NUMBER=`echo $CUDA_VERSION_STRING | sed 's/\./ /' | awk '{printf("%02u%02u", 10*int($1), 10*$2);}'`
-if [ $CUDA_VERSION_NUMBER -gt 6000 -o $CUDA_VERSION_NUMBER -lt 2030  ]; then
+if [ $CUDA_VERSION_NUMBER -gt 7500 -o $CUDA_VERSION_NUMBER -lt 2030  ]; then
 	echo "ERROR ** GPGPU-Sim version $GPGPUSIM_VERSION_STRING not tested with CUDA version $CUDA_VERSION_STRING (please see README)";
 	return
 elif [ $CUDA_VERSION_NUMBER -gt 4020 ]; then
-- 
cgit v1.3


From fabf1747e3f798c76ddc9c5fae26a8d4d3cdedec Mon Sep 17 00:00:00 2001
From: Jin Wang <wangjin.jane@gmail.com>
Date: Thu, 25 Sep 2014 16:15:09 -0400
Subject: MOD: add support for cuda65 ptx header

---
 Makefile                    | 2 ++
 libcuda/cuda_runtime_api.cc | 4 ++--
 2 files changed, 4 insertions(+), 2 deletions(-)

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index 9e2c576..a044648 100644
--- a/Makefile
+++ b/Makefile
@@ -155,7 +155,9 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.5.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.5.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.5.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.5.5; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.6.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.6.0; fi
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.6.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.6.5; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.7.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.7.5; fi
+
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
 			$(SIM_OBJ_FILES_DIR)/libcuda/*.o \
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 910bebd..e7a534e 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -1519,7 +1519,7 @@ cuobjdumpELFSection* findELFSection(const std::string identifier){
 	if (sec!=NULL)return sec;
 	sec = findELFSectionInList(libSectionList, identifier);
 	if (sec!=NULL)return sec;
-	std::cout << "Cound not find " << identifier << std::endl;
+	std::cout << "Could not find " << identifier << std::endl;
 	assert(0 && "Could not find the required ELF section");
 	return NULL;
 }
@@ -1546,7 +1546,7 @@ cuobjdumpPTXSection* findPTXSection(const std::string identifier){
 	if (sec!=NULL)return sec;
 	sec = findPTXSectionInList(libSectionList, identifier);
 	if (sec!=NULL)return sec;
-	std::cout << "Cound not find " << identifier << std::endl;
+	std::cout << "Could not find " << identifier << std::endl;
 	assert(0 && "Could not find the required PTX section");
 	return NULL;
 }
-- 
cgit v1.3


From 7b7f9050424f21a6da17f6765add0c6cd244b2a3 Mon Sep 17 00:00:00 2001
From: Ahmed ElTantawy <ahmed.mohammed.eltantawy@gmail.com>
Date: Fri, 28 Oct 2016 10:49:19 -0700
Subject: support for CUDA 8.0 PTX

---
 Makefile                    | 2 ++
 libcuda/cuda_runtime_api.cc | 7 ++++++-
 setup_environment           | 2 +-
 3 files changed, 9 insertions(+), 2 deletions(-)

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index 9e2c576..3b6e40d 100644
--- a/Makefile
+++ b/Makefile
@@ -156,6 +156,8 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.5.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.5.5; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.6.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.6.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.7.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.7.5; fi
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.8.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.8.0; fi
+
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
 			$(SIM_OBJ_FILES_DIR)/libcuda/*.o \
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 910bebd..75a2f31 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -125,7 +125,9 @@
 #include "host_defines.h"
 #include "builtin_types.h"
 #include "driver_types.h"
+#if (CUDART_VERSION < 8000)
 #include "__cudaFatFormat.h"
+#endif
 #include "../src/gpgpu-sim/gpu-sim.h"
 #include "../src/cuda-sim/ptx_loader.h"
 #include "../src/cuda-sim/cuda-sim.h"
@@ -1667,7 +1669,9 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
 		cuobjdumpRegisterFatBinary(fat_cubin_handle, filename);
 
 		return (void**)fat_cubin_handle;
-	} else {
+	} 
+	#if (CUDART_VERSION < 8000)
+	else {
 		static unsigned source_num=1;
 		unsigned long long fat_cubin_handle = next_fat_bin_handle++;
 		__cudaFatCudaBinary *info =   (__cudaFatCudaBinary *)fatCubin;
@@ -1724,6 +1728,7 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
 		}
 		return (void**)fat_cubin_handle;
 	}
+	#endif
 }
 
 void __cudaUnregisterFatBinary(void **fatCubinHandle)
diff --git a/setup_environment b/setup_environment
index 2684514..854a335 100644
--- a/setup_environment
+++ b/setup_environment
@@ -43,7 +43,7 @@ CC_VERSION=`gcc --version | head -1 | awk '{for(i=1;i<=NF;i++){ if(match($i,/^[0
 
 CUDA_VERSION_STRING=`$CUDA_INSTALL_PATH/bin/nvcc --version | awk '/release/ {print $5;}' | sed 's/,//'`;
 CUDA_VERSION_NUMBER=`echo $CUDA_VERSION_STRING | sed 's/\./ /' | awk '{printf("%02u%02u", 10*int($1), 10*$2);}'`
-if [ $CUDA_VERSION_NUMBER -gt 7500 -o $CUDA_VERSION_NUMBER -lt 2030  ]; then
+if [ $CUDA_VERSION_NUMBER -gt 8000 -o $CUDA_VERSION_NUMBER -lt 2030  ]; then
 	echo "ERROR ** GPGPU-Sim version $GPGPUSIM_VERSION_STRING not tested with CUDA version $CUDA_VERSION_STRING (please see README)";
 	return
 elif [ $CUDA_VERSION_NUMBER -gt 4020 ]; then
-- 
cgit v1.3


From 8b44f564221b9a0503ca04ccdb463bd60ae8c33c Mon Sep 17 00:00:00 2001
From: tgrogers <timrogers@gmail.com>
Date: Tue, 16 May 2017 16:56:54 -0400
Subject: Fixing a bug with where the build version is made. Moveing it after
 the right directories ect are defined

---
 Makefile             | 3 +++
 version_detection.mk | 1 -
 2 files changed, 3 insertions(+), 1 deletion(-)

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index 78ef892..7a7b187 100644
--- a/Makefile
+++ b/Makefile
@@ -54,6 +54,9 @@ ifneq ($(shell which nvcc), "")
 	endif
 endif
 
+
+$(shell mkdir -p $(SIM_OBJ_FILES_DIR)/libcuda && echo "const char *g_gpgpusim_build_string=\"$(GPGPUSIM_BUILD)\";" > $(SIM_OBJ_FILES_DIR)/detailed_version)
+
 LIBS = cuda-sim gpgpu-sim_uarch $(INTERSIM) gpgpusimlib 
 
 
diff --git a/version_detection.mk b/version_detection.mk
index 00a8b1f..f9c333f 100644
--- a/version_detection.mk
+++ b/version_detection.mk
@@ -36,7 +36,6 @@ GIT_BRANCH := $(shell git branch | grep "\*" | sed -re 's/\*\s+(.*)/\1/')
 GIT_COMMIT := $(shell git log -n 1 | head -1 | sed -re 's/commit (.*)/\1/')
 GIT_FILES_CHANGED := $(shell git diff --numstat --cached && git diff --numstat | wc | sed -re 's/^\s+([0-9]+).*/\1/')
 GPGPUSIM_BUILD := "gpgpu-sim_$(GIT_BRANCH)_$(GIT_COMMIT)_modified_$(GIT_FILES_CHANGED)"
-$(shell mkdir -p $(SIM_OBJ_FILES_DIR)/libcuda && echo "const char *g_gpgpusim_build_string=\"$(GPGPUSIM_BUILD)\";" > $(SIM_OBJ_FILES_DIR)/detailed_version)
 endif
 
 # Detect CUDA Runtime Version 
-- 
cgit v1.3


From 959808f483a43dc2371f04e18c34fcebb0425b48 Mon Sep 17 00:00:00 2001
From: tgrogers <timrogers@purdue.edu>
Date: Wed, 17 May 2017 10:15:48 -0400
Subject: Adding in the embedded .so bit

---
 Makefile | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index 7a7b187..e83de9a 100644
--- a/Makefile
+++ b/Makefile
@@ -143,7 +143,7 @@ no_opencl_support:
 	@echo "Warning: gpgpu-sim is building without opencl support. Make sure NVOPENCL_LIBDIR and NVOPENCL_INCDIR are set"
 
 $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
-	g++ -shared -Wl,-soname,libcudart.so \
+	g++ -shared -Wl,-soname,libcudart_$(GPGPUSIM_BUILD).so \
 			$(SIM_OBJ_FILES_DIR)/libcuda/*.o \
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/*.o \
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/decuda_pred_table/*.o \
@@ -174,7 +174,7 @@ $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 			-o $(SIM_LIB_DIR)/libcudart.dylib
 
 $(SIM_LIB_DIR)/libOpenCL.so: makedirs $(LIBS) opencllib
-	g++ -shared -Wl,-soname,libOpenCL.so \
+	g++ -shared -Wl,-soname,libOpenCL_$(GPGPUSIM_BUILD).so \
 			$(SIM_OBJ_FILES_DIR)/libopencl/*.o \
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/*.o \
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/decuda_pred_table/*.o \
-- 
cgit v1.3


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

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

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index e83de9a..95e4432 100644
--- a/Makefile
+++ b/Makefile
@@ -161,6 +161,7 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.6.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.6.5; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.7.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.7.5; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.8.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.8.0; fi
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.1; fi
 
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
diff --git a/setup_environment b/setup_environment
index 0895d44..8a529c3 100644
--- a/setup_environment
+++ b/setup_environment
@@ -43,7 +43,7 @@ CC_VERSION=`gcc --version | head -1 | awk '{for(i=1;i<=NF;i++){ if(match($i,/^[0
 
 CUDA_VERSION_STRING=`$CUDA_INSTALL_PATH/bin/nvcc --version | awk '/release/ {print $5;}' | sed 's/,//'`;
 CUDA_VERSION_NUMBER=`echo $CUDA_VERSION_STRING | sed 's/\./ /' | awk '{printf("%02u%02u", 10*int($1), 10*$2);}'`
-if [ $CUDA_VERSION_NUMBER -gt 8000 -o $CUDA_VERSION_NUMBER -lt 2030  ]; then
+if [ $CUDA_VERSION_NUMBER -gt 9100 -o $CUDA_VERSION_NUMBER -lt 2030  ]; then
 	echo "ERROR ** GPGPU-Sim version $GPGPUSIM_VERSION_STRING not tested with CUDA version $CUDA_VERSION_STRING (please see README)";
 	return
 elif [ $CUDA_VERSION_NUMBER -gt 4020 ]; then
-- 
cgit v1.3


From 423dc13a99e5e23d48569a8d9fa844ff5de06ca0 Mon Sep 17 00:00:00 2001
From: Mahmoud <abdallm@purdue.edu>
Date: Fri, 23 Feb 2018 10:46:13 -0500
Subject: fixing the libcardt link issue

---
 Makefile | 8 ++++++--
 1 file changed, 6 insertions(+), 2 deletions(-)

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index e83de9a..1f043fa 100644
--- a/Makefile
+++ b/Makefile
@@ -62,7 +62,7 @@ LIBS = cuda-sim gpgpu-sim_uarch $(INTERSIM) gpgpusimlib
 
 TARGETS =
 ifeq ($(shell uname),Linux)
-	TARGETS += $(SIM_LIB_DIR)/libcudart.so
+	TARGETS += $(SIM_LIB_DIR)/libcudart.so 
 else # MAC
 	TARGETS += $(SIM_LIB_DIR)/libcudart.dylib
 endif
@@ -75,6 +75,7 @@ else
 	TARGETS += $(SIM_LIB_DIR)/libOpenCL.so
 endif
 	TARGETS += cuobjdump_to_ptxplus/cuobjdump_to_ptxplus
+	TARGETS += $(SIM_LIB_DIR)/gpgpusim.out
 
 MCPAT=
 MCPAT_OBJ_DIR=
@@ -143,7 +144,7 @@ no_opencl_support:
 	@echo "Warning: gpgpu-sim is building without opencl support. Make sure NVOPENCL_LIBDIR and NVOPENCL_INCDIR are set"
 
 $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
-	g++ -shared -Wl,-soname,libcudart_$(GPGPUSIM_BUILD).so \
+	g++ -shared -Wl \
 			$(SIM_OBJ_FILES_DIR)/libcuda/*.o \
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/*.o \
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/decuda_pred_table/*.o \
@@ -162,6 +163,9 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.7.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.7.5; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.8.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.8.0; fi
 
+$(SIM_LIB_DIR)/gpgpusim.out: makedirs $(LIBS) cudalib $(SIM_LIB_DIR)/libcudart.so
+	g++ -std=c++0x -L$(SIM_LIB_DIR) -lcudart -o $(SIM_LIB_DIR)/gpgpusim.out src/trace-driven/gpgpusim_trace_driven_main.cc 
+
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
 			$(SIM_OBJ_FILES_DIR)/libcuda/*.o \
-- 
cgit v1.3


From e97f396d5892d1c7549b18e95787fb635dff4851 Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Thu, 22 Mar 2018 16:47:39 -0400
Subject: Adding support for CUDA 9.1

---
 Makefile | 2 ++
 1 file changed, 2 insertions(+)

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index e83de9a..de6264c 100644
--- a/Makefile
+++ b/Makefile
@@ -161,6 +161,8 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.6.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.6.5; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.7.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.7.5; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.8.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.8.0; fi
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.0; fi
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.1; fi
 
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
-- 
cgit v1.3


From a368a6ef96620763848f38d924a7fe041bc5728b Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Sat, 24 Mar 2018 18:08:37 -0400
Subject: The new CUDAs complain about version info not being embedded into the
 so file. Having gpgpu-sim ebed some version info to stop the error complaints

---
 Makefile             | 2 +-
 linux-so-version.txt | 2 ++
 2 files changed, 3 insertions(+), 1 deletion(-)
 create mode 100644 linux-so-version.txt

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index 95e4432..1641833 100644
--- a/Makefile
+++ b/Makefile
@@ -143,7 +143,7 @@ no_opencl_support:
 	@echo "Warning: gpgpu-sim is building without opencl support. Make sure NVOPENCL_LIBDIR and NVOPENCL_INCDIR are set"
 
 $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
-	g++ -shared -Wl,-soname,libcudart_$(GPGPUSIM_BUILD).so \
+	g++ -shared -Wl,-soname,libcudart_$(GPGPUSIM_BUILD).so -Wl,--version-script=linux-so-version.txt\
 			$(SIM_OBJ_FILES_DIR)/libcuda/*.o \
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/*.o \
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/decuda_pred_table/*.o \
diff --git a/linux-so-version.txt b/linux-so-version.txt
new file mode 100644
index 0000000..40f775d
--- /dev/null
+++ b/linux-so-version.txt
@@ -0,0 +1,2 @@
+libcudart.so.9.1{
+};
-- 
cgit v1.3


From d97022fa5fee348e15f2b07ab83ee7aba4f008e0 Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Sat, 24 Mar 2018 18:08:37 -0400
Subject: The new CUDAs complain about version info not being embedded into the
 so file. Having gpgpu-sim ebed some version info to stop the error complaints

---
 Makefile             | 2 +-
 linux-so-version.txt | 2 ++
 2 files changed, 3 insertions(+), 1 deletion(-)
 create mode 100644 linux-so-version.txt

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index de6264c..2d0466e 100644
--- a/Makefile
+++ b/Makefile
@@ -143,7 +143,7 @@ no_opencl_support:
 	@echo "Warning: gpgpu-sim is building without opencl support. Make sure NVOPENCL_LIBDIR and NVOPENCL_INCDIR are set"
 
 $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
-	g++ -shared -Wl,-soname,libcudart_$(GPGPUSIM_BUILD).so \
+	g++ -shared -Wl,-soname,libcudart_$(GPGPUSIM_BUILD).so -Wl,--version-script=linux-so-version.txt\
 			$(SIM_OBJ_FILES_DIR)/libcuda/*.o \
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/*.o \
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/decuda_pred_table/*.o \
diff --git a/linux-so-version.txt b/linux-so-version.txt
new file mode 100644
index 0000000..40f775d
--- /dev/null
+++ b/linux-so-version.txt
@@ -0,0 +1,2 @@
+libcudart.so.9.1{
+};
-- 
cgit v1.3


From 12c121a0d3ba12659288719ef5af324208d94262 Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Fri, 17 May 2019 21:27:40 -0400
Subject: 9.2 sofile

---
 Makefile | 1 +
 1 file changed, 1 insertion(+)

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index 2d0466e..3db8ce8 100644
--- a/Makefile
+++ b/Makefile
@@ -163,6 +163,7 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.8.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.8.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.1; fi
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.2 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.2; fi
 
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
-- 
cgit v1.3


From 37a709bb69fcd9f2b0fe53a189e92e548164cc4b Mon Sep 17 00:00:00 2001
From: Mahmoud <abdallm@purdue.edu>
Date: Wed, 5 Jun 2019 19:18:29 -0400
Subject: adding new cuda 9 APIs to run the deepbench workloads

---
 Makefile                       |  2 ++
 libcuda/cuda_runtime_api.cc    | 70 +++++++++++++++++++++++++++++++++++++++++-
 linux-so-version.txt           |  4 +++
 src/abstract_hardware_model.cc | 22 +++++++++++++
 src/abstract_hardware_model.h  |  1 +
 src/gpgpu-sim/gpu-sim.cc       | 10 ++++++
 src/gpgpu-sim/gpu-sim.h        |  2 ++
 7 files changed, 110 insertions(+), 1 deletion(-)

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index 2d0466e..e1e9aaa 100644
--- a/Makefile
+++ b/Makefile
@@ -159,10 +159,12 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.5.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.5.5; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.6.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.6.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.6.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.6.5; fi
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.7.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.7.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.7.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.7.5; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.8.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.8.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.1; fi
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.2 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.2; fi
 
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 3a9d613..2cc84aa 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -974,6 +974,30 @@ __host__ cudaError_t CUDARTAPI cudaGetDeviceProperties(struct cudaDeviceProp *pr
 	}
 }
 
+#if (CUDART_VERSION >= 8000)
+cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks, const char *hostFunc, int blockSize, size_t dynamicSMemSize, unsigned int flags)
+{
+	printf("GPGPU-Sim PTX: cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags %p\n", hostFunc);
+	CUctx_st *context = GPGPUSim_Context();
+	function_info *entry = context->get_kernel(hostFunc);
+	printf("Calculate Maxium Active Block with function ptr=%p, blockSize=%d, SMemSize=%d\n", hostFunc, blockSize, dynamicSMemSize);
+	if (flags == cudaOccupancyDefault) {
+		//create kernel_info based on entry
+		dim3 gridDim(context->get_device()->get_gpgpu()->max_cta_per_core()
+                        * context->get_device()->get_gpgpu()->get_config().num_shader());
+		dim3 blockDim(blockSize);
+		kernel_info_t result(gridDim, blockDim, entry);
+		*numBlocks = context->get_device()->get_gpgpu()->get_max_cta(result);
+		printf("Maximum size is %d with gridDim %d and blockDim %d\n", *numBlocks, gridDim.x, blockDim.x);
+		return g_last_cudaError = cudaSuccess;
+	} else {
+		cuda_not_implemented(__my_func__,__LINE__);
+		return g_last_cudaError = cudaErrorUnknown;
+	}
+}
+
+#endif
+
 #if (CUDART_VERSION > 5000)
 __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device)
 {
@@ -3169,6 +3193,23 @@ __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t v
     return g_last_cudaError = cudaSuccess;
 }
 
+#if CUDART_VERSION >= 9000
+__host__  cudaError_t cudaFuncSetAttribute ( const void* func, enum cudaFuncAttribute attr, int value ) {
+
+	//ignore this Attribute for now, and the default is that carveout = cudaSharedmemCarveoutDefault;   //  (-1)
+	return g_last_cudaError = cudaSuccess;
+}
+
+#if CUDART_VERSION >= 9020
+__host__ __device__ unsigned __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim,
+                                                size_t sharedMem = 0,
+                                                void *stream = 0) {
+
+	return 1;
+}
+#endif
+
+#endif
 
 #endif
 
@@ -4071,6 +4112,33 @@ CUresult CUDAAPI cuMemHostRegister(void *p, size_t bytesize, unsigned int Flags)
 	printf("WARNING: this function has not been implemented yet.");
 	return CUDA_SUCCESS;
 }
+__host__ cudaError_t cudaHostRegister(void* ptr, size_t size, unsigned int flags)
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	printf("WARNING: this function has not been implemented yet.");
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t cudaProfilerStart ( )
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	printf("WARNING: this function has not been implemented yet.");
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t cudaProfilerStop ( )
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	printf("WARNING: this function has not been implemented yet.");
+	return g_last_cudaError = cudaSuccess;
+}
+
 #endif
 #if CUDART_VERSION >= 4000
 
@@ -4952,7 +5020,7 @@ CUresult CUDAAPI cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int *numBl
 	printf("WARNING: this function has not been implemented yet.");
 	return CUDA_SUCCESS;
 }
-    
+
 CUresult CUDAAPI cuOccupancyMaxPotentialBlockSize(int *minGridSize, int *blockSize, CUfunction func, CUoccupancyB2DSize blockSizeToDynamicSMemSize, size_t dynamicSMemSize, int blockSizeLimit)
 {
 	if(g_debug_execution >= 3){
diff --git a/linux-so-version.txt b/linux-so-version.txt
index c6e03e9..a7c2d3c 100644
--- a/linux-so-version.txt
+++ b/linux-so-version.txt
@@ -1,4 +1,8 @@
+libcudart.so.9.0{
+};
 libcudart.so.9.1{
 };
+libcudart.so.9.2{
+};
 libcuda.so.1{
 };
diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index cebdb25..023f51b 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -691,6 +691,28 @@ unsigned g_kernel_launch_latency;
 
 unsigned kernel_info_t::m_next_uid = 1;
 
+kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry)
+{
+    m_kernel_entry=entry;
+    m_grid_dim=gridDim;
+    m_block_dim=blockDim;
+    m_next_cta.x=0;
+    m_next_cta.y=0;
+    m_next_cta.z=0;
+    m_next_tid=m_next_cta;
+    m_num_cores_running=0;
+    m_uid = m_next_uid++;
+    m_param_mem = new memory_space_impl<8192>("param",64*1024);
+
+    //Jin: parent and child kernel management for CDP
+    m_parent_kernel = NULL;
+
+    //Jin: launch latency management
+    m_launch_latency = g_kernel_launch_latency;
+
+    volta_cache_config_set=false;
+}
+
 /*A snapshot of the texture mappings needs to be stored in the kernel's info as 
 kernels should use the texture bindings seen at the time of launch and textures
  can be bound/unbound asynchronously with respect to streams. */
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index 22ef509..64bbaa2 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -212,6 +212,7 @@ public:
 //      m_num_cores_running=0;
 //      m_param_mem=NULL;
 //   }
+   kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry);
    kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map<std::string, const struct cudaArray*> nameToCudaArray, std::map<std::string, const struct textureInfo*> nameToTextureInfo);
    ~kernel_info_t();
 
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index c1ba934..6de5845 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -806,6 +806,16 @@ int gpgpu_sim::shader_clock() const
    return m_config.core_freq/1000;
 }
 
+int gpgpu_sim::max_cta_per_core() const
+{
+   return m_shader_config->max_cta_per_core;
+}
+
+int gpgpu_sim::get_max_cta( const kernel_info_t &k ) const
+{
+   return m_shader_config->max_cta(k);
+}
+
 void gpgpu_sim::set_prop( cudaDeviceProp *prop )
 {
    m_cuda_properties = prop;
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index c8dad89..7336cac 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -455,6 +455,8 @@ public:
    int num_registers_per_block() const;
    int wrp_size() const;
    int shader_clock() const;
+   int max_cta_per_core() const;
+   int get_max_cta( const kernel_info_t &k ) const;
    const struct cudaDeviceProp *get_prop() const;
    enum divergence_support_t simd_model() const; 
 
-- 
cgit v1.3


From ce1fabd955ee208a38e73d5a5acfba2b26369a3d Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Sun, 9 Jun 2019 22:20:15 -0400
Subject: A bunch of boilerbplate to get 10.1 to compile. Still does not yet
 run. The way CUDA calls kerenels (even on old code) has changed.

---
 Makefile                    |  2 ++
 libcuda/cuda_api.h          |  2 ++
 libcuda/cuda_runtime_api.cc | 23 +++++++++++++++++++++++
 linux-so-version.txt        |  4 ++++
 setup_environment           |  2 +-
 5 files changed, 32 insertions(+), 1 deletion(-)

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index e1e9aaa..7a3bb47 100644
--- a/Makefile
+++ b/Makefile
@@ -165,6 +165,8 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.1; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.2 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.2; fi
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.0; fi
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.1; fi
 
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
diff --git a/libcuda/cuda_api.h b/libcuda/cuda_api.h
index 3808e8a..7ee26dc 100644
--- a/libcuda/cuda_api.h
+++ b/libcuda/cuda_api.h
@@ -234,9 +234,11 @@ typedef struct CUgraphicsResource_st *CUgraphicsResource; /**< CUDA graphics int
 typedef unsigned long long CUtexObject;                   /**< An opaque value that represents a CUDA texture object */
 typedef unsigned long long CUsurfObject;                  /**< An opaque value that represents a CUDA surface object */
 
+#if __CUDA_API_VERSION < 1010
 typedef struct CUuuid_st {                                /**< CUDA definition of UUID */
     char bytes[16];
 } CUuuid;
+#endif
 
 
 #if __CUDA_API_VERSION >= 4010
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 2cc84aa..fb3e07a 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -2763,6 +2763,29 @@ cudaError_t CUDARTAPI cudaDeviceSynchronize(void){
 	return g_last_cudaError = cudaSuccess;
 }
 
+void CUDARTAPI __cudaRegisterFatBinaryEnd( void **fatCubinHandle )
+{
+
+}
+
+unsigned CUDARTAPI __cudaPushCallConfiguration(dim3 gridDim,
+                                      dim3 blockDim, 
+                                      size_t sharedMem = 0, 
+                                      struct CUstream_st *stream = 0)
+{
+
+}
+
+cudaError_t CUDARTAPI __cudaPopCallConfiguration(
+  dim3         *gridDim,
+  dim3         *blockDim,
+  size_t       *sharedMem,
+  void         *stream
+)
+{
+    return g_last_cudaError = cudaSuccess;
+}
+
 void CUDARTAPI __cudaRegisterFunction(
 		void   **fatCubinHandle,
 		const char    *hostFun,
diff --git a/linux-so-version.txt b/linux-so-version.txt
index a7c2d3c..45c40dd 100644
--- a/linux-so-version.txt
+++ b/linux-so-version.txt
@@ -4,5 +4,9 @@ libcudart.so.9.1{
 };
 libcudart.so.9.2{
 };
+libcudart.so.10.0{
+};
+libcudart.so.10.1{
+};
 libcuda.so.1{
 };
diff --git a/setup_environment b/setup_environment
index b420584..17891ce 100644
--- a/setup_environment
+++ b/setup_environment
@@ -51,7 +51,7 @@ CC_VERSION=`gcc --version | head -1 | awk '{for(i=1;i<=NF;i++){ if(match($i,/^[0
 
 CUDA_VERSION_STRING=`$CUDA_INSTALL_PATH/bin/nvcc --version | awk '/release/ {print $5;}' | sed 's/,//'`;
 export CUDA_VERSION_NUMBER=`echo $CUDA_VERSION_STRING | sed 's/\./ /' | awk '{printf("%02u%02u", 10*int($1), 10*$2);}'`
-if [ $CUDA_VERSION_NUMBER -gt 9100 -o $CUDA_VERSION_NUMBER -lt 2030  ]; then
+if [ $CUDA_VERSION_NUMBER -gt 10100 -o $CUDA_VERSION_NUMBER -lt 2030  ]; then
 	echo "ERROR ** GPGPU-Sim version $GPGPUSIM_VERSION_STRING not tested with CUDA version $CUDA_VERSION_STRING (please see README)";
 	return
 fi
-- 
cgit v1.3


From eb3b077cddb24a734bb72f5c1588dee4034214a7 Mon Sep 17 00:00:00 2001
From: Mahmoud <abdallm@purdue.edu>
Date: Wed, 4 Sep 2019 21:38:57 -0400
Subject: adding trace class and  parss kernel info

---
 Makefile                                       |   4 +-
 libcuda/gpgpu_context.h                        |   1 +
 src/cuda-sim/ptx_ir.h                          |  10 +-
 src/cuda-sim/ptx_loader.cc                     |   2 +-
 src/cuda-sim/ptx_parser.cc                     |   4 +-
 src/gpgpu-sim/gpu-sim.cc                       |   9 +
 src/gpgpu-sim/gpu-sim.h                        |   7 +
 src/gpgpusim_entrypoint.cc                     |  32 ++-
 src/gpgpusim_entrypoint.h                      |   2 +
 src/trace-driven/gpgpusim_trace_driven_main.cc | 305 +++++++++++++++----------
 src/trace-driven/trace_driven.h                |  71 ++++++
 11 files changed, 322 insertions(+), 125 deletions(-)
 create mode 100644 src/trace-driven/trace_driven.h

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index caf0a17..5cd3101 100644
--- a/Makefile
+++ b/Makefile
@@ -144,7 +144,7 @@ no_opencl_support:
 	@echo "Warning: gpgpu-sim is building without opencl support. Make sure NVOPENCL_LIBDIR and NVOPENCL_INCDIR are set"
 
 $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
-	g++ -shared -Wl,-soname,libcudart_$(GPGPUSIM_BUILD).so -Wl,--version-script=linux-so-version.txt\
+	g++ -shared -Wl,-soname,libcudart.so -Wl,--version-script=linux-so-version.txt\
 			$(SIM_OBJ_FILES_DIR)/libcuda/*.o \
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/*.o \
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/decuda_pred_table/*.o \
@@ -171,7 +171,7 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 
 
 $(SIM_LIB_DIR)/gpgpusim.out: makedirs $(LIBS) cudalib $(SIM_LIB_DIR)/libcudart.so
-	g++ -std=c++0x -L$(SIM_LIB_DIR) -lcudart -o $(SIM_LIB_DIR)/gpgpusim.out src/trace-driven/gpgpusim_trace_driven_main.cc 
+	g++ -std=c++0x -L$(SIM_LIB_DIR) -I$(CUDA_INSTALL_PATH)/include -lcudart -o $(SIM_LIB_DIR)/gpgpusim.out src/trace-driven/gpgpusim_trace_driven_main.cc 
 
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h
index 45c5cdd..d3c5d74 100644
--- a/libcuda/gpgpu_context.h
+++ b/libcuda/gpgpu_context.h
@@ -65,6 +65,7 @@ class gpgpu_context {
 	const ptx_instruction* pc_to_instruction(unsigned pc); 
 	const warp_inst_t *ptx_fetch_inst( address_type pc );
 	unsigned translate_pc_to_ptxlineno(unsigned pc);
+	class gpgpu_sim *gpgpu_trace_sim_init_perf(int argc, const char *argv[]);
 };
 gpgpu_context* GPGPU_Context();
 
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index f4c5c37..8ceb60b 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -1354,13 +1354,13 @@ public:
    void list_param( FILE *fout ) const;
    void ptx_jit_config(std::map<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) ;
 
-   const struct gpgpu_ptx_sim_info* get_kernel_info () const
+   virtual const struct gpgpu_ptx_sim_info* get_kernel_info () const
    {
       assert (m_kernel_info.maxthreads == maxnt_id);
       return &m_kernel_info;
    }
 
-   const void set_kernel_info (const struct gpgpu_ptx_sim_info &info) {
+   virtual const void set_kernel_info (const struct gpgpu_ptx_sim_info &info) {
       m_kernel_info = info;
       m_kernel_info.ptx_version = 10*get_ptx_version().ver();
       m_kernel_info.sm_target = get_ptx_version().target();
@@ -1401,6 +1401,10 @@ public:
    // backward pointer
    class gpgpu_context* gpgpu_ctx;
 
+protected:
+   //Registers/shmem/etc. used (from ptxas -v), loaded from ___.ptxinfo along with ___.ptx
+   struct gpgpu_ptx_sim_info m_kernel_info;
+
 private:
    unsigned maxnt_id;
    unsigned m_uid;
@@ -1424,8 +1428,6 @@ private:
    std::map<std::string,unsigned> labels;
    unsigned num_reconvergence_pairs;
 
-   //Registers/shmem/etc. used (from ptxas -v), loaded from ___.ptxinfo along with ___.ptx
-   struct gpgpu_ptx_sim_info m_kernel_info;
 
    symbol_table *m_symtab;
 
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index dca3cec..207fa1f 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -37,7 +37,7 @@
 
 /// extern prototypes
 
-extern int ptx_error( yyscan_t yyscanner, const char *s );
+extern int ptx_error( yyscan_t yyscanner, ptx_recognizer* recognizer, const char *s );
 extern int ptx_lex_init(yyscan_t* scanner);
 extern void ptx_set_in(FILE * _in_str ,yyscan_t yyscanner );
 extern int ptx_parse(yyscan_t scanner, ptx_recognizer* recognizer);
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 81b70af..9e9d500 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -35,7 +35,7 @@ typedef void * yyscan_t;
 
 extern int ptx_get_lineno (yyscan_t yyscanner );
 extern YYSTYPE* ptx_get_lval (yyscan_t yyscanner );
-extern int ptx_error( yyscan_t yyscanner, const char *s );
+extern int ptx_error( yyscan_t yyscanner, ptx_recognizer* recognizer, const char *s );
 extern int ptx_lex_init(yyscan_t* scanner);
 extern void ptx_set_in(FILE * _in_str ,yyscan_t yyscanner );
 extern FILE *ptx_get_in (yyscan_t yyscanner );
@@ -226,7 +226,7 @@ void ptx_recognizer::parse_error_impl( const char *file, unsigned line, const ch
 
    g_error_detected = 1;
    printf("%s:%u: Parse error: %s (%s:%u)\n\n", gpgpu_ctx->g_filename, ptx_get_lineno(scanner), buf, file, line);
-   ptx_error(scanner, NULL);
+   ptx_error(scanner, this, NULL);
    abort();
    exit(1);
 }
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index a151379..234a8d6 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -560,6 +560,15 @@ void gpgpu_sim_config::reg_options(option_parser_t opp)
     option_parser_register(opp, "-gpgpu_cdp_enabled", OPT_BOOL, 
                           &(gpgpu_ctx->device_runtime->g_cdp_enabled), "Turn on CDP",
                           "0");
+
+    //Trace driven mode parameters
+    option_parser_register(opp, "-trace_driven_mode", OPT_BOOL,
+                          &trace_driven_mode, "Turn on trace_driven_mode",
+                          "0");
+    option_parser_register(opp, "-trace", OPT_CSTR,
+                          &g_traces_filename, "traces kernel file"
+                          "traces kernel file directory",
+                          "./traces/kernelslist.g");
 }
 
 /////////////////////////////////////////////////////////////////////////////
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index 76c7a06..6b57bd8 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -348,6 +348,9 @@ public:
     size_t sync_depth_limit() const {return runtime_sync_depth_limit; }
     size_t pending_launch_count_limit() const {return runtime_pending_launch_count_limit;}
 
+    unsigned is_trace_driven_mode() const { return trace_driven_mode; }
+    char* get_traces_filename() const { return g_traces_filename; }
+
 private:
     void init_clock_domains(void ); 
 
@@ -401,6 +404,10 @@ private:
     unsigned int gpgpu_compute_capability_minor;
     unsigned long long liveness_message_freq; 
 
+    //trace driven mode options
+    bool trace_driven_mode;
+    char *g_traces_filename;
+
     friend class gpgpu_sim;
 };
 
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc
index 816159f..4da74c0 100644
--- a/src/gpgpusim_entrypoint.cc
+++ b/src/gpgpusim_entrypoint.cc
@@ -60,7 +60,6 @@ class stream_manager* g_stream_manager()  {
 	return GPGPUsim_ctx_ptr()->g_stream_manager;
 }
 
-static void print_simulation_time();
 
 void *gpgpu_sim_thread_sequential(void*)
 {
@@ -242,6 +241,37 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf()
    return GPGPUsim_ctx_ptr()->g_the_gpu;
 }
 
+gpgpu_sim *gpgpu_context::gpgpu_trace_sim_init_perf(int argc, const char *argv[])
+{
+   srand(1);
+   print_splash();
+   func_sim->read_sim_environment_variables();
+   ptx_parser->read_parser_environment_variables();
+   option_parser_t opp = option_parser_create();
+
+   ptx_reg_options(opp);
+   func_sim->ptx_opcocde_latency_options(opp);
+
+   icnt_reg_options(opp);
+   GPGPUsim_ctx_ptr()->g_the_gpu_config = new gpgpu_sim_config(this);
+   GPGPUsim_ctx_ptr()->g_the_gpu_config->reg_options(opp); // register GPU microrachitecture options
+
+   option_parser_cmdline(opp, argc, argv); // parse configuration options
+   fprintf(stdout, "GPGPU-Sim: Configuration options:\n\n");
+   option_parser_print(opp, stdout);
+   // Set the Numeric locale to a standard locale where a decimal point is a "dot" not a "comma"
+   // so it does the parsing correctly independent of the system environment variables
+   assert(setlocale(LC_NUMERIC,"C"));
+   GPGPUsim_ctx_ptr()->g_the_gpu_config->init();
+
+   GPGPUsim_ctx_ptr()->g_the_gpu = new gpgpu_sim(*(GPGPUsim_ctx_ptr()->g_the_gpu_config), this);
+   GPGPUsim_ctx_ptr()->g_stream_manager = new stream_manager((GPGPUsim_ctx_ptr()->g_the_gpu), func_sim->g_cuda_launch_blocking);
+
+   GPGPUsim_ctx_ptr()->g_simulation_starttime = time((time_t *)NULL);
+
+   return GPGPUsim_ctx_ptr()->g_the_gpu;
+}
+
 void start_sim_thread(int api)
 {
     if( GPGPUsim_ctx_ptr()->g_sim_done ) {
diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h
index 887b3c8..4a8d796 100644
--- a/src/gpgpusim_entrypoint.h
+++ b/src/gpgpusim_entrypoint.h
@@ -84,4 +84,6 @@ class stream_manager* g_stream_manager();
 
 int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid );
 
+void print_simulation_time();
+
 #endif
diff --git a/src/trace-driven/gpgpusim_trace_driven_main.cc b/src/trace-driven/gpgpusim_trace_driven_main.cc
index 6553fd7..881e351 100644
--- a/src/trace-driven/gpgpusim_trace_driven_main.cc
+++ b/src/trace-driven/gpgpusim_trace_driven_main.cc
@@ -1,138 +1,213 @@
 //developed by Mahmoud Khairy, Purdue Univ
 //abdallm@purdue.edu
 
-#include "../abstract_hardware_model.h"
+//#include "../abstract_hardware_model.h"
 #include <time.h>
 #include <stdio.h>
+#include <vector>
+#include <iostream>
+#include <fstream>
+#include <string>
+#include <sstream>
+#include <stdio.h>
+#include <math.h>
+
 
-#include "../option_parser.h"
-#include "../cuda-sim/cuda-sim.h"
-#include "../cuda-sim/ptx_ir.h"
-#include "../cuda-sim/ptx_parser.h"
+//#include "../option_parser.h"
+//#include "../cuda-sim/cuda-sim.h"
+//#include "../cuda-sim/ptx_ir.h"
+//#include "../cuda-sim/ptx_parser.h"
 #include "../gpgpu-sim/gpu-sim.h"
-#include "../gpgpu-sim/icnt_wrapper.h"
+//#include "../gpgpu-sim/icnt_wrapper.h"
+//#include "../gpgpu-sim/icnt_wrapper.h"
+#include "../../libcuda/gpgpu_context.h"
+#include "trace_driven.h"
+
 //#include "../stream_manager.h"
 
 
+void arguments_check();
+
+int main ( int argc, const char **argv )
+{
+
+	gpgpu_context* m_gpgpu_context = GPGPU_Context();
+	gpgpu_sim * m_gpgpu_sim = m_gpgpu_context->gpgpu_trace_sim_init_perf(argc,argv);
+	m_gpgpu_sim->init();
+
+	//for each kernel
+	//load file
+	//parse and create kernel info
+	//launch
+	//while loop till the end of the end kernel execution
+	//prints stats
+
+	trace_parser tracer(m_gpgpu_sim->get_config().get_traces_filename(), m_gpgpu_sim, m_gpgpu_context);
+
+	std::vector<std::string> kernellist;
+	tracer.parse_kernellist_file(kernellist);
+
+	for(unsigned i=0; i<kernellist.size(); ++i) {
+
+		trace_kernel_info_t* kernel_info  = tracer.parse_kernel_info(kernellist[i]);
+		m_gpgpu_sim->launch(kernel_info);
+
+		bool active = false;
+		bool sim_cycles = false;
+		bool break_limit = false;
+
+		do {
+			if(!m_gpgpu_sim->active())
+				break;
+
+			//performance simulation
+			if( m_gpgpu_sim->active() ) {
+				m_gpgpu_sim->cycle();
+				sim_cycles = true;
+				m_gpgpu_sim->deadlock_check();
+			}else {
+				if(m_gpgpu_sim->cycle_insn_cta_max_hit()){
+					g_stream_manager()->stop_all_running_kernels();
+					break_limit = true;
+				}
+			}
+
+			active=m_gpgpu_sim->active() ;
+
+		} while( active );
 
-gpgpu_sim_config g_the_gpu_config;
-gpgpu_sim *g_the_gpu;
-time_t g_simulation_starttime;
+		tracer.kernel_finalizer(kernel_info);
 
-#define MAX(a,b) (((a)>(b))?(a):(b))
+		m_gpgpu_sim->print_stats();
 
-static void print_simulation_time();
+		if(sim_cycles) {
+			m_gpgpu_sim->update_stats();
+			print_simulation_time();
+		}
 
+		if(break_limit) {
+			printf("GPGPU-Sim: ** break due to reaching the maximum cycles (or instructions) **\n");
+			fflush(stdout);
+			exit(1);
+		}
+	}
 
-int main ( int argc, char **argv )
+	return 1;
+}
+
+trace_parser::trace_parser(const char* kernellist_filepath, gpgpu_sim * m_gpgpu_sim, gpgpu_context* m_gpgpu_context)
 {
-   srand(1);
-
-   option_parser_t opp = option_parser_create();
-
-   icnt_reg_options(opp);
-   g_the_gpu_config.reg_options(opp); // register GPU microrachitecture options
-   ptx_reg_options(opp);
-   ptx_opcocde_latency_options(opp);
-
-   //ptx_opcocde_latency_options(opp);  //do this for trace driven
-
-   fprintf(stdout, "I am here:\n\n");
-   option_parser_cmdline(opp, argc, (const char **)argv); // parse configuration options
-
-   fprintf(stdout, "GPGPU-Sim: Configuration options:\n\n");
-   option_parser_print(opp, stdout);
-   // Set the Numeric locale to a standard locale where a decimal point is a "dot" not a "comma"
-   // so it does the parsing correctly independent of the system environment variables
-   assert(setlocale(LC_NUMERIC,"C"));
-   g_the_gpu_config.init();
-
-   g_the_gpu = new gpgpu_sim(g_the_gpu_config);
-   //g_stream_manager = new stream_manager(g_the_gpu,g_cuda_launch_blocking);
-
-   //load file
-   //create kernel info
-   //launch
-   //while loop till the end
-   //prints stats
-   //g_the_gpu->launch(grid);
-
-   g_simulation_starttime = time((time_t *)NULL);
-
-	  bool active = false;
-	  bool sim_cycles = false;
-	  bool break_limit = false;
-
-	  g_the_gpu->init();
-	  do {
-		  // check if a kernel has completed
-		  // launch operation on device if one is pending and can be run
-
-		  // Need to break this loop when a kernel completes. This was a
-		  // source of non-deterministic behaviour in GPGPU-Sim (bug 147).
-		  // If another stream operation is available, g_the_gpu remains active,
-		  // causing this loop to not break. If the next operation happens to be
-		  // another kernel, the gpu is not re-initialized and the inter-kernel
-		  // behaviour may be incorrect. Check that a kernel has finished and
-		  // no other kernel is currently running.
-		  if(!g_the_gpu->active())
-			  break;
-
-		  //performance simulation
-		  if( g_the_gpu->active() ) {
-			  g_the_gpu->cycle();
-			  sim_cycles = true;
-			  g_the_gpu->deadlock_check();
-		  }else {
-			  if(g_the_gpu->cycle_insn_cta_max_hit()){
-				  g_stream_manager->stop_all_running_kernels();
-				  break_limit = true;
-			  }
-		  }
-
-		  active=g_the_gpu->active() ;
-
-	  } while( active );
-
-		 printf("GPGPU-Sim: ** STOP simulation thread (no work) **\n");
-		 fflush(stdout);
-
-	  g_the_gpu->print_stats();
-
-	  if(sim_cycles) {
-		  g_the_gpu->update_stats();
-		  print_simulation_time();
-	  }
-
-
-		 printf("GPGPU-Sim: *** simulation thread exiting ***\n");
-		 fflush(stdout);
-
-	  if(break_limit) {
-		printf("GPGPU-Sim: ** break due to reaching the maximum cycles (or instructions) **\n");
+
+	this->m_gpgpu_sim = m_gpgpu_sim;
+	this->m_gpgpu_context = m_gpgpu_context;
+	kernellist_filename = kernellist_filepath;
+}
+
+void trace_parser::parse_kernellist_file(std::vector<std::string>& kernellist) {
+
+	ifs.open(kernellist_filename);
+
+	if (!ifs.is_open()) {
+		std::cout << "Unable to open file: " <<kernellist_filename<<std::endl;
 		exit(1);
-	  }
+	}
+
+	std::string directory(kernellist_filename);
+	const size_t last_slash_idx = directory.rfind('/');
+	if (std::string::npos != last_slash_idx)
+	{
+		directory = directory.substr(0, last_slash_idx);
+	}
+
+	std::string line, filepath;
+	while(!ifs.eof()) {
+		getline(ifs, line);
+		if(line.empty())
+			continue;
+		filepath = directory+"/"+line;
+		kernellist.push_back(filepath);
+	}
+
+	ifs.close();
+}
 
-      return 1;
+
+trace_kernel_info_t* trace_parser::parse_kernel_info(const std::string& kerneltraces_filepath) {
+
+	ifs.open(kerneltraces_filepath.c_str());
+
+	if (!ifs.is_open()) {
+		std::cout << "Unable to open file: " <<kerneltraces_filepath<<std::endl;
+		exit(1);
+	}
+
+	std::cout << "Processing kernel " <<kerneltraces_filepath<<std::endl;
+
+	unsigned grid_dim_x=0, grid_dim_y=0, grid_dim_z=0, tb_dim_x=0, tb_dim_y=0, tb_dim_z=0;
+	unsigned shmem=0, nregs=0, cuda_stream_id=0, kernel_id=0;
+	std::string line;
+	std::stringstream ss;
+	std::string string1, string2;
+	std::string  kernel_name;
+
+	while(!ifs.eof()) {
+		getline(ifs, line);
+
+		if (line.length() == 0) {
+			continue;
+		}
+		else if(line[0] == '#'){
+			break;  //the begin of the instruction stream
+		}
+		else if(line[0] == '-') {
+			ss.str(line);
+			ss.ignore();
+			ss>>string1>>string2;
+			if(string1 == "kernel" && string2 == "name") {
+				const size_t equal_idx = line.find('=');
+				kernel_name = line.substr(equal_idx+1);
+			}
+			else if(string1 == "kernel" && string2 == "id") {
+				sscanf(line.c_str(), "-kernel id = %d", &kernel_id);
+			}
+			else if(string1 == "grid" && string2 == "dim") {
+				sscanf(line.c_str(), "-grid dim = (%d,%d,%d)", &grid_dim_x, &grid_dim_y, &grid_dim_z);
+			}
+			else if (string1 == "block" && string2 == "dim") {
+				sscanf(line.c_str(), "-block dim = (%d,%d,%d)", &tb_dim_x, &tb_dim_y, &tb_dim_z);
+			}
+			else if (string1 == "shmem") {
+				sscanf(line.c_str(), "-shmem = %d", &shmem);
+			}
+			else if (string1 == "nregs") {
+				sscanf(line.c_str(), "-nregs = %d", &nregs);
+			}
+			else if (string1 == "cuda" && string2 == "stream") {
+				sscanf(line.c_str(), "-cuda stream id = %d", &cuda_stream_id);
+			}
+			continue;
+		}
+	}
+
+	gpgpu_ptx_sim_info info;
+	info.smem = shmem;
+	info.regs = nregs;
+	dim3 gridDim(grid_dim_x, grid_dim_y, grid_dim_z);
+	dim3 blockDim(tb_dim_x, tb_dim_y, tb_dim_z);
+	trace_function_info* function_info = new trace_function_info(info, m_gpgpu_context);
+	trace_kernel_info_t* kernel_info =  new trace_kernel_info_t(gridDim, blockDim, function_info, kerneltraces_filepath);
+
+	return kernel_info;
 }
 
-void print_simulation_time()
-{
-   time_t current_time, difference, d, h, m, s;
-   current_time = time((time_t *)NULL);
-   difference = MAX(current_time - g_simulation_starttime, 1);
-
-   d = difference/(3600*24);
-   h = difference/3600 - 24*d;
-   m = difference/60 - 60*(h + 24*d);
-   s = difference - 60*(m + 60*(h + 24*d));
-
-   fflush(stderr);
-   printf("\n\ngpgpu_simulation_time = %u days, %u hrs, %u min, %u sec (%u sec)\n",
-          (unsigned)d, (unsigned)h, (unsigned)m, (unsigned)s, (unsigned)difference );
-   printf("gpgpu_simulation_rate = %u (inst/sec)\n", (unsigned)(g_the_gpu->gpu_tot_sim_insn / difference) );
-   printf("gpgpu_simulation_rate = %u (cycle/sec)\n", (unsigned)(gpu_tot_sim_cycle / difference) );
-   fflush(stdout);
+
+void trace_parser::kernel_finalizer(trace_kernel_info_t* kernel_info){
+	if (ifs.is_open())
+		ifs.close();
+
+	delete kernel_info->entry();
+	delete kernel_info;
 }
 
 
+
diff --git a/src/trace-driven/trace_driven.h b/src/trace-driven/trace_driven.h
new file mode 100644
index 0000000..91ff90b
--- /dev/null
+++ b/src/trace-driven/trace_driven.h
@@ -0,0 +1,71 @@
+//developed by Mahmoud Khairy, Purdue Univ
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <assert.h>
+
+#ifndef TRACE_DRIVEN_H
+#define TRACE_DRIVEN_H
+
+#include "../abstract_hardware_model.h"
+
+
+class trace_function_info: public function_info {
+public:
+	trace_function_info(const struct gpgpu_ptx_sim_info &info, gpgpu_context* m_gpgpu_context):function_info(0, m_gpgpu_context ) {
+		m_kernel_info = info;
+	}
+
+	virtual const struct gpgpu_ptx_sim_info* get_kernel_info () const {
+		return &m_kernel_info;
+	}
+
+	virtual const void set_kernel_info (const struct gpgpu_ptx_sim_info &info) {
+		m_kernel_info = info;
+	}
+
+private:
+
+
+};
+
+class trace_kernel_info_t: public kernel_info_t {
+public:
+	trace_kernel_info_t(dim3 gridDim, dim3 blockDim, trace_function_info* m_function_info, const std::string& filepath):kernel_info_t(gridDim, blockDim, m_function_info) {
+		traces_filepath = filepath;
+	}
+
+private:
+	std::string traces_filepath;
+
+};
+
+class trace_warp_inst_t: public warp_inst_t {
+public:
+
+	trace_warp_inst_t() {
+	}
+
+private:
+
+
+};
+
+class trace_parser {
+public:
+	trace_parser(const char* kernellist_filepath, gpgpu_sim * m_gpgpu_sim, gpgpu_context* m_gpgpu_context);
+
+	void parse_kernellist_file(std::vector<std::string>& kernellist);
+	trace_kernel_info_t* parse_kernel_info(const std::string& kerneltraces_filepath);
+	void kernel_finalizer(trace_kernel_info_t* kernel_info);
+
+private:
+
+	std::string kernellist_filename;
+	std::ifstream ifs;
+	gpgpu_sim * m_gpgpu_sim;
+	gpgpu_context* m_gpgpu_context;
+
+};
+
+#endif
-- 
cgit v1.3


From 5db69b3e5b058c030075c066db64922bf1e6af02 Mon Sep 17 00:00:00 2001
From: Mahmoud <abdallm@purdue.edu>
Date: Thu, 19 Sep 2019 22:06:15 -0400
Subject: support trace-driven and changing the shader.cc

---
 Makefile                                       |  10 +-
 src/abstract_hardware_model.cc                 |   7 +-
 src/gpgpu-sim/gpu-sim.cc                       |  10 +-
 src/gpgpu-sim/mem_latency_stat.cc              |   5 +-
 src/gpgpu-sim/shader.cc                        | 170 +++++++++++------
 src/gpgpu-sim/shader.h                         |   7 +-
 src/trace-driven/gpgpusim_trace_driven_main.cc | 253 +++++++++++++++++++------
 src/trace-driven/trace_driven.h                |  53 +++++-
 src/trace-driven/trace_opcode.h                |   3 +-
 9 files changed, 386 insertions(+), 132 deletions(-)

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index 5cd3101..73f22c6 100644
--- a/Makefile
+++ b/Makefile
@@ -171,7 +171,15 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 
 
 $(SIM_LIB_DIR)/gpgpusim.out: makedirs $(LIBS) cudalib $(SIM_LIB_DIR)/libcudart.so
-	g++ -std=c++0x -L$(SIM_LIB_DIR) -I$(CUDA_INSTALL_PATH)/include -lcudart -o $(SIM_LIB_DIR)/gpgpusim.out src/trace-driven/gpgpusim_trace_driven_main.cc 
+	ar rvs   $(SIM_LIB_DIR)/libcudart_static.a\
+			$(SIM_OBJ_FILES_DIR)/libcuda/*.o \
+			$(SIM_OBJ_FILES_DIR)/cuda-sim/*.o \
+			$(SIM_OBJ_FILES_DIR)/cuda-sim/decuda_pred_table/*.o \
+			$(SIM_OBJ_FILES_DIR)/gpgpu-sim/*.o \
+			$(SIM_OBJ_FILES_DIR)/$(INTERSIM)/*.o \
+			$(SIM_OBJ_FILES_DIR)/*.o \
+			$(MCPAT)
+	g++ -std=c++0x -L$(SIM_LIB_DIR) -I$(CUDA_INSTALL_PATH)/include -lcudart -lm -lz -lGL -pthread -o $(SIM_LIB_DIR)/gpgpusim.out src/trace-driven/gpgpusim_trace_driven_main.cc 
 
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index e8e4b51..063aa8d 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -471,7 +471,8 @@ void warp_inst_t::generate_mem_accesses()
     }
 
     if ( space.get_type() == global_space ) {
-        m_config->gpgpu_ctx->stats->ptx_file_line_stats_add_uncoalesced_gmem( pc, m_accessq.size() - starting_queue_size );
+    	//TO DO: check here
+       // m_config->gpgpu_ctx->stats->ptx_file_line_stats_add_uncoalesced_gmem( pc, m_accessq.size() - starting_queue_size );
     }
     m_mem_accesses_created=true;
 }
@@ -706,7 +707,9 @@ void warp_inst_t::completed( unsigned long long cycle ) const
 {
    unsigned long long latency = cycle - issue_cycle; 
    assert(latency <= cycle); // underflow detection 
-   m_config->gpgpu_ctx->stats->ptx_file_line_stats_add_latency(pc, latency * active_count());  
+   //check the trace mode here
+   //TO DO
+   //m_config->gpgpu_ctx->stats->ptx_file_line_stats_add_latency(pc, latency * active_count());
 }
 
 
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index a70185e..ed96d42 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -66,6 +66,7 @@
 #include "stats.h"
 #include "../cuda-sim/cuda_device_runtime.h"
 #include "../../libcuda/gpgpu_context.h"
+#include "../trace-driven/trace_driven.h"
 
 #ifdef GPGPUSIM_POWER_MODEL
 #include "power_interface.h"
@@ -1482,7 +1483,12 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel )
     for (unsigned i = start_thread; i<end_thread; i++) {
         m_threadState[i].m_cta_id = free_cta_hw_id;
         unsigned warp_id = i/m_config->warp_size;
-        nthreads_in_block += ptx_sim_init_thread(kernel,&m_thread[i],m_sid,i,cta_size-(i-start_thread),m_config->n_thread_per_shader,this,free_cta_hw_id,warp_id,m_cluster->get_gpu());
+        if(m_gpu->get_config().is_trace_driven_mode()) {
+        	trace_shader_core_ctx* trace_core = static_cast<trace_shader_core_ctx*> (this);
+        	nthreads_in_block += trace_core->trace_sim_inc_thread(kernel);
+        }
+        else
+        	nthreads_in_block += ptx_sim_init_thread(kernel,&m_thread[i],m_sid,i,cta_size-(i-start_thread),m_config->n_thread_per_shader,this,free_cta_hw_id,warp_id,m_cluster->get_gpu());
         m_threadState[i].m_active = true; 
         // load thread local memory and register file
         if(m_gpu->resume_option == 1 && kernel.get_uid() == m_gpu->resume_kernel && ctaid >= m_gpu->resume_CTA && ctaid < m_gpu->checkpoint_CTA_t )
@@ -1512,7 +1518,7 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel )
     m_barriers.allocate_barrier(free_cta_hw_id,warps);
 
     // initialize the SIMT stacks and fetch hardware
-    init_warps( free_cta_hw_id, start_thread, end_thread, ctaid, cta_size, kernel.get_uid());
+    init_warps( free_cta_hw_id, start_thread, end_thread, ctaid, cta_size, kernel);
     m_n_active_cta++;
 
     shader_CTA_count_log(m_sid, 1);
diff --git a/src/gpgpu-sim/mem_latency_stat.cc b/src/gpgpu-sim/mem_latency_stat.cc
index a1b43a8..1980a3b 100644
--- a/src/gpgpu-sim/mem_latency_stat.cc
+++ b/src/gpgpu-sim/mem_latency_stat.cc
@@ -195,8 +195,9 @@ void memory_stats_t::memlatstat_dram_access(mem_fetch *mf)
       }
       mem_access_type_stats[mf->get_access_type()][dram_id][bank]++;
    }
-   if (mf->get_pc() != (unsigned)-1) 
-      m_gpu->gpgpu_ctx->stats->ptx_file_line_stats_add_dram_traffic(mf->get_pc(), mf->get_data_size());
+	 //TO DO: check here
+   //if (mf->get_pc() != (unsigned)-1)
+     // m_gpu->gpgpu_ctx->stats->ptx_file_line_stats_add_dram_traffic(mf->get_pc(), mf->get_data_size());
 }
 
 void memory_stats_t::memlatstat_icnt2mem_pop(mem_fetch *mf)
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index 86508f6..8dddc36 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -46,6 +46,7 @@
 #include "traffic_breakdown.h"
 #include "shader_trace.h"
 #include "../../libcuda/gpgpu_context.h"
+#include "../trace-driven/trace_driven.h"
 
 #define PRIORITIZE_MSHR_OVER_WB 1
 #define MAX(a,b) (((a)>(b))?(a):(b))
@@ -453,9 +454,11 @@ void shader_core_ctx::reinit(unsigned start_thread, unsigned end_thread, bool re
    }
 }
 
-void shader_core_ctx::init_warps( unsigned cta_id, unsigned start_thread, unsigned end_thread, unsigned ctaid, int cta_size, unsigned kernel_id )
+void shader_core_ctx::init_warps( unsigned cta_id, unsigned start_thread, unsigned end_thread, unsigned ctaid, int cta_size, kernel_info_t &kernel )
 {
+	//
     address_type start_pc = next_pc(start_thread);
+    unsigned kernel_id = kernel.get_uid();
     if (m_config->model == POST_DOMINATOR) {
         unsigned start_warp = start_thread / m_config->warp_size;
         unsigned warp_per_cta =  cta_size / m_config->warp_size;
@@ -482,8 +485,10 @@ void shader_core_ctx::init_warps( unsigned cta_id, unsigned start_thread, unsign
                 m_simt_stack[i]->resume(fname);
                 m_simt_stack[i]->get_pdom_stack_top_info(&pc,&rpc);
                 for (unsigned t = 0; t < m_config->warp_size; t++) {
-                  m_thread[i * m_config->warp_size + t]->set_npc(pc);
-                  m_thread[i * m_config->warp_size + t]->update_pc();
+                	if(m_thread != NULL) {
+                		m_thread[i * m_config->warp_size + t]->set_npc(pc);
+                		m_thread[i * m_config->warp_size + t]->update_pc();
+                	}
                 }   
                 start_pc=pc;
               }
@@ -493,6 +498,11 @@ void shader_core_ctx::init_warps( unsigned cta_id, unsigned start_thread, unsign
             m_not_completed += n_active;
             ++m_active_warps;
       }
+
+        if(m_gpu->get_config().is_trace_driven_mode()){
+        	trace_shader_core_ctx* trace_core = static_cast<trace_shader_core_ctx*> (this);
+            trace_core->init_traces(  start_warp,  end_warp, kernel );
+        }
    }
 }
 
@@ -741,33 +751,45 @@ void shader_core_stats::visualizer_print( gzFile visualizer_file )
                                         other memory spaces */
 void shader_core_ctx::decode()
 {
-    if( m_inst_fetch_buffer.m_valid ) {
-        // decode 1 or 2 instructions and place them into ibuffer
-        address_type pc = m_inst_fetch_buffer.m_pc;
-        const warp_inst_t* pI1 = m_gpu->gpgpu_ctx->ptx_fetch_inst(pc);
-        m_warp[m_inst_fetch_buffer.m_warp_id].ibuffer_fill(0,pI1);
-        m_warp[m_inst_fetch_buffer.m_warp_id].inc_inst_in_pipeline();
-        if( pI1 ) {
-            m_stats->m_num_decoded_insn[m_sid]++;
-            if(pI1->oprnd_type==INT_OP){
-                m_stats->m_num_INTdecoded_insn[m_sid]++;
-            }else if(pI1->oprnd_type==FP_OP) {
-            	m_stats->m_num_FPdecoded_insn[m_sid]++;
-            }
-           const warp_inst_t* pI2 = m_gpu->gpgpu_ctx->ptx_fetch_inst(pc+pI1->isize);
-           if( pI2 ) {
-               m_warp[m_inst_fetch_buffer.m_warp_id].ibuffer_fill(1,pI2);
-               m_warp[m_inst_fetch_buffer.m_warp_id].inc_inst_in_pipeline();
-               m_stats->m_num_decoded_insn[m_sid]++;
-               if(pI2->oprnd_type==INT_OP){
-                   m_stats->m_num_INTdecoded_insn[m_sid]++;
-               }else if(pI2->oprnd_type==FP_OP) {
-            	   m_stats->m_num_FPdecoded_insn[m_sid]++;
-               }
-           }
-        }
-        m_inst_fetch_buffer.m_valid = false;
-    }
+	if( m_inst_fetch_buffer.m_valid ) {
+		// decode 1 or 2 instructions and place them into ibuffer
+		address_type pc = m_inst_fetch_buffer.m_pc;
+		const warp_inst_t* pI1;
+		if(m_gpu->get_config().is_trace_driven_mode()){
+			trace_shader_core_ctx* trace_core = static_cast<trace_shader_core_ctx*> (this);
+			pI1 = trace_core->m_trace_warp[m_inst_fetch_buffer.m_warp_id].get_next_inst();
+		}
+		else
+			 pI1 = m_gpu->gpgpu_ctx->ptx_fetch_inst(pc);
+		m_warp[m_inst_fetch_buffer.m_warp_id].ibuffer_fill(0,pI1);
+		m_warp[m_inst_fetch_buffer.m_warp_id].inc_inst_in_pipeline();
+		if( pI1 ) {
+			m_stats->m_num_decoded_insn[m_sid]++;
+			if(pI1->oprnd_type==INT_OP){
+				m_stats->m_num_INTdecoded_insn[m_sid]++;
+			}else if(pI1->oprnd_type==FP_OP) {
+				m_stats->m_num_FPdecoded_insn[m_sid]++;
+			}
+			const warp_inst_t* pI2;
+			if(m_gpu->get_config().is_trace_driven_mode()){
+				trace_shader_core_ctx* trace_core = static_cast<trace_shader_core_ctx*> (this);
+				pI2 = trace_core->m_trace_warp[m_inst_fetch_buffer.m_warp_id].get_next_inst();
+			}
+			else
+				pI2 = m_gpu->gpgpu_ctx->ptx_fetch_inst(pc+pI1->isize);
+			if( pI2 ) {
+				m_warp[m_inst_fetch_buffer.m_warp_id].ibuffer_fill(1,pI2);
+				m_warp[m_inst_fetch_buffer.m_warp_id].inc_inst_in_pipeline();
+				m_stats->m_num_decoded_insn[m_sid]++;
+				if(pI2->oprnd_type==INT_OP){
+					m_stats->m_num_INTdecoded_insn[m_sid]++;
+				}else if(pI2->oprnd_type==FP_OP) {
+					m_stats->m_num_FPdecoded_insn[m_sid]++;
+				}
+			}
+		}
+		m_inst_fetch_buffer.m_valid = false;
+	}
 }
 
 void shader_core_ctx::fetch()
@@ -778,7 +800,12 @@ void shader_core_ctx::fetch()
             mem_fetch *mf = m_L1I->next_access();
             m_warp[mf->get_wid()].clear_imiss_pending();
             m_inst_fetch_buffer = ifetch_buffer_t(m_warp[mf->get_wid()].get_pc(), mf->get_access_size(), mf->get_wid());
-            assert( m_warp[mf->get_wid()].get_pc() == (mf->get_addr()-PROGRAM_MEM_START)); // Verify that we got the instruction we were expecting.
+            if(m_gpu->get_config().is_trace_driven_mode()){
+            	trace_shader_core_ctx* trace_core = static_cast<trace_shader_core_ctx*> (this);
+        		assert( trace_core->m_trace_warp[mf->get_wid()].get_pc() == (address_type)(mf->get_addr()-PROGRAM_MEM_START));
+            }
+            else
+            	assert( m_warp[mf->get_wid()].get_pc() == (mf->get_addr()-PROGRAM_MEM_START)); // Verify that we got the instruction we were expecting.
             m_inst_fetch_buffer.m_valid = true;
             m_warp[mf->get_wid()].set_last_fetch(m_gpu->gpu_sim_cycle);
             delete mf;
@@ -797,10 +824,14 @@ void shader_core_ctx::fetch()
                         if( m_threadState[tid].m_active == true ) {
                             m_threadState[tid].m_active = false; 
                             unsigned cta_id = m_warp[warp_id].get_cta_id();
-                            register_cta_thread_exit(cta_id, &(m_thread[tid]->get_kernel()));
+                            if(m_gpu->get_config().is_trace_driven_mode()) {
+                            	register_cta_thread_exit(cta_id, m_kernel);
+                            }
+                            else
+                            	register_cta_thread_exit(cta_id, &(m_thread[tid]->get_kernel()));
                             m_not_completed -= 1;
                             m_active_threads.reset(tid);
-                            assert( m_thread[tid]!= NULL );
+                            if(!m_gpu->get_config().is_trace_driven_mode()) assert( m_thread[tid]!= NULL );
                             did_exit=true;
                         }
                     }
@@ -812,9 +843,15 @@ void shader_core_ctx::fetch()
 
                 // this code fetches instructions from the i-cache or generates memory requests
                 if( !m_warp[warp_id].functional_done() && !m_warp[warp_id].imiss_pending() && m_warp[warp_id].ibuffer_empty() ) {
-                    address_type pc  = m_warp[warp_id].get_pc();
+                	address_type pc;
+                	if(m_gpu->get_config().is_trace_driven_mode()){
+                		trace_shader_core_ctx* trace_core = static_cast<trace_shader_core_ctx*> (this);
+                		pc = trace_core->m_trace_warp[warp_id].get_pc();
+                	}
+                	else
+                		pc  = m_warp[warp_id].get_pc();
                     address_type ppc = pc + PROGRAM_MEM_START;
-                    unsigned nbytes=16;
+                    unsigned nbytes= m_gpu->get_config().is_trace_driven_mode()? 32 : 16;
                     unsigned offset_in_block = pc & (m_config->m_L1I_config.get_line_sz()-1);
                     if( (offset_in_block+nbytes) > m_config->m_L1I_config.get_line_sz() )
                         nbytes = (m_config->m_L1I_config.get_line_sz()-offset_in_block);
@@ -869,25 +906,29 @@ void shader_core_ctx::func_exec_inst( warp_inst_t &inst )
 void shader_core_ctx::issue_warp( register_set& pipe_reg_set, const warp_inst_t* next_inst, const active_mask_t &active_mask, unsigned warp_id, unsigned sch_id )
 {
 	warp_inst_t** pipe_reg = pipe_reg_set.get_free(m_config->sub_core_model, sch_id);
-    assert(pipe_reg);
-
-    m_warp[warp_id].ibuffer_free();
-    assert(next_inst->valid());
-    **pipe_reg = *next_inst; // static instruction information
-    (*pipe_reg)->issue( active_mask, warp_id, m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle, m_warp[warp_id].get_dynamic_warp_id(), sch_id ); // dynamic instruction information
-    m_stats->shader_cycle_distro[2+(*pipe_reg)->active_count()]++;
-    func_exec_inst( **pipe_reg );
-    if( next_inst->op == BARRIER_OP ){
-        m_warp[warp_id].store_info_of_last_inst_at_barrier(*pipe_reg);
-        m_barriers.warp_reaches_barrier(m_warp[warp_id].get_cta_id(),warp_id,const_cast<warp_inst_t*> (next_inst));
-
-    }else if( next_inst->op == MEMORY_BARRIER_OP ){
-        m_warp[warp_id].set_membar();
-    }
+	assert(pipe_reg);
+
+	m_warp[warp_id].ibuffer_free();
+	assert(next_inst->valid());
+	**pipe_reg = *next_inst; // static instruction information
+	(*pipe_reg)->issue( active_mask, warp_id, m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle, m_warp[warp_id].get_dynamic_warp_id(), sch_id ); // dynamic instruction information
+	m_stats->shader_cycle_distro[2+(*pipe_reg)->active_count()]++;
+	func_exec_inst( **pipe_reg );
+
+	if( next_inst->op == BARRIER_OP ){
+		m_warp[warp_id].store_info_of_last_inst_at_barrier(*pipe_reg);
+		m_barriers.warp_reaches_barrier(m_warp[warp_id].get_cta_id(),warp_id,const_cast<warp_inst_t*> (next_inst));
+
+	}else if( next_inst->op == MEMORY_BARRIER_OP ){
+		m_warp[warp_id].set_membar();
+	}
+
+	if(!m_gpu->get_config().is_trace_driven_mode()) //No SIMT-stack in trace-driven  mode
+		updateSIMTStack(warp_id,*pipe_reg);
+
+	m_scoreboard->reserveRegisters(*pipe_reg);
+	m_warp[warp_id].set_next_pc(next_inst->pc + next_inst->isize);
 
-    updateSIMTStack(warp_id,*pipe_reg);
-    m_scoreboard->reserveRegisters(*pipe_reg);
-    m_warp[warp_id].set_next_pc(next_inst->pc + next_inst->isize);
 }
 
 void shader_core_ctx::issue(){
@@ -1034,7 +1075,11 @@ void scheduler_unit::cycle()
             bool valid = warp(warp_id).ibuffer_next_valid();
             bool warp_inst_issued = false;
             unsigned pc,rpc;
-            m_simt_stack[warp_id]->get_pdom_stack_top_info(&pc,&rpc);
+            if(m_shader->m_gpu->get_config().is_trace_driven_mode())
+            	pc =  pI->pc;  //assume no control hazard in trace mode. TO DO: to be fixed
+            else
+            	m_simt_stack[warp_id]->get_pdom_stack_top_info(&pc,&rpc);
+
             SCHED_DPRINTF( "Warp (warp_id %u, dynamic_warp_id %u) has valid instruction (%s)\n",
                            (*iter)->get_warp_id(), (*iter)->get_dynamic_warp_id(),
                            m_shader->m_config->gpgpu_ctx->func_sim->ptx_get_insn_str( pc).c_str() );
@@ -1052,7 +1097,11 @@ void scheduler_unit::cycle()
                         SCHED_DPRINTF( "Warp (warp_id %u, dynamic_warp_id %u) passes scoreboard\n",
                                        (*iter)->get_warp_id(), (*iter)->get_dynamic_warp_id() );
                         ready_inst = true;
-                        const active_mask_t &active_mask = m_simt_stack[warp_id]->get_active_mask();
+
+                        //For Trace-driven, the active mask already set in from traces, so just read it from the inst
+                        const active_mask_t &active_mask  = m_shader->m_gpu->get_config().is_trace_driven_mode()?
+                        		pI->get_active_mask() :  m_simt_stack[warp_id]->get_active_mask();
+
                         assert( warp(warp_id).inst_in_pipeline() );
 
                         if ( (pI->op == LOAD_OP) || (pI->op == STORE_OP) || (pI->op == MEMORY_BARRIER_OP)||(pI->op==TENSOR_CORE_LOAD_OP)||(pI->op==TENSOR_CORE_STORE_OP) ) {
@@ -3241,6 +3290,10 @@ void barrier_set_t::deallocate_barrier( unsigned cta_id )
    warp_set_t at_barrier = warps & m_warp_at_barrier;
    assert( at_barrier.any() == false ); // no warps stuck at barrier
    warp_set_t active = warps & m_warp_active;
+   std::cout<<active<<std::endl;
+   std::cout<<warps<<std::endl;
+   std::cout<<m_warp_active<<std::endl;
+
    assert( active.any() == false ); // no warps in CTA still running
    m_warp_active &= ~warps;
    m_warp_at_barrier &= ~warps;
@@ -3839,7 +3892,10 @@ simt_core_cluster::simt_core_cluster( class gpgpu_sim *gpu,
     m_core = new shader_core_ctx*[ config->n_simt_cores_per_cluster ];
     for( unsigned i=0; i < config->n_simt_cores_per_cluster; i++ ) {
         unsigned sid = m_config->cid_to_sid(i,m_cluster_id);
-        m_core[i] = new shader_core_ctx(gpu,this,sid,m_cluster_id,config,mem_config,stats);
+        if(gpu->get_config().is_trace_driven_mode())
+        	m_core[i] = new trace_shader_core_ctx(gpu,this,sid,m_cluster_id,config,mem_config,stats);
+        else
+        	m_core[i] = new shader_core_ctx(gpu,this,sid,m_cluster_id,config,mem_config,stats);
         m_core_sim_order.push_back(i); 
     }
 }
@@ -4158,3 +4214,5 @@ void shader_core_ctx::checkExecutionStatusAndUpdate(warp_inst_t &inst, unsigned
     }
 }
 
+
+
diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h
index 667cb2d..ae88b7d 100644
--- a/src/gpgpu-sim/shader.h
+++ b/src/gpgpu-sim/shader.h
@@ -1947,7 +1947,7 @@ public:
 	 }
 
     int test_res_bus(int latency);
-    void init_warps(unsigned cta_id, unsigned start_thread, unsigned end_thread,unsigned ctaid, int cta_size, unsigned kernel_id);
+    void init_warps(unsigned cta_id, unsigned start_thread, unsigned end_thread,unsigned ctaid, int cta_size, kernel_info_t &kernel);
     virtual void checkExecutionStatusAndUpdate(warp_inst_t &inst, unsigned t, unsigned tid);
     address_type next_pc( int tid ) const;
     void fetch();
@@ -1960,7 +1960,7 @@ public:
     friend class TwoLevelScheduler;
     friend class LooseRoundRobbinScheduler;
     void issue_warp( register_set& warp, const warp_inst_t *pI, const active_mask_t &active_mask, unsigned warp_id, unsigned sch_id );
-    void func_exec_inst( warp_inst_t &inst );
+    virtual void func_exec_inst( warp_inst_t &inst );
 
      // Returns numbers of addresses in translated_addrs
     unsigned translate_local_memaddr( address_type localaddr, unsigned tid, unsigned num_shader, unsigned datasize, new_addr_type* translated_addrs );
@@ -2052,7 +2052,8 @@ private:
     std::bitset<MAX_THREAD_PER_SM> m_occupied_hwtid;
     std::map<unsigned int, unsigned int> m_occupied_cta_to_hwtid; 
 
-
+    //Trace-driven simulation
+    friend class trace_shader_core_ctx;
 };
 
 class simt_core_cluster {
diff --git a/src/trace-driven/gpgpusim_trace_driven_main.cc b/src/trace-driven/gpgpusim_trace_driven_main.cc
index 2560338..fea56c1 100644
--- a/src/trace-driven/gpgpusim_trace_driven_main.cc
+++ b/src/trace-driven/gpgpusim_trace_driven_main.cc
@@ -8,30 +8,32 @@
 #include <fstream>
 #include <string>
 #include <sstream>
-#include <stdio.h>
 #include <math.h>
 
-//#include "../abstract_hardware_model.h"
-//#include "../option_parser.h"
-//#include "../cuda-sim/cuda-sim.h"
-//#include "../cuda-sim/ptx_ir.h"
-//#include "../cuda-sim/ptx_parser.h"
+#include "../abstract_hardware_model.h"
+#include "../option_parser.h"
+#include "../cuda-sim/cuda-sim.h"
+#include "../cuda-sim/ptx_ir.h"
+#include "../cuda-sim/ptx_parser.h"
 #include "../gpgpu-sim/gpu-sim.h"
 //#include "../gpgpu-sim/icnt_wrapper.h"
 //#include "../gpgpu-sim/icnt_wrapper.h"
 #include "../../libcuda/gpgpu_context.h"
 #include "trace_driven.h"
 #include "trace_opcode.h"
+#include "../gpgpusim_entrypoint.h"
+//#include "gpgpu_context.h"
 
 //#include "../stream_manager.h"
 
 
 void arguments_check();
 
+
 int main ( int argc, const char **argv )
 {
 
-	gpgpu_context* m_gpgpu_context = GPGPU_Context();
+	gpgpu_context* m_gpgpu_context = new gpgpu_context();
 	gpgpu_sim * m_gpgpu_sim = m_gpgpu_context->gpgpu_trace_sim_init_perf(argc,argv);
 	m_gpgpu_sim->init();
 
@@ -187,6 +189,7 @@ trace_kernel_info_t* trace_parser::parse_kernel_info(const std::string& kerneltr
 			else if (string1 == "cuda" && string2 == "stream") {
 				sscanf(line.c_str(), "-cuda stream id = %d", &cuda_stream_id);
 			}
+			std::cout << line << std::endl;
 			continue;
 		}
 	}
@@ -212,22 +215,49 @@ void trace_parser::kernel_finalizer(trace_kernel_info_t* kernel_info){
 	delete kernel_info;
 }
 
-bool trace_kernel_info_t::get_next_threadblock_traces(std::vector<std::vector<trace_warp_inst_t>>& threadblock_traces) {
+const trace_warp_inst_t* trace_shd_warp_t::get_next_inst(){
+	return &warp_traces[trace_pc++];
+}
+
+void trace_shd_warp_t::clear() {
+	trace_pc=0;
+	warp_traces.clear();
+}
+
+bool trace_shd_warp_t::trace_done() {
+	return trace_pc==warp_traces.size();
+}
 
-	threadblock_traces.clear();
-	unsigned warps_per_tb = ceil(float(threads_per_cta()/32));
-	threadblock_traces.resize(warps_per_tb);
+address_type trace_shd_warp_t::get_start_pc(){
+	assert(warp_traces.size() > 0);
+	return warp_traces[0].pc;
+}
+
+address_type trace_shd_warp_t::get_pc(){
+	assert(warp_traces.size() > 0);
+	return warp_traces[trace_pc].pc;
+}
+
+bool trace_kernel_info_t::get_next_threadblock_traces(std::vector<std::vector<trace_warp_inst_t>*> threadblock_traces) {
+
+	for(unsigned i=0; i<threadblock_traces.size(); ++i) {
+		threadblock_traces[i]->clear();
+	}
+	//unsigned warps_per_tb = ceil(float(threads_per_cta()/32));
+	//threadblock_traces.resize(warps_per_tb);
 
 	unsigned block_id_x=0, block_id_y=0, block_id_z=0;
 	unsigned warp_id=0;
 	unsigned insts_num=0;
-	std::string line;
-	std::stringstream ss;
-	std::string string1, string2;
+
 
 	bool start_of_tb_stream_found = false;
 
 	while(!ifs->eof()) {
+		std::string line;
+		std::stringstream ss;
+		std::string string1, string2;
+
 		getline(*ifs, line);
 
 		if (line.length() == 0) {
@@ -238,32 +268,41 @@ bool trace_kernel_info_t::get_next_threadblock_traces(std::vector<std::vector<tr
 			ss>>string1>>string2;
 			if (string1 == "#BEGIN_TB") {
 				if(!start_of_tb_stream_found)
+				{
 					start_of_tb_stream_found=true;
-				else assert(0 && "Parsing error: thread block start before the previous one finish");
+				}
+				else
+					assert(0 && "Parsing error: thread block start before the previous one finish");
+				std::cout<<line<<std::endl;
 			}
 			else if (string1 == "#END_TB") {
 				assert(start_of_tb_stream_found);
+				std::cout<<line<< std::endl;
 				break; //end of TB stream
 			}
 			else if(string1 == "thread" && string2 == "block") {
 				assert(start_of_tb_stream_found);
 				sscanf(line.c_str(), "thread block = %d,%d,%d", &block_id_x, &block_id_y, &block_id_z);
+				std::cout << line << std::endl;
 			}
 			else if (string1 == "warp") {
 				//the start of new warp stream
 				assert(start_of_tb_stream_found);
 				sscanf(line.c_str(), "warp = %d", &warp_id);
+				//std::cout << line << std::endl;
 			}
 			else if (string1 == "insts") {
 				assert(start_of_tb_stream_found);
 				sscanf(line.c_str(), "insts = %d", &insts_num);
-				threadblock_traces[warp_id].resize(insts_num);
+				threadblock_traces[warp_id]->reserve(insts_num);
+				//std::cout << line << std::endl;
 			}
 			else {
 				assert(start_of_tb_stream_found);
 				trace_warp_inst_t inst(m_gpgpu_sim->getShaderCoreConfig(), m_gpgpu_context);
+				//std::cout<<line << std::endl;
 				inst.parse_from_string(line);
-				threadblock_traces[warp_id].push_back(inst);
+				threadblock_traces[warp_id]->push_back(inst);
 			}
 		}
 	}
@@ -282,8 +321,9 @@ bool trace_warp_inst_t::parse_from_string(std::string trace){
 	unsigned threadblock_x=0, threadblock_y=0, threadblock_z=0, warpid_tb=0, sm_id=0, warpid_sm=0;
 	unsigned long long m_pc=0;
 	unsigned mask=0;
-	unsigned reg_dest=0;
+	unsigned reg_dest[4];
 	std::string opcode;
+	unsigned reg_dsts_num=0;
 	unsigned reg_srcs_num=0;
 	unsigned reg_srcs[4];
 	unsigned mem_width=0;
@@ -292,20 +332,30 @@ bool trace_warp_inst_t::parse_from_string(std::string trace){
 	ss>>std::dec>>threadblock_x>>threadblock_y>>threadblock_z>>warpid_tb>>sm_id>>warpid_sm;
 
 	ss>>std::hex>>m_pc>>mask;
+	//std::cout<<"m_pc= "<<m_pc<<std::endl;
+	//std::cout<<"mask= "<<mask<<std::endl;
+
 	std::bitset<MAX_WARP_SIZE> mask_bits(mask);
 
-	ss>>std::dec>>temp;
-	sscanf(temp.c_str(), "R%d", &reg_dest);
+	ss>>reg_dsts_num;
+
+	for(unsigned i=0; i<reg_dsts_num; ++i) {
+		ss>>std::dec>>temp;
+		sscanf(temp.c_str(), "R%d", &reg_dest[i]);
+	}
 
 	ss>>opcode;
+
 	ss>>reg_srcs_num;
 
 	for(unsigned i=0; i<reg_srcs_num; ++i) {
 		ss>>temp;
 		sscanf(temp.c_str(), "R%d", &reg_srcs[i]);
+
 	}
 
 	ss>>mem_width;
+
 	if(mem_width > 0)  //then it is a memory inst
 	{
 		for (int s = 0; s < warp_size(); s++) {
@@ -327,7 +377,9 @@ bool trace_warp_inst_t::parse_from_string(std::string trace){
 
 	//fill and initialize common params
 	m_decoded = true;
-	pc = m_pc;
+	pc = (address_type)m_pc;   //we will lose the high 32 bits from casting long to unsigned, it should be okay!
+	//std::cout<<"pc= "<<pc<<std::endl;
+
 	isize = 16;   //TO DO, change this
 	for(unsigned i=0; i<MAX_OUTPUT_VALUES; i++) {
 		out[i] = 0;
@@ -345,21 +397,24 @@ bool trace_warp_inst_t::parse_from_string(std::string trace){
 	op = ALU_OP;
 	mem_op= NOT_TEX;
 
-	//get opcode and category
-	std::unordered_map<const char*,OpcodeChar>::const_iterator it= OpcodeMap.find(opcode1.c_str());
+	std::unordered_map<std::string,OpcodeChar>::const_iterator it= OpcodeMap.find(opcode1);
 	if (it != OpcodeMap.end()) {
 		m_opcode = it->second.opcode;
 		op = (op_type)(it->second.opcode_category);
 	}
-	else
+	else {
+		std::cout<<"ERROR:  undefined instruction : "<<opcode<<" Opcode: "<<opcode1<<std::endl;
 		assert(0 && "undefined instruction");
+	}
 
 	//fill regs information
-	num_regs = reg_srcs_num+1;
+	num_regs = reg_srcs_num+reg_dsts_num;
 	num_operands = num_regs;
-	outcount=1;
-	out[0]=reg_dest;
-	arch_reg.dst[0]=reg_dest;
+	outcount=reg_dsts_num;
+	for(unsigned m=0; m<reg_dsts_num; ++m){
+		out[m]=reg_dest[m];
+		arch_reg.src[m]=reg_dest[m];
+	}
 
 	incount=reg_srcs_num;
 	for(unsigned m=0; m<reg_srcs_num; ++m){
@@ -379,7 +434,8 @@ bool trace_warp_inst_t::parse_from_string(std::string trace){
 	}
 
 	// barrier_type bar_type;
-    // reduction_type red_type;
+	// reduction_type red_type;
+
 
 	//fill memory space
 	switch(m_opcode){
@@ -389,6 +445,7 @@ bool trace_warp_inst_t::parse_from_string(std::string trace){
 		assert(mem_width>0);
 		data_size = mem_width;
 		memory_op = memory_load;
+		cache_op = CACHE_ALL;
 		if(m_opcode == OP_LDL)
 			space.set_type(local_space);
 		else
@@ -403,6 +460,7 @@ bool trace_warp_inst_t::parse_from_string(std::string trace){
 		assert(mem_width>0);
 		data_size = mem_width;
 		memory_op = memory_store;
+		cache_op = CACHE_ALL;
 		if(m_opcode == OP_STL)
 			space.set_type(local_space);
 		else
@@ -410,6 +468,14 @@ bool trace_warp_inst_t::parse_from_string(std::string trace){
 
 		if(m_opcode == OP_ATOM || m_opcode == OP_ATOMG || m_opcode == OP_RED)
 			m_isatomic = true;
+
+		for(unsigned m=0; m<reg_dsts_num; ++m){
+			out[m]=0;
+			arch_reg.src[m]=-1;
+		}
+		reg_dsts_num=0;
+		outcount=0;
+
 		break;
 	case OP_LDS:
 	case OP_STS:
@@ -418,9 +484,16 @@ bool trace_warp_inst_t::parse_from_string(std::string trace){
 		data_size = mem_width;
 		space.set_type(shared_space);
 		break;
-
+	case OP_BAR:
+		//TO DO fill this correctly
+		bar_id = 0;
+		bar_count = (unsigned)-1;
+		bar_type = SYNC;
+		//if bar_type = RED;
+		//set bar_type
+		break;
 	default:
-			break;
+		break;
 	}
 
 	return true;
@@ -482,34 +555,98 @@ void trace_warp_inst_t::set_latency(unsigned category)
 	initiation_interval = latency = 1;
 
 	switch(category){
-		case ALU_OP:
-		case INTP_OP:
-		case BRANCH_OP:
-		case CALL_OPS:
-		case RET_OPS:
-			latency = int_latency[0];
-			initiation_interval = int_init[0];
-			break;
-		case SP_OP:
-			latency = fp_latency[0];
-			initiation_interval = fp_latency[0];
-			break;
-		case DP_OP:
-			latency = dp_latency[0];
-			initiation_interval = dp_latency[0];
-			break;
-		case SFU_OP:
-			latency = sfu_latency;
-			initiation_interval = sfu_init;
-			break;
-		case TENSOR_CORE_OP:
-			latency = tensor_latency;
-			initiation_interval = tensor_init;
-			break;
-		default:
-				break;
-		}
+	case ALU_OP:
+	case INTP_OP:
+	case BRANCH_OP:
+	case CALL_OPS:
+	case RET_OPS:
+		latency = int_latency[0];
+		initiation_interval = int_init[0];
+		break;
+	case SP_OP:
+		latency = fp_latency[0];
+		initiation_interval = fp_latency[0];
+		break;
+	case DP_OP:
+		latency = dp_latency[0];
+		initiation_interval = dp_latency[0];
+		break;
+	case SFU_OP:
+		latency = sfu_latency;
+		initiation_interval = sfu_init;
+		break;
+	case TENSOR_CORE_OP:
+		latency = tensor_latency;
+		initiation_interval = tensor_init;
+		break;
+	default:
+		break;
+	}
 
 }
 
+unsigned trace_shader_core_ctx::trace_sim_inc_thread( kernel_info_t &kernel)
+{
+
+	if ( kernel.no_more_ctas_to_run() ) {
+		return 0; //finished!
+	}
+
+	if( kernel.more_threads_in_cta() ) {
+		kernel.increment_thread_id();
+	}
+
+	if( !kernel.more_threads_in_cta() )
+		kernel.increment_cta_id();
+
+	return 1;
+}
+
+void trace_shader_core_ctx::init_traces( unsigned start_warp, unsigned end_warp, kernel_info_t &kernel ) {
+
+	std::vector<std::vector<trace_warp_inst_t>*> threadblock_traces;
+	for (unsigned i = start_warp; i < end_warp; ++i) {
+		threadblock_traces.push_back(&(m_trace_warp[i].warp_traces));
+	}
+	trace_kernel_info_t& trace_kernel = static_cast<trace_kernel_info_t&> (kernel);
+	trace_kernel.get_next_threadblock_traces(threadblock_traces);
+
+	//set pc
+	for (unsigned i = start_warp; i < end_warp; ++i) {
+		m_warp[i].set_next_pc(m_trace_warp[i].get_start_pc());
+	}
+}
+
+
+void trace_shader_core_ctx::checkExecutionStatusAndUpdate(warp_inst_t &inst, unsigned t, unsigned tid)
+{
+	if(inst.isatomic())
+		m_warp[inst.warp_id()].inc_n_atomic();
+
+	if ( m_trace_warp[inst.warp_id()].trace_done() ) {
+		m_warp[inst.warp_id()].set_completed(t);
+		m_warp[inst.warp_id()].ibuffer_flush();
+	}
+
+}
+
+void trace_shader_core_ctx::func_exec_inst( warp_inst_t &inst )
+{
+	//here, we generate memory acessess and set the status if thread (done?)
+	if( inst.is_load() || inst.is_store() )
+	{
+		inst.generate_mem_accesses();
+	}
+	for ( unsigned t=0; t < m_warp_size; t++ ) {
+		if( inst.active(t) ) {
+			unsigned warpId = inst.warp_id();
+			unsigned tid=m_warp_size*warpId+t;
+
+			//virtual function
+			checkExecutionStatusAndUpdate(inst,t,tid);
+		}
+	}
+	if(m_trace_warp[inst.warp_id()].trace_done() )
+		m_barriers.warp_exit( inst.warp_id() );
+}
 
diff --git a/src/trace-driven/trace_driven.h b/src/trace-driven/trace_driven.h
index fbf7a1e..2f97958 100644
--- a/src/trace-driven/trace_driven.h
+++ b/src/trace-driven/trace_driven.h
@@ -24,6 +24,10 @@ public:
 		m_kernel_info = info;
 	}
 
+	virtual ~trace_function_info() {
+
+	}
+
 private:
 
 
@@ -33,21 +37,23 @@ class trace_warp_inst_t: public warp_inst_t {
 public:
 
 	trace_warp_inst_t() {
-
+		m_gpgpu_context=NULL;
+		m_opcode=0;
 	}
+
 	trace_warp_inst_t(const class core_config *config, gpgpu_context* gpgpu_context ):warp_inst_t(config) {
 		m_gpgpu_context = gpgpu_context;
 		m_opcode=0;
 	}
 
 	bool parse_from_string(std::string trace);
-
+	unsigned m_opcode;
 
 private:
 
 	void set_latency(unsigned cat);
 	gpgpu_context* m_gpgpu_context;
-	unsigned m_opcode;
+
 
 };
 
@@ -59,7 +65,7 @@ public:
 		m_gpgpu_context = gpgpu_context;
 	}
 
-	bool get_next_threadblock_traces(std::vector<std::vector<trace_warp_inst_t>>& threadblock_traces);
+	bool get_next_threadblock_traces(std::vector<std::vector<trace_warp_inst_t>*> threadblock_traces);
 
 private:
 	std::ifstream* ifs;
@@ -87,15 +93,48 @@ private:
 
 };
 
-class trace_shd_warp_t: public shd_warp_t {
+class trace_shd_warp_t {
 public:
-	trace_shd_warp_t(class shader_core_ctx *shader, unsigned warp_size):shd_warp_t(shader, warp_size) {
+	trace_shd_warp_t() {
+		trace_pc=0;
 	}
 
-	bool get_next_threadblock_traces(std::vector<std::vector<trace_warp_inst_t>>& threadblock_traces);
+	std::vector<trace_warp_inst_t> warp_traces;
+	const trace_warp_inst_t* get_next_inst();
+	void clear();
+	bool trace_done();
+	address_type get_start_pc();
+	address_type get_pc();
+	unsigned trace_pc;
 
 private:
 
+
+};
+
+class trace_shader_core_ctx: public shader_core_ctx {
+
+public:
+	trace_shader_core_ctx(class gpgpu_sim *gpu,
+            class simt_core_cluster *cluster,
+            unsigned shader_id,
+            unsigned tpc_id,
+            const shader_core_config *config,
+            const memory_config *mem_config,
+            shader_core_stats *stats):shader_core_ctx(gpu, cluster, shader_id, tpc_id, config, mem_config, stats) {
+
+		m_trace_warp.resize(get_config()->max_warps_per_shader);
+	}
+
+	virtual void checkExecutionStatusAndUpdate(warp_inst_t &inst, unsigned t, unsigned tid);
+	void init_traces( unsigned start_warp, unsigned end_warp, kernel_info_t &kernel );
+	unsigned trace_sim_inc_thread( kernel_info_t &kernel);
+	virtual void func_exec_inst( warp_inst_t &inst );
+	std::vector<trace_shd_warp_t> m_trace_warp;
+
+private:
+
+
 };
 
 #endif
diff --git a/src/trace-driven/trace_opcode.h b/src/trace-driven/trace_opcode.h
index f809238..827c32c 100644
--- a/src/trace-driven/trace_opcode.h
+++ b/src/trace-driven/trace_opcode.h
@@ -5,6 +5,7 @@
 
 #include "../abstract_hardware_model.h"
 #include <unordered_map>
+#include <string>
 
 
 enum TraceInstrOpcode {
@@ -60,7 +61,7 @@ struct OpcodeChar
 
 ///Volta SM_70 ISA
 //see: https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html
-static const std::unordered_map<const char*,OpcodeChar> OpcodeMap = {
+static const std::unordered_map<std::string,OpcodeChar> OpcodeMap = {
 		//Floating Point 32 Instructions
 		{"FADD", OpcodeChar(OP_FADD, SP_OP)},
 		{"FADD32I", OpcodeChar(OP_FADD32I, SP_OP)},
-- 
cgit v1.3


From 8571d5fe845eed8fe71f97e72f7988f0e7ea257d Mon Sep 17 00:00:00 2001
From: Mahmoud <abdallm@purdue.edu>
Date: Mon, 23 Sep 2019 11:57:30 -0400
Subject: updating Makefile to add trace-driven

---
 Makefile | 6 +++++-
 1 file changed, 5 insertions(+), 1 deletion(-)

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index 73f22c6..bebe6f2 100644
--- a/Makefile
+++ b/Makefile
@@ -57,7 +57,7 @@ endif
 
 $(shell mkdir -p $(SIM_OBJ_FILES_DIR)/libcuda && echo "const char *g_gpgpusim_build_string=\"$(GPGPUSIM_BUILD)\";" > $(SIM_OBJ_FILES_DIR)/detailed_version)
 
-LIBS = cuda-sim gpgpu-sim_uarch $(INTERSIM) gpgpusimlib 
+LIBS = cuda-sim gpgpu-sim_uarch $(INTERSIM) gpgpusimlib
 
 
 TARGETS =
@@ -219,6 +219,9 @@ cuda-sim: makedirs
 	$(MAKE) -C ./src/cuda-sim/ depend
 	$(MAKE) -C ./src/cuda-sim/
 
+trace-driven: makedirs
+	g++ -std=c++0x -c src/trace-driven/gpgpusim_trace_driven_main.cc 
+
 gpgpu-sim_uarch: makedirs cuda-sim
 	$(MAKE) -C ./src/gpgpu-sim/ depend
 	$(MAKE) -C ./src/gpgpu-sim/
@@ -245,6 +248,7 @@ makedirs:
 	if [ ! -d $(SIM_OBJ_FILES_DIR)/cuda-sim ]; then mkdir -p $(SIM_OBJ_FILES_DIR)/cuda-sim; fi;
 	if [ ! -d $(SIM_OBJ_FILES_DIR)/cuda-sim/decuda_pred_table ]; then mkdir -p $(SIM_OBJ_FILES_DIR)/cuda-sim/decuda_pred_table; fi;
 	if [ ! -d $(SIM_OBJ_FILES_DIR)/gpgpu-sim ]; then mkdir -p $(SIM_OBJ_FILES_DIR)/gpgpu-sim; fi;
+	if [ ! -d $(SIM_OBJ_FILES_DIR)/trace-driven ]; then mkdir -p $(SIM_OBJ_FILES_DIR)/trace-driven; fi;
 	if [ ! -d $(SIM_OBJ_FILES_DIR)/libopencl ]; then mkdir -p $(SIM_OBJ_FILES_DIR)/libopencl; fi;
 	if [ ! -d $(SIM_OBJ_FILES_DIR)/libopencl/bin ]; then mkdir -p $(SIM_OBJ_FILES_DIR)/libopencl/bin; fi;
 	if [ ! -d $(SIM_OBJ_FILES_DIR)/$(INTERSIM) ]; then mkdir -p $(SIM_OBJ_FILES_DIR)/$(INTERSIM); fi;
-- 
cgit v1.3


From a113dfbbd5393146e778acd66d37ac1dd84bf0ec Mon Sep 17 00:00:00 2001
From: Mahmoud <abdallm@purdue.edu>
Date: Mon, 30 Sep 2019 10:52:51 -0400
Subject: fixing the execution-driven linking error

---
 Makefile                                       |   5 +-
 src/trace-driven/gpgpusim_trace_driven_main.cc | 545 -------------------------
 2 files changed, 3 insertions(+), 547 deletions(-)

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index bebe6f2..acef015 100644
--- a/Makefile
+++ b/Makefile
@@ -57,7 +57,7 @@ endif
 
 $(shell mkdir -p $(SIM_OBJ_FILES_DIR)/libcuda && echo "const char *g_gpgpusim_build_string=\"$(GPGPUSIM_BUILD)\";" > $(SIM_OBJ_FILES_DIR)/detailed_version)
 
-LIBS = cuda-sim gpgpu-sim_uarch $(INTERSIM) gpgpusimlib
+LIBS = cuda-sim gpgpu-sim_uarch $(INTERSIM) gpgpusimlib trace-driven
 
 
 TARGETS =
@@ -150,6 +150,7 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/decuda_pred_table/*.o \
 			$(SIM_OBJ_FILES_DIR)/gpgpu-sim/*.o \
 			$(SIM_OBJ_FILES_DIR)/$(INTERSIM)/*.o \
+			$(SIM_OBJ_FILES_DIR)/trace-driven/*.o \
 			$(SIM_OBJ_FILES_DIR)/*.o -lm -lz -lGL -pthread \
 			$(MCPAT) \
 			-o $(SIM_LIB_DIR)/libcudart.so
@@ -220,7 +221,7 @@ cuda-sim: makedirs
 	$(MAKE) -C ./src/cuda-sim/
 
 trace-driven: makedirs
-	g++ -std=c++0x -c src/trace-driven/gpgpusim_trace_driven_main.cc 
+	g++ -fPIC -std=c++0x -I$(CUDA_INSTALL_PATH)/include -c src/trace-driven/trace_driven.cc -o $(SIM_OBJ_FILES_DIR)/trace-driven/trace_driven.o
 
 gpgpu-sim_uarch: makedirs cuda-sim
 	$(MAKE) -C ./src/gpgpu-sim/ depend
diff --git a/src/trace-driven/gpgpusim_trace_driven_main.cc b/src/trace-driven/gpgpusim_trace_driven_main.cc
index 56e493c..642b91a 100644
--- a/src/trace-driven/gpgpusim_trace_driven_main.cc
+++ b/src/trace-driven/gpgpusim_trace_driven_main.cc
@@ -13,8 +13,6 @@
 #include "../abstract_hardware_model.h"
 #include "../option_parser.h"
 #include "../cuda-sim/cuda-sim.h"
-#include "../cuda-sim/ptx_ir.h"
-#include "../cuda-sim/ptx_parser.h"
 #include "../gpgpu-sim/gpu-sim.h"
 #include "../../libcuda/gpgpu_context.h"
 #include "trace_driven.h"
@@ -87,546 +85,3 @@ int main ( int argc, const char **argv )
 
 	return 1;
 }
-
-trace_parser::trace_parser(const char* kernellist_filepath, gpgpu_sim * m_gpgpu_sim, gpgpu_context* m_gpgpu_context)
-{
-
-	this->m_gpgpu_sim = m_gpgpu_sim;
-	this->m_gpgpu_context = m_gpgpu_context;
-	kernellist_filename = kernellist_filepath;
-}
-
-std::vector<std::string> trace_parser::parse_kernellist_file() {
-
-	ifs.open(kernellist_filename);
-
-	if (!ifs.is_open()) {
-		std::cout << "Unable to open file: " <<kernellist_filename<<std::endl;
-		exit(1);
-	}
-
-	std::string directory(kernellist_filename);
-	const size_t last_slash_idx = directory.rfind('/');
-	if (std::string::npos != last_slash_idx)
-	{
-		directory = directory.substr(0, last_slash_idx);
-	}
-
-	std::string line, filepath;
-	std::vector<std::string> kernellist;
-	while(!ifs.eof()) {
-		getline(ifs, line);
-		if(line.empty())
-			continue;
-		filepath = directory+"/"+line;
-		kernellist.push_back(filepath);
-	}
-
-	ifs.close();
-	return kernellist;
-}
-
-
-trace_kernel_info_t* trace_parser::parse_kernel_info(const std::string& kerneltraces_filepath) {
-
-	ifs.open(kerneltraces_filepath.c_str());
-
-	if (!ifs.is_open()) {
-		std::cout << "Unable to open file: " <<kerneltraces_filepath<<std::endl;
-		exit(1);
-	}
-
-	std::cout << "Processing kernel " <<kerneltraces_filepath<<std::endl;
-
-	unsigned grid_dim_x=0, grid_dim_y=0, grid_dim_z=0, tb_dim_x=0, tb_dim_y=0, tb_dim_z=0;
-	unsigned shmem=0, nregs=0, cuda_stream_id=0, kernel_id=0;
-	std::string line;
-	std::stringstream ss;
-	std::string string1, string2;
-	std::string  kernel_name;
-
-	while(!ifs.eof()) {
-		getline(ifs, line);
-
-		if (line.length() == 0) {
-			continue;
-		}
-		else if(line[0] == '#'){
-			//the trace format, ignore this and assume fixed format for now
-			break;  //the begin of the instruction stream
-		}
-		else if(line[0] == '-') {
-			ss.str(line);
-			ss.ignore();
-			ss>>string1>>string2;
-			if(string1 == "kernel" && string2 == "name") {
-				const size_t equal_idx = line.find('=');
-				kernel_name = line.substr(equal_idx+1);
-			}
-			else if(string1 == "kernel" && string2 == "id") {
-				sscanf(line.c_str(), "-kernel id = %d", &kernel_id);
-			}
-			else if(string1 == "grid" && string2 == "dim") {
-				sscanf(line.c_str(), "-grid dim = (%d,%d,%d)", &grid_dim_x, &grid_dim_y, &grid_dim_z);
-			}
-			else if (string1 == "block" && string2 == "dim") {
-				sscanf(line.c_str(), "-block dim = (%d,%d,%d)", &tb_dim_x, &tb_dim_y, &tb_dim_z);
-			}
-			else if (string1 == "shmem") {
-				sscanf(line.c_str(), "-shmem = %d", &shmem);
-			}
-			else if (string1 == "nregs") {
-				sscanf(line.c_str(), "-nregs = %d", &nregs);
-			}
-			else if (string1 == "cuda" && string2 == "stream") {
-				sscanf(line.c_str(), "-cuda stream id = %d", &cuda_stream_id);
-			}
-			std::cout << line << std::endl;
-			continue;
-		}
-	}
-
-	gpgpu_ptx_sim_info info;
-	info.smem = shmem;
-	info.regs = nregs;
-	dim3 gridDim(grid_dim_x, grid_dim_y, grid_dim_z);
-	dim3 blockDim(tb_dim_x, tb_dim_y, tb_dim_z);
-	trace_function_info* function_info = new trace_function_info(info, m_gpgpu_context);
-	function_info->set_name(kernel_name.c_str());
-	trace_kernel_info_t* kernel_info =  new trace_kernel_info_t(gridDim, blockDim, function_info, &ifs, m_gpgpu_sim, m_gpgpu_context);
-
-	return kernel_info;
-}
-
-
-void trace_parser::kernel_finalizer(trace_kernel_info_t* kernel_info){
-	if (ifs.is_open())
-		ifs.close();
-
-	delete kernel_info->entry();
-	delete kernel_info;
-}
-
-const trace_warp_inst_t* trace_shd_warp_t::get_next_inst(){
-	if(trace_pc < warp_traces.size())
-		return &warp_traces[trace_pc++];
-	else
-		return NULL;
-}
-
-void trace_shd_warp_t::clear() {
-	trace_pc=0;
-	warp_traces.clear();
-}
-
-bool trace_shd_warp_t::trace_done() {
-	return trace_pc==(warp_traces.size());
-}
-
-address_type trace_shd_warp_t::get_start_pc(){
-	assert(warp_traces.size() > 0);
-	return warp_traces[0].pc;
-}
-
-address_type trace_shd_warp_t::get_pc(){
-	assert(warp_traces.size() > 0 );
-	assert(trace_pc < warp_traces.size());
-	return warp_traces[trace_pc].pc;
-}
-
-bool trace_kernel_info_t::get_next_threadblock_traces(std::vector<std::vector<trace_warp_inst_t>*> threadblock_traces) {
-
-	for(unsigned i=0; i<threadblock_traces.size(); ++i) {
-		threadblock_traces[i]->clear();
-	}
-
-	unsigned block_id_x=0, block_id_y=0, block_id_z=0;
-	unsigned warp_id=0;
-	unsigned insts_num=0;
-
-
-	bool start_of_tb_stream_found = false;
-
-	while(!ifs->eof()) {
-		std::string line;
-		std::stringstream ss;
-		std::string string1, string2;
-
-		getline(*ifs, line);
-
-		if (line.length() == 0) {
-			continue;
-		}
-		else {
-			ss.str(line);
-			ss>>string1>>string2;
-			if (string1 == "#BEGIN_TB") {
-				if(!start_of_tb_stream_found)
-				{
-					start_of_tb_stream_found=true;
-				}
-				else
-					assert(0 && "Parsing error: thread block start before the previous one finish");
-			}
-			else if (string1 == "#END_TB") {
-				assert(start_of_tb_stream_found);
-				break; //end of TB stream
-			}
-			else if(string1 == "thread" && string2 == "block") {
-				assert(start_of_tb_stream_found);
-				sscanf(line.c_str(), "thread block = %d,%d,%d", &block_id_x, &block_id_y, &block_id_z);
-				std::cout << line << std::endl;
-			}
-			else if (string1 == "warp") {
-				//the start of new warp stream
-				assert(start_of_tb_stream_found);
-				sscanf(line.c_str(), "warp = %d", &warp_id);
-			}
-			else if (string1 == "insts") {
-				assert(start_of_tb_stream_found);
-				sscanf(line.c_str(), "insts = %d", &insts_num);
-				threadblock_traces[warp_id]->reserve(insts_num);
-			}
-			else {
-				assert(start_of_tb_stream_found);
-				trace_warp_inst_t inst(m_gpgpu_sim->getShaderCoreConfig(), m_gpgpu_context);
-				inst.parse_from_string(line);
-				threadblock_traces[warp_id]->push_back(inst);
-			}
-		}
-	}
-
-	return true;
-}
-
-
-bool trace_warp_inst_t::parse_from_string(std::string trace){
-
-	std::stringstream ss;
-	ss.str(trace);
-
-
-	std::string temp;
-	unsigned threadblock_x=0, threadblock_y=0, threadblock_z=0, warpid_tb=0, sm_id=0, warpid_sm=0;
-	unsigned long long m_pc=0;
-	unsigned mask=0;
-	unsigned reg_dest[4];
-	std::string opcode;
-	unsigned reg_dsts_num=0;
-	unsigned reg_srcs_num=0;
-	unsigned reg_srcs[4];
-	unsigned mem_width=0;
-	unsigned long long mem_addresses[warp_size()];
-
-	//Start Parsing
-	ss>>std::dec>>threadblock_x>>threadblock_y>>threadblock_z>>warpid_tb>>sm_id>>warpid_sm;
-
-	ss>>std::hex>>m_pc;
-	ss>>std::hex>>mask;
-
-	std::bitset<MAX_WARP_SIZE> mask_bits(mask);
-
-	ss>>std::dec>>reg_dsts_num;
-	for(unsigned i=0; i<reg_dsts_num; ++i) {
-		ss>>std::dec>>temp;
-		sscanf(temp.c_str(), "R%d", &reg_dest[i]);
-	}
-
-	ss>>opcode;
-
-	ss>>reg_srcs_num;
-	for(unsigned i=0; i<reg_srcs_num; ++i) {
-		ss>>temp;
-		sscanf(temp.c_str(), "R%d", &reg_srcs[i]);
-
-	}
-
-	ss>>mem_width;
-
-	if(mem_width > 0)  //then it is a memory inst
-	{
-		for (int s = 0; s < warp_size(); s++) {
-			if(mask_bits.test(s))
-				ss>>std::hex>>mem_addresses[s];
-			else
-				mem_addresses[s]=0;
-		}
-	}
-	//Finish Parsing
-	//After parsing, fill the inst_t and warp_inst_t params
-
-	//fill active mask
-	active_mask_t active_mask = mask_bits;
-	set_active( active_mask );
-
-	//get the opcode
-	std::string opcode1 = opcode.substr(0, opcode.find("."));
-
-	//fill and initialize common params
-	m_decoded = true;
-	pc = (address_type)m_pc;   //we will lose the high 32 bits from casting long to unsigned, it should be okay!
-
-	isize = 16;   //TO DO, change this
-	for(unsigned i=0; i<MAX_OUTPUT_VALUES; i++) {
-		out[i] = 0;
-	}
-	for(unsigned i=0; i<MAX_INPUT_VALUES; i++) {
-		in[i] = 0;
-	}
-
-	is_vectorin = 0;
-	is_vectorout = 0;
-	ar1 = 0;
-	ar2 = 0;
-	memory_op = no_memory_op;
-	data_size = 0;
-	op = ALU_OP;
-	mem_op= NOT_TEX;
-
-	std::unordered_map<std::string,OpcodeChar>::const_iterator it= OpcodeMap.find(opcode1);
-	if (it != OpcodeMap.end()) {
-		m_opcode = it->second.opcode;
-		op = (op_type)(it->second.opcode_category);
-	}
-	else {
-		std::cout<<"ERROR:  undefined instruction : "<<opcode<<" Opcode: "<<opcode1<<std::endl;
-		assert(0 && "undefined instruction");
-	}
-
-	//fill regs information
-	num_regs = reg_srcs_num+reg_dsts_num;
-	num_operands = num_regs;
-	outcount=reg_dsts_num;
-	for(unsigned m=0; m<reg_dsts_num; ++m){
-		out[m]=reg_dest[m]+1;         //Increment by one because GPGPU-sim starts from R1, while SASS starts from R0
-		arch_reg.src[m]=reg_dest[m]+1;
-	}
-
-	incount=reg_srcs_num;
-	for(unsigned m=0; m<reg_srcs_num; ++m){
-		in[m]=reg_srcs[m]+1;	     //Increment by one because GPGPU-sim starts from R1, while SASS starts from R0
-		arch_reg.src[m]=reg_srcs[m]+1;
-	}
-	//handle: vector, store insts have no output, double inst and hmma, and 64 bit address
-
-
-	//fill latency and initl
-	set_latency(op);
-
-	//fill addresses
-	if(mem_width > 0) {
-		for(unsigned i=0; i<warp_size(); ++i)
-			set_addr(i, mem_addresses[i]);
-	}
-
-
-	//fill memory space
-	switch(m_opcode){
-	case OP_LD:
-	case OP_LDG:
-	case OP_LDL:
-		assert(mem_width>0);
-		data_size = mem_width;
-		memory_op = memory_load;
-		cache_op = CACHE_ALL;
-		if(m_opcode == OP_LDL)
-			space.set_type(local_space);
-		else
-			space.set_type(global_space);
-		break;
-	case OP_ST:
-	case OP_STG:
-	case OP_STL:
-	case OP_ATOM:
-	case OP_ATOMG:
-	case OP_RED:
-		assert(mem_width>0);
-		data_size = mem_width;
-		memory_op = memory_store;
-		cache_op = CACHE_ALL;
-		if(m_opcode == OP_STL)
-			space.set_type(local_space);
-		else
-			space.set_type(global_space);
-
-		if(m_opcode == OP_ATOM || m_opcode == OP_ATOMG || m_opcode == OP_RED)
-			m_isatomic = true;
-
-		break;
-	case OP_LDS:
-	case OP_STS:
-	case OP_ATOMS:
-		assert(mem_width>0);
-		data_size = mem_width;
-		space.set_type(shared_space);
-		break;
-	case OP_BAR:
-		//TO DO fill this correctly
-		bar_id = 0;
-		bar_count = (unsigned)-1;
-		bar_type = SYNC;
-		//TO DO
-		//if bar_type = RED;
-		//set bar_type
-		// barrier_type bar_type;
-		// reduction_type red_type;
-		break;
-	default:
-		break;
-	}
-
-	return true;
-}
-
-void trace_warp_inst_t::set_latency(unsigned category)
-{
-	unsigned int_latency[5];
-	unsigned fp_latency[5];
-	unsigned dp_latency[5];
-	unsigned sfu_latency;
-	unsigned tensor_latency;
-	unsigned int_init[5];
-	unsigned fp_init[5];
-	unsigned dp_init[5];
-	unsigned sfu_init;
-	unsigned tensor_init;
-
-	/*
-	 * [0] ADD,SUB
-	 * [1] MAX,Min
-	 * [2] MUL
-	 * [3] MAD
-	 * [4] DIV
-	 */
-	sscanf(m_gpgpu_context->func_sim->opcode_latency_int, "%u,%u,%u,%u,%u",
-			&int_latency[0],&int_latency[1],&int_latency[2],
-			&int_latency[3],&int_latency[4]);
-	sscanf(m_gpgpu_context->func_sim->opcode_latency_fp, "%u,%u,%u,%u,%u",
-			&fp_latency[0],&fp_latency[1],&fp_latency[2],
-			&fp_latency[3],&fp_latency[4]);
-	sscanf(m_gpgpu_context->func_sim->opcode_latency_dp, "%u,%u,%u,%u,%u",
-			&dp_latency[0],&dp_latency[1],&dp_latency[2],
-			&dp_latency[3],&dp_latency[4]);
-	sscanf(m_gpgpu_context->func_sim->opcode_latency_sfu, "%u",
-			&sfu_latency);
-	sscanf(m_gpgpu_context->func_sim->opcode_latency_tensor, "%u",
-			&tensor_latency);
-	sscanf(m_gpgpu_context->func_sim->opcode_initiation_int, "%u,%u,%u,%u,%u",
-			&int_init[0],&int_init[1],&int_init[2],
-			&int_init[3],&int_init[4]);
-	sscanf(m_gpgpu_context->func_sim->opcode_initiation_fp, "%u,%u,%u,%u,%u",
-			&fp_init[0],&fp_init[1],&fp_init[2],
-			&fp_init[3],&fp_init[4]);
-	sscanf(m_gpgpu_context->func_sim->opcode_initiation_dp, "%u,%u,%u,%u,%u",
-			&dp_init[0],&dp_init[1],&dp_init[2],
-			&dp_init[3],&dp_init[4]);
-	sscanf(m_gpgpu_context->func_sim->opcode_initiation_sfu, "%u",
-			&sfu_init);
-	sscanf(m_gpgpu_context->func_sim->opcode_initiation_tensor, "%u",
-			&tensor_init);
-	sscanf(m_gpgpu_context->func_sim->cdp_latency_str, "%u,%u,%u,%u,%u",
-			&m_gpgpu_context->func_sim->cdp_latency[0],
-			&m_gpgpu_context->func_sim->cdp_latency[1],
-			&m_gpgpu_context->func_sim->cdp_latency[2],
-			&m_gpgpu_context->func_sim->cdp_latency[3],
-			&m_gpgpu_context->func_sim->cdp_latency[4]);
-
-	initiation_interval = latency = 1;
-
-	switch(category){
-	case ALU_OP:
-	case INTP_OP:
-	case BRANCH_OP:
-	case CALL_OPS:
-	case RET_OPS:
-		latency = int_latency[0];
-		initiation_interval = int_init[0];
-		break;
-	case SP_OP:
-		latency = fp_latency[0];
-		initiation_interval = fp_latency[0];
-		break;
-	case DP_OP:
-		latency = dp_latency[0];
-		initiation_interval = dp_latency[0];
-		break;
-	case SFU_OP:
-		latency = sfu_latency;
-		initiation_interval = sfu_init;
-		break;
-	case TENSOR_CORE_OP:
-		latency = tensor_latency;
-		initiation_interval = tensor_init;
-		break;
-	default:
-		break;
-	}
-
-}
-
-unsigned trace_shader_core_ctx::trace_sim_inc_thread( kernel_info_t &kernel)
-{
-
-	if ( kernel.no_more_ctas_to_run() ) {
-		return 0; //finished!
-	}
-
-	if( kernel.more_threads_in_cta() ) {
-		kernel.increment_thread_id();
-	}
-
-	if( !kernel.more_threads_in_cta() )
-		kernel.increment_cta_id();
-
-	return 1;
-}
-
-void trace_shader_core_ctx::init_traces( unsigned start_warp, unsigned end_warp, kernel_info_t &kernel ) {
-
-	std::vector<std::vector<trace_warp_inst_t>*> threadblock_traces;
-	for (unsigned i = start_warp; i < end_warp; ++i) {
-		m_trace_warp[i].clear();
-		threadblock_traces.push_back(&(m_trace_warp[i].warp_traces));
-	}
-	trace_kernel_info_t& trace_kernel = static_cast<trace_kernel_info_t&> (kernel);
-	trace_kernel.get_next_threadblock_traces(threadblock_traces);
-
-	//set pc
-	for (unsigned i = start_warp; i < end_warp; ++i) {
-		m_warp[i].set_next_pc(m_trace_warp[i].get_start_pc());
-	}
-}
-
-
-void trace_shader_core_ctx::checkExecutionStatusAndUpdate(warp_inst_t &inst, unsigned t, unsigned tid)
-{
-	if(inst.isatomic())
-		m_warp[inst.warp_id()].inc_n_atomic();
-
-	if ( m_trace_warp[inst.warp_id()].trace_done() ) {
-		m_warp[inst.warp_id()].set_completed(t);
-		m_warp[inst.warp_id()].ibuffer_flush();
-	}
-
-}
-
-void trace_shader_core_ctx::func_exec_inst( warp_inst_t &inst )
-{
-	//here, we generate memory acessess and set the status if thread (done?)
-	if( inst.is_load() || inst.is_store() )
-	{
-		inst.generate_mem_accesses();
-	}
-	for ( unsigned t=0; t < m_warp_size; t++ ) {
-		if( inst.active(t) ) {
-			unsigned warpId = inst.warp_id();
-			unsigned tid=m_warp_size*warpId+t;
-
-			//virtual function
-			checkExecutionStatusAndUpdate(inst,t,tid);
-		}
-	}
-	if(m_trace_warp[inst.warp_id()].trace_done() )
-		m_barriers.warp_exit( inst.warp_id() );
-}
-
-- 
cgit v1.3


From a0aa4823f46b7040f0458d0d84d085b47ece8c8d Mon Sep 17 00:00:00 2001
From: Mahmoud <abdallm@purdue.edu>
Date: Mon, 30 Sep 2019 18:34:34 -0400
Subject: trace driven static linking and fixing the inst count bug

---
 Makefile                         | 13 ++++++++++---
 src/abstract_hardware_model.h    |  3 ++-
 src/gpgpu-sim/shader.cc          |  5 +++--
 src/trace-driven/trace_driven.cc |  9 +++++----
 src/trace-driven/trace_opcode.h  |  2 +-
 5 files changed, 21 insertions(+), 11 deletions(-)

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index acef015..ee48514 100644
--- a/Makefile
+++ b/Makefile
@@ -172,15 +172,16 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 
 
 $(SIM_LIB_DIR)/gpgpusim.out: makedirs $(LIBS) cudalib $(SIM_LIB_DIR)/libcudart.so
-	ar rvs   $(SIM_LIB_DIR)/libcudart_static.a\
+	ar rvs   $(SIM_LIB_DIR)/libgpgpusim_static.a\
 			$(SIM_OBJ_FILES_DIR)/libcuda/*.o \
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/*.o \
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/decuda_pred_table/*.o \
 			$(SIM_OBJ_FILES_DIR)/gpgpu-sim/*.o \
 			$(SIM_OBJ_FILES_DIR)/$(INTERSIM)/*.o \
+			$(SIM_OBJ_FILES_DIR)/trace-driven/*.o \
 			$(SIM_OBJ_FILES_DIR)/*.o \
 			$(MCPAT)
-	g++ -std=c++0x -L$(SIM_LIB_DIR) -I$(CUDA_INSTALL_PATH)/include -lcudart -lm -lz -lGL -pthread -o $(SIM_LIB_DIR)/gpgpusim.out src/trace-driven/gpgpusim_trace_driven_main.cc 
+	g++ -std=c++0x -o $(SIM_LIB_DIR)/gpgpusim.out src/trace-driven/gpgpusim_trace_driven_main.cc  -L$(SIM_LIB_DIR) -I$(CUDA_INSTALL_PATH)/include -lgpgpusim_static -lm -lz -lGL -pthread 
 
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
@@ -220,8 +221,14 @@ cuda-sim: makedirs
 	$(MAKE) -C ./src/cuda-sim/ depend
 	$(MAKE) -C ./src/cuda-sim/
 
+TFLAGS = -std=c++0x -I$(CUDA_INSTALL_PATH)/include
+ifneq ($(DEBUG),1)
+	TFLAGS += -O3
+endif
+TFLAGS += -g3 -fPIC
+
 trace-driven: makedirs
-	g++ -fPIC -std=c++0x -I$(CUDA_INSTALL_PATH)/include -c src/trace-driven/trace_driven.cc -o $(SIM_OBJ_FILES_DIR)/trace-driven/trace_driven.o
+	g++ $(TFLAGS) -c src/trace-driven/trace_driven.cc -o $(SIM_OBJ_FILES_DIR)/trace-driven/trace_driven.o
 
 gpgpu-sim_uarch: makedirs cuda-sim
 	$(MAKE) -C ./src/gpgpu-sim/ depend
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index 341e44c..a267d38 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -103,7 +103,8 @@ enum uarch_op_t {
    BARRIER_OP,
    MEMORY_BARRIER_OP,
    CALL_OPS,
-   RET_OPS
+   RET_OPS,
+   EXIT_OPS
 };
 typedef enum uarch_op_t op_type;
 
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index cc85f3f..e5a1e7d 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -1083,6 +1083,7 @@ void scheduler_unit::cycle()
             SCHED_DPRINTF( "Warp (warp_id %u, dynamic_warp_id %u) has valid instruction (%s)\n",
                            (*iter)->get_warp_id(), (*iter)->get_dynamic_warp_id(),
                            m_shader->m_config->gpgpu_ctx->func_sim->ptx_get_insn_str( pc).c_str() );
+
             if( pI ) {
                 assert(valid);
                 if( pc != pI->pc ) {
@@ -1559,8 +1560,8 @@ void shader_core_ctx::warp_inst_complete(const warp_inst_t &inst)
 {
 
   #if 0
-      printf("[warp_inst_complete] uid=%u core=%u warp=%u pc=%#x @ time=%llu issued@%llu\n",
-             inst.get_uid(), m_sid, inst.warp_id(), inst.pc, gpu_tot_sim_cycle + gpu_sim_cycle, inst.get_issue_cycle());
+      printf("[warp_inst_complete] uid=%u core=%u warp=%u pc=%#x @ time=%llu \n",
+             inst.get_uid(), m_sid, inst.warp_id(), inst.pc,  m_gpu->gpu_tot_sim_cycle +  m_gpu->gpu_sim_cycle);
   #endif
 
   if(inst.op_pipe==SP__OP)
diff --git a/src/trace-driven/trace_driven.cc b/src/trace-driven/trace_driven.cc
index b73f6cd..8ebc5ca 100644
--- a/src/trace-driven/trace_driven.cc
+++ b/src/trace-driven/trace_driven.cc
@@ -153,6 +153,7 @@ void trace_shd_warp_t::clear() {
 	warp_traces.clear();
 }
 
+//functional_done
 bool trace_shd_warp_t::trace_done() {
 	return trace_pc==(warp_traces.size());
 }
@@ -537,10 +538,8 @@ void trace_shader_core_ctx::checkExecutionStatusAndUpdate(warp_inst_t &inst, uns
 	if(inst.isatomic())
 		m_warp[inst.warp_id()].inc_n_atomic();
 
-	if ( m_trace_warp[inst.warp_id()].trace_done() ) {
+	if ( inst.op == EXIT_OPS )
 		m_warp[inst.warp_id()].set_completed(t);
-		m_warp[inst.warp_id()].ibuffer_flush();
-	}
 
 }
 
@@ -560,7 +559,9 @@ void trace_shader_core_ctx::func_exec_inst( warp_inst_t &inst )
 			checkExecutionStatusAndUpdate(inst,t,tid);
 		}
 	}
-	if(m_trace_warp[inst.warp_id()].trace_done() )
+	if(m_trace_warp[inst.warp_id()].trace_done() && m_warp[inst.warp_id()].functional_done()) {
+		m_warp[inst.warp_id()].ibuffer_flush();
 		m_barriers.warp_exit( inst.warp_id() );
+	}
 }
 
diff --git a/src/trace-driven/trace_opcode.h b/src/trace-driven/trace_opcode.h
index e74dec4..2138e1c 100644
--- a/src/trace-driven/trace_opcode.h
+++ b/src/trace-driven/trace_opcode.h
@@ -186,7 +186,7 @@ static const std::unordered_map<std::string,OpcodeChar> OpcodeMap = {
 		{"BSSY", OpcodeChar(OP_BSSY, BRANCH_OP)},
 		{"BSYNC", OpcodeChar(OP_BSYNC, BRANCH_OP)},
 		{"CALL", OpcodeChar(OP_CALL, CALL_OPS)},
-		{"EXIT", OpcodeChar(OP_EXIT, BRANCH_OP)},
+		{"EXIT", OpcodeChar(OP_EXIT, EXIT_OPS)},
 		{"JMP", OpcodeChar(OP_JMP, BRANCH_OP)},
 		{"JMX", OpcodeChar(OP_JMX, BRANCH_OP)},
 		{"KILL", OpcodeChar(OP_KILL, BRANCH_OP)},
-- 
cgit v1.3


From ee3c46354c5ce758bf71cfb7bd2e7da7718bcc8a Mon Sep 17 00:00:00 2001
From: Mahmoud <abdallm@purdue.edu>
Date: Mon, 1 Jun 2020 15:09:34 -0400
Subject: renaming the executable

---
 Makefile                                       |   6 +-
 configs/tested-cfgs/SM7_TITANV/gpgpusim.config |   2 +-
 src/trace-driven/gpgpusim_trace_driven_main.cc | 163 -------------------------
 src/trace-driven/main.cc                       | 163 +++++++++++++++++++++++++
 4 files changed, 167 insertions(+), 167 deletions(-)
 delete mode 100644 src/trace-driven/gpgpusim_trace_driven_main.cc
 create mode 100644 src/trace-driven/main.cc

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index ee48514..feb9dd0 100644
--- a/Makefile
+++ b/Makefile
@@ -75,7 +75,7 @@ else
 	TARGETS += $(SIM_LIB_DIR)/libOpenCL.so
 endif
 	TARGETS += cuobjdump_to_ptxplus/cuobjdump_to_ptxplus
-	TARGETS += $(SIM_LIB_DIR)/gpgpusim.out
+	TARGETS += $(SIM_LIB_DIR)/accelsim.out
 
 MCPAT=
 MCPAT_OBJ_DIR=
@@ -171,7 +171,7 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.1; fi
 
 
-$(SIM_LIB_DIR)/gpgpusim.out: makedirs $(LIBS) cudalib $(SIM_LIB_DIR)/libcudart.so
+$(SIM_LIB_DIR)/accelsim.out: makedirs $(LIBS) cudalib $(SIM_LIB_DIR)/libcudart.so
 	ar rvs   $(SIM_LIB_DIR)/libgpgpusim_static.a\
 			$(SIM_OBJ_FILES_DIR)/libcuda/*.o \
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/*.o \
@@ -181,7 +181,7 @@ $(SIM_LIB_DIR)/gpgpusim.out: makedirs $(LIBS) cudalib $(SIM_LIB_DIR)/libcudart.s
 			$(SIM_OBJ_FILES_DIR)/trace-driven/*.o \
 			$(SIM_OBJ_FILES_DIR)/*.o \
 			$(MCPAT)
-	g++ -std=c++0x -o $(SIM_LIB_DIR)/gpgpusim.out src/trace-driven/gpgpusim_trace_driven_main.cc  -L$(SIM_LIB_DIR) -I$(CUDA_INSTALL_PATH)/include -lgpgpusim_static -lm -lz -lGL -pthread 
+	g++ -std=c++0x -o $(SIM_LIB_DIR)/accelsim.out src/trace-driven/main.cc  -L$(SIM_LIB_DIR) -I$(CUDA_INSTALL_PATH)/include -lgpgpusim_static -lm -lz -lGL -pthread 
 
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
diff --git a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
index 64edc67..3fa51ee 100644
--- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
+++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
@@ -131,7 +131,7 @@
 -gpgpu_cache:dl2_texture_only 0
 -gpgpu_dram_partition_queues 64:64:64:64
 -gpgpu_perf_sim_memcpy 1
--gpgpu_memory_partition_indexing 0
+-gpgpu_memory_partition_indexing 2
 
 # 128 KB Inst.
 -gpgpu_cache:il1 N:64:128:16,L:R:f:N:L,S:2:48,4
diff --git a/src/trace-driven/gpgpusim_trace_driven_main.cc b/src/trace-driven/gpgpusim_trace_driven_main.cc
deleted file mode 100644
index f12d39a..0000000
--- a/src/trace-driven/gpgpusim_trace_driven_main.cc
+++ /dev/null
@@ -1,163 +0,0 @@
-// developed by Mahmoud Khairy, Purdue Univ
-// abdallm@purdue.edu
-
-#include <math.h>
-#include <stdio.h>
-#include <time.h>
-#include <fstream>
-#include <iostream>
-#include <sstream>
-#include <string>
-#include <vector>
-
-#include "../../libcuda/gpgpu_context.h"
-#include "../abstract_hardware_model.h"
-#include "../cuda-sim/cuda-sim.h"
-#include "../gpgpu-sim/gpu-sim.h"
-#include "../gpgpu-sim/icnt_wrapper.h"
-#include "../gpgpusim_entrypoint.h"
-#include "../option_parser.h"
-#include "ISA_Def/trace_opcode.h"
-#include "trace_driven.h"
-
-/* TO DO:
- * NOTE: the current version of trace-driven is functionally working fine,
- * but we still need to improve traces compression and simulation speed.
- * This includes:
- * 1- Prefetch concurrent thread that prefetches traces from disk (to not be
- * limited by disk speed) 2- traces compression format a. cfg format and remove
- * thread/block Id from the head b. using zlib library to save in binary format
- *
- * 3- Efficient memory improvement (save string not objects - parse only 10 in
- * the buffer) 4- Seeking capability - thread scheduler (save tb index and warp
- * index info in the traces header) 5- Get rid off traces intermediate files -
- * change the tracer
- */
-gpgpu_sim* gpgpu_trace_sim_init_perf_model(int argc, const char* argv[],
-                                           gpgpu_context* m_gpgpu_context,
-                                           class trace_config* m_config);
-
-int main(int argc, const char** argv) {
-  gpgpu_context* m_gpgpu_context = new gpgpu_context();
-  trace_config tconfig;
-
-  gpgpu_sim* m_gpgpu_sim =
-      gpgpu_trace_sim_init_perf_model(argc, argv, m_gpgpu_context, &tconfig);
-  m_gpgpu_sim->init();
-
-  // for each kernel
-  // load file
-  // parse and create kernel info
-  // launch
-  // while loop till the end of the end kernel execution
-  // prints stats
-
-  trace_parser tracer(tconfig.get_traces_filename(), m_gpgpu_sim,
-                      m_gpgpu_context);
-  tconfig.parse_config();
-
-  std::vector<std::string> commandlist = tracer.parse_kernellist_file();
-
-  for (unsigned i = 0; i < commandlist.size(); ++i) {
-    trace_kernel_info_t* kernel_info = NULL;
-    if (commandlist[i].substr(0, 6) == "Memcpy") {
-      size_t addre, Bcount;
-      tracer.parse_memcpy_info(commandlist[i], addre, Bcount);
-      std::cout << commandlist[i] << std::endl;
-      m_gpgpu_sim->perf_memcpy_to_gpu(addre, Bcount);
-      continue;
-    } else {
-      kernel_info = tracer.parse_kernel_info(commandlist[i], &tconfig);
-      m_gpgpu_sim->launch(kernel_info);
-    }
-
-    bool active = false;
-    bool sim_cycles = false;
-    bool break_limit = false;
-
-    do {
-      if (!m_gpgpu_sim->active()) break;
-
-      // performance simulation
-      if (m_gpgpu_sim->active()) {
-        m_gpgpu_sim->cycle();
-        sim_cycles = true;
-        m_gpgpu_sim->deadlock_check();
-      } else {
-        if (m_gpgpu_sim->cycle_insn_cta_max_hit()) {
-          m_gpgpu_context->the_gpgpusim->g_stream_manager
-              ->stop_all_running_kernels();
-          break_limit = true;
-        }
-      }
-
-      active = m_gpgpu_sim->active();
-
-    } while (active);
-
-    if (kernel_info) {
-      tracer.kernel_finalizer(kernel_info);
-      m_gpgpu_sim->print_stats();
-    }
-
-    if (sim_cycles) {
-      m_gpgpu_sim->update_stats();
-      m_gpgpu_context->print_simulation_time();
-    }
-
-    if (break_limit) {
-      printf(
-          "GPGPU-Sim: ** break due to reaching the maximum cycles (or "
-          "instructions) **\n");
-      fflush(stdout);
-      exit(1);
-    }
-  }
-
-  // we print this message to inform the gpgpu-simulation stats_collect script
-  // that we are done
-  printf("GPGPU-Sim: *** simulation thread exiting ***\n");
-  printf("GPGPU-Sim: *** exit detected ***\n");
-
-  return 1;
-}
-
-gpgpu_sim* gpgpu_trace_sim_init_perf_model(int argc, const char* argv[],
-                                           gpgpu_context* m_gpgpu_context,
-                                           trace_config* m_config) {
-  srand(1);
-  print_splash();
-
-  option_parser_t opp = option_parser_create();
-
-  m_gpgpu_context->ptx_reg_options(opp);
-  m_gpgpu_context->func_sim->ptx_opcocde_latency_options(opp);
-
-  icnt_reg_options(opp);
-
-  m_gpgpu_context->the_gpgpusim->g_the_gpu_config =
-      new gpgpu_sim_config(m_gpgpu_context);
-  m_gpgpu_context->the_gpgpusim->g_the_gpu_config->reg_options(
-      opp);  // register GPU microrachitecture options
-  m_config->reg_options(opp);
-
-  option_parser_cmdline(opp, argc, argv);  // parse configuration options
-  fprintf(stdout, "GPGPU-Sim: Configuration options:\n\n");
-  option_parser_print(opp, stdout);
-  // Set the Numeric locale to a standard locale where a decimal point is a
-  // "dot" not a "comma" so it does the parsing correctly independent of the
-  // system environment variables
-  assert(setlocale(LC_NUMERIC, "C"));
-  m_gpgpu_context->the_gpgpusim->g_the_gpu_config->init();
-
-  m_gpgpu_context->the_gpgpusim->g_the_gpu = new trace_gpgpu_sim(
-      *(m_gpgpu_context->the_gpgpusim->g_the_gpu_config), m_gpgpu_context);
-
-  m_gpgpu_context->the_gpgpusim->g_stream_manager =
-      new stream_manager((m_gpgpu_context->the_gpgpusim->g_the_gpu),
-                         m_gpgpu_context->func_sim->g_cuda_launch_blocking);
-
-  m_gpgpu_context->the_gpgpusim->g_simulation_starttime = time((time_t*)NULL);
-
-  return m_gpgpu_context->the_gpgpusim->g_the_gpu;
-}
diff --git a/src/trace-driven/main.cc b/src/trace-driven/main.cc
new file mode 100644
index 0000000..f12d39a
--- /dev/null
+++ b/src/trace-driven/main.cc
@@ -0,0 +1,163 @@
+// developed by Mahmoud Khairy, Purdue Univ
+// abdallm@purdue.edu
+
+#include <math.h>
+#include <stdio.h>
+#include <time.h>
+#include <fstream>
+#include <iostream>
+#include <sstream>
+#include <string>
+#include <vector>
+
+#include "../../libcuda/gpgpu_context.h"
+#include "../abstract_hardware_model.h"
+#include "../cuda-sim/cuda-sim.h"
+#include "../gpgpu-sim/gpu-sim.h"
+#include "../gpgpu-sim/icnt_wrapper.h"
+#include "../gpgpusim_entrypoint.h"
+#include "../option_parser.h"
+#include "ISA_Def/trace_opcode.h"
+#include "trace_driven.h"
+
+/* TO DO:
+ * NOTE: the current version of trace-driven is functionally working fine,
+ * but we still need to improve traces compression and simulation speed.
+ * This includes:
+ * 1- Prefetch concurrent thread that prefetches traces from disk (to not be
+ * limited by disk speed) 2- traces compression format a. cfg format and remove
+ * thread/block Id from the head b. using zlib library to save in binary format
+ *
+ * 3- Efficient memory improvement (save string not objects - parse only 10 in
+ * the buffer) 4- Seeking capability - thread scheduler (save tb index and warp
+ * index info in the traces header) 5- Get rid off traces intermediate files -
+ * change the tracer
+ */
+gpgpu_sim* gpgpu_trace_sim_init_perf_model(int argc, const char* argv[],
+                                           gpgpu_context* m_gpgpu_context,
+                                           class trace_config* m_config);
+
+int main(int argc, const char** argv) {
+  gpgpu_context* m_gpgpu_context = new gpgpu_context();
+  trace_config tconfig;
+
+  gpgpu_sim* m_gpgpu_sim =
+      gpgpu_trace_sim_init_perf_model(argc, argv, m_gpgpu_context, &tconfig);
+  m_gpgpu_sim->init();
+
+  // for each kernel
+  // load file
+  // parse and create kernel info
+  // launch
+  // while loop till the end of the end kernel execution
+  // prints stats
+
+  trace_parser tracer(tconfig.get_traces_filename(), m_gpgpu_sim,
+                      m_gpgpu_context);
+  tconfig.parse_config();
+
+  std::vector<std::string> commandlist = tracer.parse_kernellist_file();
+
+  for (unsigned i = 0; i < commandlist.size(); ++i) {
+    trace_kernel_info_t* kernel_info = NULL;
+    if (commandlist[i].substr(0, 6) == "Memcpy") {
+      size_t addre, Bcount;
+      tracer.parse_memcpy_info(commandlist[i], addre, Bcount);
+      std::cout << commandlist[i] << std::endl;
+      m_gpgpu_sim->perf_memcpy_to_gpu(addre, Bcount);
+      continue;
+    } else {
+      kernel_info = tracer.parse_kernel_info(commandlist[i], &tconfig);
+      m_gpgpu_sim->launch(kernel_info);
+    }
+
+    bool active = false;
+    bool sim_cycles = false;
+    bool break_limit = false;
+
+    do {
+      if (!m_gpgpu_sim->active()) break;
+
+      // performance simulation
+      if (m_gpgpu_sim->active()) {
+        m_gpgpu_sim->cycle();
+        sim_cycles = true;
+        m_gpgpu_sim->deadlock_check();
+      } else {
+        if (m_gpgpu_sim->cycle_insn_cta_max_hit()) {
+          m_gpgpu_context->the_gpgpusim->g_stream_manager
+              ->stop_all_running_kernels();
+          break_limit = true;
+        }
+      }
+
+      active = m_gpgpu_sim->active();
+
+    } while (active);
+
+    if (kernel_info) {
+      tracer.kernel_finalizer(kernel_info);
+      m_gpgpu_sim->print_stats();
+    }
+
+    if (sim_cycles) {
+      m_gpgpu_sim->update_stats();
+      m_gpgpu_context->print_simulation_time();
+    }
+
+    if (break_limit) {
+      printf(
+          "GPGPU-Sim: ** break due to reaching the maximum cycles (or "
+          "instructions) **\n");
+      fflush(stdout);
+      exit(1);
+    }
+  }
+
+  // we print this message to inform the gpgpu-simulation stats_collect script
+  // that we are done
+  printf("GPGPU-Sim: *** simulation thread exiting ***\n");
+  printf("GPGPU-Sim: *** exit detected ***\n");
+
+  return 1;
+}
+
+gpgpu_sim* gpgpu_trace_sim_init_perf_model(int argc, const char* argv[],
+                                           gpgpu_context* m_gpgpu_context,
+                                           trace_config* m_config) {
+  srand(1);
+  print_splash();
+
+  option_parser_t opp = option_parser_create();
+
+  m_gpgpu_context->ptx_reg_options(opp);
+  m_gpgpu_context->func_sim->ptx_opcocde_latency_options(opp);
+
+  icnt_reg_options(opp);
+
+  m_gpgpu_context->the_gpgpusim->g_the_gpu_config =
+      new gpgpu_sim_config(m_gpgpu_context);
+  m_gpgpu_context->the_gpgpusim->g_the_gpu_config->reg_options(
+      opp);  // register GPU microrachitecture options
+  m_config->reg_options(opp);
+
+  option_parser_cmdline(opp, argc, argv);  // parse configuration options
+  fprintf(stdout, "GPGPU-Sim: Configuration options:\n\n");
+  option_parser_print(opp, stdout);
+  // Set the Numeric locale to a standard locale where a decimal point is a
+  // "dot" not a "comma" so it does the parsing correctly independent of the
+  // system environment variables
+  assert(setlocale(LC_NUMERIC, "C"));
+  m_gpgpu_context->the_gpgpusim->g_the_gpu_config->init();
+
+  m_gpgpu_context->the_gpgpusim->g_the_gpu = new trace_gpgpu_sim(
+      *(m_gpgpu_context->the_gpgpusim->g_the_gpu_config), m_gpgpu_context);
+
+  m_gpgpu_context->the_gpgpusim->g_stream_manager =
+      new stream_manager((m_gpgpu_context->the_gpgpusim->g_the_gpu),
+                         m_gpgpu_context->func_sim->g_cuda_launch_blocking);
+
+  m_gpgpu_context->the_gpgpusim->g_simulation_starttime = time((time_t*)NULL);
+
+  return m_gpgpu_context->the_gpgpusim->g_the_gpu;
+}
-- 
cgit v1.3


From 183467b5d3063bf94a1b0ea8459f4d70a4da6942 Mon Sep 17 00:00:00 2001
From: Mahmoud <abdallm@purdue.edu>
Date: Mon, 1 Jun 2020 17:48:53 -0400
Subject: adding Makefile

---
 Makefile | 21 +--------------------
 1 file changed, 1 insertion(+), 20 deletions(-)

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index feb9dd0..4f35cc3 100644
--- a/Makefile
+++ b/Makefile
@@ -57,7 +57,7 @@ endif
 
 $(shell mkdir -p $(SIM_OBJ_FILES_DIR)/libcuda && echo "const char *g_gpgpusim_build_string=\"$(GPGPUSIM_BUILD)\";" > $(SIM_OBJ_FILES_DIR)/detailed_version)
 
-LIBS = cuda-sim gpgpu-sim_uarch $(INTERSIM) gpgpusimlib trace-driven
+LIBS = cuda-sim gpgpu-sim_uarch $(INTERSIM) gpgpusimlib
 
 
 TARGETS =
@@ -75,7 +75,6 @@ else
 	TARGETS += $(SIM_LIB_DIR)/libOpenCL.so
 endif
 	TARGETS += cuobjdump_to_ptxplus/cuobjdump_to_ptxplus
-	TARGETS += $(SIM_LIB_DIR)/accelsim.out
 
 MCPAT=
 MCPAT_OBJ_DIR=
@@ -150,7 +149,6 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 			$(SIM_OBJ_FILES_DIR)/cuda-sim/decuda_pred_table/*.o \
 			$(SIM_OBJ_FILES_DIR)/gpgpu-sim/*.o \
 			$(SIM_OBJ_FILES_DIR)/$(INTERSIM)/*.o \
-			$(SIM_OBJ_FILES_DIR)/trace-driven/*.o \
 			$(SIM_OBJ_FILES_DIR)/*.o -lm -lz -lGL -pthread \
 			$(MCPAT) \
 			-o $(SIM_LIB_DIR)/libcudart.so
@@ -170,19 +168,6 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.1; fi
 
-
-$(SIM_LIB_DIR)/accelsim.out: makedirs $(LIBS) cudalib $(SIM_LIB_DIR)/libcudart.so
-	ar rvs   $(SIM_LIB_DIR)/libgpgpusim_static.a\
-			$(SIM_OBJ_FILES_DIR)/libcuda/*.o \
-			$(SIM_OBJ_FILES_DIR)/cuda-sim/*.o \
-			$(SIM_OBJ_FILES_DIR)/cuda-sim/decuda_pred_table/*.o \
-			$(SIM_OBJ_FILES_DIR)/gpgpu-sim/*.o \
-			$(SIM_OBJ_FILES_DIR)/$(INTERSIM)/*.o \
-			$(SIM_OBJ_FILES_DIR)/trace-driven/*.o \
-			$(SIM_OBJ_FILES_DIR)/*.o \
-			$(MCPAT)
-	g++ -std=c++0x -o $(SIM_LIB_DIR)/accelsim.out src/trace-driven/main.cc  -L$(SIM_LIB_DIR) -I$(CUDA_INSTALL_PATH)/include -lgpgpusim_static -lm -lz -lGL -pthread 
-
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
 			$(SIM_OBJ_FILES_DIR)/libcuda/*.o \
@@ -227,9 +212,6 @@ ifneq ($(DEBUG),1)
 endif
 TFLAGS += -g3 -fPIC
 
-trace-driven: makedirs
-	g++ $(TFLAGS) -c src/trace-driven/trace_driven.cc -o $(SIM_OBJ_FILES_DIR)/trace-driven/trace_driven.o
-
 gpgpu-sim_uarch: makedirs cuda-sim
 	$(MAKE) -C ./src/gpgpu-sim/ depend
 	$(MAKE) -C ./src/gpgpu-sim/
@@ -256,7 +238,6 @@ makedirs:
 	if [ ! -d $(SIM_OBJ_FILES_DIR)/cuda-sim ]; then mkdir -p $(SIM_OBJ_FILES_DIR)/cuda-sim; fi;
 	if [ ! -d $(SIM_OBJ_FILES_DIR)/cuda-sim/decuda_pred_table ]; then mkdir -p $(SIM_OBJ_FILES_DIR)/cuda-sim/decuda_pred_table; fi;
 	if [ ! -d $(SIM_OBJ_FILES_DIR)/gpgpu-sim ]; then mkdir -p $(SIM_OBJ_FILES_DIR)/gpgpu-sim; fi;
-	if [ ! -d $(SIM_OBJ_FILES_DIR)/trace-driven ]; then mkdir -p $(SIM_OBJ_FILES_DIR)/trace-driven; fi;
 	if [ ! -d $(SIM_OBJ_FILES_DIR)/libopencl ]; then mkdir -p $(SIM_OBJ_FILES_DIR)/libopencl; fi;
 	if [ ! -d $(SIM_OBJ_FILES_DIR)/libopencl/bin ]; then mkdir -p $(SIM_OBJ_FILES_DIR)/libopencl/bin; fi;
 	if [ ! -d $(SIM_OBJ_FILES_DIR)/$(INTERSIM) ]; then mkdir -p $(SIM_OBJ_FILES_DIR)/$(INTERSIM); fi;
-- 
cgit v1.3


From 03f6d7eb714a93c0ddff024fb92288ae9131f01c Mon Sep 17 00:00:00 2001
From: tgrogers <tgrogers@purdue.edu>
Date: Wed, 10 Jun 2020 22:01:35 -0400
Subject: CUDA 11 support

---
 Makefile             | 1 +
 linux-so-version.txt | 2 ++
 setup_environment    | 2 +-
 3 files changed, 4 insertions(+), 1 deletion(-)

(limited to 'Makefile')

diff --git a/Makefile b/Makefile
index f4e8e4b..5e964b8 100644
--- a/Makefile
+++ b/Makefile
@@ -167,6 +167,7 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.2 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.2; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.1; fi
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.11.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.11.0; fi
 
 
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
diff --git a/linux-so-version.txt b/linux-so-version.txt
index 45c40dd..3db07ac 100644
--- a/linux-so-version.txt
+++ b/linux-so-version.txt
@@ -8,5 +8,7 @@ libcudart.so.10.0{
 };
 libcudart.so.10.1{
 };
+libcudart.so.11.0{
+};
 libcuda.so.1{
 };
diff --git a/setup_environment b/setup_environment
index daf7fab..07d0788 100644
--- a/setup_environment
+++ b/setup_environment
@@ -48,7 +48,7 @@ CC_VERSION=`gcc --version | head -1 | awk '{for(i=1;i<=NF;i++){ if(match($i,/^[0
 
 CUDA_VERSION_STRING=`$CUDA_INSTALL_PATH/bin/nvcc --version | awk '/release/ {print $5;}' | sed 's/,//'`;
 export CUDA_VERSION_NUMBER=`echo $CUDA_VERSION_STRING | sed 's/\./ /' | awk '{printf("%02u%02u", 10*int($1), 10*$2);}'`
-if [ $CUDA_VERSION_NUMBER -gt 10100 -o $CUDA_VERSION_NUMBER -lt 2030  ]; then
+if [ $CUDA_VERSION_NUMBER -gt 11100 -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)";
 	echo $CUDA_VERSION_NUMBER
   return
-- 
cgit v1.3