adding traces generator

7 files changed, 882 insertions, 5 deletions

diff --git a/src/trace-driven/gpgpusim_trace_driven_main.cc b/src/trace-driven/gpgpusim_trace_driven_main.cc
index 76c2cda..82f284c 100644
--- a/src/trace-driven/gpgpusim_trace_driven_main.cc
+++ b/src/trace-driven/gpgpusim_trace_driven_main.cc
@@ -394,14 +394,14 @@ bool trace_warp_inst_t::parse_from_string(std::string trace){
 	num_operands = num_regs;
 	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];
+		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];
-		arch_reg.src[m]=reg_srcs[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
 
diff --git a/src/trace-driven/trace_opcode.h b/src/trace-driven/trace_opcode.h
index 827c32c..e74dec4 100644
--- a/src/trace-driven/trace_opcode.h
+++ b/src/trace-driven/trace_opcode.h
@@ -75,7 +75,8 @@ static const std::unordered_map<std::string,OpcodeChar> OpcodeMap = {
 		{"FSET", OpcodeChar(OP_FSET, SP_OP)},
 		{"FSETP", OpcodeChar(OP_FSETP, SP_OP)},
 		{"FSWZADD", OpcodeChar(OP_FSWZADD, SP_OP)},
-		{"MUFU", OpcodeChar(OP_MUFU, SP_OP)},
+		//SFU
+		{"MUFU", OpcodeChar(OP_MUFU, SFU_OP)},
 
 		//Floating Point 16 Instructions
 		{"HADD2", OpcodeChar(OP_HADD2, SP_OP)},
diff --git a/src/trace-driven/traces-generator/Gbit_tool/Makefile b/src/trace-driven/traces-generator/Gbit_tool/Makefile
new file mode 100644
index 0000000..fc0209a
--- /dev/null
+++ b/src/trace-driven/traces-generator/Gbit_tool/Makefile
@@ -0,0 +1,23 @@
+NVCC=nvcc -ccbin=`which gcc` -D_FORCE_INLINES
+NVBIT_PATH=../../core
+INCLUDES=-I$(NVBIT_PATH)
+LIBS=-L$(NVBIT_PATH) -lnvbit
+NVCC_PATH=-L $(subst bin/nvcc,lib64,$(shell which nvcc | tr -s /))
+SOURCES=$(wildcard *.cu)
+OBJECTS=$(SOURCES:.cu=.o)
+ARCH=35
+
+mkfile_path := $(abspath $(lastword $(MAKEFILE_LIST)))
+current_dir := $(notdir $(patsubst %/,%,$(dir $(mkfile_path))))
+
+all: $(OBJECTS) $(NVBIT_PATH)/libnvbit.a
+	$(NVCC) -arch=sm_$(ARCH) -O3 *.o $(LIBS) $(NVCC_PATH) -lcuda -lcudart_static -shared -o ${current_dir}.so
+
+%.o: %.cu
+	$(NVCC) -dc -c -std=c++11 $(INCLUDES) -Xptxas -cloning=no -maxrregcount=16 -Xcompiler -Wall -arch=sm_$(ARCH) -O3 -Xcompiler -fPIC $< -o $@
+
+$(NVBIT_PATH)/libnvbit.a:
+	make -C $(NVBIT_PATH)
+
+clean:
+	rm -f *.so *.o
diff --git a/src/trace-driven/traces-generator/Gbit_tool/traceall.cu b/src/trace-driven/traces-generator/Gbit_tool/traceall.cu
new file mode 100644
index 0000000..1426528
--- /dev/null
+++ b/src/trace-driven/traces-generator/Gbit_tool/traceall.cu
@@ -0,0 +1,646 @@
+/* Copyright (c) 2017, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ *  * Redistributions of source code must retain the above copyright
+ *    notice, this list of conditions and the following disclaimer.
+ *  * Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimer in the
+ *    documentation and/or other materials provided with the distribution.
+ *  * Neither the name of NVIDIA CORPORATION nor the names of its
+ *    contributors may be used to endorse or promote products derived
+ *    from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
+ * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+/* Author: Oreste Villa, [email protected] - 2018 */
+
+#include <assert.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <unistd.h>
+#include <string>
+#include <map>
+#include <inttypes.h>
+#include <unistd.h>
+#include <vector>
+#include <iostream>
+#include <bitset>
+#include <sys/stat.h>
+#include <sstream>
+#include <algorithm>
+#include <iterator>
+/* every tool needs to include this once */
+#include "nvbit_tool.h"
+
+/* nvbit interface file */
+#include "nvbit.h"
+
+/* for channel */
+#include "utils/channel.hpp"
+
+/* for _cuda_safe and GET_VAR* macros */
+#include "macros.h"
+
+/* Channel used to communicate from GPU to CPU receiving thread */
+#define CHANNEL_SIZE (1l << 20)
+static __managed__ ChannelDev channel_dev;
+static ChannelHost channel_host;
+
+/* receiving thread and its control variables */
+pthread_t recv_thread;
+volatile bool recv_thread_started = false;
+volatile bool recv_thread_receiving = false;
+
+/* skip flag used to avoid re-entry on the nvbit_callback when issuing
+ * flush_channel kernel call */
+bool skip_flag = false;
+
+/* global control variables for this tool */
+uint32_t instr_begin_interval = 0;
+uint32_t instr_end_interval = UINT32_MAX;
+int verbose = 0;
+
+/* opcode to id map and reverse map  */
+std::map<std::string, int> opcode_to_id_map;
+std::map<int, std::string> id_to_opcode_map;
+
+/* kernel instruction counter, updated by the GPU */
+static __managed__ uint64_t total_dynamic_instr_counter = 0;
+static __managed__ uint64_t reported_dynamic_instr_counter = 0;
+static __managed__ uint64_t dynamic_instr_limit = 0;
+uint64_t dynamic_instr_limit_input = 0;  //0 means no limit
+
+#define MAX_SRC 4
+/* information collected in the instrumentation function */
+typedef struct {
+	int cta_id_x;
+	int cta_id_y;
+	int cta_id_z;
+	int warpid_tb;
+	int warpid_sm;
+	int sm_id;
+	int opcode_id;
+	uint64_t addrs[32];
+	uint32_t vpc;
+	bool is_mem;
+	int32_t GPRDst;
+	int32_t GPRSrcs[MAX_SRC];
+	int32_t numSrcs;
+	int32_t width;
+	uint32_t active_mask;
+
+} mem_access_t;
+
+/* Instrumentation function that we want to inject, please note the use of
+ * 1. extern "C" __device__ __noinline__
+ *    To prevent "dead"-code elimination by the compiler.
+ * 2. NVBIT_EXPORT_FUNC(dev_func)
+ *    To notify nvbit the name of the function we want to inject.
+ *    This name must match exactly the function name.
+ */
+extern "C" __device__ __noinline__ void instrument_mem(int pred, int opcode_id, int32_t vpc,
+		uint32_t reg_high,
+		uint32_t reg_low,
+		int32_t imm,
+		int32_t srcReg1, int32_t srcReg2, int32_t desReg, int32_t width) {
+	if (!pred) {
+		return;
+	}
+
+	uint32_t active_mask = __ballot(1);
+	const int laneid = get_laneid();
+	const int first_laneid = __ffs(active_mask) - 1;
+
+	if (dynamic_instr_limit && total_dynamic_instr_counter >= dynamic_instr_limit)
+		if (first_laneid == laneid) {
+			atomicAdd((unsigned long long*)&total_dynamic_instr_counter, 1);
+			return;
+		}
+
+	mem_access_t ma;
+
+	/* collect memory address information */
+	int64_t base_addr = (((uint64_t)reg_high) << 32) | ((uint64_t)reg_low);
+	uint64_t addr = base_addr + imm;
+	for (int i = 0; i < 32; i++) {
+		ma.addrs[i] = __shfl(addr, i);
+	}
+
+	int4 cta = get_ctaid();
+	int uniqe_threadId = threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x;
+	ma.warpid_tb = uniqe_threadId/32;
+
+	ma.cta_id_x = cta.x;
+	ma.cta_id_y = cta.y;
+	ma.cta_id_z = cta.z;
+	ma.warpid_sm = get_warpid();
+	ma.opcode_id = opcode_id;
+	ma.is_mem = true;
+	ma.vpc = vpc;
+	ma.width = width;
+	ma.GPRDst = desReg;
+	ma.GPRSrcs[0] = srcReg1;
+	ma.GPRSrcs[1] = srcReg2;
+	ma.GPRSrcs[2] = -1;
+	ma.GPRSrcs[3] = -1;
+	ma.numSrcs = 2;
+	ma.active_mask = active_mask;
+	ma.sm_id =  get_smid();
+
+	/* first active lane pushes information on the channel */
+	if (first_laneid == laneid) {
+		channel_dev.push(&ma, sizeof(mem_access_t));
+		atomicAdd((unsigned long long*)&total_dynamic_instr_counter, 1);
+		atomicAdd((unsigned long long*)&reported_dynamic_instr_counter, 1);
+	}
+}
+NVBIT_EXPORT_FUNC(instrument_mem);
+
+
+extern "C" __device__ __noinline__ void instrument_inst(int pred, int opcode_id,
+		uint32_t vpc, int desReg, int srcReg1, int srcReg2, int srcReg3, int srcReg4, int srcNum) {
+	if (!pred) {
+		return;
+	}
+
+	int active_mask = __ballot(1);
+	const int laneid = get_laneid();
+	const int first_laneid = __ffs(active_mask) - 1;
+
+	if (dynamic_instr_limit && total_dynamic_instr_counter >= dynamic_instr_limit)
+		if (first_laneid == laneid) {
+			atomicAdd((unsigned long long*)&total_dynamic_instr_counter, 1);
+			return;
+		}
+
+
+	mem_access_t ma;
+
+	int4 cta = get_ctaid();
+	int uniqe_threadId = threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x;
+	ma.warpid_tb = uniqe_threadId/32;
+
+	ma.cta_id_x = cta.x;
+	ma.cta_id_y = cta.y;
+	ma.cta_id_z = cta.z;
+	ma.warpid_sm = get_warpid();
+	ma.opcode_id = opcode_id;
+	ma.is_mem = false;
+	ma.vpc = vpc;
+
+	ma.GPRDst = desReg;
+	ma.numSrcs = srcNum;    //this is the total src number including the register and others
+	ma.GPRSrcs[0] = srcReg1;
+	ma.GPRSrcs[1] = srcReg2;
+	ma.GPRSrcs[2] = srcReg3;
+	ma.GPRSrcs[3] = srcReg4;
+
+	ma.active_mask = active_mask;
+	ma.sm_id =  get_smid();
+
+	/* first active lane pushes information on the channel */
+	if (first_laneid == laneid) {
+		channel_dev.push(&ma, sizeof(mem_access_t));
+		atomicAdd((unsigned long long*)&total_dynamic_instr_counter, 1);
+		atomicAdd((unsigned long long*)&reported_dynamic_instr_counter, 1);
+	}
+}
+
+NVBIT_EXPORT_FUNC(instrument_inst);
+
+void nvbit_at_init() {
+	setenv("CUDA_MANAGED_FORCE_DEVICE_ALLOC", "1", 1);
+	GET_VAR_INT(
+			instr_begin_interval, "INSTR_BEGIN", 0,
+			"Beginning of the instruction interval where to apply instrumentation");
+	GET_VAR_INT(
+			instr_end_interval, "INSTR_END", UINT32_MAX,
+			"End of the instruction interval where to apply instrumentation");
+	GET_VAR_LONG(
+			dynamic_instr_limit_input, "DYNAMIC_INSTR_LIMIT", 0,
+			"Limit of the number instructions to be printed, 0 means no limit");
+	GET_VAR_INT(verbose, "TOOL_VERBOSE", 0, "Enable verbosity inside the tool");
+	std::string pad(100, '-');
+	printf("%s\n", pad.c_str());
+}
+
+/* instrument each memory instruction adding a call to the above instrumentation
+ * function */
+void nvbit_at_function_first_load(CUcontext ctx, CUfunction f) {
+
+	dynamic_instr_limit = dynamic_instr_limit_input;
+
+	const std::vector<Instr *> &instrs = nvbit_get_instrs(ctx, f);
+	if (verbose) {
+		printf("Inspecting function %s at address 0x%lx\n",
+				nvbit_get_func_name(ctx, f), nvbit_get_func_addr(f));
+	}
+
+	uint32_t cnt = 0;
+	/* iterate on all the static instructions in the function */
+	for (auto instr : instrs) {
+		if (cnt < instr_begin_interval || cnt >= instr_end_interval ) {
+			cnt++;
+			continue;
+		}
+		//if (verbose) {
+		instr->printDecoded();
+		//}
+
+		if (opcode_to_id_map.find(instr->getOpcode()) ==
+				opcode_to_id_map.end()) {
+			int opcode_id = opcode_to_id_map.size();
+			opcode_to_id_map[instr->getOpcode()] = opcode_id;
+			id_to_opcode_map[opcode_id] = instr->getOpcode();
+		}
+
+		int opcode_id = opcode_to_id_map[instr->getOpcode()];
+
+		//TO DO: handle generic and TEX memory space
+		if(instr->isLoad() && !instr->isStore() && instr->getMemOpType() != Instr::CONSTANT) {   //Mem load inst //ignore constant for now
+			assert(instr->getNumOperands() == 2);
+
+			/* get the operand */
+			const Instr::operand_t *dst = instr->getOperand(0);
+			const Instr::operand_t *src = instr->getOperand(1);
+
+			assert(dst->type == Instr::REG);
+			assert(src->type == Instr::MREF);
+
+			/* insert call to the instrumentation function with its
+			 * arguments */
+			nvbit_insert_call(instr, "instrument_mem", IPOINT_BEFORE);
+			nvbit_add_call_arg_pred_val(instr);
+			nvbit_add_call_arg_const_val32(instr, opcode_id);
+			nvbit_add_call_arg_const_val32(instr, (int)instr->getOffset());
+			if (instr->isExtended()) {
+				nvbit_add_call_arg_reg_val(instr, (int)src->value[0] + 1);
+			} else {
+				nvbit_add_call_arg_reg_val(instr, (int)Instr::RZ);
+			}
+			nvbit_add_call_arg_reg_val(instr, (int)src->value[0]);
+			nvbit_add_call_arg_const_val32(instr, (int)src->value[1]);
+			nvbit_add_call_arg_const_val32(instr, (int)src->value[0]);
+			nvbit_add_call_arg_const_val32(instr, -1);
+			nvbit_add_call_arg_const_val32(instr, (int)dst->value[0]);
+			nvbit_add_call_arg_const_val32(instr, (int)instr->getSize());
+		}
+		else if(instr->isStore() && !instr->isLoad() && instr->getMemOpType() != Instr::CONSTANT) {   //Mem store inst //ignore constant for now
+			assert(instr->getNumOperands() == 2);
+
+			/* get the operand */
+			const Instr::operand_t *dst = instr->getOperand(0);
+			const Instr::operand_t *src = instr->getOperand(1);
+
+			assert(dst->type == Instr::MREF);
+			assert(src->type == Instr::REG);
+
+			/* insert call to the instrumentation function with its
+			 * arguments */
+			nvbit_insert_call(instr, "instrument_mem", IPOINT_BEFORE);
+			nvbit_add_call_arg_pred_val(instr);
+			nvbit_add_call_arg_const_val32(instr, opcode_id);
+			nvbit_add_call_arg_const_val32(instr, (int)instr->getOffset());
+			if (instr->isExtended()) {
+				nvbit_add_call_arg_reg_val(instr, (int)dst->value[0] + 1);
+			} else {
+				nvbit_add_call_arg_reg_val(instr, (int)Instr::RZ);
+			}
+			nvbit_add_call_arg_reg_val(instr, (int)dst->value[0]);
+			nvbit_add_call_arg_const_val32(instr, (int)dst->value[1]);
+			nvbit_add_call_arg_const_val32(instr, (int)dst->value[0]);
+			nvbit_add_call_arg_const_val32(instr, (int)src->value[0]);
+			nvbit_add_call_arg_const_val32(instr, -1);
+			nvbit_add_call_arg_const_val32(instr, (int)instr->getSize());
+		}
+		else if(instr->isLoad() && instr->isStore() && instr->getMemOpType() != Instr::CONSTANT) {   //if it is load and store i.e. atomic inst 
+			assert(instr->getNumOperands() == 2);
+
+			/* get the operand */
+			const Instr::operand_t *dst = instr->getOperand(0);
+			const Instr::operand_t *src = instr->getOperand(1);
+
+			assert(dst->type == Instr::MREF);
+			assert(src->type == Instr::REG);
+
+			/* insert call to the instrumentation function with its
+			 * arguments */
+			nvbit_insert_call(instr, "instrument_mem", IPOINT_BEFORE);
+			nvbit_add_call_arg_pred_val(instr);
+			nvbit_add_call_arg_const_val32(instr, opcode_id);
+			nvbit_add_call_arg_const_val32(instr, (int)instr->getOffset());
+			if (instr->isExtended()) {
+				nvbit_add_call_arg_reg_val(instr, (int)dst->value[0] + 1);
+			} else {
+				nvbit_add_call_arg_reg_val(instr, (int)Instr::RZ);
+			}
+			nvbit_add_call_arg_reg_val(instr, (int)dst->value[0]);
+			nvbit_add_call_arg_const_val32(instr, (int)dst->value[1]);
+			nvbit_add_call_arg_const_val32(instr, (int)dst->value[0]);
+			nvbit_add_call_arg_const_val32(instr, (int)src->value[0]);
+			nvbit_add_call_arg_const_val32(instr, -1);
+			nvbit_add_call_arg_const_val32(instr, (int)instr->getSize());
+		}
+		else    //Other ALU, FP, DP insts
+		{
+
+			nvbit_insert_call(instr, "instrument_inst", IPOINT_BEFORE);
+			nvbit_add_call_arg_pred_val(instr);
+			nvbit_add_call_arg_const_val32(instr, opcode_id);
+			nvbit_add_call_arg_const_val32(instr, (int)instr->getOffset());
+			int srcNum = 0;
+			for (int i = 0; i < MAX_SRC+1; i++) {
+				/* get the operand "i" */
+				if(i < instr->getNumOperands()) {
+					const Instr::operand_t *op = instr->getOperand(i);
+					if (op->type == Instr::REG)
+						nvbit_add_call_arg_const_val32(instr, (int)op->value[0]);
+					else
+						nvbit_add_call_arg_const_val32(instr, -1);
+
+					srcNum++;
+				}
+				else
+					nvbit_add_call_arg_const_val32(instr, -1);
+			}
+			nvbit_add_call_arg_const_val32(instr, srcNum);
+		}
+		cnt++;
+	}
+}
+
+__global__ void flush_channel() {
+	/* push memory access with negative cta id to communicate the kernel is
+	 * completed */
+	mem_access_t ma;
+	ma.cta_id_x = -1;
+	channel_dev.push(&ma, sizeof(mem_access_t));
+
+	/* flush channel */
+	channel_dev.flush();
+}
+
+static FILE *resultsFile = NULL;
+static FILE *kernelsFile= NULL;
+static FILE *statsFile= NULL;
+static int kernelid = 1;
+
+unsigned old_total_insts = 0;
+unsigned old_total_reported_insts = 0;
+
+
+void nvbit_at_cuda_event(CUcontext ctx, int is_exit, nvbit_api_cuda_t cbid,
+		const char *name, void *params, CUresult *pStatus) {
+	if (skip_flag) return;
+
+	if (cbid == API_CUDA_cuLaunchKernel_ptsz ||
+			cbid == API_CUDA_cuLaunchKernel) {
+		cuLaunchKernel_params *p = (cuLaunchKernel_params *)params;
+
+		if (!is_exit) {
+
+
+			if (mkdir("traces", S_IRWXU | S_IRWXG | S_IROTH | S_IXOTH) == -1) {
+				if( errno == EEXIST ) {
+					// alredy exists
+				} else {
+					// something else
+					std::cout << "cannot create folder error:" << strerror(errno) << std::endl;
+					return;
+				}
+			}
+
+			int nregs;
+			_cuda_safe(
+					cuFuncGetAttribute(&nregs, CU_FUNC_ATTRIBUTE_NUM_REGS, p->f));
+
+			int shmem_static_nbytes;
+			_cuda_safe(cuFuncGetAttribute(&shmem_static_nbytes,
+					CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES,
+					p->f));
+
+
+
+			std::string func_name(nvbit_get_func_name(ctx, p->f));
+			std::string::size_type end_pos = func_name.find('(');
+			if (end_pos != std::string::npos)
+			{
+				// std::string::size_type pos = func_name.find('<');
+				//if (pos != std::string::npos)
+				//	end_pos = pos;
+
+				//std::string::size_type start_pos = func_name.find(' ');
+				//if (start_pos == std::string::npos)
+				//	start_pos = 0;
+				//else
+				//	start_pos++;
+
+				func_name = func_name.substr(0, end_pos);
+			}
+
+			char buffer[1024];
+			sprintf (buffer, "./traces/%d-%s.trace", kernelid, func_name.c_str());
+
+			resultsFile = fopen(buffer, "w");
+
+			printf("Writing results to %s\n", buffer);
+
+			fprintf(resultsFile, "-kernel name = %s",  nvbit_get_func_name(ctx, p->f));
+			fprintf(resultsFile, "\n");
+			fprintf(resultsFile, "-kernel id = %d",  kernelid);
+			fprintf(resultsFile, "\n");
+			fprintf(resultsFile, "-grid dim = (%d,%d,%d)",  p->gridDimX, p->gridDimY, p->gridDimZ);
+			fprintf(resultsFile, "\n");
+			fprintf(resultsFile, "-block dim = (%d,%d,%d)",  p->blockDimX, p->blockDimY, p->blockDimZ);
+			fprintf(resultsFile, "\n");
+			fprintf(resultsFile, "-shmem = %d",   shmem_static_nbytes + p->sharedMemBytes);
+			fprintf(resultsFile, "\n");
+			fprintf(resultsFile, "-nregs = %d",   nregs);
+			fprintf(resultsFile, "\n");
+			fprintf(resultsFile, "-cuda stream id = %d",  (uint64_t)p->hStream);
+			fprintf(resultsFile, "\n\n");
+
+			fprintf(resultsFile, "#traces format = threadblock_x threadblock_y threadblock_z warpid_tb sm_id warpid_sm PC mask dest_num reg_dests opcode src_num reg_srcs mem_width mem_addresses");
+			fprintf(resultsFile, "\n");
+
+			if (kernelid == 1) {
+				kernelsFile = fopen("./traces/kernelslist", "w");
+				statsFile = fopen("./traces/stats.csv", "w");
+				fprintf(statsFile, "kernel name,total_insts,total_reported_insts\n");
+			}
+			else {
+				kernelsFile = fopen("./traces/kernelslist", "a");
+				statsFile = fopen("./traces/stats.csv", "a");
+			}
+
+			sprintf (buffer, "%d-%s.trace", kernelid, func_name.c_str());
+			fprintf(kernelsFile, buffer);
+			fprintf(kernelsFile, "\n");
+			fclose(kernelsFile);
+
+			fprintf(statsFile, buffer);
+			fprintf(statsFile, ",");
+
+			kernelid++;
+			recv_thread_receiving = true;
+
+		} else {
+			/* make sure current kernel is completed */
+			cudaDeviceSynchronize();
+			assert(cudaGetLastError() == cudaSuccess);
+
+			/* make sure we prevent re-entry on the nvbit_callback when issuing
+			 * the flush_channel kernel */
+			skip_flag = true;
+
+			/* issue flush of channel so we are sure all the memory accesses
+			 * have been pushed */
+			flush_channel<<<1, 1>>>();
+			cudaDeviceSynchronize();
+			assert(cudaGetLastError() == cudaSuccess);
+
+			/* unset the skip flag */
+			skip_flag = false;
+
+			/* wait here until the receiving thread has not finished with the
+			 * current kernel */
+			while (recv_thread_receiving) {
+				pthread_yield();
+			}
+
+			unsigned total_insts_per_kernel =  total_dynamic_instr_counter - old_total_insts;
+			old_total_insts = total_dynamic_instr_counter;
+
+			unsigned reported_insts_per_kernel =  reported_dynamic_instr_counter - old_total_reported_insts;
+			old_total_reported_insts = reported_dynamic_instr_counter;
+
+			fprintf(statsFile, "");
+			fprintf(statsFile, "%d,%d",total_insts_per_kernel,reported_insts_per_kernel);
+			fprintf(statsFile, "\n");
+
+
+			fclose(resultsFile);
+			fclose(statsFile);
+		}
+	}
+}
+
+bool is_number(const std::string& s)
+{
+	std::string::const_iterator it = s.begin();
+	while (it != s.end() && std::isdigit(*it)) ++it;
+	return !s.empty() && it == s.end();
+}
+
+void *recv_thread_fun(void *) {
+	char *recv_buffer = (char *)malloc(CHANNEL_SIZE);
+
+	while (recv_thread_started) {
+		uint32_t num_recv_bytes = 0;
+		if (recv_thread_receiving &&
+				(num_recv_bytes = channel_host.recv(recv_buffer, CHANNEL_SIZE)) >
+		0) {
+			uint32_t num_processed_bytes = 0;
+			while (num_processed_bytes < num_recv_bytes) {
+				mem_access_t *ma =
+						(mem_access_t *)&recv_buffer[num_processed_bytes];
+
+				/* when we get this cta_id_x it means the kernel has completed
+				 */
+				if (ma->cta_id_x == -1) {
+					recv_thread_receiving = false;
+					break;
+				}
+
+				fprintf(resultsFile, "%d ", ma->cta_id_x);
+				fprintf(resultsFile, "%d ", ma->cta_id_y);
+				fprintf(resultsFile, "%d ", ma->cta_id_z);
+				fprintf(resultsFile, "%d ", ma->warpid_tb);
+				fprintf(resultsFile, "%d ", ma->sm_id);
+				fprintf(resultsFile, "%d ", ma->warpid_sm);
+				fprintf(resultsFile, "0x%016lx ", ma->vpc); // Print the virtual PC.
+				fprintf(resultsFile, "%-8.8" PRIx32 " ", ma->active_mask);
+				if(ma->GPRDst >= 0) {
+					fprintf(resultsFile, "1 ");
+					fprintf(resultsFile, "R%d ", ma->GPRDst);
+				}
+				else
+					fprintf(resultsFile, "0 ");
+
+					// Print the opcode.
+					fprintf(resultsFile, "%s ", id_to_opcode_map[ma->opcode_id].c_str());
+				unsigned src_count=0;
+				for (int s = 0; s < MAX_SRC; s++)      // GPR srcs count.
+					if(ma->GPRSrcs[s] >= 0)  src_count++;
+				fprintf(resultsFile, "%d ", src_count);
+
+				for (int s = 0; s < MAX_SRC; s++)      // GPR srcs.
+					if(ma->GPRSrcs[s] >= 0)  fprintf(resultsFile, "R%d ", ma->GPRSrcs[s]);
+
+				//print addresses
+				std::bitset<32> mask(ma->active_mask);
+				if(ma->is_mem) {
+					//fprintf(resultsFile, "%d ", ma->width);
+					std::istringstream iss(id_to_opcode_map[ma->opcode_id]);
+					std::vector<std::string> tokens;
+					std::string token;
+					while (std::getline(iss, token, '.')) {
+						if (!token.empty())
+							tokens.push_back(token);
+					}
+					if (tokens.size()>=3){
+						if (is_number(tokens[2])){
+							fprintf(resultsFile, "%d ", (std::stoi(tokens[2],nullptr)/8));
+						}
+						else{
+							fprintf(resultsFile, "%d ", 4);
+						}
+					}
+					else{
+						fprintf(resultsFile, "%d ", 4);
+					}
+
+					for (int s = 0; s < 32; s++)
+						if(mask.test(s))
+							fprintf(resultsFile, "0x%016lx ", ma->addrs[s]);
+				}
+				else
+				{
+					fprintf(resultsFile, "0 ");
+				}
+
+				fprintf(resultsFile, "\n");
+
+				num_processed_bytes += sizeof(mem_access_t);
+			}
+		}
+	}
+	free(recv_buffer);
+	return NULL;
+}
+
+void nvbit_at_ctx_init(CUcontext ctx) {
+	recv_thread_started = true;
+	channel_host.init(0, CHANNEL_SIZE, &channel_dev, NULL);
+	pthread_create(&recv_thread, NULL, recv_thread_fun, NULL);
+}
+
+void nvbit_at_ctx_term(CUcontext ctx) {
+	if (recv_thread_started) {
+		recv_thread_started = false;
+		pthread_join(recv_thread, NULL);
+	}
+}
diff --git a/src/trace-driven/traces-generator/README b/src/trace-driven/traces-generator/README
new file mode 100644
index 0000000..d79d075
--- /dev/null
+++ b/src/trace-driven/traces-generator/README
@@ -0,0 +1 @@
+TO DO
diff --git a/src/trace-driven/traces-generator/traces-post-processing/Makefile b/src/trace-driven/traces-generator/traces-post-processing/Makefile
new file mode 100755
index 0000000..63c3f38
--- /dev/null
+++ b/src/trace-driven/traces-generator/traces-post-processing/Makefile
@@ -0,0 +1,11 @@
+TARGET := post-traces-processing
+
+$(TARGET): post-traces-processing.cpp
+	g++ -o $@ $^
+
+run: $(TARGET)
+	./$(TARGET)
+
+clean: 
+	rm -f $(TARGET) *.o
+
diff --git a/src/trace-driven/traces-generator/traces-post-processing/post-traces-processing.cpp b/src/trace-driven/traces-generator/traces-post-processing/post-traces-processing.cpp
new file mode 100644
index 0000000..0dbb4e1
--- /dev/null
+++ b/src/trace-driven/traces-generator/traces-post-processing/post-traces-processing.cpp
@@ -0,0 +1,195 @@
+#include <vector>

+#include <iostream>

+#include <fstream>

+#include <string>

+#include <sstream>

+#include <stdio.h>

+#include <math.h>   

+using namespace std;

+

+struct threadblock_info

+{

+	bool initialized;

+	unsigned tb_id_x, tb_id_y, tb_id_z;

+	vector< vector< string > > warp_insts_array;

+	threadblock_info() {

+		initialized = false;

+		tb_id_x = tb_id_y = tb_id_z = 0;

+	}

+};

+

+void group_per_block(const char* filepath);

+void group_per_core(const char* filepath);

+

+//////////////////////////////////////////////////////////////////////////////////////////////////////////////////

+int main(int argc, char** argv)

+{ 

+

+	string kernellist_filepath;

+	bool is_per_core;

+	if(argc == 1)

+	{

+		cout << "File path is missing\n";

+		return 0;

+	} else if(argc == 2)

+	{

+		kernellist_filepath = argv[1];

+		is_per_core = true;

+

+	} else if(argc == 3) {

+		kernellist_filepath = argv[1];

+		is_per_core = bool(argv[2]);

+	}

+	else {

+		cout << "Too Many Arguemnts!\n";

+		return 0;

+	}

+

+	ifstream ifs;

+	ofstream ofs;

+

+	ifs.open(kernellist_filepath.c_str());

+	ofs.open((string(kernellist_filepath) + ".g").c_str());

+

+	if (!ifs.is_open()) {

+		cout << "Unable to open file: " <<kernellist_filepath<<endl;

+		return 0;

+	}

+

+

+	string directory(kernellist_filepath);

+	const size_t last_slash_idx = directory.rfind('/');

+	if (std::string::npos != last_slash_idx)

+	{

+		directory = directory.substr(0, last_slash_idx);

+	}

+

+	string line;

+	string filepath;

+	while(!ifs.eof()) {

+		getline(ifs, line);

+		if(line.empty())

+			continue;

+		filepath = directory+"/"+line;

+		group_per_block(filepath.c_str());

+		ofs<<line + "g"<<endl;

+	}

+

+	ifs.close();

+	ofs.close();

+	return 0;

+}

+

+

+void group_per_block(const char* filepath) {

+

+	ofstream ofs;

+	ifstream ifs;

+

+	ifs.open(filepath);

+

+	if (!ifs.is_open()) {

+		cout << "Unable to open file: " <<filepath<<endl; 

+		return;

+	}

+

+	cout << "Processing file " <<filepath<<endl;

+	ofs.open((string(filepath) + "g").c_str());

+

+	vector<threadblock_info> insts;

+	unsigned grid_dim_x, grid_dim_y, grid_dim_z, tb_dim_x, tb_dim_y, tb_dim_z;

+	unsigned tb_id_x, tb_id_y, tb_id_z, tb_id, warpid_tb;

+	string line;

+	stringstream ss;

+	string string1, string2;

+	bool found_grid_dim = false, found_block_dim = false;

+

+	while(!ifs.eof()) {

+		getline(ifs, line);

+

+		if (line.length() == 0 || line[0] == '#') {

+			ofs<<line<<endl;

+			continue;

+		}

+

+		else if(line[0] == '-') {

+			ss.str(line);

+			ss.ignore();

+			ss>>string1>>string2;

+			if(string1 == "grid" && string2 == "dim") {

+				sscanf(line.c_str(), "-grid dim = (%d,%d,%d)", &grid_dim_x, &grid_dim_y, &grid_dim_z);

+				found_grid_dim = true;

+			}

+			else if (string1 == "block" && string2 == "dim") {

+				sscanf(line.c_str(), "-block dim = (%d,%d,%d)", &tb_dim_x, &tb_dim_y, &tb_dim_z);

+				found_block_dim = true;

+			}

+

+			if(found_grid_dim && found_block_dim) {

+				insts.resize(grid_dim_x*grid_dim_y*grid_dim_z);

+				for(unsigned i = 0; i<insts.size(); ++i) {

+					insts[i].warp_insts_array.resize(ceil(float(tb_dim_x*tb_dim_y*tb_dim_z)/32));

+				}

+			}

+			ofs<<line<<endl;

+			continue;

+		}

+		else {

+

+			ss.str(line);

+			ss>>tb_id_x>>tb_id_y>>tb_id_z>>warpid_tb;

+			tb_id = tb_id_z * grid_dim_y * grid_dim_x + tb_id_y * grid_dim_x + tb_id_x;

+			if(!insts[tb_id].initialized) {

+				insts[tb_id].tb_id_x = tb_id_x;

+				insts[tb_id].tb_id_y = tb_id_y;

+				insts[tb_id].tb_id_z = tb_id_z;

+				insts[tb_id].initialized = true;

+			}

+			insts[tb_id].warp_insts_array[warpid_tb].push_back(line);

+		}

+

+	}  

+

+

+	for(unsigned i=0; i<insts.size(); ++i) {

+		//ofs<<string<<endl;

+		if(insts[i].initialized && insts[i].warp_insts_array.size() > 0) {

+			ofs<<endl<<"#BEGIN_TB"<<endl;

+			ofs<<endl<<"thread block = "<<insts[i].tb_id_x<<","<<insts[i].tb_id_y<<","<<insts[i].tb_id_z<<endl;

+		}

+		else {

+			cout<<"Warning: Thread block "<<insts[i].tb_id_x<<","<<insts[i].tb_id_y<<","<<insts[i].tb_id_z<<" is empty"<<endl;

+			continue;

+			//ofs.close();

+			//return;

+		}

+		for(unsigned j=0; j<insts[i].warp_insts_array.size(); ++j) {

+			ofs<<endl<<"warp = "<<j<<endl;

+			ofs<<"insts = "<<insts[i].warp_insts_array[j].size()<<endl;

+			if(insts[i].warp_insts_array[j].size() == 0) {

+				cout<<"Warning: Warp "<<j<<" in thread block"<<insts[i].tb_id_x<<","<<insts[i].tb_id_y<<","<<insts[i].tb_id_z<<" is empty"<<endl;

+				//	ofs.close();

+				//	return;

+			}

+			for(unsigned k=0; k<insts[i].warp_insts_array[j].size(); ++k) {

+				ofs<<insts[i].warp_insts_array[j][k]<<endl;

+			}

+		}

+		ofs<<endl<<"#END_TB"<<endl;

+	}

+

+

+

+	ofs.close();

+	ifs.close();

+

+}

+

+

+void group_per_core(const char* filepath)

+{

+

+	//TO DO

+

+

+}