From 06d6f3bb15f4c1dc56943304696ac83b36413907 Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Sun, 8 Aug 2010 23:57:43 -0800 Subject: refactor: shader.cc free of extern declarations [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7176] --- libcuda/cuda_runtime_api.cc | 2 - src/cuda-sim/cuda-sim.cc | 2 +- src/cuda-sim/cuda-sim.h | 55 ++++++++++++++- src/cuda-sim/ptx_sim.cc | 6 +- src/gpgpu-sim/delayqueue.cc | 3 +- src/gpgpu-sim/dram.cc | 8 +-- src/gpgpu-sim/dram.h | 2 +- src/gpgpu-sim/dram_sched.cc | 1 - src/gpgpu-sim/dwf.cc | 1 - src/gpgpu-sim/gpu-sim.cc | 7 +- src/gpgpu-sim/gpu-sim.h | 39 +++++++++++ src/gpgpu-sim/l2cache.cc | 2 - src/gpgpu-sim/mem_latency_stat.cc | 1 - src/gpgpu-sim/mem_latency_stat.h | 1 - src/gpgpu-sim/shader.cc | 117 +++++--------------------------- src/gpgpu-sim/visualizer.cc | 8 +-- src/gpgpu-sim/visualizer.h | 77 +++++++++++++++++++++ src/gpgpu-sim/warp_tracker.cc | 2 - src/gpgpusim_entrypoint.cc | 1 - src/intersim/interconnect_interface.cpp | 2 - 20 files changed, 194 insertions(+), 143 deletions(-) create mode 100644 src/gpgpu-sim/visualizer.h diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index b6e99ca..55432b4 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -888,8 +888,6 @@ __host__ cudaError_t CUDARTAPI cudaStreamQuery(cudaStream_t stream) * * *******************************************************************************/ -extern signed long long gpu_tot_sim_cycle; - struct timer_event { int m_uid; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 737b94a..a55fb79 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -79,6 +79,7 @@ #include "ptx-stats.h" #include "ptx_loader.h" #include "ptx_parser.h" +#include "../gpgpu-sim/gpu-sim.h" extern bool g_interactive_debugger_enabled; @@ -759,7 +760,6 @@ unsigned datatype2size( unsigned data_type ) return data_size; } -extern unsigned long long gpu_sim_cycle; unsigned g_warp_active_mask; void function_info::ptx_exec_inst( ptx_thread_info *thread, diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 2c5e0a7..7e05f7e 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -2,6 +2,7 @@ #define CUDASIM_H_INCLUDED #include "../abstract_hardware_model.h" +#include "dram_callback.h" #include #include #include @@ -17,10 +18,16 @@ extern int g_debug_execution; extern int g_debug_thread_uid; extern std::map *g_kernel_name_to_function_lookup; -extern void gpgpu_cuda_ptx_sim_init_grid(const char *kernel_key,struct gpgpu_ptx_sim_arg *args, struct dim3 gridDim, struct dim3 blockDim ); -extern void gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,struct gpgpu_ptx_sim_arg *args, struct dim3 gridDim, struct dim3 blockDim ); +extern void gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, + struct gpgpu_ptx_sim_arg *args, + struct dim3 gridDim, + struct dim3 blockDim ); +extern void gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, + struct gpgpu_ptx_sim_arg *args, + struct dim3 gridDim, + struct dim3 blockDim ); extern void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, struct gpgpu_ptx_sim_arg *); -extern void print_splash(); +extern void print_splash(); extern void* gpgpu_ptx_sim_malloc( size_t count ); extern void* gpgpu_ptx_sim_mallocarray( size_t count ); extern void gpgpu_ptx_sim_memcpy_to_gpu( size_t dst_start_addr, const void *src, size_t count ); @@ -43,5 +50,47 @@ extern void read_sim_environment_variables(); extern void register_function_implementation( const char *name, function_info *impl ); extern void ptxinfo_cuda_addinfo(); extern void ptxinfo_opencl_addinfo( std::map &kernels ); +extern void ptx_sim_free_sm( class ptx_thread_info** thread_info ); +unsigned ptx_sim_init_thread( class ptx_thread_info** thread_info, + int sid, + unsigned tid, + unsigned threads_left, + unsigned num_threads, + class core_t *core, + unsigned hw_cta_id, + unsigned hw_warp_id ); +unsigned ptx_sim_cta_size(); +const struct gpgpu_ptx_sim_kernel_info* ptx_sim_kernel_info(); +void set_option_gpgpu_spread_blocks_across_cores(int option); +int ptx_thread_done( void *thd ); +unsigned ptx_thread_donecycle( void *thr ); +int ptx_thread_get_next_pc( void *thd ); +void* ptx_thread_get_next_finfo( void *thd ); +int ptx_thread_at_barrier( void *thd ); +int ptx_thread_all_at_barrier( void *thd ); +unsigned long long ptx_thread_get_cta_uid( void *thd ); +void ptx_thread_reset_barrier( void *thd ); +void ptx_thread_release_barrier( void *thd ); +void ptx_print_insn( address_type pc, FILE *fp ); +unsigned int ptx_set_tex_cache_linesize( unsigned linesize); +void ptx_decode_inst( void *thd, + unsigned *op, + int *i1, + int *i2, + int *i3, + int *i4, + int *o1, + int *o2, + int *o3, + int *o4, + int *vectorin, + int *vectorout, + int *arch_reg ); +void ptx_exec_inst( void *thd, + address_type *addr, + memory_space_t *space, + unsigned *data_size, + dram_callback_t* callback, + unsigned warp_active_mask ); #endif diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc index 31c8c81..2156bd6 100644 --- a/src/cuda-sim/ptx_sim.cc +++ b/src/cuda-sim/ptx_sim.cc @@ -64,6 +64,7 @@ #include "ptx_sim.h" #include #include "ptx_ir.h" +#include "../gpgpu-sim/gpu-sim.h" void feature_not_implemented( const char *f ); @@ -249,8 +250,6 @@ const ptx_version &ptx_thread_info::get_ptx_version() const return m_func_info->get_ptx_version(); } -extern unsigned long long gpu_sim_cycle; - void ptx_thread_info::set_done() { assert( !m_at_barrier ); @@ -258,9 +257,6 @@ void ptx_thread_info::set_done() m_cycle_done = gpu_sim_cycle; } -extern unsigned long long gpu_sim_cycle; -extern signed long long gpu_tot_sim_cycle; - unsigned ptx_thread_info::get_builtin( int builtin_id, unsigned dim_mod ) { assert( m_valid ); diff --git a/src/gpgpu-sim/delayqueue.cc b/src/gpgpu-sim/delayqueue.cc index dc870d0..659018b 100644 --- a/src/gpgpu-sim/delayqueue.cc +++ b/src/gpgpu-sim/delayqueue.cc @@ -65,10 +65,9 @@ #include "delayqueue.h" #include "gpu-misc.h" +#include "gpu-sim.h" #include "../intersim/statwraper.h" -extern unsigned long long gpu_sim_cycle; //for stat collection - unsigned char dq_full( delay_queue* dq ) { if (dq->max_len && dq->length >= dq->max_len) diff --git a/src/gpgpu-sim/dram.cc b/src/gpgpu-sim/dram.cc index 9f76b32..a4e69f8 100644 --- a/src/gpgpu-sim/dram.cc +++ b/src/gpgpu-sim/dram.cc @@ -64,18 +64,14 @@ * Vancouver, BC V6T 1Z4 */ -//#include "gpu-sim.h" +#include "gpu-sim.h" #include "gpu-misc.h" //#include "shader.h" #include "dram.h" +#include "mem_latency_stat.h" -extern unsigned long long gpu_sim_cycle; -extern signed long long gpu_tot_sim_cycle; -extern int gpgpu_memlatency_stat; extern unsigned max_dq_latency; -extern unsigned dq_lat_table[24]; extern unsigned max_dq_latency; -extern unsigned dq_lat_table[24]; unsigned int gpu_n_warps; extern unsigned int gpu_n_mem_per_ctrlr; extern unsigned int recent_dram_util; diff --git a/src/gpgpu-sim/dram.h b/src/gpgpu-sim/dram.h index fff23f6..cd92fa9 100644 --- a/src/gpgpu-sim/dram.h +++ b/src/gpgpu-sim/dram.h @@ -126,7 +126,7 @@ typedef struct { unsigned int n_idle; } bank_t; -typedef struct { +typedef struct dram_timing { unsigned int id; unsigned int tCCD; //column to column delay diff --git a/src/gpgpu-sim/dram_sched.cc b/src/gpgpu-sim/dram_sched.cc index 7e8b38c..30c8752 100644 --- a/src/gpgpu-sim/dram_sched.cc +++ b/src/gpgpu-sim/dram_sched.cc @@ -69,7 +69,6 @@ #include "gpu-sim.h" #include "../abstract_hardware_model.h" -extern unsigned long long gpu_sim_cycle; extern unsigned max_mrq_latency; extern unsigned mrq_lat_table[24]; extern int gpgpu_memlatency_stat; diff --git a/src/gpgpu-sim/dwf.cc b/src/gpgpu-sim/dwf.cc index cd66084..0cad2b6 100644 --- a/src/gpgpu-sim/dwf.cc +++ b/src/gpgpu-sim/dwf.cc @@ -2323,7 +2323,6 @@ void g_print_max_heap(int sid) { #include "stat-tool.cc" unsigned gpgpu_thread_swizzling = 0; -unsigned long long gpu_sim_cycle = 0; int regfile_hash(signed istream_number, unsigned simd_size, unsigned n_banks) { if (gpgpu_thread_swizzling) { diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 9a16320..0cef489 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -94,6 +94,7 @@ #include "../abstract_hardware_model.h" #include "../debug.h" #include "../gpgpusim_entrypoint.h" +#include "../cuda-sim/cuda-sim.h" #include #include @@ -229,7 +230,6 @@ double icnt_time=0; double dram_time=0; double l2_time=0; -#define MhZ *1000000 double core_freq=2 MhZ; double icnt_freq=2 MhZ; double dram_freq=2 MhZ; @@ -327,14 +327,11 @@ unsigned char single_check_icnt_has_buffer(int chip, int sid, unsigned char is_w unsigned char fq_pop(int tpc_id); void fill_shd_L1_with_new_line(shader_core_ctx_t * sc, mem_fetch_t * mf); -void set_option_gpgpu_spread_blocks_across_cores(int option); void set_param_gpgpu_num_shaders(int num_shaders); unsigned ptx_sim_grid_size(); void icnt_init_grid(); void interconnect_stats(); void icnt_overal_stat(); -unsigned ptx_sim_cta_size(); -unsigned ptx_sim_init_thread( ptx_thread_info** thread_info, int sid, unsigned tid,unsigned threads_left,unsigned num_threads, core_t *core, unsigned hw_cta_id, unsigned hw_warp_id ); void gpu_sim_loop( int grid_num ); @@ -1690,8 +1687,6 @@ void dump_regs(unsigned sid, unsigned tid) ptx_dump_regs( s->thread[tid].ptx_thd_info ); } -int ptx_thread_done( void *thr ); - void shader_dump_istream_state(shader_core_ctx_t *shader, FILE *fout ) { fprintf( fout, "\n"); diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index d962111..58ff4de 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -126,12 +126,51 @@ extern unsigned int gpu_n_shader; extern unsigned int gpu_n_mem; extern bool gpgpu_reg_bankconflict; extern int gpgpu_dram_sched_queue_size; +extern unsigned long long gpu_sim_cycle; extern unsigned long long gpu_tot_sim_cycle; +extern unsigned long long gpu_sim_insn; extern unsigned int gpu_n_warp_per_shader; extern unsigned int **max_conc_access2samerow; extern unsigned int **max_servicetime2samerow; extern unsigned int **row_access; extern unsigned int **num_activates; +extern struct dram_timing **dram; +extern int *num_warps_issuable; +extern int *num_warps_issuable_pershader; +extern unsigned long long gpu_sim_insn_no_ld_const; +extern unsigned long long gpu_sim_insn_last_update; +extern unsigned long long gpu_completed_thread; +extern class shader_core_ctx **sc; +extern unsigned int gpgpu_pre_mem_stages; +extern unsigned int gpgpu_no_divg_load; +extern unsigned int gpgpu_thread_swizzling; +extern unsigned int gpgpu_strict_simd_wrbk; +extern unsigned int warp_conflict_at_writeback; +extern unsigned int gpgpu_commit_pc_beyond_two; +extern int gpgpu_spread_blocks_across_cores; +extern int gpgpu_cflog_interval; +extern unsigned int gpu_stall_by_MSHRwb; +extern unsigned int gpu_stall_shd_mem; +extern unsigned int gpu_stall_sh2icnt; +extern int gpgpu_operand_collector; +extern int gpu_runtime_stat_flag; +extern unsigned int *max_return_queue_length; +extern int gpgpu_partial_write_mask; +extern int gpgpu_n_mem_write_local; +extern int gpgpu_n_mem_write_global; +extern int gpgpu_cache_wt_through; +extern double core_freq; +extern double icnt_freq; +extern double dram_freq; +extern double l2_freq; +extern int pdom_sched_type; +extern int n_pdom_sc_orig_stat; +extern int n_pdom_sc_single_stat; +extern int gpgpu_cuda_sim; + +#ifndef MhZ + #define MhZ *1000000 +#endif extern void check_time_vector_update(unsigned int uid,int slot ,long int latency,int type); diff --git a/src/gpgpu-sim/l2cache.cc b/src/gpgpu-sim/l2cache.cc index de8de47..48dcda1 100644 --- a/src/gpgpu-sim/l2cache.cc +++ b/src/gpgpu-sim/l2cache.cc @@ -47,9 +47,7 @@ struct L2cacheblk // external dependencies extern unsigned long long int addrdec_mask[5]; -extern dram_t **dram; extern int gpgpu_dram_sched_queue_size; -extern unsigned long long gpu_sim_cycle; extern unsigned made_write_mfs; extern unsigned freed_L1write_mfs; extern unsigned freed_L2write_mfs; diff --git a/src/gpgpu-sim/mem_latency_stat.cc b/src/gpgpu-sim/mem_latency_stat.cc index 16b71b0..81589de 100644 --- a/src/gpgpu-sim/mem_latency_stat.cc +++ b/src/gpgpu-sim/mem_latency_stat.cc @@ -76,7 +76,6 @@ #include #include -extern unsigned long long gpu_sim_cycle; extern unsigned int gpu_n_mem; extern unsigned int gpu_n_shader; extern int gpgpu_dram_sched_queue_size; diff --git a/src/gpgpu-sim/mem_latency_stat.h b/src/gpgpu-sim/mem_latency_stat.h index 7a251ec..3b266ef 100644 --- a/src/gpgpu-sim/mem_latency_stat.h +++ b/src/gpgpu-sim/mem_latency_stat.h @@ -68,7 +68,6 @@ #ifndef MEM_LATENCY_STAT_H #define MEM_LATENCY_STAT_H -extern unsigned long long gpu_sim_cycle; extern int gpgpu_dram_scheduler; extern unsigned int gpu_mem_n_bk; extern bool gpgpu_memlatency_stat; diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 5c3d698..ef3a8c0 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -79,36 +79,13 @@ #include "../cuda-sim/cuda-sim.h" #include "gpu-sim.h" #include "mem_fetch.h" +#include "mem_latency_stat.h" +#include "visualizer.h" #include #define PRIORITIZE_MSHR_OVER_WB 1 #define MAX(a,b) (((a)>(b))?(a):(b)) -extern int gpgpu_memlatency_stat; -extern dram_t **dram; -extern int *num_warps_issuable; -extern int *num_warps_issuable_pershader; - -extern unsigned long long gpu_sim_insn; -extern unsigned long long gpu_sim_insn_no_ld_const; -extern unsigned long long gpu_sim_insn_last_update; -extern unsigned long long gpu_completed_thread; -extern unsigned long long gpu_sim_cycle; -extern shader_core_ctx_t **sc; -extern unsigned int gpgpu_pre_mem_stages; -extern int gpgpu_no_divg_load; -extern unsigned int gpgpu_thread_swizzling; -extern unsigned int gpgpu_strict_simd_wrbk; -extern unsigned int warp_conflict_at_writeback; -extern unsigned int gpgpu_commit_pc_beyond_two; -extern int gpgpu_spread_blocks_across_cores; -extern int gpgpu_cflog_interval; - -extern unsigned int gpu_stall_by_MSHRwb; -extern unsigned int gpu_stall_shd_mem; -extern unsigned int gpu_stall_sh2icnt; -extern int gpgpu_operand_collector; - enum mem_stage_access_type { C_MEM, T_MEM, @@ -130,23 +107,15 @@ enum mem_stage_stall_type { WB_CACHE_RSRV_FAIL, N_MEM_STAGE_STALL_TYPE }; -unsigned int gpu_stall_shd_mem_breakdown[N_MEM_STAGE_ACCESS_TYPE][N_MEM_STAGE_STALL_TYPE] = { {0} }; +unsigned int gpu_stall_shd_mem_breakdown[N_MEM_STAGE_ACCESS_TYPE][N_MEM_STAGE_STALL_TYPE] = { {0} }; unsigned warp_size = 4; int pipe_simd_width; -extern unsigned int **totalbankaccesses; //bankaccesses[shader id][dram chip id][bank id] -extern unsigned int *MCB_accesses; //upon cache miss, tracks which memory controllers accessed by a warp -extern unsigned int *num_MCBs_accessed; //tracks how many memory controllers are accessed whenever any thread in a warp misses in cache -extern unsigned int *max_return_queue_length; - unsigned int *shader_cycle_distro; - unsigned int g_waiting_at_barrier = 0; - unsigned int gpgpu_shmem_size = 16384; unsigned int gpgpu_shader_registers = 8192; unsigned int gpgpu_shader_cta = 8; - unsigned int gpgpu_n_load_insn = 0; unsigned int gpgpu_n_store_insn = 0; unsigned int gpgpu_n_shmem_insn = 0; @@ -154,77 +123,25 @@ unsigned int gpgpu_n_tex_insn = 0; unsigned int gpgpu_n_const_insn = 0; unsigned int gpgpu_n_param_insn = 0; unsigned int gpgpu_multi_unq_fetches = 0; - -extern int gpgpu_cache_wt_through; - -int gpgpu_shmem_bkconflict = 0; +int gpgpu_shmem_bkconflict = 0; unsigned int gpgpu_n_shmem_bkconflict = 0; -int gpgpu_n_shmem_bank = 16; - -int gpgpu_cache_bkconflict = 0; +int gpgpu_n_shmem_bank = 16; +int gpgpu_cache_bkconflict = 0; unsigned int gpgpu_n_cache_bkconflict = 0; unsigned int gpgpu_n_cmem_portconflict = 0; -int gpgpu_n_cache_bank = 16; - -extern int gpu_runtime_stat_flag; -int gpgpu_warpdistro_shader = -1; - -int gpgpu_interwarp_mshr_merge = 0; -int gpgpu_n_intrawarp_mshr_merge = 0; - -extern int gpgpu_partial_write_mask; -int gpgpu_n_partial_writes = 0; - -extern int gpgpu_n_mem_write_local; -extern int gpgpu_n_mem_write_global; - -#ifndef MhZ - #define MhZ *1000000 -#endif -extern double core_freq; -extern double icnt_freq; -extern double dram_freq; -extern double l2_freq; - -int gpgpu_shmem_port_per_bank = 4; -int gpgpu_cache_port_per_bank = 4; -int gpgpu_const_port_per_bank = 4; -int gpgpu_shmem_pipe_speedup = 2; - +int gpgpu_n_cache_bank = 16; +int gpgpu_warpdistro_shader = -1; +int gpgpu_interwarp_mshr_merge = 0; +int gpgpu_n_intrawarp_mshr_merge = 0; +int gpgpu_n_partial_writes = 0; +int gpgpu_shmem_port_per_bank = 4; +int gpgpu_cache_port_per_bank = 4; +int gpgpu_const_port_per_bank = 4; +int gpgpu_shmem_pipe_speedup = 2; unsigned int gpu_max_cta_per_shader = 8; unsigned int gpu_padded_cta_size = 32; -int gpgpu_local_mem_map = 1; - - -int gpgpu_operand_collector_num_units = 4; - -extern int pdom_sched_type; -extern int n_pdom_sc_orig_stat; -extern int n_pdom_sc_single_stat; -extern int gpgpu_cuda_sim; -extern unsigned long long gpu_tot_sim_cycle; - -void ptx_decode_inst( void *thd, unsigned *op, int *i1, int *i2, int *i3, int *i4, int *o1, int *o2, int *o3, int *o4, int *vectorin, int *vectorout, int *arch_reg ); -void ptx_exec_inst( void *thd, address_type *addr, memory_space_t *space, unsigned *data_size, dram_callback_t* callback, unsigned warp_active_mask); -void ptx_sim_free_sm( class ptx_thread_info** thread_info ); -unsigned ptx_sim_init_thread( class ptx_thread_info** thread_info, int sid, unsigned tid,unsigned threads_left,unsigned num_threads, core_t *core, unsigned hw_cta_id, unsigned hw_warp_id); -unsigned ptx_sim_cta_size(); -const struct gpgpu_ptx_sim_kernel_info* ptx_sim_kernel_info(); -void set_option_gpgpu_spread_blocks_across_cores(int option); -int ptx_thread_done( void *thr ); -unsigned ptx_thread_donecycle( void *thr ); -int ptx_thread_get_next_pc( void *thd ); -void* ptx_thread_get_next_finfo( void *thd ); -int ptx_thread_at_barrier( void *thd ); -int ptx_thread_all_at_barrier( void *thd ); -unsigned long long ptx_thread_get_cta_uid( void *thd ); -void ptx_thread_reset_barrier( void *thd ); -void ptx_thread_release_barrier( void *thd ); -void ptx_print_insn( address_type pc, FILE *fp ); -int ptx_set_tex_cache_linesize( unsigned linesize); -void time_vector_update(unsigned int uid,int slot ,long int cycle,int type); - - +int gpgpu_local_mem_map = 1; +int gpgpu_operand_collector_num_units = 4; ///////////////////////////////////////////////////////////////////////////// /*-------------------------------------------------------------------------*/ diff --git a/src/gpgpu-sim/visualizer.cc b/src/gpgpu-sim/visualizer.cc index a0733fd..d74c5fb 100644 --- a/src/gpgpu-sim/visualizer.cc +++ b/src/gpgpu-sim/visualizer.cc @@ -59,6 +59,8 @@ * Vancouver, BC V6T 1Z4 */ +#include "visualizer.h" + #include "gpu-sim.h" #include "l2cache.h" #include "shader.h" @@ -68,19 +70,13 @@ #include extern unsigned int gpu_mem_n_bk; -extern shader_core_ctx_t **sc; -extern dram_t **dram; extern unsigned long long int mf_total_lat; extern unsigned num_mfs; -extern unsigned long long gpu_sim_cycle; -extern unsigned long long gpu_sim_insn; extern unsigned long long gpu_tot_sim_insn; -extern unsigned long long gpu_completed_thread; extern unsigned int gpgpu_n_sent_writes; extern unsigned int gpgpu_n_processed_writes; extern unsigned int gpgpu_n_cache_bkconflict; extern unsigned int gpgpu_n_shmem_bkconflict; -extern unsigned int gpu_stall_by_MSHRwb; extern unsigned int *max_return_queue_length; extern unsigned max_mrq_latency; extern unsigned max_dq_latency; diff --git a/src/gpgpu-sim/visualizer.h b/src/gpgpu-sim/visualizer.h new file mode 100644 index 0000000..bbc94bf --- /dev/null +++ b/src/gpgpu-sim/visualizer.h @@ -0,0 +1,77 @@ +/* + * Copyright © 2009 by Tor M. Aamodt, Wilson W. L. Fung and the University of + * British Columbia, Vancouver, BC V6T 1Z4, All Rights Reserved. + * + * THIS IS A LEGAL DOCUMENT BY DOWNLOADING GPGPU-SIM, YOU ARE AGREEING TO THESE + * TERMS AND CONDITIONS. + * + * 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 OWNERS 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. + * + * NOTE: The files libcuda/cuda_runtime_api.c and src/cuda-sim/cuda-math.h + * are derived from the CUDA Toolset available from http://www.nvidia.com/cuda + * (property of NVIDIA). The files benchmarks/BlackScholes/ and + * benchmarks/template/ are derived from the CUDA SDK available from + * http://www.nvidia.com/cuda (also property of NVIDIA). The files from + * src/intersim/ are derived from Booksim (a simulator provided with the + * textbook "Principles and Practices of Interconnection Networks" available + * from http://cva.stanford.edu/books/ppin/). As such, those files are bound by + * the corresponding legal terms and conditions set forth separately (original + * copyright notices are left in files from these sources and where we have + * modified a file our copyright notice appears before the original copyright + * notice). + * + * Using this version of GPGPU-Sim requires a complete installation of CUDA + * which is distributed seperately by NVIDIA under separate terms and + * conditions. To use this version of GPGPU-Sim with OpenCL requires a + * recent version of NVIDIA's drivers which support OpenCL. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * 1. Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimer. + * + * 2. 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. + * + * 3. 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. + * + * 4. This version of GPGPU-SIM is distributed freely for non-commercial use only. + * + * 5. No nonprofit user may place any restrictions on the use of this software, + * including as modified by the user, by any other authorized user. + * + * 6. GPGPU-SIM was developed primarily by Tor M. Aamodt, Wilson W. L. Fung, + * Ali Bakhoda, George L. Yuan, at the University of British Columbia, + * Vancouver, BC V6T 1Z4 + */ + +#ifndef VISUALIZER_H_INCLUDED +#define VISUALIZER_H_INCLUDED + +#include +#include + +void visualizer_options(class OptionParser *opp); +void visualizer_printstat(); +void time_vector_create(int ld_size,int st_size); +void time_vector_print(void); +void time_vector_print_interval2file(FILE *outfile); +void time_vector_print_interval2gzfile(gzFile outfile); +void time_vector_update(unsigned int uid,int slot ,long int cycle,int type); +void check_time_vector_update(unsigned int uid,int slot ,long int latency,int type); + +#endif diff --git a/src/gpgpu-sim/warp_tracker.cc b/src/gpgpu-sim/warp_tracker.cc index 811961d..138011c 100644 --- a/src/gpgpu-sim/warp_tracker.cc +++ b/src/gpgpu-sim/warp_tracker.cc @@ -207,8 +207,6 @@ void init_warp_tracker( ) printf("%zd\n", free_wpt.size()); } -extern signed long long gpu_sim_cycle; - void wpt_register_warp( int *tid_in, shader_core_ctx_t *shd ) { int sid = shd->sid; diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 29d9787..7707b35 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -103,7 +103,6 @@ void gpgpu_ptx_sim_init_perf() } extern unsigned long long gpu_tot_sim_insn; -extern unsigned long long gpu_tot_sim_cycle; static void print_simulation_time() { diff --git a/src/intersim/interconnect_interface.cpp b/src/intersim/interconnect_interface.cpp index 6981e08..34fd39d 100644 --- a/src/intersim/interconnect_interface.cpp +++ b/src/intersim/interconnect_interface.cpp @@ -20,7 +20,6 @@ #include "../gpgpu-sim/gpu-sim.h" #include -extern unsigned long long gpu_sim_cycle; int _flit_size ; extern unsigned int warp_size; @@ -558,7 +557,6 @@ void transfer2boundary_buf(int output) { } void time_vector_update(unsigned int uid, int slot , long int cycle, int type); -extern unsigned long long gpu_tot_sim_cycle; void time_vector_update_icnt_injected(void* data, int input) { mem_fetch_t* mf = (mem_fetch_t*) data; unsigned int uid = mf->write? mf->request_uid : mf->mshr->insts[0].uid; -- cgit v1.3