Merge branch 'dev' of https://github.com/gpgpu-sim/gpgpu-sim_distribution into dev

16 files changed, 272 insertions, 69 deletions

diff --git a/CHANGES b/CHANGES
index c5b97fc..2a4222b 100644
--- a/CHANGES
+++ b/CHANGES
@@ -14,11 +14,12 @@ Version 4.0.0 (development branch) versus 3.2.3
 4- debug support for wmma instruction using debug_tensorcore flag
 -GPU Core Performance Simulation:
 1- Fermi/Pascal coalescer: coalescing on 32-bytes sectors.
-2- Adding separate dp unit pipeline.
+2- Adding separate int and dp units pipeline.
 3- diff dual issue: allow scheduler to issue diff insts at a time
 4- Fair memory issue from multiple schedulers.
 5- Added tensorcore unit pipeline.
 6- Corrected the bug in vector load instruction
+7- Volta sub-core model (scheduler isolation)
 -Cache System:
 1- Sector L1/L2 cache
 2- Fetch-on-write and lazy-fetch-on-read write allocation policy.
diff --git a/Jenkinsfile b/Jenkinsfile
index 84d847e..4bdbddf 100644
--- a/Jenkinsfile
+++ b/Jenkinsfile
@@ -24,9 +24,8 @@ pipeline {
         stage('simulations-build'){
             steps{
                 sh 'rm -rf gpgpu-sim_simulations'
-                sh 'git clone [email protected]:TimRogersGroup/gpgpu-sim_simulations.git && \
+                sh 'git clone [email protected]:gpgpu-sim/gpgpu-sim_simulations.git && \
                     cd gpgpu-sim_simulations && \
-                    git checkout purdue-cluster && \
                     git pull && \
                     ln -s /home/tgrogers-raid/a/common/data_dirs benchmarks/'
                 sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/4.2_env_setup.sh &&\
diff --git a/README b/README
index 8b8939c..998cb51 100644
--- a/README
+++ b/README
@@ -18,15 +18,23 @@ Analyzing CUDA Workloads Using a Detailed GPU Simulator, in IEEE International
 Symposium on Performance Analysis of Systems and Software (ISPASS), Boston, MA,
 April 19-21, 2009.
 
+If you use cuDNN and Pytorch support, the Checkpoint function or the Debigging tool for functional simulation error in GPGPU-Sim for your research,
+please cite:
+Jonathan Lew, Deval Shah, Suchita Pati, Shaylin Cattell, Mengchi Zhang, Amruth Sandhupatla, Christopher Ng, Negar Goli, Matthew D. Sinclair, Timothy G. Rogers, Tor M. Aamodt
+Analyzing Machine Learning Workloads Using a Detailed GPU Simulator, arXiv:1811.08933,
+https://arxiv.org/abs/1811.08933
+
 If you use the memory system in GPGPU-Sim, or the Volta/Pascal models,
 please cite:
 Mahmoud Khairy, Jain Akshay, Tor Aamodt, Timothy G Rogers,
 Exploring Modern GPU Memory System Design Challenges through Accurate Modeling, arXiv:1810.07269, 
 https://arxiv.org/abs/1810.07269
 
-If you use the tensorcore in GPGPU-Sim or CUTLASS Library in your research
+If you use the Tensor Core in GPGPU-Sim or CUTLASS Library for your research
 please cite:
-	add the arxiv link here
+Md Aamir Raihan, Negar Goli, Tor Aamodt,
+Modeling Deep Learning Accelerator Enabled GPUs, arXiv:1811.08309, 
+https://arxiv.org/abs/1811.08309
 
 If you use the GPUWattch energy model in your research, please cite:
 
diff --git a/configs/deprecated-cfgs/SM6_GTX1080/gpgpusim.config b/configs/deprecated-cfgs/SM6_GTX1080/gpgpusim.config
index fb044c6..3261d5a 100644
--- a/configs/deprecated-cfgs/SM6_GTX1080/gpgpusim.config
+++ b/configs/deprecated-cfgs/SM6_GTX1080/gpgpusim.config
@@ -31,12 +31,14 @@
 -gpgpu_simd_model 1 
 
 # Pipeline widths and number of FUs
-# ID_OC_SP,ID_OC_SFU,ID_OC_MEM,OC_EX_SP,OC_EX_SFU,OC_EX_MEM,EX_WB
+# ID_OC_SP, ID_OC_DP, ID_OC_INT, ID_OC_SFU, ID_OC_MEM, OC_EX_SP, OC_EX_DP, OC_EX_INT, OC_EX_SFU, OC_EX_MEM, EX_WB
 ## Pascal GP102 has 4 SP SIMD units and 1 SFU unit
 ## we need to scale the number of pipeline registers to be equal to the number of SP units
--gpgpu_pipeline_widths 4,1,1,4,1,1,6
+-gpgpu_pipeline_widths 4,0,0,1,1,4,0,0,1,1,6
 -gpgpu_num_sp_units 4
 -gpgpu_num_sfu_units 1
+-gpgpu_tensor_core_avail 0
+-gpgpu_num_tensor_core_units 0
 
 # Instruction latencies and initiation intervals
 # "ADD,MAX,MUL,MAD,DIV"
@@ -54,20 +56,20 @@
 # Pascal GP102 has 96KB Shared memory
 # Pascal GP102 has 64KB L1 cache
 # The default is to disable the L1 cache, unless cache modifieres is used
--gpgpu_cache:dl1  64:128:6,L:L:m:N:H,A:128:8,8
+-gpgpu_cache:dl1  N:64:128:6,L:L:m:N:H,A:128:8,8
 -gpgpu_shmem_size 98304
 -gmem_skip_L1D 1
 
 # 64 sets, each 128 bytes 16-way for each memory sub partition (128 KB per memory sub partition). This gives 3MB L2 cache
--gpgpu_cache:dl2 64:128:16,L:B:m:W:L,A:1024:1024,4:0,32 # used to be 128:4
+-gpgpu_cache:dl2 N:64:128:16,L:B:m:W:L,A:1024:1024,4:0,32 # used to be 128:4
 -gpgpu_cache:dl2_texture_only 0 
 
 # 4 KB Inst.
--gpgpu_cache:il1 8:128:4,L:R:f:N:L,A:2:48,4
+-gpgpu_cache:il1 N:8:128:4,L:R:f:N:L,A:2:48,4
 # 48 KB Tex 
--gpgpu_tex_cache:l1 16:128:24,L:R:m:N:L,F:128:4,128:2
+-gpgpu_tex_cache:l1 N:16:128:24,L:R:m:N:L,F:128:4,128:2
 # 12 KB Const
--gpgpu_const_cache:l1 128:64:2,L:R:f:N:L,A:2:64,4
+-gpgpu_const_cache:l1 N:128:64:2,L:R:f:N:L,A:2:64,4
 
 # enable operand collector 
 ## larger operand collectors and reg_banks are needed for the 4 warp schedulers and 4 SIMD units
diff --git a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config
index d71b2fd..cf3627b 100644
--- a/configs/tested-cfgs/SM2_GTX480/gpgpusim.config
+++ b/configs/tested-cfgs/SM2_GTX480/gpgpusim.config
@@ -3,6 +3,13 @@
 -gpgpu_ptx_sim_mode 0
 -gpgpu_ptx_force_max_capability 20 
 
+# Device Limits
+-gpgpu_stack_size_limit 1024
+-gpgpu_heap_size_limit 8388608
+
+# Compute Capability
+-gpgpu_compute_capability_major 2
+-gpgpu_compute_capability_minor 0
 
 # SASS execution (only supported with CUDA >= 4.0)
 -gpgpu_ptx_convert_to_ptxplus 0
diff --git a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config
index 73e3951..2fe898a 100644
--- a/configs/tested-cfgs/SM6_TITANX/gpgpusim.config
+++ b/configs/tested-cfgs/SM6_TITANX/gpgpusim.config
@@ -8,6 +8,16 @@
 -gpgpu_ptx_force_max_capability 61 
 -gpgpu_ignore_resources_limitation 1
 
+# Device Limits
+-gpgpu_stack_size_limit 1024
+-gpgpu_heap_size_limit 8388608
+-gpgpu_runtime_sync_depth_limit 2
+-gpgpu_runtime_pending_launch_count_limit 2048
+
+# Compute Capability
+-gpgpu_compute_capability_major 6
+-gpgpu_compute_capability_minor 1
+
 # SASS execution (only supported with CUDA >= 4.0)
 -gpgpu_ptx_convert_to_ptxplus 0
 -gpgpu_ptx_save_converted_ptxplus 0
diff --git a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
index f48e897..13d8c8f 100644
--- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
+++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
@@ -13,6 +13,17 @@
 -gpgpu_ptx_sim_mode 0
 -gpgpu_ptx_force_max_capability 70 
 
+
+# Device Limits
+-gpgpu_stack_size_limit 1024
+-gpgpu_heap_size_limit 8388608
+-gpgpu_runtime_sync_depth_limit 2
+-gpgpu_runtime_pending_launch_count_limit 2048
+
+# Compute Capability
+-gpgpu_compute_capability_major 7
+-gpgpu_compute_capability_minor 0
+
 # SASS execution (only supported with CUDA >= 4.0)
 -gpgpu_ptx_convert_to_ptxplus 0
 -gpgpu_ptx_save_converted_ptxplus 0
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 95a3c24..3a9d613 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -354,8 +354,8 @@ struct _cuda_device_id *GPGPUSim_Init()
 
 		cudaDeviceProp *prop = (cudaDeviceProp *) calloc(sizeof(cudaDeviceProp),1);
 		snprintf(prop->name,256,"GPGPU-Sim_v%s", g_gpgpusim_version_string );
-		prop->major = 5;
-		prop->minor = 2;
+		prop->major = the_gpu->compute_capability_major();
+		prop->minor = the_gpu->compute_capability_minor();
 		prop->totalGlobalMem = 0x80000000 /* 2 GB */;
 		prop->memPitch = 0;
 		if(prop->major >= 2) {
@@ -1120,11 +1120,11 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic
                 case 41:
                         *value= 0;
                         break;
-                case 75:
-                        *value= 9 ;
+                case 75://cudaDevAttrComputeCapabilityMajor
+                        *value= prop->major ;
                         break;
-                case 76:
-                        *value= 3 ;
+                case 76://cudaDevAttrComputeCapabilityMinor
+                        *value= prop->minor ;
                         break;
                 case 78:
                         *value= 0 ; //TODO: as of now, we dont support stream priorities.
@@ -1200,6 +1200,59 @@ __host__ cudaError_t CUDARTAPI cudaGetDevice(int *device)
 	return g_last_cudaError = cudaSuccess;
 }
 
+__host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit limit )
+{
+	if(g_debug_execution >= 3){
+            announce_call(__my_func__);
+   	 }
+        _cuda_device_id *dev = GPGPUSim_Init();
+	const struct cudaDeviceProp *prop = dev->get_prop();
+	const gpgpu_sim_config& config=dev->get_gpgpu()->get_config();
+        switch(limit) {
+        case 0:  // cudaLimitStackSize
+        	*pValue=config.stack_limit();
+                break;
+        case 2:  // cudaLimitMallocHeapSize
+                *pValue=config.heap_limit();
+                break;
+#if (CUDART_VERSION > 5050)
+        case 3: // cudaLimitDevRuntimeSyncDepth
+		if(prop->major > 2){
+			*pValue=config.sync_depth_limit();
+	                break;
+		}
+		else{
+			printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+			abort();
+		}
+        case 4: // cudaLimitDevRuntimePendingLaunchCount
+		if(prop->major > 2){
+     	        	*pValue=config.pending_launch_count_limit();
+	                break;
+		}
+		else{
+			printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+			abort();
+		}
+#endif
+        default:
+                        printf("ERROR:Limit %s unimplemented \n",limit);
+                        abort();
+        }
+	return g_last_cudaError = cudaSuccess;
+
+}
+
+__host__ cudaError_t CUDARTAPI cudaStreamGetPriority ( cudaStream_t hStream, int* priority )
+{
+        if(g_debug_execution >= 3){
+            announce_call(__my_func__);
+   	 }
+        cuda_not_implemented(__my_func__,__LINE__);
+        return g_last_cudaError = cudaSuccess;
+
+}
+
 __host__ cudaError_t CUDARTAPI cudaDeviceGetPCIBusId (
 		char *pciBusId,
 		int len,
@@ -1235,6 +1288,16 @@ __host__ cudaError_t cudaIpcOpenMemHandle(
 	return g_last_cudaError = cudaErrorUnknown;
 }
 
+__host__ cudaError_t CUDARTAPI cudaDestroyTextureObject(cudaTextureObject_t texObject)
+{
+        if(g_debug_execution >= 3){
+            announce_call(__my_func__);
+   	 }
+        cuda_not_implemented(__my_func__,__LINE__);
+        return g_last_cudaError = cudaErrorUnknown;
+}
+
+
 /*******************************************************************************
  *                                                                              *
  *                                                                              *
@@ -1469,42 +1532,26 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun )
 	return g_last_cudaError = cudaSuccess;
 }
 
-
 __host__ cudaError_t CUDARTAPI cudaLaunchKernel ( const char* hostFun, dim3 gridDim, dim3 blockDim, const void** args, size_t sharedMem, cudaStream_t stream )
 {
-	struct CUstream_st *s = (struct CUstream_st *)stream;
-	g_cuda_launch_stack.push_back( kernel_config(gridDim,blockDim,sharedMem,s) );
-
-	//printf("cudaLaunchKernel:sizeof(Arg[0])=%d)\n ",sizeof(args[0]));
-	kernel_config &config = g_cuda_launch_stack.back();
-	config.set_arg(args[0],432,0);//standard interface for cutlass library #TODO Implementing a generalized kernel
-
-	CUctx_st* context = GPGPUSim_Context();
-	char *mode = getenv("PTX_SIM_MODE_FUNC");
-	if( mode )
-		sscanf(mode,"%u", &g_ptx_sim_mode);
-	gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
-	kernel_config config1 = g_cuda_launch_stack.back();
-	struct CUstream_st *stream1 = config1.get_stream();
-	printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun,
-			g_ptx_sim_mode?"functional simulation":"performance simulation", stream1?stream1->get_uid():0 );
-	kernel_info_t *grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config1.get_args(),config1.grid_dim(),config1.block_dim(),context);
-	std::string kname = grid->name();
-	dim3 gridDim1 = config1.grid_dim();
-	dim3 blockDim1 = config1.block_dim();
-	printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n",
-			kname.c_str(), stream1?stream1->get_uid():0, gridDim1.x,gridDim1.y,gridDim1.z,blockDim1.x,blockDim1.y,blockDim1.z );
 
-	/*Kernel is hardcoded to enable the cutlass library*/
-	std::string cutlass("cutlass");
-	assert(kname.find(cutlass) != std::string::npos);
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    	}
+        CUctx_st *context = GPGPUSim_Context();
+        function_info *entry = context->get_kernel(hostFun);
+    
+	cudaConfigureCall(gridDim, blockDim, sharedMem, stream);
+    	for(unsigned i = 0; i < entry->num_args(); i++){
+        	std::pair<size_t, unsigned> p = entry->get_param_config(i);
+        	cudaSetupArgument(args[i], p.first, p.second);
+    	}  
 
-	stream_operation op(grid,g_ptx_sim_mode,stream1);
-	g_stream_manager->push(op);
-	g_cuda_launch_stack.pop_back();
+	cudaLaunch(hostFun);
 	return g_last_cudaError = cudaSuccess;
 }
 
+
 /*******************************************************************************
  *                                                                              *
  *                                                                              *
@@ -1555,6 +1602,9 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
 	    announce_call(__my_func__);
     }
 #if (CUDART_VERSION >= 3000)
+	//per-stream synchronization required for application using external libraries without explicit synchronization in the code to 
+	//avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. 
+	stream->synchronize();
 	g_stream_manager->destroy_stream(stream);
 #endif
 	return g_last_cudaError = cudaSuccess;
@@ -2552,8 +2602,17 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
         // Making this a runtime variable based on the app, enables GPGPU-Sim compiled
         // with a newer version of CUDA to run apps compiled with older versions of
         // CUDA. This is especially useful for PTXPLUS execution.
-        int app_cuda_version = get_app_cuda_version();
-        assert( app_cuda_version == CUDART_VERSION / 1000  && "The app must be compiled with same major version as the simulator." );
+        //Skip cuda version check for pytorch application
+	std::string app_binary_path =  get_app_binary();
+        int pos = app_binary_path.find("python");
+        if (pos==std::string::npos){
+        	// Not pytorch app : checking cuda version
+		int app_cuda_version = get_app_cuda_version();
+        	assert( app_cuda_version == CUDART_VERSION / 1000  && "The app must be compiled with same major version as the simulator." );
+        }
+
+	//int app_cuda_version = get_app_cuda_version();
+        //assert( app_cuda_version == CUDART_VERSION / 1000  && "The app must be compiled with same major version as the simulator." );
         const char* filename;
 #if CUDART_VERSION < 6000
             // FatBin handle from the .fatbin.c file (one of the intermediate files generated by NVCC)
@@ -2672,14 +2731,14 @@ cudaError_t cudaDeviceReset ( void ) {
 	return g_last_cudaError = cudaSuccess;
 }
 cudaError_t CUDARTAPI cudaDeviceSynchronize(void){
-	// I don't know what this should do
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
+	//Blocks until the device has completed all preceding requested tasks
+	synchronize();
 	return g_last_cudaError = cudaSuccess;
 }
 
-
 void CUDARTAPI __cudaRegisterFunction(
 		void   **fatCubinHandle,
 		const char    *hostFun,
@@ -3109,6 +3168,8 @@ __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t v
     }
     return g_last_cudaError = cudaSuccess;
 }
+
+
 #endif
 
 #endif
@@ -3287,7 +3348,12 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
 	    announce_call(__my_func__);
     }
 	function_info *entry = context->get_kernel(hostFun);
-	kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry);
+	gpgpu_t* gpu= context->get_device()->get_gpgpu();
+	/*
+	Passing a snapshot of the GPU's current texture mapping to the kernel's info
+	as kernels should use texture bindings present at the time of their launch.
+	*/
+	kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping());
 	if( entry == NULL ) {
 		printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found for %p\n", hostFun);
 		abort();
@@ -4447,6 +4513,16 @@ CUresult CUDAAPI cuPointerGetAttribute(void *data, CUpointer_attribute attribute
 #endif /* CUDART_VERSION >= 4000 */
 
 #if CUDART_VERSION >= 8000
+__host__ cudaError_t CUDARTAPI cudaCreateTextureObject ( cudaTextureObject_t* pTexObject, const cudaResourceDesc* pResDesc, const cudaTextureDesc* pTexDesc, const cudaResourceViewDesc* pResViewDesc )
+{
+         if(g_debug_execution >= 3){
+            announce_call(__my_func__);
+        }
+        cuda_not_implemented(__my_func__,__LINE__);
+        return g_last_cudaError = cudaSuccess;
+
+}
+
 CUresult CUDAAPI cuMemPrefetchAsync(CUdeviceptr devPtr, size_t count, CUdevice dstDevice, CUstream hStream)
 {
 	if(g_debug_execution >= 3){
diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index 307df40..cebdb25 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -88,8 +88,6 @@ void checkpoint::load_global_mem(class memory_space *temp_mem, char * f1name)
       fclose ( fp2 );
 }
 
-
-
 void checkpoint::store_global_mem(class memory_space * mem, char *fname, char * format)
 {
 
@@ -192,6 +190,11 @@ gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config )
    checkpoint_CTA_t = m_function_model_config.get_checkpoint_CTA_t();
    checkpoint_insn_Y = m_function_model_config.get_checkpoint_insn_Y();
 
+   // initialize texture mappings to empty
+   m_NameToTextureInfo.clear();
+   m_NameToCudaArray.clear();
+   m_TextureRefToName.clear();
+   m_NameToAttribute.clear();
 
    if(m_function_model_config.get_ptx_inst_debug_to_file() != 0) 
       ptx_inst_debug_file = fopen(m_function_model_config.get_ptx_inst_debug_file(), "w");
@@ -688,7 +691,10 @@ unsigned g_kernel_launch_latency;
 
 unsigned kernel_info_t::m_next_uid = 1;
 
-kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry )
+/*A snapshot of the texture mappings needs to be stored in the kernel's info as 
+kernels should use the texture bindings seen at the time of launch and textures
+ can be bound/unbound asynchronously with respect to streams. */
+kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map<std::string, const struct cudaArray*> nameToCudaArray, std::map<std::string, const struct textureInfo*> nameToTextureInfo)   
 {
     m_kernel_entry=entry;
     m_grid_dim=gridDim;
@@ -708,6 +714,8 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *
     m_launch_latency = g_kernel_launch_latency;
 
     volta_cache_config_set=false;
+    m_NameToCudaArray = nameToCudaArray;
+    m_NameToTextureInfo = nameToTextureInfo;
 }
 
 kernel_info_t::~kernel_info_t()
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index 08aa88c..e8716ab 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -73,6 +73,7 @@ enum FuncCache
 #include <set>
 
 typedef unsigned long long new_addr_type;
+typedef unsigned long long cudaTextureObject_t;
 typedef unsigned address_type;
 typedef unsigned addr_t;
 
@@ -211,7 +212,7 @@ public:
 //      m_num_cores_running=0;
 //      m_param_mem=NULL;
 //   }
-   kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry );
+   kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map<std::string, const struct cudaArray*> nameToCudaArray, std::map<std::string, const struct textureInfo*> nameToTextureInfo);
    ~kernel_info_t();
 
    void inc_running() { m_num_cores_running++; }
@@ -274,6 +275,23 @@ public:
    std::list<class ptx_thread_info *> &active_threads() { return m_active_threads; }
    class memory_space *get_param_memory() { return m_param_mem; }
 
+   
+   //The following functions access texture bindings present at the kernel's launch
+   
+   const struct cudaArray* get_texarray( const std::string &texname ) const
+   {
+      std::map<std::string,const struct cudaArray*>::const_iterator t=m_NameToCudaArray.find(texname);
+      assert(t != m_NameToCudaArray.end());
+      return t->second;
+   }
+
+   const struct textureInfo* get_texinfo( const std::string &texname ) const
+   {
+      std::map<std::string, const struct textureInfo*>::const_iterator t=m_NameToTextureInfo.find(texname);
+      assert(t != m_NameToTextureInfo.end());
+      return t->second;
+   }
+
 private:
    kernel_info_t( const kernel_info_t & ); // disable copy constructor
    void operator=( const kernel_info_t & ); // disable copy operator
@@ -282,6 +300,10 @@ private:
 
    unsigned m_uid;
    static unsigned m_next_uid;
+   
+   //These maps contain the snapshot of the texture mappings at kernel launch
+   std::map<std::string, const struct cudaArray*> m_NameToCudaArray;
+   std::map<std::string, const struct textureInfo*> m_NameToTextureInfo;
 
    dim3 m_grid_dim;
    dim3 m_block_dim;
@@ -579,8 +601,8 @@ public:
 
     const struct textureInfo* get_texinfo( const std::string &texname ) const
     {
-        std::map<std::string, const struct textureInfo*>::const_iterator t=m_NameToTexureInfo.find(texname);
-        assert(t != m_NameToTexureInfo.end());
+        std::map<std::string, const struct textureInfo*>::const_iterator t=m_NameToTextureInfo.find(texname);
+        assert(t != m_NameToTextureInfo.end());
         return t->second;
     }
 
@@ -593,6 +615,10 @@ public:
 
     const gpgpu_functional_sim_config &get_config() const { return m_function_model_config; }
     FILE* get_ptx_inst_debug_file() { return ptx_inst_debug_file; }
+    
+    //  These maps return the current texture mappings for the GPU at any given time.
+    std::map<std::string, const struct cudaArray*> getNameArrayMapping() {return m_NameToCudaArray;}
+    std::map<std::string, const struct textureInfo*> getNameInfoMapping() {return m_NameToTextureInfo;}
 
 protected:
     const gpgpu_functional_sim_config &m_function_model_config;
@@ -603,11 +629,11 @@ protected:
     class memory_space *m_surf_mem;
 
     unsigned long long m_dev_malloc;
-    
+    //  These maps contain the current texture mappings for the GPU at any given time. 
     std::map<std::string, std::set<const struct textureReference*> > m_NameToTextureRef;
     std::map<const struct textureReference*, std::string> m_TextureRefToName;
     std::map<std::string, const struct cudaArray*> m_NameToCudaArray;
-    std::map<std::string, const struct textureInfo*> m_NameToTexureInfo;
+    std::map<std::string, const struct textureInfo*> m_NameToTextureInfo;
     std::map<std::string, const struct textureReferenceAttr*> m_NameToAttribute;
 };
 
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index ec68b5b..f7bb9cc 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -218,7 +218,7 @@ void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* te
    texInfo->Ty_numbits = intLOGB2(Ty);
    texInfo->texel_size = texel_size;
    texInfo->texel_size_numbits = intLOGB2(texel_size);
-   m_NameToTexureInfo[texname] = texInfo;
+   m_NameToTextureInfo[texname] = texInfo;
 }
 
 void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref)
@@ -226,7 +226,7 @@ void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref)
    //assumes bind-use-unbind-bind-use-unbind pattern
    std::string texname = gpgpu_ptx_sim_findNamefromTexture(texref);
    m_NameToCudaArray.erase(texname);
-   m_NameToTexureInfo.erase(texname);
+   m_NameToTextureInfo.erase(texname);
 }
 
 unsigned g_assemble_code_next_pc=0; 
@@ -1507,7 +1507,15 @@ static unsigned get_tex_datasize( const ptx_instruction *pI, ptx_thread_info *th
    std::string texname = src1.name();
 
    gpgpu_t *gpu = thread->get_gpu();
-   const struct textureInfo* texInfo = gpu->get_texinfo(texname);
+   /*
+     For programs with many streams, textures can be bound and unbound
+     asynchronously.  This means we need to use the kernel's "snapshot" of
+     the state of the texture mappings when it was launched (so that we
+     don't try to access the incorrect texture mapping if it's been updated,
+     or that we don't access a mapping that has been unbound).
+    */
+   kernel_info_t& k = thread->get_kernel();
+   const struct textureInfo* texInfo = k.get_texinfo(texname);
 
    unsigned data_size = texInfo->texel_size;
    return data_size; 
@@ -1918,7 +1926,7 @@ kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,
                                              struct dim3 blockDim,
                                              gpgpu_t *gpu )
 {
-   kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry);
+   kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping());
    unsigned argcount=args.size();
    unsigned argn=1;
    for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) {
diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc
index b399133..917e7a8 100644
--- a/src/cuda-sim/cuda_device_runtime.cc
+++ b/src/cuda-sim/cuda_device_runtime.cc
@@ -198,7 +198,8 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info *
             memory_space *device_kernel_param_mem;
 
             //create child kernel_info_t and index it with parameter_buffer address
-            device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry);
+	    gpgpu_t* gpu=thread->get_gpu();
+            device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry, gpu->getNameArrayMapping(), gpu->getNameInfoMapping());
             device_grid->launch_cycle = gpu_sim_cycle + gpu_tot_sim_cycle;
             kernel_info_t & parent_grid = thread->get_kernel();
             DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" <<
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index c85654c..11001d7 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -5258,11 +5258,18 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread )
    if (!ptx_tex_regs) ptx_tex_regs = new ptx_reg_t[4];
    unsigned nelem = src2.get_vect_nelem();
    thread->get_vector_operand_values(src2, ptx_tex_regs, nelem); //ptx_reg should be 4 entry vector type...coordinates into texture
-
+   /*
+     For programs with many streams, textures can be bound and unbound
+     asynchronously.  This means we need to use the kernel's "snapshot" of
+     the state of the texture mappings when it was launched (so that we
+     don't try to access the incorrect texture mapping if it's been updated,
+     or that we don't access a mapping that has been unbound).
+   */
    gpgpu_t *gpu = thread->get_gpu();
+   kernel_info_t &k = thread->get_kernel();
    const struct textureReference* texref = gpu->get_texref(texname);
-   const struct cudaArray* cuArray = gpu->get_texarray(texname); 
-   const struct textureInfo* texInfo = gpu->get_texinfo(texname);
+   const struct cudaArray* cuArray = k.get_texarray(texname); 
+   const struct textureInfo* texInfo = k.get_texinfo(texname);
    const struct textureReferenceAttr* texAttr = gpu->get_texattr(texname);
 
    //assume always 2D f32 input
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index d4bf631..3232361 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -61,6 +61,7 @@ addc    TC; ptx_lval.int_value = ADDC_OP; return OPCODE;
 and	TC; ptx_lval.int_value = AND_OP; return OPCODE;
 andn	TC; ptx_lval.int_value = ANDN_OP; return OPCODE;
 atom	TC; ptx_lval.int_value = ATOM_OP; return OPCODE;
+bar.warp 	TC; ptx_lval.int_value = NOP_OP; return OPCODE;
 bar 	TC; ptx_lval.int_value = BAR_OP; return OPCODE;
 bfe     TC; ptx_lval.int_value = BFE_OP; return OPCODE;
 bfi     TC; ptx_lval.int_value = BFI_OP; return OPCODE;
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index d837526..3f1fc7e 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -505,6 +505,12 @@ void gpgpu_sim_config::reg_options(option_parser_t opp)
    option_parser_register(opp, "-liveness_message_freq", OPT_INT64, &liveness_message_freq, 
                "Minimum number of seconds between simulation liveness messages (0 = always print)",
                "1");
+   option_parser_register(opp, "-gpgpu_compute_capability_major", OPT_UINT32, &gpgpu_compute_capability_major,
+                 "Major compute capability version number",
+                 "7");
+   option_parser_register(opp, "-gpgpu_compute_capability_minor", OPT_UINT32, &gpgpu_compute_capability_minor,
+                 "Minor compute capability version number",
+                 "0");
    option_parser_register(opp, "-gpgpu_flush_l1_cache", OPT_BOOL, &gpgpu_flush_l1_cache,
                 "Flush L1 cache at the end of each kernel call",
                 "0");
@@ -538,6 +544,14 @@ void gpgpu_sim_config::reg_options(option_parser_t opp)
    option_parser_register(opp, "-visualizer_zlevel", OPT_INT32,
                           &g_visualizer_zlevel, "Compression level of the visualizer output log (0=no comp, 9=highest)",
                           "6");
+   option_parser_register(opp, "-gpgpu_stack_size_limit", OPT_INT32, &stack_size_limit,
+                          "GPU thread stack size", "1024" );
+   option_parser_register(opp, "-gpgpu_heap_size_limit", OPT_INT32, &heap_size_limit,
+                          "GPU malloc heap size ", "8388608" );
+   option_parser_register(opp, "-gpgpu_runtime_sync_depth_limit", OPT_INT32, &runtime_sync_depth_limit,
+                          "GPU device runtime synchronize depth", "2" );
+   option_parser_register(opp, "-gpgpu_runtime_pending_launch_count_limit", OPT_INT32, &runtime_pending_launch_count_limit,
+                          "GPU device runtime pending launch count", "2048" );
     option_parser_register(opp, "-trace_enabled", OPT_BOOL, 
                           &Trace::enabled, "Turn on traces",
                           "0");
@@ -800,6 +814,16 @@ void gpgpu_sim::set_prop( cudaDeviceProp *prop )
    m_cuda_properties = prop;
 }
 
+int gpgpu_sim::compute_capability_major() const
+{
+   return m_config.gpgpu_compute_capability_major;
+}
+
+int gpgpu_sim::compute_capability_minor() const
+{
+   return m_config.gpgpu_compute_capability_minor;
+}
+
 const struct cudaDeviceProp *gpgpu_sim::get_prop() const
 {
    return m_cuda_properties;
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index efda65a..c8dad89 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -336,6 +336,11 @@ public:
     unsigned get_max_concurrent_kernel() const { return max_concurrent_kernel; }
     unsigned checkpoint_option;
 
+    size_t stack_limit() const {return stack_size_limit; }
+    size_t heap_limit() const {return heap_size_limit; }
+    size_t sync_depth_limit() const {return runtime_sync_depth_limit; }
+    size_t pending_launch_count_limit() const {return runtime_pending_launch_count_limit;}
+
 private:
     void init_clock_domains(void ); 
 
@@ -376,8 +381,15 @@ private:
     int gpu_stat_sample_freq;
     int gpu_runtime_stat_flag;
 
+    // Device Limits
+    size_t stack_size_limit;
+    size_t heap_size_limit;
+    size_t runtime_sync_depth_limit;
+    size_t runtime_pending_launch_count_limit;	
 
-
+ //gpu compute capability options
+    unsigned int gpgpu_compute_capability_major;
+    unsigned int gpgpu_compute_capability_minor;
     unsigned long long liveness_message_freq; 
 
     friend class gpgpu_sim;
@@ -437,6 +449,8 @@ public:
 
    int shared_mem_size() const;
    int shared_mem_per_block() const;
+   int compute_capability_major() const;
+   int compute_capability_minor() const;
    int num_registers_per_core() const;
    int num_registers_per_block() const;
    int wrp_size() const;