summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--configs/tested-cfgs/SM2_GTX480/gpgpusim.config1
-rw-r--r--configs/tested-cfgs/SM6_TITANX/gpgpusim.config1
-rw-r--r--configs/tested-cfgs/SM7_QV100/gpgpusim.config11
-rw-r--r--configs/tested-cfgs/SM7_TITANV/gpgpusim.config3
-rw-r--r--libcuda/cuda_runtime_api.cc21
-rw-r--r--src/gpgpu-sim/gpu-sim.cc4
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();
}
}
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 21:22:59 +0000