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

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index df59506..27b4bd3 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -167,9 +167,20 @@ struct gpgpu_ptx_sim_arg {
 struct gpgpu_ptx_sim_arg *g_ptx_sim_params;
 cudaError_t g_last_cudaError;
 
+#include "../src/cuda-sim/ptx_loader.h"
 #include "../src/cuda-sim/cuda-sim.h"
 #include "../src/gpgpusim_entrypoint.h"
 
+void register_ptx_function( const char *name, function_info *impl )
+{
+   register_function_implementation( name, impl );
+}
+
+extern "C" void ptxinfo_addinfo()
+{
+   ptxinfo_cuda_addinfo();
+}
+
 #if defined __APPLE__
 #   define __my_func__    __PRETTY_FUNCTION__
 #else
@@ -822,9 +833,9 @@ __host__ cudaError_t CUDARTAPI cudaLaunch(const char *symbol )
    printf("GPGPU-Sim PTX: cudaLaunch for %p (mode=%s)\n", symbol,
           g_ptx_sim_mode?"functional simulation":"performance simulation");
    if( g_ptx_sim_mode )
-      gpgpu_ptx_sim_main_func( symbol, g_cudaGridDim, g_cudaBlockDim, g_ptx_sim_params );
+      gpgpu_cuda_ptx_sim_main_func( symbol, g_cudaGridDim, g_cudaBlockDim, g_ptx_sim_params );
    else
-      gpgpu_ptx_sim_main_perf( symbol, g_cudaGridDim, g_cudaBlockDim, g_ptx_sim_params );
+      gpgpu_cuda_ptx_sim_main_perf( symbol, g_cudaGridDim, g_cudaBlockDim, g_ptx_sim_params );
    g_ptx_sim_params=NULL;
    return g_last_cudaError = cudaSuccess;
 }