Merge pull request #4 from purdue-aalp/dev

Merge aalp dev into Mengchi

43 files changed, 1055 insertions, 479 deletions

diff --git a/Jenkinsfile b/Jenkinsfile
index a345bc0..8a3ad19 100644
--- a/Jenkinsfile
+++ b/Jenkinsfile
@@ -13,11 +13,11 @@ pipeline {
                 parallel "4.2": {
                     sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/4.2_env_setup.sh &&\
                     source `pwd`/setup_environment &&\
-                    make -j'
+                    make -j 10'
                 }, "10.1" : {
                     sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/10.1_env_setup.sh &&\
                     source `pwd`/setup_environment &&\
-                    make -j'
+                    make -j 10'
                 }
             }
         }
@@ -32,13 +32,13 @@ pipeline {
                     source `pwd`/setup_environment &&\
                     cd gpgpu-sim_simulations && \
                     source ./benchmarks/src/setup_environment && \
-                    make -j -C ./benchmarks/src rodinia_2.0-ft sdk-4.2 && \
+                    make -j 10 -C ./benchmarks/src rodinia_2.0-ft sdk-4.2 && \
                     make -C ./benchmarks/src data'
                 sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/10.1_env_setup.sh &&\
                     source `pwd`/setup_environment &&\
                     cd gpgpu-sim_simulations && \
                     source ./benchmarks/src/setup_environment && \
-                    make -j -C ./benchmarks/src/ rodinia_2.0-ft sdk-4.2 && \
+                    make -j 10 -C ./benchmarks/src/ rodinia_2.0-ft sdk-4.2 && \
                     make -C ./benchmarks/src data'
             }
         }
@@ -77,8 +77,8 @@ pipeline {
                         ./gpgpu-sim_simulations/util/plotting/plot-get-stats.py -c per-app-merge-10.1.csv -P cuda-10.1 &&\
                         ./gpgpu-sim_simulations/util/plotting/merge-stats.py -c ./gpgpu-sim-results-repo/${JOB_NAME}/stats-per-kernel-4.2-ptxplus.csv,./stats-per-kernel-4.2-ptxplus.csv -R > per-kernel-merge-4.2-ptxplus.csv &&\
                         ./gpgpu-sim_simulations/util/plotting/merge-stats.py -c ./gpgpu-sim-results-repo/${JOB_NAME}/stats-per-kernel-10.1.csv,./stats-per-kernel-10.1.csv -R > per-kernel-merge-10.1.csv &&\
-                        ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c per-kernel-merge-4.2-ptxplus.csv -p cuda-4.2 | grep "Correl=" | tee correl.4.2.txt &&\
-                        ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c per-kernel-merge-10.1.csv -p cuda-10.1 | grep "Correl=" | tee correl.10.1.txt &&\
+                        ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c per-kernel-merge-4.2-ptxplus.csv -p cuda-4.2 | grep -B 1 "Correl=" | tee correl.4.2.txt &&\
+                        ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c per-kernel-merge-10.1.csv -p cuda-10.1 | grep -B 1 "Correl=" | tee correl.10.1.txt &&\
                         mkdir -p ./gpgpu-sim-results-repo/${JOB_NAME}/ && cp stats-per-*.csv ./gpgpu-sim-results-repo/${JOB_NAME}/ &&\
                         cd ./gpgpu-sim-results-repo &&\
                         git diff --quiet && git diff --staged --quiet || git commit -am "Jenkins automated checkin ${JOB_NAME} Build:${BUILD_NUMBER}" &&\
diff --git a/Makefile b/Makefile
index a69130c..f4e8e4b 100644
--- a/Makefile
+++ b/Makefile
@@ -159,6 +159,7 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.5.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.5.5; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.6.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.6.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.6.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.6.5; fi
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.7.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.7.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.7.5 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.7.5; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.8.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.8.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.0; fi
@@ -167,6 +168,7 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.1; fi
 
+
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
 			$(SIM_OBJ_FILES_DIR)/libcuda/*.o \
diff --git a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config
index cf3627b..4a7a3c3 100644
--- a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config
+++ b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config
@@ -61,6 +61,7 @@
 # Note: Hashing set index function (H) only applies to a set size of 32 or 64. 
 -gpgpu_cache:dl1  N:32:128:4,L:L:m:N:H,S:64:8,8
 -gpgpu_shmem_size 49152
+-gpgpu_shmem_sizeDefault 49152
 -icnt_flit_size 40
 -gmem_skip_L1D 0
 -gpgpu_n_cluster_ejection_buffer_size 32
diff --git a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config
index 2fe898a..e6d8f1d 100644
--- a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config
+++ b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config
@@ -81,6 +81,7 @@
 -gpgpu_cache:dl1PrefL1  S:4:128:48,L:L:s:N:L,A:256:8,16:0,32
 -gpgpu_cache:dl1PrefShared  S:4:128:48,L:L:s:N:L,A:256:8,16:0,32
 -gpgpu_shmem_size 49152
+-gpgpu_shmem_sizeDefault 49152
 -gpgpu_shmem_size_PrefL1 49152
 -gpgpu_shmem_size_PrefShared 49152
 # By default, L1 cache is disabled in Pascal P102.
diff --git a/configs/tested-cfgs/SM7_QV100/config_volta_islip.icnt b/configs/tested-cfgs/SM7_QV100/config_volta_islip.icnt
new file mode 100644
index 0000000..5ad7ecd
--- /dev/null
+++ b/configs/tested-cfgs/SM7_QV100/config_volta_islip.icnt
@@ -0,0 +1,74 @@
+//21*1 fly with 32 flits per packet under gpgpusim injection mode
+use_map = 0;
+flit_size = 40; 
+
+// currently we do not use this, see subnets below
+network_count = 2;
+
+// Topology
+topology = fly;
+k = 144;
+n = 1;
+
+// Routing
+
+routing_function = dest_tag;
+
+
+// Flow control
+
+num_vcs     = 1;
+vc_buf_size = 256;
+input_buffer_size = 256;
+ejection_buffer_size = 256;
+boundary_buffer_size = 256;
+
+wait_for_tail_credit = 0;
+
+// Router architecture
+
+vc_allocator = islip; //separable_input_first;
+sw_allocator = islip; //separable_input_first;
+alloc_iters  = 1;
+
+credit_delay   = 0;
+routing_delay  = 0;
+vc_alloc_delay = 1;
+sw_alloc_delay = 1;
+
+input_speedup     = 1;
+output_speedup    = 1;
+internal_speedup  = 2.0;
+
+// Traffic, GPGPU-Sim does not use this
+
+traffic                = uniform;
+packet_size ={{1,2,3,4},{10,20}};
+packet_size_rate={{1,1,1,1},{2,1}};
+
+// Simulation - Don't change
+
+sim_type       = gpgpusim;
+//sim_type = latency;
+injection_rate = 0.1;
+
+subnets = 2;
+
+// Always use read and write no matter following line
+//use_read_write = 1;
+
+
+read_request_subnet = 0;
+read_reply_subnet = 1;
+write_request_subnet = 0;
+write_reply_subnet = 1;
+
+read_request_begin_vc = 0;
+read_request_end_vc = 0;
+write_request_begin_vc = 0;
+write_request_end_vc = 0;
+read_reply_begin_vc = 0;
+read_reply_end_vc = 0;
+write_reply_begin_vc = 0;
+write_reply_end_vc = 0;
+
diff --git a/configs/tested-cfgs/SM7_QV100/gpgpusim.config b/configs/tested-cfgs/SM7_QV100/gpgpusim.config
new file mode 100644
index 0000000..f807e11
--- /dev/null
+++ b/configs/tested-cfgs/SM7_QV100/gpgpusim.config
@@ -0,0 +1,205 @@
+# This config models the Volta
+# 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 
+
+
+# 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
+
+# high level architecture configuration
+-gpgpu_n_clusters 80
+-gpgpu_n_cores_per_cluster 1
+-gpgpu_n_mem 32
+-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 1132.0:1132.0:1132.0:850.0
+# boost mode
+# -gpgpu_clock_domains 1628.0:1628.0:1628.0:850.0
+
+# shader core pipeline config
+-gpgpu_shader_registers 65536
+-gpgpu_registers_per_block 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 GV100 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_cache_config 1
+# Volta unified cache has four banks
+-l1_banks 4
+#-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
+-gpgpu_shmem_sizeDefault 98304
+-gpgpu_shmem_per_block 65536
+-gmem_skip_L1D 0
+-icnt_flit_size 40
+-gpgpu_n_cluster_ejection_buffer_size 32
+-l1_latency 20
+-smem_latency 20
+-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 6MB 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 4
+
+# 128 KB Inst.
+-gpgpu_cache:il1 N:64:128:16,L:R:f:N:L,S:2:48,4
+# 128 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:4:128:256,L:R:m:N:L,T:512:8,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
+
+# interconnection
+-network_mode 1 
+-inter_config_file config_volta_islip.icnt
+# for local xbar, use:
+# "-network_mode 2 -inct_in_buffer_limit 512  -inct_out_buffer_limit 512  -inct_subnets 2"
+
+# memory partition latency config 
+-rop_latency 160
+-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, 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/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
index ebd442f..888ce71 100644
--- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
+++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
@@ -44,6 +44,7 @@
 
 # shader core pipeline config
 -gpgpu_shader_registers 65536
+-gpgpu_registers_per_block 65536
 -gpgpu_occupancy_sm_number 70
 
 # This implies a maximum of 64 warps/SM
@@ -86,16 +87,19 @@
 # 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
+-adaptive_cache_config 1
+# Volta unified cache has four banks
+-l1_banks 4
+#-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
+-gpgpu_shmem_sizeDefault 98304
+-gpgpu_shmem_per_block 65536
 -gmem_skip_L1D 0
 -icnt_flit_size 40
 -gpgpu_n_cluster_ejection_buffer_size 32
--l1_latency 28
--smem_latency 19
+-l1_latency 20
+-smem_latency 20
 -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
@@ -103,13 +107,13 @@
 -gpgpu_cache:dl2_texture_only 0
 -gpgpu_dram_partition_queues 64:64:64:64
 -perf_sim_memcpy 1
--memory_partition_indexing 0
+-memory_partition_indexing 4
 
 # 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
+-gpgpu_tex_cache:l1 N:4:128:256,L:R:m:N:L,T:512:8,128:2
 # 64 KB Const
 -gpgpu_const_cache:l1 N:128:64:8,L:R:f:N:L,S:2:64,4
 
@@ -137,9 +141,11 @@
 # interconnection
 -network_mode 1 
 -inter_config_file config_volta_islip.icnt
+# for local xbar, use:
+# "-network_mode 2 -inct_in_buffer_limit 512  -inct_out_buffer_limit 512  -inct_subnets 2"
 
 # memory partition latency config 
--rop_latency 120
+-rop_latency 160
 -dram_latency 100
 
 # dram model config
@@ -161,7 +167,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 HBM runs at 850 MHZ
+# Timing for 850 MHZ, 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"
 
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 10a651a..716e297 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -235,7 +235,7 @@ struct _cuda_device_id *gpgpu_context::GPGPUSim_Init()
   	        prop->sharedMemPerMultiprocessor = the_gpu->shared_mem_size();
 #endif	
 		prop->sharedMemPerBlock = the_gpu->shared_mem_per_block();
-		prop->regsPerBlock = the_gpu->num_registers_per_core();
+		prop->regsPerBlock = the_gpu->num_registers_per_block();
 		prop->warpSize = the_gpu->wrp_size();
 		prop->clockRate = the_gpu->shader_clock();
 #if (CUDART_VERSION >= 2010)
@@ -548,7 +548,7 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimitInternal( size_t* pValue, cudaL
 		break;
 	    }
 	    else{
-		printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+		printf("ERROR:Limit %d is not supported on this architecture \n", limit);
 		abort();
 	    }
 	case 4: // cudaLimitDevRuntimePendingLaunchCount
@@ -557,12 +557,12 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimitInternal( size_t* pValue, cudaL
 		break;
 	    }
 	    else{
-		printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+		printf("ERROR:Limit %d is not supported on this architecture \n",limit);
 		abort();
 	    }
 #endif
 	default:
-	    printf("ERROR:Limit %s unimplemented \n",limit);
+	    printf("ERROR:Limit %d unimplemented \n",limit);
 	    abort();
     }
     return g_last_cudaError = cudaSuccess;
@@ -825,207 +825,6 @@ __host__ cudaError_t CUDARTAPI cudaGetDevicePropertiesInternal(struct cudaDevice
     }
 }
 
-#if (CUDART_VERSION > 5000)
-__host__ cudaError_t CUDARTAPI cudaDeviceGetAttributeInternal(int *value, enum cudaDeviceAttr attr, int device, gpgpu_context* gpgpu_ctx = NULL)
-{
-    gpgpu_context *ctx;
-    if (gpgpu_ctx){
-	ctx = gpgpu_ctx;
-    } else {
-	ctx = GPGPU_Context();
-    }
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-        const struct cudaDeviceProp *prop;
-        _cuda_device_id *dev = ctx->GPGPUSim_Init();
-        if (device <= dev->num_devices() )  {
-                prop = dev->get_prop();
-                switch (attr) {
-                case 1:
-                        *value= prop->maxThreadsDim[0] * prop->maxThreadsDim[1] * prop->maxThreadsDim[2] * prop->maxGridSize[0] * prop->maxGridSize[1] * prop->maxGridSize[2];
-                        break;
-                case 2:
-                        *value= prop->maxThreadsDim[0];
-                        break;
-                case 3:
-                        *value= prop->maxThreadsDim[1];
-                        break;
-                case 4:
-                        *value= prop->maxThreadsDim[2];
-                        break; 
-		case 5:
-                        *value= prop->maxGridSize[0];
-                        break;
-                case 6:
-                        *value= prop->maxGridSize[1];
-                        break;
-                case 7:
-                        *value= prop->maxGridSize[2];
-                        break;
-		case 8:
-                        *value= prop->sharedMemPerBlock;
-                        break;
-		case 9:
-                        *value= prop->totalConstMem;
-                        break;
-                case 10:
-                        *value= prop->warpSize;
-                        break;
-                case 11:
-                        *value= 16;//dummy value
-                        break;
-                case 12:
-                        *value= prop->regsPerBlock;
-                        break;
-                case 13:
-                        *value= 1480000;//for 1080ti
-                        break;
-                case 14:
-                        *value= prop->textureAlignment ;
-                        break;
-                case 15:
-                        *value = 0;
-                        break;
-                case 16:
-                        *value= prop->multiProcessorCount ;
-                        break;
-                case 17:
-                case 18:
-                case 19:
-                        *value = 0;
-                        break;
-                case 21:
-                case 22:
-                case 23:
-                case 24:
-                case 25:
-                case 26:
-                case 27:
-                case 28:
-                case 42:
-                case 45:
-                case 46:
-                case 47:
-                case 48:
-                case 49:
-                case 52:
-                case 53:
-                case 55:
-                case 56:
-                case 57:
-                case 58:
-                case 59:
-                case 60:
-                case 61:
-                case 62:
-                case 63:
-                case 64:
-                case 66:
-                case 67:
-                case 69:
-                case 70:
-                case 71:
-                case 73:
-                case 74:
-                case 77:
-                        *value = 1000;//dummy value
-                        break;
-                case 29:
-                case 43:
-                case 54:
-                case 65:
-                case 68:
-                case 72:
-                        *value = 10;//dummy value
-                        break;
-                case 30:
-                case 51:
-                        *value = 128;//dummy value
-                        break;
-                case 31:
-                        *value = 1;
-                        break;
-                case 32:
-                        *value = 0;
-                        break;
-                case 33:
-                case 50:
-                        *value = 0;//dummy value
-                        break;
-		        case 34:
-                        *value= 0;
-                        break;
-                case 35:
-                        *value = 0;
-                        break;
-                case 36:
-                        *value = 1250000;//CK value for 1080ti
-                        break;
-                case 37:
-                        *value = 352;//value for 1080ti
-                        break;
-                case 38:
-                        *value = 3000000;//value for 1080ti
-                        break;
-                case 39:
-                        *value= dev->get_gpgpu()->threads_per_core();
-                        break;
-                case 40:
-                        *value= 0;
-                        break;
-                case 41:
-                        *value= 0;
-                        break;
-                case 75://cudaDevAttrComputeCapabilityMajor
-                        *value= prop->major ;
-                        break;
-                case 76://cudaDevAttrComputeCapabilityMinor
-                        *value= prop->minor ;
-                        break;
-                case 78:
-                        *value= 0 ; //TODO: as of now, we dont support stream priorities.
-                        break;
-                case 79:
-                        *value= 0;
-                        break;
-                case 80:
-                        *value= 0;
-                        break;
-		#if (CUDART_VERSION > 5050)
-		case 81:
-                        *value= prop->sharedMemPerMultiprocessor;
-                        break;
-                case 82:
-                        *value= prop->regsPerMultiprocessor;
-                        break;
-		#endif
-                case 83:
-                case 84:
-                case 85:
-                case 86:
-                        *value= 0;
-                        break;
-                case 87:
-                        *value= 4;//dummy value
-                        break;
-                case 88:
-                case 89:
-                case 90:
-                case 91:
-                case 95:
-                        *value= 0;
-                        break;
-		default:
-			printf("ERROR: Attribute number %d unimplemented \n",attr);
-			abort();
-                }
-                return g_last_cudaError = cudaSuccess;
-        } else {
-                return g_last_cudaError = cudaErrorInvalidDevice;
-        }
-}
-#endif
 
 __host__ cudaError_t CUDARTAPI cudaChooseDeviceInternal(int *device, const struct cudaDeviceProp *prop, gpgpu_context* gpgpu_ctx = NULL)
 {
@@ -1091,6 +890,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx =
 		}
 	}
 	struct CUstream_st *stream = config.get_stream();
+
 	printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun,
 			(ctx->func_sim->g_ptx_sim_mode)?"functional simulation":"performance simulation", stream?stream->get_uid():0 );
 	kernel_info_t *grid = ctx->api->gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context);
@@ -1370,9 +1170,12 @@ size_t getMaxThreadsPerBlock(struct cudaFuncAttributes *attr, gpgpu_context *ctx
 
   size_t max = prop.maxThreadsPerBlock;
 
-  if ((prop.regsPerBlock / attr->numRegs) < max)
+  if (attr->numRegs && (prop.regsPerBlock / attr->numRegs) < max)
     max = prop.regsPerBlock / attr->numRegs;
 
+  if (attr->sharedSizeBytes && (prop.sharedMemPerBlock / attr->sharedSizeBytes) < max)
+      max = prop.sharedMemPerBlock / attr->sharedSizeBytes;
+
   return max;
 }
 
@@ -1776,6 +1579,34 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArrayAsync(void *dst, size_t dpit
 	return g_last_cudaError = cudaErrorUnknown;
 }
 
+#if (CUDART_VERSION >= 8000)
+cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks, const char *hostFunc, int blockSize, size_t dynamicSMemSize, unsigned int flags)
+{
+	printf("GPGPU-Sim PTX: cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags %p\n", hostFunc);
+	CUctx_st *context = GPGPUSim_Context();
+	function_info *entry = context->get_kernel(hostFunc);
+	printf("Calculate Maxium Active Block with function ptr=%p, blockSize=%d, SMemSize=%d\n", hostFunc, blockSize, dynamicSMemSize);
+	if (flags == cudaOccupancyDefault) {
+		//create kernel_info based on entry
+		dim3 gridDim(context->get_device()->get_gpgpu()->max_cta_per_core()
+                        * context->get_device()->get_gpgpu()->get_config().num_shader());
+		dim3 blockDim(blockSize);
+		kernel_info_t result(gridDim, blockDim, entry);
+		//if(entry == NULL){
+		//	*numBlocks = 1;
+		//	return g_last_cudaError = cudaErrorUnknown;
+		//}
+		*numBlocks = context->get_device()->get_gpgpu()->get_max_cta(result);
+		printf("Maximum size is %d with gridDim %d and blockDim %d\n", *numBlocks, gridDim.x, blockDim.x);
+		return g_last_cudaError = cudaSuccess;
+	} else {
+		cuda_not_implemented(__my_func__,__LINE__);
+		return g_last_cudaError = cudaErrorUnknown;
+	}
+}
+
+#endif
+
 
 
 /*******************************************************************************
@@ -1783,7 +1614,6 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArrayAsync(void *dst, size_t dpit
  *                                                                              *
  *                                                                              *
  *******************************************************************************/
-
 __host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count)
 {
 	if(g_debug_execution >= 3){
@@ -1795,6 +1625,207 @@ __host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count)
 	return g_last_cudaError = cudaSuccess;
 }
 
+#if (CUDART_VERSION > 5000)
+__host__ cudaError_t CUDARTAPI cudaDeviceGetAttributeInternal(int *value, enum cudaDeviceAttr attr, int device, gpgpu_context* gpgpu_ctx = NULL)
+{
+	    gpgpu_context *ctx;
+		if (gpgpu_ctx){
+		ctx = gpgpu_ctx;
+		} else {
+		ctx = GPGPU_Context();
+		}
+		if(g_debug_execution >= 3){
+			announce_call(__my_func__);
+		}
+
+        const struct cudaDeviceProp *prop;
+        _cuda_device_id *dev = ctx->GPGPUSim_Init();
+
+        if (device <= dev->num_devices() )  {
+                prop = dev->get_prop();
+                switch (attr) {
+                case 1:
+                        *value= prop->maxThreadsPerBlock;
+                        break;
+                case 2:
+                        *value= prop->maxThreadsDim[0];
+                        break;
+                case 3:
+                        *value= prop->maxThreadsDim[1];
+                        break;
+                case 4:
+                        *value= prop->maxThreadsDim[2];
+                        break; 
+		case 5:
+                        *value= prop->maxGridSize[0];
+                        break;
+                case 6:
+                        *value= prop->maxGridSize[1];
+                        break;
+                case 7:
+                        *value= prop->maxGridSize[2];
+                        break;
+		case 8:
+                        *value= prop->sharedMemPerBlock;
+                        break;
+		case 9:
+                        *value= prop->totalConstMem;
+                        break;
+                case 10:
+                        *value= prop->warpSize;
+                        break;
+                case 11:
+                        *value= 16;//dummy value
+                        break;
+                case 12:
+                        *value= prop->regsPerBlock;
+                        break;
+                case 13:
+                        *value= 1480000;//for 1080ti
+                        break;
+                case 14:
+                        *value= prop->textureAlignment ;
+                        break;
+                case 15:
+                        *value = 0;
+                        break;
+                case 16:
+                        *value= prop->multiProcessorCount ;
+                        break;
+                case 17:
+                case 18:
+                case 19:
+                        *value = 0;
+                        break;
+                case 21:
+                case 22:
+                case 23:
+                case 24:
+                case 25:
+                case 26:
+                case 27:
+                case 28:
+                case 42:
+                case 45:
+                case 46:
+                case 47:
+                case 48:
+                case 49:
+                case 52:
+                case 53:
+                case 55:
+                case 56:
+                case 57:
+                case 58:
+                case 59:
+                case 60:
+                case 61:
+                case 62:
+                case 63:
+                case 64:
+                case 66:
+                case 67:
+                case 69:
+                case 70:
+                case 71:
+                case 73:
+                case 74:
+                case 77:
+                        *value = 1000;//dummy value
+                        break;
+                case 29:
+                case 43:
+                case 54:
+                case 65:
+                case 68:
+                case 72:
+                        *value = 10;//dummy value
+                        break;
+                case 30:
+                case 51:
+                        *value = 128;//dummy value
+                        break;
+                case 31:
+                        *value = 1;
+                        break;
+                case 32:
+                        *value = 0;
+                        break;
+                case 33:
+                case 50:
+                        *value = 0;//dummy value
+                        break;
+		        case 34:
+                        *value= 0;
+                        break;
+                case 35:
+                        *value = 0;
+                        break;
+                case 36:
+                        *value = 1250000;//CK value for 1080ti
+                        break;
+                case 37:
+                        *value = 352;//value for 1080ti
+                        break;
+                case 38:
+                        *value = 3000000;//value for 1080ti
+                        break;
+                case 39:
+                        *value= dev->get_gpgpu()->threads_per_core();
+                        break;
+                case 40:
+                        *value= 0;
+                        break;
+                case 41:
+                        *value= 0;
+                        break;
+                case 75://cudaDevAttrComputeCapabilityMajor
+                        *value= prop->major ;
+                        break;
+                case 76://cudaDevAttrComputeCapabilityMinor
+                        *value= prop->minor ;
+                        break;
+                case 78:
+                        *value= 0 ; //TODO: as of now, we dont support stream priorities.
+                        break;
+                case 79:
+                        *value= 0;
+                        break;
+                case 80:
+                        *value= 0;
+                        break;
+		#if (CUDART_VERSION > 5050)
+		case 81:
+                        *value= prop->sharedMemPerMultiprocessor;
+                        break;
+                case 82:
+                        *value= prop->regsPerMultiprocessor;
+                        break;
+		#endif
+                case 83:
+                case 84:
+                case 85:
+                case 86:
+                        *value= 0;
+                        break;
+                case 87:
+                        *value= 4;//dummy value
+                        break;
+                case 88:
+                case 89:
+                        *value= 0;
+                        break;
+		default:
+			printf("ERROR: Attribute number %d unimplemented \n",attr);
+			abort();
+                }
+                return g_last_cudaError = cudaSuccess;
+        } else {
+                return g_last_cudaError = cudaErrorInvalidDevice;
+        }
+}
+#endif
+
 //memset operation is done but i think its not async?
 __host__ cudaError_t CUDARTAPI cudaMemsetAsync(void *mem, int c, size_t count, 	cudaStream_t stream=0)
 {
@@ -2471,7 +2502,6 @@ void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){
     while (std::getline(infile, line))
     {
          //int pos = line.find(std::string(get_app_binary_name(app_binary)));
-         const char *ptx_file = line.c_str();
          int pos1 = line.find("sm_");
          int pos2 = line.find_last_of(".");
          if (pos1==std::string::npos&&pos2==std::string::npos){
@@ -2499,11 +2529,10 @@ void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){
  * */
 void cuda_runtime_api::extract_code_using_cuobjdump(){
     CUctx_st *context = GPGPUSim_Context();
-    unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
 
     //prevent the dumping by cuobjdump everytime we execute the code!
     const char *override_cuobjdump = getenv("CUOBJDUMP_SIM_FILE"); 
-    char command[1000], ptx_file[1000];
+    char command[1000];
     std::string app_binary = get_app_binary(); 
     //Running cuobjdump using dynamic link to current process
     snprintf(command,1000,"md5sum %s ", app_binary.c_str());
@@ -2998,6 +3027,7 @@ cudaError_t CUDARTAPI __cudaPopCallConfiguration(
     return g_last_cudaError = cudaSuccess;
 }
 
+
 void CUDARTAPI __cudaRegisterFunction(
 		void   **fatCubinHandle,
 		const char    *hostFun,
@@ -3265,6 +3295,7 @@ cudaError_t CUDARTAPI cudaSetDeviceFlags( int flags )
     }
 }
 
+
 cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, const char *hostFun )
 {
     return cudaFuncGetAttributesInternal(attr, hostFun );
@@ -3307,6 +3338,13 @@ __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t v
     return g_last_cudaError = cudaSuccess;
 }
 
+//#if CUDART_VERSION >= 9000
+//__host__  cudaError_t cudaFuncSetAttribute ( const void* func, enum cudaFuncAttribute attr, int value ) {
+
+	//ignore this Attribute for now, and the default is that carveout = cudaSharedmemCarveoutDefault;   //  (-1)
+//	return g_last_cudaError = cudaSuccess;
+//}
+
 
 #endif
 
@@ -4222,6 +4260,33 @@ CUresult CUDAAPI cuMemHostRegister(void *p, size_t bytesize, unsigned int Flags)
 	printf("WARNING: this function has not been implemented yet.");
 	return CUDA_SUCCESS;
 }
+__host__ cudaError_t cudaHostRegister(void* ptr, size_t size, unsigned int flags)
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	printf("WARNING: this function has not been implemented yet.");
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t cudaProfilerStart ( )
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	printf("WARNING: this function has not been implemented yet.");
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t cudaProfilerStop ( )
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	printf("WARNING: this function has not been implemented yet.");
+	return g_last_cudaError = cudaSuccess;
+}
+
 #endif
 #if CUDART_VERSION >= 4000
 
@@ -5103,7 +5168,7 @@ CUresult CUDAAPI cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int *numBl
 	printf("WARNING: this function has not been implemented yet.");
 	return CUDA_SUCCESS;
 }
-    
+
 CUresult CUDAAPI cuOccupancyMaxPotentialBlockSize(int *minGridSize, int *blockSize, CUfunction func, CUoccupancyB2DSize blockSizeToDynamicSMemSize, size_t dynamicSMemSize, int blockSizeLimit)
 {
 	if(g_debug_execution >= 3){
diff --git a/nightly.jenkinsfile b/nightly.jenkinsfile
index e5ffa57..c841fb0 100644
--- a/nightly.jenkinsfile
+++ b/nightly.jenkinsfile
@@ -17,7 +17,7 @@ pipeline {
             steps {
                 sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
                 source `pwd`/setup_environment &&\
-                make -j'
+                make -j 10'
             }
         }
         stage('nightly-simulations-build'){
@@ -34,7 +34,7 @@ pipeline {
                     source `pwd`/setup_environment &&\
                     cd gpgpu-sim_simulations && \
                     source ./benchmarks/src/setup_environment && \
-                    make -i -j -C ./benchmarks/src/ all && \
+                    make -i -j 10 -C ./benchmarks/src/ all && \
                     make -C ./benchmarks/src data'
             }
         }
@@ -46,7 +46,7 @@ pipeline {
                     sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
                         source `pwd`/setup_environment &&\
                         ./gpgpu-sim_simulations/util/job_launching/run_simulations.py \
-                            -B `cat ./gpgpu-sim_simulations/util/job_launching/apps/all-apps.list` -C TITANV-2B -N nightly-$$ && \
+                            -B `cat ./gpgpu-sim_simulations/util/job_launching/apps/correlation-apps.list` -C QV100-2B -N nightly-$$ && \
                         ./gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -I -S 1800 -v \
                             -s stats-per-app-9.1.csv -T 12 -K -N nightly-$$ && \
                         ./gpgpu-sim_simulations/util/plotting/plot-get-stats.py -c stats-per-app-9.1.csv -P cuda-9.1.nightly'
@@ -65,8 +65,8 @@ pipeline {
 
                 sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
                     ./gpgpu-sim_simulations/util/job_launching/get_stats.py -R -K -k \
-                        -B `cat ./gpgpu-sim_simulations/util/job_launching/apps/all-apps.list` \
-                        -C TITANV-2B > stats-per-kernel-9.1.csv &&\
+                        -B `cat ./gpgpu-sim_simulations/util/job_launching/apps/correlation-apps.list` \
+                        -C QV100-2B > stats-per-kernel-9.1.csv &&\
                     ./gpgpu-sim_simulations/util/plotting/merge-stats.py \
                         -c ./gpgpu-sim-results-repo/${JOB_NAME}/stats-per-app-9.1.csv,./stats-per-app-9.1.csv -R \
                         > per-app-merge-9.1.csv &&\
@@ -76,7 +76,7 @@ pipeline {
                     source `pwd`/setup_environment &&\
                     ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c per-kernel-merge-9.1.csv \
                         -p cuda-9.1.nightly -b ./gpgpu-sim_simulations/util/plotting/known.correlation.outliers.list \
-                        | grep "Correl=" | tee correl.9.1.txt &&\
+                        | grep "Correl=" -B 1 | tee correl.9.1.txt &&\
                     cp stats-per-*.csv ./gpgpu-sim-results-repo/${JOB_NAME}/ &&\
                     cd ./gpgpu-sim-results-repo &&\
                     git pull &&\
diff --git a/setup_environment b/setup_environment
index 58f4713..ca60d6b 100644
--- a/setup_environment
+++ b/setup_environment
@@ -7,13 +7,10 @@ export GPGPUSIM_ROOT="$( cd "$( dirname "$BASH_SOURCE" )" && pwd )"
 
 GPGPUSIM_VERSION_STRING=`cat $GPGPUSIM_ROOT/version | awk '/Version/ {print $8}'`
 #Detect Git branch and commit #
-CURRENT_PWD=`pwd`
-cd $GPGPUSIM_ROOT
-GIT_COMMIT=`git log -n 1 | head -1 | sed -re 's/commit (.*)/\1/'`
-GIT_FILES_CHANGED=`git diff --numstat | wc | sed -re 's/^\s+([0-9]+).*/\1./'`
-GIT_FILES_CHANGED+=`git diff --numstat --cached | wc | sed -re 's/^\s+([0-9]+).*/\1/'`
+GIT_COMMIT=`git --git-dir=$GPGPUSIM_ROOT/.git log -n 1 | head -1 | sed -re 's/commit (.*)/\1/'`
+GIT_FILES_CHANGED=`git --git-dir=$GPGPUSIM_ROOT/.git diff --numstat | wc | sed -re 's/^\s+([0-9]+).*/\1./'`
+GIT_FILES_CHANGED+=`git --git-dir=$GPGPUSIM_ROOT/.git diff --numstat --cached | wc | sed -re 's/^\s+([0-9]+).*/\1/'`
 GPGPUSIM_BUILD_STRING="gpgpu-sim_git-commit-$GIT_COMMIT-modified_$GIT_FILES_CHANGED"
-cd $CURRENT_PWD
 
 echo -n "GPGPU-Sim version $GPGPUSIM_VERSION_STRING (build $GPGPUSIM_BUILD_STRING) ";
 
diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index d8d5fbd..9a91818 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -710,6 +710,28 @@ void warp_inst_t::completed( unsigned long long cycle ) const
 }
 
 
+kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry)
+{
+    m_kernel_entry=entry;
+    m_grid_dim=gridDim;
+    m_block_dim=blockDim;
+    m_next_cta.x=0;
+    m_next_cta.y=0;
+    m_next_cta.z=0;
+    m_next_tid=m_next_cta;
+    m_num_cores_running=0;
+    m_uid = (entry->gpgpu_ctx->kernel_info_m_next_uid)++;
+    m_param_mem = new memory_space_impl<8192>("param",64*1024);
+
+    //Jin: parent and child kernel management for CDP
+    m_parent_kernel = NULL;
+
+    //Jin: launch latency management
+    m_launch_latency = entry->gpgpu_ctx->device_runtime->g_kernel_launch_latency;
+
+    volta_cache_config_set=false;
+}
+
 /*A snapshot of the texture mappings needs to be stored in the kernel's info as 
 kernels should use the texture bindings seen at the time of launch and textures
  can be bound/unbound asynchronously with respect to streams. */
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index d13b8c6..29e4a30 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -212,6 +212,7 @@ public:
 //      m_num_cores_running=0;
 //      m_param_mem=NULL;
 //   }
+   kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry);
    kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map<std::string, const struct cudaArray*> nameToCudaArray, std::map<std::string, const struct textureInfo*> nameToTextureInfo);
    ~kernel_info_t();
 
@@ -441,19 +442,25 @@ protected:
     class gpgpu_sim * m_gpu;
 };
 
-#define GLOBAL_HEAP_START 0xC0000000
-   // start allocating from this address (lower values used for allocating globals in .ptx file)
-#define SHARED_MEM_SIZE_MAX (64*1024)
-#define LOCAL_MEM_SIZE_MAX (8*1024)
-#define MAX_STREAMING_MULTIPROCESSORS 64
-#define MAX_THREAD_PER_SM 2048
-#define MAX_WARP_PER_SM 64
-#define TOTAL_LOCAL_MEM_PER_SM (MAX_THREAD_PER_SM*LOCAL_MEM_SIZE_MAX)
-#define TOTAL_SHARED_MEM (MAX_STREAMING_MULTIPROCESSORS*SHARED_MEM_SIZE_MAX)
-#define TOTAL_LOCAL_MEM (MAX_STREAMING_MULTIPROCESSORS*MAX_THREAD_PER_SM*LOCAL_MEM_SIZE_MAX)
-#define SHARED_GENERIC_START (GLOBAL_HEAP_START-TOTAL_SHARED_MEM)
-#define LOCAL_GENERIC_START (SHARED_GENERIC_START-TOTAL_LOCAL_MEM)
-#define STATIC_ALLOC_LIMIT (GLOBAL_HEAP_START - (TOTAL_LOCAL_MEM+TOTAL_SHARED_MEM))
+// Let's just upgrade to C++11 so we can use constexpr here...
+// start allocating from this address (lower values used for allocating globals in .ptx file)
+const unsigned long long GLOBAL_HEAP_START = 0xC0000000;
+// Volta max shmem size is 96kB
+const unsigned long long SHARED_MEM_SIZE_MAX = 96 * (1 << 10);
+// Volta max local mem is 16kB
+const unsigned long long LOCAL_MEM_SIZE_MAX = 1 << 14;
+// Volta Titan V has 80 SMs
+const unsigned MAX_STREAMING_MULTIPROCESSORS = 80;
+// Max 2048 threads / SM
+const unsigned MAX_THREAD_PER_SM = 1 << 11;
+// MAX 64 warps / SM
+const unsigned MAX_WARP_PER_SM = 1 << 6;
+const unsigned long long TOTAL_LOCAL_MEM_PER_SM = MAX_THREAD_PER_SM * LOCAL_MEM_SIZE_MAX;
+const unsigned long long TOTAL_SHARED_MEM = MAX_STREAMING_MULTIPROCESSORS * SHARED_MEM_SIZE_MAX;
+const unsigned long long TOTAL_LOCAL_MEM = MAX_STREAMING_MULTIPROCESSORS * MAX_THREAD_PER_SM * LOCAL_MEM_SIZE_MAX;
+const unsigned long long SHARED_GENERIC_START = GLOBAL_HEAP_START - TOTAL_SHARED_MEM;
+const unsigned long long LOCAL_GENERIC_START = SHARED_GENERIC_START - TOTAL_LOCAL_MEM;
+const unsigned long long STATIC_ALLOC_LIMIT = GLOBAL_HEAP_START - (TOTAL_LOCAL_MEM + TOTAL_SHARED_MEM);
 
 #if !defined(__CUDA_RUNTIME_API_H__)
 
@@ -519,10 +526,10 @@ private:
     int checkpoint_option;
     int checkpoint_kernel;
     int checkpoint_CTA;
-    int resume_option;
-    int resume_kernel;
-    int resume_CTA;
-    int checkpoint_CTA_t;
+    unsigned resume_option;
+    unsigned resume_kernel;
+    unsigned resume_CTA;
+    unsigned checkpoint_CTA_t;
     int checkpoint_insn_Y;
     int   g_ptx_inst_debug_to_file;
     char* g_ptx_inst_debug_file;
@@ -540,10 +547,10 @@ public:
     int checkpoint_option;
     int checkpoint_kernel;
     int checkpoint_CTA;
-    int resume_option;
-    int resume_kernel;
-    int resume_CTA;
-    int checkpoint_CTA_t;
+    unsigned resume_option;
+    unsigned resume_kernel;
+    unsigned resume_CTA;
+    unsigned checkpoint_CTA_t;
     int checkpoint_insn_Y;
 
     //Move some cycle core stats here instead of being global
@@ -992,7 +999,7 @@ public:
 			printf("Printing mem access generated\n");
 			std::list<mem_access_t>::iterator it;	
 			for (it = m_accessq.begin(); it != m_accessq.end(); ++it){
-   				 printf("MEM_TXN_GEN:%s:%x, Size:%d \n",mem_access_type_str(it->get_type()), it->get_addr(),it->get_size());
+   				 printf("MEM_TXN_GEN:%s:%llx, Size:%d \n",mem_access_type_str(it->get_type()), it->get_addr(),it->get_size());
 			}	
 		}
     }   
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index b9e6552..28b4bf9 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -858,7 +858,7 @@ void ptx_instruction::set_opcode_and_latency()
        op=TENSOR_CORE_OP;
 	   break;
    case SHFL_OP:
-	   latency = 32;
+	   latency = 4;
 	   initiation_interval = 4;
 	   break;
    default: 
@@ -1309,7 +1309,7 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt
     char buff[1024];
     std::string filename_c(filename+"_c");
     snprintf(buff,1024,"c++filt %s > %s", get_name().c_str(), filename_c.c_str());
-    system(buff);
+    assert(system(buff) != NULL);
     FILE *fp = fopen(filename_c.c_str(), "r");
     fgets(buff, 1024, fp);
     fclose(fp);
@@ -1432,13 +1432,13 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt
     fout = fopen(ptx_config_fn.c_str(), "a");
     assert(fout!=NULL);
     for (unsigned i = 0; i<line_number; i++){
-        fgets(buff, 1024, fin);
+        assert(fgets(buff, 1024, fin) != NULL);
         assert(!feof(fin));
     }
     fprintf(fout, "\n\n");
     do{
         fprintf(fout, "%s", buff);
-        fgets(buff, 1024, fin);
+        assert(fgets(buff, 1024, fin) != NULL);
         if(feof(fin)){
             break;
         }
@@ -1491,7 +1491,6 @@ static unsigned get_tex_datasize( const ptx_instruction *pI, ptx_thread_info *th
    const operand_info &src1 = pI->src1(); //the name of the texture
    std::string texname = src1.name();
 
-   gpgpu_t *gpu = thread->get_gpu();
    /*
      For programs with many streams, textures can be bound and unbound
      asynchronously.  This means we need to use the kernel's "snapshot" of
@@ -1577,7 +1576,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
       }
       
       //Tensorcore is warp synchronous operation. So these instructions needs to be executed only once. To make the simulation faster removing the redundant tensorcore operation
-      if(!tensorcore_op(inst_opcode)||(tensorcore_op(inst_opcode))&&(lane_id==0)){
+      if(!tensorcore_op(inst_opcode)||((tensorcore_op(inst_opcode))&&(lane_id==0))){
 	      switch ( inst_opcode ) {
 	#define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,this); op_classification = CLASSIFICATION; break;
 	#define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,get_core(),inst); op_classification = CLASSIFICATION; break;
@@ -2141,13 +2140,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL
     unsigned max_cta_tot = max_cta(kernel_info,kernel.threads_per_cta(), g_the_gpu()->getShaderCoreConfig()->warp_size, g_the_gpu()->getShaderCoreConfig()->n_thread_per_shader, g_the_gpu()->getShaderCoreConfig()->gpgpu_shmem_size, g_the_gpu()->getShaderCoreConfig()->gpgpu_shader_registers, g_the_gpu()->getShaderCoreConfig()->max_cta_per_core);
     printf("Max CTA : %d\n",max_cta_tot);
 
-    
-
-
-      
-    int inst_count=50;
     int cp_op= g_the_gpu()->checkpoint_option;
-    int cp_CTA = g_the_gpu()->checkpoint_CTA;
     int cp_kernel= g_the_gpu()->checkpoint_kernel;
     cp_count= g_the_gpu()->checkpoint_insn_Y;
     cp_cta_resume= g_the_gpu()->checkpoint_CTA_t;
@@ -2184,7 +2177,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL
 	{
       char f1name[2048];
       snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", kernel.get_uid() );
-      g_checkpoint->store_global_mem(g_the_gpu()->get_global_memory(), f1name , "%08x");
+      g_checkpoint->store_global_mem(g_the_gpu()->get_global_memory(), f1name , (char *)"%08x");
 	}
 
 
@@ -2312,18 +2305,15 @@ void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp)
     checkpoint *g_checkpoint;
     g_checkpoint = new checkpoint();
     
-    symbol * sym;
     ptx_reg_t regval;
     regval.u64= 123;
-    symbol_table * symtab= m_kernel->entry()->get_symtab();
-
 
     unsigned ctaid =m_kernel->get_next_cta_id_single();
     if(m_gpu->checkpoint_option==1 && (m_kernel->get_uid()==m_gpu->checkpoint_kernel) && (ctaid_cp>=m_gpu->checkpoint_CTA) && (ctaid_cp<m_gpu->checkpoint_CTA_t))
    {
        char fname[2048];
        snprintf(fname,2048,"checkpoint_files/shared_mem_%d.txt",ctaid-1 );
-       g_checkpoint->store_global_mem(m_thread[0]->m_shared_mem, fname , "%08x");
+       g_checkpoint->store_global_mem(m_thread[0]->m_shared_mem, fname , (char *)"%08x");
       for(int i=0; i<32*m_warp_count;i++)
       {
          char fname[2048];
@@ -2331,7 +2321,7 @@ void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp)
           m_thread[i]->print_reg_thread(fname);
           char f1name[2048];
          snprintf(f1name,2048,"checkpoint_files/local_mem_thread_%d_%d_reg.txt",i,ctaid-1 );
-         g_checkpoint->store_global_mem(m_thread[i]->m_local_mem, f1name , "%08x");
+         g_checkpoint->store_global_mem(m_thread[i]->m_local_mem, f1name , (char *)"%08x");
          m_thread[i]->set_done();
          m_thread[i]->exitCore();
          m_thread[i]->registerExit();
diff --git a/src/cuda-sim/half.h b/src/cuda-sim/half.h
index 8f1a8eb..9f74bb7 100644
--- a/src/cuda-sim/half.h
+++ b/src/cuda-sim/half.h
@@ -642,10 +642,10 @@ namespace half_float
 			if(exp > 16)

 			{

 				if(R == std::round_toward_infinity)

-					return hbits | 0x7C00 - (hbits>>15);

+					return hbits | (0x7C00 - (hbits>>15));

 				else if(R == std::round_toward_neg_infinity)

-					return hbits | 0x7BFF + (hbits>>15);

-				return hbits | 0x7BFF + (R!=std::round_toward_zero);

+					return hbits | (0x7BFF + (hbits>>15));

+				return hbits | (0x7BFF + (R!=std::round_toward_zero));

 			}

 			if(exp < -13)

 				value = std::ldexp(value, 24);

diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 58a077e..014e588 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -203,7 +203,7 @@ void ptx_thread_info::print_reg_thread(char * fname)
           const std::string &name = it->first->name();
           const std::string &dec= it->first->decl_location();
           unsigned size = it->first->get_size_in_bytes();
-          fprintf(fp,"%s %llu %s %d\n",name.c_str(),it->second, dec.c_str(),size );
+          fprintf(fp,"%s %llu %s %d\n", name.c_str(), it->second, dec.c_str(), size);
           
         }
    //m_regs.pop_back();     
@@ -224,7 +224,6 @@ void ptx_thread_info::resume_reg_thread(char * fname, symbol_table * symtab)
       {
           symbol *reg;
           char * pch;
-          unsigned size;
           pch = strtok (line," ");
           char * name =pch;
           reg= symtab->lookup(name);
@@ -232,11 +231,7 @@ void ptx_thread_info::resume_reg_thread(char * fname, symbol_table * symtab)
           pch = strtok (NULL," "); 
           data = atoi(pch);
           pch = strtok (NULL," ");
-          char * decl= pch;
           pch = strtok (NULL," ");
-          size = atoi(pch);
-
-
           m_regs.back()[reg] = data;
       }
       fclose ( fp2 );
@@ -1819,9 +1814,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
    	ptx_reg_t matrix_d[16][16];
    	ptx_reg_t src_data;
 	ptx_thread_info *thread;
-	int stride;
 
-	unsigned wmma_type = pI->get_wmma_type();
 	unsigned a_layout = pI->get_wmma_layout(0);
 	unsigned b_layout = pI->get_wmma_layout(1);
 	unsigned type = pI->get_type();
@@ -1833,7 +1826,6 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
          	tid= inst.warp_id_func()*core->get_warp_size();
         else
          	tid= inst.warp_id()*core->get_warp_size();
-	unsigned thread_group_index;
 	float temp;
 	half temp2; 	
 	
@@ -1847,9 +1839,9 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
          		ptx_reg_t v[8];
          		thread->get_vector_operand_values( src_a, v, nelem );
 			if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){
-				printf("Thread%d_Iteration=%d\n:",thrd,operand_num);
-				for(k=0;k<nelem;k++){
-					printf("%x ",v[k].u64);
+				printf("Thread%d_Iteration=%d\n:", thrd, operand_num);
+				for(k = 0; k < nelem; k++){
+					printf("%llx ",v[k].u64);
 				}
 				printf("\n");
 			}
@@ -2027,7 +2019,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
 
 				printf("thread%d:",thrd);
 				for(k=0;k<8;k++){
-					printf("%x ",matrix_d[row_t[k]][col_t[k]].f16);
+					printf("%x ", (unsigned int)matrix_d[row_t[k]][col_t[k]].f16);
 				}
 				printf("\n");
 			}
@@ -2038,7 +2030,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
 			nw_data4.s64=((matrix_d[row_t[6]][col_t[6]].s64   & 0xffff))|((matrix_d[row_t[7]][col_t[7]].s64&0xffff)<<16);
    			thread->set_vector_operand_values(dst,nw_data1,nw_data2,nw_data3,nw_data4);
 			if(core->get_gpu()->gpgpu_ctx->debug_tensorcore)
-		 		printf("thread%d=%x,%x,%x,%x",thrd,nw_data1.s64,nw_data2.s64,nw_data3.s64,nw_data4.s64);
+		 		printf("thread%d=%llx,%llx,%llx,%llx", thrd, nw_data1.s64, nw_data2.s64, nw_data3.s64, nw_data4.s64);
 		
 		}
 		else{
@@ -2298,9 +2290,8 @@ unsigned int saturatei(unsigned int a, unsigned int max)
 
 ptx_reg_t f2x( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign, int rounding_mode, int saturation_mode )
 {
- half mytemp;
- float myfloat;
-  half_float::half tmp_h;
+   half mytemp;
+   half_float::half tmp_h;
    //assert( from_width == 32); 
 
    enum cudaRoundMode mode = cudaRoundZero;
@@ -3085,7 +3076,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
    size_t size;
    unsigned smid;
    int t;
-   int thrd,odd,inx,k;
+   int thrd, k;
    ptx_thread_info *thread;
    
    const operand_info &src  = pI->operand_lookup(1);
@@ -3105,15 +3096,13 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
    _memory_op_t insn_memory_op = pI->has_memory_read() ? memory_load : memory_store;
    for (thrd=0; thrd < core->get_warp_size(); thrd++) {
 	thread = core->get_thread_info()[tid+thrd];
-	odd=thrd%2;
-	inx=thrd/2;
-    	ptx_reg_t addr_reg = thread->get_operand_value(src1, src, type, thread, 1);
+    ptx_reg_t addr_reg = thread->get_operand_value(src1, src, type, thread, 1);
    	ptx_reg_t src2_data = thread->get_operand_value(src2, src, type, thread, 1);
 	const operand_info &src_a=  pI->operand_lookup(1);
         unsigned nelem = src_a.get_vect_nelem();
         ptx_reg_t* v= new ptx_reg_t[8]; 
        	thread->get_vector_operand_values( src_a, v, nelem );
-  	stride=src2_data.u32;
+  	stride = src2_data.u32;
  
    	memory_space_t space = pI->get_space();
 
@@ -3130,9 +3119,9 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
 	}
    	decode_space(space,thread,src1,mem,addr);
 
-   	type_info_key::type_decode(type,size,t);
+   	type_info_key::type_decode(type, size, t);
    	if(core->get_gpu()->gpgpu_ctx->debug_tensorcore)
-		printf("mma_st: thrd=%d,addr=%x, fp(size=%d), stride=%d\n",thrd,addr_reg.u32,size,src2_data.u32);
+		printf("mma_st: thrd=%d, addr=%x, fp(size=%zu), stride=%d\n", thrd, addr_reg.u32, size, src2_data.u32);
 	addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type,stride)*size/8;  
 	addr_t push_addr;
 
@@ -3152,7 +3141,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
 			mem_txn_addr[num_mem_txn++]=push_addr;
 	
 			if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){
-				printf("wmma:store:thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,v[0].s64,v[1].s64,v[2].s64,v[3].s64,v[4].s64,v[5].s64,v[6].s64,v[7].s64);   
+				printf("wmma:store:thread%d=%llx,%llx,%llx,%llx,%llx,%llx,%llx,%llx\n",thrd,v[0].s64,v[1].s64,v[2].s64,v[3].s64,v[4].s64,v[5].s64,v[6].s64,v[7].s64);   
 				float temp;
 				int l;
 				printf("thread=%d:",thrd);
@@ -3179,7 +3168,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
 			}
 	
 			if(core->get_gpu()->gpgpu_ctx->debug_tensorcore)
-				printf("wmma:store:thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,nw_v[0].s64,nw_v[1].s64,nw_v[2].s64,nw_v[3].s64,nw_v[4].s64,nw_v[5].s64,nw_v[6].s64,nw_v[7].s64);   
+				printf("wmma:store:thread%d=%llx,%llx,%llx,%llx,%llx,%llx,%llx,%llx\n",thrd,nw_v[0].s64,nw_v[1].s64,nw_v[2].s64,nw_v[3].s64,nw_v[4].s64,nw_v[5].s64,nw_v[6].s64,nw_v[7].s64);   
 		}
 	}
    	
@@ -3238,11 +3227,11 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
 	}
 
 	decode_space(space,thread,src1,mem,addr);
-   	type_info_key::type_decode(type,size,t);
+   	type_info_key::type_decode(type, size, t);
 	
 	ptx_reg_t data[16];
    	if(core->get_gpu()->gpgpu_ctx->debug_tensorcore)	
-		printf("mma_ld: thrd=%d,addr=%x, fpsize=%d, stride=%d\n",thrd,src1_data.u32,size,src2_data.u32);
+		printf("mma_ld: thrd=%d,addr=%x, fpsize=%zu, stride=%d\n", thrd, src1_data.u32, size, src2_data.u32);
 	
 	addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type,stride)*size/8;  
 	addr_t fetch_addr;
@@ -3341,7 +3330,7 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
 		if(type==F16_TYPE){
 			printf("\nmma_ld:thread%d= ",thrd);
 			for(i=0;i<16;i++){
-				printf("%x ",data[i].u64);
+				printf("%llx ",data[i].u64);
 			}
 			printf("\n");
 			
@@ -3361,7 +3350,7 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
 			printf("\n");
 			printf("\nmma_ld:thread%d= ",thrd);
 			for(i=0;i<8;i++){
-				printf("%x ",data[i].u64);
+				printf("%llx ",data[i].u64);
 			}
 			printf("\n");
 		}
@@ -3388,15 +3377,15 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
 		else
    			thread->set_wmma_vector_operand_values(dst,nw_data[0],nw_data[1],nw_data[2],nw_data[3],nw_data[4],nw_data[5],nw_data[6],nw_data[7]);
 		if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){	
-			printf("mma_ld:data[0].s64=%x,data[1].s64=%x,new_data[0].s64=%x\n",data[0].u64,data[1].u64,nw_data[0].u64);	
-			printf("mma_ld:data[2].s64=%x,data[3].s64=%x,new_data[1].s64=%x\n",data[2].u64,data[3].u64,nw_data[1].u64);	
-			printf("mma_ld:data[4].s64=%x,data[5].s64=%x,new_data[2].s64=%x\n",data[4].u64,data[5].u64,nw_data[2].u64);	
-			printf("mma_ld:data[6].s64=%x,data[7].s64=%x,new_data[3].s64=%x\n",data[6].u64,data[7].u64,nw_data[3].u64);	
+			printf("mma_ld:data[0].s64=%llx,data[1].s64=%llx,new_data[0].s64=%llx\n",data[0].u64,data[1].u64,nw_data[0].u64);	
+			printf("mma_ld:data[2].s64=%llx,data[3].s64=%llx,new_data[1].s64=%llx\n",data[2].u64,data[3].u64,nw_data[1].u64);	
+			printf("mma_ld:data[4].s64=%llx,data[5].s64=%llx,new_data[2].s64=%llx\n",data[4].u64,data[5].u64,nw_data[2].u64);	
+			printf("mma_ld:data[6].s64=%llx,data[7].s64=%llx,new_data[3].s64=%llx\n",data[6].u64,data[7].u64,nw_data[3].u64);	
 			if(wmma_type!=LOAD_C){
-			printf("mma_ld:data[8].s64=%x,data[9].s64=%x,new_data[4].s64=%x\n",data[8].u64,data[9].u64,nw_data[4].s64);	
-			printf("mma_ld:data[10].s64=%x,data[11].s64=%x,new_data[5].s64=%x\n",data[10].u64,data[11].u64,nw_data[5].u64);	
-			printf("mma_ld:data[12].s64=%x,data[13].s64=%x,new_data[6].s64=%x\n",data[12].u64,data[13].u64,nw_data[6].u64);	
-			printf("mma_ld:data[14].s64=%x,data[15].s64=%x,new_data[7].s64=%x\n",data[14].u64,data[15].u64,nw_data[3].u64);	
+			printf("mma_ld:data[8].s64=%llx,data[9].s64=%llx,new_data[4].s64=%llx\n",data[8].u64,data[9].u64,nw_data[4].s64);	
+			printf("mma_ld:data[10].s64=%llx,data[11].s64=%llx,new_data[5].s64=%llx\n",data[10].u64,data[11].u64,nw_data[5].u64);	
+			printf("mma_ld:data[12].s64=%llx,data[13].s64=%llx,new_data[6].s64=%llx\n",data[12].u64,data[13].u64,nw_data[6].u64);	
+			printf("mma_ld:data[14].s64=%llx,data[15].s64=%llx,new_data[7].s64=%llx\n",data[14].u64,data[15].u64,nw_data[3].u64);	
 			}
 		}
 	}
@@ -4132,9 +4121,9 @@ int prmt_mode_present(int mode)
 	}
 	return returnval;
 }
-int read_byte(int mode,int control,int d_sel_index,signed long long value){
+int read_byte(int mode, int control, int d_sel_index, signed long long value){
 
-	int returnval;
+	int returnval = 0;
 	int prmt_f4e_mode[4][4]={{0,1,2,3},{1,2,3,4},{2,3,4,5},{3,4,5,6}};
 	int prmt_b4e_mode[4][4]={{0,7,6,5},{1,0,7,6},{2,1,0,7},{3,2,1,0}};
 	int prmt_rc8_mode[4][4]={{0,0,0,0},{1,1,1,1},{2,2,2,2},{3,3,3,3}};
@@ -4157,11 +4146,12 @@ int read_byte(int mode,int control,int d_sel_index,signed long long value){
 			case PRMT_RC8_MODE:	returnval=prmt_rc8_mode[control][d_sel_index];break;
 			case PRMT_ECL_MODE:	returnval=prmt_ecl_mode[control][d_sel_index];break;
 			case PRMT_ECR_MODE:	returnval=prmt_ecr_mode[control][d_sel_index];break;
-			case PRMT_RC16_MODE:	returnval=prmt_rc16_mode[control][d_sel_index];break;
-			default: printf("ERROR\n");break;
+			case PRMT_RC16_MODE: returnval=prmt_rc16_mode[control][d_sel_index];break;
+            // Change the default from printing "ERROR" to just asserting
+			default: assert(false);
 		}
 	}	
-	return (returnval<<8*d_sel_index);
+	return (returnval << 8 * d_sel_index);
 }
 
 void prmt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { 
diff --git a/src/cuda-sim/opcodes.h b/src/cuda-sim/opcodes.h
index b91d92f..86d3b99 100644
--- a/src/cuda-sim/opcodes.h
+++ b/src/cuda-sim/opcodes.h
@@ -68,6 +68,10 @@ enum wmma_type{
 	MMA,
 	ROW,
 	COL,
-	M16N16K16
+	M16N16K16,
+	M32N8K16,
+	M8N32K16
+
+
 };
 #endif
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index 3a2a839..2dadda4 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -175,6 +175,9 @@ breakaddr  TC; yylval->int_value = BREAKADDR_OP; return OPCODE;
 \.row TC; yylval->int_value = ROW; return LAYOUT;
 \.col TC; yylval->int_value = COL; return LAYOUT;
 \.m16n16k16 TC; yylval->int_value = M16N16K16; return CONFIGURATION;
+\.m32n8k16 TC; yylval->int_value = M32N8K16; return CONFIGURATION;
+\.m8n32k16 TC;  yylval->int_value = M8N32K16; return CONFIGURATION;
+
 \.f4e TC;  return PRMT_F4E_MODE;
 \.b4e TC;  return PRMT_B4E_MODE;
 \.rc8 TC;  return PRMT_RC8_MODE;
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index 3384d49..d8943d2 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -186,8 +186,8 @@ void symbol_table::add_function( function_info *func, const char *filename, unsi
 
 //Jin: handle instruction group for cdp
 symbol_table* symbol_table::start_inst_group() {
-   char inst_group_name[1024];
-   snprintf(inst_group_name, 1024, "%s_inst_group_%u", m_scope_name.c_str(), m_inst_group_id);
+   char inst_group_name[4096];
+   snprintf(inst_group_name, 4096, "%s_inst_group_%u", m_scope_name.c_str(), m_inst_group_id);
 
    //previous added
    assert(m_inst_group_symtab.find(std::string(inst_group_name)) == m_inst_group_symtab.end());
@@ -1154,6 +1154,8 @@ ptx_instruction::ptx_instruction( int opcode,
       		  m_wmma_layout[rr++]=last_ptx_inst_option;
       		  break;
       		case M16N16K16:
+      		case M32N8K16:
+      		case M8N32K16:
 			break;
       		default:
       		   assert(0);
@@ -1413,7 +1415,7 @@ unsigned function_info::print_insn( unsigned pc, FILE * fp ) const
    snprintf(command,1024,"c++filt -p %s",m_name.c_str());
    FILE *p = popen(command,"r");
    buffer[0]=0;
-   fgets(buffer, 1023, p);
+   assert(fgets(buffer, 1023, p) != NULL);
    // Remove trailing "\n" in buffer
    char *c;
    if ((c=strchr(buffer, '\n')) != NULL) *c = '\0';
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index dca3cec..33bcf45 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -214,7 +214,8 @@ void fix_duplicate_errors(char fname2[1024]) {
 	long filesize = ftell(ptxsource);
 	rewind(ptxsource);
 	char *ptxdata = (char*)malloc((filesize+1)*sizeof(char));
-	fread(ptxdata, filesize, 1, ptxsource);
+    // Fail if we do not read the file
+	assert(fread(ptxdata, filesize, 1, ptxsource) == 1);
 	fclose(ptxsource);
 
 	FILE *ptxdest = fopen(fname2,"w");
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 81b70af..a4f4a0c 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -421,7 +421,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi
          assert( (num_bits%8) == 0  );
          addr = g_current_symbol_table->get_sstarr_next();
          addr_pad = pad_address(addr, num_bits/8, 128);
-         printf("from 0x%x to 0x%lx (sstarr memory space)\n",
+         printf("from 0x%llx to 0x%llx (sstarr memory space)\n",
                  addr+addr_pad,
                  addr+addr_pad + num_bits/8);
             fflush(stdout);
diff --git a/src/gpgpu-sim/addrdec.cc b/src/gpgpu-sim/addrdec.cc
index ca88ec9..e3713f3 100644
--- a/src/gpgpu-sim/addrdec.cc
+++ b/src/gpgpu-sim/addrdec.cc
@@ -167,6 +167,8 @@ void linear_to_raw_address_translation::addrdec_tlx(new_addr_type addr, addrdec_
 		}
 		case RANDOM:
 		{
+			//This is an unrealistic hashing using software hashtable
+			//we generate a random set for each memory address and save the value in a big hashtable for future reuse
 			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() ) {
@@ -182,7 +184,7 @@ void linear_to_raw_address_translation::addrdec_tlx(new_addr_type addr, addrdec_
 			  }
 
 			  assert(tlx->chip < m_n_channel);
-			  assert(tlx->sub_partition < m_n_channel*m_n_sub_partition_in_channel);
+			  assert(tlx->sub_partition < m_n_channel * m_n_sub_partition_in_channel);
 			  return;
 			break;
 		}
diff --git a/src/gpgpu-sim/addrdec.h b/src/gpgpu-sim/addrdec.h
index a5333fb..c9a1420 100644
--- a/src/gpgpu-sim/addrdec.h
+++ b/src/gpgpu-sim/addrdec.h
@@ -92,7 +92,7 @@ private:
    new_addr_type sub_partition_id_mask; 
 
    unsigned int gap;
-   int m_n_channel;
+   unsigned m_n_channel;
    int m_n_sub_partition_in_channel; 
 
 };
diff --git a/src/gpgpu-sim/dram.cc b/src/gpgpu-sim/dram.cc
index d443d79..9c33822 100644
--- a/src/gpgpu-sim/dram.cc
+++ b/src/gpgpu-sim/dram.cc
@@ -482,7 +482,6 @@ void dram_t::cycle()
 
         bool memory_pending_rw_found=false;
         for (unsigned j=0;j<m_config->nbk;j++) {
-     	   unsigned grp = get_bankgrp_number(j);
      	   if (bk[j]->mrq && (((bk[j]->curr_row == bk[j]->mrq->row) &&
      		  (bk[j]->mrq->rw == READ)  &&
      		  (bk[j]->state == BANK_ACTIVE))
@@ -817,10 +816,10 @@ void dram_t::visualize() const
 
 void dram_t::print_stat( FILE* simFile ) 
 {
-   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 ",
+   fprintf(simFile,"DRAM (%u): 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);
+   fprintf(simFile, "mrqq: %d %.4g mrqsmax=%llu ", max_mrqs, (float)ave_mrqs/n_cmd, max_mrqs_temp);
    fprintf(simFile, "\n");
    fprintf(simFile, "dram_util_bins:");
    for (unsigned i=0;i<10;i++) fprintf(simFile, " %d", dram_util_bins[i]);
@@ -899,10 +898,10 @@ void dram_t::set_dram_power_stats(	unsigned &cmd,
 unsigned dram_t::get_bankgrp_number(unsigned i)
 {
 	if(m_config->dram_bnkgrp_indexing_policy == HIGHER_BITS) { //higher bits
-		return i>>m_config->bk_tag_length;
+		return i >> m_config->bk_tag_length;
 	}
 	else if (m_config->dram_bnkgrp_indexing_policy == LOWER_BITS) { //lower bits
-		return i&((m_config->nbkgrp-1));
+		return i & ((m_config->nbkgrp - 1));
 	}
 	else {
 		assert(1);
diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc
index 1705821..1e99fec 100644
--- a/src/gpgpu-sim/gpu-cache.cc
+++ b/src/gpgpu-sim/gpu-cache.cc
@@ -64,6 +64,14 @@ const char * cache_fail_status_str(enum cache_reservation_fail_reason status)
    return static_cache_reservation_fail_reason_str[status];
 }
 
+unsigned l1d_cache_config::set_bank(new_addr_type addr) const{
+
+	if(m_cache_type == SECTOR)
+		return (addr >> m_sector_sz_log2) & (l1_banks-1);
+	else
+		return (addr >> m_line_sz_log2) & (l1_banks-1);
+}
+
 unsigned l1d_cache_config::set_index(new_addr_type addr) const{
     unsigned set_index = m_nset; // Default to linear set index function
     unsigned lower_xor = 0;
@@ -777,7 +785,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] = %llu\n",
+    	  fprintf(fout, "\t%s[%s][%s] = %u\n",
 				m_cache_name.c_str(),
 				mem_access_type_str((enum mem_access_type)type),
 				"TOTAL_ACCESS",
@@ -790,7 +798,7 @@ void cache_stats::print_fail_stats(FILE *fout, const char *cache_name) const{
 	    for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) {
 	        for (unsigned fail = 0; fail < NUM_CACHE_RESERVATION_FAIL_STATUS; ++fail) {
 	            if(m_fail_stats[type][fail] > 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),
 						cache_fail_status_str((enum cache_reservation_fail_reason)fail),
@@ -1417,8 +1425,6 @@ data_cache::wr_miss_wa_lazy_fetch_on_read( new_addr_type addr,
 {
 
 	    new_addr_type block_addr = m_config.block_addr(addr);
-	    new_addr_type mshr_addr = m_config.mshr_addr(mf->get_addr());
-
 
 		//if the request writes to the whole cache line/sector, then, write and set cache line Modified.
 		//and no need to send read request to memory or reserve mshr
diff --git a/src/gpgpu-sim/gpu-cache.h b/src/gpgpu-sim/gpu-cache.h
index be33d96..6f39221 100644
--- a/src/gpgpu-sim/gpu-cache.h
+++ b/src/gpgpu-sim/gpu-cache.h
@@ -584,6 +584,7 @@ public:
         m_nset_log2 = LOGB2(m_nset);
         m_valid = true;
         m_atom_sz = (m_cache_type == SECTOR)? SECTOR_SIZE : m_line_sz;
+        m_sector_sz_log2 = LOGB2(SECTOR_SIZE);
         original_m_assoc = m_assoc;
 
         //For more details about difference between FETCH_ON_WRITE and WRITE VALIDAE policies
@@ -734,6 +735,7 @@ protected:
     unsigned m_nset_log2;
     unsigned m_assoc;
     unsigned m_atom_sz;
+    unsigned m_sector_sz_log2;
     unsigned original_m_assoc;
     bool m_is_streaming;
 
@@ -775,7 +777,9 @@ class l1d_cache_config : public cache_config{
 public:
 	l1d_cache_config() : cache_config(){}
 	virtual unsigned set_index(new_addr_type addr) const;
+    unsigned set_bank(new_addr_type addr) const;
 	unsigned l1_latency;
+	unsigned l1_banks;
 };
 
 class l2_cache_config : public cache_config {
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index e4ae04f..d236c74 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -83,6 +83,7 @@ class  gpgpu_sim_wrapper {};
 
 
 bool g_interactive_debugger_enabled=false;
+
 tr1_hash_map<new_addr_type,unsigned> address_random_interleaving;
 
 /* Clock Domains */
@@ -137,6 +138,8 @@ void memory_config::reg_options(class OptionParser * opp)
 {
     option_parser_register(opp, "-perf_sim_memcpy", OPT_BOOL, &m_perf_sim_memcpy, 
                                 "Fill the L2 cache on memcpy", "1");
+    option_parser_register(opp, "-simple_dram_model", OPT_BOOL, &simple_dram_model,
+                                "simple_dram_model with fixed latency and BW", "0");
     option_parser_register(opp, "-gpgpu_dram_scheduler", OPT_INT32, &scheduler_type, 
                                 "0 = fifo, 1 = FR-FCFS (defaul)", "1");
     option_parser_register(opp, "-gpgpu_dram_partition_queues", OPT_CSTR, &gpgpu_L2_queue_config, 
@@ -236,9 +239,12 @@ void shader_core_config::reg_options(class OptionParser * opp)
                    "per-shader L1 data cache config "
                    " {<nsets>:<bsize>:<assoc>,<rep>:<wr>:<alloc>:<wr_alloc>,<mshr>:<N>:<merge>,<mq> | none}",
                    "none" );
+    option_parser_register(opp, "-l1_banks", OPT_UINT32, &m_L1D_config.l1_banks,
+                 "The number of L1 cache banks",
+                 "1");
     option_parser_register(opp, "-l1_latency", OPT_UINT32, &m_L1D_config.l1_latency,
                  "L1 Hit Latency",
-                 "0");
+                 "1");
     option_parser_register(opp, "-smem_latency", OPT_UINT32, &smem_latency,
                  "smem Latency",
                  "3");
@@ -299,10 +305,10 @@ 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, "-adaptive_volta_cache_config", OPT_BOOL, &adaptive_volta_cache_config,
-                 "adaptive_volta_cache_config",
+    option_parser_register(opp, "-adaptive_cache_config", OPT_BOOL, &adaptive_volta_cache_config,
+                 "adaptive_cache_config",
                  "0");
-    option_parser_register(opp, "-gpgpu_shmem_size", OPT_UINT32, &gpgpu_shmem_sizeDefault,
+    option_parser_register(opp, "-gpgpu_shmem_sizeDefault", OPT_UINT32, &gpgpu_shmem_sizeDefault,
                  "Size of shared memory per shader core (default 16kB)",
                  "16384");
     option_parser_register(opp, "-gpgpu_shmem_size_PrefL1", OPT_UINT32, &gpgpu_shmem_sizePrefL1,
@@ -469,10 +475,10 @@ void gpgpu_sim_config::reg_options(option_parser_t opp)
     m_shader_config.reg_options(opp);
     m_memory_config.reg_options(opp);
     power_config::reg_options(opp);
-   option_parser_register(opp, "-gpgpu_max_cycle", OPT_INT32, &gpu_max_cycle_opt, 
+   option_parser_register(opp, "-gpgpu_max_cycle", OPT_INT64, &gpu_max_cycle_opt, 
                "terminates gpu simulation early (0 = no limit)",
                "0");
-   option_parser_register(opp, "-gpgpu_max_insn", OPT_INT32, &gpu_max_insn_opt, 
+   option_parser_register(opp, "-gpgpu_max_insn", OPT_INT64, &gpu_max_insn_opt, 
                "terminates gpu simulation early (0 = no limit)",
                "0");
    option_parser_register(opp, "-gpgpu_max_cta", OPT_INT32, &gpu_max_cta_opt, 
@@ -795,6 +801,16 @@ int gpgpu_sim::shader_clock() const
    return m_config.core_freq/1000;
 }
 
+int gpgpu_sim::max_cta_per_core() const
+{
+   return m_shader_config->max_cta_per_core;
+}
+
+int gpgpu_sim::get_max_cta( const kernel_info_t &k ) const
+{
+   return m_shader_config->max_cta(k);
+}
+
 void gpgpu_sim::set_prop( cudaDeviceProp *prop )
 {
    m_cuda_properties = prop;
@@ -1044,7 +1060,7 @@ void gpgpu_sim::change_cache_config(FuncCache cache_config)
 	if(cache_config != m_shader_config->m_L1D_config.get_cache_status()){
 		printf("FLUSH L1 Cache at configuration change between kernels\n");
 		for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++) {
-			m_cluster[i]->cache_flush();
+			m_cluster[i]->cache_invalidate();
 	    }
 	}
 
@@ -1101,8 +1117,8 @@ void gpgpu_sim::gpu_print_stat()
    printf("gpu_tot_sim_insn = %lld\n", gpu_tot_sim_insn+gpu_sim_insn);
    printf("gpu_tot_ipc = %12.4f\n", (float)(gpu_tot_sim_insn+gpu_sim_insn) / (gpu_tot_sim_cycle+gpu_sim_cycle));
    printf("gpu_tot_issued_cta = %lld\n", gpu_tot_issued_cta + m_total_cta_launched);
-   printf("gpu_occupancy = %.4f\% \n", gpu_occupancy.get_occ_fraction() * 100);
-   printf("gpu_tot_occupancy = %.4f\% \n", (gpu_occupancy + gpu_tot_occupancy).get_occ_fraction() * 100);
+   printf("gpu_occupancy = %.4f%% \n", gpu_occupancy.get_occ_fraction() * 100);
+   printf("gpu_tot_occupancy = %.4f%% \n", (gpu_occupancy + gpu_tot_occupancy).get_occ_fraction() * 100);
 
 
    fprintf(statfout, "max_total_param_size = %llu\n", gpgpu_ctx->device_runtime->g_max_total_param_size);
@@ -1343,7 +1359,7 @@ bool shader_core_ctx::occupy_shader_resource_1block(kernel_info_t & k, bool occu
        m_occupied_regs += (padded_cta_size * ((kernel_info->regs+3)&~3));
        m_occupied_ctas++;
 
-      SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Occupied %d threads, %d shared mem, %d registers, %d ctas\n",
+      SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Occupied %u threads, %u shared mem, %u registers, %u ctas\n",
             m_occupied_n_threads, m_occupied_shmem, m_occupied_regs, m_occupied_ctas);  
    }
 
@@ -1460,7 +1476,7 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel )
         nthreads_in_block += ptx_sim_init_thread(kernel,&m_thread[i],m_sid,i,cta_size-(i-start_thread),m_config->n_thread_per_shader,this,free_cta_hw_id,warp_id,m_cluster->get_gpu());
         m_threadState[i].m_active = true; 
         // load thread local memory and register file
-        if(m_gpu->resume_option==1 && kernel.get_uid()==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaid<m_gpu->checkpoint_CTA_t )
+        if(m_gpu->resume_option == 1 && kernel.get_uid() == m_gpu->resume_kernel && ctaid >= m_gpu->resume_CTA && ctaid < m_gpu->checkpoint_CTA_t )
         {
             char fname[2048];
             snprintf(fname,2048,"checkpoint_files/thread_%d_%d_reg.txt",i%cta_size,ctaid );
@@ -1475,7 +1491,7 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel )
     assert( nthreads_in_block > 0 && nthreads_in_block <= m_config->n_thread_per_shader); // should be at least one, but less than max
     m_cta_status[free_cta_hw_id]=nthreads_in_block;
 
-    if(m_gpu->resume_option==1 && kernel.get_uid()==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaid<m_gpu->checkpoint_CTA_t )
+    if(m_gpu->resume_option == 1 && kernel.get_uid() == m_gpu->resume_kernel && ctaid >= m_gpu->resume_CTA && ctaid < m_gpu->checkpoint_CTA_t )
     {
         char f1name[2048];
         snprintf(f1name,2048,"checkpoint_files/shared_mem_%d.txt", ctaid);
@@ -1582,7 +1598,10 @@ void gpgpu_sim::cycle()
 
    if (clock_mask & DRAM) {
       for (unsigned i=0;i<m_memory_config->m_n_mem;i++){
-         m_memory_partition_unit[i]->dram_cycle(); // Issue the dram command (scheduler + delay model)
+    	  if(m_memory_config->simple_dram_model)
+    		  m_memory_partition_unit[i]->simple_dram_model_cycle();
+    	  else
+    		  m_memory_partition_unit[i]->dram_cycle(); // Issue the dram command (scheduler + delay model)
          // Update performance counters for DRAM
          m_memory_partition_unit[i]->set_dram_power_stats(m_power_stats->pwr_mem_stat->n_cmd[CURRENT_STAT_IDX][i], m_power_stats->pwr_mem_stat->n_activity[CURRENT_STAT_IDX][i],
                         m_power_stats->pwr_mem_stat->n_nop[CURRENT_STAT_IDX][i], m_power_stats->pwr_mem_stat->n_act[CURRENT_STAT_IDX][i], m_power_stats->pwr_mem_stat->n_pre[CURRENT_STAT_IDX][i],
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index 19e1eb3..fba770d 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -293,9 +293,12 @@ class memory_config {
    unsigned write_high_watermark;
    unsigned write_low_watermark;
    bool m_perf_sim_memcpy;
+   bool simple_dram_model;
+
    gpgpu_context* gpgpu_ctx;
 };
 
+
 extern bool g_interactive_debugger_enabled;
 
 class gpgpu_sim_config : public power_config, public gpgpu_functional_sim_config {
@@ -365,8 +368,8 @@ private:
     double l2_period;
 
     // GPGPU-Sim timing model options
-    unsigned gpu_max_cycle_opt;
-    unsigned gpu_max_insn_opt;
+    unsigned long long gpu_max_cycle_opt;
+    unsigned long long gpu_max_insn_opt;
     unsigned gpu_max_cta_opt;
     char *gpgpu_runtime_stat;
     bool  gpgpu_flush_l1_cache;
@@ -482,6 +485,8 @@ public:
    int num_registers_per_block() const;
    int wrp_size() const;
    int shader_clock() const;
+   int max_cta_per_core() const;
+   int get_max_cta( const kernel_info_t &k ) const;
    const struct cudaDeviceProp *get_prop() const;
    enum divergence_support_t simd_model() const; 
 
diff --git a/src/gpgpu-sim/icnt_wrapper.cc b/src/gpgpu-sim/icnt_wrapper.cc
index 6e0950c..67724d0 100644
--- a/src/gpgpu-sim/icnt_wrapper.cc
+++ b/src/gpgpu-sim/icnt_wrapper.cc
@@ -182,6 +182,8 @@ void icnt_reg_options( class OptionParser * opp )
    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");
+   option_parser_register(opp, "-arbiter_algo", OPT_UINT32, &g_inct_config.arbiter_algo, "arbiter_algo", "1");
+
 
 }
 
diff --git a/src/gpgpu-sim/l2cache.cc b/src/gpgpu-sim/l2cache.cc
index 6540b52..fb4ce32 100644
--- a/src/gpgpu-sim/l2cache.cc
+++ b/src/gpgpu-sim/l2cache.cc
@@ -82,7 +82,7 @@ void memory_partition_unit::handle_memcpy_to_gpu( size_t addr, unsigned global_s
     unsigned p = global_sub_partition_id_to_local_id(global_subpart_id);
     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()); 
+    MEMPART_DPRINTF("Copy Engine Request Received For Address=%zx, 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,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle, mask);
 }
 
@@ -206,7 +206,67 @@ int memory_partition_unit::global_sub_partition_id_to_local_id(int global_sub_pa
     return (global_sub_partition_id - m_id * m_config->m_n_sub_partition_per_memory_channel); 
 }
 
-void memory_partition_unit::dram_cycle() 
+void memory_partition_unit::simple_dram_model_cycle()
+{
+
+	// pop completed memory request from dram and push it to dram-to-L2 queue
+	// of the original sub partition
+	if (!m_dram_latency_queue.empty() && ( (m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle) >= m_dram_latency_queue.front().ready_cycle )) {
+		 mem_fetch* mf_return = m_dram_latency_queue.front().req;
+		 if( mf_return->get_access_type() != L1_WRBK_ACC && mf_return->get_access_type() != L2_WRBK_ACC ) {
+			 mf_return->set_reply();
+
+		    unsigned dest_global_spid = mf_return->get_sub_partition_id();
+			int dest_spid = global_sub_partition_id_to_local_id(dest_global_spid);
+			assert(m_sub_partition[dest_spid]->get_id() == dest_global_spid);
+			if (!m_sub_partition[dest_spid]->dram_L2_queue_full()) {
+				if( mf_return->get_access_type() == L1_WRBK_ACC ) {
+					m_sub_partition[dest_spid]->set_done(mf_return);
+					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,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);
+				}
+				m_dram_latency_queue.pop_front();
+			}
+
+		  } else {
+			 this->set_done(mf_return);
+			 delete mf_return;
+			 m_dram_latency_queue.pop_front();
+		  }
+	}
+
+   // mem_fetch *mf = m_sub_partition[spid]->L2_dram_queue_top();
+	//if( !m_dram->full(mf->is_write()) ) {
+		// L2->DRAM queue to DRAM latency queue
+		// Arbitrate among multiple L2 subpartitions
+		int last_issued_partition = m_arbitration_metadata.last_borrower();
+		for (unsigned p = 0; p < m_config->m_n_sub_partition_per_memory_channel; p++) {
+			int spid = (p + last_issued_partition + 1) % m_config->m_n_sub_partition_per_memory_channel;
+			if (!m_sub_partition[spid]->L2_dram_queue_empty() && can_issue_to_dram(spid)) {
+				mem_fetch *mf = m_sub_partition[spid]->L2_dram_queue_top();
+				if(m_dram->full(mf->is_write()) )
+					break;
+
+				m_sub_partition[spid]->L2_dram_queue_pop();
+				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 = 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,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
+			}
+		}
+	//}
+
+}
+
+void memory_partition_unit::dram_cycle()
 { 
     // pop completed memory request from dram and push it to dram-to-L2 queue 
     // of the original sub partition 
@@ -231,8 +291,8 @@ void memory_partition_unit::dram_cycle()
         m_dram->return_queue_pop(); 
     }
     
-    m_dram->cycle(); 
-    m_dram->dram_log(SAMPLELOG);   
+    m_dram->cycle();
+    m_dram->dram_log(SAMPLELOG);
 
    // mem_fetch *mf = m_sub_partition[spid]->L2_dram_queue_top();
     //if( !m_dram->full(mf->is_write()) ) {
@@ -260,7 +320,6 @@ void memory_partition_unit::dram_cycle()
     //}
 
     // DRAM latency queue
-
     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();
@@ -622,7 +681,7 @@ std::vector<mem_fetch*> memory_sub_partition::breakdown_request_to_sector_reques
 			}
 		} else
 			{
-			    printf("Invalid sector received, address = 0x%06x, sector mask = %s, data size = %d",
+			    printf("Invalid sector received, address = 0x%06llx, sector mask = %s, data size = %d",
 			    		mf->get_addr(), mf->get_access_sector_mask(), mf->get_data_size());
 				assert(0 && "Undefined sector mask is received");
 			}
@@ -657,7 +716,7 @@ std::vector<mem_fetch*> memory_sub_partition::breakdown_request_to_sector_reques
 			 byte_sector_mask <<= SECTOR_SIZE;
 		}
 	} else {
-		 printf("Invalid sector received, address = 0x%06x, sector mask = %d, byte mask = , data size = %d",
+		 printf("Invalid sector received, address = 0x%06llx, sector mask = %d, byte mask = , data size = %u",
 					    		mf->get_addr(), mf->get_access_sector_mask().count(), mf->get_data_size());
 		 assert(0 && "Undefined data size is received");
 	}
diff --git a/src/gpgpu-sim/l2cache.h b/src/gpgpu-sim/l2cache.h
index 1f74c47..0f6fe32 100644
--- a/src/gpgpu-sim/l2cache.h
+++ b/src/gpgpu-sim/l2cache.h
@@ -65,6 +65,7 @@ public:
 
    void cache_cycle( unsigned cycle );
    void dram_cycle();
+   void simple_dram_model_cycle();
 
    void set_done( mem_fetch *mf );
 
diff --git a/src/gpgpu-sim/local_interconnect.cc b/src/gpgpu-sim/local_interconnect.cc
index 66d6648..c70477c 100644
--- a/src/gpgpu-sim/local_interconnect.cc
+++ b/src/gpgpu-sim/local_interconnect.cc
@@ -36,7 +36,7 @@
 #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)
+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, enum Arbiteration_type m_arbit_type)
 {
 	m_id=router_id;
 	router_type=m_type;
@@ -45,9 +45,11 @@ xbar_router::xbar_router(unsigned router_id, enum Interconnect_type m_type, unsi
 	total_nodes = n_shader+n_mem;
 	in_buffers.resize(total_nodes);
 	out_buffers.resize(total_nodes);
-	next_node=0;
+	next_node.resize(total_nodes,0);
 	in_buffer_limit = m_in_buffer_limit;
 	out_buffer_limit = m_out_buffer_limit;
+	arbit_type = m_arbit_type;
+	next_node_id=0;
 	if(m_type == REQ_NET) {
 		active_in_buffers=n_shader;
 		active_out_buffers=n_mem;
@@ -110,12 +112,23 @@ bool xbar_router::Has_Buffer_Out(unsigned output_deviceID, unsigned size){
 }
 
 void xbar_router::Advance() {
+
+	if(arbit_type == NAIVE_RR)
+		RR_Advance();
+	else if(arbit_type == iSLIP)
+		iSLIP_Advance();
+	else
+		assert(0);
+
+}
+
+void xbar_router::RR_Advance() {
 	cycles++;
 
 	vector<bool> issued(total_nodes, false);
 
 	for(unsigned i=0; i<total_nodes; ++i){
-		unsigned node_id = (i+next_node)%total_nodes;
+		unsigned node_id = (i+next_node_id)%total_nodes;
 
 		if(!in_buffers[node_id].empty()) {
 			Packet _packet = in_buffers[node_id].front();
@@ -129,12 +142,16 @@ void xbar_router::Advance() {
 				else
 					conflicts++;
 			}
-			else
+			else {
 				out_buffer_full++;
+
+				if(issued[_packet.output_deviceID])
+					conflicts++;
+			}
 		}
 	}
 
-	next_node = (++next_node % total_nodes);
+	next_node_id = (++next_node_id % total_nodes);
 
 	//collect some stats about buffer util
 	for(unsigned i=0; i<total_nodes; ++i){
@@ -143,6 +160,65 @@ void xbar_router::Advance() {
 	}
 }
 
+//iSLIP algorithm
+//McKeown, Nick. "The iSLIP scheduling algorithm for input-queued switches." IEEE/ACM transactions on networking 2 (1999): 188-201.
+//https://www.cs.rutgers.edu/~sn624/552-F18/papers/islip.pdf
+void xbar_router::iSLIP_Advance() {
+	cycles++;
+
+	vector<unsigned> node_tmp;
+
+
+	//calcaulte how many conflicts are there for stats
+	for (unsigned i=0; i<total_nodes; ++i){
+		
+		if(!in_buffers[i].empty()){
+			Packet _packet_tmp = in_buffers[i].front();
+			if (!node_tmp.empty()){
+				if (std::find(node_tmp.begin(), node_tmp.end(), _packet_tmp.output_deviceID)!=node_tmp.end()){
+					conflicts++;
+				}
+				else 
+					node_tmp.push_back(_packet_tmp.output_deviceID);
+			}
+			else{
+				node_tmp.push_back(_packet_tmp.output_deviceID);
+			}
+		}
+	}
+
+
+	//do iSLIP
+	for(unsigned i=0; i<total_nodes; ++i){
+
+		if(Has_Buffer_Out(i, 1)) {
+			for(unsigned j=0; j<total_nodes; ++j){
+				unsigned node_id = (j+next_node[i])%total_nodes;
+
+				if(!in_buffers[node_id].empty()) {
+					Packet _packet = in_buffers[node_id].front();
+					if(_packet.output_deviceID==i){
+						out_buffers[_packet.output_deviceID].push(_packet);
+						in_buffers[node_id].pop();
+						next_node[i] = (++node_id % total_nodes);
+						break;
+					}
+				}
+			}
+		}
+		else
+			out_buffer_full++;
+	}
+
+	//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){
@@ -177,7 +253,7 @@ LocalInterconnect::LocalInterconnect(const struct inct_config& m_localinct_confi
 }
 
 LocalInterconnect::~LocalInterconnect(){
-	for (int i=0; i<m_inct_config.subnets; ++i) {
+	for (unsigned i = 0; i < m_inct_config.subnets; ++i) {
 		delete net[i];
 	}
 }
@@ -188,7 +264,7 @@ void LocalInterconnect::CreateInterconnect(unsigned m_n_shader, unsigned 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 );
+		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,m_inct_config.arbiter_algo);
 	}
 
 }
diff --git a/src/gpgpu-sim/local_interconnect.h b/src/gpgpu-sim/local_interconnect.h
index 502c80d..a784da8 100644
--- a/src/gpgpu-sim/local_interconnect.h
+++ b/src/gpgpu-sim/local_interconnect.h
@@ -35,27 +35,35 @@
 using namespace std;
 
 
+enum Interconnect_type {
+	REQ_NET=0,
+	REPLY_NET=1
+};
+
+enum Arbiteration_type {
+	NAIVE_RR=0,
+	iSLIP=1
+};
+
 struct inct_config
 {
-
 	//config for local interconnect
 	unsigned   in_buffer_limit;
 	unsigned   out_buffer_limit;
 	unsigned   subnets;
+	Arbiteration_type   arbiter_algo;
 };
 
-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(unsigned router_id, enum Interconnect_type m_type, unsigned n_shader, unsigned n_mem, unsigned m_in_buffer_limit, unsigned m_out_buffer_limit, enum Arbiteration_type m_arbit_type);
 	~xbar_router();
 	void Push(unsigned input_deviceID, unsigned output_deviceID, void* data, unsigned int size);
 	void* Pop(unsigned ouput_deviceID);
-	void Advance();
+	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);
@@ -70,6 +78,9 @@ public:
 	unsigned long long packets_num;
 
 private:
+	void iSLIP_Advance();
+	void RR_Advance();
+
 	struct Packet{
 		Packet(void* m_data, unsigned m_output_deviceID) {
 			data = m_data;
@@ -82,10 +93,12 @@ private:
 	vector<queue<Packet> > out_buffers;
 	unsigned _n_shader, _n_mem, total_nodes;
 	unsigned in_buffer_limit, out_buffer_limit;
-	unsigned next_node;
+	vector<unsigned> next_node;  //used for iSLIP arbit
+	unsigned next_node_id;   //used for RR arbit
 	unsigned m_id;
 	enum Interconnect_type router_type;
 	unsigned active_in_buffers,active_out_buffers;
+	Arbiteration_type arbit_type;
 
 	friend class LocalInterconnect;
 
diff --git a/src/gpgpu-sim/scoreboard.cc b/src/gpgpu-sim/scoreboard.cc
index 80f95c6..1017e75 100644
--- a/src/gpgpu-sim/scoreboard.cc
+++ b/src/gpgpu-sim/scoreboard.cc
@@ -140,10 +140,10 @@ bool Scoreboard::checkCollision( unsigned wid, const class inst_t *inst ) const
 	// Get list of all input and output registers
 	std::set<int> inst_regs;
 
-	for(int iii=0;iii<inst->outcount;iii++)
+	for(unsigned iii=0; iii < inst->outcount; iii++)
 		inst_regs.insert(inst->out[iii]);
 
-	for(int jjj=0;jjj<inst->incount;jjj++)
+	for(unsigned jjj=0;jjj<inst->incount;jjj++)
 		inst_regs.insert(inst->in[jjj]);
 
         if(inst->pred > 0) inst_regs.insert(inst->pred);
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index c697450..6a0e3d6 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -87,7 +87,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu,
                                   shader_core_stats *stats )
    : core_t( gpu, NULL, config->warp_size, config->n_thread_per_shader ),
      m_barriers( this, config->max_warps_per_shader, config->max_cta_per_core, config->max_barriers_per_cta, config->warp_size ),
-     m_dynamic_warp_id(0), m_active_warps(0)
+     m_active_warps(0), m_dynamic_warp_id(0)
 {
     m_cluster = cluster;
     m_config = config;
@@ -164,7 +164,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu,
                                          NUM_CONCRETE_SCHEDULERS;
     assert ( scheduler != NUM_CONCRETE_SCHEDULERS );
     
-    for (int i = 0; i < m_config->gpgpu_num_sched_per_core; i++) {
+    for (unsigned i = 0; i < m_config->gpgpu_num_sched_per_core; i++) {
         switch( scheduler )
         {
             case CONCRETE_SCHEDULER_LRR:
@@ -263,7 +263,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu,
         //distribute i's evenly though schedulers;
         schedulers[i%m_config->gpgpu_num_sched_per_core]->add_supervised_warp_id(i);
     }
-    for ( int i = 0; i < m_config->gpgpu_num_sched_per_core; ++i ) {
+    for ( unsigned i = 0; i < m_config->gpgpu_num_sched_per_core; ++i ) {
         schedulers[i]->done_adding_supervised_warps();
     }
     
@@ -474,7 +474,7 @@ void shader_core_ctx::init_warps( unsigned cta_id, unsigned start_thread, unsign
             }
             m_simt_stack[i]->launch(start_pc,active_threads);
 
-              if(m_gpu->resume_option==1 && kernel_id==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaid<m_gpu->checkpoint_CTA_t )
+              if(m_gpu->resume_option == 1 && kernel_id == m_gpu->resume_kernel && ctaid >= m_gpu->resume_CTA && ctaid < m_gpu->checkpoint_CTA_t )
                { 
                 char fname[2048];
                 snprintf(fname,2048,"checkpoint_files/warp_%d_%d_simt.txt",i%warp_per_cta,ctaid );
@@ -868,7 +868,7 @@ void shader_core_ctx::func_exec_inst( warp_inst_t &inst )
 
 void shader_core_ctx::issue_warp( register_set& pipe_reg_set, const warp_inst_t* next_inst, const active_mask_t &active_mask, unsigned warp_id, unsigned sch_id )
 {
-	warp_inst_t** pipe_reg = pipe_reg = pipe_reg_set.get_free(m_config->sub_core_model, sch_id);
+	warp_inst_t** pipe_reg = pipe_reg_set.get_free(m_config->sub_core_model, sch_id);
     assert(pipe_reg);
 
     m_warp[warp_id].ibuffer_free();
@@ -1659,35 +1659,46 @@ 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(),m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle);
-
     if(m_config->m_L1D_config.l1_latency > 0)
 	{
-    	if((l1_latency_queue[m_config->m_L1D_config.l1_latency-1]) == NULL)
-    	{
-    		l1_latency_queue[m_config->m_L1D_config.l1_latency-1] = mf;
+    	for(int j=0; j<m_config->m_L1D_config.l1_banks; j++) {  //We can handle at max l1_banks reqs per cycle
 
-    		if( mf->get_inst().is_store() ) {
-				unsigned inc_ack = (m_config->m_L1D_config.get_mshr_type() == SECTOR_ASSOC)?
-						(mf->get_data_size()/SECTOR_SIZE) : 1;
+    		if( inst.accessq_empty() )
+    		        return result;
 
-				for(unsigned i=0; i< inc_ack; ++i)
-					m_core->inc_store_req( inst.warp_id() );
-			}
+    	    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);
+    	    unsigned bank_id = m_config->m_L1D_config.set_bank(mf->get_addr());
+    	    assert(bank_id < m_config->m_L1D_config.l1_banks);
+
+			if((l1_latency_queue[bank_id][m_config->m_L1D_config.l1_latency-1]) == NULL)
+			{
+				l1_latency_queue[bank_id][m_config->m_L1D_config.l1_latency-1] = mf;
+
+				if( mf->get_inst().is_store() ) {
+					unsigned inc_ack = (m_config->m_L1D_config.get_mshr_type() == SECTOR_ASSOC)?
+							(mf->get_data_size()/SECTOR_SIZE) : 1;
+
+					for(unsigned i=0; i< inc_ack; ++i)
+						m_core->inc_store_req( inst.warp_id() );
+				}
 
-    		inst.accessq_pop_back();
+				inst.accessq_pop_back();
+			}
+			else
+			{
+				result = BK_CONF;
+				delete mf;
+				break;   //do not try again, just break from the loop and try the next cycle
+			}
     	}
-    	else
-        {
-        	result = BK_CONF;
-        	delete mf;
-        }
-        if( !inst.accessq_empty() &&  result !=BK_CONF)
+		if( !inst.accessq_empty() &&  result !=BK_CONF)
 		   result = COAL_STALL;
-	   return result;
+
+	     return result;
 	}
     else
     {
+	    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,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 );
@@ -1696,63 +1707,64 @@ mem_stage_stall_type ldst_unit::process_memory_access_queue_l1cache( l1_cache *c
 
 void ldst_unit::L1_latency_queue_cycle()
 {
-	//std::deque< std::pair<mem_fetch*,bool> >::iterator it = m_latency_queue.begin();
-	if((l1_latency_queue[0]) != NULL)
-    {
-		    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,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle,events);
+	for(int j=0; j<m_config->m_L1D_config.l1_banks; j++) {
+		if((l1_latency_queue[j][0]) != NULL)
+		{
+				mem_fetch* mf_next = l1_latency_queue[j][0];
+				std::list<cache_event> 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);
+			   bool write_sent = was_write_sent(events);
+			   bool read_sent = was_read_sent(events);
 
-		   if ( status == HIT ) {
-			   assert( !read_sent );
-			   l1_latency_queue[0] = NULL;
-			   if ( mf_next->get_inst().is_load() ) {
-				   for ( unsigned r=0; r < MAX_OUTPUT_VALUES; r++)
-					   if (mf_next->get_inst().out[r] > 0)
-					   {
-						   assert(m_pending_writes[mf_next->get_inst().warp_id()][mf_next->get_inst().out[r]]>0);
-						   unsigned still_pending = --m_pending_writes[mf_next->get_inst().warp_id()][mf_next->get_inst().out[r]];
-						   if(!still_pending)
+			   if ( status == HIT ) {
+				   assert( !read_sent );
+				   l1_latency_queue[j][0] = NULL;
+				   if ( mf_next->get_inst().is_load() ) {
+					   for ( unsigned r=0; r < MAX_OUTPUT_VALUES; r++)
+						   if (mf_next->get_inst().out[r] > 0)
 						   {
-							m_pending_writes[mf_next->get_inst().warp_id()].erase(mf_next->get_inst().out[r]);
-							m_scoreboard->releaseRegister(mf_next->get_inst().warp_id(),mf_next->get_inst().out[r]);
-							m_core->warp_inst_complete(mf_next->get_inst());
+							   assert(m_pending_writes[mf_next->get_inst().warp_id()][mf_next->get_inst().out[r]]>0);
+							   unsigned still_pending = --m_pending_writes[mf_next->get_inst().warp_id()][mf_next->get_inst().out[r]];
+							   if(!still_pending)
+							   {
+								m_pending_writes[mf_next->get_inst().warp_id()].erase(mf_next->get_inst().out[r]);
+								m_scoreboard->releaseRegister(mf_next->get_inst().warp_id(),mf_next->get_inst().out[r]);
+								m_core->warp_inst_complete(mf_next->get_inst());
+							   }
 						   }
-					   }
-			   }
+				   }
 
-			   //For write hit in WB policy
-			   if(mf_next->get_inst().is_store() && !write_sent)
-			   {
-				   unsigned dec_ack = (m_config->m_L1D_config.get_mshr_type() == SECTOR_ASSOC)?
-				   						(mf_next->get_data_size()/SECTOR_SIZE) : 1;
+				   //For write hit in WB policy
+				   if(mf_next->get_inst().is_store() && !write_sent)
+				   {
+					   unsigned dec_ack = (m_config->m_L1D_config.get_mshr_type() == SECTOR_ASSOC)?
+											(mf_next->get_data_size()/SECTOR_SIZE) : 1;
 
-				   mf_next->set_reply();
+					   mf_next->set_reply();
 
-				   for(unsigned i=0; i< dec_ack; ++i)
-				      m_core->store_ack(mf_next);
-			   }
+					   for(unsigned i=0; i< dec_ack; ++i)
+						  m_core->store_ack(mf_next);
+				   }
 
-			   if( !write_sent )
-				   delete mf_next;
+				   if( !write_sent )
+					   delete mf_next;
 
-		   } else if ( status == RESERVATION_FAIL ) {
-			   assert( !read_sent );
-			   assert( !write_sent );
-		   } else {
-			   assert( status == MISS || status == HIT_RESERVED );
-			   l1_latency_queue[0] = NULL;
-	   }
-    }
+			   } else if ( status == RESERVATION_FAIL ) {
+				   assert( !read_sent );
+				   assert( !write_sent );
+			   } else {
+				   assert( status == MISS || status == HIT_RESERVED );
+				   l1_latency_queue[j][0] = NULL;
+		   }
+		}
 
-	 for( unsigned stage = 0; stage<m_config->m_L1D_config.l1_latency-1; ++stage)
-	  if( l1_latency_queue[stage] == NULL) {
-		   l1_latency_queue[stage] = l1_latency_queue[stage+1] ;
-		   l1_latency_queue[stage+1] = NULL;
-	   }
+		 for( unsigned stage = 0; stage<m_config->m_L1D_config.l1_latency-1; ++stage)
+		  if( l1_latency_queue[j][stage] == NULL) {
+			   l1_latency_queue[j][stage] = l1_latency_queue[j][stage+1] ;
+			   l1_latency_queue[j][stage+1] = NULL;
+		   }
+	}
 
 }
 
@@ -2132,11 +2144,12 @@ ldst_unit::ldst_unit( mem_fetch_interface *icnt,
                               IN_L1D_MISS_QUEUE,
 							  core->get_gpu());
 
-        if(m_config->m_L1D_config.l1_latency > 0)
-	    {
-        	for(int i=0; i<m_config->m_L1D_config.l1_latency; i++ )
-        		l1_latency_queue.push_back((mem_fetch*)NULL);
-	    }
+        l1_latency_queue.resize(m_config->m_L1D_config.l1_banks);
+        assert(m_config->m_L1D_config.l1_latency > 0);
+
+		for(unsigned j = 0; j < m_config->m_L1D_config.l1_banks; j++ )
+			l1_latency_queue[j].resize(m_config->m_L1D_config.l1_latency,(mem_fetch*)NULL);
+
     }
     m_name = "MEM ";
 }
@@ -2446,7 +2459,7 @@ void shader_core_ctx::register_cta_thread_exit( unsigned cta_num, kernel_info_t
       m_barriers.deallocate_barrier(cta_num);
       shader_CTA_count_unlog(m_sid, 1);
 
-     SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Finished CTA #%d (%lld,%lld), %u CTAs running\n",
+     SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Finished CTA #%u (%lld,%lld), %u CTAs running\n",
         cta_num, m_gpu->gpu_sim_cycle, m_gpu->gpu_tot_sim_cycle, m_n_active_cta);
 
       if( m_n_active_cta == 0 ) {
diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h
index b0d7f7f..667cb2d 100644
--- a/src/gpgpu-sim/shader.h
+++ b/src/gpgpu-sim/shader.h
@@ -1326,7 +1326,7 @@ protected:
    unsigned long long m_last_inst_gpu_sim_cycle;
    unsigned long long m_last_inst_gpu_tot_sim_cycle;
 
-   std::deque<mem_fetch* > l1_latency_queue;
+   std::vector<std::deque<mem_fetch* >> l1_latency_queue;
    void L1_latency_queue_cycle();
 };
 
@@ -1392,9 +1392,9 @@ class shader_core_config : public core_config
 	 	If we won't remove it, old regression will be broken.
 		So to support the legacy config files it's best to handle in this way.
          */  
-	int num_config_to_read=N_PIPELINE_STAGES-2*(!gpgpu_tensor_core_avail);
+	int num_config_to_read= N_PIPELINE_STAGES - 2 * (!gpgpu_tensor_core_avail);
 
-        for (unsigned i = 0; i <num_config_to_read; i++) { 
+        for (int i = 0; i < num_config_to_read; i++) { 
 	    assert(toks);
 	    ntok = sscanf(toks,"%d", &pipe_widths[i]);
 	    assert(ntok == 1); 
@@ -1455,7 +1455,7 @@ class shader_core_config : public core_config
 
     bool gpgpu_dwf_reg_bankconflict;
 
-    int gpgpu_num_sched_per_core;
+    unsigned gpgpu_num_sched_per_core;
     int gpgpu_max_insn_issue_per_warp;
     bool gpgpu_dual_issue_diff_exec_units;
 
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc
index 816159f..d3deb24 100644
--- a/src/gpgpusim_entrypoint.cc
+++ b/src/gpgpusim_entrypoint.cc
@@ -269,7 +269,9 @@ 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)(GPGPUsim_ctx_ptr()->g_the_gpu->gpu_tot_sim_insn / difference) );
-   printf("gpgpu_simulation_rate = %u (cycle/sec)\n", (unsigned)(GPGPUsim_ctx_ptr()->g_the_gpu->gpu_tot_sim_cycle / difference) );
+   const unsigned cycles_per_sec = (unsigned)(GPGPUsim_ctx_ptr()->g_the_gpu->gpu_tot_sim_cycle / difference);
+   printf("gpgpu_simulation_rate = %u (cycle/sec)\n", cycles_per_sec );
+   printf("gpgpu_silicon_slowdown = %ux\n", GPGPUsim_ctx_ptr()->g_the_gpu->shader_clock() * 1000 / cycles_per_sec);
    fflush(stdout);
 }
 
diff --git a/src/intersim2/flit.hpp b/src/intersim2/flit.hpp
index fd48306..1c58c68 100644
--- a/src/intersim2/flit.hpp
+++ b/src/intersim2/flit.hpp
@@ -57,8 +57,8 @@ public:
   int  itime;
   int  atime;
 
-  int  id;
-  int  pid;
+  unsigned long long  id;
+  unsigned long long  pid;
 
   bool record;
 
diff --git a/src/intersim2/gputrafficmanager.cpp b/src/intersim2/gputrafficmanager.cpp
index bf422d6..6897a22 100644
--- a/src/intersim2/gputrafficmanager.cpp
+++ b/src/intersim2/gputrafficmanager.cpp
@@ -105,7 +105,7 @@ void GPUTrafficManager::_RetireFlit( Flit *f, int dest )
     if(f->head) {
       head = f;
     } else {
-      map<int, Flit *>::iterator iter = _retired_packets[f->cl].find(f->pid);
+      map<unsigned long long, Flit *>::iterator iter = _retired_packets[f->cl].find(f->pid);
       assert(iter != _retired_packets[f->cl].end());
       head = iter->second;
       _retired_packets[f->cl].erase(iter);
@@ -195,8 +195,8 @@ void GPUTrafficManager::_GeneratePacket(int source, int stype, int cl, int time,
   
   //  Flit::FlitType packet_type = Flit::ANY_TYPE;
   int size = packet_size; //input size
-  int pid = _cur_pid++;
-  assert(_cur_pid);
+  unsigned long long pid = _cur_pid++;
+  assert(_cur_pid > 0);
   int packet_destination = dest;
   bool record = false;
   bool watch = gWatchOut && (_packets_to_watch.count(pid) > 0);
diff --git a/src/intersim2/routers/iq_router.cpp b/src/intersim2/routers/iq_router.cpp
index d97f485..7dffb3a 100644
--- a/src/intersim2/routers/iq_router.cpp
+++ b/src/intersim2/routers/iq_router.cpp
@@ -306,7 +306,7 @@ bool IQRouter::_ReceiveFlits( )
 
       if(f->watch) {
 	*gWatchOut << GetSimTime() << " | " << FullName() << " | "
-		   << "Received flit " << f->id
+		   << "Received flit " << (unsigned) f->id
 		   << " from channel at input " << input
 		   << "." << endl;
       }
diff --git a/src/intersim2/stats.hpp b/src/intersim2/stats.hpp
index 1aaf013..e186f4d 100644
--- a/src/intersim2/stats.hpp
+++ b/src/intersim2/stats.hpp
@@ -62,6 +62,9 @@ public:
   inline void AddSample( int val ) {
     AddSample( (double)val );
   }
+  inline void AddSample( unsigned long long val ) {
+    AddSample( (double)val );
+  }
 
   int GetBin(int b){ return _hist[b];}
 
diff --git a/src/intersim2/trafficmanager.cpp b/src/intersim2/trafficmanager.cpp
index 8a015bb..7a20d07 100644
--- a/src/intersim2/trafficmanager.cpp
+++ b/src/intersim2/trafficmanager.cpp
@@ -679,7 +679,7 @@ void TrafficManager::_RetireFlit( Flit *f, int dest )
         if(f->head) {
             head = f;
         } else {
-            map<int, Flit *>::iterator iter = _retired_packets[f->cl].find(f->pid);
+            map<unsigned long long, Flit *>::iterator iter = _retired_packets[f->cl].find(f->pid);
             assert(iter != _retired_packets[f->cl].end());
             head = iter->second;
             _retired_packets[f->cl].erase(iter);
@@ -1380,7 +1380,7 @@ void TrafficManager::_DisplayRemaining( ostream & os ) const
 {
     for(int c = 0; c < _classes; ++c) {
 
-        map<int, Flit *>::const_iterator iter;
+        map<unsigned long long, Flit *>::const_iterator iter;
         int i;
 
         os << "Class " << c << ":" << endl;
@@ -1463,7 +1463,7 @@ bool TrafficManager::_SingleSim( )
             double latency = (double)_plat_stats[c]->Sum();
             double count = (double)_plat_stats[c]->NumSamples();
       
-            map<int, Flit *>::const_iterator iter;
+            map<unsigned long long, Flit *>::const_iterator iter;
             for(iter = _total_in_flight_flits[c].begin(); 
                 iter != _total_in_flight_flits[c].end(); 
                 iter++) {
@@ -1568,7 +1568,7 @@ bool TrafficManager::_SingleSim( )
                         double acc_latency = _plat_stats[c]->Sum();
                         double acc_count = (double)_plat_stats[c]->NumSamples();
 	    
-                        map<int, Flit *>::const_iterator iter;
+                        map<unsigned long long, Flit *>::const_iterator iter;
                         for(iter = _total_in_flight_flits[c].begin(); 
                             iter != _total_in_flight_flits[c].end(); 
                             iter++) {
diff --git a/src/intersim2/trafficmanager.hpp b/src/intersim2/trafficmanager.hpp
index 9694df4..97564ea 100644
--- a/src/intersim2/trafficmanager.hpp
+++ b/src/intersim2/trafficmanager.hpp
@@ -113,9 +113,9 @@ protected:
   vector<vector<bool> > _qdrained;
   vector<vector<list<Flit *> > > _partial_packets;
 
-  vector<map<int, Flit *> > _total_in_flight_flits;
-  vector<map<int, Flit *> > _measured_in_flight_flits;
-  vector<map<int, Flit *> > _retired_packets;
+  vector<map<unsigned long long, Flit *> > _total_in_flight_flits;
+  vector<map<unsigned long long, Flit *> > _measured_in_flight_flits;
+  vector<map<unsigned long long, Flit *> > _retired_packets;
   bool _empty_network;
 
   bool _hold_switch_for_packet;
@@ -229,12 +229,12 @@ protected:
   vector<double> _warmup_threshold;
   vector<double> _acc_warmup_threshold;
 
-  int _cur_id;
-  int _cur_pid;
+  unsigned long long _cur_id;
+  unsigned long long _cur_pid;
   int _time;
 
-  set<int> _flits_to_watch;
-  set<int> _packets_to_watch;
+  set<unsigned long long> _flits_to_watch;
+  set<unsigned long long> _packets_to_watch;
 
   bool _print_csv_results;
 
diff --git a/src/stream_manager.h b/src/stream_manager.h
index 91d1b36..3fbdbaf 100644
--- a/src/stream_manager.h
+++ b/src/stream_manager.h
@@ -258,6 +258,8 @@ public:
     void pushCudaStreamWaitEventToAllStreams( CUevent_st *e, unsigned int flags );
     bool operation(bool * sim);
     void stop_all_running_kernels();
+    unsigned size() {return m_streams.size(); };
+    bool is_blocking() {return m_cuda_launch_blocking; };
 private:
     void print_impl( FILE *fp);