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

diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h
index 6662b4c..2c5e0a7 100644
--- a/src/cuda-sim/cuda-sim.h
+++ b/src/cuda-sim/cuda-sim.h
@@ -16,10 +16,11 @@ extern memory_space *g_global_mem;
 extern int g_debug_execution;
 extern int g_debug_thread_uid;
 extern std::map<std::string,function_info*> *g_kernel_name_to_function_lookup;
-extern std::map<std::string,symbol_table*> g_kernel_name_to_symtab_lookup;
 
-extern void   gpgpu_ptx_sim_add_ptxstring( const char *ptx, const char *source_fname );
-extern void   gpgpu_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, struct gpgpu_ptx_sim_arg *);
+extern void   gpgpu_cuda_ptx_sim_init_grid(const char *kernel_key,struct gpgpu_ptx_sim_arg *args, struct dim3 gridDim, struct dim3 blockDim );
+extern void   gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,struct gpgpu_ptx_sim_arg *args, struct dim3 gridDim, struct dim3 blockDim );
+extern void   gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, struct gpgpu_ptx_sim_arg *);
+extern void print_splash();
 extern void*  gpgpu_ptx_sim_malloc( size_t count );
 extern void*  gpgpu_ptx_sim_mallocarray( size_t count );
 extern void   gpgpu_ptx_sim_memcpy_to_gpu( size_t dst_start_addr, const void *src, size_t count );
@@ -39,6 +40,8 @@ extern int    gpgpu_ptx_sim_sizeofTexture(const char* name);
 extern const char* gpgpu_ptx_sim_findNamefromTexture(const struct textureReference* texref);
 extern const struct textureReference* gpgpu_ptx_sim_accessTextureofName(const char* name); 
 extern void read_sim_environment_variables();
-extern void register_ptx_function( const char *name, function_info *impl, symbol_table *symtab );
+extern void register_function_implementation( const char *name, function_info *impl );
+extern void ptxinfo_cuda_addinfo();
+extern void ptxinfo_opencl_addinfo( std::map<std::string,function_info*> &kernels );
 
 #endif