From fe6506a14839b8614b5f4daa2d8331441feb6902 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Mon, 6 May 2019 20:23:39 -0400 Subject: adding volta configs --- .../tested-cfgs/SM7_V100/config_volta_islip.icnt | 74 ++++++++ configs/tested-cfgs/SM7_V100/gpgpusim.config | 188 ++++++++++++++++++++ .../SM7_V100_fastlocalxbar/gpgpusim.config | 191 +++++++++++++++++++++ 3 files changed, 453 insertions(+) create mode 100644 configs/tested-cfgs/SM7_V100/config_volta_islip.icnt create mode 100644 configs/tested-cfgs/SM7_V100/gpgpusim.config create mode 100644 configs/tested-cfgs/SM7_V100_fastlocalxbar/gpgpusim.config (limited to 'configs') diff --git a/configs/tested-cfgs/SM7_V100/config_volta_islip.icnt b/configs/tested-cfgs/SM7_V100/config_volta_islip.icnt new file mode 100644 index 0000000..f6e0833 --- /dev/null +++ b/configs/tested-cfgs/SM7_V100/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 = 164; +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_V100/gpgpusim.config b/configs/tested-cfgs/SM7_V100/gpgpusim.config new file mode 100644 index 0000000..d7ebe39 --- /dev/null +++ b/configs/tested-cfgs/SM7_V100/gpgpusim.config @@ -0,0 +1,188 @@ +# 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 + +# 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 ::: +# Volta NVIDIA V100 clock domains are adopted from +# https://en.wikipedia.org/wiki/Nvidia_Tesla +-gpgpu_clock_domains 1200.0:1200.0:1200.0:875.0 +# boost mode +# -gpgpu_clock_domains 1455.0:1455.0:1455.0:875.0 + +# shader core pipeline config +-gpgpu_shader_registers 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 + +# ::,::::,::,:** +# ** 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 +-adpative_volta_cache_config 1 +# Volta unified cache has four ports +-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 +-gmem_skip_L1D 0 +-icnt_flit_size 40 +-gpgpu_n_cluster_ejection_buffer_size 32 +-l1_latency 28 +-smem_latency 19 +-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 +-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 +# 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 +# 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 reisiter 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 + +# memory partition latency config +-rop_latency 120 +-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, Tesla TITANV 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_V100_fastlocalxbar/gpgpusim.config b/configs/tested-cfgs/SM7_V100_fastlocalxbar/gpgpusim.config new file mode 100644 index 0000000..401aafd --- /dev/null +++ b/configs/tested-cfgs/SM7_V100_fastlocalxbar/gpgpusim.config @@ -0,0 +1,191 @@ +# 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 + +# 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 ::: +# Volta NVIDIA V100 clock domains are adopted from +# https://en.wikipedia.org/wiki/Nvidia_Tesla +-gpgpu_clock_domains 1200.0:1200.0:1200.0:875.0 +# boost mode +# -gpgpu_clock_domains 1455.0:1455.0:1455.0:875.0 + +# shader core pipeline config +-gpgpu_shader_registers 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 + +# ::,::::,::,:** +# ** 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 +-adpative_volta_cache_config 1 +# Volta unified cache has four ports +-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 +-gmem_skip_L1D 0 +-icnt_flit_size 40 +-gpgpu_n_cluster_ejection_buffer_size 32 +-l1_latency 28 +-smem_latency 19 +-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 +-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 +# 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 +# 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 reisiter 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 + +# use local xbar interconnection +-network_mode 2 +-inct_in_buffer_limit 512 +-inct_out_buffer_limit 512 +-inct_subnets 2 +-fast_execution_mode 1 + +# memory partition latency config +-rop_latency 120 +-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, Tesla 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 + -- cgit v1.3 From d82c8bab428ebbb0d3240bf039dd81da21807a13 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Wed, 8 May 2019 11:39:11 -0400 Subject: fixing volta config --- configs/tested-cfgs/SM7_V100/gpgpusim.config | 2 +- configs/tested-cfgs/SM7_V100_fastlocalxbar/gpgpusim.config | 3 +-- 2 files changed, 2 insertions(+), 3 deletions(-) (limited to 'configs') diff --git a/configs/tested-cfgs/SM7_V100/gpgpusim.config b/configs/tested-cfgs/SM7_V100/gpgpusim.config index d7ebe39..ef93cea 100644 --- a/configs/tested-cfgs/SM7_V100/gpgpusim.config +++ b/configs/tested-cfgs/SM7_V100/gpgpusim.config @@ -75,7 +75,7 @@ # 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 --adpative_volta_cache_config 1 +-adaptive_volta_cache_config 1 # Volta unified cache has four ports -mem_unit_ports 4 -gpgpu_cache:dl1 S:4:128:64,L:L:s:N:L,A:256:8,16:0,32 diff --git a/configs/tested-cfgs/SM7_V100_fastlocalxbar/gpgpusim.config b/configs/tested-cfgs/SM7_V100_fastlocalxbar/gpgpusim.config index 401aafd..66f5251 100644 --- a/configs/tested-cfgs/SM7_V100_fastlocalxbar/gpgpusim.config +++ b/configs/tested-cfgs/SM7_V100_fastlocalxbar/gpgpusim.config @@ -75,7 +75,7 @@ # 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 --adpative_volta_cache_config 1 +-adaptive_volta_cache_config 1 # Volta unified cache has four ports -mem_unit_ports 4 -gpgpu_cache:dl1 S:4:128:64,L:L:s:N:L,A:256:8,16:0,32 @@ -128,7 +128,6 @@ -inct_in_buffer_limit 512 -inct_out_buffer_limit 512 -inct_subnets 2 --fast_execution_mode 1 # memory partition latency config -rop_latency 120 -- cgit v1.3 From 24659c165247c94eedab5615096f3384f9207e20 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Tue, 28 May 2019 18:04:34 -0400 Subject: adding QV100 config --- .../tested-cfgs/SM7_QV100/config_volta_islip.icnt | 74 ++++++++ configs/tested-cfgs/SM7_QV100/gpgpusim.config | 201 +++++++++++++++++++++ configs/tested-cfgs/SM7_TITANV/gpgpusim.config | 4 +- .../SM7_TITANV_fastlocalxbar/gpgpusim.config | 190 ------------------- .../tested-cfgs/SM7_V100/config_volta_islip.icnt | 74 -------- configs/tested-cfgs/SM7_V100/gpgpusim.config | 188 ------------------- .../SM7_V100_fastlocalxbar/gpgpusim.config | 190 ------------------- 7 files changed, 278 insertions(+), 643 deletions(-) create mode 100644 configs/tested-cfgs/SM7_QV100/config_volta_islip.icnt create mode 100644 configs/tested-cfgs/SM7_QV100/gpgpusim.config delete mode 100644 configs/tested-cfgs/SM7_TITANV_fastlocalxbar/gpgpusim.config delete mode 100644 configs/tested-cfgs/SM7_V100/config_volta_islip.icnt delete mode 100644 configs/tested-cfgs/SM7_V100/gpgpusim.config delete mode 100644 configs/tested-cfgs/SM7_V100_fastlocalxbar/gpgpusim.config (limited to 'configs') 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..f6e0833 --- /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 = 164; +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..16c131b --- /dev/null +++ b/configs/tested-cfgs/SM7_QV100/gpgpusim.config @@ -0,0 +1,201 @@ +# 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 ::: +# 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_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 + +# ::,::::,::,:** +# ** 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_volta_cache_config 1 +# Volta unified cache has four ports +-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 +-gmem_skip_L1D 0 +-icnt_flit_size 40 +-gpgpu_n_cluster_ejection_buffer_size 32 +-l1_latency 28 +-smem_latency 19 +-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 +-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 +# 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 +# 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 120 +-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..b06f048 100644 --- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config +++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config @@ -137,6 +137,8 @@ # 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 @@ -161,7 +163,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/configs/tested-cfgs/SM7_TITANV_fastlocalxbar/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV_fastlocalxbar/gpgpusim.config deleted file mode 100644 index 4625ac4..0000000 --- a/configs/tested-cfgs/SM7_TITANV_fastlocalxbar/gpgpusim.config +++ /dev/null @@ -1,190 +0,0 @@ -# This config models the Volta Titan V -# 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 - -# 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 40 --gpgpu_n_cores_per_cluster 2 --gpgpu_n_mem 24 --gpgpu_n_sub_partition_per_mchannel 2 - -# volta clock domains -#-gpgpu_clock_domains ::: -# Volta NVIDIA TITANV clock domains are adopted from -# https://en.wikipedia.org/wiki/Volta_(microarchitecture) --gpgpu_clock_domains 1200.0:1200.0:1200.0:850.0 -# boost mode -# -gpgpu_clock_domains 1455.0:1455.0:1455.0:850.0 - -# shader core pipeline config --gpgpu_shader_registers 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 TITANV 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 - -# ::,::::,::,:** -# ** 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_volta_cache_config 1 -# Volta unified cache has four ports --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 --gmem_skip_L1D 0 --icnt_flit_size 40 --gpgpu_n_cluster_ejection_buffer_size 32 --l1_latency 28 --smem_latency 19 --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 --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 0 - -# 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 -# 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 - -# use local xbar interconnection --network_mode 2 --inct_in_buffer_limit 512 --inct_out_buffer_limit 512 --inct_subnets 2 - -# memory partition latency config --rop_latency 120 --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, Tesla TITANV 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_V100/config_volta_islip.icnt b/configs/tested-cfgs/SM7_V100/config_volta_islip.icnt deleted file mode 100644 index f6e0833..0000000 --- a/configs/tested-cfgs/SM7_V100/config_volta_islip.icnt +++ /dev/null @@ -1,74 +0,0 @@ -//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 = 164; -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_V100/gpgpusim.config b/configs/tested-cfgs/SM7_V100/gpgpusim.config deleted file mode 100644 index ef93cea..0000000 --- a/configs/tested-cfgs/SM7_V100/gpgpusim.config +++ /dev/null @@ -1,188 +0,0 @@ -# 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 - -# 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 ::: -# Volta NVIDIA V100 clock domains are adopted from -# https://en.wikipedia.org/wiki/Nvidia_Tesla --gpgpu_clock_domains 1200.0:1200.0:1200.0:875.0 -# boost mode -# -gpgpu_clock_domains 1455.0:1455.0:1455.0:875.0 - -# shader core pipeline config --gpgpu_shader_registers 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 - -# ::,::::,::,:** -# ** 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_volta_cache_config 1 -# Volta unified cache has four ports --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 --gmem_skip_L1D 0 --icnt_flit_size 40 --gpgpu_n_cluster_ejection_buffer_size 32 --l1_latency 28 --smem_latency 19 --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 --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 -# 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 -# 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 reisiter 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 - -# memory partition latency config --rop_latency 120 --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, Tesla TITANV 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_V100_fastlocalxbar/gpgpusim.config b/configs/tested-cfgs/SM7_V100_fastlocalxbar/gpgpusim.config deleted file mode 100644 index 66f5251..0000000 --- a/configs/tested-cfgs/SM7_V100_fastlocalxbar/gpgpusim.config +++ /dev/null @@ -1,190 +0,0 @@ -# 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 - -# 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 ::: -# Volta NVIDIA V100 clock domains are adopted from -# https://en.wikipedia.org/wiki/Nvidia_Tesla --gpgpu_clock_domains 1200.0:1200.0:1200.0:875.0 -# boost mode -# -gpgpu_clock_domains 1455.0:1455.0:1455.0:875.0 - -# shader core pipeline config --gpgpu_shader_registers 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 - -# ::,::::,::,:** -# ** 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_volta_cache_config 1 -# Volta unified cache has four ports --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 --gmem_skip_L1D 0 --icnt_flit_size 40 --gpgpu_n_cluster_ejection_buffer_size 32 --l1_latency 28 --smem_latency 19 --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 --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 -# 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 -# 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 reisiter 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 - -# use local xbar interconnection --network_mode 2 --inct_in_buffer_limit 512 --inct_out_buffer_limit 512 --inct_subnets 2 - -# memory partition latency config --rop_latency 120 --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, Tesla 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 - -- cgit v1.3 From dc1d8114a3088de791833c336a2d2484a00a361f Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Tue, 28 May 2019 18:23:28 -0400 Subject: fixing quadro file --- configs/tested-cfgs/SM7_QV100/config_volta_islip.icnt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'configs') diff --git a/configs/tested-cfgs/SM7_QV100/config_volta_islip.icnt b/configs/tested-cfgs/SM7_QV100/config_volta_islip.icnt index f6e0833..5ad7ecd 100644 --- a/configs/tested-cfgs/SM7_QV100/config_volta_islip.icnt +++ b/configs/tested-cfgs/SM7_QV100/config_volta_islip.icnt @@ -7,7 +7,7 @@ network_count = 2; // Topology topology = fly; -k = 164; +k = 144; n = 1; // Routing -- cgit v1.3 From 17815dfdd7b6e9b558997bc5a9ba157a0493da16 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Fri, 12 Jul 2019 21:31:54 -0400 Subject: fixing device and function parameters config --- configs/tested-cfgs/SM2_GTX480/gpgpusim.config | 1 + configs/tested-cfgs/SM6_TITANX/gpgpusim.config | 1 + configs/tested-cfgs/SM7_QV100/gpgpusim.config | 11 +++++++---- configs/tested-cfgs/SM7_TITANV/gpgpusim.config | 3 +++ libcuda/cuda_runtime_api.cc | 21 ++++++++++++--------- src/gpgpu-sim/gpu-sim.cc | 4 ++-- 6 files changed, 26 insertions(+), 15 deletions(-) (limited to 'configs') 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;in_simt_clusters;i++) { - m_cluster[i]->cache_flush(); + m_cluster[i]->cache_invalidate(); } } -- cgit v1.3 From e5e8fd34aa037aa510fa3936fa133f6b345e1825 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Thu, 18 Jul 2019 17:41:58 -0400 Subject: increase the tex cache in Volta --- configs/tested-cfgs/SM7_QV100/gpgpusim.config | 6 +++--- configs/tested-cfgs/SM7_TITANV/gpgpusim.config | 2 +- 2 files changed, 4 insertions(+), 4 deletions(-) (limited to 'configs') diff --git a/configs/tested-cfgs/SM7_QV100/gpgpusim.config b/configs/tested-cfgs/SM7_QV100/gpgpusim.config index 5f64908..b2fe5c0 100644 --- a/configs/tested-cfgs/SM7_QV100/gpgpusim.config +++ b/configs/tested-cfgs/SM7_QV100/gpgpusim.config @@ -101,7 +101,7 @@ -smem_latency 19 -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 +# 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 @@ -110,9 +110,9 @@ # 128 KB Inst. -gpgpu_cache:il1 N:64:128:16,L:R:f:N:L,S:2:48,4 -# 48 KB Tex +# 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: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 diff --git a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config index 6c21dcb..786e75e 100644 --- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config +++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config @@ -112,7 +112,7 @@ -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 -- cgit v1.3 From 8127368f2611b0472dc729f3460ff091e87f51c9 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Sun, 21 Jul 2019 20:54:12 -0400 Subject: fixing cache/mem latency in V100 config --- configs/tested-cfgs/SM7_QV100/gpgpusim.config | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) (limited to 'configs') diff --git a/configs/tested-cfgs/SM7_QV100/gpgpusim.config b/configs/tested-cfgs/SM7_QV100/gpgpusim.config index b2fe5c0..f70ad55 100644 --- a/configs/tested-cfgs/SM7_QV100/gpgpusim.config +++ b/configs/tested-cfgs/SM7_QV100/gpgpusim.config @@ -97,7 +97,7 @@ -gmem_skip_L1D 0 -icnt_flit_size 40 -gpgpu_n_cluster_ejection_buffer_size 32 --l1_latency 28 +-l1_latency 80 -smem_latency 19 -gpgpu_flush_l1_cache 1 @@ -144,7 +144,7 @@ # "-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 -- cgit v1.3 From c2a1e3a668f9a88239184e13460f7e1725b15c90 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Mon, 26 Aug 2019 12:17:52 -0400 Subject: Banked L1, adding iSLIP and RR arbiteratio and adding some comments --- configs/tested-cfgs/SM7_QV100/gpgpusim.config | 11 +- configs/tested-cfgs/SM7_TITANV/gpgpusim.config | 13 +- src/gpgpu-sim/addrdec.cc | 2 + src/gpgpu-sim/gpu-cache.cc | 8 ++ src/gpgpu-sim/gpu-cache.h | 4 + src/gpgpu-sim/icnt_wrapper.cc | 2 + src/gpgpu-sim/local_interconnect.cc | 40 ++++-- src/gpgpu-sim/shader.cc | 168 +++++++++++++------------ src/gpgpu-sim/shader.h | 2 +- 9 files changed, 152 insertions(+), 98 deletions(-) (limited to 'configs') diff --git a/configs/tested-cfgs/SM7_QV100/gpgpusim.config b/configs/tested-cfgs/SM7_QV100/gpgpusim.config index f70ad55..f807e11 100644 --- a/configs/tested-cfgs/SM7_QV100/gpgpusim.config +++ b/configs/tested-cfgs/SM7_QV100/gpgpusim.config @@ -87,9 +87,10 @@ # 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 @@ -97,8 +98,8 @@ -gmem_skip_L1D 0 -icnt_flit_size 40 -gpgpu_n_cluster_ejection_buffer_size 32 --l1_latency 80 --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 6MB L2 cache diff --git a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config index 786e75e..f8e4afe 100644 --- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config +++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config @@ -88,8 +88,9 @@ # 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 +# 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 @@ -97,8 +98,8 @@ -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 @@ -106,7 +107,7 @@ -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 @@ -144,7 +145,7 @@ # "-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 diff --git a/src/gpgpu-sim/addrdec.cc b/src/gpgpu-sim/addrdec.cc index ca88ec9..b0db034 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::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 370f6e6..db9701d 100644 --- a/src/gpgpu-sim/gpu-cache.cc +++ b/src/gpgpu-sim/gpu-cache.cc @@ -63,6 +63,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 337f710..90adbb5 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/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/local_interconnect.cc b/src/gpgpu-sim/local_interconnect.cc index 1416b2c..da8a65c 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 issued(total_nodes, false); for(unsigned i=0; i node_tmp; + + //calcaulte how many conflicts are there for stats for (unsigned i=0; i(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(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/shader.cc b/src/gpgpu-sim/shader.cc index e38eefd..ffd3035 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -1645,35 +1645,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()); - 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; jm_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()); + 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; - inst.accessq_pop_back(); + 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(); + } + 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()); std::list events; enum cache_request_status status = cache->access(mf->get_addr(),mf,gpu_sim_cycle+gpu_tot_sim_cycle,events); return process_cache_access( cache, mf->get_addr(), inst, events, mf, status ); @@ -1683,62 +1694,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 >::iterator it = m_latency_queue.begin(); - if((l1_latency_queue[0]) != NULL) - { - mem_fetch* mf_next = l1_latency_queue[0]; - std::list events; - enum cache_request_status status = m_L1D->access(mf_next->get_addr(),mf_next,gpu_sim_cycle+gpu_tot_sim_cycle,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) + for(int j=0; jm_L1D_config.l1_banks; j++) { + if((l1_latency_queue[j][0]) != NULL) + { + mem_fetch* mf_next = l1_latency_queue[j][0]; + std::list events; + enum cache_request_status status = m_L1D->access(mf_next->get_addr(),mf_next,gpu_sim_cycle+gpu_tot_sim_cycle,events); + + bool write_sent = was_write_sent(events); + bool read_sent = was_read_sent(events); + + 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; - - mf_next->set_reply(); - - for(unsigned i=0; i< dec_ack; ++i) - m_core->store_ack(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; - } - } + } - for( unsigned stage = 0; stagem_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 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(); + + for(unsigned i=0; i< dec_ack; ++i) + m_core->store_ack(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[j][0] = NULL; + } + } + + for( unsigned stage = 0; stagem_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; + } + } } @@ -2117,11 +2130,12 @@ ldst_unit::ldst_unit( mem_fetch_interface *icnt, m_mf_allocator, IN_L1D_MISS_QUEUE ); - if(m_config->m_L1D_config.l1_latency > 0) - { - for(int i=0; im_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(int j=0; jm_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 a0c2b63..05476f4 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1325,7 +1325,7 @@ protected: unsigned long long m_last_inst_gpu_sim_cycle; unsigned long long m_last_inst_gpu_tot_sim_cycle; - std::deque l1_latency_queue; + std::vector> l1_latency_queue; void L1_latency_queue_cycle(); }; -- cgit v1.3 From 2e2b0911f8cc9d172e4602765ce4c20ffb583cf0 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Mon, 26 Aug 2019 14:11:16 -0400 Subject: update titanV config file --- configs/tested-cfgs/SM7_TITANV/gpgpusim.config | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'configs') diff --git a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config index f8e4afe..888ce71 100644 --- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config +++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config @@ -87,7 +87,7 @@ # 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 +-adaptive_cache_config 1 # Volta unified cache has four banks -l1_banks 4 #-mem_unit_ports 4 -- cgit v1.3 From 93597743b24b7d7726096d337e6c9c4516cb2273 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Fri, 30 Aug 2019 18:43:21 -0400 Subject: update the adaptive cache behaviour and make L1 fully assoc in Volta --- configs/tested-cfgs/SM7_QV100/gpgpusim.config | 2 +- configs/tested-cfgs/SM7_TITANV/gpgpusim.config | 2 +- src/abstract_hardware_model.cc | 4 +- src/abstract_hardware_model.h | 9 ++++- src/gpgpu-sim/addrdec.cc | 4 +- src/gpgpu-sim/gpu-cache.cc | 7 +++- src/gpgpu-sim/gpu-cache.h | 5 +++ src/gpgpu-sim/gpu-sim.cc | 2 +- src/gpgpu-sim/shader.cc | 54 ++++++++++++++++---------- 9 files changed, 59 insertions(+), 30 deletions(-) (limited to 'configs') diff --git a/configs/tested-cfgs/SM7_QV100/gpgpusim.config b/configs/tested-cfgs/SM7_QV100/gpgpusim.config index f807e11..c0d22ee 100644 --- a/configs/tested-cfgs/SM7_QV100/gpgpusim.config +++ b/configs/tested-cfgs/SM7_QV100/gpgpusim.config @@ -91,7 +91,7 @@ # 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_cache:dl1 S:1:128:256,L:L:s:N:L,A:256:8,16:0,32 -gpgpu_shmem_size 98304 -gpgpu_shmem_sizeDefault 98304 -gpgpu_shmem_per_block 65536 diff --git a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config index 888ce71..0339b0d 100644 --- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config +++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config @@ -91,7 +91,7 @@ # 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_cache:dl1 S:1:128:256,L:L:s:N:L,A:256:8,16:0,32 -gpgpu_shmem_size 98304 -gpgpu_shmem_sizeDefault 98304 -gpgpu_shmem_per_block 65536 diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 023f51b..35a3984 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -710,7 +710,7 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * //Jin: launch latency management m_launch_latency = g_kernel_launch_latency; - volta_cache_config_set=false; + cache_config_set=false; } /*A snapshot of the texture mappings needs to be stored in the kernel's info as @@ -735,7 +735,7 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * //Jin: launch latency management m_launch_latency = g_kernel_launch_latency; - volta_cache_config_set=false; + cache_config_set=false; m_NameToCudaArray = nameToCudaArray; m_NameToTextureInfo = nameToTextureInfo; } diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 27a1ba6..231b6a2 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -65,6 +65,11 @@ enum FuncCache FuncCachePreferL1 = 2 }; +enum AdaptiveCache +{ + FIXED = 0, + VOLTA = 1 +}; #ifdef __cplusplus @@ -345,7 +350,7 @@ public: unsigned long long end_cycle; unsigned m_launch_latency; - mutable bool volta_cache_config_set; + mutable bool cache_config_set; }; struct core_config { @@ -388,7 +393,7 @@ struct core_config { unsigned gpgpu_max_insn_issue_per_warp; bool gmem_skip_L1D; // on = global memory access always skip the L1 cache - bool adaptive_volta_cache_config; + unsigned adaptive_cache_config; }; // bounded stack that implements simt reconvergence using pdom mechanism from MICRO'07 paper diff --git a/src/gpgpu-sim/addrdec.cc b/src/gpgpu-sim/addrdec.cc index b0db034..3262456 100644 --- a/src/gpgpu-sim/addrdec.cc +++ b/src/gpgpu-sim/addrdec.cc @@ -124,8 +124,8 @@ void linear_to_raw_address_translation::addrdec_tlx(new_addr_type addr, addrdec_ * Rau, B. R et al. * ISCA 1991 * - * equations are adopted from: - * "Sacat: streaming-aware conflict-avoiding thrashing-resistant gpgpu cache management scheme." + * equations are corresponding to IPOLY(37) and are adopted from: + * "SACAT: streaming-aware conflict-avoiding thrashing-resistant gpgpu cache management scheme." * Khairy et al. * IEEE TPDS 2017. */ diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc index db9701d..8d00ea9 100644 --- a/src/gpgpu-sim/gpu-cache.cc +++ b/src/gpgpu-sim/gpu-cache.cc @@ -65,8 +65,13 @@ const char * cache_fail_status_str(enum cache_reservation_fail_reason status) unsigned l1d_cache_config::set_bank(new_addr_type addr) const{ - if(m_cache_type == SECTOR) + //For sector cache, we select one sector per bank (sector interleaving) + //This is what was found in Volta (one sector per bank, sector interleaving) + //otherwise, line interleaving + if(m_cache_type == SECTOR) { + //assert(l1_banks == SECTOR_CHUNCK_SIZE); return (addr >> m_sector_sz_log2) & (l1_banks-1); + } else return (addr >> m_line_sz_log2) & (l1_banks-1); } diff --git a/src/gpgpu-sim/gpu-cache.h b/src/gpgpu-sim/gpu-cache.h index 90adbb5..dd22886 100644 --- a/src/gpgpu-sim/gpu-cache.h +++ b/src/gpgpu-sim/gpu-cache.h @@ -656,6 +656,11 @@ public: assert( m_valid ); return MAX_DEFAULT_CACHE_SIZE_MULTIBLIER * m_nset * original_m_assoc; } + unsigned get_max_assoc() const + { + assert( m_valid ); + return MAX_DEFAULT_CACHE_SIZE_MULTIBLIER * original_m_assoc; + } void print( FILE *fp ) const { fprintf( fp, "Size = %d B (%d Set x %d-way x %d byte line)\n", diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 4f071c7..ed94865 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -323,7 +323,7 @@ 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_cache_config", OPT_BOOL, &adaptive_volta_cache_config, + option_parser_register(opp, "-adaptive_cache_config", OPT_UINT32, &adaptive_cache_config, "adaptive_cache_config", "0"); option_parser_register(opp, "-gpgpu_shmem_sizeDefault", OPT_UINT32, &gpgpu_shmem_sizeDefault, diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index ffd3035..ff2bf3f 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -2980,33 +2980,47 @@ unsigned int shader_core_config::max_cta( const kernel_info_t &k ) const abort(); } - if(adaptive_volta_cache_config && !k.volta_cache_config_set) { - //For Volta, we assign the remaining shared memory to L1 cache - //For more info, see https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory-7-x + if(adaptive_cache_config && !k.cache_config_set) { + //For more info about adaptive cache, see https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory-7-x unsigned total_shmed = kernel_info->smem * result; assert(total_shmed >=0 && total_shmed <= gpgpu_shmem_size); - assert(gpgpu_shmem_size == 98304); //Volta has 96 KB shared - assert(m_L1D_config.get_nset() == 4); //Volta L1 has four sets + //assert(gpgpu_shmem_size == 98304); //Volta has 96 KB shared + //assert(m_L1D_config.get_nset() == 4); //Volta L1 has four sets if(total_shmed < gpgpu_shmem_size){ - if(total_shmed == 0) - m_L1D_config.set_assoc(256); //L1 is 128KB ans shd=0 - else if(total_shmed > 0 && total_shmed <= 8192) - m_L1D_config.set_assoc(240); //L1 is 120KB ans shd=8KB - else if(total_shmed > 8192 && total_shmed <= 16384) - m_L1D_config.set_assoc(224); //L1 is 112KB ans shd=16KB - else if(total_shmed > 16384 && total_shmed <= 32768) - m_L1D_config.set_assoc(192); //L1 is 96KB ans shd=32KB - else if(total_shmed > 32768 && total_shmed <= 65536) - m_L1D_config.set_assoc(128); //L1 is 64KB ans shd=64KB - else if(total_shmed > 65536 && total_shmed <= gpgpu_shmem_size) - m_L1D_config.set_assoc(64); //L1 is 32KB and shd=96KB - else + switch (adaptive_cache_config) { + case FIXED: + break; + case VOLTA: { + //For Volta, we assign the remaining shared memory to L1 cache + //For more info about adaptive cache, see https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory-7-x + assert(gpgpu_shmem_size == 98304); //Volta has 96 KB shared + + //To Do: make it flexible and not tuned to 9KB share memory + unsigned max_assoc = m_L1D_config.get_max_assoc(); + if(total_shmed == 0) + m_L1D_config.set_assoc(max_assoc); //L1 is 128KB and shd=0 + else if(total_shmed > 0 && total_shmed <= 8192) + m_L1D_config.set_assoc(0.9375 * max_assoc); //L1 is 120KB and shd=8KB + else if(total_shmed > 8192 && total_shmed <= 16384) + m_L1D_config.set_assoc(0.875 * max_assoc); //L1 is 112KB and shd=16KB + else if(total_shmed > 16384 && total_shmed <= 32768) + m_L1D_config.set_assoc(0.75 * max_assoc); //L1 is 96KB and shd=32KB + else if(total_shmed > 32768 && total_shmed <= 65536) + m_L1D_config.set_assoc(0.5 * max_assoc); //L1 is 64KB and shd=64KB + else if(total_shmed > 65536 && total_shmed <= gpgpu_shmem_size) + m_L1D_config.set_assoc(0.25 * max_assoc); //L1 is 32KB and shd=96KB + else + assert(0); + break; + } + default: assert(0); + } - printf ("GPGPU-Sim: Reconfigure L1 cache in Volta Archi to %uKB\n", m_L1D_config.get_total_size_inKB()); + printf ("GPGPU-Sim: Reconfigure L1 cache to %uKB\n", m_L1D_config.get_total_size_inKB()); } - k.volta_cache_config_set = true; + k.cache_config_set = true; } return result; -- cgit v1.3 From 326a038bc3b7386ccd2f16ebee6e3cc44eb677e5 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Fri, 4 Oct 2019 21:32:38 -0400 Subject: increase RF banks --- configs/tested-cfgs/SM7_QV100/gpgpusim.config | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) (limited to 'configs') diff --git a/configs/tested-cfgs/SM7_QV100/gpgpusim.config b/configs/tested-cfgs/SM7_QV100/gpgpusim.config index c0d22ee..23a57fa 100644 --- a/configs/tested-cfgs/SM7_QV100/gpgpusim.config +++ b/configs/tested-cfgs/SM7_QV100/gpgpusim.config @@ -122,11 +122,13 @@ -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_units_gen 32 -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 +# Volta has 8 banks, 4 schedulers, two banks per scheduler +# However, since we do not model register file cache (RFC), we increase the number of banks to 32 to +# resuce banks conflits, mitigate RFC impact, and matche the hardware performance. To Do: RFC +-gpgpu_num_reg_banks 32 # shared memory bankconflict detection -gpgpu_shmem_num_banks 32 -- cgit v1.3