Merge branch 'dev' of github.com:purdue-aalp/gpgpu-sim_distribution into dev

50 files changed, 1479 insertions, 404 deletions

diff --git a/aerialvision/organizedata.py b/aerialvision/organizedata.py
index 0d6cd91..090b90f 100644
--- a/aerialvision/organizedata.py
+++ b/aerialvision/organizedata.py
@@ -97,7 +97,7 @@ def organizedata(fileVars):
         'sparse':OrganizeSparse,        # Vector data with 2D index  (used by DRAM access stats)
         'custom':0
     }
-       data_type_char = {int:'I', float:'f'}
+    data_type_char = {int:'I', float:'f'}
 
     print "Organizing data into internal format..."
 
diff --git a/configs/deprecated-cfgs/SM6_GTX1080/gpgpusim.config b/configs/deprecated-cfgs/SM6_GTX1080/gpgpusim.config
index fb044c6..3261d5a 100644
--- a/configs/deprecated-cfgs/SM6_GTX1080/gpgpusim.config
+++ b/configs/deprecated-cfgs/SM6_GTX1080/gpgpusim.config
@@ -31,12 +31,14 @@
 -gpgpu_simd_model 1 
 
 # Pipeline widths and number of FUs
-# ID_OC_SP,ID_OC_SFU,ID_OC_MEM,OC_EX_SP,OC_EX_SFU,OC_EX_MEM,EX_WB
+# ID_OC_SP, ID_OC_DP, ID_OC_INT, ID_OC_SFU, ID_OC_MEM, OC_EX_SP, OC_EX_DP, OC_EX_INT, OC_EX_SFU, OC_EX_MEM, EX_WB
 ## Pascal GP102 has 4 SP SIMD units and 1 SFU unit
 ## we need to scale the number of pipeline registers to be equal to the number of SP units
--gpgpu_pipeline_widths 4,1,1,4,1,1,6
+-gpgpu_pipeline_widths 4,0,0,1,1,4,0,0,1,1,6
 -gpgpu_num_sp_units 4
 -gpgpu_num_sfu_units 1
+-gpgpu_tensor_core_avail 0
+-gpgpu_num_tensor_core_units 0
 
 # Instruction latencies and initiation intervals
 # "ADD,MAX,MUL,MAD,DIV"
@@ -54,20 +56,20 @@
 # Pascal GP102 has 96KB Shared memory
 # Pascal GP102 has 64KB L1 cache
 # The default is to disable the L1 cache, unless cache modifieres is used
--gpgpu_cache:dl1  64:128:6,L:L:m:N:H,A:128:8,8
+-gpgpu_cache:dl1  N:64:128:6,L:L:m:N:H,A:128:8,8
 -gpgpu_shmem_size 98304
 -gmem_skip_L1D 1
 
 # 64 sets, each 128 bytes 16-way for each memory sub partition (128 KB per memory sub partition). This gives 3MB L2 cache
--gpgpu_cache:dl2 64:128:16,L:B:m:W:L,A:1024:1024,4:0,32 # used to be 128:4
+-gpgpu_cache:dl2 N:64:128:16,L:B:m:W:L,A:1024:1024,4:0,32 # used to be 128:4
 -gpgpu_cache:dl2_texture_only 0 
 
 # 4 KB Inst.
--gpgpu_cache:il1 8:128:4,L:R:f:N:L,A:2:48,4
+-gpgpu_cache:il1 N:8:128:4,L:R:f:N:L,A:2:48,4
 # 48 KB Tex 
--gpgpu_tex_cache:l1 16:128:24,L:R:m:N:L,F:128:4,128:2
+-gpgpu_tex_cache:l1 N:16:128:24,L:R:m:N:L,F:128:4,128:2
 # 12 KB Const
--gpgpu_const_cache:l1 128:64:2,L:R:f:N:L,A:2:64,4
+-gpgpu_const_cache:l1 N:128:64:2,L:R:f:N:L,A:2:64,4
 
 # enable operand collector 
 ## larger operand collectors and reg_banks are needed for the 4 warp schedulers and 4 SIMD units
diff --git a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config
index d71b2fd..cf3627b 100644
--- a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config
+++ b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config
@@ -3,6 +3,13 @@
 -gpgpu_ptx_sim_mode 0
 -gpgpu_ptx_force_max_capability 20 
 
+# Device Limits
+-gpgpu_stack_size_limit 1024
+-gpgpu_heap_size_limit 8388608
+
+# Compute Capability
+-gpgpu_compute_capability_major 2
+-gpgpu_compute_capability_minor 0
 
 # SASS execution (only supported with CUDA >= 4.0)
 -gpgpu_ptx_convert_to_ptxplus 0
diff --git a/configs/tested-cfgs/SM6_TITANX/config_fermi_islip.icnt b/configs/tested-cfgs/SM6_TITANX/config_pascal_islip.icnt
index dec4789..dec4789 100644
--- a/configs/tested-cfgs/SM6_TITANX/config_fermi_islip.icnt
+++ b/configs/tested-cfgs/SM6_TITANX/config_pascal_islip.icnt
diff --git a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config
index cb23ab3..2fe898a 100644
--- a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config
+++ b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config
@@ -8,6 +8,16 @@
 -gpgpu_ptx_force_max_capability 61 
 -gpgpu_ignore_resources_limitation 1
 
+# Device Limits
+-gpgpu_stack_size_limit 1024
+-gpgpu_heap_size_limit 8388608
+-gpgpu_runtime_sync_depth_limit 2
+-gpgpu_runtime_pending_launch_count_limit 2048
+
+# Compute Capability
+-gpgpu_compute_capability_major 6
+-gpgpu_compute_capability_minor 1
+
 # SASS execution (only supported with CUDA >= 4.0)
 -gpgpu_ptx_convert_to_ptxplus 0
 -gpgpu_ptx_save_converted_ptxplus 0
@@ -92,7 +102,7 @@
 # 4 KB Inst.
 -gpgpu_cache:il1 N:8:128:4,L:R:f:N:L,S:2:48,4
 # 48 KB Tex 
-# Note, TEX is deprected in Volta, It is used for legacy apps only. Use L1D cache instead with .nc modifier or __ldg mehtod
+# Note, TEX is deprected in Pascal, It is used for legacy apps only. Use L1D cache instead with .nc modifier or __ldg mehtod
 -gpgpu_tex_cache:l1 N:16:128:24,L:R:m:N:L,T:128:4,128:2
 # 12 KB Const
 -gpgpu_const_cache:l1 N:128:64:2,L:R:f:N:L,S:2:64,4
diff --git a/configs/tested-cfgs/SM7_TITANV/config_fermi_islip.icnt b/configs/tested-cfgs/SM7_TITANV/config_volta_islip.icnt
index 615d0a9..615d0a9 100644
--- a/configs/tested-cfgs/SM7_TITANV/config_fermi_islip.icnt
+++ b/configs/tested-cfgs/SM7_TITANV/config_volta_islip.icnt
diff --git a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
index 8ed4cd0..ebd442f 100644
--- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
+++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
@@ -1,4 +1,4 @@
-# This config models the Volta Titan X
+# This config models the Volta Titan V
 # For more info about volta architecture:
 # http://images.nvidia.com/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf
 # https://ieeexplore.ieee.org/stamp/stamp.jsp?tp=&arnumber=8344474&tag=1# 
@@ -13,6 +13,17 @@
 -gpgpu_ptx_sim_mode 0
 -gpgpu_ptx_force_max_capability 70 
 
+
+# Device Limits
+-gpgpu_stack_size_limit 1024
+-gpgpu_heap_size_limit 8388608
+-gpgpu_runtime_sync_depth_limit 2
+-gpgpu_runtime_pending_launch_count_limit 2048
+
+# Compute Capability
+-gpgpu_compute_capability_major 7
+-gpgpu_compute_capability_minor 0
+
 # SASS execution (only supported with CUDA >= 4.0)
 -gpgpu_ptx_convert_to_ptxplus 0
 -gpgpu_ptx_save_converted_ptxplus 0
@@ -25,7 +36,7 @@
 
 # volta clock domains
 #-gpgpu_clock_domains <Core Clock>:<Interconnect Clock>:<L2 Clock>:<DRAM Clock>
-# Volta NVIDIA V100  clock domains are adopted from 
+# Volta NVIDIA TITANV  clock domains are adopted from 
 # https://en.wikipedia.org/wiki/Volta_(microarchitecture)
 -gpgpu_clock_domains 1200.0:1200.0:1200.0:850.0
 # boost mode
@@ -42,7 +53,7 @@
 
 # Pipeline widths and number of FUs
 # ID_OC_SP,ID_OC_DP,ID_OC_INT,ID_OC_SFU,ID_OC_MEM,OC_EX_SP,OC_EX_DP,OC_EX_INT,OC_EX_SFU,OC_EX_MEM,EX_WB,ID_OC_TENSOR_CORE,OC_EX_TENSOR_CORE
-## Volta GV100 has 4 SP SIMD units, 4 SFU units, 4 DP units per core
+## Volta TITANV has 4 SP SIMD units, 4 INT units, 4 SFU units, 4 DP units per core, 4 Tensor core units
 ## we need to scale the number of pipeline registers to be equal to the number of SP units
 -gpgpu_pipeline_widths 4,4,4,4,4,4,4,4,4,4,8,4,4
 -gpgpu_num_sp_units 4
@@ -75,7 +86,7 @@
 # if the assigned shd mem = 0, then L1 cache = 128KB
 # For more info, see https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory-7-x 
 # disable this mode in case of multi kernels/apps execution
--adpative_volta_cache_config 1
+-adaptive_volta_cache_config 1
 # Volta unified cache has four ports
 -mem_unit_ports 4
 -gpgpu_cache:dl1  S:4:128:64,L:L:s:N:L,A:256:8,16:0,32
@@ -102,9 +113,9 @@
 # 64 KB Const
 -gpgpu_const_cache:l1 N:128:64:8,L:R:f:N:L,S:2:64,4
 
-# Volta has sub core model, in which each scheduler has its own reisiter file and EUs
+# Volta has sub core model, in which each scheduler has its own register file and EUs
 # i.e. schedulers are isolated
--sub_core_model 0
+-sub_core_model 1
 # disable specialized operand collectors and use generic operand collectors instead
 -enable_specialized_operand_collector 0
 -gpgpu_operand_collector_num_units_gen 8
@@ -125,7 +136,7 @@
 
 # interconnection
 -network_mode 1 
--inter_config_file config_fermi_islip.icnt
+-inter_config_file config_volta_islip.icnt
 
 # memory partition latency config 
 -rop_latency 120
@@ -150,7 +161,7 @@
 #-gpgpu_dram_timing_opt "nbk=16:CCD=1:RRD=4:RCD=14:RAS=33:RP=14:RC=47:
 #                        CL=14:WL=2:CDLR=3:WR=12:nbkgrp=4:CCDL=2:RTPL=4"
 
-# Timing for 850 MHZ, Tesla TITANV V100 HBM runs at 850 MHZ
+# Timing for 850 MHZ, Tesla TITANV HBM runs at 850 MHZ
 -gpgpu_dram_timing_opt "nbk=16:CCD=1:RRD=3:RCD=12:RAS=28:RP=12:RC=40:
                         CL=12:WL=2:CDLR=3:WR=10:nbkgrp=4:CCDL=2:RTPL=3"
 
@@ -163,7 +174,7 @@
 #-Seperate_Write_Queue_Enable 1
 #-Write_Queue_Size 64:56:32
 
-# Pascal has two schedulers per core
+# Volta has four schedulers per core
 -gpgpu_num_sched_per_core 4
 # Two Level Scheduler with active and pending pools
 #-gpgpu_scheduler two_level_active:6:0:1
@@ -178,7 +189,7 @@
 -enable_ptx_file_line_stats 1
 -visualizer_enabled 0
 
-# power model configs, disable it untill we create a real energy model for Pascal 100
+# power model configs, disable it untill we create a real energy model for Volta
 -power_simulation_enabled 0
 
 # tracing functionality
diff --git a/configs/tested-cfgs/SM7_TITANV_fastlocalxbar/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV_fastlocalxbar/gpgpusim.config
new file mode 100644
index 0000000..4625ac4
--- /dev/null
+++ b/configs/tested-cfgs/SM7_TITANV_fastlocalxbar/gpgpusim.config
@@ -0,0 +1,190 @@
+# This config models the Volta Titan V
+# For more info about volta architecture:
+# http://images.nvidia.com/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf
+# https://ieeexplore.ieee.org/stamp/stamp.jsp?tp=&arnumber=8344474&tag=1# 
+# http://on-demand.gputechconf.com/gtc/2018/presentation/s8122-dissecting-the-volta-gpu-architecture-through-microbenchmarking.pdf
+# https://en.wikipedia.org/wiki/Volta_(microarchitecture)
+# https://www.hotchips.org/wp-content/uploads/hc_archives/hc29/HC29.21-Monday-Pub/HC29.21.10-GPU-Gaming-Pub/HC29.21.132-Volta-Choquette-NVIDIA-Final3.pdf
+# https://devblogs.nvidia.com/inside-volta/
+# http://on-demand.gputechconf.com/gtc/2017/presentation/s7798-luke-durant-inside-volta.pdf
+
+# functional simulator specification
+-gpgpu_ptx_instruction_classification 0
+-gpgpu_ptx_sim_mode 0
+-gpgpu_ptx_force_max_capability 70 
+
+# SASS execution (only supported with CUDA >= 4.0)
+-gpgpu_ptx_convert_to_ptxplus 0
+-gpgpu_ptx_save_converted_ptxplus 0
+
+# high level architecture configuration
+-gpgpu_n_clusters 40
+-gpgpu_n_cores_per_cluster 2
+-gpgpu_n_mem 24
+-gpgpu_n_sub_partition_per_mchannel 2 
+
+# volta clock domains
+#-gpgpu_clock_domains <Core Clock>:<Interconnect Clock>:<L2 Clock>:<DRAM Clock>
+# Volta NVIDIA TITANV clock domains are adopted from 
+# https://en.wikipedia.org/wiki/Volta_(microarchitecture)
+-gpgpu_clock_domains 1200.0:1200.0:1200.0:850.0
+# boost mode
+# -gpgpu_clock_domains 1455.0:1455.0:1455.0:850.0
+
+# shader core pipeline config
+-gpgpu_shader_registers 65536
+-gpgpu_occupancy_sm_number 70
+
+# This implies a maximum of 64 warps/SM
+-gpgpu_shader_core_pipeline 2048:32 
+-gpgpu_shader_cta 32
+-gpgpu_simd_model 1 
+
+# Pipeline widths and number of FUs
+# ID_OC_SP,ID_OC_DP,ID_OC_INT,ID_OC_SFU,ID_OC_MEM,OC_EX_SP,OC_EX_DP,OC_EX_INT,OC_EX_SFU,OC_EX_MEM,EX_WB,ID_OC_TENSOR_CORE,OC_EX_TENSOR_CORE
+## Volta TITANV has 4 SP SIMD units, 4 SFU units, 4 DP units per core, 4 Tensor core units
+## we need to scale the number of pipeline registers to be equal to the number of SP units
+-gpgpu_pipeline_widths 4,4,4,4,4,4,4,4,4,4,8,4,4
+-gpgpu_num_sp_units 4
+-gpgpu_num_sfu_units 4
+-gpgpu_num_dp_units 4
+-gpgpu_num_int_units 4
+-gpgpu_tensor_core_avail 1
+-gpgpu_num_tensor_core_units 4
+
+# Instruction latencies and initiation intervals
+# "ADD,MAX,MUL,MAD,DIV"
+# All Div operations are executed on SFU unit
+# Throughput (initiation latency) are adopted from 
+# http://on-demand.gputechconf.com/gtc/2018/presentation/s8122-dissecting-the-volta-gpu-architecture-through-microbenchmarking.pdf
+-ptx_opcode_latency_int 4,13,4,5,145
+-ptx_opcode_initiation_int 2,2,2,2,8
+-ptx_opcode_latency_fp 4,13,4,5,39
+-ptx_opcode_initiation_fp 2,2,2,2,4
+-ptx_opcode_latency_dp 8,19,8,8,330
+-ptx_opcode_initiation_dp 4,4,4,4,130
+-ptx_opcode_latency_sfu 100
+-ptx_opcode_initiation_sfu 8
+-ptx_opcode_latency_tesnor 64
+-ptx_opcode_initiation_tensor 64
+
+# <nsets>:<bsize>:<assoc>,<rep>:<wr>:<alloc>:<wr_alloc>:<set_index_fn>,<mshr>:<N>:<merge>,<mq>:**<fifo_entry>
+# ** Optional parameter - Required when mshr_type==Texture Fifo
+# Defualt config is 32KB DL1 and 96KB shared memory
+# In Volta, we assign the remaining shared memory to L1 cache 
+# if the assigned shd mem = 0, then L1 cache = 128KB
+# For more info, see https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory-7-x 
+# disable this mode in case of multi kernels/apps execution
+-adaptive_volta_cache_config 1
+# Volta unified cache has four ports
+-mem_unit_ports 4
+-gpgpu_cache:dl1  S:4:128:64,L:L:s:N:L,A:256:8,16:0,32
+-gpgpu_shmem_size 98304
+-gmem_skip_L1D 0
+-icnt_flit_size 40
+-gpgpu_n_cluster_ejection_buffer_size 32
+-l1_latency 28
+-smem_latency 19
+-gpgpu_flush_l1_cache 1
+
+# 32 sets, each 128 bytes 24-way for each memory sub partition (96 KB per memory sub partition). This gives us 4.5MB L2 cache
+-gpgpu_cache:dl2 S:32:128:24,L:B:m:L:L,A:192:4,32:0,32
+-gpgpu_cache:dl2_texture_only 0
+-gpgpu_dram_partition_queues 64:64:64:64
+-perf_sim_memcpy 1
+-memory_partition_indexing 0
+
+# 128 KB Inst.
+-gpgpu_cache:il1 N:64:128:16,L:R:f:N:L,S:2:48,4
+# 48 KB Tex
+# Note, TEX is deprected in Volta, It is used for legacy apps only. Use L1D cache instead with .nc modifier or __ldg mehtod
+-gpgpu_tex_cache:l1 N:16:128:24,L:R:m:N:L,T:128:4,128:2
+# 64 KB Const
+-gpgpu_const_cache:l1 N:128:64:8,L:R:f:N:L,S:2:64,4
+
+# Volta has sub core model, in which each scheduler has its own register file and EUs
+# i.e. schedulers are isolated
+-sub_core_model 1
+# disable specialized operand collectors and use generic operand collectors instead
+-enable_specialized_operand_collector 0
+-gpgpu_operand_collector_num_units_gen 8
+-gpgpu_operand_collector_num_in_ports_gen 8
+-gpgpu_operand_collector_num_out_ports_gen 8
+# volta has 8 banks, 4 schedulers, two banks per scheduler
+-gpgpu_num_reg_banks 8
+
+# shared memory bankconflict detection 
+-gpgpu_shmem_num_banks 32
+-gpgpu_shmem_limited_broadcast 0
+-gpgpu_shmem_warp_parts 1
+-gpgpu_coalesce_arch 60
+
+## In Volta, a warp scheduler can issue 1 inst per cycle
+-gpgpu_max_insn_issue_per_warp 1
+-gpgpu_dual_issue_diff_exec_units 1
+
+# use local xbar interconnection
+-network_mode 2 
+-inct_in_buffer_limit 512
+-inct_out_buffer_limit 512
+-inct_subnets 2
+
+# memory partition latency config 
+-rop_latency 120
+-dram_latency 100
+
+# dram model config
+-gpgpu_dram_scheduler 1
+-gpgpu_frfcfs_dram_sched_queue_size 64
+-gpgpu_dram_return_queue_size 192
+
+# for HBM, three stacks, 24 channles, each (128 bits) 16 bytes width
+-gpgpu_n_mem_per_ctrlr 1
+-gpgpu_dram_buswidth 16
+-gpgpu_dram_burst_length 2
+-dram_data_command_freq_ratio 2  # HBM is DDR
+-gpgpu_mem_address_mask 1
+-gpgpu_mem_addr_mapping dramid@8;00000000.00000000.00000000.00000000.0000RRRR.RRRRRRRR.RBBBCCCB.CCCSSSSS
+
+# HBM timing are adopted from hynix JESD235 standered and nVidia HPCA 2017 paper (http://www.cs.utah.edu/~nil/pubs/hpca17.pdf)
+# Timing for 1 GHZ
+# tRRDl and tWTR are missing, need to be added
+#-gpgpu_dram_timing_opt "nbk=16:CCD=1:RRD=4:RCD=14:RAS=33:RP=14:RC=47:
+#                        CL=14:WL=2:CDLR=3:WR=12:nbkgrp=4:CCDL=2:RTPL=4"
+
+# Timing for 850 MHZ, Tesla TITANV V100 HBM runs at 850 MHZ
+-gpgpu_dram_timing_opt "nbk=16:CCD=1:RRD=3:RCD=12:RAS=28:RP=12:RC=40:
+                        CL=12:WL=2:CDLR=3:WR=10:nbkgrp=4:CCDL=2:RTPL=3"
+
+# HBM has dual bus interface, in which it can issue two col and row commands at a time
+-dual_bus_interface 1
+# select lower bits for bnkgrp to increase bnkgrp parallelism
+-dram_bnk_indexing_policy 0
+-dram_bnkgrp_indexing_policy 1
+
+#-Seperate_Write_Queue_Enable 1
+#-Write_Queue_Size 64:56:32
+
+# Volta has four schedulers per core
+-gpgpu_num_sched_per_core 4
+# Two Level Scheduler with active and pending pools
+#-gpgpu_scheduler two_level_active:6:0:1
+# Loose round robbin scheduler
+#-gpgpu_scheduler lrr
+# Greedy then oldest scheduler
+-gpgpu_scheduler gto
+
+# stat collection
+-gpgpu_memlatency_stat 14 
+-gpgpu_runtime_stat 500
+-enable_ptx_file_line_stats 1
+-visualizer_enabled 0
+
+# power model configs, disable it untill we create a real energy model for Volta
+-power_simulation_enabled 0
+
+# tracing functionality
+#-trace_enabled 1
+#-trace_components WARP_SCHEDULER,SCOREBOARD
+#-trace_sampling_core 0
+
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 27644b3..11dcd5a 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -354,8 +354,8 @@ struct _cuda_device_id *GPGPUSim_Init()
 
 		cudaDeviceProp *prop = (cudaDeviceProp *) calloc(sizeof(cudaDeviceProp),1);
 		snprintf(prop->name,256,"GPGPU-Sim_v%s", g_gpgpusim_version_string );
-		prop->major = 5;
-		prop->minor = 2;
+		prop->major = the_gpu->compute_capability_major();
+		prop->minor = the_gpu->compute_capability_minor();
 		prop->totalGlobalMem = 0x80000000 /* 2 GB */;
 		prop->memPitch = 0;
 		if(prop->major >= 2) {
@@ -1120,11 +1120,11 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic
                 case 41:
                         *value= 0;
                         break;
-                case 75:
-                        *value= 9 ;
+                case 75://cudaDevAttrComputeCapabilityMajor
+                        *value= prop->major ;
                         break;
-                case 76:
-                        *value= 3 ;
+                case 76://cudaDevAttrComputeCapabilityMinor
+                        *value= prop->minor ;
                         break;
                 case 78:
                         *value= 0 ; //TODO: as of now, we dont support stream priorities.
@@ -1199,17 +1199,50 @@ __host__ cudaError_t CUDARTAPI cudaGetDevice(int *device)
 	*device = g_active_device;
 	return g_last_cudaError = cudaSuccess;
 }
+
 __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit limit )
 {
 	if(g_debug_execution >= 3){
             announce_call(__my_func__);
    	 }
-        cuda_not_implemented(__my_func__,__LINE__);
-        return g_last_cudaError = cudaSuccess;
+        _cuda_device_id *dev = GPGPUSim_Init();
+	const struct cudaDeviceProp *prop = dev->get_prop();
+	const gpgpu_sim_config& config=dev->get_gpgpu()->get_config();
+        switch(limit) {
+        case 0:  // cudaLimitStackSize
+        	*pValue=config.stack_limit();
+                break;
+        case 2:  // cudaLimitMallocHeapSize
+                *pValue=config.heap_limit();
+                break;
+#if (CUDART_VERSION > 5050)
+        case 3: // cudaLimitDevRuntimeSyncDepth
+		if(prop->major > 2){
+			*pValue=config.sync_depth_limit();
+	                break;
+		}
+		else{
+			printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+			abort();
+		}
+        case 4: // cudaLimitDevRuntimePendingLaunchCount
+		if(prop->major > 2){
+     	        	*pValue=config.pending_launch_count_limit();
+	                break;
+		}
+		else{
+			printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+			abort();
+		}
+#endif
+        default:
+                        printf("ERROR:Limit %s unimplemented \n",limit);
+                        abort();
+        }
+	return g_last_cudaError = cudaSuccess;
 
 }
 
-
 __host__ cudaError_t CUDARTAPI cudaStreamGetPriority ( cudaStream_t hStream, int* priority )
 {
         if(g_debug_execution >= 3){
@@ -1569,6 +1602,9 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
 	    announce_call(__my_func__);
     }
 #if (CUDART_VERSION >= 3000)
+	//per-stream synchronization required for application using external libraries without explicit synchronization in the code to 
+	//avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. 
+	stream->synchronize();
 	g_stream_manager->destroy_stream(stream);
 #endif
 	return g_last_cudaError = cudaSuccess;
@@ -1935,8 +1971,11 @@ void setCuobjdumpsassfilename(const char* filename){
 	}
 	(dynamic_cast<cuobjdumpELFSection*>(cuobjdumpSectionList.front()))->setSASSfilename(filename);
 }
-extern int cuobjdump_parse();
-extern FILE *cuobjdump_in;
+typedef void * yyscan_t;
+extern int cuobjdump_lex_init(yyscan_t* scanner);
+extern void cuobjdump_set_in  (FILE * _in_str ,yyscan_t yyscanner );
+extern int cuobjdump_parse(yyscan_t scanner);
+extern int cuobjdump_lex_destroy(yyscan_t scanner);
 
 //! Return the executable file of the process containing the PTX/SASS code 
 //!
@@ -2145,9 +2184,14 @@ void extract_code_using_cuobjdump(){
 
         if (parse_output) {
             printf("Parsing file %s\n", fname);
+	    FILE *cuobjdump_in;
             cuobjdump_in = fopen(fname, "r");
 
-            cuobjdump_parse();
+	    yyscan_t scanner;
+	    cuobjdump_lex_init(&scanner);
+	    cuobjdump_set_in(cuobjdump_in, scanner);
+	    cuobjdump_parse(scanner);
+	    cuobjdump_lex_destroy(scanner);
             fclose(cuobjdump_in);
             printf("Done parsing!!!\n");
         } else {
@@ -2195,8 +2239,13 @@ void extract_code_using_cuobjdump(){
                     std::cout << "Done" << std::endl;
 
                     std::cout << "Trying to parse " << libcodfn.str() << std::endl;
+		    FILE *cuobjdump_in;
                     cuobjdump_in = fopen(libcodfn.str().c_str(), "r");
-                    cuobjdump_parse();
+		    yyscan_t scanner;
+		    cuobjdump_lex_init(&scanner);
+		    cuobjdump_set_in(cuobjdump_in, scanner);
+		    cuobjdump_parse(scanner);
+		    cuobjdump_lex_destroy(scanner);
                     fclose(cuobjdump_in);
                     std::getline(libsf, line);
             }
@@ -2695,14 +2744,14 @@ cudaError_t cudaDeviceReset ( void ) {
 	return g_last_cudaError = cudaSuccess;
 }
 cudaError_t CUDARTAPI cudaDeviceSynchronize(void){
-	// I don't know what this should do
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
+	//Blocks until the device has completed all preceding requested tasks
+	synchronize();
 	return g_last_cudaError = cudaSuccess;
 }
 
-
 void CUDARTAPI __cudaRegisterFunction(
 		void   **fatCubinHandle,
 		const char    *hostFun,
@@ -3132,6 +3181,8 @@ __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t v
     }
     return g_last_cudaError = cudaSuccess;
 }
+
+
 #endif
 
 #endif
@@ -3310,7 +3361,12 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
 	    announce_call(__my_func__);
     }
 	function_info *entry = context->get_kernel(hostFun);
-	kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry);
+	gpgpu_t* gpu= context->get_device()->get_gpgpu();
+	/*
+	Passing a snapshot of the GPU's current texture mapping to the kernel's info
+	as kernels should use texture bindings present at the time of their launch.
+	*/
+	kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping());
 	if( entry == NULL ) {
 		printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found for %p\n", hostFun);
 		abort();
diff --git a/libcuda/cuobjdump.l b/libcuda/cuobjdump.l
index 2b0dac8..9359281 100644
--- a/libcuda/cuobjdump.l
+++ b/libcuda/cuobjdump.l
@@ -38,9 +38,7 @@
 
 #define YYDEBUG 1
 
-#define yylval cuobjdump_lval
-
-void cuobjdump_error(const char*);
+void cuobjdump_error(yyscan_t yyscanner, const char* msg);
 %}
 
 %option stack
@@ -48,6 +46,8 @@ void cuobjdump_error(const char*);
 %option yylineno
 %option nounput
 
+%option bison-bridge
+%option reentrant
 
 %s ptxcode
 %s sasscode
@@ -69,54 +69,54 @@ newlines	{newline}+
 
 
 "ptxasOptions"{notnewline}*{newline}
-[1-9]{numeric}* yylval.string_value = strdup(yytext); return DECIMAL;
+[1-9]{numeric}* yylval->string_value = strdup(yytext); return DECIMAL;
 
 "has debug info"{newline}
 
 
 "Fatbin ptx code:"{newline}	{
-	yy_push_state(ptxcode);
-	yy_push_state(identifier);
-	yy_push_state(ptxheader);
-	yylval.string_value = strdup(yytext);
+	yy_push_state(ptxcode, yyscanner);
+	yy_push_state(identifier, yyscanner);
+	yy_push_state(ptxheader, yyscanner);
+	yylval->string_value = strdup(yytext);
 	return PTXHEADER;
 }
 "Fatbin elf code:"{newline}	{
-	yy_push_state(elfcode);
-	yy_push_state(identifier);
-	yy_push_state(elfheader);
-	yylval.string_value = strdup(yytext);
+	yy_push_state(elfcode, yyscanner);
+	yy_push_state(identifier, yyscanner);
+	yy_push_state(elfheader, yyscanner);
+	yylval->string_value = strdup(yytext);
 	return ELFHEADER;
 }
 	/*PTX code tokens*/
-<ptxcode>{notnewline}*{newline}	yylval.string_value = strdup(yytext); return PTXLINE;
+<ptxcode>{notnewline}*{newline}	yylval->string_value = strdup(yytext); return PTXLINE;
 
 	/*ELF code tokens*/
 <elfcode>{whitespace}*"code for sm_"{numeric}+{newline}	{
 	BEGIN(sasscode); 
-	yylval.string_value = strdup(yytext); 
+	yylval->string_value = strdup(yytext);
 	return SASSLINE;
 }
-<elfcode>{notnewline}*{newline} yylval.string_value = strdup(yytext); return ELFLINE;
+<elfcode>{notnewline}*{newline} yylval->string_value = strdup(yytext); return ELFLINE;
 
 	/*SASS code tokens*/
-<sasscode>{notnewline}*{newline} yylval.string_value = strdup(yytext); return SASSLINE;
+<sasscode>{notnewline}*{newline} yylval->string_value = strdup(yytext); return SASSLINE;
 
 <identifier>{newline}"compressed"{newline}	BEGIN(conidentifier); return H_COMPRESSED;
 <identifier>{newline}"identifier = "	BEGIN(endidentifier); return H_IDENTIFIER;
-<identifier>{newline}{newline} yy_pop_state();
+<identifier>{newline}{newline} yy_pop_state(yyscanner);
 
 
 <conidentifier>"identifier = "	BEGIN(endidentifier); return H_IDENTIFIER;
-<conidentifier>{newline} yy_pop_state();
+<conidentifier>{newline} yy_pop_state(yyscanner);
 
 
-<endidentifier>{notnewline}+	yylval.string_value = strdup(yytext); yy_pop_state(); return FILENAME;
+<endidentifier>{notnewline}+	yylval->string_value = strdup(yytext); yy_pop_state(yyscanner); return FILENAME;
 
 
 
 	/*Header tokens*/
-<elfheader>[[:alnum:]_]+ yylval.string_value = strdup(yytext); return IDENTIFIER;
+<elfheader>[[:alnum:]_]+ yylval->string_value = strdup(yytext); return IDENTIFIER;
 <elfheader>"================"	return H_SEPARATOR;
 <elfheader>"arch = "		return H_ARCH;
 <elfheader>"code version = "	return H_CODEVERSION;
@@ -128,7 +128,7 @@ newlines	{newline}+
 
 
 	/*Header tokens*/
-<ptxheader>[[:alnum:]_]+ yylval.string_value = strdup(yytext); return IDENTIFIER;
+<ptxheader>[[:alnum:]_]+ yylval->string_value = strdup(yytext); return IDENTIFIER;
 <ptxheader>"================"	return H_SEPARATOR;
 <ptxheader>"arch = "		return H_ARCH;
 <ptxheader>"code version = "	return H_CODEVERSION;
@@ -143,7 +143,7 @@ newlines	{newline}+
 
 
 	/* Looking for the identifier (filename) then the header is done */
-<endheader>{notnewline}+	yylval.string_value = strdup(yytext); yy_pop_state(); return IDENTIFIER;
+<endheader>{notnewline}+	yylval->string_value = strdup(yytext); yy_pop_state(yyscanner); return IDENTIFIER;
 
 
 
@@ -153,11 +153,12 @@ newlines	{newline}+
 <<EOF>> BEGIN(INITIAL);return 0;
 
 	/*No other rule matched. Throw an error*/
-.	cuobjdump_error("Invalid token");
+.	cuobjdump_error(yyscanner, "Invalid token");
 
 %%
-void cuobjdump_error(const char* message)
+void cuobjdump_error(yyscan_t yyscanner, const char* msg)
 {
-	printf("   %s near \"%s\"",message, yytext);
-	printf(" on line %i\n",yylineno); 
+    struct yyguts_t * yyg = (struct yyguts_t*)yyscanner;
+	printf("   %s near \"%s\"",msg, yytext);
+	printf(" on line %i\n",yylineno);
 }
diff --git a/libcuda/cuobjdump.y b/libcuda/cuobjdump.y
index 31760f7..66cbace 100644
--- a/libcuda/cuobjdump.y
+++ b/libcuda/cuobjdump.y
@@ -29,8 +29,8 @@
 %{
 #include <stdio.h>
 
-int yylex(void);
-void yyerror(const char*);
+typedef void * yyscan_t;
+
 extern void addCuobjdumpSection(int sectiontype);
 void setCuobjdumparch(const char* arch);
 void setCuobjdumpidentifier(const char* identifier);
@@ -44,9 +44,17 @@ FILE *elffile;
 FILE *sassfile;
 char filename [1024];
 %}
+%define api.pure full
+%parse-param {yyscan_t scanner}
+%lex-param {yyscan_t scanner}
+
 %union {
 	char* string_value;
 }
+%{
+int yylex(YYSTYPE * yylval_param, yyscan_t yyscanner);
+void yyerror(yyscan_t yyscanner, const char* msg);
+%}
 %token <string_value> H_SEPARATOR H_ARCH H_CODEVERSION H_PRODUCER H_HOST H_COMPILESIZE H_IDENTIFIER H_UNKNOWN H_COMPRESSED
 %token <string_value> CODEVERSION
 %token <string_value> STRING
diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index 307df40..63b139e 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -88,8 +88,6 @@ void checkpoint::load_global_mem(class memory_space *temp_mem, char * f1name)
       fclose ( fp2 );
 }
 
-
-
 void checkpoint::store_global_mem(class memory_space * mem, char *fname, char * format)
 {
 
@@ -192,9 +190,17 @@ gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config )
    checkpoint_CTA_t = m_function_model_config.get_checkpoint_CTA_t();
    checkpoint_insn_Y = m_function_model_config.get_checkpoint_insn_Y();
 
+   // initialize texture mappings to empty
+   m_NameToTextureInfo.clear();
+   m_NameToCudaArray.clear();
+   m_TextureRefToName.clear();
+   m_NameToAttribute.clear();
 
    if(m_function_model_config.get_ptx_inst_debug_to_file() != 0) 
       ptx_inst_debug_file = fopen(m_function_model_config.get_ptx_inst_debug_file(), "w");
+
+   gpu_sim_cycle=0;
+   gpu_tot_sim_cycle=0;
 }
 
 address_type line_size_based_tag_func(new_addr_type address, new_addr_type line_size)
@@ -688,7 +694,10 @@ 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 )
+/*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. */
+kernel_info_t::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)   
 {
     m_kernel_entry=entry;
     m_grid_dim=gridDim;
@@ -708,6 +717,8 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *
     m_launch_latency = g_kernel_launch_latency;
 
     volta_cache_config_set=false;
+    m_NameToCudaArray = nameToCudaArray;
+    m_NameToTextureInfo = nameToTextureInfo;
 }
 
 kernel_info_t::~kernel_info_t()
@@ -822,10 +833,11 @@ void kernel_info_t::destroy_cta_streams() {
      m_cta_streams.clear();
 }
 
-simt_stack::simt_stack( unsigned wid, unsigned warpSize)
+simt_stack::simt_stack( unsigned wid, unsigned warpSize,  class gpgpu_sim * gpu)
 {
     m_warp_id=wid;
     m_warp_size = warpSize;
+    m_gpu=gpu;
     reset();
 }
 
@@ -1025,7 +1037,7 @@ void simt_stack::update( simt_mask_t &thread_done, addr_vector_t &next_pc, addre
     		simt_stack_entry new_stack_entry;
     		new_stack_entry.m_pc = tmp_next_pc;
     		new_stack_entry.m_active_mask = tmp_active_mask;
-    		new_stack_entry.m_branch_div_cycle = gpu_sim_cycle+gpu_tot_sim_cycle;
+    		new_stack_entry.m_branch_div_cycle = m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle;
     		new_stack_entry.m_type = STACK_ENTRY_TYPE_CALL;
     		m_stack.push_back(new_stack_entry);
     		return;
@@ -1057,7 +1069,7 @@ void simt_stack::update( simt_mask_t &thread_done, addr_vector_t &next_pc, addre
             new_recvg_pc = recvg_pc;
             if (new_recvg_pc != top_recvg_pc) {
                 m_stack.back().m_pc = new_recvg_pc;
-                m_stack.back().m_branch_div_cycle = gpu_sim_cycle+gpu_tot_sim_cycle;
+                m_stack.back().m_branch_div_cycle = m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle;
 
                 m_stack.push_back(simt_stack_entry());
             }
@@ -1149,7 +1161,7 @@ void core_t::initilizeSIMTStack(unsigned warp_count, unsigned warp_size)
 { 
     m_simt_stack = new simt_stack*[warp_count];
     for (unsigned i = 0; i < warp_count; ++i) 
-        m_simt_stack[i] = new simt_stack(i,warp_size);
+        m_simt_stack[i] = new simt_stack(i,warp_size,m_gpu);
     m_warp_size = warp_size;
     m_warp_count = warp_count;
 }
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index 28a22e0..1735c2f 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -212,7 +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();
 
    void inc_running() { m_num_cores_running++; }
@@ -275,6 +275,23 @@ public:
    std::list<class ptx_thread_info *> &active_threads() { return m_active_threads; }
    class memory_space *get_param_memory() { return m_param_mem; }
 
+   
+   //The following functions access texture bindings present at the kernel's launch
+   
+   const struct cudaArray* get_texarray( const std::string &texname ) const
+   {
+      std::map<std::string,const struct cudaArray*>::const_iterator t=m_NameToCudaArray.find(texname);
+      assert(t != m_NameToCudaArray.end());
+      return t->second;
+   }
+
+   const struct textureInfo* get_texinfo( const std::string &texname ) const
+   {
+      std::map<std::string, const struct textureInfo*>::const_iterator t=m_NameToTextureInfo.find(texname);
+      assert(t != m_NameToTextureInfo.end());
+      return t->second;
+   }
+
 private:
    kernel_info_t( const kernel_info_t & ); // disable copy constructor
    void operator=( const kernel_info_t & ); // disable copy operator
@@ -283,6 +300,10 @@ private:
 
    unsigned m_uid;
    static unsigned m_next_uid;
+   
+   //These maps contain the snapshot of the texture mappings at kernel launch
+   std::map<std::string, const struct cudaArray*> m_NameToCudaArray;
+   std::map<std::string, const struct textureInfo*> m_NameToTextureInfo;
 
    dim3 m_grid_dim;
    dim3 m_block_dim;
@@ -365,6 +386,8 @@ struct core_config {
 
 	unsigned gpgpu_max_insn_issue_per_warp;
 	bool gmem_skip_L1D; // on = global memory access always skip the L1 cache
+
+	bool adaptive_volta_cache_config;
 };
 
 // bounded stack that implements simt reconvergence using pdom mechanism from MICRO'07 paper
@@ -376,7 +399,7 @@ typedef std::vector<address_type> addr_vector_t;
 
 class simt_stack {
 public:
-    simt_stack( unsigned wid,  unsigned warpSize);
+    simt_stack( unsigned wid,  unsigned warpSize, class gpgpu_sim * gpu);
 
     void reset();
     void launch( address_type start_pc, const simt_mask_t &active_mask );
@@ -393,6 +416,7 @@ protected:
     unsigned m_warp_id;
     unsigned m_warp_size;
 
+
     enum stack_entry_type {
         STACK_ENTRY_TYPE_NORMAL = 0,
         STACK_ENTRY_TYPE_CALL
@@ -410,6 +434,8 @@ protected:
     };
 
     std::deque<simt_stack_entry> m_stack;
+
+    class gpgpu_sim * m_gpu;
 };
 
 #define GLOBAL_HEAP_START 0xC0000000
@@ -548,6 +574,12 @@ public:
     int resume_CTA;
     int checkpoint_CTA_t;
     int checkpoint_insn_Y;
+
+    //Move some cycle core stats here instead of being global
+    unsigned long long  gpu_sim_cycle;
+    unsigned long long  gpu_tot_sim_cycle;
+
+
     void* gpu_malloc( size_t size );
     void* gpu_mallocarray( size_t count );
     void  gpu_memset( size_t dst_start_addr, int c, size_t count );
@@ -580,8 +612,8 @@ public:
 
     const struct textureInfo* get_texinfo( const std::string &texname ) const
     {
-        std::map<std::string, const struct textureInfo*>::const_iterator t=m_NameToTexureInfo.find(texname);
-        assert(t != m_NameToTexureInfo.end());
+        std::map<std::string, const struct textureInfo*>::const_iterator t=m_NameToTextureInfo.find(texname);
+        assert(t != m_NameToTextureInfo.end());
         return t->second;
     }
 
@@ -594,6 +626,10 @@ public:
 
     const gpgpu_functional_sim_config &get_config() const { return m_function_model_config; }
     FILE* get_ptx_inst_debug_file() { return ptx_inst_debug_file; }
+    
+    //  These maps return the current texture mappings for the GPU at any given time.
+    std::map<std::string, const struct cudaArray*> getNameArrayMapping() {return m_NameToCudaArray;}
+    std::map<std::string, const struct textureInfo*> getNameInfoMapping() {return m_NameToTextureInfo;}
 
 protected:
     const gpgpu_functional_sim_config &m_function_model_config;
@@ -604,11 +640,11 @@ protected:
     class memory_space *m_surf_mem;
 
     unsigned long long m_dev_malloc;
-    
+    //  These maps contain the current texture mappings for the GPU at any given time. 
     std::map<std::string, std::set<const struct textureReference*> > m_NameToTextureRef;
     std::map<const struct textureReference*, std::string> m_TextureRefToName;
     std::map<std::string, const struct cudaArray*> m_NameToCudaArray;
-    std::map<std::string, const struct textureInfo*> m_NameToTexureInfo;
+    std::map<std::string, const struct textureInfo*> m_NameToTextureInfo;
     std::map<std::string, const struct textureReferenceAttr*> m_NameToAttribute;
 };
 
@@ -808,8 +844,8 @@ public:
 
 class mem_fetch_allocator {
 public:
-    virtual mem_fetch *alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr ) const = 0;
-    virtual mem_fetch *alloc( const class warp_inst_t &inst, const mem_access_t &access ) const = 0;
+    virtual mem_fetch *alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const = 0;
+    virtual mem_fetch *alloc( const class warp_inst_t &inst, const mem_access_t &access, unsigned long long cycle ) const = 0;
 };
 
 // the maximum number of destination, source, or address uarch operands in a instruction
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index ec68b5b..f7bb9cc 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -218,7 +218,7 @@ void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* te
    texInfo->Ty_numbits = intLOGB2(Ty);
    texInfo->texel_size = texel_size;
    texInfo->texel_size_numbits = intLOGB2(texel_size);
-   m_NameToTexureInfo[texname] = texInfo;
+   m_NameToTextureInfo[texname] = texInfo;
 }
 
 void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref)
@@ -226,7 +226,7 @@ void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref)
    //assumes bind-use-unbind-bind-use-unbind pattern
    std::string texname = gpgpu_ptx_sim_findNamefromTexture(texref);
    m_NameToCudaArray.erase(texname);
-   m_NameToTexureInfo.erase(texname);
+   m_NameToTextureInfo.erase(texname);
 }
 
 unsigned g_assemble_code_next_pc=0; 
@@ -1507,7 +1507,15 @@ static unsigned get_tex_datasize( const ptx_instruction *pI, ptx_thread_info *th
    std::string texname = src1.name();
 
    gpgpu_t *gpu = thread->get_gpu();
-   const struct textureInfo* texInfo = gpu->get_texinfo(texname);
+   /*
+     For programs with many streams, textures can be bound and unbound
+     asynchronously.  This means we need to use the kernel's "snapshot" of
+     the state of the texture mappings when it was launched (so that we
+     don't try to access the incorrect texture mapping if it's been updated,
+     or that we don't access a mapping that has been unbound).
+    */
+   kernel_info_t& k = thread->get_kernel();
+   const struct textureInfo* texInfo = k.get_texinfo(texname);
 
    unsigned data_size = texInfo->texel_size;
    return data_size; 
@@ -1918,7 +1926,7 @@ kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,
                                              struct dim3 blockDim,
                                              gpgpu_t *gpu )
 {
-   kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry);
+   kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping());
    unsigned argcount=args.size();
    unsigned argn=1;
    for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) {
diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h
index abd32f9..e690356 100644
--- a/src/cuda-sim/cuda-sim.h
+++ b/src/cuda-sim/cuda-sim.h
@@ -47,6 +47,8 @@ extern int g_debug_thread_uid;
 extern void ** g_inst_classification_stat;
 extern void ** g_inst_op_classification_stat;
 extern int g_ptx_kernel_count; // used for classification stat collection purposes 
+extern char *opcode_latency_int, *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor;
+
 
 void ptx_opcocde_latency_options (option_parser_t opp);
 extern class kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,
diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc
index b399133..86e8147 100644
--- a/src/cuda-sim/cuda_device_runtime.cc
+++ b/src/cuda-sim/cuda_device_runtime.cc
@@ -198,8 +198,9 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info *
             memory_space *device_kernel_param_mem;
 
             //create child kernel_info_t and index it with parameter_buffer address
-            device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry);
-            device_grid->launch_cycle = gpu_sim_cycle + gpu_tot_sim_cycle;
+	    gpgpu_t* gpu=thread->get_gpu();
+            device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry, gpu->getNameArrayMapping(), gpu->getNameInfoMapping());
+            device_grid->launch_cycle = gpu->gpu_sim_cycle + gpu->gpu_tot_sim_cycle;
             kernel_info_t & parent_grid = thread->get_kernel();
             DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" <<
                 thread->get_ctaid().x << ", " << thread->get_ctaid().y << ", " << thread->get_ctaid().z <<
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index c85654c..11001d7 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -5258,11 +5258,18 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread )
    if (!ptx_tex_regs) ptx_tex_regs = new ptx_reg_t[4];
    unsigned nelem = src2.get_vect_nelem();
    thread->get_vector_operand_values(src2, ptx_tex_regs, nelem); //ptx_reg should be 4 entry vector type...coordinates into texture
-
+   /*
+     For programs with many streams, textures can be bound and unbound
+     asynchronously.  This means we need to use the kernel's "snapshot" of
+     the state of the texture mappings when it was launched (so that we
+     don't try to access the incorrect texture mapping if it's been updated,
+     or that we don't access a mapping that has been unbound).
+   */
    gpgpu_t *gpu = thread->get_gpu();
+   kernel_info_t &k = thread->get_kernel();
    const struct textureReference* texref = gpu->get_texref(texname);
-   const struct cudaArray* cuArray = gpu->get_texarray(texname); 
-   const struct textureInfo* texInfo = gpu->get_texinfo(texname);
+   const struct cudaArray* cuArray = k.get_texarray(texname); 
+   const struct textureInfo* texInfo = k.get_texinfo(texname);
    const struct textureReferenceAttr* texAttr = gpu->get_texattr(texname);
 
    //assume always 2D f32 input
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index d4bf631..3232361 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -61,6 +61,7 @@ addc    TC; ptx_lval.int_value = ADDC_OP; return OPCODE;
 and	TC; ptx_lval.int_value = AND_OP; return OPCODE;
 andn	TC; ptx_lval.int_value = ANDN_OP; return OPCODE;
 atom	TC; ptx_lval.int_value = ATOM_OP; return OPCODE;
+bar.warp 	TC; ptx_lval.int_value = NOP_OP; return OPCODE;
 bar 	TC; ptx_lval.int_value = BAR_OP; return OPCODE;
 bfe     TC; ptx_lval.int_value = BFE_OP; return OPCODE;
 bfi     TC; ptx_lval.int_value = BFI_OP; return OPCODE;
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 8c6f361..da7c32d 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -50,9 +50,12 @@ extern int ptx__scan_string(const char*);
 extern std::map<unsigned,const char*> get_duplicate();
 
 const char *g_ptxinfo_filename;
-extern int ptxinfo_parse();
-extern int ptxinfo_debug;
-extern FILE *ptxinfo_in;
+
+typedef void * yyscan_t;
+extern int ptxinfo_lex_init(yyscan_t* scanner);
+extern void ptxinfo_set_in  (FILE * _in_str ,yyscan_t yyscanner );
+extern int ptxinfo_parse(yyscan_t scanner);
+extern int ptxinfo_lex_destroy(yyscan_t scanner);
 
 static bool g_save_embedded_ptx;
 static int g_occupancy_sm_number;
@@ -347,8 +350,13 @@ void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_versio
 	}
 
 	g_ptxinfo_filename = strdup(ptxas_filename.c_str());
+    FILE *ptxinfo_in;
     ptxinfo_in = fopen(g_ptxinfo_filename,"r");
-    ptxinfo_parse();
+    yyscan_t scanner;
+    ptxinfo_lex_init(&scanner);
+    ptxinfo_set_in(ptxinfo_in, scanner);
+    ptxinfo_parse(scanner);
+    ptxinfo_lex_destroy(scanner);
     fclose(ptxinfo_in);
 }
 
@@ -410,9 +418,15 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
     if( result != 0 ) {
     	// 65280 = duplicate errors
     	if (result == 65280) {
+		FILE *ptxinfo_in;
     		ptxinfo_in = fopen(tempfile_ptxinfo,"r");
 		g_ptxinfo_filename = tempfile_ptxinfo;
-		ptxinfo_parse();
+		yyscan_t scanner;
+		ptxinfo_lex_init(&scanner);
+		ptxinfo_set_in(ptxinfo_in, scanner);
+		ptxinfo_parse(scanner);
+		ptxinfo_lex_destroy(scanner);
+		fclose(ptxinfo_in);
 
     		fix_duplicate_errors(fname2);
     		snprintf(commandline,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file  /dev/null 2> %s",
@@ -495,9 +509,15 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
         g_ptxinfo_filename = final_tempfile_ptxinfo;
     else
 	g_ptxinfo_filename = tempfile_ptxinfo;
+    FILE *ptxinfo_in;
     ptxinfo_in = fopen(g_ptxinfo_filename,"r");
 
-    ptxinfo_parse();
+    yyscan_t scanner;
+    ptxinfo_lex_init(&scanner);
+    ptxinfo_set_in(ptxinfo_in, scanner);
+    ptxinfo_parse(scanner);
+    ptxinfo_lex_destroy(scanner);
+    fclose(ptxinfo_in);
 
     snprintf(commandline,1024,"rm -f *info");
     if( system(commandline) != 0 ) {
@@ -515,4 +535,4 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num
     	    exit(1);
         }
     }
-}
\ No newline at end of file
+}
diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc
index 820287d..8ad651e 100644
--- a/src/cuda-sim/ptx_sim.cc
+++ b/src/cuda-sim/ptx_sim.cc
@@ -222,7 +222,7 @@ void ptx_thread_info::set_done()
 {
    assert( !m_at_barrier );
    m_thread_done = true;
-   m_cycle_done = gpu_sim_cycle; 
+   m_cycle_done = m_gpu->gpu_sim_cycle;
 }
 
 unsigned ptx_thread_info::get_builtin( int builtin_id, unsigned dim_mod ) 
@@ -230,15 +230,15 @@ unsigned ptx_thread_info::get_builtin( int builtin_id, unsigned dim_mod )
    assert( m_valid );
    switch ((builtin_id&0xFFFF)) {
    case CLOCK_REG:
-      return (unsigned)(gpu_sim_cycle + gpu_tot_sim_cycle);
+      return (unsigned)(m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle);
    case CLOCK64_REG:
       abort(); // change return value to unsigned long long?
 	  // GPGPUSim clock is 4 times slower - multiply by 4
-	   return (gpu_sim_cycle + gpu_tot_sim_cycle)*4;
+	   return (m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle)*4;
    case HALFCLOCK_ID:
       // GPGPUSim clock is 4 times slower - multiply by 4
 	  // Hardware clock counter is incremented at half the shader clock frequency - divide by 2 (Henry '10)
-      return (gpu_sim_cycle + gpu_tot_sim_cycle)*2;
+      return (m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle)*2;
    case CTAID_REG:
       assert( dim_mod < 3 );
       if( dim_mod == 0 ) return m_ctaid.x;
diff --git a/src/cuda-sim/ptxinfo.l b/src/cuda-sim/ptxinfo.l
index 33c2748..a190e6d 100644
--- a/src/cuda-sim/ptxinfo.l
+++ b/src/cuda-sim/ptxinfo.l
@@ -31,17 +31,19 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 %option noyywrap
 %option yylineno
 %option prefix="ptxinfo_"
+
+%option bison-bridge
+%option reentrant
+
 %{
 #include "ptxinfo.tab.h"
 #include <string.h>
 
 #define LINEBUF_SIZE 1024
-char ptxinfo_linebuf[LINEBUF_SIZE];
-unsigned ptxinfo_col = 0;
-#define TC if( (ptxinfo_lineno == 1) && ((ptxinfo_col + strlen(ptxinfo_text)) < LINEBUF_SIZE) ) { \
-		strncpy(ptxinfo_linebuf+ptxinfo_col,ptxinfo_text,strlen(ptxinfo_text)); \
+#define TC if( (yylineno == 1) && (yylval->col + strlen(yytext) < LINEBUF_SIZE) ) { \
+		strncpy(yylval->linebuf+yylval->col,yytext,strlen(yytext)); \
 	   } \
-	   ptxinfo_col+=strlen(ptxinfo_text); 
+	   yylval->col+=strlen(yytext);
 %}
 
 %%
@@ -61,12 +63,12 @@ unsigned ptxinfo_col = 0;
 "for"		TC; return FOR;
 "textures"	TC; return TEXTURES;
 "error   : Duplicate definition of"	TC; return DUPLICATE;
-"function"	TC; ptxinfo_lval.string_value = strdup(yytext); return FUNCTION;
-"variable"	TC; ptxinfo_lval.string_value = strdup(yytext); return VARIABLE;
+"function"	TC; yylval->string_value = strdup(yytext); return FUNCTION;
+"variable"	TC; yylval->string_value = strdup(yytext); return VARIABLE;
 "fatal   : Ptx assembly aborted due to errors"	TC; return FATAL;
 
-[_A-Za-z$%][_0-9A-Za-z$]*  TC; ptxinfo_lval.string_value = strdup(yytext); return IDENTIFIER;
-[-]{0,1}[0-9]+	 TC; ptxinfo_lval.int_value =  atoi(yytext); return INT_OPERAND;
+[_A-Za-z$%][_0-9A-Za-z$]*  TC; yylval->string_value = strdup(yytext); return IDENTIFIER;
+[-]{0,1}[0-9]+	 TC; yylval->int_value =  atoi(yytext); return INT_OPERAND;
 
 "+"	TC; return PLUS;
 ","     TC; return COMMA;
@@ -78,7 +80,7 @@ unsigned ptxinfo_col = 0;
 " " TC;
 "\t" TC;
 
-\n.*  ptxinfo_col=0; strncpy(ptxinfo_linebuf, yytext + 1, 1024); yyless( 1 );
+\n.*  yylval->col=0; strncpy(yylval->linebuf, yytext + 1, 1024); yyless( 1 );
 
 %%
 
@@ -86,18 +88,19 @@ extern int g_ptxinfo_error_detected;
 extern const char *g_filename;
 extern const char *g_ptxinfo_filename;
 
-int ptxinfo_error( const char *s )
+int ptxinfo_error(yyscan_t yyscanner, const char* msg)
 {
+    struct yyguts_t * yyg = (struct yyguts_t*)yyscanner;
 	int i;
 	g_ptxinfo_error_detected = 1;
 	fflush(stdout);
 	printf("GPGPU-Sim: ERROR while parsing output of ptxas (used to capture resource usage information)\n");
-	if( s != NULL )
-		printf("GPGPU-Sim:     %s (%s:%u) Syntax error:\n\n", g_filename, g_ptxinfo_filename, ptxinfo_lineno );
-	printf("   %s\n", ptxinfo_linebuf );
+	if( msg != NULL )
+		printf("GPGPU-Sim:     %s (%s:%u) Syntax error:\n\n", g_filename, g_ptxinfo_filename, yylineno );
+	printf("   %s\n", yylval->linebuf );
 	printf("   ");
-	for( i=0; i < ptxinfo_col-1; i++ ) {
-		if( ptxinfo_linebuf[i] == '\t' ) printf("\t");
+	for( i=0; i < yylval->col-1; i++ ) {
+		if( yylval->linebuf[i] == '\t' ) printf("\t");
 		else printf(" ");
 	}
 			
diff --git a/src/cuda-sim/ptxinfo.y b/src/cuda-sim/ptxinfo.y
index d241d8c..eed304c 100644
--- a/src/cuda-sim/ptxinfo.y
+++ b/src/cuda-sim/ptxinfo.y
@@ -27,9 +27,20 @@ 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.
 */
 
+%{
+typedef void * yyscan_t;
+%}
+
+%define api.pure full
+%parse-param {yyscan_t scanner}
+%lex-param {yyscan_t scanner}
+
 %union {
+#define LINEBUF_SIZE 1024
   int    int_value;
   char * string_value;
+  char linebuf[LINEBUF_SIZE];
+  unsigned col;
 }
 
 %token <int_value> INT_OPERAND
@@ -66,7 +77,8 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 	
 	static unsigned g_declared;
 	static unsigned g_system;
-	int ptxinfo_lex(void);
+	int ptxinfo_lex(YYSTYPE * yylval_param, yyscan_t yyscanner);
+	void yyerror(yyscan_t yyscanner, const char* msg);
 	void ptxinfo_addinfo();
 	void ptxinfo_function(const char *fname );
 	void ptxinfo_regs( unsigned nregs );
@@ -74,7 +86,6 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 	void ptxinfo_gmem( unsigned declared, unsigned system );
 	void ptxinfo_smem( unsigned declared, unsigned system );
 	void ptxinfo_cmem( unsigned nbytes, unsigned bank );
-	int ptxinfo_error(const char*);
 	void ptxinfo_linenum( unsigned );
 	void ptxinfo_dup_type( const char* );
 %}
diff --git a/src/gpgpu-sim/addrdec.cc b/src/gpgpu-sim/addrdec.cc
index 8651869..ca88ec9 100644
--- a/src/gpgpu-sim/addrdec.cc
+++ b/src/gpgpu-sim/addrdec.cc
@@ -165,6 +165,27 @@ void linear_to_raw_address_translation::addrdec_tlx(new_addr_type addr, addrdec_
 			assert(tlx->chip < m_n_channel);
 			break;
 		}
+		case RANDOM:
+		{
+			new_addr_type chip_address = (addr>>ADDR_CHIP_S);
+			tr1_hash_map<new_addr_type,unsigned>::const_iterator got = address_random_interleaving.find (chip_address);
+			  if ( got == address_random_interleaving.end() ) {
+				  unsigned new_chip_id = rand() % (m_n_channel*m_n_sub_partition_in_channel);
+				  address_random_interleaving[chip_address] = new_chip_id;
+				  tlx->chip = new_chip_id/m_n_sub_partition_in_channel;
+				  tlx->sub_partition = new_chip_id;
+			  }
+			  else {
+				  unsigned new_chip_id = got->second;
+				  tlx->chip = new_chip_id/m_n_sub_partition_in_channel;
+				  tlx->sub_partition = new_chip_id;
+			  }
+
+			  assert(tlx->chip < m_n_channel);
+			  assert(tlx->sub_partition < m_n_channel*m_n_sub_partition_in_channel);
+			  return;
+			break;
+		}
 		case CUSTOM:
 			/* No custom set function implemented */
 			//Do you custom index here
@@ -175,9 +196,9 @@ void linear_to_raw_address_translation::addrdec_tlx(new_addr_type addr, addrdec_
 	}
 
    // combine the chip address and the lower bits of DRAM bank address to form the subpartition ID
-   unsigned sub_partition_addr_mask = m_n_sub_partition_in_channel - 1; 
+   unsigned sub_partition_addr_mask = m_n_sub_partition_in_channel - 1;
    tlx->sub_partition = tlx->chip * m_n_sub_partition_in_channel
-                        + (tlx->bk & sub_partition_addr_mask); 
+                        + (tlx->bk & sub_partition_addr_mask);
 }
 
 void linear_to_raw_address_translation::addrdec_parseoption(const char *option)
@@ -396,6 +417,10 @@ void linear_to_raw_address_translation::init(unsigned int n_channel, unsigned in
    if (run_test) {
       sweep_test(); 
    }
+
+   if(memory_partition_indexing == RANDOM)
+     srand (1);
+
 }
 
 #include "../tr1_hash_map.h" 
diff --git a/src/gpgpu-sim/addrdec.h b/src/gpgpu-sim/addrdec.h
index bdc5fec..a5333fb 100644
--- a/src/gpgpu-sim/addrdec.h
+++ b/src/gpgpu-sim/addrdec.h
@@ -40,6 +40,7 @@ enum partition_index_function{
 	BITWISE_PERMUTATION,
 	IPOLY,
 	PAE,
+	RANDOM,
     CUSTOM
 };
 
@@ -55,6 +56,7 @@ struct addrdec_t {
    unsigned sub_partition; 
 };
 
+
 class linear_to_raw_address_translation {
 public:
    linear_to_raw_address_translation();
@@ -62,7 +64,7 @@ public:
    void init(unsigned int n_channel, unsigned int n_sub_partition_in_channel); 
 
    // accessors
-   void addrdec_tlx(new_addr_type addr, addrdec_t *tlx) const; 
+   void addrdec_tlx(new_addr_type addr, addrdec_t *tlx) const;
    new_addr_type partition_address( new_addr_type addr ) const;
 
 private:
@@ -92,6 +94,7 @@ private:
    unsigned int gap;
    int m_n_channel;
    int m_n_sub_partition_in_channel; 
+
 };
 
 #endif
diff --git a/src/gpgpu-sim/dram.cc b/src/gpgpu-sim/dram.cc
index 6c11b43..5e36d4b 100644
--- a/src/gpgpu-sim/dram.cc
+++ b/src/gpgpu-sim/dram.cc
@@ -42,12 +42,13 @@ template class fifo_pipeline<mem_fetch>;
 template class fifo_pipeline<dram_req_t>;
 
 dram_t::dram_t( unsigned int partition_id, const struct memory_config *config, memory_stats_t *stats,
-                memory_partition_unit *mp )
+                memory_partition_unit *mp, gpgpu_sim* gpu )
 {
    id = partition_id;
    m_memory_partition_unit = mp;
    m_stats = stats;
    m_config = config;
+   m_gpu = gpu;
 
    //rowblp
    access_num=0;
@@ -191,11 +192,12 @@ unsigned int dram_t::queue_limit() const
 }
 
 
-dram_req_t::dram_req_t( class mem_fetch *mf, unsigned banks, unsigned dram_bnk_indexing_policy)
+dram_req_t::dram_req_t( class mem_fetch *mf, unsigned banks, unsigned dram_bnk_indexing_policy, class gpgpu_sim* gpu)
 {
    txbytes = 0;
    dqbytes = 0;
    data = mf;
+   m_gpu = gpu;
 
    const addrdec_t &tlx = mf->get_tlx_addr();
 
@@ -226,9 +228,9 @@ dram_req_t::dram_req_t( class mem_fetch *mf, unsigned banks, unsigned dram_bnk_i
    col = tlx.col; 
    nbytes = mf->get_data_size();
 
-   timestamp = gpu_tot_sim_cycle + gpu_sim_cycle;
+   timestamp = m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle;
    addr = mf->get_addr();
-   insertion_time = (unsigned) gpu_sim_cycle;
+   insertion_time = (unsigned) m_gpu->gpu_sim_cycle;
    rw = data->get_is_write()?WRITE:READ;
 }
 
@@ -236,9 +238,9 @@ void dram_t::push( class mem_fetch *data )
 {
    assert(id == data->get_tlx_addr().chip); // Ensure request is in correct memory partition
 
-   dram_req_t *mrq = new dram_req_t(data,m_config->nbk,m_config->dram_bnk_indexing_policy);
+   dram_req_t *mrq = new dram_req_t(data,m_config->nbk,m_config->dram_bnk_indexing_policy,m_memory_partition_unit->get_mgpu());
 
-   data->set_status(IN_PARTITION_MC_INTERFACE_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle);
+   data->set_status(IN_PARTITION_MC_INTERFACE_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
 	   mrqq->push(mrq);
 
    // stats...
@@ -259,7 +261,7 @@ void dram_t::scheduler_fifo()
    if (!mrqq->empty()) {
       unsigned int bkn;
       dram_req_t *head_mrqq = mrqq->top();
-      head_mrqq->data->set_status(IN_PARTITION_MC_BANK_ARB_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle);
+      head_mrqq->data->set_status(IN_PARTITION_MC_BANK_ARB_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
       bkn = head_mrqq->bk;
       if (!bk[bkn]->mrq) 
          bk[bkn]->mrq = mrqq->pop();
@@ -283,7 +285,7 @@ void dram_t::cycle()
 
            if (cmd->dqbytes >= cmd->nbytes) {
               mem_fetch *data = cmd->data; 
-              data->set_status(IN_PARTITION_MC_RETURNQ,gpu_sim_cycle+gpu_tot_sim_cycle); 
+              data->set_status(IN_PARTITION_MC_RETURNQ,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
               if( data->get_access_type() != L1_WRBK_ACC && data->get_access_type() != L2_WRBK_ACC ) {
                  data->set_reply();
                  returnq->push(data);
@@ -566,7 +568,7 @@ bool dram_t::issue_col_command(int j)
 	bool issued = false;
 	unsigned grp = get_bankgrp_number(j);
     if (bk[j]->mrq) { //if currently servicing a memory request
-        bk[j]->mrq->data->set_status(IN_PARTITION_DRAM,gpu_sim_cycle+gpu_tot_sim_cycle);
+        bk[j]->mrq->data->set_status(IN_PARTITION_DRAM,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
        // correct row activated for a READ
        if ( !issued && !CCDc && !bk[j]->RCDc &&
             !(bkgrp[grp]->CCDLc) &&
@@ -654,7 +656,7 @@ bool dram_t::issue_row_command(int j)
 	bool issued = false;
 	unsigned grp = get_bankgrp_number(j);
     if (bk[j]->mrq) { //if currently servicing a memory request
-        bk[j]->mrq->data->set_status(IN_PARTITION_DRAM,gpu_sim_cycle+gpu_tot_sim_cycle);
+        bk[j]->mrq->data->set_status(IN_PARTITION_DRAM,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
       //     bank is idle
     //else
   	  if ( !issued && !RRDc &&
@@ -723,10 +725,10 @@ void dram_t::print( FILE* simFile) const
            id, m_config->nbk, m_config->busW, m_config->BL, m_config->CL );
    fprintf(simFile,"tRRD=%d tCCD=%d, tRCD=%d tRAS=%d tRP=%d tRC=%d\n",
            m_config->tRRD, m_config->tCCD, m_config->tRCD, m_config->tRAS, m_config->tRP, m_config->tRC );
-   fprintf(simFile,"n_cmd=%d n_nop=%d n_act=%d n_pre=%d n_ref_event=%d n_req=%d n_rd=%d n_rd_L2_A=%d n_write=%d n_wr_bk=%d bw_util=%.4g\n",
+   fprintf(simFile,"n_cmd=%llu n_nop=%llu n_act=%llu n_pre=%llu n_ref_event=%llu n_req=%llu n_rd=%llu n_rd_L2_A=%llu n_write=%llu n_wr_bk=%llu bw_util=%.4g\n",
            n_cmd, n_nop, n_act, n_pre, n_ref, n_req, n_rd, n_rd_L2_A, n_wr, n_wr_WB,
            (float)bwutil/n_cmd);
-   fprintf(simFile,"n_activity=%d dram_eff=%.4g\n",
+   fprintf(simFile,"n_activity=%llu dram_eff=%.4g\n",
            n_activity, (float)bwutil/n_activity);
    for (i=0;i<m_config->nbk;i++) {
       fprintf(simFile, "bk%d: %da %di ",i,bk[i]->n_access,bk[i]->n_idle);
@@ -745,39 +747,39 @@ void dram_t::print( FILE* simFile) const
 
    printf("\nBW Util details:\n");
    printf("bwutil = %.6f \n", (float)bwutil/n_cmd);
-   printf("total_CMD = %d \n", n_cmd);
-   printf("util_bw = %d \n", util_bw);
-   printf("Wasted_Col = %d \n", wasted_bw_col);
-   printf("Wasted_Row = %d \n", wasted_bw_row);
-   printf("Idle = %d \n", idle_bw);
+   printf("total_CMD = %llu \n", n_cmd);
+   printf("util_bw = %llu \n", util_bw);
+   printf("Wasted_Col = %llu \n", wasted_bw_col);
+   printf("Wasted_Row = %llu \n", wasted_bw_row);
+   printf("Idle = %llu \n", idle_bw);
 
    printf("\nBW Util Bottlenecks: \n");
-   printf("RCDc_limit = %d \n", RCDc_limit);
-   printf("RCDWRc_limit = %d \n", RCDWRc_limit);
-   printf("WTRc_limit = %d \n", WTRc_limit);
-   printf("RTWc_limit = %d \n", RTWc_limit);
-   printf("CCDLc_limit = %d \n", CCDLc_limit);
-   printf("rwq = %d \n", rwq_limit);
-   printf("CCDLc_limit_alone = %d \n", CCDLc_limit_alone);
-   printf("WTRc_limit_alone = %d \n", WTRc_limit_alone);
-   printf("RTWc_limit_alone = %d \n", RTWc_limit_alone);
+   printf("RCDc_limit = %llu \n", RCDc_limit);
+   printf("RCDWRc_limit = %llu \n", RCDWRc_limit);
+   printf("WTRc_limit = %llu \n", WTRc_limit);
+   printf("RTWc_limit = %llu \n", RTWc_limit);
+   printf("CCDLc_limit = %llu \n", CCDLc_limit);
+   printf("rwq = %llu \n", rwq_limit);
+   printf("CCDLc_limit_alone = %llu \n", CCDLc_limit_alone);
+   printf("WTRc_limit_alone = %llu \n", WTRc_limit_alone);
+   printf("RTWc_limit_alone = %llu \n", RTWc_limit_alone);
 
    printf("\nCommands details: \n");
-   printf("total_CMD = %d \n", n_cmd);
-   printf("n_nop = %d \n", n_nop);
-   printf("Read = %d \n", n_rd);
-   printf("Write = %d \n",n_wr);
-   printf("L2_Alloc = %d \n", n_rd_L2_A);
-   printf("L2_WB = %d \n", n_wr_WB);
-   printf("n_act = %d \n", n_act);
-   printf("n_pre = %d \n", n_pre);
-   printf("n_ref = %d \n", n_ref);
-   printf("n_req = %d \n", n_req );
-   printf("total_req = %d \n", n_rd+n_wr+n_rd_L2_A+n_wr_WB);
+   printf("total_CMD = %llu \n", n_cmd);
+   printf("n_nop = %llu \n", n_nop);
+   printf("Read = %llu \n", n_rd);
+   printf("Write = %llu \n",n_wr);
+   printf("L2_Alloc = %llu \n", n_rd_L2_A);
+   printf("L2_WB = %llu \n", n_wr_WB);
+   printf("n_act = %llu \n", n_act);
+   printf("n_pre = %llu \n", n_pre);
+   printf("n_ref = %llu \n", n_ref);
+   printf("n_req = %llu \n", n_req );
+   printf("total_req = %llu \n", n_rd+n_wr+n_rd_L2_A+n_wr_WB);
 
    printf("\nDual Bus Interface Util: \n");
-   printf("issued_total_row = %lu \n", issued_total_row);
-   printf("issued_total_col = %lu \n", issued_total_col);
+   printf("issued_total_row = %llu \n", issued_total_row);
+   printf("issued_total_col = %llu \n", issued_total_col);
    printf("Row_Bus_Util =  %.6f \n", (float)issued_total_row / n_cmd);
    printf("CoL_Bus_Util = %.6f \n", (float)issued_total_col / n_cmd);
    printf("Either_Row_CoL_Bus_Util = %.6f \n",  (float)issued_total / n_cmd);
@@ -815,7 +817,7 @@ void dram_t::visualize() const
 
 void dram_t::print_stat( FILE* simFile ) 
 {
-   fprintf(simFile,"DRAM (%d): n_cmd=%d n_nop=%d n_act=%d n_pre=%d n_ref=%d n_req=%d n_rd=%d n_write=%d bw_util=%.4g ",
+   fprintf(simFile,"DRAM (%llu): n_cmd=%llu n_nop=%llu n_act=%llu n_pre=%llu n_ref=%llu n_req=%llu n_rd=%llu n_write=%llu bw_util=%.4g ",
            id, n_cmd, n_nop, n_act, n_pre, n_ref, n_req, n_rd, n_wr,
            (float)bwutil/n_cmd);
    fprintf(simFile, "mrqq: %d %.4g mrqsmax=%d ", max_mrqs, (float)ave_mrqs/n_cmd, max_mrqs_temp);
diff --git a/src/gpgpu-sim/dram.h b/src/gpgpu-sim/dram.h
index bee5b7b..7a3a2da 100644
--- a/src/gpgpu-sim/dram.h
+++ b/src/gpgpu-sim/dram.h
@@ -48,7 +48,7 @@
 
 class dram_req_t {
 public:
-   dram_req_t( class mem_fetch *data , unsigned banks, unsigned dram_bnk_indexing_policy);
+   dram_req_t( class mem_fetch *data , unsigned banks, unsigned dram_bnk_indexing_policy, class gpgpu_sim* gpu);
 
    unsigned int row;
    unsigned int col;
@@ -62,6 +62,7 @@ public:
    unsigned long long int addr;
    unsigned int insertion_time;
    class mem_fetch * data;
+   class gpgpu_sim * m_gpu;
 };
 
 struct bankgrp_t
@@ -110,7 +111,7 @@ class dram_t
 {
 public:
    dram_t( unsigned int parition_id, const struct memory_config *config, class memory_stats_t *stats, 
-           class memory_partition_unit *mp );
+           class memory_partition_unit *mp, class gpgpu_sim* gpu );
 
    bool full(bool is_write) const;
    void print( FILE* simFile ) const;
@@ -129,6 +130,7 @@ public:
    void dram_log (int task);
 
    class memory_partition_unit *m_memory_partition_unit;
+   class gpgpu_sim* m_gpu;
    unsigned int id;
 
    // Power Model
@@ -178,39 +180,39 @@ private:
    unsigned int dram_eff_bins[10];
    unsigned int last_n_cmd, last_n_activity, last_bwutil;
 
-   unsigned int n_cmd;
-   unsigned int n_activity;
-   unsigned int n_nop;
-   unsigned int n_act;
-   unsigned int n_pre;
-   unsigned int n_ref;
-   unsigned int n_rd;
-   unsigned int n_rd_L2_A;
-   unsigned int n_wr;
-   unsigned int n_wr_WB;
-   unsigned int n_req;
-   unsigned int max_mrqs_temp;
+   unsigned long long n_cmd;
+   unsigned long long n_activity;
+   unsigned long long n_nop;
+   unsigned long long n_act;
+   unsigned long long n_pre;
+   unsigned long long n_ref;
+   unsigned long long n_rd;
+   unsigned long long n_rd_L2_A;
+   unsigned long long n_wr;
+   unsigned long long n_wr_WB;
+   unsigned long long n_req;
+   unsigned long long max_mrqs_temp;
 
-   //some statistics to collect to see where BW is wasted?
-   unsigned wasted_bw_row;
-   unsigned wasted_bw_col;
-   unsigned util_bw;
-   unsigned idle_bw;
-   unsigned RCDc_limit;
-   unsigned CCDLc_limit;
-   unsigned CCDLc_limit_alone;
-   unsigned CCDc_limit;
-   unsigned WTRc_limit;
-   unsigned WTRc_limit_alone;
-   unsigned RCDWRc_limit;
-   unsigned RTWc_limit;
-   unsigned RTWc_limit_alone;
-   unsigned rwq_limit;
+   //some statistics to see where BW is wasted?
+   unsigned long long wasted_bw_row;
+   unsigned long long wasted_bw_col;
+   unsigned long long util_bw;
+   unsigned long long idle_bw;
+   unsigned long long RCDc_limit;
+   unsigned long long CCDLc_limit;
+   unsigned long long CCDLc_limit_alone;
+   unsigned long long CCDc_limit;
+   unsigned long long WTRc_limit;
+   unsigned long long WTRc_limit_alone;
+   unsigned long long RCDWRc_limit;
+   unsigned long long RTWc_limit;
+   unsigned long long RTWc_limit_alone;
+   unsigned long long rwq_limit;
 
    //row locality, BLP and other statistics
-   unsigned long  access_num;
-   unsigned long  read_num;
-   unsigned long  write_num;
+   unsigned long long access_num;
+   unsigned long long read_num;
+   unsigned long long write_num;
    unsigned long long hits_num;
    unsigned long long hits_read_num;
    unsigned long long hits_write_num;
diff --git a/src/gpgpu-sim/dram_sched.cc b/src/gpgpu-sim/dram_sched.cc
index ff50050..6ee6271 100644
--- a/src/gpgpu-sim/dram_sched.cc
+++ b/src/gpgpu-sim/dram_sched.cc
@@ -84,13 +84,13 @@ void frfcfs_scheduler::add_req( dram_req_t *req )
 
 void frfcfs_scheduler::data_collection(unsigned int bank)
 {
-   if (gpu_sim_cycle > row_service_timestamp[bank]) {
-      curr_row_service_time[bank] = gpu_sim_cycle - row_service_timestamp[bank];
+   if (m_dram->m_gpu->gpu_sim_cycle > row_service_timestamp[bank]) {
+      curr_row_service_time[bank] = m_dram->m_gpu->gpu_sim_cycle - row_service_timestamp[bank];
       if (curr_row_service_time[bank] > m_stats->max_servicetime2samerow[m_dram->id][bank])
          m_stats->max_servicetime2samerow[m_dram->id][bank] = curr_row_service_time[bank];
    }
    curr_row_service_time[bank] = 0;
-   row_service_timestamp[bank] = gpu_sim_cycle;
+   row_service_timestamp[bank] = m_dram->m_gpu->gpu_sim_cycle;
    if (m_stats->concurrent_row_access[m_dram->id][bank] > m_stats->max_conc_access2samerow[m_dram->id][bank]) {
       m_stats->max_conc_access2samerow[m_dram->id][bank] = m_stats->concurrent_row_access[m_dram->id][bank];
    }
@@ -215,7 +215,7 @@ void dram_t::scheduler_frfcfs()
     	  m_stats->total_n_reads++;
       }
 
-      req->data->set_status(IN_PARTITION_MC_INPUT_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle);
+      req->data->set_status(IN_PARTITION_MC_INPUT_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
       sched->add_req(req);
    }
 
@@ -228,14 +228,14 @@ void dram_t::scheduler_frfcfs()
          req = sched->schedule(b, bk[b]->curr_row);
 
          if ( req ) {
-            req->data->set_status(IN_PARTITION_MC_BANK_ARB_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle);
+            req->data->set_status(IN_PARTITION_MC_BANK_ARB_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
             prio = (prio+1)%m_config->nbk;
             bk[b]->mrq = req;
             if (m_config->gpgpu_memlatency_stat) {
-               mrq_latency = gpu_sim_cycle + gpu_tot_sim_cycle - bk[b]->mrq->timestamp;
+               mrq_latency = m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle - bk[b]->mrq->timestamp;
                m_stats->tot_mrq_latency += mrq_latency;
                m_stats->tot_mrq_num++;
-               bk[b]->mrq->timestamp = gpu_tot_sim_cycle + gpu_sim_cycle;
+               bk[b]->mrq->timestamp =m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle;
                m_stats->mrq_lat_table[LOGB2(mrq_latency)]++;
                if (mrq_latency > m_stats->max_mrq_latency) {
                   m_stats->max_mrq_latency = mrq_latency;
diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc
index ba81440..62849f8 100644
--- a/src/gpgpu-sim/gpu-cache.cc
+++ b/src/gpgpu-sim/gpu-cache.cc
@@ -26,6 +26,7 @@
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
 #include "gpu-cache.h"
+#include "gpu-sim.h"
 #include "stat-tool.h"
 #include <assert.h>
 
@@ -256,7 +257,7 @@ enum cache_request_status tag_array::probe( new_addr_type addr, unsigned &idx, m
 
     unsigned invalid_line = (unsigned)-1;
     unsigned valid_line = (unsigned)-1;
-    unsigned valid_timestamp = (unsigned)-1;
+    unsigned long long valid_timestamp = (unsigned)-1;
 
     bool all_reserved = true;
 
@@ -654,7 +655,7 @@ enum cache_request_status cache_stats::select_stats_status(enum cache_request_st
 		return access;
 }
 
-unsigned &cache_stats::operator()(int access_type, int access_outcome, bool fail_outcome){
+unsigned long long &cache_stats::operator()(int access_type, int access_outcome, bool fail_outcome){
     ///
     /// Simple method to read/modify the stat corresponding to (access_type, access_outcome)
     /// Used overloaded () to avoid the need for separate read/write member functions
@@ -673,7 +674,7 @@ unsigned &cache_stats::operator()(int access_type, int access_outcome, bool fail
 	}
 }
 
-unsigned cache_stats::operator()(int access_type, int access_outcome, bool fail_outcome) const{
+unsigned long long cache_stats::operator()(int access_type, int access_outcome, bool fail_outcome) const{
     ///
     /// Const accessor into m_stats.
     ///
@@ -740,7 +741,7 @@ void cache_stats::print_stats(FILE *fout, const char *cache_name) const{
     std::string m_cache_name = cache_name;
     for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) {
         for (unsigned status = 0; status < NUM_CACHE_REQUEST_STATUS; ++status) {
-            fprintf(fout, "\t%s[%s][%s] = %u\n",
+            fprintf(fout, "\t%s[%s][%s] = %llu\n",
                 m_cache_name.c_str(),
                 mem_access_type_str((enum mem_access_type)type),
                 cache_request_status_str((enum cache_request_status)status),
@@ -751,7 +752,7 @@ void cache_stats::print_stats(FILE *fout, const char *cache_name) const{
     }
     for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) {
     	 if(total_access[type] > 0)
-    	  fprintf(fout, "\t%s[%s][%s] = %u\n",
+    	  fprintf(fout, "\t%s[%s][%s] = %llu\n",
 				m_cache_name.c_str(),
 				mem_access_type_str((enum mem_access_type)type),
 				"TOTAL_ACCESS",
@@ -788,13 +789,13 @@ void cache_sub_stats::print_port_stats(FILE *fout, const char *cache_name) const
     fprintf(fout, "%s_fill_port_util = %.3f\n", cache_name, fill_port_util); 
 }
 
-unsigned cache_stats::get_stats(enum mem_access_type *access_type, unsigned num_access_type, enum cache_request_status *access_status, unsigned num_access_status) const{
+unsigned long long cache_stats::get_stats(enum mem_access_type *access_type, unsigned num_access_type, enum cache_request_status *access_status, unsigned num_access_status) const{
     ///
     /// Returns a sum of the stats corresponding to each "access_type" and "access_status" pair.
     /// "access_type" is an array of "num_access_type" mem_access_types.
     /// "access_status" is an array of "num_access_status" cache_request_statuses.
     ///
-    unsigned total=0;
+    unsigned long long total=0;
     for(unsigned type =0; type < num_access_type; ++type){
         for(unsigned status=0; status < num_access_status; ++status){
             if(!check_valid((int)access_type[type], (int)access_status[status]))
@@ -1183,7 +1184,8 @@ data_cache::wr_miss_wa_naive( new_addr_type addr,
                     mf->get_wid(),
                     mf->get_sid(),
                     mf->get_tpc(),
-                    mf->get_mem_config());
+                    mf->get_mem_config(),
+					m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle);
 
     bool do_miss = false;
     bool wb = false;
@@ -1201,7 +1203,7 @@ data_cache::wr_miss_wa_naive( new_addr_type addr,
         if( wb && (m_config.m_write_policy != WRITE_THROUGH) ) { 
         	assert(status == MISS);   //SECTOR_MISS and HIT_RESERVED should not send write back
             mem_fetch *wb = m_memfetch_creator->alloc(evicted.m_block_addr,
-                m_wrbk_type,evicted.m_modified_size,true);
+                m_wrbk_type,evicted.m_modified_size,true,m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle);
             send_write_request(wb, cache_event(WRITE_BACK_REQUEST_SENT, evicted), time, events);
         }
         return MISS;
@@ -1245,7 +1247,7 @@ data_cache::wr_miss_wa_fetch_on_write( new_addr_type addr,
 			   // (already modified lower level)
 			   if( wb && (m_config.m_write_policy != WRITE_THROUGH) ) {
 				   mem_fetch *wb = m_memfetch_creator->alloc(evicted.m_block_addr,
-					   m_wrbk_type,evicted.m_modified_size,true);
+					   m_wrbk_type,evicted.m_modified_size,true,m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle);
 				   send_write_request(wb, cache_event(WRITE_BACK_REQUEST_SENT, evicted), time, events);
 			   }
 			   return MISS;
@@ -1297,6 +1299,7 @@ data_cache::wr_miss_wa_fetch_on_write( new_addr_type addr,
 								mf->get_sid(),
 								mf->get_tpc(),
 								mf->get_mem_config(),
+								m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle,
 								NULL,
 								mf);
 
@@ -1320,7 +1323,7 @@ data_cache::wr_miss_wa_fetch_on_write( new_addr_type addr,
 				// (already modified lower level)
 				if(wb && (m_config.m_write_policy != WRITE_THROUGH) ){
 					mem_fetch *wb = m_memfetch_creator->alloc(evicted.m_block_addr,
-						m_wrbk_type,evicted.m_modified_size,true);
+						m_wrbk_type,evicted.m_modified_size,true,m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle);
 					send_write_request(wb, cache_event(WRITE_BACK_REQUEST_SENT, evicted), time, events);
 			}
 				return MISS;
@@ -1373,7 +1376,7 @@ data_cache::wr_miss_wa_lazy_fetch_on_read( new_addr_type addr,
 			   // (already modified lower level)
 			   if( wb && (m_config.m_write_policy != WRITE_THROUGH) ) {
 				   mem_fetch *wb = m_memfetch_creator->alloc(evicted.m_block_addr,
-					   m_wrbk_type,evicted.m_modified_size,true);
+					   m_wrbk_type,evicted.m_modified_size,true,m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle);
 				   send_write_request(wb, cache_event(WRITE_BACK_REQUEST_SENT, evicted), time, events);
 			   }
 			   return MISS;
@@ -1458,7 +1461,7 @@ data_cache::rd_miss_base( new_addr_type addr,
         // (already modified lower level)
         if(wb && (m_config.m_write_policy != WRITE_THROUGH) ){ 
             mem_fetch *wb = m_memfetch_creator->alloc(evicted.m_block_addr,
-                m_wrbk_type,evicted.m_modified_size,true);
+                m_wrbk_type,evicted.m_modified_size,true,m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle);
         send_write_request(wb, WRITE_BACK_REQUEST_SENT, time, events);
     }
         return MISS;
diff --git a/src/gpgpu-sim/gpu-cache.h b/src/gpgpu-sim/gpu-cache.h
index e663cf6..85e534e 100644
--- a/src/gpgpu-sim/gpu-cache.h
+++ b/src/gpgpu-sim/gpu-cache.h
@@ -119,9 +119,9 @@ struct cache_block_t {
     virtual enum cache_block_state get_status( mem_access_sector_mask_t sector_mask) = 0;
     virtual void set_status(enum cache_block_state m_status, mem_access_sector_mask_t sector_mask) = 0;
 
-    virtual unsigned get_last_access_time() = 0;
-    virtual void set_last_access_time(unsigned time, mem_access_sector_mask_t sector_mask) = 0;
-    virtual unsigned get_alloc_time() = 0;
+    virtual unsigned long long get_last_access_time() = 0;
+    virtual void set_last_access_time(unsigned long long time, mem_access_sector_mask_t sector_mask) = 0;
+    virtual unsigned long long get_alloc_time() = 0;
     virtual void set_ignore_on_fill(bool m_ignore, mem_access_sector_mask_t sector_mask) = 0;
     virtual void set_modified_on_fill(bool m_modified, mem_access_sector_mask_t sector_mask) = 0;
     virtual unsigned get_modified_size() = 0;
@@ -192,15 +192,15 @@ struct line_cache_block: public cache_block_t  {
 	    {
 	    	m_status = status;
 	    }
-		virtual unsigned get_last_access_time()
+		virtual unsigned long long get_last_access_time()
 		{
 			return m_last_access_time;
 		}
-		virtual void set_last_access_time(unsigned time, mem_access_sector_mask_t sector_mask)
+		virtual void set_last_access_time(unsigned long long time, mem_access_sector_mask_t sector_mask)
 	    {
 	    	m_last_access_time = time;
 	    }
-		virtual unsigned get_alloc_time()
+		virtual unsigned long long get_alloc_time()
 	    {
 	    	return m_alloc_time;
 	    }
@@ -229,9 +229,9 @@ struct line_cache_block: public cache_block_t  {
 
 
 private:
-	    unsigned         m_alloc_time;
-	    unsigned         m_last_access_time;
-	    unsigned         m_fill_time;
+	    unsigned long long     m_alloc_time;
+	    unsigned long long     m_last_access_time;
+	    unsigned long long     m_fill_time;
 	    cache_block_state    m_status;
 	    bool m_ignore_on_fill_status;
 	    bool m_set_modified_on_fill;
@@ -364,12 +364,12 @@ struct sector_cache_block : public cache_block_t {
 		m_status[sidx] = status;
 	}
 
-    virtual unsigned get_last_access_time()
+    virtual unsigned long long get_last_access_time()
 	{
 		return m_line_last_access_time;
 	}
 
-    virtual void set_last_access_time(unsigned time, mem_access_sector_mask_t sector_mask)
+    virtual void set_last_access_time(unsigned long long time, mem_access_sector_mask_t sector_mask)
 	{
 		unsigned sidx = get_sector_index(sector_mask);
 
@@ -377,7 +377,7 @@ struct sector_cache_block : public cache_block_t {
 		m_line_last_access_time = time;
 	}
 
-    virtual unsigned get_alloc_time()
+    virtual unsigned long long get_alloc_time()
 	{
 		return m_line_alloc_time;
 	}
@@ -915,10 +915,10 @@ private:
 /// Simple struct to maintain cache accesses, misses, pending hits, and reservation fails.
 ///
 struct cache_sub_stats{
-    unsigned accesses;
-    unsigned misses;
-    unsigned pending_hits;
-    unsigned res_fails;
+    unsigned long long accesses;
+    unsigned long long misses;
+    unsigned long long pending_hits;
+    unsigned long long res_fails;
 
     unsigned long long port_available_cycles; 
     unsigned long long data_port_busy_cycles; 
@@ -981,14 +981,14 @@ public:
     void inc_stats(int access_type, int access_outcome);
     void inc_fail_stats(int access_type, int fail_outcome);
     enum cache_request_status select_stats_status(enum cache_request_status probe, enum cache_request_status access) const;
-    unsigned &operator()(int access_type, int access_outcome, bool fail_outcome);
-    unsigned operator()(int access_type, int access_outcome, bool fail_outcome) const;
+    unsigned long long &operator()(int access_type, int access_outcome, bool fail_outcome);
+    unsigned long long operator()(int access_type, int access_outcome, bool fail_outcome) const;
     cache_stats operator+(const cache_stats &cs);
     cache_stats &operator+=(const cache_stats &cs);
     void print_stats(FILE *fout, const char *cache_name = "Cache_stats") const;
     void print_fail_stats(FILE *fout, const char *cache_name = "Cache_fail_stats") const;
 
-    unsigned get_stats(enum mem_access_type *access_type, unsigned num_access_type, enum cache_request_status *access_status, unsigned num_access_status)  const;
+    unsigned long long get_stats(enum mem_access_type *access_type, unsigned num_access_type, enum cache_request_status *access_status, unsigned num_access_status)  const;
     void get_sub_stats(struct cache_sub_stats &css) const;
 
     void sample_cache_port_utility(bool data_port_busy, bool fill_port_busy); 
@@ -996,8 +996,8 @@ private:
     bool check_valid(int type, int status) const;
     bool check_fail_valid(int type, int fail) const;
 
-    std::vector< std::vector<unsigned> > m_stats;
-    std::vector< std::vector<unsigned> > m_fail_stats;
+    std::vector< std::vector<unsigned long long> > m_stats;
+    std::vector< std::vector<unsigned long long> > m_fail_stats;
 
     unsigned long long m_cache_port_available_cycles; 
     unsigned long long m_cache_data_port_busy_cycles; 
@@ -1212,12 +1212,13 @@ public:
     data_cache( const char *name, cache_config &config,
     			int core_id, int type_id, mem_fetch_interface *memport,
                 mem_fetch_allocator *mfcreator, enum mem_fetch_status status,
-                mem_access_type wr_alloc_type, mem_access_type wrbk_type )
+                mem_access_type wr_alloc_type, mem_access_type wrbk_type, class gpgpu_sim* gpu )
     			: baseline_cache(name,config,core_id,type_id,memport,status)
     {
         init( mfcreator );
         m_wr_alloc_type = wr_alloc_type;
         m_wrbk_type = wrbk_type;
+        m_gpu=gpu;
     }
 
     virtual ~data_cache() {}
@@ -1275,16 +1276,19 @@ protected:
                 enum mem_fetch_status status,
                 tag_array* new_tag_array,
                 mem_access_type wr_alloc_type,
-                mem_access_type wrbk_type)
+                mem_access_type wrbk_type,
+				class gpgpu_sim* gpu )
     : baseline_cache(name, config, core_id, type_id, memport,status, new_tag_array)
     {
         init( mfcreator );
         m_wr_alloc_type = wr_alloc_type;
         m_wrbk_type = wrbk_type;
+        m_gpu=gpu;
     }
 
     mem_access_type m_wr_alloc_type; // Specifies type of write allocate request (e.g., L1 or L2)
     mem_access_type m_wrbk_type; // Specifies type of writeback request (e.g., L1 or L2)
+    class gpgpu_sim* m_gpu;
 
     //! A general function that takes the result of a tag_array probe
     //  and performs the correspding functions based on the cache configuration
@@ -1441,8 +1445,8 @@ class l1_cache : public data_cache {
 public:
     l1_cache(const char *name, cache_config &config,
             int core_id, int type_id, mem_fetch_interface *memport,
-            mem_fetch_allocator *mfcreator, enum mem_fetch_status status )
-            : data_cache(name,config,core_id,type_id,memport,mfcreator,status, L1_WR_ALLOC_R, L1_WRBK_ACC){}
+            mem_fetch_allocator *mfcreator, enum mem_fetch_status status, class gpgpu_sim* gpu )
+            : data_cache(name,config,core_id,type_id,memport,mfcreator,status, L1_WR_ALLOC_R, L1_WRBK_ACC, gpu){}
 
     virtual ~l1_cache(){}
 
@@ -1460,10 +1464,11 @@ protected:
               mem_fetch_interface *memport,
               mem_fetch_allocator *mfcreator,
               enum mem_fetch_status status,
-              tag_array* new_tag_array )
+              tag_array* new_tag_array,
+			  class gpgpu_sim* gpu)
     : data_cache( name,
                   config,
-                  core_id,type_id,memport,mfcreator,status, new_tag_array, L1_WR_ALLOC_R, L1_WRBK_ACC ){}
+                  core_id,type_id,memport,mfcreator,status, new_tag_array, L1_WR_ALLOC_R, L1_WRBK_ACC, gpu ){}
 
 };
 
@@ -1473,8 +1478,8 @@ class l2_cache : public data_cache {
 public:
     l2_cache(const char *name,  cache_config &config,
             int core_id, int type_id, mem_fetch_interface *memport,
-            mem_fetch_allocator *mfcreator, enum mem_fetch_status status )
-            : data_cache(name,config,core_id,type_id,memport,mfcreator,status, L2_WR_ALLOC_R, L2_WRBK_ACC){}
+            mem_fetch_allocator *mfcreator, enum mem_fetch_status status, class gpgpu_sim* gpu )
+            : data_cache(name,config,core_id,type_id,memport,mfcreator,status, L2_WR_ALLOC_R, L2_WRBK_ACC, gpu){}
 
     virtual ~l2_cache() {}
 
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index ec570bf..6f19640 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -82,22 +82,7 @@ class  gpgpu_sim_wrapper {};
 
 
 bool g_interactive_debugger_enabled=false;
-
-unsigned long long  gpu_sim_cycle = 0;
-unsigned long long  gpu_tot_sim_cycle = 0;
-
-
-// performance counter for stalls due to congestion.
-unsigned int gpu_stall_dramfull = 0; 
-unsigned int gpu_stall_icnt2sh = 0;
-unsigned long long partiton_reqs_in_parallel = 0;
-unsigned long long partiton_reqs_in_parallel_total = 0;
-unsigned long long partiton_reqs_in_parallel_util = 0;
-unsigned long long partiton_reqs_in_parallel_util_total = 0;
-unsigned long long  gpu_sim_cycle_parition_util = 0;
-unsigned long long  gpu_tot_sim_cycle_parition_util = 0;
-unsigned long long partiton_replys_in_parallel = 0;
-unsigned long long partiton_replys_in_parallel_total = 0;
+tr1_hash_map<new_addr_type,unsigned> address_random_interleaving;
 
 /* Clock Domains */
 
@@ -110,8 +95,6 @@ unsigned long long partiton_replys_in_parallel_total = 0;
 #define MEM_LATENCY_STAT_IMPL
 
 
-
-
 #include "mem_latency_stat.h"
 
 void power_config::reg_options(class OptionParser * opp)
@@ -315,8 +298,8 @@ void shader_core_config::reg_options(class OptionParser * opp)
     option_parser_register(opp, "-gpgpu_shmem_size", OPT_UINT32, &gpgpu_shmem_size,
                  "Size of shared memory per shader core (default 16kB)",
                  "16384");
-    option_parser_register(opp, "-adpative_volta_cache_config", OPT_BOOL, &adpative_volta_cache_config,
-                 "adpative_volta_cache_config",
+    option_parser_register(opp, "-adaptive_volta_cache_config", OPT_BOOL, &adaptive_volta_cache_config,
+                 "adaptive_volta_cache_config",
                  "0");
     option_parser_register(opp, "-gpgpu_shmem_size", OPT_UINT32, &gpgpu_shmem_sizeDefault,
                  "Size of shared memory per shader core (default 16kB)",
@@ -476,6 +459,7 @@ void shader_core_config::reg_options(class OptionParser * opp)
     option_parser_register(opp, "-gpgpu_concurrent_kernel_sm", OPT_BOOL, &gpgpu_concurrent_kernel_sm, 
                 "Support concurrent kernels on a SM (default = disabled)", 
                 "0");
+
 }
 
 void gpgpu_sim_config::reg_options(option_parser_t opp)
@@ -499,6 +483,12 @@ void gpgpu_sim_config::reg_options(option_parser_t opp)
    option_parser_register(opp, "-liveness_message_freq", OPT_INT64, &liveness_message_freq, 
                "Minimum number of seconds between simulation liveness messages (0 = always print)",
                "1");
+   option_parser_register(opp, "-gpgpu_compute_capability_major", OPT_UINT32, &gpgpu_compute_capability_major,
+                 "Major compute capability version number",
+                 "7");
+   option_parser_register(opp, "-gpgpu_compute_capability_minor", OPT_UINT32, &gpgpu_compute_capability_minor,
+                 "Minor compute capability version number",
+                 "0");
    option_parser_register(opp, "-gpgpu_flush_l1_cache", OPT_BOOL, &gpgpu_flush_l1_cache,
                 "Flush L1 cache at the end of each kernel call",
                 "0");
@@ -532,6 +522,14 @@ void gpgpu_sim_config::reg_options(option_parser_t opp)
    option_parser_register(opp, "-visualizer_zlevel", OPT_INT32,
                           &g_visualizer_zlevel, "Compression level of the visualizer output log (0=no comp, 9=highest)",
                           "6");
+   option_parser_register(opp, "-gpgpu_stack_size_limit", OPT_INT32, &stack_size_limit,
+                          "GPU thread stack size", "1024" );
+   option_parser_register(opp, "-gpgpu_heap_size_limit", OPT_INT32, &heap_size_limit,
+                          "GPU malloc heap size ", "8388608" );
+   option_parser_register(opp, "-gpgpu_runtime_sync_depth_limit", OPT_INT32, &runtime_sync_depth_limit,
+                          "GPU device runtime synchronize depth", "2" );
+   option_parser_register(opp, "-gpgpu_runtime_pending_launch_count_limit", OPT_INT32, &runtime_pending_launch_count_limit,
+                          "GPU device runtime pending launch count", "2048" );
     option_parser_register(opp, "-trace_enabled", OPT_BOOL, 
                           &Trace::enabled, "Turn on traces",
                           "0");
@@ -714,7 +712,7 @@ gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config )
 #endif
 
     m_shader_stats = new shader_core_stats(m_shader_config);
-    m_memory_stats = new memory_stats_t(m_config.num_shader(),m_shader_config,m_memory_config);
+    m_memory_stats = new memory_stats_t(m_config.num_shader(),m_shader_config,m_memory_config,this);
     average_pipeline_duty_cycle = (float *)malloc(sizeof(float));
     active_sms=(float *)malloc(sizeof(float));
     m_power_stats = new power_stat_t(m_shader_config,average_pipeline_duty_cycle,active_sms,m_shader_stats,m_memory_config,m_memory_stats);
@@ -725,6 +723,16 @@ gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config )
     m_total_cta_launched = 0;
     gpu_deadlock = false;
 
+    gpu_stall_dramfull = 0;
+    gpu_stall_icnt2sh = 0;
+    partiton_reqs_in_parallel = 0;
+    partiton_reqs_in_parallel_total = 0;
+    partiton_reqs_in_parallel_util = 0;
+    partiton_reqs_in_parallel_util_total = 0;
+    gpu_sim_cycle_parition_util = 0;
+    gpu_tot_sim_cycle_parition_util = 0;
+    partiton_replys_in_parallel = 0;
+    partiton_replys_in_parallel_total = 0;
 
     m_cluster = new simt_core_cluster*[m_shader_config->n_simt_clusters];
     for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++) 
@@ -733,7 +741,7 @@ gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config )
     m_memory_partition_unit = new memory_partition_unit*[m_memory_config->m_n_mem];
     m_memory_sub_partition = new memory_sub_partition*[m_memory_config->m_n_mem_sub_partition];
     for (unsigned i=0;i<m_memory_config->m_n_mem;i++) {
-        m_memory_partition_unit[i] = new memory_partition_unit(i, m_memory_config, m_memory_stats);
+        m_memory_partition_unit[i] = new memory_partition_unit(i, m_memory_config, m_memory_stats, this);
         for (unsigned p = 0; p < m_memory_config->m_n_sub_partition_per_memory_channel; p++) {
             unsigned submpid = i * m_memory_config->m_n_sub_partition_per_memory_channel + p; 
             m_memory_sub_partition[submpid] = m_memory_partition_unit[i]->get_sub_partition(p); 
@@ -794,6 +802,16 @@ void gpgpu_sim::set_prop( cudaDeviceProp *prop )
    m_cuda_properties = prop;
 }
 
+int gpgpu_sim::compute_capability_major() const
+{
+   return m_config.gpgpu_compute_capability_major;
+}
+
+int gpgpu_sim::compute_capability_minor() const
+{
+   return m_config.gpgpu_compute_capability_minor;
+}
+
 const struct cudaDeviceProp *gpgpu_sim::get_prop() const
 {
    return m_cuda_properties;
@@ -1157,19 +1175,19 @@ void gpgpu_sim::gpu_print_stat()
            m_memory_sub_partition[i]->accumulate_L2cache_stats(l2_stats);
            m_memory_sub_partition[i]->get_L2cache_sub_stats(l2_css);
 
-           fprintf( stdout, "L2_cache_bank[%d]: Access = %u, Miss = %u, Miss_rate = %.3lf, Pending_hits = %u, Reservation_fails = %u\n",
+           fprintf( stdout, "L2_cache_bank[%d]: Access = %llu, Miss = %llu, Miss_rate = %.3lf, Pending_hits = %llu, Reservation_fails = %llu\n",
                     i, l2_css.accesses, l2_css.misses, (double)l2_css.misses / (double)l2_css.accesses, l2_css.pending_hits, l2_css.res_fails);
 
            total_l2_css += l2_css;
        }
        if (!m_memory_config->m_L2_config.disabled() && m_memory_config->m_L2_config.get_num_lines()) {
           //L2c_print_cache_stat();
-          printf("L2_total_cache_accesses = %u\n", total_l2_css.accesses);
-          printf("L2_total_cache_misses = %u\n", total_l2_css.misses);
+          printf("L2_total_cache_accesses = %llu\n", total_l2_css.accesses);
+          printf("L2_total_cache_misses = %llu\n", total_l2_css.misses);
           if(total_l2_css.accesses > 0)
               printf("L2_total_cache_miss_rate = %.4lf\n", (double)total_l2_css.misses/(double)total_l2_css.accesses);
-          printf("L2_total_cache_pending_hits = %u\n", total_l2_css.pending_hits);
-          printf("L2_total_cache_reservation_fails = %u\n", total_l2_css.res_fails);
+          printf("L2_total_cache_pending_hits = %llu\n", total_l2_css.pending_hits);
+          printf("L2_total_cache_reservation_fails = %llu\n", total_l2_css.res_fails);
           printf("L2_total_cache_breakdown:\n");
           l2_stats.print_stats(stdout, "L2_cache_stats_breakdown");
           printf("L2_total_cache_reservation_fail_breakdown:\n");
@@ -1477,7 +1495,7 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel )
 
     shader_CTA_count_log(m_sid, 1);
     SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: cta:%2u, start_tid:%4u, end_tid:%4u, initialized @(%lld,%lld)\n", 
-        free_cta_hw_id, start_thread, end_thread, gpu_sim_cycle, gpu_tot_sim_cycle );
+        free_cta_hw_id, start_thread, end_thread, m_gpu->gpu_sim_cycle, m_gpu->gpu_tot_sim_cycle );
 
 }
 
@@ -1694,7 +1712,7 @@ void gpgpu_sim::cycle()
             for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++) {
                 m_cluster[i]->get_current_occupancy(active, total);
             }
-            DPRINTF(LIVENESS, "uArch: inst.: %lld (ipc=%4.1f, occ=%0.4f\% [%llu / %llu]) sim_rate=%u (inst/sec) elapsed = %u:%u:%02u:%02u / %s", 
+            DPRINTFG(LIVENESS, "uArch: inst.: %lld (ipc=%4.1f, occ=%0.4f\% [%llu / %llu]) sim_rate=%u (inst/sec) elapsed = %u:%u:%02u:%02u / %s",
                    gpu_tot_sim_insn + gpu_sim_insn, 
                    (double)gpu_sim_insn/(double)gpu_sim_cycle,
                    float(active)/float(total) * 100, active, total,
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index 6ce5524..5ea5765 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -62,9 +62,7 @@
 #define SAMPLELOG 222
 #define DUMPLOG 333
 
-
-
-
+extern tr1_hash_map<new_addr_type,unsigned> address_random_interleaving;
 
 enum dram_ctrl_t {
    DRAM_FIFO=0,
@@ -293,9 +291,6 @@ struct memory_config {
    bool m_perf_sim_memcpy;
 };
 
-// global counters and flags (please try not to add to this list!!!)
-extern unsigned long long  gpu_sim_cycle;
-extern unsigned long long  gpu_tot_sim_cycle;
 extern bool g_interactive_debugger_enabled;
 
 class gpgpu_sim_config : public power_config, public gpgpu_functional_sim_config {
@@ -337,6 +332,11 @@ public:
     unsigned get_max_concurrent_kernel() const { return max_concurrent_kernel; }
     unsigned checkpoint_option;
 
+    size_t stack_limit() const {return stack_size_limit; }
+    size_t heap_limit() const {return heap_size_limit; }
+    size_t sync_depth_limit() const {return runtime_sync_depth_limit; }
+    size_t pending_launch_count_limit() const {return runtime_pending_launch_count_limit;}
+
 private:
     void init_clock_domains(void ); 
 
@@ -377,8 +377,15 @@ private:
     int gpu_stat_sample_freq;
     int gpu_runtime_stat_flag;
 
+    // Device Limits
+    size_t stack_size_limit;
+    size_t heap_size_limit;
+    size_t runtime_sync_depth_limit;
+    size_t runtime_pending_launch_count_limit;	
 
-
+ //gpu compute capability options
+    unsigned int gpgpu_compute_capability_major;
+    unsigned int gpgpu_compute_capability_minor;
     unsigned long long liveness_message_freq; 
 
     friend class gpgpu_sim;
@@ -438,6 +445,8 @@ public:
 
    int shared_mem_size() const;
    int shared_mem_per_block() const;
+   int compute_capability_major() const;
+   int compute_capability_minor() const;
    int num_registers_per_core() const;
    int num_registers_per_block() const;
    int wrp_size() const;
@@ -553,6 +562,18 @@ public:
    occupancy_stats gpu_occupancy;
    occupancy_stats gpu_tot_occupancy;
 
+   // performance counter for stalls due to congestion.
+   unsigned int gpu_stall_dramfull;
+   unsigned int gpu_stall_icnt2sh;
+   unsigned long long partiton_reqs_in_parallel;
+   unsigned long long partiton_reqs_in_parallel_total;
+   unsigned long long partiton_reqs_in_parallel_util;
+   unsigned long long partiton_reqs_in_parallel_util_total;
+   unsigned long long  gpu_sim_cycle_parition_util;
+   unsigned long long  gpu_tot_sim_cycle_parition_util;
+   unsigned long long partiton_replys_in_parallel;
+   unsigned long long partiton_replys_in_parallel_total;
+
 
    FuncCache get_cache_config(std::string kernel_name);
    void set_cache_config(std::string kernel_name, FuncCache cacheConfig );
diff --git a/src/gpgpu-sim/icnt_wrapper.cc b/src/gpgpu-sim/icnt_wrapper.cc
index ee58ece..6e0950c 100644
--- a/src/gpgpu-sim/icnt_wrapper.cc
+++ b/src/gpgpu-sim/icnt_wrapper.cc
@@ -29,6 +29,8 @@
 #include <assert.h>
 #include "../intersim2/globals.hpp"
 #include "../intersim2/interconnect_interface.hpp"
+#include "local_interconnect.h"
+
 
 icnt_create_p                icnt_create;
 icnt_init_p                  icnt_init;
@@ -42,9 +44,13 @@ icnt_display_overall_stats_p icnt_display_overall_stats;
 icnt_display_state_p         icnt_display_state;
 icnt_get_flit_size_p         icnt_get_flit_size;
 
-int   g_network_mode;
+unsigned   g_network_mode;
 char* g_network_config_filename;
 
+
+struct inct_config g_inct_config;
+LocalInterconnect *g_localicnt_interface;
+
 #include "../option_parser.h"
 
 // Wrapper to intersim2 to accompany old icnt_wrapper
@@ -105,10 +111,78 @@ static unsigned intersim2_get_flit_size()
    return g_icnt_interface->GetFlitSize();
 }
 
+
+//////////////////////////////////////////////////////
+
+static void LocalInterconnect_create(unsigned int n_shader, unsigned int n_mem)
+{
+   g_localicnt_interface->CreateInterconnect(n_shader, n_mem);
+}
+
+static void LocalInterconnect_init()
+{
+   g_localicnt_interface->Init();
+}
+
+static bool LocalInterconnect_has_buffer(unsigned input, unsigned int size)
+{
+   return g_localicnt_interface->HasBuffer(input, size);
+}
+
+static void LocalInterconnect_push(unsigned input, unsigned output, void* data, unsigned int size)
+{
+   g_localicnt_interface->Push(input, output, data, size);
+}
+
+static void* LocalInterconnect_pop(unsigned output)
+{
+   return g_localicnt_interface->Pop(output);
+}
+
+static void LocalInterconnect_transfer()
+{
+   g_localicnt_interface->Advance();
+}
+
+static bool LocalInterconnect_busy()
+{
+   return g_localicnt_interface->Busy();
+}
+
+static void LocalInterconnect_display_stats()
+{
+   g_localicnt_interface->DisplayStats();
+}
+
+static void LocalInterconnect_display_overall_stats()
+{
+   g_localicnt_interface->DisplayOverallStats();
+}
+
+static void LocalInterconnect_display_state(FILE *fp)
+{
+   g_localicnt_interface->DisplayState(fp);
+}
+
+static unsigned LocalInterconnect_get_flit_size()
+{
+   return g_localicnt_interface->GetFlitSize();
+}
+
+
+///////////////////////////
+
 void icnt_reg_options( class OptionParser * opp )
 {
    option_parser_register(opp, "-network_mode", OPT_INT32, &g_network_mode, "Interconnection network mode", "1");
    option_parser_register(opp, "-inter_config_file", OPT_CSTR, &g_network_config_filename, "Interconnection network config file", "mesh");
+
+
+   //parameters for local xbar
+   option_parser_register(opp, "-inct_in_buffer_limit", OPT_UINT32, &g_inct_config.in_buffer_limit, "in_buffer_limit", "64");
+   option_parser_register(opp, "-inct_out_buffer_limit", OPT_UINT32, &g_inct_config.out_buffer_limit, "out_buffer_limit", "64");
+   option_parser_register(opp, "-inct_subnets", OPT_UINT32, &g_inct_config.subnets, "subnets", "2");
+
 }
 
 void icnt_wrapper_init()
@@ -129,6 +203,20 @@ void icnt_wrapper_init()
          icnt_display_state = intersim2_display_state;
          icnt_get_flit_size = intersim2_get_flit_size;
          break;
+      case LOCAL_XBAR:
+         g_localicnt_interface = LocalInterconnect::New(g_inct_config);
+         icnt_create     = LocalInterconnect_create;
+         icnt_init       = LocalInterconnect_init;
+         icnt_has_buffer = LocalInterconnect_has_buffer;
+         icnt_push       = LocalInterconnect_push;
+         icnt_pop        = LocalInterconnect_pop;
+         icnt_transfer   = LocalInterconnect_transfer;
+         icnt_busy       = LocalInterconnect_busy;
+         icnt_display_stats = LocalInterconnect_display_stats;
+         icnt_display_overall_stats = LocalInterconnect_display_overall_stats;
+         icnt_display_state = LocalInterconnect_display_state;
+         icnt_get_flit_size = LocalInterconnect_get_flit_size;
+         break;
       default:
          assert(0);
          break;
diff --git a/src/gpgpu-sim/icnt_wrapper.h b/src/gpgpu-sim/icnt_wrapper.h
index a4d123e..e1086f9 100644
--- a/src/gpgpu-sim/icnt_wrapper.h
+++ b/src/gpgpu-sim/icnt_wrapper.h
@@ -57,13 +57,16 @@ extern icnt_display_stats_p icnt_display_stats;
 extern icnt_display_overall_stats_p icnt_display_overall_stats;
 extern icnt_display_state_p icnt_display_state;
 extern icnt_get_flit_size_p icnt_get_flit_size;
-extern int g_network_mode;
+extern unsigned g_network_mode;
 
 enum network_mode {
    INTERSIM = 1,
+   LOCAL_XBAR = 2,
    N_NETWORK_MODE
 };
 
+
+
 void icnt_wrapper_init();
 void icnt_reg_options( class OptionParser * opp );
 
diff --git a/src/gpgpu-sim/l2cache.cc b/src/gpgpu-sim/l2cache.cc
index 25da107..526e999 100644
--- a/src/gpgpu-sim/l2cache.cc
+++ b/src/gpgpu-sim/l2cache.cc
@@ -46,7 +46,7 @@
 #include "l2cache_trace.h"
 
 
-mem_fetch * partition_mf_allocator::alloc(new_addr_type addr, mem_access_type type, unsigned size, bool wr ) const 
+mem_fetch * partition_mf_allocator::alloc(new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const
 {
     assert( wr );
     mem_access_t access( type, addr, size, wr );
@@ -56,22 +56,25 @@ mem_fetch * partition_mf_allocator::alloc(new_addr_type addr, mem_access_type ty
                                    -1, 
                                    -1, 
                                    -1,
-                                   m_memory_config );
+                                   m_memory_config,
+								   cycle);
     return mf;
 }
 
 memory_partition_unit::memory_partition_unit( unsigned partition_id, 
                                               const struct memory_config *config,
-                                              class memory_stats_t *stats )
-: m_id(partition_id), m_config(config), m_stats(stats), m_arbitration_metadata(config) 
+                                              class memory_stats_t *stats,
+											  class gpgpu_sim* gpu)
+: m_id(partition_id), m_config(config), m_stats(stats), m_arbitration_metadata(config), m_gpu(gpu)
 {
-    m_dram = new dram_t(m_id,m_config,m_stats,this);
+    m_dram = new dram_t(m_id,m_config,m_stats,this,gpu);
 
     m_sub_partition = new memory_sub_partition*[m_config->m_n_sub_partition_per_memory_channel]; 
     for (unsigned p = 0; p < m_config->m_n_sub_partition_per_memory_channel; p++) {
         unsigned sub_partition_id = m_id * m_config->m_n_sub_partition_per_memory_channel + p; 
-        m_sub_partition[p] = new memory_sub_partition(sub_partition_id, m_config, stats); 
+        m_sub_partition[p] = new memory_sub_partition(sub_partition_id, m_config, stats, gpu);
     }
+
 }
 
 void memory_partition_unit::handle_memcpy_to_gpu( size_t addr, unsigned global_subpart_id, mem_access_sector_mask_t mask )
@@ -80,7 +83,7 @@ void memory_partition_unit::handle_memcpy_to_gpu( size_t addr, unsigned global_s
     std::string mystring =
         mask.to_string<char,std::string::traits_type,std::string::allocator_type>();
     MEMPART_DPRINTF("Copy Engine Request Received For Address=%llx, local_subpart=%u, global_subpart=%u, sector_mask=%s \n", addr, p, global_subpart_id, mystring.c_str()); 
-    m_sub_partition[p]->force_l2_tag_update(addr,gpu_sim_cycle+gpu_tot_sim_cycle, mask);
+    m_sub_partition[p]->force_l2_tag_update(addr,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle, mask);
 }
 
 memory_partition_unit::~memory_partition_unit() 
@@ -218,7 +221,7 @@ void memory_partition_unit::dram_cycle()
                 delete mf_return;
             } else {
                 m_sub_partition[dest_spid]->dram_L2_queue_push(mf_return);
-                mf_return->set_status(IN_PARTITION_DRAM_TO_L2_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle);
+                mf_return->set_status(IN_PARTITION_DRAM_TO_L2_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
                 m_arbitration_metadata.return_credit(dest_spid); 
                 MEMPART_DPRINTF("mem_fetch request %p return from dram to sub partition %d\n", mf_return, dest_spid); 
             }
@@ -247,9 +250,9 @@ void memory_partition_unit::dram_cycle()
                 MEMPART_DPRINTF("Issue mem_fetch request %p from sub partition %d to dram\n", mf, spid); 
                 dram_delay_t d;
                 d.req = mf;
-                d.ready_cycle = gpu_sim_cycle+gpu_tot_sim_cycle + m_config->dram_latency;
+                d.ready_cycle = m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle + m_config->dram_latency;
                 m_dram_latency_queue.push_back(d);
-                mf->set_status(IN_PARTITION_DRAM_LATENCY_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle);
+                mf->set_status(IN_PARTITION_DRAM_LATENCY_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
                 m_arbitration_metadata.borrow_credit(spid); 
                 break;  // the DRAM should only accept one request per cycle 
             }
@@ -258,7 +261,7 @@ void memory_partition_unit::dram_cycle()
 
     // DRAM latency queue
 
-    if( !m_dram_latency_queue.empty() && ( (gpu_sim_cycle+gpu_tot_sim_cycle) >= m_dram_latency_queue.front().ready_cycle ) && !m_dram->full(m_dram_latency_queue.front().req->is_write()) ) {
+    if( !m_dram_latency_queue.empty() && ( (m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle) >= m_dram_latency_queue.front().ready_cycle ) && !m_dram->full(m_dram_latency_queue.front().req->is_write()) ) {
     	mem_fetch* mf = m_dram_latency_queue.front().req;
     	m_dram_latency_queue.pop_front();
         m_dram->push(mf);
@@ -310,11 +313,13 @@ void memory_partition_unit::print( FILE *fp ) const
 
 memory_sub_partition::memory_sub_partition( unsigned sub_partition_id, 
                                             const struct memory_config *config,
-                                            class memory_stats_t *stats )
+                                            class memory_stats_t *stats,
+											class gpgpu_sim* gpu)
 {
     m_id = sub_partition_id;
     m_config=config;
     m_stats=stats;
+    m_gpu = gpu;
     m_memcpy_cycle_offset = 0;
 
     assert(m_id < m_config->m_n_mem_sub_partition); 
@@ -325,7 +330,7 @@ memory_sub_partition::memory_sub_partition( unsigned sub_partition_id,
     m_mf_allocator = new partition_mf_allocator(config);
 
     if(!m_config->m_L2_config.disabled())
-       m_L2cache = new l2_cache(L2c_name,m_config->m_L2_config,-1,-1,m_L2interface,m_mf_allocator,IN_PARTITION_L2_MISS_QUEUE);
+       m_L2cache = new l2_cache(L2c_name,m_config->m_L2_config,-1,-1,m_L2interface,m_mf_allocator,IN_PARTITION_L2_MISS_QUEUE, gpu);
 
     unsigned int icnt_L2;
     unsigned int L2_dram;
@@ -357,7 +362,7 @@ void memory_sub_partition::cache_cycle( unsigned cycle )
            mem_fetch *mf = m_L2cache->next_access();
            if(mf->get_access_type() != L2_WR_ALLOC_R){ // Don't pass write allocate read request back to upper level cache
 				mf->set_reply();
-				mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle);
+				mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
 				m_L2_icnt_queue->push(mf);
            }else{
         	    if(m_config->m_L2_config.m_write_alloc_policy == FETCH_ON_WRITE)
@@ -365,7 +370,7 @@ void memory_sub_partition::cache_cycle( unsigned cycle )
         	    	mem_fetch* original_wr_mf = mf->get_original_wr_mf();
 					assert(original_wr_mf);
 					original_wr_mf->set_reply();
-					original_wr_mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle);
+					original_wr_mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
 					m_L2_icnt_queue->push(original_wr_mf);
         	    }
 				m_request_tracker.erase(mf);
@@ -379,13 +384,13 @@ void memory_sub_partition::cache_cycle( unsigned cycle )
         mem_fetch *mf = m_dram_L2_queue->top();
         if ( !m_config->m_L2_config.disabled() && m_L2cache->waiting_for_fill(mf) ) {
             if (m_L2cache->fill_port_free()) {
-                mf->set_status(IN_PARTITION_L2_FILL_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle);
-                m_L2cache->fill(mf,gpu_sim_cycle+gpu_tot_sim_cycle+m_memcpy_cycle_offset);
+                mf->set_status(IN_PARTITION_L2_FILL_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
+                m_L2cache->fill(mf,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle+m_memcpy_cycle_offset);
                 m_dram_L2_queue->pop();
             }
         } else if ( !m_L2_icnt_queue->full() ) {
         	if(mf->is_write() && mf->get_type() == WRITE_ACK)
-            mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle);
+            mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
             m_L2_icnt_queue->push(mf);
             m_dram_L2_queue->pop();
         }
@@ -406,7 +411,7 @@ void memory_sub_partition::cache_cycle( unsigned cycle )
             bool port_free = m_L2cache->data_port_free(); 
             if ( !output_full && port_free ) {
                 std::list<cache_event> events;
-                enum cache_request_status status = m_L2cache->access(mf->get_addr(),mf,gpu_sim_cycle+gpu_tot_sim_cycle+m_memcpy_cycle_offset,events);
+                enum cache_request_status status = m_L2cache->access(mf->get_addr(),mf,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle+m_memcpy_cycle_offset,events);
                 bool write_sent = was_write_sent(events);
                 bool read_sent = was_read_sent(events);
                 MEM_SUBPART_DPRINTF("Probing L2 cache Address=%llx, status=%u\n", mf->get_addr(), status); 
@@ -420,7 +425,7 @@ void memory_sub_partition::cache_cycle( unsigned cycle )
                             delete mf;
                         } else {
                             mf->set_reply();
-                            mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle);
+                            mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
                             m_L2_icnt_queue->push(mf);
                         }
                         m_icnt_L2_queue->pop();
@@ -431,7 +436,7 @@ void memory_sub_partition::cache_cycle( unsigned cycle )
                 } else if ( status != RESERVATION_FAIL ) {
                 	if(mf->is_write() && (m_config->m_L2_config.m_write_alloc_policy == FETCH_ON_WRITE || m_config->m_L2_config.m_write_alloc_policy == LAZY_FETCH_ON_READ) && !was_writeallocate_sent(events)) {
                 		mf->set_reply();
-                		mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle);
+                		mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
                 		m_L2_icnt_queue->push(mf);
                 	}
                     // L2 cache accepted request
@@ -444,7 +449,7 @@ void memory_sub_partition::cache_cycle( unsigned cycle )
             }
         } else {
             // L2 is disabled or non-texture access to texture-only L2
-            mf->set_status(IN_PARTITION_L2_TO_DRAM_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle);
+            mf->set_status(IN_PARTITION_L2_TO_DRAM_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
             m_L2_dram_queue->push(mf);
             m_icnt_L2_queue->pop();
         }
@@ -455,7 +460,7 @@ void memory_sub_partition::cache_cycle( unsigned cycle )
         mem_fetch* mf = m_rop.front().req;
         m_rop.pop();
         m_icnt_L2_queue->push(mf);
-        mf->set_status(IN_PARTITION_ICNT_TO_L2_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle);
+        mf->set_status(IN_PARTITION_ICNT_TO_L2_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
     }
 }
 
@@ -635,6 +640,7 @@ std::vector<mem_fetch*> memory_sub_partition::breakdown_request_to_sector_reques
 								mf->get_sid(),
 								mf->get_tpc(),
 								mf->get_mem_config(),
+								m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle,
 								mf);
 
 			 result.push_back(n_mf);
@@ -664,13 +670,13 @@ void memory_sub_partition::push( mem_fetch* m_req, unsigned long long cycle )
 			m_request_tracker.insert(req);
 			if( req->istexture() ) {
 				m_icnt_L2_queue->push(req);
-				req->set_status(IN_PARTITION_ICNT_TO_L2_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle);
+				req->set_status(IN_PARTITION_ICNT_TO_L2_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
 			} else {
 				rop_delay_t r;
 				r.req = req;
 				r.ready_cycle = cycle + m_config->rop_latency;
 				m_rop.push(r);
-				req->set_status(IN_PARTITION_ROP_DELAY,gpu_sim_cycle+gpu_tot_sim_cycle);
+				req->set_status(IN_PARTITION_ROP_DELAY,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
 			}
     	}
     }
diff --git a/src/gpgpu-sim/l2cache.h b/src/gpgpu-sim/l2cache.h
index 18c0a8b..c8a213c 100644
--- a/src/gpgpu-sim/l2cache.h
+++ b/src/gpgpu-sim/l2cache.h
@@ -42,12 +42,12 @@ public:
     {
         m_memory_config = config;
     }
-    virtual mem_fetch * alloc(const class warp_inst_t &inst, const mem_access_t &access) const 
+    virtual mem_fetch * alloc(const class warp_inst_t &inst, const mem_access_t &access, unsigned long long cycle) const
     {
         abort();
         return NULL;
     }
-    virtual mem_fetch * alloc(new_addr_type addr, mem_access_type type, unsigned size, bool wr) const;
+    virtual mem_fetch * alloc(new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle) const;
 private:
     const memory_config *m_memory_config;
 };
@@ -58,7 +58,7 @@ private:
 class memory_partition_unit
 {
 public: 
-   memory_partition_unit( unsigned partition_id, const struct memory_config *config, class memory_stats_t *stats );
+   memory_partition_unit( unsigned partition_id, const struct memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu );
    ~memory_partition_unit(); 
 
    bool busy() const;
@@ -93,6 +93,8 @@ public:
 
    unsigned get_mpid() const { return m_id; }
 
+   class gpgpu_sim* get_mgpu() const { return m_gpu; }
+
 private: 
 
    unsigned m_id;
@@ -140,12 +142,14 @@ private:
       class mem_fetch* req;
    };
    std::list<dram_delay_t> m_dram_latency_queue;
+
+   class gpgpu_sim* m_gpu;
 };
 
 class memory_sub_partition
 {
 public:
-   memory_sub_partition( unsigned sub_partition_id, const struct memory_config *config, class memory_stats_t *stats );
+   memory_sub_partition( unsigned sub_partition_id, const struct memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu );
    ~memory_sub_partition(); 
 
    unsigned get_id() const { return m_id; } 
@@ -192,6 +196,7 @@ private:
    const struct memory_config *m_config;
    class l2_cache *m_L2cache;
    class L2interface *m_L2interface;
+   class gpgpu_sim* m_gpu;
    partition_mf_allocator *m_mf_allocator;
 
    // model delay of ROP units with a fixed latency
diff --git a/src/gpgpu-sim/l2cache_trace.h b/src/gpgpu-sim/l2cache_trace.h
index 2235cdc..d2dd948 100644
--- a/src/gpgpu-sim/l2cache_trace.h
+++ b/src/gpgpu-sim/l2cache_trace.h
@@ -42,7 +42,7 @@
 #define MEMPART_DPRINTF(...) do {\
     if (MEMPART_DTRACE(MEMORY_PARTITION_UNIT)) {\
         printf( MEMPART_PRINT_STR,\
-                gpu_sim_cycle + gpu_tot_sim_cycle,\
+                m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle,\
                 Trace::trace_streams_str[Trace::MEMORY_PARTITION_UNIT],\
                 get_mpid() );\
         printf(__VA_ARGS__);\
@@ -52,7 +52,7 @@
 #define MEM_SUBPART_DPRINTF(...) do {\
     if (MEM_SUBPART_DTRACE(MEMORY_PARTITION_UNIT)) {\
         printf( MEM_SUBPART_PRINT_STR,\
-                gpu_sim_cycle + gpu_tot_sim_cycle,\
+                m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle,\
                 Trace::trace_streams_str[Trace::MEMORY_SUBPARTITION_UNIT],\
                 m_id );\
         printf(__VA_ARGS__);\
diff --git a/src/gpgpu-sim/local_interconnect.cc b/src/gpgpu-sim/local_interconnect.cc
new file mode 100644
index 0000000..66d6648
--- /dev/null
+++ b/src/gpgpu-sim/local_interconnect.cc
@@ -0,0 +1,301 @@
+// Copyright (c) 2019, Mahmoud Khairy
+// Purdue University
+// 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 The University of British Columbia 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 AND CONTRIBUTORS "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 HOLDER 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.
+
+#include <fstream>
+#include <iostream>
+#include <sstream>
+#include <iomanip>
+#include <cmath>
+#include <utility>
+#include <algorithm>
+
+#include "local_interconnect.h"
+#include "mem_fetch.h"
+
+xbar_router::xbar_router(unsigned router_id, enum Interconnect_type m_type, unsigned n_shader, unsigned n_mem, unsigned m_in_buffer_limit, unsigned m_out_buffer_limit)
+{
+	m_id=router_id;
+	router_type=m_type;
+	_n_mem = n_mem;
+	_n_shader = n_shader;
+	total_nodes = n_shader+n_mem;
+	in_buffers.resize(total_nodes);
+	out_buffers.resize(total_nodes);
+	next_node=0;
+	in_buffer_limit = m_in_buffer_limit;
+	out_buffer_limit = m_out_buffer_limit;
+	if(m_type == REQ_NET) {
+		active_in_buffers=n_shader;
+		active_out_buffers=n_mem;
+	}
+	else if(m_type == REPLY_NET) {
+		active_in_buffers=n_mem;
+		active_out_buffers=n_shader;
+	}
+
+	cycles = 0;
+	conflicts= 0;
+	out_buffer_full=0;
+	in_buffer_full=0;
+	out_buffer_util=0;
+	in_buffer_util=0;
+	packets_num=0;
+}
+
+
+xbar_router::~xbar_router()
+{
+
+}
+
+void xbar_router::Push(unsigned input_deviceID, unsigned output_deviceID, void* data, unsigned int size)
+{
+	assert(input_deviceID < total_nodes);
+	in_buffers[input_deviceID].push(Packet(data, output_deviceID));
+	packets_num++;
+}
+
+void* xbar_router::Pop(unsigned ouput_deviceID)
+{
+	assert(ouput_deviceID < total_nodes);
+	void* data = NULL;
+
+	if(!out_buffers[ouput_deviceID].empty()) {
+		data = out_buffers[ouput_deviceID].front().data;
+		out_buffers[ouput_deviceID].pop();
+	}
+
+	return data;
+}
+
+
+bool xbar_router::Has_Buffer_In(unsigned input_deviceID, unsigned size, bool update_counter){
+
+	assert(input_deviceID < total_nodes);
+
+	bool has_buffer = (in_buffers[input_deviceID].size() + size <= in_buffer_limit);
+	if(update_counter && !has_buffer)
+		in_buffer_full++;
+
+	return has_buffer;
+
+}
+
+bool xbar_router::Has_Buffer_Out(unsigned output_deviceID, unsigned size){
+	return (out_buffers[output_deviceID].size() + size <= out_buffer_limit);
+}
+
+void xbar_router::Advance() {
+	cycles++;
+
+	vector<bool> issued(total_nodes, false);
+
+	for(unsigned i=0; i<total_nodes; ++i){
+		unsigned node_id = (i+next_node)%total_nodes;
+
+		if(!in_buffers[node_id].empty()) {
+			Packet _packet = in_buffers[node_id].front();
+			//ensure that the outbuffer has space and not issued before in this cycle
+			if(Has_Buffer_Out(_packet.output_deviceID, 1)){
+				if(!issued[_packet.output_deviceID]) {
+					out_buffers[_packet.output_deviceID].push(_packet);
+					in_buffers[node_id].pop();
+					issued[_packet.output_deviceID]=true;
+				}
+				else
+					conflicts++;
+			}
+			else
+				out_buffer_full++;
+		}
+	}
+
+	next_node = (++next_node % total_nodes);
+
+	//collect some stats about buffer util
+	for(unsigned i=0; i<total_nodes; ++i){
+		in_buffer_util+=in_buffers[i].size();
+		out_buffer_util+=out_buffers[i].size();
+	}
+}
+
+bool xbar_router::Busy() const {
+
+	for(unsigned i=0; i<total_nodes; ++i){
+		if(!in_buffers[i].empty())
+			return true;
+
+		if(!out_buffers[i].empty())
+			return true;
+	}
+	return false;
+}
+
+
+////////////////////////////////////////////////////
+/////////////LocalInterconnect/////////////////////
+
+//assume all the packets are one flit
+#define LOCAL_INCT_FLIT_SIZE 40
+
+LocalInterconnect* LocalInterconnect::New(const struct inct_config& m_localinct_config)
+{
+
+	LocalInterconnect* icnt_interface = new LocalInterconnect(m_localinct_config);
+
+	return icnt_interface;
+}
+
+LocalInterconnect::LocalInterconnect(const struct inct_config& m_localinct_config): m_inct_config(m_localinct_config){
+	n_shader=0;
+	n_mem=0;
+	n_subnets = m_localinct_config.subnets;
+}
+
+LocalInterconnect::~LocalInterconnect(){
+	for (int i=0; i<m_inct_config.subnets; ++i) {
+		delete net[i];
+	}
+}
+
+void LocalInterconnect::CreateInterconnect(unsigned m_n_shader, unsigned m_n_mem){
+	n_shader = m_n_shader;
+	n_mem = m_n_mem;
+
+	net.resize(n_subnets);
+	for (unsigned i = 0; i < n_subnets; ++i) {
+		net[i] = new xbar_router( i, static_cast<Interconnect_type>(i), m_n_shader, m_n_mem, m_inct_config.in_buffer_limit, m_inct_config.out_buffer_limit );
+	}
+
+}
+
+
+void LocalInterconnect::Init() {
+	//empty
+	//there is nothing to do
+
+}
+
+void LocalInterconnect::Push(unsigned input_deviceID, unsigned output_deviceID, void* data, unsigned int size){
+
+	unsigned subnet;
+	if (n_subnets == 1) {
+		subnet = 0;
+	} else {
+		if (input_deviceID < n_shader ) {
+			subnet = 0;
+		} else {
+			subnet = 1;
+		}
+	}
+
+	// it should have free buffer
+	//assume all the packets have size of one
+	//no flits are implemented
+	assert(net[subnet]->Has_Buffer_In(input_deviceID, 1));
+
+	net[subnet]->Push(input_deviceID, output_deviceID, data, size);
+
+}
+
+void* LocalInterconnect::Pop(unsigned ouput_deviceID){
+
+	// 0-_n_shader-1 indicates reply(network 1), otherwise request(network 0)
+	int subnet = 0;
+	if (ouput_deviceID < n_shader)
+		subnet = 1;
+
+	return net[subnet]->Pop(ouput_deviceID);
+
+}
+
+void LocalInterconnect::Advance(){
+
+	for (unsigned i = 0; i < n_subnets; ++i) {
+		net[i]->Advance();
+	}
+
+}
+
+bool  LocalInterconnect::Busy() const{
+
+	for (unsigned i = 0; i < n_subnets; ++i) {
+		if(net[i]->Busy())
+			return true;
+	}
+	return false;
+}
+
+bool  LocalInterconnect::HasBuffer(unsigned deviceID, unsigned int size) const{
+
+	bool has_buffer = false;
+
+	if ((n_subnets>1) && deviceID >= n_shader) // deviceID is memory node
+		has_buffer = net[REPLY_NET]->Has_Buffer_In(deviceID, 1, true);
+	else
+		has_buffer = net[REQ_NET]->Has_Buffer_In(deviceID, 1, true);
+
+	return has_buffer;
+
+}
+
+void  LocalInterconnect::DisplayStats() const{
+
+	cout<<"Req_Network_injected_packets_num = "<<net[REQ_NET]->packets_num<<endl;
+	cout<<"Req_Network_cycles = "<<net[REQ_NET]->cycles<<endl;
+	cout<<"Req_Network_injected_packets_per_cycle = "<<(float)(net[REQ_NET]->packets_num) / (net[REQ_NET]->cycles)<<endl;
+	cout<<"Req_Network_conflicts_per_cycle = "<<(float)(net[REQ_NET]->conflicts) / (net[REQ_NET]->cycles)<<endl;
+	cout<<"Req_Network_in_buffer_full_per_cycle = "<<(float)(net[REQ_NET]->in_buffer_full) / (net[REQ_NET]->cycles)<<endl;
+	cout<<"Req_Network_in_buffer_avg_util = "<<((float)(net[REQ_NET]->in_buffer_util) / (net[REQ_NET]->cycles) / net[REQ_NET]->active_in_buffers)<<endl;
+	cout<<"Req_Network_out_buffer_full_per_cycle = "<<(float)(net[REQ_NET]->out_buffer_full) / (net[REQ_NET]->cycles)<<endl;
+	cout<<"Req_Network_out_buffer_avg_util = "<<((float)(net[REQ_NET]->out_buffer_util) / (net[REQ_NET]->cycles) / net[REQ_NET]->active_out_buffers)<<endl;
+
+	cout<<endl;
+	cout<<"Reply_Network_injected_packets_num = "<<net[REPLY_NET]->packets_num<<endl;
+	cout<<"Reply_Network_cycles = "<<net[REPLY_NET]->cycles<<endl;
+	cout<<"Reply_Network_injected_packets_per_cycle = "<<(float)(net[REPLY_NET]->packets_num) / (net[REPLY_NET]->cycles)<<endl;
+	cout<<"Reply_Network_conflicts_per_cycle = "<<(float)(net[REPLY_NET]->conflicts) / (net[REPLY_NET]->cycles)<<endl;
+	cout<<"Reply_Network_in_buffer_full_per_cycle = "<<(float)(net[REPLY_NET]->in_buffer_full) / (net[REPLY_NET]->cycles)<<endl;
+	cout<<"Reply_Network_in_buffer_avg_util = "<<((float)(net[REPLY_NET]->in_buffer_util) / (net[REPLY_NET]->cycles) / net[REPLY_NET]->active_in_buffers)<<endl;
+	cout<<"Reply_Network_out_buffer_full_per_cycle = "<<(float)(net[REPLY_NET]->out_buffer_full) / (net[REPLY_NET]->cycles)<<endl;
+	cout<<"Reply_Network_out_buffer_avg_util= "<<((float)(net[REPLY_NET]->out_buffer_util) / (net[REPLY_NET]->cycles) / net[REPLY_NET]->active_out_buffers)<<endl;
+
+}
+
+void  LocalInterconnect::DisplayOverallStats() const{
+
+}
+
+unsigned  LocalInterconnect::GetFlitSize() const{
+	return LOCAL_INCT_FLIT_SIZE;
+}
+
+void  LocalInterconnect::DisplayState(FILE* fp) const{
+
+	fprintf(fp, "GPGPU-Sim uArch: ICNT:Display State: Under implementation\n");
+}
+
diff --git a/src/gpgpu-sim/local_interconnect.h b/src/gpgpu-sim/local_interconnect.h
new file mode 100644
index 0000000..502c80d
--- /dev/null
+++ b/src/gpgpu-sim/local_interconnect.h
@@ -0,0 +1,127 @@
+// Copyright (c) 2019, Mahmoud Khairy
+// Purdue University
+// 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 The University of British Columbia 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 AND CONTRIBUTORS "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 HOLDER 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.
+
+#ifndef _LOCAL_INTERCONNECT_HPP_
+#define _LOCAL_INTERCONNECT_HPP_
+
+#include <vector>
+#include <queue>
+#include <iostream>
+#include <map>
+using namespace std;
+
+
+struct inct_config
+{
+
+	//config for local interconnect
+	unsigned   in_buffer_limit;
+	unsigned   out_buffer_limit;
+	unsigned   subnets;
+};
+
+enum Interconnect_type {
+	REQ_NET=0,
+	REPLY_NET=1
+};
+class xbar_router {
+
+public:
+	xbar_router(unsigned router_id, enum Interconnect_type m_type, unsigned n_shader, unsigned n_mem, unsigned m_in_buffer_limit, unsigned m_out_buffer_limit);
+	~xbar_router();
+	void Push(unsigned input_deviceID, unsigned output_deviceID, void* data, unsigned int size);
+	void* Pop(unsigned ouput_deviceID);
+	void Advance();
+	bool Busy() const;
+	bool Has_Buffer_In(unsigned input_deviceID, unsigned size, bool update_counter=false);
+	bool Has_Buffer_Out(unsigned output_deviceID, unsigned size);
+
+	//some stats
+	unsigned long long cycles;
+	unsigned long long conflicts;
+	unsigned long long out_buffer_full;
+	unsigned long long out_buffer_util;
+	unsigned long long in_buffer_full;
+	unsigned long long in_buffer_util;
+	unsigned long long packets_num;
+
+private:
+	struct Packet{
+		Packet(void* m_data, unsigned m_output_deviceID) {
+			data = m_data;
+			output_deviceID = m_output_deviceID;
+		}
+		void* data;
+		unsigned output_deviceID;
+	};
+	vector<queue<Packet> > in_buffers;
+	vector<queue<Packet> > out_buffers;
+	unsigned _n_shader, _n_mem, total_nodes;
+	unsigned in_buffer_limit, out_buffer_limit;
+	unsigned next_node;
+	unsigned m_id;
+	enum Interconnect_type router_type;
+	unsigned active_in_buffers,active_out_buffers;
+
+	friend class LocalInterconnect;
+
+};
+
+class LocalInterconnect {
+public:
+	LocalInterconnect(const struct inct_config& m_localinct_config);
+	~LocalInterconnect();
+	static LocalInterconnect* New(const struct inct_config& m_inct_config);
+	void CreateInterconnect(unsigned n_shader, unsigned n_mem);
+
+	//node side functions
+	void Init();
+	void Push(unsigned input_deviceID, unsigned output_deviceID, void* data, unsigned int size);
+	void* Pop(unsigned ouput_deviceID);
+	void Advance();
+	bool Busy() const;
+	bool HasBuffer(unsigned deviceID, unsigned int size) const;
+	void DisplayStats() const;
+	void DisplayOverallStats() const;
+	unsigned GetFlitSize() const;
+
+	void DisplayState(FILE* fp) const;
+
+
+protected:
+
+	const inct_config& m_inct_config;
+
+	unsigned n_shader, n_mem;
+	unsigned n_subnets;
+	vector<xbar_router *> net;
+
+};
+
+#endif
+
+
diff --git a/src/gpgpu-sim/mem_fetch.cc b/src/gpgpu-sim/mem_fetch.cc
index a260a35..c9b0484 100644
--- a/src/gpgpu-sim/mem_fetch.cc
+++ b/src/gpgpu-sim/mem_fetch.cc
@@ -40,6 +40,7 @@ mem_fetch::mem_fetch( const mem_access_t &access,
                       unsigned sid, 
                       unsigned tpc, 
                       const struct memory_config *config,
+					  unsigned long long cycle,
 					  mem_fetch *m_original_mf,
 					  mem_fetch *m_original_wr_mf)
 
@@ -58,10 +59,10 @@ mem_fetch::mem_fetch( const mem_access_t &access,
    config->m_address_mapping.addrdec_tlx(access.get_addr(),&m_raw_addr);
    m_partition_addr = config->m_address_mapping.partition_address(access.get_addr());
    m_type = m_access.is_write()?WRITE_REQUEST:READ_REQUEST;
-   m_timestamp = gpu_sim_cycle + gpu_tot_sim_cycle;
+   m_timestamp = cycle;
    m_timestamp2 = 0;
    m_status = MEM_FETCH_INITIALIZED;
-   m_status_change = gpu_sim_cycle + gpu_tot_sim_cycle;
+   m_status_change = cycle;
    m_mem_config = config;
    icnt_flit_size = config->icnt_flit_size;
    original_mf = m_original_mf;
diff --git a/src/gpgpu-sim/mem_fetch.h b/src/gpgpu-sim/mem_fetch.h
index e5efffd..4eb3a52 100644
--- a/src/gpgpu-sim/mem_fetch.h
+++ b/src/gpgpu-sim/mem_fetch.h
@@ -56,6 +56,7 @@ public:
                unsigned sid, 
                unsigned tpc, 
                const struct memory_config *config,
+			   unsigned long long cycle,
 			   mem_fetch *original_mf = NULL,
 			   mem_fetch *original_wr_mf = NULL);
    ~mem_fetch();
diff --git a/src/gpgpu-sim/mem_latency_stat.cc b/src/gpgpu-sim/mem_latency_stat.cc
index c5452b9..c7d20d1 100644
--- a/src/gpgpu-sim/mem_latency_stat.cc
+++ b/src/gpgpu-sim/mem_latency_stat.cc
@@ -42,7 +42,7 @@
 #include <stdlib.h>
 #include <stdio.h>
 
-memory_stats_t::memory_stats_t( unsigned n_shader, const struct shader_core_config *shader_config, const struct memory_config *mem_config )
+memory_stats_t::memory_stats_t( unsigned n_shader, const struct shader_core_config *shader_config, const struct memory_config *mem_config, const class gpgpu_sim* gpu )
 {
    assert( mem_config->m_valid );
    assert( shader_config->m_valid );
@@ -67,6 +67,7 @@ memory_stats_t::memory_stats_t( unsigned n_shader, const struct shader_core_conf
 
    m_n_shader=n_shader;
    m_memory_config=mem_config;
+   m_gpu=gpu;
    total_n_access=0;
    total_n_reads=0;
    total_n_writes=0;
@@ -141,7 +142,7 @@ memory_stats_t::memory_stats_t( unsigned n_shader, const struct shader_core_conf
 unsigned memory_stats_t::memlatstat_done(mem_fetch *mf )
 {
    unsigned mf_latency;
-   mf_latency = (gpu_sim_cycle+gpu_tot_sim_cycle) - mf->get_timestamp();
+   mf_latency = (m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle) - mf->get_timestamp();
    mf_num_lat_pw++;
    mf_tot_lat_pw += mf_latency;
    unsigned idx = LOGB2(mf_latency);
@@ -161,7 +162,7 @@ void memory_stats_t::memlatstat_read_done(mem_fetch *mf)
       if (mf_latency > mf_max_lat_table[mf->get_tlx_addr().chip][mf->get_tlx_addr().bk]) 
          mf_max_lat_table[mf->get_tlx_addr().chip][mf->get_tlx_addr().bk] = mf_latency;
       unsigned icnt2sh_latency;
-      icnt2sh_latency = (gpu_tot_sim_cycle+gpu_sim_cycle) - mf->get_return_timestamp();
+      icnt2sh_latency = (m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle) - mf->get_return_timestamp();
       tot_icnt2sh_latency += icnt2sh_latency;
       icnt2sh_lat_table[LOGB2(icnt2sh_latency)]++;
       if (icnt2sh_latency > max_icnt2sh_latency)
@@ -195,7 +196,7 @@ void memory_stats_t::memlatstat_icnt2mem_pop(mem_fetch *mf)
 {
    if (m_memory_config->gpgpu_memlatency_stat) {
       unsigned icnt2mem_latency;
-      icnt2mem_latency = (gpu_tot_sim_cycle+gpu_sim_cycle) - mf->get_timestamp();
+      icnt2mem_latency = (m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle) - mf->get_timestamp();
       tot_icnt2mem_latency += icnt2mem_latency;
       icnt2mem_lat_table[LOGB2(icnt2mem_latency)]++;
       if (icnt2mem_latency > max_icnt2mem_latency)
@@ -366,7 +367,7 @@ void memory_stats_t::memlatstat_print( unsigned n_mem, unsigned gpu_mem_n_bk )
          m = 0;
          printf("\n");
       }
-      printf("total reads: %d\n", k);
+      printf("total dram reads = %d\n", k);
       if (min_bank_accesses)
          printf("bank skew: %d/%d = %4.2f\n", max_bank_accesses, min_bank_accesses, (float)max_bank_accesses/min_bank_accesses);
       else
@@ -404,7 +405,7 @@ void memory_stats_t::memlatstat_print( unsigned n_mem, unsigned gpu_mem_n_bk )
          m = 0;
          printf("\n");
       }
-      printf("total reads: %d\n", k);
+      printf("total dram writes = %d\n", k);
       if (min_bank_accesses)
          printf("bank skew: %d/%d = %4.2f\n", max_bank_accesses, min_bank_accesses, (float)max_bank_accesses/min_bank_accesses);
       else
diff --git a/src/gpgpu-sim/mem_latency_stat.h b/src/gpgpu-sim/mem_latency_stat.h
index 5b89202..b86740d 100644
--- a/src/gpgpu-sim/mem_latency_stat.h
+++ b/src/gpgpu-sim/mem_latency_stat.h
@@ -36,7 +36,8 @@ class memory_stats_t {
 public:
    memory_stats_t( unsigned n_shader, 
                    const struct shader_core_config *shader_config, 
-                   const struct memory_config *mem_config );
+                   const struct memory_config *mem_config,
+				   const class gpgpu_sim* gpu);
 
    unsigned memlatstat_done( class mem_fetch *mf );
    void memlatstat_read_done( class mem_fetch *mf );
@@ -51,6 +52,7 @@ public:
 
    const struct shader_core_config *m_shader_config;
    const struct memory_config *m_memory_config;
+   const class gpgpu_sim* m_gpu;
 
    unsigned max_mrq_latency;
    unsigned max_dq_latency;
diff --git a/src/gpgpu-sim/scoreboard.cc b/src/gpgpu-sim/scoreboard.cc
index ebec891..80f95c6 100644
--- a/src/gpgpu-sim/scoreboard.cc
+++ b/src/gpgpu-sim/scoreboard.cc
@@ -32,13 +32,15 @@
 
 
 //Constructor
-Scoreboard::Scoreboard( unsigned sid, unsigned n_warps )
+Scoreboard::Scoreboard( unsigned sid, unsigned n_warps, class gpgpu_t* gpu )
 : longopregs()
 {
 	m_sid = sid;
 	//Initialize size of table
 	reg_table.resize(n_warps);
 	longopregs.resize(n_warps);
+
+	m_gpu = gpu;
 }
 
 // Print scoreboard contents
diff --git a/src/gpgpu-sim/scoreboard.h b/src/gpgpu-sim/scoreboard.h
index 4a76ea3..a4baa19 100644
--- a/src/gpgpu-sim/scoreboard.h
+++ b/src/gpgpu-sim/scoreboard.h
@@ -38,7 +38,7 @@
 
 class Scoreboard {
 public:
-    Scoreboard( unsigned sid, unsigned n_warps );
+    Scoreboard( unsigned sid, unsigned n_warps, class gpgpu_t* gpu );
 
     void reserveRegisters(const warp_inst_t *inst);
     void releaseRegisters(const warp_inst_t *inst);
@@ -59,6 +59,8 @@ private:
     std::vector< std::set<unsigned> > reg_table;
     //Register that depend on a long operation (global, local or tex memory)
     std::vector< std::set<unsigned> > longopregs;
+
+    class gpgpu_t* m_gpu;
 };
 
 
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index 3db988b..69b619a 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -133,7 +133,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu,
     m_L1I = new read_only_cache( name,m_config->m_L1I_config,m_sid,get_shader_instruction_cache_id(),m_icnt,IN_L1I_MISS_QUEUE);
     
     m_warp.resize(m_config->max_warps_per_shader, shd_warp_t(this, warp_size));
-    m_scoreboard = new Scoreboard(m_sid, m_config->max_warps_per_shader);
+    m_scoreboard = new Scoreboard(m_sid, m_config->max_warps_per_shader, gpu);
     
     //scedulers
     //must currently occur after all inputs have been initialized.
@@ -767,7 +767,7 @@ void shader_core_ctx::fetch()
             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.
             m_inst_fetch_buffer.m_valid = true;
-            m_warp[mf->get_wid()].set_last_fetch(gpu_sim_cycle);
+            m_warp[mf->get_wid()].set_last_fetch(m_gpu->gpu_sim_cycle);
             delete mf;
         }
         else {
@@ -815,17 +815,19 @@ void shader_core_ctx::fetch()
                             warp_id,
                             m_sid,
                             m_tpc,
-                            m_memory_config );
+                            m_memory_config,
+							m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle
+							);
                     std::list<cache_event> events;
-                    enum cache_request_status status = m_L1I->access( (new_addr_type)ppc, mf, gpu_sim_cycle+gpu_tot_sim_cycle,events);
+                    enum cache_request_status status = m_L1I->access( (new_addr_type)ppc, mf, m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle,events);
                     if( status == MISS ) {
                         m_last_warp_fetched=warp_id;
                         m_warp[warp_id].set_imiss_pending();
-                        m_warp[warp_id].set_last_fetch(gpu_sim_cycle);
+                        m_warp[warp_id].set_last_fetch(m_gpu->gpu_sim_cycle);
                     } else if( status == HIT ) {
                         m_last_warp_fetched=warp_id;
                         m_inst_fetch_buffer = ifetch_buffer_t(pc,nbytes,warp_id);
-                        m_warp[warp_id].set_last_fetch(gpu_sim_cycle);
+                        m_warp[warp_id].set_last_fetch(m_gpu->gpu_sim_cycle);
                         delete mf;
                     } else {
                         m_last_warp_fetched=warp_id;
@@ -859,7 +861,7 @@ void shader_core_ctx::issue_warp( register_set& pipe_reg_set, const warp_inst_t*
     m_warp[warp_id].ibuffer_free();
     assert(next_inst->valid());
     **pipe_reg = *next_inst; // static instruction information
-    (*pipe_reg)->issue( active_mask, warp_id, gpu_tot_sim_cycle + gpu_sim_cycle, m_warp[warp_id].get_dynamic_warp_id(), sch_id ); // dynamic 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 ){
@@ -1514,7 +1516,7 @@ void shader_core_ctx::warp_inst_complete(const warp_inst_t &inst)
 
   m_stats->m_num_sim_winsn[m_sid]++;
   m_gpu->gpu_sim_insn += inst.active_count();
-  inst.completed(gpu_tot_sim_cycle + gpu_sim_cycle);
+  inst.completed(m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle);
 }
 
 void shader_core_ctx::writeback()
@@ -1552,9 +1554,9 @@ void shader_core_ctx::writeback()
         m_warp[warp_id].dec_inst_in_pipeline();
         warp_inst_complete(*pipe_reg);
         m_gpu->gpu_sim_insn_last_update_sid = m_sid;
-        m_gpu->gpu_sim_insn_last_update = gpu_sim_cycle;
-        m_last_inst_gpu_sim_cycle = gpu_sim_cycle;
-        m_last_inst_gpu_tot_sim_cycle = gpu_tot_sim_cycle;
+        m_gpu->gpu_sim_insn_last_update = m_gpu->gpu_sim_cycle;
+        m_last_inst_gpu_sim_cycle = m_gpu->gpu_sim_cycle;
+        m_last_inst_gpu_tot_sim_cycle = m_gpu->gpu_tot_sim_cycle;
         pipe_reg->clear();
         preg = m_pipeline_reg[EX_WB].get_ready();
         pipe_reg = (preg==NULL)? NULL:*preg;
@@ -1633,9 +1635,9 @@ mem_stage_stall_type ldst_unit::process_memory_access_queue( cache_t *cache, war
         return DATA_PORT_STALL; 
 
     //const mem_access_t &access = inst.accessq_back();
-    mem_fetch *mf = m_mf_allocator->alloc(inst,inst.accessq_back());
+    mem_fetch *mf = m_mf_allocator->alloc(inst,inst.accessq_back(),m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle);
     std::list<cache_event> events;
-    enum cache_request_status status = cache->access(mf->get_addr(),mf,gpu_sim_cycle+gpu_tot_sim_cycle,events);
+    enum cache_request_status status = cache->access(mf->get_addr(),mf,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle,events);
     return process_cache_access( cache, mf->get_addr(), inst, events, mf, status );
 }
 
@@ -1645,7 +1647,7 @@ mem_stage_stall_type ldst_unit::process_memory_access_queue_l1cache( l1_cache *c
     if( inst.accessq_empty() )
         return result;
 
-    mem_fetch *mf = m_mf_allocator->alloc(inst,inst.accessq_back());
+    mem_fetch *mf = m_mf_allocator->alloc(inst,inst.accessq_back(),m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle);
 
     if(m_config->m_L1D_config.l1_latency > 0)
 	{
@@ -1675,7 +1677,7 @@ mem_stage_stall_type ldst_unit::process_memory_access_queue_l1cache( l1_cache *c
     else
     {
 		std::list<cache_event> events;
-		enum cache_request_status status = cache->access(mf->get_addr(),mf,gpu_sim_cycle+gpu_tot_sim_cycle,events);
+		enum cache_request_status status = cache->access(mf->get_addr(),mf,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle,events);
 		return process_cache_access( cache, mf->get_addr(), inst, events, mf, status );
     }
 }
@@ -1687,7 +1689,7 @@ void ldst_unit::L1_latency_queue_cycle()
     {
 		    mem_fetch* mf_next = l1_latency_queue[0];
 			std::list<cache_event> events;
-			enum cache_request_status status = m_L1D->access(mf_next->get_addr(),mf_next,gpu_sim_cycle+gpu_tot_sim_cycle,events);
+			enum cache_request_status status = m_L1D->access(mf_next->get_addr(),mf_next,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle,events);
 
 		   bool write_sent = was_write_sent(events);
 		   bool read_sent = was_read_sent(events);
@@ -1804,7 +1806,7 @@ bool ldst_unit::memory_cycle( warp_inst_t &inst, mem_stage_stall_type &stall_rea
        if( m_icnt->full(size, inst.is_store() || inst.isatomic()) ) {
            stall_cond = ICNT_RC_FAIL;
        } else {
-           mem_fetch *mf = m_mf_allocator->alloc(inst,access);
+           mem_fetch *mf = m_mf_allocator->alloc(inst,access,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle);
            m_icnt->push(mf);
            inst.accessq_pop_back();
            //inst.clear_active( access.get_warp_mask() );
@@ -1840,7 +1842,7 @@ bool ldst_unit::response_buffer_full() const
 
 void ldst_unit::fill( mem_fetch *mf )
 {
-    mf->set_status(IN_SHADER_LDST_RESPONSE_FIFO,gpu_sim_cycle+gpu_tot_sim_cycle);
+    mf->set_status(IN_SHADER_LDST_RESPONSE_FIFO,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle);
     m_response_fifo.push_back(mf);
 }
 
@@ -1893,6 +1895,17 @@ void tensor_core::issue( register_set& source_reg )
 	pipelined_simd_unit::issue(source_reg);
 }
 
+unsigned pipelined_simd_unit::get_active_lanes_in_pipeline(){
+	active_mask_t active_lanes;
+	active_lanes.reset();
+	 if(m_core->get_gpu()->get_config().g_power_simulation_enabled){
+		for( unsigned stage=0; (stage+1)<m_pipeline_depth; stage++ ){
+			if( !m_pipeline_reg[stage]->empty() )
+				active_lanes|=m_pipeline_reg[stage]->get_active_mask();
+		}
+	 }
+	return active_lanes.count();
+}
 
 void ldst_unit::active_lanes_in_pipeline(){
 	unsigned active_count=pipelined_simd_unit::get_active_lanes_in_pipeline();
@@ -1946,13 +1959,13 @@ sp_unit::sp_unit( register_set* result_port, const shader_core_config *config,sh
 }
 
 dp_unit::dp_unit( register_set* result_port, const shader_core_config *config,shader_core_ctx *core)
-    : pipelined_simd_unit(result_port,config,config->max_sfu_latency,core)
+    : pipelined_simd_unit(result_port,config,config->max_dp_latency,core)
 {
     m_name = "DP ";
 }
 
 int_unit::int_unit( register_set* result_port, const shader_core_config *config,shader_core_ctx *core)
-    : pipelined_simd_unit(result_port,config,config->max_sp_latency,core)
+    : pipelined_simd_unit(result_port,config,config->max_int_latency,core)
 {
     m_name = "INT ";
 }
@@ -1993,19 +2006,25 @@ pipelined_simd_unit::pipelined_simd_unit( register_set* result_port, const shade
     for( unsigned i=0; i < m_pipeline_depth; i++ ) 
 	m_pipeline_reg[i] = new warp_inst_t( config );
     m_core=core;
+    active_insts_in_pipeline=0;
 }
 
 void pipelined_simd_unit::cycle()
 {
     if( !m_pipeline_reg[0]->empty() ){
         m_result_port->move_in(m_pipeline_reg[0]);
+        assert(active_insts_in_pipeline > 0);
+        active_insts_in_pipeline--;
+    }
+    if(active_insts_in_pipeline){
+		for( unsigned stage=0; (stage+1)<m_pipeline_depth; stage++ )
+			move_warp(m_pipeline_reg[stage], m_pipeline_reg[stage+1]);
     }
-    for( unsigned stage=0; (stage+1)<m_pipeline_depth; stage++ )
-        move_warp(m_pipeline_reg[stage], m_pipeline_reg[stage+1]);
     if( !m_dispatch_reg->empty() ) {
         if( !m_dispatch_reg->dispatch_delay()){
             int start_stage = m_dispatch_reg->latency - m_dispatch_reg->initiation_interval;
             move_warp(m_pipeline_reg[start_stage],m_dispatch_reg);
+            active_insts_in_pipeline++;
         }
     }
     occupied >>=1;
@@ -2098,7 +2117,8 @@ ldst_unit::ldst_unit( mem_fetch_interface *icnt,
                               get_shader_normal_cache_id(),
                               m_icnt,
                               m_mf_allocator,
-                              IN_L1D_MISS_QUEUE );
+                              IN_L1D_MISS_QUEUE,
+							  core->get_gpu());
 
         if(m_config->m_L1D_config.l1_latency > 0)
 	    {
@@ -2185,8 +2205,8 @@ void ldst_unit::writeback()
                 m_core->warp_inst_complete(m_next_wb);
             }
             m_next_wb.clear();
-            m_last_inst_gpu_sim_cycle = gpu_sim_cycle;
-            m_last_inst_gpu_tot_sim_cycle = gpu_tot_sim_cycle;
+            m_last_inst_gpu_sim_cycle = m_core->get_gpu()->gpu_sim_cycle;
+            m_last_inst_gpu_tot_sim_cycle = m_core->get_gpu()->gpu_tot_sim_cycle;
         }
     }
 
@@ -2294,13 +2314,13 @@ void ldst_unit::cycle()
        mem_fetch *mf = m_response_fifo.front();
        if (mf->get_access_type() == TEXTURE_ACC_R) {
            if (m_L1T->fill_port_free()) {
-               m_L1T->fill(mf,gpu_sim_cycle+gpu_tot_sim_cycle);
+               m_L1T->fill(mf,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle);
                m_response_fifo.pop_front(); 
            }
        } else if (mf->get_access_type() == CONST_ACC_R)  {
            if (m_L1C->fill_port_free()) {
-               mf->set_status(IN_SHADER_FETCHED,gpu_sim_cycle+gpu_tot_sim_cycle);
-               m_L1C->fill(mf,gpu_sim_cycle+gpu_tot_sim_cycle);
+               mf->set_status(IN_SHADER_FETCHED,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle);
+               m_L1C->fill(mf,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle);
                m_response_fifo.pop_front(); 
            }
        } else {
@@ -2320,13 +2340,13 @@ void ldst_unit::cycle()
                }
                if( bypassL1D ) {
                    if ( m_next_global == NULL ) {
-                       mf->set_status(IN_SHADER_FETCHED,gpu_sim_cycle+gpu_tot_sim_cycle);
+                       mf->set_status(IN_SHADER_FETCHED,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle);
                        m_response_fifo.pop_front();
                        m_next_global = mf;
                    }
                } else {
                    if (m_L1D->fill_port_free()) {
-                       m_L1D->fill(mf,gpu_sim_cycle+gpu_tot_sim_cycle);
+                       m_L1D->fill(mf,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle);
                        m_response_fifo.pop_front();
                    }
                }
@@ -2415,7 +2435,7 @@ void shader_core_ctx::register_cta_thread_exit( unsigned cta_num, kernel_info_t
       shader_CTA_count_unlog(m_sid, 1);
 
      SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Finished CTA #%d (%lld,%lld), %u CTAs running\n",
-        cta_num, gpu_sim_cycle, gpu_tot_sim_cycle, m_n_active_cta);
+        cta_num, m_gpu->gpu_sim_cycle, m_gpu->gpu_tot_sim_cycle, m_n_active_cta);
 
       if( m_n_active_cta == 0 ) {
         SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Empty (last released kernel %u \'%s\').\n",
@@ -2525,13 +2545,13 @@ void gpgpu_sim::shader_print_cache_stats( FILE *fout ) const{
             m_cluster[i]->get_L1I_sub_stats(css);
             total_css += css;
         }
-        fprintf(fout, "\tL1I_total_cache_accesses = %u\n", total_css.accesses);
-        fprintf(fout, "\tL1I_total_cache_misses = %u\n", total_css.misses);
+        fprintf(fout, "\tL1I_total_cache_accesses = %llu\n", total_css.accesses);
+        fprintf(fout, "\tL1I_total_cache_misses = %llu\n", total_css.misses);
         if(total_css.accesses > 0){
             fprintf(fout, "\tL1I_total_cache_miss_rate = %.4lf\n", (double)total_css.misses / (double)total_css.accesses);
         }
-        fprintf(fout, "\tL1I_total_cache_pending_hits = %u\n", total_css.pending_hits);
-        fprintf(fout, "\tL1I_total_cache_reservation_fails = %u\n", total_css.res_fails);
+        fprintf(fout, "\tL1I_total_cache_pending_hits = %llu\n", total_css.pending_hits);
+        fprintf(fout, "\tL1I_total_cache_reservation_fails = %llu\n", total_css.res_fails);
     }
 
     // L1D
@@ -2542,18 +2562,18 @@ void gpgpu_sim::shader_print_cache_stats( FILE *fout ) const{
         for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++){
             m_cluster[i]->get_L1D_sub_stats(css);
 
-            fprintf( stdout, "\tL1D_cache_core[%d]: Access = %d, Miss = %d, Miss_rate = %.3lf, Pending_hits = %u, Reservation_fails = %u\n",
+            fprintf( stdout, "\tL1D_cache_core[%d]: Access = %llu, Miss = %llu, Miss_rate = %.3lf, Pending_hits = %llu, Reservation_fails = %llu\n",
                      i, css.accesses, css.misses, (double)css.misses / (double)css.accesses, css.pending_hits, css.res_fails);
 
             total_css += css;
         }
-        fprintf(fout, "\tL1D_total_cache_accesses = %u\n", total_css.accesses);
-        fprintf(fout, "\tL1D_total_cache_misses = %u\n", total_css.misses);
+        fprintf(fout, "\tL1D_total_cache_accesses = %llu\n", total_css.accesses);
+        fprintf(fout, "\tL1D_total_cache_misses = %llu\n", total_css.misses);
         if(total_css.accesses > 0){
             fprintf(fout, "\tL1D_total_cache_miss_rate = %.4lf\n", (double)total_css.misses / (double)total_css.accesses);
         }
-        fprintf(fout, "\tL1D_total_cache_pending_hits = %u\n", total_css.pending_hits);
-        fprintf(fout, "\tL1D_total_cache_reservation_fails = %u\n", total_css.res_fails);
+        fprintf(fout, "\tL1D_total_cache_pending_hits = %llu\n", total_css.pending_hits);
+        fprintf(fout, "\tL1D_total_cache_reservation_fails = %llu\n", total_css.res_fails);
         total_css.print_port_stats(fout, "\tL1D_cache"); 
     }
 
@@ -2566,13 +2586,13 @@ void gpgpu_sim::shader_print_cache_stats( FILE *fout ) const{
             m_cluster[i]->get_L1C_sub_stats(css);
             total_css += css;
         }
-        fprintf(fout, "\tL1C_total_cache_accesses = %u\n", total_css.accesses);
-        fprintf(fout, "\tL1C_total_cache_misses = %u\n", total_css.misses);
+        fprintf(fout, "\tL1C_total_cache_accesses = %llu\n", total_css.accesses);
+        fprintf(fout, "\tL1C_total_cache_misses = %llu\n", total_css.misses);
         if(total_css.accesses > 0){
             fprintf(fout, "\tL1C_total_cache_miss_rate = %.4lf\n", (double)total_css.misses / (double)total_css.accesses);
         }
-        fprintf(fout, "\tL1C_total_cache_pending_hits = %u\n", total_css.pending_hits);
-        fprintf(fout, "\tL1C_total_cache_reservation_fails = %u\n", total_css.res_fails);
+        fprintf(fout, "\tL1C_total_cache_pending_hits = %llu\n", total_css.pending_hits);
+        fprintf(fout, "\tL1C_total_cache_reservation_fails = %llu\n", total_css.res_fails);
     }
 
     // L1T
@@ -2584,13 +2604,13 @@ void gpgpu_sim::shader_print_cache_stats( FILE *fout ) const{
             m_cluster[i]->get_L1T_sub_stats(css);
             total_css += css;
         }
-        fprintf(fout, "\tL1T_total_cache_accesses = %u\n", total_css.accesses);
-        fprintf(fout, "\tL1T_total_cache_misses = %u\n", total_css.misses);
+        fprintf(fout, "\tL1T_total_cache_accesses = %llu\n", total_css.accesses);
+        fprintf(fout, "\tL1T_total_cache_misses = %llu\n", total_css.misses);
         if(total_css.accesses > 0){
             fprintf(fout, "\tL1T_total_cache_miss_rate = %.4lf\n", (double)total_css.misses / (double)total_css.accesses);
         }
-        fprintf(fout, "\tL1T_total_cache_pending_hits = %u\n", total_css.pending_hits);
-        fprintf(fout, "\tL1T_total_cache_reservation_fails = %u\n", total_css.res_fails);
+        fprintf(fout, "\tL1T_total_cache_pending_hits = %llu\n", total_css.pending_hits);
+        fprintf(fout, "\tL1T_total_cache_reservation_fails = %llu\n", total_css.res_fails);
     }
 }
 
@@ -2806,7 +2826,7 @@ void shader_core_ctx::display_pipeline(FILE *fout, int print_mem, int mask ) con
 {
    fprintf(fout, "=================================================\n");
    fprintf(fout, "shader %u at cycle %Lu+%Lu (%u threads running)\n", m_sid, 
-           gpu_tot_sim_cycle, gpu_sim_cycle, m_not_completed);
+           m_gpu->gpu_tot_sim_cycle, m_gpu->gpu_sim_cycle, m_not_completed);
    fprintf(fout, "=================================================\n");
 
    dump_warp_state(fout);
@@ -2949,7 +2969,7 @@ unsigned int shader_core_config::max_cta( const kernel_info_t &k ) const
        abort();
     }
 
-    if(adpative_volta_cache_config && !k.volta_cache_config_set) {
+    if(adaptive_volta_cache_config && !k.volta_cache_config_set) {
     	//For Volta, we assign the remaining shared memory to L1 cache
     	//For more info, see https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory-7-x
     	unsigned total_shmed = kernel_info->smem * result;
@@ -2981,8 +3001,53 @@ unsigned int shader_core_config::max_cta( const kernel_info_t &k ) const
     return result;
 }
 
+void shader_core_config::set_pipeline_latency() {
+
+		//calculate the max latency  based on the input
+
+		unsigned int_latency[5];
+		unsigned fp_latency[5];
+		unsigned dp_latency[5];
+		unsigned sfu_latency;
+		unsigned tensor_latency;
+
+			/*
+			 * [0] ADD,SUB
+			 * [1] MAX,Min
+			 * [2] MUL
+			 * [3] MAD
+			 * [4] DIV
+			 */
+			sscanf(opcode_latency_int, "%u,%u,%u,%u,%u",
+					&int_latency[0],&int_latency[1],&int_latency[2],
+					&int_latency[3],&int_latency[4]);
+			sscanf(opcode_latency_fp, "%u,%u,%u,%u,%u",
+					&fp_latency[0],&fp_latency[1],&fp_latency[2],
+					&fp_latency[3],&fp_latency[4]);
+			sscanf(opcode_latency_dp, "%u,%u,%u,%u,%u",
+					&dp_latency[0],&dp_latency[1],&dp_latency[2],
+					&dp_latency[3],&dp_latency[4]);
+			sscanf(opcode_latency_sfu, "%u",
+					&sfu_latency);
+			sscanf(opcode_latency_tensor, "%u",
+					&tensor_latency);
+
+		//all div operation are executed on sfu
+		//assume that the max latency are dp div or normal sfu_latency
+		max_sfu_latency = std::max(dp_latency[4],sfu_latency);
+		//assume that the max operation has the max latency
+		max_sp_latency = fp_latency[1];
+		max_int_latency = int_latency[1];
+		max_dp_latency = dp_latency[1];
+		max_tensor_core_latency = tensor_latency;
+
+}
+
 void shader_core_ctx::cycle()
 {
+	if(!isactive() && get_not_completed() == 0)
+		return;
+
 	m_stats->shader_cycles[m_sid]++;
     writeback();
     execute();
@@ -3159,7 +3224,7 @@ void barrier_set_t::warp_reaches_barrier(unsigned cta_id,unsigned warp_id,warp_i
    cta_to_warp_t::iterator w=m_cta_to_warps.find(cta_id);
 
    if( w == m_cta_to_warps.end() ) { // cta is active
-      printf("ERROR ** cta_id %u not found in barrier set on cycle %llu+%llu...\n", cta_id, gpu_tot_sim_cycle, gpu_sim_cycle );
+      printf("ERROR ** cta_id %u not found in barrier set on cycle %llu+%llu...\n", cta_id, m_shader->get_gpu()->gpu_tot_sim_cycle, m_shader->get_gpu()->gpu_sim_cycle );
       dump();
       abort();
    }
@@ -3332,8 +3397,8 @@ bool shader_core_ctx::fetch_unit_response_buffer_full() const
 
 void shader_core_ctx::accept_fetch_response( mem_fetch *mf )
 {
-    mf->set_status(IN_SHADER_FETCHED,gpu_sim_cycle+gpu_tot_sim_cycle);
-    m_L1I->fill(mf,gpu_sim_cycle+gpu_tot_sim_cycle);
+    mf->set_status(IN_SHADER_FETCHED,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
+    m_L1I->fill(mf,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
 }
 
 bool shader_core_ctx::ldst_unit_response_buffer_full() const
@@ -3895,7 +3960,7 @@ void simt_core_cluster::icnt_inject_request_packet(class mem_fetch *mf)
    }
    m_stats->m_outgoing_traffic_stats->record_traffic(mf, packet_size); 
    unsigned destination = mf->get_sub_partition_id();
-   mf->set_status(IN_ICNT_TO_MEM,gpu_sim_cycle+gpu_tot_sim_cycle);
+   mf->set_status(IN_ICNT_TO_MEM,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
    if (!mf->get_is_write() && !mf->isatomic())
       ::icnt_push(m_cluster_id, m_config->mem2device(destination), (void*)mf, mf->get_ctrl_size() );
    else 
@@ -3934,7 +3999,7 @@ void simt_core_cluster::icnt_cycle()
         // - For write-ack, the packet only has control metadata
         unsigned int packet_size = (mf->get_is_write())? mf->get_ctrl_size() : mf->size(); 
         m_stats->m_incoming_traffic_stats->record_traffic(mf, packet_size); 
-        mf->set_status(IN_CLUSTER_TO_SHADER_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle);
+        mf->set_status(IN_CLUSTER_TO_SHADER_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle);
         //m_memory_stats->memlatstat_read_done(mf,m_shader_config->max_warps_per_shader);
         m_response_fifo.push_back(mf);
         m_stats->n_mem_to_simt[m_cluster_id] += mf->get_num_flits(false);
diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h
index 9abd223..fde87b6 100644
--- a/src/gpgpu-sim/shader.h
+++ b/src/gpgpu-sim/shader.h
@@ -55,7 +55,6 @@
 #include "traffic_breakdown.h"
 
 
-
 #define NO_OP_FLAG            0xFF
 
 /* READ_PACKET_SIZE:
@@ -1073,16 +1072,8 @@ public:
     //modifiers
     virtual void cycle();
     virtual void issue( register_set& source_reg );
-    virtual unsigned get_active_lanes_in_pipeline()
-    {
-    	active_mask_t active_lanes;
-    	active_lanes.reset();
-        for( unsigned stage=0; (stage+1)<m_pipeline_depth; stage++ ){
-        	if( !m_pipeline_reg[stage]->empty() )
-        		active_lanes|=m_pipeline_reg[stage]->get_active_mask();
-        }
-        return active_lanes.count();
-    }
+    virtual unsigned get_active_lanes_in_pipeline();
+
     virtual void active_lanes_in_pipeline() = 0;
 /*
     virtual void issue( register_set& source_reg )
@@ -1113,6 +1104,9 @@ protected:
     warp_inst_t **m_pipeline_reg;
     register_set *m_result_port;
     class shader_core_ctx *m_core;
+
+    unsigned active_insts_in_pipeline;
+
 };
 
 class sfu : public pipelined_simd_unit
@@ -1413,10 +1407,8 @@ struct shader_core_config : public core_config
         }
         max_warps_per_shader =  n_thread_per_shader/warp_size;
         assert( !(n_thread_per_shader % warp_size) );
-        max_sfu_latency = 512;
-        max_sp_latency = 32;
-        
-	max_tensor_core_latency = 64;
+
+        set_pipeline_latency();
         
 	m_L1I_config.init(m_L1I_config.m_config_string,FuncCachePreferNone);
         m_L1T_config.init(m_L1T_config.m_config_string,FuncCachePreferNone);
@@ -1432,6 +1424,7 @@ struct shader_core_config : public core_config
     unsigned sid_to_cluster( unsigned sid ) const { return sid / n_simt_cores_per_cluster; }
     unsigned sid_to_cid( unsigned sid )     const { return sid % n_simt_cores_per_cluster; }
     unsigned cid_to_sid( unsigned cid, unsigned cluster_id ) const { return cluster_id*n_simt_cores_per_cluster + cid; }
+    void set_pipeline_latency();
 
 // data
     char *gpgpu_shader_core_pipeline_opt;
@@ -1506,7 +1499,9 @@ struct shader_core_config : public core_config
     bool sub_core_model;
     
     unsigned max_sp_latency;
+    unsigned max_int_latency;
     unsigned max_sfu_latency;
+    unsigned max_dp_latency;
     unsigned max_tensor_core_latency;
     
     unsigned n_simt_cores_per_cluster;
@@ -1524,6 +1519,7 @@ struct shader_core_config : public core_config
     bool gpgpu_concurrent_kernel_sm;
 
     bool adpative_volta_cache_config;
+
 };
 
 struct shader_core_stats_pod {
@@ -1734,7 +1730,7 @@ public:
     	m_cluster_id = cluster_id;
     	m_memory_config = config;
     }
-    mem_fetch *alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr ) const 
+    mem_fetch *alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const
     {
     	mem_access_t access( type, addr, size, wr );
     	mem_fetch *mf = new mem_fetch( access, 
@@ -1743,11 +1739,12 @@ public:
     				       -1, 
     				       m_core_id, 
     				       m_cluster_id,
-    				       m_memory_config );
+    				       m_memory_config,
+						   cycle);
     	return mf;
     }
     
-    mem_fetch *alloc( const warp_inst_t &inst, const mem_access_t &access ) const
+    mem_fetch *alloc( const warp_inst_t &inst, const mem_access_t &access, unsigned long long cycle ) const
     {
         warp_inst_t inst_copy = inst;
         mem_fetch *mf = new mem_fetch(access, 
@@ -1756,7 +1753,8 @@ public:
                                       inst.warp_id(),
                                       m_core_id, 
                                       m_cluster_id, 
-                                      m_memory_config);
+                                      m_memory_config,
+									  cycle);
         return mf;
     }
 
diff --git a/src/gpgpu-sim/shader_trace.h b/src/gpgpu-sim/shader_trace.h
index de3e059..ac4e894 100644
--- a/src/gpgpu-sim/shader_trace.h
+++ b/src/gpgpu-sim/shader_trace.h
@@ -44,7 +44,7 @@
 #define SHADER_DPRINTF(x, ...) do {\
     if (SHADER_DTRACE(x)) {\
         printf( SHADER_PRINT_STR,\
-                gpu_sim_cycle + gpu_tot_sim_cycle,\
+                m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle,\
                 Trace::trace_streams_str[Trace::x],\
                 get_sid() );\
         printf(__VA_ARGS__);\
@@ -56,7 +56,7 @@
 #define SCHED_DPRINTF(...) do {\
     if (SHADER_DTRACE(WARP_SCHEDULER)) {\
         printf( SCHED_PRINT_STR,\
-                gpu_sim_cycle + gpu_tot_sim_cycle,\
+                m_shader->get_gpu()->gpu_sim_cycle + m_shader->get_gpu()->gpu_tot_sim_cycle,\
                 Trace::trace_streams_str[Trace::WARP_SCHEDULER],\
                 get_sid(),\
                 m_id );\
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc
index 270bace..777f9bb 100644
--- a/src/gpgpusim_entrypoint.cc
+++ b/src/gpgpusim_entrypoint.cc
@@ -222,10 +222,12 @@ gpgpu_sim *gpgpu_ptx_sim_init_perf()
    read_parser_environment_variables();
    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);
+
+   icnt_reg_options(opp);
+   g_the_gpu_config.reg_options(opp); // register GPU microrachitecture options
+
    option_parser_cmdline(opp, sg_argc, sg_argv); // parse configuration options
    fprintf(stdout, "GPGPU-Sim: Configuration options:\n\n");
    option_parser_print(opp, stdout);
@@ -273,7 +275,7 @@ void print_simulation_time()
    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) );
+   printf("gpgpu_simulation_rate = %u (cycle/sec)\n", (unsigned)(g_the_gpu->gpu_tot_sim_cycle / difference) );
    fflush(stdout);
 }
 
diff --git a/src/stream_manager.cc b/src/stream_manager.cc
index 6cd62a2..e07f4e4 100644
--- a/src/stream_manager.cc
+++ b/src/stream_manager.cc
@@ -190,7 +190,7 @@ bool stream_operation::do_operation( gpgpu_sim *gpu )
     case stream_event: {
         printf("event update\n");
         time_t wallclock = time((time_t *)NULL);
-        m_event->update( gpu_tot_sim_cycle, wallclock );
+        m_event->update( gpu->gpu_tot_sim_cycle, wallclock );
         m_stream->record_next_done();
         } 
         break;
diff --git a/src/trace.h b/src/trace.h
index a79b4a0..0b96dcf 100644
--- a/src/trace.h
+++ b/src/trace.h
@@ -31,9 +31,6 @@
 #ifndef __TRACE_H__
 #define __TRACE_H__
 
-extern unsigned long long  gpu_sim_cycle;
-extern unsigned long long  gpu_tot_sim_cycle;
-
 namespace Trace {
 
 #define TS_TUP_BEGIN(X) enum X {
@@ -63,17 +60,26 @@ namespace Trace {
 #define DPRINTF(x, ...) do {\
     if (DTRACE(x)) {\
         printf( SIM_PRINT_STR,\
-                gpu_sim_cycle + gpu_tot_sim_cycle,\
+        		m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle,\
                 Trace::trace_streams_str[Trace::x] );\
         printf(__VA_ARGS__);\
     }\
 } while (0)
 
+#define DPRINTFG(x, ...) do {\
+    if (DTRACE(x)) {\
+        printf( SIM_PRINT_STR,\
+        		gpu_sim_cycle + gpu_tot_sim_cycle,\
+                Trace::trace_streams_str[Trace::x] );\
+        printf(__VA_ARGS__);\
+    }\
+} while (0)
 
 #else 
 
 #define DTRACE(x) (false)
 #define DPRINTF(x, ...) do {} while (0)
+#define DPRINTFG(x, ...) do {} while (0)
 
 #endif