** 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]

18 files changed, 457 insertions, 296 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;
 }
diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index 2790e53..51ae840 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -78,6 +78,7 @@
 #include "../src/abstract_hardware_model.h"
 #include "../src/cuda-sim/cuda-sim.h"
 #include "../src/gpgpusim_entrypoint.h"
+#include "../src/cuda-sim/ptx_loader.h"
 
 struct gpgpu_ptx_sim_arg {
    const void *m_start;
@@ -198,6 +199,13 @@ private:
    size_t m_size;
 };
 
+struct pgm_info {
+   std::string   m_source;
+   std::string   m_asm;
+   class symbol_table *m_symtab;
+   std::map<std::string,function_info*> m_kernels;
+};
+
 struct _cl_program {
    _cl_program( cl_context context,
                 cl_uint           count, 
@@ -211,12 +219,11 @@ struct _cl_program {
 
 private:
    cl_context m_context;
-   std::map<cl_uint,std::string> m_strings;
-   std::map<cl_uint,std::string> m_ptx;
+   std::map<cl_uint,pgm_info> m_pgm;
 };
 
 struct _cl_kernel {
-   _cl_kernel( cl_program prog, const char* kernel_name, void *kernel_impl );
+   _cl_kernel( cl_program prog, const char* kernel_name, class function_info *kernel_impl );
    void SetKernelArg(
       cl_uint      arg_index,
       size_t       arg_size,
@@ -225,6 +232,7 @@ struct _cl_kernel {
    std::string name() const { return m_kernel_name; }
    size_t get_workgroup_size();
    cl_program get_program() { return m_prog; }
+   class function_info *get_implementation() { return m_kernel_impl; }
 private:
    unsigned m_uid;
    static unsigned sm_context_uid;
@@ -238,7 +246,7 @@ private:
    };
    
    std::map<unsigned, arg_info> m_args;
-   void *m_kernel_impl;
+   class function_info *m_kernel_impl;
 };
 
 struct _cl_platform_id {
@@ -247,9 +255,6 @@ struct _cl_platform_id {
 
 struct _cl_platform_id g_gpgpu_sim_platform_id;
 
-void *gpgpusim_opencl_getkernel_Object( const char *kernel_name );
-void gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num );
-
 void gpgpusim_exit()
 {
    abort();
@@ -266,7 +271,7 @@ void gpgpusim_opencl_error( const char* func, unsigned line, const char *desc )
    gpgpusim_exit();
 }
 
-_cl_kernel::_cl_kernel( cl_program prog, const char* kernel_name, void *kernel_impl )
+_cl_kernel::_cl_kernel( cl_program prog, const char* kernel_name, class function_info *kernel_impl )
 {
    m_uid = sm_context_uid++;
    m_kernel_name = std::string(kernel_name);
@@ -435,19 +440,33 @@ _cl_program::_cl_program( cl_context        context,
       char *tmp = (char*)malloc(len+1);
       memcpy(tmp,strings[i],len);
       tmp[len] = 0;
-      m_strings[i] = tmp;
+      m_pgm[i].m_source = tmp;
       free(tmp);
    }
 }
 
 extern const char *g_filename;
-unsigned g_source_num;
+
+static pgm_info *sg_info;
+
+void register_ptx_function( const char *name, function_info *impl )
+{
+   sg_info->m_kernels[name] = impl;
+}
+
+extern "C" void ptxinfo_addinfo()
+{
+   ptxinfo_opencl_addinfo( sg_info->m_kernels );
+}
 
 void _cl_program::Build(const char *options)
 {
    printf("GPGPU-Sim OpenCL API: compiling OpenCL kernels...\n"); 
-   std::map<cl_uint,std::string>::iterator i;
-   for( i = m_strings.begin(); i!= m_strings.end(); i++ ) {
+   std::map<cl_uint,pgm_info>::iterator i;
+   for( i = m_pgm.begin(); i!= m_pgm.end(); i++ ) {
+      pgm_info &info=i->second;
+      sg_info = &info;
+      unsigned source_num=i->first;
       char ptx_fname[1024];
       char *use_extracted_ptx = getenv("PTX_SIM_USE_PTX_FILE");
       if( use_extracted_ptx == NULL ) {
@@ -468,8 +487,7 @@ void _cl_program::Build(const char *options)
             exit(1); 
 
          char cl_fname[1024];
-         std::string src = i->second;
-         const char *source = src.c_str();
+         const char *source = info.m_source.c_str();
 
          // call wrapper
          char *ld_library_path_orig = getenv("LD_LIBRARY_PATH");
@@ -512,7 +530,7 @@ void _cl_program::Build(const char *options)
          if( result != 0 ) 
             printf("GPGPU-Sim OpenCL API: could not remove temporary files generated while generating PTX\n");
       } else {
-         snprintf(ptx_fname,1024,"_%u.ptx", g_source_num);
+         snprintf(ptx_fname,1024,"_%u.ptx", source_num);
       }
 
       // read in PTX generated by wrapper
@@ -545,9 +563,8 @@ void _cl_program::Build(const char *options)
             len--;
          }
       }
-      m_ptx[g_source_num] = tmp;
-      gpgpu_ptx_sim_load_ptx_from_string( tmp, g_source_num );
-      g_source_num++;
+      info.m_asm = tmp;
+      info.m_symtab = gpgpu_ptx_sim_load_ptx_from_string( tmp, source_num );
       free(tmp);
    }
    printf("GPGPU-Sim OpenCL API: finished compiling OpenCL kernels.\n"); 
@@ -556,30 +573,39 @@ void _cl_program::Build(const char *options)
 cl_kernel _cl_program::CreateKernel( const char *kernel_name, cl_int *errcode_ret )
 {
    cl_kernel result = NULL;
-   void *kernel_impl = gpgpusim_opencl_getkernel_Object( kernel_name );
-   if( kernel_impl == NULL ) 
-      *errcode_ret = CL_INVALID_PROGRAM_EXECUTABLE;
-   else {
-      result = new _cl_kernel(this,kernel_name,kernel_impl);
-      gpgpu_ptx_sim_register_kernel((const char*)result,kernel_name);
+   class function_info *finfo=NULL;
+   std::map<cl_uint,pgm_info>::iterator f;
+   for( f = m_pgm.begin(); f!= m_pgm.end(); f++ ) {
+      pgm_info &info=f->second;
+      std::map<std::string,function_info*>::iterator k = info.m_kernels.find(kernel_name);
+      if( k != info.m_kernels.end() ) {
+         assert( finfo == NULL ); // kernels with same name in different .cl files
+         finfo = k->second;
+      }
    }
+
+   if( finfo == NULL ) 
+      *errcode_ret = CL_INVALID_PROGRAM_EXECUTABLE;
+   else 
+      result = new _cl_kernel(this,kernel_name,finfo);
    return result;
 }
 
 char *_cl_program::get_ptx()
 {
-   if( m_ptx.empty() ) {
-      printf("GPGPU-Sim PTX OpenCL API: Cannot get PTX before having built program\n");
+   if( m_pgm.empty() ) {
+      printf("GPGPU-Sim PTX OpenCL API: Cannot get PTX before building program\n");
       abort();
    }
    size_t buffer_length= get_ptx_size();
    char *tmp = (char*)calloc(buffer_length,1);
    unsigned n=0;
-   std::map<cl_uint,std::string>::iterator p;
-   for( p=m_ptx.begin(); p != m_ptx.end(); p++ ) {
-      unsigned len = strlen( p->second.c_str() ) + 1;
+   std::map<cl_uint,pgm_info>::iterator p;
+   for( p=m_pgm.begin(); p != m_pgm.end(); p++ ) {
+      const char *ptx = p->second.m_asm.c_str();
+      unsigned len = strlen( ptx ) + 1;
       assert( (n+len-1) < buffer_length );
-      memcpy(tmp+n,p->second.c_str(),len);
+      memcpy(tmp+n,ptx,len);
       n+=len;
    }
    assert( n < buffer_length );
@@ -590,9 +616,9 @@ char *_cl_program::get_ptx()
 size_t _cl_program::get_ptx_size()
 {
    size_t buffer_length=0;
-   std::map<cl_uint,std::string>::iterator p;
-   for( p=m_ptx.begin(); p != m_ptx.end(); p++ ) {
-      buffer_length += p->second.length();
+   std::map<cl_uint,pgm_info>::iterator p;
+   for( p=m_pgm.begin(); p != m_pgm.end(); p++ ) {
+      buffer_length += p->second.m_asm.length();
       buffer_length++;
    }
    buffer_length++;
@@ -800,19 +826,40 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
    printf("GPGPU-Sim OpenCL API: clEnqueueNDRangeKernel '%s' (mode=%s)\n", kernel->name().c_str(),
           g_ptx_sim_mode?"functional simulation":"performance simulation");
    if ( !work_dim || work_dim > 3 ) return CL_INVALID_WORK_DIMENSION;
+   size_t _local_size[3];
+   if( local_work_size != NULL ) {
+      for ( unsigned d=0; d < work_dim; d++ ) 
+         _local_size[d]=local_work_size[d];
+   } else {
+      printf("GPGPU-Sim OpenCL API: clEnqueueNDRangeKernel automatic local work size selection:\n");
+      for ( unsigned d=0; d < work_dim; d++ ) {
+          if( d==0 ) {
+             if( global_work_size[d] <= gpu_n_thread_per_shader ) {
+                _local_size[d] = global_work_size[d];
+             } else { 
+                printf("GPGPU-Sim OpenCL API: ERROR clEnqueueNDRangeKernel does not know how to divide work\n" );
+                return CL_INVALID_WORK_GROUP_SIZE;
+             }
+          } else {
+             _local_size[d] = 1;
+          }
+          printf("GPGPU-Sim OpenCL API: clEnqueueNDRangeKernel global_work_size[%u] = %zu\n", d, global_work_size[d] );
+          printf("GPGPU-Sim OpenCL API: clEnqueueNDRangeKernel local_work_size[%u]  = %zu\n", d, _local_size[d] );
+      }
+   }
    for ( unsigned d=0; d < work_dim; d++ ) {
       _global_size[d] = (int)global_work_size[d];
-      if ( (global_work_size[d] % local_work_size[d]) != 0 )
+      if ( (global_work_size[d] % _local_size[d]) != 0 )
          return CL_INVALID_WORK_GROUP_SIZE;
    }
 
-   assert( global_work_size[0] == local_work_size[0] * (global_work_size[0]/local_work_size[0]) ); // i.e., we can divide into equal CTAs
-   g_cudaGridDim.x = global_work_size[0]/local_work_size[0];
-   g_cudaGridDim.y = (work_dim < 2)?1:(global_work_size[1]/local_work_size[1]);
-   g_cudaGridDim.z = (work_dim < 3)?1:(global_work_size[2]/local_work_size[2]);
-   g_cudaBlockDim.x = local_work_size[0];
-   g_cudaBlockDim.y = (work_dim < 2)?1:local_work_size[1];
-   g_cudaBlockDim.z = (work_dim < 3)?1:local_work_size[2];
+   assert( global_work_size[0] == _local_size[0] * (global_work_size[0]/_local_size[0]) ); // i.e., we can divide into equal CTAs
+   g_cudaGridDim.x = global_work_size[0]/_local_size[0];
+   g_cudaGridDim.y = (work_dim < 2)?1:(global_work_size[1]/_local_size[1]);
+   g_cudaGridDim.z = (work_dim < 3)?1:(global_work_size[2]/_local_size[2]);
+   g_cudaBlockDim.x = _local_size[0];
+   g_cudaBlockDim.y = (work_dim < 2)?1:_local_size[1];
+   g_cudaBlockDim.z = (work_dim < 3)?1:_local_size[2];
 
    cl_int err_val = kernel->bind_args(&g_ptx_sim_params);
    if ( err_val != CL_SUCCESS ) {
@@ -826,9 +873,9 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
    gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1  );
 
    if ( g_ptx_sim_mode )
-      gpgpu_ptx_sim_main_func( (const char*)kernel, g_cudaGridDim, g_cudaBlockDim, g_ptx_sim_params );
+      gpgpu_opencl_ptx_sim_main_func( kernel->get_implementation(), g_cudaGridDim, g_cudaBlockDim, g_ptx_sim_params );
    else
-      gpgpu_ptx_sim_main_perf( (const char*)kernel, g_cudaGridDim, g_cudaBlockDim, g_ptx_sim_params );
+      gpgpu_opencl_ptx_sim_main_perf( kernel->get_implementation(), g_cudaGridDim, g_cudaBlockDim, g_ptx_sim_params );
    g_ptx_sim_params=NULL;
    return CL_SUCCESS;
 }
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 6a0ee36..737b94a 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1114,11 +1114,10 @@ void init_inst_classification_stat() {
 
 
 std::map<std::string,function_info*> *g_kernel_name_to_function_lookup=NULL;
-std::map<std::string,symbol_table*> g_kernel_name_to_symtab_lookup;
 std::map<const void*,std::string> *g_host_to_kernel_entrypoint_name_lookup=NULL;
 extern unsigned g_ptx_thread_info_uid_next;
 
-void gpgpu_ptx_sim_init_grid( const char *kernel_key, struct gpgpu_ptx_sim_arg* args,
+void gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, struct gpgpu_ptx_sim_arg* args,
                                          struct dim3 gridDim, struct dim3 blockDim ) 
 {
    g_gx=0;
@@ -1152,7 +1151,39 @@ void gpgpu_ptx_sim_init_grid( const char *kernel_key, struct gpgpu_ptx_sim_arg*
       abort();
    }
    g_entrypoint_func_info = g_func_info = (*g_kernel_name_to_function_lookup)[kname];
-   g_entrypoint_symbol_table = g_kernel_name_to_symtab_lookup[kname];
+
+   unsigned argcount=0;
+   struct gpgpu_ptx_sim_arg *tmparg = args;
+   while (tmparg) {
+      tmparg = tmparg->m_next;
+      argcount++;
+   }
+
+   unsigned argn=1;
+   while (args) {
+      g_func_info->add_param_data(argcount-argn,args);
+      args = args->m_next;
+      argn++;
+   }
+   g_func_info->finalize(g_param_mem);
+   g_ptx_kernel_count++; 
+   if ( gpgpu_ptx_instruction_classification ) {
+      init_inst_classification_stat();
+   }
+   fflush(stdout);
+}
+
+void gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,struct gpgpu_ptx_sim_arg *args, struct dim3 gridDim, struct dim3 blockDim )
+{
+   g_gx=0;
+   g_gy=0;
+   g_gz=0;
+   g_cudaGridDim = gridDim;
+   g_cudaBlockDim = blockDim;
+   g_sm_idx_offset_next.clear();
+   g_sm_next_index = 0;  
+
+   g_entrypoint_func_info = g_func_info = entry;
 
    unsigned argcount=0;
    struct gpgpu_ptx_sim_arg *tmparg = args;
@@ -1209,20 +1240,6 @@ void gpgpu_ptx_sim_register_kernel(const char *hostFun, const char *deviceFun)
    printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : 0x%Lx\n", deviceFun, (unsigned long long)hostFun);
 }
 
-void register_ptx_function( const char *name, function_info *impl, symbol_table *symtab )
-{
-   printf("GPGPU-Sim PTX: parsing function %s\n", name );
-   if( g_kernel_name_to_function_lookup == NULL )
-      g_kernel_name_to_function_lookup = new std::map<std::string,function_info*>;
-
-   std::map<std::string,function_info*>::iterator i_kernel = g_kernel_name_to_function_lookup->find(name);
-   if (i_kernel != g_kernel_name_to_function_lookup->end() && i_kernel->second != NULL) {
-      printf("GPGPU-Sim PTX: WARNING: Function already parsed once. Overwriting.\n");
-   }
-   (*g_kernel_name_to_function_lookup)[name] = impl;
-   g_kernel_name_to_symtab_lookup[name] = symtab;
-}
-
 std::map<const void*,std::string>   g_const_name_lookup; // indexed by hostVar
 std::map<const void*,std::string>   g_global_name_lookup; // indexed by hostVar
 std::set<std::string>   g_globals;
@@ -1314,9 +1331,7 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co
    fflush(stdout);
 }
 
-int g_ptx_sim_mode=0; 
-// used by libcuda.a if non-zer cudaLaunch() will call gpgpu_ptx_sim_main_func()
-// if zero it calls gpgpu_ptx_sim_main_perf()
+int g_ptx_sim_mode=0; // if non-zero run functional simulation only (i.e., no notion of a clock cycle)
 
 extern "C" int ptx_debug;
 
@@ -1378,13 +1393,13 @@ void read_sim_environment_variables()
 
 
 
-extern time_t simulation_starttime;
+extern time_t g_simulation_starttime;
 
 ptx_cta_info *g_func_cta_info = NULL;
 
 #define MAX(a,b) (((a)>(b))?(a):(b))
 
-void gpgpu_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, struct gpgpu_ptx_sim_arg *args)
+void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, struct gpgpu_ptx_sim_arg *args)
 {
    printf("GPGPU-Sim: Performing Functional Simulation...\n");
 
@@ -1396,7 +1411,7 @@ void gpgpu_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDi
    int i1, i2, i3, i4, o1, o2, o3, o4;
    int vectorin, vectorout;
 
-   gpgpu_ptx_sim_init_grid(kernel_key, args,gridDim,blockDim);
+   gpgpu_cuda_ptx_sim_init_grid(kernel_key, args,gridDim,blockDim);
 
    memory_space *shared_mem = new memory_space_impl<16*1024>("shared",4);
 
@@ -1499,7 +1514,7 @@ void gpgpu_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDi
       StatDisp ( g_inst_op_classification_stat[g_ptx_kernel_count]);
    }
    end_time = time((time_t *)NULL);
-   elapsed_time = MAX(end_time - simulation_starttime, 1);
+   elapsed_time = MAX(end_time - g_simulation_starttime, 1);
 
    days    = elapsed_time/(3600*24);
    hrs     = elapsed_time/3600 - 24*days;
@@ -1609,21 +1624,8 @@ extern "C" void ptxinfo_cmem( unsigned nbytes, unsigned bank )
     g_ptxinfo_kinfo.cmem+=nbytes;
 }
 
-extern "C" void ptxinfo_addinfo()
+void clear_ptxinfo()
 {
-    if ( g_kernel_name_to_function_lookup ) {
-        std::map<std::string,function_info*>::iterator i=g_kernel_name_to_function_lookup->find(g_ptxinfo_kname);
-        if ( (g_kernel_name_to_function_lookup == NULL) || (i == g_kernel_name_to_function_lookup->end()) ) {
-           printf ("GPGPU-Sim PTX: Kernel '%s' in %s not found. Ignoring.\n", g_ptxinfo_kname, g_filename);
-        } else {
-           printf ("GPGPU-Sim PTX: Kernel %s\n", g_ptxinfo_kname);
-           function_info *fi = i->second;
-           fi->set_kernel_info(&g_ptxinfo_kinfo);
-        }
-    } else {
-        printf ("GPGPU-Sim PTX: Kernel '%s' in %s not found (no kernels registered).\n", g_ptxinfo_kname, g_filename);
-    }
-
     free(g_ptxinfo_kname);
     g_ptxinfo_kname=NULL;
     g_ptxinfo_kinfo.regs=0;
@@ -1632,6 +1634,62 @@ extern "C" void ptxinfo_addinfo()
     g_ptxinfo_kinfo.cmem=0;
 }
 
+void ptxinfo_opencl_addinfo( std::map<std::string,function_info*> &kernels )
+{
+   if( !strcmp("__cuda_dummy_entry__",g_ptxinfo_kname) ) {
+      // this string produced by ptxas for empty ptx files (e.g., bandwidth test)
+      clear_ptxinfo();
+      return;
+   }
+   std::map<std::string,function_info*>::iterator k=kernels.find(g_ptxinfo_kname);
+   if( k==kernels.end() ) {
+      printf ("GPGPU-Sim PTX: ERROR ** implementation for '%s' not found.\n", g_ptxinfo_kname );
+      abort();
+   } else {
+      printf ("GPGPU-Sim PTX: Kernel \'%s\' : regs=%u, lmem=%u, smem=%u, cmem=%u\n", 
+              g_ptxinfo_kname,
+              g_ptxinfo_kinfo.regs,
+              g_ptxinfo_kinfo.lmem,
+              g_ptxinfo_kinfo.smem,
+              g_ptxinfo_kinfo.cmem );
+      function_info *finfo = k->second;
+      assert(finfo!=NULL);
+      finfo->set_kernel_info( g_ptxinfo_kinfo );
+   }
+   clear_ptxinfo();
+}
+
+void ptxinfo_cuda_addinfo()
+{
+   if( !strcmp("__cuda_dummy_entry__",g_ptxinfo_kname) ) {
+      // this string produced by ptxas for empty ptx files (e.g., bandwidth test)
+      clear_ptxinfo();
+      return;
+   }
+   if ( g_kernel_name_to_function_lookup ) {
+      std::map<std::string,function_info*>::iterator i=g_kernel_name_to_function_lookup->find(g_ptxinfo_kname);
+      if ( (g_kernel_name_to_function_lookup == NULL) || (i == g_kernel_name_to_function_lookup->end()) ) {
+        printf ("GPGPU-Sim PTX: ERROR ** implementation for '%s' not found.\n", g_ptxinfo_kname );
+        abort();
+      } else {
+        printf ("GPGPU-Sim PTX: Kernel \'%s\' : regs=%u, lmem=%u, smem=%u, cmem=%u\n", 
+                g_ptxinfo_kname,
+                g_ptxinfo_kinfo.regs,
+                g_ptxinfo_kinfo.lmem,
+                g_ptxinfo_kinfo.smem,
+                g_ptxinfo_kinfo.cmem );
+        function_info *fi = i->second;
+        assert(fi!=NULL);
+        fi->set_kernel_info(g_ptxinfo_kinfo);
+      }
+   } else {
+      printf ("GPGPU-Sim PTX: ERROR ** Kernel '%s' not found (no kernels registered).\n", g_ptxinfo_kname );
+      abort();
+   }
+   clear_ptxinfo();
+}
+
+
 void dwf_insert_reconv_pt(address_type pc); 
 
 struct rec_pts {
@@ -1704,6 +1762,19 @@ void dwf_process_reconv_pts()
    }
 }
 
+void register_function_implementation( const char *name, function_info *impl )
+{
+   printf("GPGPU-Sim PTX: parsing function %s\n", name );
+   if( g_kernel_name_to_function_lookup == NULL )
+      g_kernel_name_to_function_lookup = new std::map<std::string,function_info*>;
+
+   std::map<std::string,function_info*>::iterator i_kernel = g_kernel_name_to_function_lookup->find(name);
+   if (i_kernel != g_kernel_name_to_function_lookup->end() && i_kernel->second != NULL) {
+      printf("GPGPU-Sim PTX: WARNING: Function already parsed once. Overwriting.\n");
+   }
+   (*g_kernel_name_to_function_lookup)[name] = impl;
+}
+
 void *gpgpusim_opencl_getkernel_Object( const char *kernel_name )
 {
    std::map<std::string,function_info*>::iterator i=g_kernel_name_to_function_lookup->find(kernel_name);
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
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index cd70daf..5945b72 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -191,6 +191,8 @@ void symbol_table::add_function( function_info *func, const char *filename, unsi
    m_symbols[ func->get_name() ] = s;
 }
 
+void register_ptx_function( const char *name, function_info *impl ); // either libcuda or libopencl
+
 bool symbol_table::add_function_decl( const char *name, int entry_point, function_info **func_info, symbol_table **sym_table )
 {
    std::string key = std::string(name);
@@ -209,13 +211,14 @@ bool symbol_table::add_function_decl( const char *name, int entry_point, functio
       *sym_table = m_function_symtab_lookup[key];
    } else {
       assert( !prior_decl );
-      *sym_table = new symbol_table( "", entry_point, g_global_symbol_table );
+      *sym_table = new symbol_table( "", entry_point, this );
       symbol *null_reg = (*sym_table)->add_variable("_",NULL,0,"",0); 
       null_reg->set_regno(0, 0);
       (*sym_table)->set_name(name);
       (*func_info)->set_symtab(*sym_table);
       m_function_symtab_lookup[key] = *sym_table;
-      register_ptx_function(name,*func_info,*sym_table);
+      assert( (*func_info)->get_symtab() == *sym_table );
+      register_ptx_function(name,*func_info);
    }
    return prior_decl;
 }
@@ -1010,8 +1013,6 @@ void function_info::print_insn( unsigned pc, FILE * fp ) const
    }
 }
 
-static int g_save_embedded_ptx = 0;
-
 extern "C" int ptx_parse();
 extern "C" int ptx__scan_string(const char*);
 extern "C" FILE *ptx_in;
@@ -1021,128 +1022,6 @@ extern "C" int ptxinfo_parse();
 extern "C" int ptxinfo_debug;
 extern "C" FILE *ptxinfo_in;
 
-static void print_ptx_file( const char *p, unsigned source_num, const char *filename );
-
-void ptx_reg_options(option_parser_t opp)
-{
-   option_parser_register(opp, "-save_embedded_ptx", OPT_BOOL, &g_save_embedded_ptx, 
-                "saves ptx files embedded in binary as <n>.ptx",
-                "0");
-}
-
-void gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ) 
-{
-    char buf[1024];
-    snprintf(buf,1024,"_%u.ptx", source_num );
-    if( g_save_embedded_ptx ) {
-       FILE *fp = fopen(buf,"w");
-       fprintf(fp,"%s",p);
-       fclose(fp);
-    }
-    init_parser(buf);
-    ptx__scan_string(p);
-    int errors = ptx_parse ();
-    if ( errors ) {
-        char fname[1024];
-        snprintf(fname,1024,"_ptx_XXXXXX");
-        int fd=mkstemp(fname); 
-        close(fd);
-        printf("GPGPU-Sim PTX: parser error detected, exiting... but first extracting .ptx to \"%s\"\n", fname);
-        FILE *ptxfile = fopen(fname,"w");
-        fprintf(ptxfile,"%s", p );
-        fclose(ptxfile);
-        abort();
-        exit(40);
-    }
-
-    if ( g_debug_execution >= 100 ) 
-       print_ptx_file(p,source_num,buf);
-
-    printf("GPGPU-Sim PTX: finished parsing EMBEDDED .ptx file %s\n",buf);
-
-    char fname[1024];
-    snprintf(fname,1024,"_ptx_XXXXXX");
-    int fd=mkstemp(fname); 
-    close(fd);
-
-    printf("GPGPU-Sim PTX: extracting embedded .ptx to temporary file \"%s\"\n", fname);
-    FILE *ptxfile = fopen(fname,"w");
-    fprintf(ptxfile,"%s",p);
-    fclose(ptxfile);
-
-    char fname2[1024];
-    snprintf(fname2,1024,"_ptx2_XXXXXX");
-    fd=mkstemp(fname2); 
-    close(fd);
-    char commandline2[4096];
-    snprintf(commandline2,4096,"cat %s | sed 's/.version 1.5/.version 1.4/' | sed 's/, texmode_independent//' | sed 's/\\(\\.extern \\.const\\[1\\] .b8 \\w\\+\\)\\[\\]/\\1\\[1\\]/' | sed 's/const\\[.\\]/const\\[0\\]/g' > %s", fname, fname2);
-    int result = system(commandline2);
-    if( result != 0 ) {
-       printf("GPGPU-Sim PTX: ERROR ** while loading PTX (a) %d\n", result);
-       printf("               Ensure you have write access to simulation directory\n");
-       printf("               and have \'cat\' and \'sed\' in your path.\n");
-       exit(1);
-    }
-
-    char tempfile_ptxinfo[1024];
-    snprintf(tempfile_ptxinfo,1024,"%sinfo",fname);
-    char commandline[1024];
-    char extra_flags[1024];
-    extra_flags[0]=0;
-#if CUDART_VERSION >= 3000
-    snprintf(extra_flags,1024,"--gpu-name=sm_20");
-#endif
-    snprintf(commandline,1024,"ptxas %s -v %s --output-file /dev/null 2> %s", 
-             extra_flags, fname2, tempfile_ptxinfo);
-    printf("GPGPU-Sim PTX: generating ptxinfo using \"%s\"\n", commandline);
-    result = system(commandline);
-    if( result != 0 ) {
-       printf("GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result);
-       printf("               Ensure ptxas is in your path.\n");
-       exit(1);
-    }
-
-    ptxinfo_in = fopen(tempfile_ptxinfo,"r");
-    g_ptxinfo_filename = tempfile_ptxinfo;
-    ptxinfo_parse();
-    snprintf(commandline,1024,"rm -f %s %s %s", fname, fname2, tempfile_ptxinfo);
-    printf("GPGPU-Sim PTX: removing ptxinfo using \"%s\"\n", commandline);
-    result = system(commandline);
-    if( result != 0 ) {
-       printf("GPGPU-Sim PTX: ERROR ** while loading PTX (c) %d\n", result);
-       exit(1);
-    }
-}
-
-
-const ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned linenumber );
-
-void print_ptx_file( const char *p, unsigned source_num, const char *filename )
-{
-   printf("\nGPGPU-Sim PTX: file _%u.ptx contents:\n\n", source_num );
-   char *s = strdup(p);
-   char *t = s;
-   unsigned n=1;
-   while ( *t != '\0'  ) {
-      char *u = t;
-      while ( (*u != '\n') && (*u != '\0') ) u++;
-      unsigned last = (*u == '\0');
-      *u = '\0';
-      const ptx_instruction *pI = ptx_instruction_lookup(filename,n);
-      char pc[64];
-      if( pI && pI->get_PC() )
-         snprintf(pc,64,"%4u", pI->get_PC() );
-      else 
-         snprintf(pc,64,"    ");
-      printf("    _%u.ptx  %4u (pc=%s):  %s\n", source_num, n, pc, t );
-      if ( last ) break;
-      t = u+1;
-      n++;
-   }
-   free(s);
-   fflush(stdout);
-}
-
 void gpgpu_ptx_assemble( std::string kname, void *kinfo )
 {
     function_info *func_info = (function_info *)kinfo;
@@ -1150,12 +1029,6 @@ void gpgpu_ptx_assemble( std::string kname, void *kinfo )
        printf("GPGPU-Sim PTX: skipping assembly for extern declared function \'%s\'\n", func_info->get_name().c_str() );
        return;
     }
-    if( g_kernel_name_to_symtab_lookup[ kname ] == NULL ) {
-       printf("\nGPGPU-Sim PTX: ERROR no information for kernel \'%s\'\n"
-              "               this can happen for kernels contained in CUDA\n"
-              "               libraries (such as CUBLAS, CUFFT, or CUDPP)\n", kname.c_str() );
-       exit(1);
-    }
     g_func_info = func_info;
     func_info->ptx_assemble();
 }
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index fc7f192..1b1d864 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -1015,8 +1015,8 @@ public:
       return &m_kernel_info;
    }
 
-   const void set_kernel_info (const struct gpgpu_ptx_sim_kernel_info *info) {
-      m_kernel_info = *info;
+   const void set_kernel_info (const struct gpgpu_ptx_sim_kernel_info &info) {
+      m_kernel_info = info;
    }
    symbol_table *get_symtab()
    {
@@ -1254,14 +1254,11 @@ struct textureInfo {
 
 extern function_info *g_func_info;
 
-extern int g_error_detected;
-extern symbol_table *g_entrypoint_symbol_table;
 extern function_info *g_entrypoint_func_info;
-extern symbol_table *g_global_symbol_table;
 extern std::map<std::string,symbol_table*> g_sym_name_to_symbol_table;
+
 #define GLOBAL_HEAP_START 0x10000000
    // start allocating from this address (lower values used for allocating globals in .ptx file)
-
 #define SHARED_MEM_SIZE_MAX (64*1024)
 #define LOCAL_MEM_SIZE_MAX 1024
 #define MAX_STREAMING_MULTIPROCESSORS 64
@@ -1274,7 +1271,6 @@ extern std::map<std::string,symbol_table*> g_sym_name_to_symbol_table;
 
 #define STATIC_ALLOC_LIMIT (GLOBAL_HEAP_START - (TOTAL_LOCAL_MEM+TOTAL_SHARED_MEM))
 
-void gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num );
 void gpgpu_ptx_assemble( std::string kname, void *kinfo );
 #include "../option_parser.h"
 void ptx_reg_options(option_parser_t opp);
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index fb4e8dd..713bdb3 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -212,20 +212,20 @@ void gpgpu_ptx_sim_load_gpu_kernels()
         ptx_in = fopen( g_filename, "r" );
     gpgpu_ptx_sim_init_memory();
     if (ptx_in) {
-        init_parser(g_filename);
+        symbol_table *symtab=init_parser(g_filename);
         ptx_parse();
         ptxinfo_in = open_ptxinfo(g_filename);
         ptxinfo_parse();
-        load_static_globals(g_global_symbol_table,STATIC_ALLOC_LIMIT,0xFFFFFFFF);
-        load_constants(g_global_symbol_table,STATIC_ALLOC_LIMIT);
+        load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF);
+        load_constants(symtab,STATIC_ALLOC_LIMIT);
     } else {
         if (!g_override_embedded_ptx) {
             printf("GPGPU-Sim PTX: USING EMBEDDED .ptx files...\n"); 
             ptx_info_t *s;
             for ( s=g_ptx_source_array; s!=NULL; s=s->next ) {
-                 gpgpu_ptx_sim_load_ptx_from_string(s->str, ++source_num);
-                 load_static_globals(g_global_symbol_table,STATIC_ALLOC_LIMIT,0xFFFFFFFF);
-                 load_constants(g_global_symbol_table,STATIC_ALLOC_LIMIT);
+                 symbol_table *symtab=gpgpu_ptx_sim_load_ptx_from_string(s->str, ++source_num);
+                 load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF);
+                 load_constants(symtab,STATIC_ALLOC_LIMIT);
             }
         } else {
             const char *filename = NULL;
@@ -243,12 +243,12 @@ void gpgpu_ptx_sim_load_gpu_kernels()
                     printf("Parsing %s..\n", filename);
                     ptx_in = fopen( filename, "r" );
                     free(namelist[n]);
-                    init_parser(filename);
+                    symbol_table *symtab=init_parser(filename);
                     ptx_parse ();
                     ptxinfo_in = open_ptxinfo(filename);
                     ptxinfo_parse();
-                    load_static_globals(g_global_symbol_table,STATIC_ALLOC_LIMIT,0xFFFFFFFF);
-                    load_constants(g_global_symbol_table,STATIC_ALLOC_LIMIT);
+                    load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF);
+                    load_constants(symtab,STATIC_ALLOC_LIMIT);
                 }
                 free(namelist);
             }
@@ -298,3 +298,124 @@ void gpgpu_ptx_sim_add_ptxstring( const char *ptx_string, const char *sourcefnam
         l_ptx_source->next = t;
     }
 }
+
+static int g_save_embedded_ptx = 0;
+
+void ptx_reg_options(option_parser_t opp)
+{
+   option_parser_register(opp, "-save_embedded_ptx", OPT_BOOL, &g_save_embedded_ptx, 
+                "saves ptx files embedded in binary as <n>.ptx",
+                "0");
+}
+
+void print_ptx_file( const char *p, unsigned source_num, const char *filename )
+{
+   printf("\nGPGPU-Sim PTX: file _%u.ptx contents:\n\n", source_num );
+   char *s = strdup(p);
+   char *t = s;
+   unsigned n=1;
+   while ( *t != '\0'  ) {
+      char *u = t;
+      while ( (*u != '\n') && (*u != '\0') ) u++;
+      unsigned last = (*u == '\0');
+      *u = '\0';
+      const ptx_instruction *pI = ptx_instruction_lookup(filename,n);
+      char pc[64];
+      if( pI && pI->get_PC() )
+         snprintf(pc,64,"%4u", pI->get_PC() );
+      else 
+         snprintf(pc,64,"    ");
+      printf("    _%u.ptx  %4u (pc=%s):  %s\n", source_num, n, pc, t );
+      if ( last ) break;
+      t = u+1;
+      n++;
+   }
+   free(s);
+   fflush(stdout);
+}
+
+symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ) 
+{
+    char buf[1024];
+    snprintf(buf,1024,"_%u.ptx", source_num );
+    if( g_save_embedded_ptx ) {
+       FILE *fp = fopen(buf,"w");
+       fprintf(fp,"%s",p);
+       fclose(fp);
+    }
+    symbol_table *symtab=init_parser(buf);
+    ptx__scan_string(p);
+    int errors = ptx_parse ();
+    if ( errors ) {
+        char fname[1024];
+        snprintf(fname,1024,"_ptx_XXXXXX");
+        int fd=mkstemp(fname); 
+        close(fd);
+        printf("GPGPU-Sim PTX: parser error detected, exiting... but first extracting .ptx to \"%s\"\n", fname);
+        FILE *ptxfile = fopen(fname,"w");
+        fprintf(ptxfile,"%s", p );
+        fclose(ptxfile);
+        abort();
+        exit(40);
+    }
+
+    if ( g_debug_execution >= 100 ) 
+       print_ptx_file(p,source_num,buf);
+
+    printf("GPGPU-Sim PTX: finished parsing EMBEDDED .ptx file %s\n",buf);
+
+    char fname[1024];
+    snprintf(fname,1024,"_ptx_XXXXXX");
+    int fd=mkstemp(fname); 
+    close(fd);
+
+    printf("GPGPU-Sim PTX: extracting embedded .ptx to temporary file \"%s\"\n", fname);
+    FILE *ptxfile = fopen(fname,"w");
+    fprintf(ptxfile,"%s",p);
+    fclose(ptxfile);
+
+    char fname2[1024];
+    snprintf(fname2,1024,"_ptx2_XXXXXX");
+    fd=mkstemp(fname2); 
+    close(fd);
+    char commandline2[4096];
+    snprintf(commandline2,4096,"cat %s | sed 's/.version 1.5/.version 1.4/' | sed 's/, texmode_independent//' | sed 's/\\(\\.extern \\.const\\[1\\] .b8 \\w\\+\\)\\[\\]/\\1\\[1\\]/' | sed 's/const\\[.\\]/const\\[0\\]/g' > %s", fname, fname2);
+    int result = system(commandline2);
+    if( result != 0 ) {
+       printf("GPGPU-Sim PTX: ERROR ** while loading PTX (a) %d\n", result);
+       printf("               Ensure you have write access to simulation directory\n");
+       printf("               and have \'cat\' and \'sed\' in your path.\n");
+       exit(1);
+    }
+
+    char tempfile_ptxinfo[1024];
+    snprintf(tempfile_ptxinfo,1024,"%sinfo",fname);
+    char commandline[1024];
+    char extra_flags[1024];
+    extra_flags[0]=0;
+#if CUDART_VERSION >= 3000
+    snprintf(extra_flags,1024,"--gpu-name=sm_20");
+#endif
+    snprintf(commandline,1024,"ptxas %s -v %s --output-file /dev/null 2> %s", 
+             extra_flags, fname2, tempfile_ptxinfo);
+    printf("GPGPU-Sim PTX: generating ptxinfo using \"%s\"\n", commandline);
+    result = system(commandline);
+    if( result != 0 ) {
+       printf("GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result);
+       printf("               Ensure ptxas is in your path.\n");
+       exit(1);
+    }
+
+    ptxinfo_in = fopen(tempfile_ptxinfo,"r");
+    g_ptxinfo_filename = tempfile_ptxinfo;
+    ptxinfo_parse();
+    snprintf(commandline,1024,"rm -f %s %s %s", fname, fname2, tempfile_ptxinfo);
+    printf("GPGPU-Sim PTX: removing ptxinfo using \"%s\"\n", commandline);
+    result = system(commandline);
+    if( result != 0 ) {
+       printf("GPGPU-Sim PTX: ERROR ** while loading PTX (c) %d\n", result);
+       exit(1);
+    }
+    return symtab;
+}
+
diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h
index 33c6dc9..95130f1 100644
--- a/src/cuda-sim/ptx_loader.h
+++ b/src/cuda-sim/ptx_loader.h
@@ -76,5 +76,6 @@ extern bool g_override_embedded_ptx;
  
 void gpgpu_ptx_sim_load_gpu_kernels();
 void gpgpu_ptx_sim_add_ptxstring( const char *ptx_string, const char *sourcefname );
+class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num );
 
 #endif
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,
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();
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index cd4851b..bb228a5 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -66,6 +66,16 @@
 
 #include "gpu-sim.h"
 
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include "zlib.h"
+
+#include "../option_parser.h"
+#include "shader.h"
+#include "dram.h"
+#include "mem_fetch.h"
+
 #include <time.h>
 #include "gpu-cache.h"
 #include "gpu-misc.h"
@@ -83,6 +93,7 @@
 #include "../intersim/statwraper.h"
 #include "../abstract_hardware_model.h"
 #include "../debug.h"
+#include "../gpgpusim_entrypoint.h"
 
 #include <stdio.h>
 #include <string.h>
@@ -299,10 +310,8 @@ unsigned int more_thread = 1;
 extern unsigned int n_regconflict_stall;
 unsigned int warp_conflict_at_writeback = 0;
 unsigned int gpgpu_commit_pc_beyond_two = 0;
-extern int g_network_mode;
 int gpgpu_cache_wt_through = 0;
 
-
 //memory access classification
 int gpgpu_n_mem_read_local = 0;
 int gpgpu_n_mem_write_local = 0;
@@ -595,6 +604,7 @@ inline int mem2device(int memid) {
 /* Allocate memory for uArch structures */
 void init_gpu () 
 { 
+   // initialize the GPU microarchitecture model
    int i;
 
    gpu_max_cycle = gpu_max_cycle_opt;
@@ -724,6 +734,7 @@ void init_once(void ) {
 // return the number of cycle required to run all the trace on the gpu 
 unsigned int run_gpu_sim(int grid_num) 
 {
+   // run a CUDA grid on the GPU microarchitecture simulator
 
    int not_completed;
    int mem_busy;
@@ -1432,7 +1443,6 @@ inline int next_clock_domain(void)
    return mask;
 }
 
-extern time_t simulation_starttime;
 void gpu_sim_loop( int grid_num ) 
 {
    int clock_mask = next_clock_domain();
@@ -1598,7 +1608,7 @@ void gpu_sim_loop( int grid_num )
          time_t days, hrs, minutes, sec;
          time_t curr_time;
          time(&curr_time);
-         unsigned long long  elapsed_time = MAX(curr_time - simulation_starttime, 1);
+         unsigned long long  elapsed_time = MAX(curr_time - g_simulation_starttime, 1);
          days    = elapsed_time/(3600*24);
          hrs     = elapsed_time/3600 - 24*days;
          minutes = elapsed_time/60 - 60*(hrs + 24*days);
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index 27d83ad..2e9794f 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -66,16 +66,6 @@
  * Vancouver, BC V6T 1Z4
  */
 
-#include <stdio.h>
-#include <stdlib.h>
-#include <math.h>
-#include "zlib.h"
-
-#include "../option_parser.h"
-#include "shader.h"
-#include "dram.h"
-#include "mem_fetch.h"
-
 #ifndef GPU_SIM_H
 #define GPU_SIM_H
 
@@ -113,5 +103,8 @@ int mem_ctrl_full( int mc_id );
 
 void dramqueue_latency_log_dump();
 void dump_pipeline_impl( int mask, int s, int m );
+unsigned int run_gpu_sim(int grid_num);
+extern void gpu_reg_options(class OptionParser * opp);
+extern void init_gpu();
 
 #endif
diff --git a/src/gpgpu-sim/icnt_wrapper.cc b/src/gpgpu-sim/icnt_wrapper.cc
index 3e38269..aa55110 100644
--- a/src/gpgpu-sim/icnt_wrapper.cc
+++ b/src/gpgpu-sim/icnt_wrapper.cc
@@ -74,8 +74,16 @@ icnt_pop_p        icnt_pop;
 icnt_transfer_p   icnt_transfer;
 icnt_busy_p       icnt_busy;
 
-extern int   g_network_mode;
-extern char* g_network_config_filename;
+int   g_network_mode;
+char* g_network_config_filename;
+
+#include "../option_parser.h"
+
+void icnt_reg_options( class OptionParser * opp )
+{
+   option_parser_register(opp, "-network_mode", OPT_INT32, &g_network_mode, "Interconnection network mode", "1");
+   option_parser_register(opp, "-inter_config_file", OPT_CSTR, &g_network_config_filename, "Interconnection network config file", "mesh");
+}
 
 void icnt_init( unsigned int n_shader, unsigned int n_mem )
 {
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
diff --git a/src/gpgpu-sim/visualizer.cc b/src/gpgpu-sim/visualizer.cc
index d0f406c..0a2d920 100644
--- a/src/gpgpu-sim/visualizer.cc
+++ b/src/gpgpu-sim/visualizer.cc
@@ -60,6 +60,7 @@
  */
 
 #include "gpu-sim.h"
+#include "shader.h"
 #include "../option_parser.h"
 #include <time.h>
 #include <string.h>
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc
index 6ce564b..29d9787 100644
--- a/src/gpgpusim_entrypoint.cc
+++ b/src/gpgpusim_entrypoint.cc
@@ -64,69 +64,52 @@
 
 #include "gpgpusim_entrypoint.h"
 #include <stdio.h>
-#include <time.h>
 
 #include "option_parser.h"
 #include "cuda-sim/cuda-sim.h"
 #include "cuda-sim/ptx_ir.h"
 #include "cuda-sim/ptx_parser.h"
+#include "gpgpu-sim/gpu-sim.h"
+#include "gpgpu-sim/icnt_wrapper.h"
 
 #define MAX(a,b) (((a)>(b))?(a):(b))
 
 struct gpgpu_ptx_sim_arg *grid_params;
-int g_grid_num=0;
-int g_argc = 3;
-const char *g_argv[] = {"", "-config","gpgpusim.config"};
+static int sg_grid_num=0;
+static int sg_argc = 3;
+static const char *sg_argv[] = {"", "-config","gpgpusim.config"};
 
-unsigned int run_gpu_sim(int grid_num);
-void gpgpu_ptx_sim_init_grid(const char *kernel_key,struct gpgpu_ptx_sim_arg *args, struct dim3 gridDim, struct dim3 blockDim );
-
-int   g_network_mode = 0;
-char* g_network_config_filename;
-option_parser_t opp;
-extern void print_splash();
-
-extern void gpu_reg_options(option_parser_t opp);
-extern void init_gpu();
-
-time_t simulation_starttime;
+time_t g_simulation_starttime;
 
 void gpgpu_ptx_sim_init_perf()
 {
    print_splash();
    read_sim_environment_variables();
    read_parser_environment_variables();
-   opp = option_parser_create();
-   option_parser_register(opp, "-network_mode", OPT_INT32, &g_network_mode, "Interconnection network mode", "1");
-   option_parser_register(opp, "-inter_config_file", OPT_CSTR, &g_network_config_filename, "Interconnection network config file", "mesh");
+   option_parser_t opp = option_parser_create();
+   icnt_reg_options(opp);
    gpu_reg_options(opp); // register GPU microrachitecture options
    ptx_reg_options(opp);
-   option_parser_cmdline(opp, g_argc, g_argv); // parse configuration options
+   option_parser_cmdline(opp, sg_argc, sg_argv); // parse configuration options
 
    srand(1); 
 
    fprintf(stdout, "GPGPU-Sim: Configuration options:\n\n");
    option_parser_print(opp, stdout);
-   init_gpu(); // initialize the GPU microarchitecture model
+   init_gpu();
    fprintf(stdout, "GPU performance model initialization complete.\n");
 
-   simulation_starttime = time((time_t *)NULL);
+   g_simulation_starttime = time((time_t *)NULL);
 }
 
 extern unsigned long long  gpu_tot_sim_insn;
 extern unsigned long long  gpu_tot_sim_cycle;
 
-int gpgpu_ptx_sim_main_perf( const char *kernel_key, struct dim3 gridDim, struct dim3 blockDim, struct gpgpu_ptx_sim_arg *grid_params )
+static void print_simulation_time()
 {
    time_t current_time, difference, d, h, m, s;
-   gpgpu_ptx_sim_init_grid(kernel_key,grid_params,gridDim,blockDim);
-
-   run_gpu_sim(g_grid_num); // run a CUDA grid on the GPU microarchitecture simulator
-
-   g_grid_num++;
-
    current_time = time((time_t *)NULL);
-   difference = MAX(current_time - simulation_starttime, 1);
+   difference = MAX(current_time - g_simulation_starttime, 1);
 
    d = difference/(3600*24);
    h = difference/3600 - 24*d;
@@ -139,6 +122,35 @@ int gpgpu_ptx_sim_main_perf( const char *kernel_key, struct dim3 gridDim, struct
    printf("gpgpu_simulation_rate = %u (inst/sec)\n", (unsigned)(gpu_tot_sim_insn / difference) );
    printf("gpgpu_simulation_rate = %u (cycle/sec)\n", (unsigned)(gpu_tot_sim_cycle / difference) );
    fflush(stdout);
+}
 
+int gpgpu_cuda_ptx_sim_main_perf( const char *kernel_key, 
+                                  struct dim3 gridDim, 
+                                  struct dim3 blockDim, 
+                                  struct gpgpu_ptx_sim_arg *grid_params )
+{
+   gpgpu_cuda_ptx_sim_init_grid(kernel_key,grid_params,gridDim,blockDim);
+   run_gpu_sim(sg_grid_num++);
+   print_simulation_time();
    return 0;
 }
+
+int gpgpu_opencl_ptx_sim_main_perf( class function_info *entry, 
+                                  struct dim3 gridDim, 
+                                  struct dim3 blockDim, 
+                                  struct gpgpu_ptx_sim_arg *grid_params )
+{
+   gpgpu_opencl_ptx_sim_init_grid(entry,grid_params,gridDim,blockDim);
+   run_gpu_sim(sg_grid_num++);
+   print_simulation_time();
+   return 0;
+}
+
+int gpgpu_opencl_ptx_sim_main_func( class function_info *entry, 
+                                  struct dim3 gridDim, 
+                                  struct dim3 blockDim, 
+                                  struct gpgpu_ptx_sim_arg *grid_params )
+{
+   printf("GPGPU-Sim PTX API: OpenCL functional-only simulation not yet implemented (use performance simulation)\n");
+   exit(1);
+}
diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h
index 8de05c7..2563c5e 100644
--- a/src/gpgpusim_entrypoint.h
+++ b/src/gpgpusim_entrypoint.h
@@ -67,7 +67,23 @@
 
 #include "abstract_hardware_model.h"
 
+#include <time.h>
+extern time_t g_simulation_starttime;
+
 void gpgpu_ptx_sim_init_perf();
-int gpgpu_ptx_sim_main_perf( const char *kernel_key, struct dim3 gridDim, struct dim3 blockDim, struct gpgpu_ptx_sim_arg *grid_params );
 
+int  gpgpu_cuda_ptx_sim_main_perf( const char *kernel_key, 
+                                   struct dim3 gridDim, 
+                                   struct dim3 blockDim, 
+                                   struct gpgpu_ptx_sim_arg *grid_params );
+
+int gpgpu_opencl_ptx_sim_main_perf( class function_info *entry, 
+                                  struct dim3 gridDim, 
+                                  struct dim3 blockDim, 
+                                  struct gpgpu_ptx_sim_arg *grid_params );
+
+int gpgpu_opencl_ptx_sim_main_func( class function_info *entry, 
+                                  struct dim3 gridDim, 
+                                  struct dim3 blockDim, 
+                                  struct gpgpu_ptx_sim_arg *grid_params );
 #endif
diff --git a/src/option_parser.h b/src/option_parser.h
index e14c08e..5e27ef1 100644
--- a/src/option_parser.h
+++ b/src/option_parser.h
@@ -67,7 +67,7 @@
 #include <stdlib.h>
 
 // pointer to C++ class
-typedef void* option_parser_t;
+typedef class OptionParser *option_parser_t;
 
 // data type of the option
 enum option_dtype {