** 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]

5 files changed, 30 insertions, 24 deletions

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 <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include "zlib.h"
+
+#include "../option_parser.h"
+#include "shader.h"
+#include "dram.h"
+#include "mem_fetch.h"
+
 #include <time.h>
 #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 <stdio.h>
 #include <string.h>
@@ -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);
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index 27d83ad..2e9794f 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -66,16 +66,6 @@
  * Vancouver, BC V6T 1Z4
  */
 
-#include <stdio.h>
-#include <stdlib.h>
-#include <math.h>
-#include "zlib.h"
-
-#include "../option_parser.h"
-#include "shader.h"
-#include "dram.h"
-#include "mem_fetch.h"
-
 #ifndef GPU_SIM_H
 #define GPU_SIM_H
 
@@ -113,5 +103,8 @@ int mem_ctrl_full( int mc_id );
 
 void dramqueue_latency_log_dump();
 void dump_pipeline_impl( int mask, int s, int m );
+unsigned int run_gpu_sim(int grid_num);
+extern void gpu_reg_options(class OptionParser * opp);
+extern void init_gpu();
 
 #endif
diff --git a/src/gpgpu-sim/icnt_wrapper.cc b/src/gpgpu-sim/icnt_wrapper.cc
index 3e38269..aa55110 100644
--- a/src/gpgpu-sim/icnt_wrapper.cc
+++ b/src/gpgpu-sim/icnt_wrapper.cc
@@ -74,8 +74,16 @@ icnt_pop_p        icnt_pop;
 icnt_transfer_p   icnt_transfer;
 icnt_busy_p       icnt_busy;
 
-extern int   g_network_mode;
-extern char* g_network_config_filename;
+int   g_network_mode;
+char* g_network_config_filename;
+
+#include "../option_parser.h"
+
+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");
+}
 
 void icnt_init( unsigned int n_shader, unsigned int n_mem )
 {
diff --git a/src/gpgpu-sim/icnt_wrapper.h b/src/gpgpu-sim/icnt_wrapper.h
index 998a4f3..0961fc4 100644
--- a/src/gpgpu-sim/icnt_wrapper.h
+++ b/src/gpgpu-sim/icnt_wrapper.h
@@ -75,26 +75,20 @@ typedef void (*icnt_transfer_p)( );
 typedef unsigned (*icnt_busy_p)( );
 typedef void (*icnt_drain_p)( );
 
-
 extern icnt_has_buffer_p icnt_has_buffer;
 extern icnt_push_p       icnt_push;
 extern icnt_pop_p        icnt_pop;
 extern icnt_transfer_p   icnt_transfer;
 extern icnt_busy_p       icnt_busy;
 extern icnt_drain_p      icnt_drain;
+extern int g_network_mode;
 
-/*
-// definition of valid gpu network mode.
-extern enum {
-   INTERSIM = 1,
-   N_NETWORK_MODE
-} gpu_network_mode;
-*/
 enum network_mode {
    INTERSIM = 1,
    N_NETWORK_MODE
 };
 
 void icnt_init( unsigned int n_shader, unsigned int n_mem );
+void icnt_reg_options( class OptionParser * opp );
 
 #endif
diff --git a/src/gpgpu-sim/visualizer.cc b/src/gpgpu-sim/visualizer.cc
index d0f406c..0a2d920 100644
--- a/src/gpgpu-sim/visualizer.cc
+++ b/src/gpgpu-sim/visualizer.cc
@@ -60,6 +60,7 @@
  */
 
 #include "gpu-sim.h"
+#include "shader.h"
 #include "../option_parser.h"
 #include <time.h>
 #include <string.h>