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

diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index 076fa11..8ffe593 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -66,14 +66,17 @@
 #define ptx_parser_INCLUDED
 
 #include "../abstract_hardware_model.h"
-
 extern const char *g_filename;
+extern int g_error_detected;
+
 #ifdef __cplusplus 
+class symbol_table* init_parser(const char*);
+const class ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned linenumber );
 extern "C" {
 #endif
+
 const char *decode_token( int type );
 void read_parser_environment_variables();
-void init_parser(const char*);
 void start_function( int entry_point );
 void add_function_name( const char *fname );
 void init_directive_state();