** 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, 2 insertions, 8 deletions

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