fixing device and function parameters config

6 files changed, 26 insertions, 15 deletions

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/gpgpusim.config b/configs/tested-cfgs/SM7_QV100/gpgpusim.config
index 1a34d0f..5f64908 100644
--- a/configs/tested-cfgs/SM7_QV100/gpgpusim.config
+++ b/configs/tested-cfgs/SM7_QV100/gpgpusim.config
@@ -11,7 +11,7 @@
 # functional simulator specification
 -gpgpu_ptx_instruction_classification 0
 -gpgpu_ptx_sim_mode 0
--gpgpu_ptx_force_max_capability 60 
+-gpgpu_ptx_force_max_capability 70 
 
 
 # Device Limits
@@ -21,7 +21,7 @@
 -gpgpu_runtime_pending_launch_count_limit 2048
 
 # Compute Capability
--gpgpu_compute_capability_major 6
+-gpgpu_compute_capability_major 7
 -gpgpu_compute_capability_minor 0
 
 # SASS execution (only supported with CUDA >= 4.0)
@@ -44,12 +44,13 @@
 
 # shader core pipeline config
 -gpgpu_shader_registers 65536
--gpgpu_occupancy_sm_number 60
+-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 
+-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
@@ -91,6 +92,8 @@
 -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
diff --git a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config
index b06f048..6c21dcb 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
@@ -91,6 +92,8 @@
 -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
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index c70a570..45511d4 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -383,7 +383,7 @@ struct _cuda_device_id *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)
@@ -1014,7 +1014,7 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic
                 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];
+                        *value= prop->maxThreadsPerBlock;
                         break;
                 case 2:
                         *value= prop->maxThreadsDim[0];
@@ -1504,13 +1504,13 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun )
 	{
 		dim3 gridDim = config.grid_dim();
 		dim3 blockDim = config.block_dim();
-		if (gridDim.x * gridDim.y * gridDim.z == 0 || blockDim.x * blockDim.y * blockDim.z == 0)
-		{
+		//if (gridDim.x * gridDim.y * gridDim.z == 0 || blockDim.x * blockDim.y * blockDim.z == 0)
+		//{
 			//can't launch
-			printf("can't launch a empty kernel\n");
-			g_cuda_launch_stack.pop_back();
-			return g_last_cudaError = cudaErrorInvalidConfiguration;
-		}
+		//	printf("can't launch a empty kernel\n");
+		//	g_cuda_launch_stack.pop_back();
+		//	return g_last_cudaError = cudaErrorInvalidConfiguration;
+		//}
 	}
 	struct CUstream_st *stream = config.get_stream();
 	if(g_stream_manager->is_blocking())
@@ -3151,9 +3151,12 @@ size_t getMaxThreadsPerBlock(struct cudaFuncAttributes *attr) {
 
   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;
 }
 
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index 6de5845..72cb32b 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -320,7 +320,7 @@ void shader_core_config::reg_options(class OptionParser * opp)
     option_parser_register(opp, "-adaptive_volta_cache_config", OPT_BOOL, &adaptive_volta_cache_config,
                  "adaptive_volta_cache_config",
                  "0");
-    option_parser_register(opp, "-gpgpu_shmem_size", OPT_UINT32, &gpgpu_shmem_sizeDefault,
+    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,
@@ -1065,7 +1065,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();
 	    }
 	}