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

1 files changed, 10 insertions, 2 deletions

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 )
 {