1. refactoring cuda api code for loading PTX (removing external PTX loading entirely)

2. some bug fixes for warp_inst_t
3. creating a new class, gpgpu_t, which contains the functional "memory" state visible
   to all threads running on a GPU (doing this as part of my continuing effort to hunt 
   down and eradicate every global variable that is not the top level "the gpu")
4. other misc. changes

Almost passing CUDA 3.1 regression?  oclHistogram keeps failing under torque, but
does not fail when run on the command line from the same directory.

[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7827]

1 files changed, 26 insertions, 19 deletions

diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index c4b9c2c..b331f92 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -148,7 +148,7 @@ private:
 };
 
 struct _cl_mem {
-   _cl_mem( cl_mem_flags flags, size_t size , void *host_ptr, cl_int *errcode_ret );
+   _cl_mem( cl_mem_flags flags, size_t size , void *host_ptr, cl_int *errcode_ret, cl_device_id gpu );
    cl_mem device_ptr();
    void* host_ptr();
    bool is_on_host() { return m_is_on_host; }
@@ -295,7 +295,8 @@ _cl_mem::_cl_mem(
    cl_mem_flags flags,
    size_t       size ,
    void *       host_ptr,
-   cl_int *     errcode_ret )
+   cl_int *     errcode_ret,
+   cl_device_id gpu )
 {
    if( errcode_ret ) 
       *errcode_ret = CL_SUCCESS;
@@ -305,7 +306,6 @@ _cl_mem::_cl_mem(
    m_size = size;
    m_host_ptr = host_ptr;
    m_device_ptr = 0;
-   gpgpu_ptx_sim_init_memory();
 
    if( (flags & (CL_MEM_USE_HOST_PTR|CL_MEM_COPY_HOST_PTR)) && host_ptr == NULL ) {
       if( errcode_ret != NULL ) 
@@ -330,9 +330,9 @@ _cl_mem::_cl_mem(
    }
    if( !(flags & (CL_MEM_USE_HOST_PTR|CL_MEM_ALLOC_HOST_PTR)) ) {
       // if not allocating on host, then allocate GPU memory and make a copy
-      m_device_ptr = (size_t) gpgpu_ptx_sim_malloc(size);
+      m_device_ptr = (size_t) gpu->the_device()->gpgpu_ptx_sim_malloc(size);
       if( host_ptr )
-         gpgpu_ptx_sim_memcpy_to_gpu( m_device_ptr, host_ptr, size );
+         gpu->the_device()->gpgpu_ptx_sim_memcpy_to_gpu( m_device_ptr, host_ptr, size );
    }
 }
 
@@ -356,7 +356,7 @@ cl_mem _cl_context::CreateBuffer(
    if( host_ptr && (m_hostptr_to_cl_mem.find(host_ptr) != m_hostptr_to_cl_mem.end()) ) {
       printf("GPGPU-Sim OpenCL API: WARNING ** clCreateBuffer - buffer already created for this host variable\n");
    }
-   cl_mem result = new _cl_mem(flags,size,host_ptr,errcode_ret);
+   cl_mem result = new _cl_mem(flags,size,host_ptr,errcode_ret,m_gpu);
    m_devptr_to_cl_mem[result->device_ptr()] = result;
    if( host_ptr ) 
       m_hostptr_to_cl_mem[host_ptr] = result;
@@ -516,7 +516,8 @@ void _cl_program::Build(const char *options)
          }
       }
       info.m_asm = tmp;
-      info.m_symtab = gpgpu_ptx_sim_load_ptx_from_string( tmp, tmp, source_num );
+      info.m_symtab = gpgpu_ptx_sim_load_ptx_from_string( tmp, source_num );
+      gpgpu_ptxinfo_load_from_string( tmp, source_num );
       free(tmp);
    }
    printf("GPGPU-Sim OpenCL API: finished compiling OpenCL kernels.\n"); 
@@ -828,16 +829,19 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
    if ( err_val != CL_SUCCESS ) 
       return err_val;
 
-   gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1 );
-   gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1  );
-   gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1  );
-   gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1  );
-   gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1  );
+   gpgpu_t *gpu = command_queue->get_device()->the_device();
 
+   gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1, gpu );
+   gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1, gpu  );
+   gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1, gpu );
+   gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
+   gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
+
+   kernel_info_t grid = gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu);
    if ( g_ptx_sim_mode )
-      gpgpu_opencl_ptx_sim_main_func( kernel->get_implementation(), GridDim, BlockDim, params );
+      gpgpu_opencl_ptx_sim_main_func( grid, GridDim, BlockDim, params );
    else
-      gpgpu_opencl_ptx_sim_main_perf( kernel->get_implementation(), GridDim, BlockDim, params );
+      gpgpu_opencl_ptx_sim_main_perf( grid, GridDim, BlockDim, params );
    return CL_SUCCESS;
 }
 
@@ -854,7 +858,8 @@ clEnqueueReadBuffer(cl_command_queue    command_queue,
 {
    if( !blocking_read ) 
       gpgpusim_opencl_warning(__my_func__,__LINE__, "non-blocking read treated as blocking read");
-   gpgpu_ptx_sim_memcpy_from_gpu( ptr, (size_t)buffer, cb );
+   gpgpu_t *gpu = command_queue->get_device()->the_device();
+   gpu->gpgpu_ptx_sim_memcpy_from_gpu( ptr, (size_t)buffer, cb );
    return CL_SUCCESS;
 }
 
@@ -871,7 +876,8 @@ clEnqueueWriteBuffer(cl_command_queue   command_queue,
 {
    if( !blocking_write ) 
       gpgpusim_opencl_warning(__my_func__,__LINE__, "non-blocking write treated as blocking write");
-   gpgpu_ptx_sim_memcpy_to_gpu( (size_t)buffer, ptr, cb );
+   gpgpu_t *gpu = command_queue->get_device()->the_device();
+   gpu->gpgpu_ptx_sim_memcpy_to_gpu( (size_t)buffer, ptr, cb );
    return CL_SUCCESS;
 }
 
@@ -1143,12 +1149,13 @@ clEnqueueCopyBuffer(cl_command_queue    command_queue,
    if( src == NULL || dst == NULL ) 
       return CL_INVALID_MEM_OBJECT;
 
+   gpgpu_t *gpu = command_queue->get_device()->the_device();
    if( src->is_on_host() && !dst->is_on_host() )
-      gpgpu_ptx_sim_memcpy_to_gpu( ((size_t)dst->device_ptr())+dst_offset, ((char*)src->host_ptr())+src_offset, cb );
+      gpu->gpgpu_ptx_sim_memcpy_to_gpu( ((size_t)dst->device_ptr())+dst_offset, ((char*)src->host_ptr())+src_offset, cb );
    else if( !src->is_on_host() && dst->is_on_host() ) 
-      gpgpu_ptx_sim_memcpy_from_gpu( ((char*)dst->host_ptr())+dst_offset, ((size_t)src->device_ptr())+src_offset, cb );
+      gpu->gpgpu_ptx_sim_memcpy_from_gpu( ((char*)dst->host_ptr())+dst_offset, ((size_t)src->device_ptr())+src_offset, cb );
    else if( !src->is_on_host() && !dst->is_on_host() ) 
-      gpgpu_ptx_sim_memcpy_gpu_to_gpu( ((size_t)dst->device_ptr())+dst_offset, ((size_t)src->device_ptr())+src_offset, cb );
+      gpu->gpgpu_ptx_sim_memcpy_gpu_to_gpu( ((size_t)dst->device_ptr())+dst_offset, ((size_t)src->device_ptr())+src_offset, cb );
    else
       opencl_not_implemented(__my_func__,__LINE__);
    return CL_SUCCESS;