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

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;
 }