From c2917b4c93a84a6645aa7d7c5ac1c389b3369586 Mon Sep 17 00:00:00 2001 From: tgrogers Date: Tue, 20 Nov 2018 15:34:03 -0500 Subject: Fixing the compile in gcc 4.4.x --- src/cuda-sim/cuda-sim.cc | 2 +- src/cuda-sim/instructions.cc | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 3773f6f..ec68b5b 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1308,7 +1308,7 @@ void function_info::list_param( FILE *fout ) const void function_info::ptx_jit_config(std::map mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) { - static unsigned long counter = 0; + static unsigned long long counter = 0; std::vector< std::pair > param_data; std::vector offsets; std::vector paramIsPointer; diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 31a33c6..c85654c 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -195,7 +195,7 @@ void ptx_thread_info::print_reg_thread(char * fname) { reg_map_t reg = m_regs.back(); - typename reg_map_t::const_iterator it; + reg_map_t::const_iterator it; for (it = reg.begin(); it != reg.end(); ++it) { const std::string &name = it->first->name(); -- cgit v1.3 From 8b8e35b9e8cb7b1e18ea2faf06450f8052e48b7e Mon Sep 17 00:00:00 2001 From: Deval Shah Date: Sun, 24 Mar 2019 15:05:10 -0700 Subject: Parse bar.warp.sync as NOP operation ( which is acceptable in Stack based reconvergence ) --- src/cuda-sim/ptx.l | 1 + 1 file changed, 1 insertion(+) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index d4bf631..3232361 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -61,6 +61,7 @@ addc TC; ptx_lval.int_value = ADDC_OP; return OPCODE; and TC; ptx_lval.int_value = AND_OP; return OPCODE; andn TC; ptx_lval.int_value = ANDN_OP; return OPCODE; atom TC; ptx_lval.int_value = ATOM_OP; return OPCODE; +bar.warp TC; ptx_lval.int_value = NOP_OP; return OPCODE; bar TC; ptx_lval.int_value = BAR_OP; return OPCODE; bfe TC; ptx_lval.int_value = BFE_OP; return OPCODE; bfi TC; ptx_lval.int_value = BFI_OP; return OPCODE; -- cgit v1.3 From 73050b730b31e9f58f8c2a0bfc4b3a8cd92be7ff Mon Sep 17 00:00:00 2001 From: Suchita Pati Date: Wed, 10 Apr 2019 23:27:42 -0500 Subject: Adding texture Fix --- libcuda/cuda_runtime_api.cc | 28 ++++++++++++++++++++++++---- src/abstract_hardware_model.cc | 4 +++- src/abstract_hardware_model.h | 22 +++++++++++++++++++++- src/cuda-sim/cuda-sim.cc | 5 +++-- src/cuda-sim/cuda_device_runtime.cc | 3 ++- src/cuda-sim/instructions.cc | 5 +++-- 6 files changed, 56 insertions(+), 11 deletions(-) (limited to 'src/cuda-sim') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index cf16d9b..7569b20 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1204,8 +1204,24 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit li if(g_debug_execution >= 3){ announce_call(__my_func__); } - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaSuccess; + switch(limit) { + case 0: + *pValue=1024; + break; + case 2: + *pValue=8388608; + break; + case 3: + *pValue=2; + break; + case 4: + *pValue=2048; + break; + default: + printf("ERROR:Limit %s unimplemented \n",limit); + abort(); + } + return g_last_cudaError = cudaSuccess; } @@ -1571,7 +1587,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) #if (CUDART_VERSION >= 3000) //synchronization required for application using external libraries without explicit synchronization in the code to //avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. - synchronize(); + stream->synchronize(); g_stream_manager->destroy_stream(stream); #endif return g_last_cudaError = cudaSuccess; @@ -2702,6 +2718,7 @@ cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ if(g_debug_execution >= 3){ announce_call(__my_func__); } + synchronize(); return g_last_cudaError = cudaSuccess; } @@ -3135,6 +3152,8 @@ __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t v } return g_last_cudaError = cudaSuccess; } + + #endif #endif @@ -3313,7 +3332,8 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, announce_call(__my_func__); } function_info *entry = context->get_kernel(hostFun); - kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry); + gpgpu_t* gpu= context->get_device()->get_gpgpu(); + kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping()); if( entry == NULL ) { printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found for %p\n", hostFun); abort(); diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 307df40..3db87c9 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -688,7 +688,7 @@ unsigned g_kernel_launch_latency; unsigned kernel_info_t::m_next_uid = 1; -kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry ) +kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map NameToCudaArray, std::map NameToTexureInfo) { m_kernel_entry=entry; m_grid_dim=gridDim; @@ -708,6 +708,8 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * m_launch_latency = g_kernel_launch_latency; volta_cache_config_set=false; + m_NameToCudaArray = NameToCudaArray; + m_NameToTexureInfo = NameToTexureInfo; } kernel_info_t::~kernel_info_t() diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 28a22e0..bbc31e5 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -212,7 +212,7 @@ public: // m_num_cores_running=0; // m_param_mem=NULL; // } - kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry ); + kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map m_NameToCudaArray, std::map m_NameToTexureInfo); ~kernel_info_t(); void inc_running() { m_num_cores_running++; } @@ -275,6 +275,20 @@ public: std::list &active_threads() { return m_active_threads; } class memory_space *get_param_memory() { return m_param_mem; } + const struct cudaArray* get_texarray( const std::string &texname ) const + { + std::map::const_iterator t=m_NameToCudaArray.find(texname); + assert(t != m_NameToCudaArray.end()); + return t->second; + } + + const struct textureInfo* get_texinfo( const std::string &texname ) const + { + std::map::const_iterator t=m_NameToTexureInfo.find(texname); + assert(t != m_NameToTexureInfo.end()); + return t->second; + } + private: kernel_info_t( const kernel_info_t & ); // disable copy constructor void operator=( const kernel_info_t & ); // disable copy operator @@ -284,6 +298,9 @@ private: unsigned m_uid; static unsigned m_next_uid; + std::map m_NameToCudaArray; + std::map m_NameToTexureInfo; + dim3 m_grid_dim; dim3 m_block_dim; dim3 m_next_cta; @@ -595,6 +612,9 @@ public: const gpgpu_functional_sim_config &get_config() const { return m_function_model_config; } FILE* get_ptx_inst_debug_file() { return ptx_inst_debug_file; } + std::map getNameArrayMapping() {return m_NameToCudaArray;} + std::map getNameInfoMapping() {return m_NameToTexureInfo;} + protected: const gpgpu_functional_sim_config &m_function_model_config; FILE* ptx_inst_debug_file; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index ec68b5b..a51ea3e 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1507,7 +1507,8 @@ static unsigned get_tex_datasize( const ptx_instruction *pI, ptx_thread_info *th std::string texname = src1.name(); gpgpu_t *gpu = thread->get_gpu(); - const struct textureInfo* texInfo = gpu->get_texinfo(texname); + kernel_info_t& k = thread->get_kernel(); + const struct textureInfo* texInfo = k.get_texinfo(texname); unsigned data_size = texInfo->texel_size; return data_size; @@ -1918,7 +1919,7 @@ kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, struct dim3 blockDim, gpgpu_t *gpu ) { - kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry); + kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping()); unsigned argcount=args.size(); unsigned argn=1; for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) { diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index b399133..917e7a8 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -198,7 +198,8 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * memory_space *device_kernel_param_mem; //create child kernel_info_t and index it with parameter_buffer address - device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry); + gpgpu_t* gpu=thread->get_gpu(); + device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry, gpu->getNameArrayMapping(), gpu->getNameInfoMapping()); device_grid->launch_cycle = gpu_sim_cycle + gpu_tot_sim_cycle; kernel_info_t & parent_grid = thread->get_kernel(); DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" << diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index c85654c..d29fc8b 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -5260,9 +5260,10 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread ) thread->get_vector_operand_values(src2, ptx_tex_regs, nelem); //ptx_reg should be 4 entry vector type...coordinates into texture gpgpu_t *gpu = thread->get_gpu(); + kernel_info_t &k = thread->get_kernel(); const struct textureReference* texref = gpu->get_texref(texname); - const struct cudaArray* cuArray = gpu->get_texarray(texname); - const struct textureInfo* texInfo = gpu->get_texinfo(texname); + const struct cudaArray* cuArray = k.get_texarray(texname); + const struct textureInfo* texInfo = k.get_texinfo(texname); const struct textureReferenceAttr* texAttr = gpu->get_texattr(texname); //assume always 2D f32 input -- cgit v1.3 From bb7ed349e146ea6b9ef1100d81fa7ac5ef2751cb Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Thu, 11 Apr 2019 19:53:23 -0400 Subject: adding local xbar and fast execution mode --- .../SM7_TITANV_fastlocalxbar/gpgpusim.config | 192 ++++++++++++++ src/cuda-sim/cuda-sim.h | 2 + src/gpgpu-sim/gpu-sim.cc | 4 + src/gpgpu-sim/icnt_wrapper.cc | 90 ++++++- src/gpgpu-sim/icnt_wrapper.h | 5 +- src/gpgpu-sim/local_interconnect.cc | 277 +++++++++++++++++++++ src/gpgpu-sim/local_interconnect.h | 116 +++++++++ src/gpgpu-sim/shader.cc | 78 +++++- src/gpgpu-sim/shader.h | 26 +- src/gpgpusim_entrypoint.cc | 6 +- 10 files changed, 773 insertions(+), 23 deletions(-) create mode 100644 configs/tested-cfgs/SM7_TITANV_fastlocalxbar/gpgpusim.config create mode 100644 src/gpgpu-sim/local_interconnect.cc create mode 100644 src/gpgpu-sim/local_interconnect.h (limited to 'src/cuda-sim') diff --git a/configs/tested-cfgs/SM7_TITANV_fastlocalxbar/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV_fastlocalxbar/gpgpusim.config new file mode 100644 index 0000000..73516e1 --- /dev/null +++ b/configs/tested-cfgs/SM7_TITANV_fastlocalxbar/gpgpusim.config @@ -0,0 +1,192 @@ +# This config models the Volta Titan X +# 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 V100 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 GV100 has 4 SP SIMD units, 4 SFU units, 4 DP units per core +## 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 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 reisiter file and EUs +# i.e. schedulers are isolated +-sub_core_model 0 +# 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 2 +-inter_config_file config_fermi_islip.icnt +-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 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 + +# Pascal has two 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 Pascal 100 +-power_simulation_enabled 0 + +# tracing functionality +#-trace_enabled 1 +#-trace_components WARP_SCHEDULER,SCOREBOARD +#-trace_sampling_core 0 + diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index abd32f9..e690356 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -47,6 +47,8 @@ extern int g_debug_thread_uid; extern void ** g_inst_classification_stat; extern void ** g_inst_op_classification_stat; extern int g_ptx_kernel_count; // used for classification stat collection purposes +extern char *opcode_latency_int, *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor; + void ptx_opcocde_latency_options (option_parser_t opp); extern class kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index ec570bf..81d9b69 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -476,6 +476,10 @@ void shader_core_config::reg_options(class OptionParser * opp) option_parser_register(opp, "-gpgpu_concurrent_kernel_sm", OPT_BOOL, &gpgpu_concurrent_kernel_sm, "Support concurrent kernels on a SM (default = disabled)", "0"); + + option_parser_register(opp, "-fast_execution_mode", OPT_BOOL, &fast_execution_mode, + "fast_execution_mode (default = disabled)", + "0"); } void gpgpu_sim_config::reg_options(option_parser_t opp) diff --git a/src/gpgpu-sim/icnt_wrapper.cc b/src/gpgpu-sim/icnt_wrapper.cc index ee58ece..075e371 100644 --- a/src/gpgpu-sim/icnt_wrapper.cc +++ b/src/gpgpu-sim/icnt_wrapper.cc @@ -29,6 +29,8 @@ #include #include "../intersim2/globals.hpp" #include "../intersim2/interconnect_interface.hpp" +#include "local_interconnect.h" + icnt_create_p icnt_create; icnt_init_p icnt_init; @@ -42,9 +44,13 @@ icnt_display_overall_stats_p icnt_display_overall_stats; icnt_display_state_p icnt_display_state; icnt_get_flit_size_p icnt_get_flit_size; -int g_network_mode; +unsigned g_network_mode; char* g_network_config_filename; + +struct inct_config g_inct_config; +LocalInterconnect *g_localicnt_interface; + #include "../option_parser.h" // Wrapper to intersim2 to accompany old icnt_wrapper @@ -105,10 +111,78 @@ static unsigned intersim2_get_flit_size() return g_icnt_interface->GetFlitSize(); } + +////////////////////////////////////////////////////// + +static void LocalInterconnect_create(unsigned int n_shader, unsigned int n_mem) +{ + g_localicnt_interface->CreateInterconnect(n_shader, n_mem); +} + +static void LocalInterconnect_init() +{ + g_localicnt_interface->Init(); +} + +static bool LocalInterconnect_has_buffer(unsigned input, unsigned int size) +{ + return g_localicnt_interface->HasBuffer(input, size); +} + +static void LocalInterconnect_push(unsigned input, unsigned output, void* data, unsigned int size) +{ + g_localicnt_interface->Push(input, output, data, size); +} + +static void* LocalInterconnect_pop(unsigned output) +{ + return g_localicnt_interface->Pop(output); +} + +static void LocalInterconnect_transfer() +{ + g_localicnt_interface->Advance(); +} + +static bool LocalInterconnect_busy() +{ + return g_localicnt_interface->Busy(); +} + +static void LocalInterconnect_display_stats() +{ + g_localicnt_interface->DisplayStats(); +} + +static void LocalInterconnect_display_overall_stats() +{ + g_localicnt_interface->DisplayOverallStats(); +} + +static void LocalInterconnect_display_state(FILE *fp) +{ + g_localicnt_interface->DisplayState(fp); +} + +static unsigned LocalInterconnect_get_flit_size() +{ + return g_localicnt_interface->GetFlitSize(); +} + + +/////////////////////////// + void icnt_reg_options( class OptionParser * opp ) { option_parser_register(opp, "-network_mode", OPT_INT32, &g_network_mode, "Interconnection network mode", "1"); option_parser_register(opp, "-inter_config_file", OPT_CSTR, &g_network_config_filename, "Interconnection network config file", "mesh"); + + + //parameters for local xbar + 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"); + } void icnt_wrapper_init() @@ -129,6 +203,20 @@ void icnt_wrapper_init() icnt_display_state = intersim2_display_state; icnt_get_flit_size = intersim2_get_flit_size; break; + case LOCAL_XBAR: + g_localicnt_interface = LocalInterconnect::New(g_inct_config); + icnt_create = LocalInterconnect_create; + icnt_init = LocalInterconnect_init; + icnt_has_buffer = LocalInterconnect_has_buffer; + icnt_push = LocalInterconnect_push; + icnt_pop = LocalInterconnect_pop; + icnt_transfer = LocalInterconnect_transfer; + icnt_busy = LocalInterconnect_busy; + icnt_display_stats = LocalInterconnect_display_stats; + icnt_display_overall_stats = LocalInterconnect_display_overall_stats; + icnt_display_state = LocalInterconnect_display_state; + icnt_get_flit_size = LocalInterconnect_get_flit_size; + break; default: assert(0); break; diff --git a/src/gpgpu-sim/icnt_wrapper.h b/src/gpgpu-sim/icnt_wrapper.h index a4d123e..e1086f9 100644 --- a/src/gpgpu-sim/icnt_wrapper.h +++ b/src/gpgpu-sim/icnt_wrapper.h @@ -57,13 +57,16 @@ extern icnt_display_stats_p icnt_display_stats; extern icnt_display_overall_stats_p icnt_display_overall_stats; extern icnt_display_state_p icnt_display_state; extern icnt_get_flit_size_p icnt_get_flit_size; -extern int g_network_mode; +extern unsigned g_network_mode; enum network_mode { INTERSIM = 1, + LOCAL_XBAR = 2, N_NETWORK_MODE }; + + void icnt_wrapper_init(); void icnt_reg_options( class OptionParser * opp ); diff --git a/src/gpgpu-sim/local_interconnect.cc b/src/gpgpu-sim/local_interconnect.cc new file mode 100644 index 0000000..1d93fe6 --- /dev/null +++ b/src/gpgpu-sim/local_interconnect.cc @@ -0,0 +1,277 @@ +// Copyright (c) 2009-2013, Tor M. Aamodt, Dongdong Li, Ali Bakhoda +// The University of British Columbia +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// +// Redistributions of source code must retain the above copyright notice, this +// list of conditions and the following disclaimer. +// Redistributions in binary form must reproduce the above copyright notice, this +// list of conditions and the following disclaimer in the documentation and/or +// other materials provided with the distribution. +// Neither the name of The University of British Columbia nor the names of its +// contributors may be used to endorse or promote products derived from this +// software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND +// ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED +// WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +#include +#include +#include +#include +#include +#include +#include + +#include "local_interconnect.h" +#include "mem_fetch.h" + +xbar_router::xbar_router(unsigned router_id, unsigned n_shader, unsigned n_mem, unsigned m_in_buffer_limit, unsigned m_out_buffer_limit) +{ + m_id=router_id; + _n_mem = n_mem; + _n_shader = n_shader; + total_nodes = n_shader+n_mem; + in_buffers.resize(total_nodes); + out_buffers.resize(total_nodes); + next_node=0; + in_buffer_limit = m_in_buffer_limit; + out_buffer_limit = m_out_buffer_limit; + + cycles = 0; + conflicts= 0; + out_buffer_full=0; + packets_num=0; +} + + +xbar_router::~xbar_router() +{ + +} + +void xbar_router::Push(unsigned input_deviceID, unsigned output_deviceID, void* data, unsigned int size) +{ + assert(input_deviceID < total_nodes); + in_buffers[input_deviceID].push(Packet(data, output_deviceID)); + packets_num++; +} + +void* xbar_router::Pop(unsigned ouput_deviceID) +{ + assert(ouput_deviceID < total_nodes); + void* data = NULL; + + if(!out_buffers[ouput_deviceID].empty()) { + data = out_buffers[ouput_deviceID].front().data; + out_buffers[ouput_deviceID].pop(); + } + + return data; +} + + +bool xbar_router::Has_Buffer_In(unsigned input_deviceID, unsigned size) const{ + + assert(input_deviceID < total_nodes); + return (in_buffers[input_deviceID].size() + size <= in_buffer_limit); + +} + +bool xbar_router::Has_Buffer_Out(unsigned output_deviceID, unsigned size) const{ + return (out_buffers[output_deviceID].size() + size <= out_buffer_limit); +} + +void xbar_router::Advance() { + cycles++; + + vector issued(total_nodes, false); + + for(unsigned i=0; iHas_Buffer_In(input_deviceID, 1)); + + net[subnet]->Push(input_deviceID, output_deviceID, data, size); + +} + +void* LocalInterconnect::Pop(unsigned ouput_deviceID){ + + // 0-_n_shader-1 indicates reply(network 1), otherwise request(network 0) + int subnet = 0; + if (ouput_deviceID < n_shader) + subnet = 1; + + return net[subnet]->Pop(ouput_deviceID); + +} + +void LocalInterconnect::Advance(){ + + for (unsigned i = 0; i < n_subnets; ++i) { + net[i]->Advance(); + } + +} + +bool LocalInterconnect::Busy() const{ + + for (unsigned i = 0; i < n_subnets; ++i) { + if(net[i]->Busy()) + return true; + } + return false; +} + +bool LocalInterconnect::HasBuffer(unsigned deviceID, unsigned int size) const{ + + bool has_buffer = false; + + has_buffer = net[0]->Has_Buffer_In(deviceID, size); + + if ((n_subnets>1) && deviceID >= n_shader) // deviceID is memory node + has_buffer = net[1]->Has_Buffer_In(deviceID, size); + + return has_buffer; + +} + +void LocalInterconnect::DisplayStats() const{ + + cout<<"Req_Network_injected_packets_num = "<packets_num< +#include +#include +#include +using namespace std; + + +struct inct_config +{ + + //config for local interconnect + unsigned in_buffer_limit; + unsigned out_buffer_limit; + unsigned subnets; +}; + +class xbar_router { + +public: + xbar_router(unsigned router_id, unsigned n_shader, unsigned n_mem, unsigned m_in_buffer_limit, unsigned m_out_buffer_limit); + ~xbar_router(); + void Push(unsigned input_deviceID, unsigned output_deviceID, void* data, unsigned int size); + void* Pop(unsigned ouput_deviceID); + void Advance(); + bool Busy() const; + bool Has_Buffer_In(unsigned input_deviceID, unsigned size) const; + bool Has_Buffer_Out(unsigned output_deviceID, unsigned size) const; + + //some stats + unsigned long long cycles; + unsigned long long conflicts; + unsigned long long out_buffer_full; + unsigned long long packets_num; + +private: + struct Packet{ + Packet(void* m_data, unsigned m_output_deviceID) { + data = m_data; + output_deviceID = m_output_deviceID; + } + void* data; + unsigned output_deviceID; + }; + vector > in_buffers; + vector > out_buffers; + unsigned _n_shader, _n_mem, total_nodes; + unsigned in_buffer_limit, out_buffer_limit; + unsigned next_node; + unsigned m_id; + +}; + +class LocalInterconnect { +public: + LocalInterconnect(const struct inct_config& m_localinct_config); + ~LocalInterconnect(); + static LocalInterconnect* New(const struct inct_config& m_inct_config); + void CreateInterconnect(unsigned n_shader, unsigned n_mem); + + //node side functions + void Init(); + void Push(unsigned input_deviceID, unsigned output_deviceID, void* data, unsigned int size); + void* Pop(unsigned ouput_deviceID); + void Advance(); + bool Busy() const; + bool HasBuffer(unsigned deviceID, unsigned int size) const; + void DisplayStats() const; + void DisplayOverallStats() const; + unsigned GetFlitSize() const; + + void DisplayState(FILE* fp) const; + + +protected: + + const inct_config& m_inct_config; + + unsigned n_shader, n_mem; + unsigned n_subnets; + vector net; + +}; + +#endif + + diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 3db988b..25a765a 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -1893,6 +1893,17 @@ void tensor_core::issue( register_set& source_reg ) pipelined_simd_unit::issue(source_reg); } +unsigned pipelined_simd_unit::get_active_lanes_in_pipeline(){ + active_mask_t active_lanes; + active_lanes.reset(); + if(!m_config->fast_execution_mode || active_insts_in_pipeline){ + for( unsigned stage=0; (stage+1)empty() ) + active_lanes|=m_pipeline_reg[stage]->get_active_mask(); + } + } + return active_lanes.count(); +} void ldst_unit::active_lanes_in_pipeline(){ unsigned active_count=pipelined_simd_unit::get_active_lanes_in_pipeline(); @@ -1946,13 +1957,13 @@ sp_unit::sp_unit( register_set* result_port, const shader_core_config *config,sh } dp_unit::dp_unit( register_set* result_port, const shader_core_config *config,shader_core_ctx *core) - : pipelined_simd_unit(result_port,config,config->max_sfu_latency,core) + : pipelined_simd_unit(result_port,config,config->max_dp_latency,core) { m_name = "DP "; } int_unit::int_unit( register_set* result_port, const shader_core_config *config,shader_core_ctx *core) - : pipelined_simd_unit(result_port,config,config->max_sp_latency,core) + : pipelined_simd_unit(result_port,config,config->max_int_latency,core) { m_name = "INT "; } @@ -1993,19 +2004,25 @@ pipelined_simd_unit::pipelined_simd_unit( register_set* result_port, const shade for( unsigned i=0; i < m_pipeline_depth; i++ ) m_pipeline_reg[i] = new warp_inst_t( config ); m_core=core; + active_insts_in_pipeline=0; } void pipelined_simd_unit::cycle() { if( !m_pipeline_reg[0]->empty() ){ m_result_port->move_in(m_pipeline_reg[0]); + assert(active_insts_in_pipeline > 0); + active_insts_in_pipeline--; + } + if(!m_config->fast_execution_mode || active_insts_in_pipeline){ + for( unsigned stage=0; (stage+1)empty() ) { if( !m_dispatch_reg->dispatch_delay()){ int start_stage = m_dispatch_reg->latency - m_dispatch_reg->initiation_interval; move_warp(m_pipeline_reg[start_stage],m_dispatch_reg); + active_insts_in_pipeline++; } } occupied >>=1; @@ -2981,8 +2998,61 @@ unsigned int shader_core_config::max_cta( const kernel_info_t &k ) const return result; } +void shader_core_config::set_pipeline_latency() { + + if(fast_execution_mode) { + //calculate the max latency based on the input + + unsigned int_latency[5]; + unsigned fp_latency[5]; + unsigned dp_latency[5]; + unsigned sfu_latency; + unsigned tensor_latency; + + /* + * [0] ADD,SUB + * [1] MAX,Min + * [2] MUL + * [3] MAD + * [4] DIV + */ + sscanf(opcode_latency_int, "%u,%u,%u,%u,%u", + &int_latency[0],&int_latency[1],&int_latency[2], + &int_latency[3],&int_latency[4]); + sscanf(opcode_latency_fp, "%u,%u,%u,%u,%u", + &fp_latency[0],&fp_latency[1],&fp_latency[2], + &fp_latency[3],&fp_latency[4]); + sscanf(opcode_latency_dp, "%u,%u,%u,%u,%u", + &dp_latency[0],&dp_latency[1],&dp_latency[2], + &dp_latency[3],&dp_latency[4]); + sscanf(opcode_latency_sfu, "%u", + &sfu_latency); + sscanf(opcode_latency_tensor, "%u", + &tensor_latency); + + //all div operation are executed on sfu + //assume that the max latency are dp div or normal sfu_latency + max_sfu_latency = std::max(dp_latency[4],sfu_latency); + //assume that the max operation has the max latency + max_sp_latency = fp_latency[1]; + max_int_latency = int_latency[1]; + max_dp_latency = dp_latency[1]; + max_tensor_core_latency = tensor_latency; + } else { + max_sfu_latency = 512; + max_sp_latency = 32; + max_int_latency = 32; + max_dp_latency = 512; + max_tensor_core_latency = 64; + } + +} + void shader_core_ctx::cycle() { + if(m_config->fast_execution_mode && !isactive() && get_not_completed() == 0) + return; + m_stats->shader_cycles[m_sid]++; writeback(); execute(); diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index 9abd223..2204697 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -55,7 +55,6 @@ #include "traffic_breakdown.h" - #define NO_OP_FLAG 0xFF /* READ_PACKET_SIZE: @@ -1073,16 +1072,8 @@ public: //modifiers virtual void cycle(); virtual void issue( register_set& source_reg ); - virtual unsigned get_active_lanes_in_pipeline() - { - active_mask_t active_lanes; - active_lanes.reset(); - for( unsigned stage=0; (stage+1)empty() ) - active_lanes|=m_pipeline_reg[stage]->get_active_mask(); - } - return active_lanes.count(); - } + virtual unsigned get_active_lanes_in_pipeline(); + virtual void active_lanes_in_pipeline() = 0; /* virtual void issue( register_set& source_reg ) @@ -1113,6 +1104,9 @@ protected: warp_inst_t **m_pipeline_reg; register_set *m_result_port; class shader_core_ctx *m_core; + + unsigned active_insts_in_pipeline; + }; class sfu : public pipelined_simd_unit @@ -1413,10 +1407,8 @@ struct shader_core_config : public core_config } max_warps_per_shader = n_thread_per_shader/warp_size; assert( !(n_thread_per_shader % warp_size) ); - max_sfu_latency = 512; - max_sp_latency = 32; - - max_tensor_core_latency = 64; + + set_pipeline_latency(); m_L1I_config.init(m_L1I_config.m_config_string,FuncCachePreferNone); m_L1T_config.init(m_L1T_config.m_config_string,FuncCachePreferNone); @@ -1432,6 +1424,7 @@ struct shader_core_config : public core_config unsigned sid_to_cluster( unsigned sid ) const { return sid / n_simt_cores_per_cluster; } unsigned sid_to_cid( unsigned sid ) const { return sid % n_simt_cores_per_cluster; } unsigned cid_to_sid( unsigned cid, unsigned cluster_id ) const { return cluster_id*n_simt_cores_per_cluster + cid; } + void set_pipeline_latency(); // data char *gpgpu_shader_core_pipeline_opt; @@ -1506,7 +1499,9 @@ struct shader_core_config : public core_config bool sub_core_model; unsigned max_sp_latency; + unsigned max_int_latency; unsigned max_sfu_latency; + unsigned max_dp_latency; unsigned max_tensor_core_latency; unsigned n_simt_cores_per_cluster; @@ -1524,6 +1519,7 @@ struct shader_core_config : public core_config bool gpgpu_concurrent_kernel_sm; bool adpative_volta_cache_config; + bool fast_execution_mode; }; struct shader_core_stats_pod { diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 270bace..9e2bfa2 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -222,10 +222,12 @@ gpgpu_sim *gpgpu_ptx_sim_init_perf() read_parser_environment_variables(); option_parser_t opp = option_parser_create(); - icnt_reg_options(opp); - g_the_gpu_config.reg_options(opp); // register GPU microrachitecture options ptx_reg_options(opp); ptx_opcocde_latency_options(opp); + + icnt_reg_options(opp); + g_the_gpu_config.reg_options(opp); // register GPU microrachitecture options + option_parser_cmdline(opp, sg_argc, sg_argv); // parse configuration options fprintf(stdout, "GPGPU-Sim: Configuration options:\n\n"); option_parser_print(opp, stdout); -- cgit v1.3 From 69af057daf1999f17b81d761c7a616f67f56adbb Mon Sep 17 00:00:00 2001 From: Suchita Pati Date: Tue, 16 Apr 2019 23:36:02 -0500 Subject: Adding additional fixes for the texture bugs --- libcuda/cuda_runtime_api.cc | 22 ++++++++++++---------- src/abstract_hardware_model.cc | 13 ++++++++++--- src/abstract_hardware_model.h | 36 +++++++++++++++++++++--------------- src/cuda-sim/cuda-sim.cc | 7 +++++++ src/cuda-sim/instructions.cc | 8 +++++++- 5 files changed, 57 insertions(+), 29 deletions(-) (limited to 'src/cuda-sim') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 7569b20..f6e5e6f 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1121,10 +1121,10 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic *value= 0; break; case 75: - *value= 9 ; + *value= 7 ; break; case 76: - *value= 3 ; + *value= 0 ; break; case 78: *value= 0 ; //TODO: as of now, we dont support stream priorities. @@ -1199,22 +1199,23 @@ __host__ cudaError_t CUDARTAPI cudaGetDevice(int *device) *device = g_active_device; return g_last_cudaError = cudaSuccess; } + __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit limit ) { if(g_debug_execution >= 3){ announce_call(__my_func__); } switch(limit) { - case 0: + case 0: // cudaLimitStackSize *pValue=1024; break; - case 2: + case 2: // cudaLimitMallocHeapSize *pValue=8388608; break; - case 3: + case 3: // cudaLimitDevRuntimeSyncDepth *pValue=2; break; - case 4: + case 4: // cudaLimitDevRuntimePendingLaunchCount *pValue=2048; break; default: @@ -1225,7 +1226,6 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit li } - __host__ cudaError_t CUDARTAPI cudaStreamGetPriority ( cudaStream_t hStream, int* priority ) { if(g_debug_execution >= 3){ @@ -1585,7 +1585,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) announce_call(__my_func__); } #if (CUDART_VERSION >= 3000) - //synchronization required for application using external libraries without explicit synchronization in the code to + //per-stream synchronization required for application using external libraries without explicit synchronization in the code to //avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. stream->synchronize(); g_stream_manager->destroy_stream(stream); @@ -2714,7 +2714,6 @@ cudaError_t cudaDeviceReset ( void ) { return g_last_cudaError = cudaSuccess; } cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ - // I don't know what this should do if(g_debug_execution >= 3){ announce_call(__my_func__); } @@ -2722,7 +2721,6 @@ cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ return g_last_cudaError = cudaSuccess; } - void CUDARTAPI __cudaRegisterFunction( void **fatCubinHandle, const char *hostFun, @@ -3333,6 +3331,10 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, } function_info *entry = context->get_kernel(hostFun); gpgpu_t* gpu= context->get_device()->get_gpgpu(); + /* + Passing a snapshot of the GPU's current texture mapping to the kernel's info + as kernels should use texture bindings present at the time of their launch. + */ kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping()); if( entry == NULL ) { printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found for %p\n", hostFun); diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 3db87c9..820f5fd 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -192,6 +192,9 @@ gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config ) checkpoint_CTA_t = m_function_model_config.get_checkpoint_CTA_t(); checkpoint_insn_Y = m_function_model_config.get_checkpoint_insn_Y(); + // initialize texture mappings to empty + m_NameToTextureInfo.clear(); + m_NameToCudaArray.clear(); if(m_function_model_config.get_ptx_inst_debug_to_file() != 0) ptx_inst_debug_file = fopen(m_function_model_config.get_ptx_inst_debug_file(), "w"); @@ -688,7 +691,11 @@ unsigned g_kernel_launch_latency; unsigned kernel_info_t::m_next_uid = 1; -kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map NameToCudaArray, std::map NameToTexureInfo) +/*A snapshot of the texture mappings needs to be stored in the kernel's info as +kernels should use the texture bindings seen at the time of launch and textures + can be bound/unbound asynchronously with respect to streams. */ + +kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTexureInfo) { m_kernel_entry=entry; m_grid_dim=gridDim; @@ -708,8 +715,8 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * m_launch_latency = g_kernel_launch_latency; volta_cache_config_set=false; - m_NameToCudaArray = NameToCudaArray; - m_NameToTexureInfo = NameToTexureInfo; + m_NameToCudaArray = nameToCudaArray; + m_NameToTexureInfo = nameToTexureInfo; } kernel_info_t::~kernel_info_t() diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index bbc31e5..201b21f 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -212,7 +212,7 @@ public: // m_num_cores_running=0; // m_param_mem=NULL; // } - kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map m_NameToCudaArray, std::map m_NameToTexureInfo); + kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTexureInfo); ~kernel_info_t(); void inc_running() { m_num_cores_running++; } @@ -275,19 +275,22 @@ public: std::list &active_threads() { return m_active_threads; } class memory_space *get_param_memory() { return m_param_mem; } + + //The following functions access texture bindings present at the kernel's launch + const struct cudaArray* get_texarray( const std::string &texname ) const - { - std::map::const_iterator t=m_NameToCudaArray.find(texname); - assert(t != m_NameToCudaArray.end()); - return t->second; - } + { + std::map::const_iterator t=m_NameToCudaArray.find(texname); + assert(t != m_NameToCudaArray.end()); + return t->second; + } - const struct textureInfo* get_texinfo( const std::string &texname ) const - { - std::map::const_iterator t=m_NameToTexureInfo.find(texname); - assert(t != m_NameToTexureInfo.end()); - return t->second; - } + const struct textureInfo* get_texinfo( const std::string &texname ) const + { + std::map::const_iterator t=m_NameToTexureInfo.find(texname); + assert(t != m_NameToTexureInfo.end()); + return t->second; + } private: kernel_info_t( const kernel_info_t & ); // disable copy constructor @@ -297,7 +300,8 @@ private: unsigned m_uid; static unsigned m_next_uid; - + + //These maps contain the snapshot of the texture mappings at kernel launch std::map m_NameToCudaArray; std::map m_NameToTexureInfo; @@ -611,7 +615,9 @@ public: const gpgpu_functional_sim_config &get_config() const { return m_function_model_config; } FILE* get_ptx_inst_debug_file() { return ptx_inst_debug_file; } - + + // These maps return the current texture mappings for the GPU at any given time. + std::map getNameArrayMapping() {return m_NameToCudaArray;} std::map getNameInfoMapping() {return m_NameToTexureInfo;} @@ -624,7 +630,7 @@ protected: class memory_space *m_surf_mem; unsigned long long m_dev_malloc; - + // These maps contain the current texture mappings for the GPU at any given time. std::map > m_NameToTextureRef; std::map m_TextureRefToName; std::map m_NameToCudaArray; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index a51ea3e..9d59411 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1507,6 +1507,13 @@ static unsigned get_tex_datasize( const ptx_instruction *pI, ptx_thread_info *th std::string texname = src1.name(); gpgpu_t *gpu = thread->get_gpu(); + /* + For programs with many streams, textures can be bound and unbound + asynchronously. This means we need to use the kernel's "snapshot" of + the state of the texture mappings when it was launched (so that we + don't try to access the incorrect texture mapping if it's been updated, + or that we don't access a mapping that has been unbound). + */ kernel_info_t& k = thread->get_kernel(); const struct textureInfo* texInfo = k.get_texinfo(texname); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index d29fc8b..11001d7 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -5258,7 +5258,13 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread ) if (!ptx_tex_regs) ptx_tex_regs = new ptx_reg_t[4]; unsigned nelem = src2.get_vect_nelem(); thread->get_vector_operand_values(src2, ptx_tex_regs, nelem); //ptx_reg should be 4 entry vector type...coordinates into texture - + /* + For programs with many streams, textures can be bound and unbound + asynchronously. This means we need to use the kernel's "snapshot" of + the state of the texture mappings when it was launched (so that we + don't try to access the incorrect texture mapping if it's been updated, + or that we don't access a mapping that has been unbound). + */ gpgpu_t *gpu = thread->get_gpu(); kernel_info_t &k = thread->get_kernel(); const struct textureReference* texref = gpu->get_texref(texname); -- cgit v1.3 From 49c6eeff7752b5816fb08c0d6baf2b47c789bc18 Mon Sep 17 00:00:00 2001 From: Suchita Pati Date: Wed, 17 Apr 2019 22:09:17 -0500 Subject: Adding configs for device limits and fixed other misc bugs --- libcuda/cuda_runtime_api.cc | 11 +++++++---- src/abstract_hardware_model.cc | 9 ++++----- src/abstract_hardware_model.h | 17 ++++++++--------- src/cuda-sim/cuda-sim.cc | 4 ++-- src/gpgpu-sim/gpu-sim.cc | 8 ++++++++ src/gpgpu-sim/gpu-sim.h | 11 ++++++++++- 6 files changed, 39 insertions(+), 21 deletions(-) (limited to 'src/cuda-sim') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index f6e5e6f..86435b9 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1205,18 +1205,20 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit li if(g_debug_execution >= 3){ announce_call(__my_func__); } + _cuda_device_id *dev = GPGPUSim_Init(); + const gpgpu_sim_config& config=dev->get_gpgpu()->get_config(); switch(limit) { case 0: // cudaLimitStackSize - *pValue=1024; + *pValue=config.stack_limit(); break; case 2: // cudaLimitMallocHeapSize - *pValue=8388608; + *pValue=config.heap_limit(); break; case 3: // cudaLimitDevRuntimeSyncDepth - *pValue=2; + *pValue=config.sync_depth_limit(); break; case 4: // cudaLimitDevRuntimePendingLaunchCount - *pValue=2048; + *pValue=config.pending_launch_count_limit(); break; default: printf("ERROR:Limit %s unimplemented \n",limit); @@ -2717,6 +2719,7 @@ cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ if(g_debug_execution >= 3){ announce_call(__my_func__); } + //Blocks until the device has completed all preceding requested tasks synchronize(); return g_last_cudaError = cudaSuccess; } diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 820f5fd..cebdb25 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -88,8 +88,6 @@ void checkpoint::load_global_mem(class memory_space *temp_mem, char * f1name) fclose ( fp2 ); } - - void checkpoint::store_global_mem(class memory_space * mem, char *fname, char * format) { @@ -195,6 +193,8 @@ gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config ) // initialize texture mappings to empty m_NameToTextureInfo.clear(); m_NameToCudaArray.clear(); + m_TextureRefToName.clear(); + m_NameToAttribute.clear(); if(m_function_model_config.get_ptx_inst_debug_to_file() != 0) ptx_inst_debug_file = fopen(m_function_model_config.get_ptx_inst_debug_file(), "w"); @@ -694,8 +694,7 @@ unsigned kernel_info_t::m_next_uid = 1; /*A snapshot of the texture mappings needs to be stored in the kernel's info as kernels should use the texture bindings seen at the time of launch and textures can be bound/unbound asynchronously with respect to streams. */ - -kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTexureInfo) +kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTextureInfo) { m_kernel_entry=entry; m_grid_dim=gridDim; @@ -716,7 +715,7 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * volta_cache_config_set=false; m_NameToCudaArray = nameToCudaArray; - m_NameToTexureInfo = nameToTexureInfo; + m_NameToTextureInfo = nameToTextureInfo; } kernel_info_t::~kernel_info_t() diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 201b21f..e8716ab 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -212,7 +212,7 @@ public: // m_num_cores_running=0; // m_param_mem=NULL; // } - kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTexureInfo); + kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTextureInfo); ~kernel_info_t(); void inc_running() { m_num_cores_running++; } @@ -287,8 +287,8 @@ public: const struct textureInfo* get_texinfo( const std::string &texname ) const { - std::map::const_iterator t=m_NameToTexureInfo.find(texname); - assert(t != m_NameToTexureInfo.end()); + std::map::const_iterator t=m_NameToTextureInfo.find(texname); + assert(t != m_NameToTextureInfo.end()); return t->second; } @@ -303,7 +303,7 @@ private: //These maps contain the snapshot of the texture mappings at kernel launch std::map m_NameToCudaArray; - std::map m_NameToTexureInfo; + std::map m_NameToTextureInfo; dim3 m_grid_dim; dim3 m_block_dim; @@ -601,8 +601,8 @@ public: const struct textureInfo* get_texinfo( const std::string &texname ) const { - std::map::const_iterator t=m_NameToTexureInfo.find(texname); - assert(t != m_NameToTexureInfo.end()); + std::map::const_iterator t=m_NameToTextureInfo.find(texname); + assert(t != m_NameToTextureInfo.end()); return t->second; } @@ -617,9 +617,8 @@ public: FILE* get_ptx_inst_debug_file() { return ptx_inst_debug_file; } // These maps return the current texture mappings for the GPU at any given time. - std::map getNameArrayMapping() {return m_NameToCudaArray;} - std::map getNameInfoMapping() {return m_NameToTexureInfo;} + std::map getNameInfoMapping() {return m_NameToTextureInfo;} protected: const gpgpu_functional_sim_config &m_function_model_config; @@ -634,7 +633,7 @@ protected: std::map > m_NameToTextureRef; std::map m_TextureRefToName; std::map m_NameToCudaArray; - std::map m_NameToTexureInfo; + std::map m_NameToTextureInfo; std::map m_NameToAttribute; }; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 9d59411..f7bb9cc 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -218,7 +218,7 @@ void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* te texInfo->Ty_numbits = intLOGB2(Ty); texInfo->texel_size = texel_size; texInfo->texel_size_numbits = intLOGB2(texel_size); - m_NameToTexureInfo[texname] = texInfo; + m_NameToTextureInfo[texname] = texInfo; } void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) @@ -226,7 +226,7 @@ void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) //assumes bind-use-unbind-bind-use-unbind pattern std::string texname = gpgpu_ptx_sim_findNamefromTexture(texref); m_NameToCudaArray.erase(texname); - m_NameToTexureInfo.erase(texname); + m_NameToTextureInfo.erase(texname); } unsigned g_assemble_code_next_pc=0; diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index ec570bf..2a862a4 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -532,6 +532,14 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) option_parser_register(opp, "-visualizer_zlevel", OPT_INT32, &g_visualizer_zlevel, "Compression level of the visualizer output log (0=no comp, 9=highest)", "6"); + option_parser_register(opp, "-gpgpu_stack_size_limit", OPT_INT32, &stack_size_limit, + "GPU thread stack size", "1024" ); + option_parser_register(opp, "-gpgpu_heap_size_limit", OPT_INT32, &heap_size_limit, + "GPU malloc heap size ", "8388608" ); + option_parser_register(opp, "-gpgpu_runtime_sync_depth_limit", OPT_INT32, &runtime_sync_depth_limit, + "GPU device runtime synchronize depth", "2" ); + option_parser_register(opp, "-gpgpu_runtime_pending_launch_count_limit", OPT_INT32, &runtime_pending_launch_count_limit, + "GPU device runtime pending launch count", "2048" ); option_parser_register(opp, "-trace_enabled", OPT_BOOL, &Trace::enabled, "Turn on traces", "0"); diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 6ce5524..11ffe49 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -337,6 +337,11 @@ public: unsigned get_max_concurrent_kernel() const { return max_concurrent_kernel; } unsigned checkpoint_option; + size_t stack_limit() const {return stack_size_limit; } + size_t heap_limit() const {return heap_size_limit; } + size_t sync_depth_limit() const {return runtime_sync_depth_limit; } + size_t pending_launch_count_limit() const {return runtime_pending_launch_count_limit;} + private: void init_clock_domains(void ); @@ -377,7 +382,11 @@ private: int gpu_stat_sample_freq; int gpu_runtime_stat_flag; - + // Device Limits + size_t stack_size_limit; + size_t heap_size_limit; + size_t runtime_sync_depth_limit; + size_t runtime_pending_launch_count_limit; unsigned long long liveness_message_freq; -- cgit v1.3