** 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, 4 insertions, 3 deletions

diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 17e4ca2..05ce397 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -73,7 +73,7 @@ const char *g_filename;
 unsigned g_max_regs_per_thread = 0;
 
 // the program intermediate representation...
-symbol_table *g_global_symbol_table = NULL;
+static symbol_table *g_global_symbol_table = NULL;
 std::map<std::string,symbol_table*> g_sym_name_to_symbol_table;
 static symbol_table *g_current_symbol_table = NULL;
 static std::list<ptx_instruction*> g_instructions;
@@ -133,7 +133,7 @@ void read_parser_environment_variables()
    }
 }
 
-void init_parser( const char *ptx_filename )
+symbol_table *init_parser( const char *ptx_filename )
 {
    g_filename = strdup(ptx_filename);
    g_global_symbol_table = g_current_symbol_table = new symbol_table("global",0,NULL);
@@ -142,6 +142,8 @@ void init_parser( const char *ptx_filename )
 #define DEF(X,Y) g_ptx_token_decode[X] = Y;
 #include "ptx_parser_decode.def"
 #undef DEF
+
+   return g_global_symbol_table;
 }
 
 void init_directive_state()
@@ -192,7 +194,6 @@ void add_function_name( const char *name )
    bool prior_decl = g_global_symbol_table->add_function_decl( name, g_entry_point, &g_func_info, &g_current_symbol_table );
    if( g_entry_point ) {
       g_entrypoint_func_info = g_func_info;
-      g_entrypoint_symbol_table = g_current_symbol_table;
    }
    if( g_add_identifier_cached__identifier ) {
       add_identifier( g_add_identifier_cached__identifier,