From b400820185db543ff98fd823a15958444e12fee6 Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Sun, 8 Aug 2010 21:15:00 -0800 Subject: ** fully decouple kernel compilation from kernel launch for opencl changes required to do this (this enables oclReduction SDK 3.1 to pass and is a prerequisite to Fermi style concurrent kernel launching) - separate kernel launch for cuda and opencl (and init_grid functions) - add pgm_info structure for holding kernel function_info pointer in _cl_program object after ptx compilation ** added support for clEnqueueNDRangeKernel with local_work_size==NULL which is part of OpenCL spec and used in oclSortingNetworks (however, this is a braindead implementation that only handles the case where global_work_size[0] is smaller than the max number of threads per shader; moreover oclSortingNetworks is still not working but for what looks like another reason) ** refactoring / cleanup - g_global_symbol_table made static to ptx_parser.cc - remove g_kernel_name_to_symtab_lookup (not really being used) - moving various function prototypes into headers (e.g., if foo() defined in bar.cc, then place prototype in bar.h) - adding icnt_reg_options() [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7168] --- src/gpgpu-sim/gpu-sim.cc | 18 ++++++++++++++---- 1 file changed, 14 insertions(+), 4 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index cd4851b..bb228a5 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -66,6 +66,16 @@ #include "gpu-sim.h" +#include +#include +#include +#include "zlib.h" + +#include "../option_parser.h" +#include "shader.h" +#include "dram.h" +#include "mem_fetch.h" + #include #include "gpu-cache.h" #include "gpu-misc.h" @@ -83,6 +93,7 @@ #include "../intersim/statwraper.h" #include "../abstract_hardware_model.h" #include "../debug.h" +#include "../gpgpusim_entrypoint.h" #include #include @@ -299,10 +310,8 @@ unsigned int more_thread = 1; extern unsigned int n_regconflict_stall; unsigned int warp_conflict_at_writeback = 0; unsigned int gpgpu_commit_pc_beyond_two = 0; -extern int g_network_mode; int gpgpu_cache_wt_through = 0; - //memory access classification int gpgpu_n_mem_read_local = 0; int gpgpu_n_mem_write_local = 0; @@ -595,6 +604,7 @@ inline int mem2device(int memid) { /* Allocate memory for uArch structures */ void init_gpu () { + // initialize the GPU microarchitecture model int i; gpu_max_cycle = gpu_max_cycle_opt; @@ -724,6 +734,7 @@ void init_once(void ) { // return the number of cycle required to run all the trace on the gpu unsigned int run_gpu_sim(int grid_num) { + // run a CUDA grid on the GPU microarchitecture simulator int not_completed; int mem_busy; @@ -1432,7 +1443,6 @@ inline int next_clock_domain(void) return mask; } -extern time_t simulation_starttime; void gpu_sim_loop( int grid_num ) { int clock_mask = next_clock_domain(); @@ -1598,7 +1608,7 @@ void gpu_sim_loop( int grid_num ) time_t days, hrs, minutes, sec; time_t curr_time; time(&curr_time); - unsigned long long elapsed_time = MAX(curr_time - simulation_starttime, 1); + unsigned long long elapsed_time = MAX(curr_time - g_simulation_starttime, 1); days = elapsed_time/(3600*24); hrs = elapsed_time/3600 - 24*days; minutes = elapsed_time/60 - 60*(hrs + 24*days); -- cgit v1.3