Merge branch 'dev' into fix_warnings

34 files changed, 884 insertions, 344 deletions

diff --git a/Jenkinsfile b/Jenkinsfile
index a345bc0..9a32af1 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'
             }
         }
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 43a5864..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)
@@ -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)
 {
@@ -2996,6 +3027,7 @@ cudaError_t CUDARTAPI __cudaPopCallConfiguration(
     return g_last_cudaError = cudaSuccess;
 }
 
+
 void CUDARTAPI __cudaRegisterFunction(
 		void   **fatCubinHandle,
 		const char    *hostFun,
@@ -3263,6 +3295,7 @@ cudaError_t CUDARTAPI cudaSetDeviceFlags( int flags )
     }
 }
 
+
 cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, const char *hostFun )
 {
     return cudaFuncGetAttributesInternal(attr, hostFun );
@@ -3305,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
 
@@ -4220,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
 
@@ -5101,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..79e2719 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'
             }
         }
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 1982e04..fdc4dc3 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();
 
@@ -443,9 +444,9 @@ protected:
 
 #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 SHARED_MEM_SIZE_MAX (96*1024)
+#define LOCAL_MEM_SIZE_MAX (16*1024)
+#define MAX_STREAMING_MULTIPROCESSORS 80   //scale it to Volta
 #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)
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index f8d0b3e..fbce75b 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: 
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..6978cc1 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);
diff --git a/src/gpgpu-sim/addrdec.cc b/src/gpgpu-sim/addrdec.cc
index 09bbc3c..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() ) {
diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc
index dec61db..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;
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 1f9a422..1720836 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,
@@ -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();
 	    }
 	}
 
@@ -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..76c7a06 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 {
@@ -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 862461f..fb4ce32 100644
--- a/src/gpgpu-sim/l2cache.cc
+++ b/src/gpgpu-sim/l2cache.cc
@@ -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();
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 bb09d44..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;
@@ -46,9 +46,10 @@ xbar_router::xbar_router(unsigned router_id, enum Interconnect_type m_type, unsi
 	in_buffers.resize(total_nodes);
 	out_buffers.resize(total_nodes);
 	next_node.resize(total_nodes,0);
-//	next_node = 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;
@@ -109,14 +110,25 @@ bool xbar_router::Has_Buffer_In(unsigned input_deviceID, unsigned size, bool upd
 bool xbar_router::Has_Buffer_Out(unsigned output_deviceID, unsigned size){
 	return (out_buffers[output_deviceID].size() + size <= out_buffer_limit);
 }
-/*
+
 void xbar_router::Advance() {
+
+	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();
@@ -130,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,13 +159,17 @@ void xbar_router::Advance() {
 		out_buffer_util+=out_buffers[i].size();
 	}
 }
-*/
 
-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()){
@@ -167,6 +187,8 @@ void xbar_router::Advance() {
 		}
 	}
 
+
+	//do iSLIP
 	for(unsigned i=0; i<total_nodes; ++i){
 
 		if(Has_Buffer_Out(i, 1)) {
@@ -242,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 f4a2af1..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,11 +93,12 @@ private:
 	vector<queue<Packet> > out_buffers;
 	unsigned _n_shader, _n_mem, total_nodes;
 	unsigned in_buffer_limit, out_buffer_limit;
-	vector<unsigned> next_node;
-//	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/shader.cc b/src/gpgpu-sim/shader.cc
index c365ebb..6a0e3d6 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -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(unsigned 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 ";
 }
diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h
index dbe2285..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();
 };
 
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);