OpenCL newer diver fix

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

6 files changed, 62 insertions, 13 deletions

diff --git a/CHANGES b/CHANGES
index 1057c1d..8119f29 100644
--- a/CHANGES
+++ b/CHANGES
@@ -1,4 +1,4 @@
-CHANGE LOG:
+LOG:
 
 Version 3.1.1+edits (development branch) versus 3.1.1
 - Added support for CUDA 4.1 and 4.2 (PTX 3.0). 
@@ -25,6 +25,7 @@ Version 3.1.1+edits (development branch) versus 3.1.1
   '=' and ';'.  This allow options with a long string of sub-options (e.g. DRAM
   timing) to have a more friendly format. 
 - Changed the DRAM timing option to a more reader-friendly format. 
+- Add support for the PTX v3 codes generated by new OpenCL drivers. 
 - Bug Fixes:
     - Fixed a compile error that happens with newer gcc/g++ versions (4.7.1)
     - Fixed a bug with the association between cuobjdump output and cubin
diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index 5506ea4..1b98881 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -894,7 +894,14 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
       if ( (global_work_size[d] % _local_size[d]) != 0 )
          return CL_INVALID_WORK_GROUP_SIZE;
    }
-
+   if (global_work_offset != NULL){
+	   for ( unsigned d=0; d < work_dim; d++ ) {
+		   if (global_work_offset[d] != 0){
+			   printf("GPGPU-Sim: global id offset is not supported\n");
+			   abort();
+		   }
+	   }
+   }
    assert( global_work_size[0] == _local_size[0] * (global_work_size[0]/_local_size[0]) ); // i.e., we can divide into equal CTAs
    dim3 GridDim;
    GridDim.x = global_work_size[0]/_local_size[0];
@@ -911,13 +918,13 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
       return err_val;
 
    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 );
-
+   if (kernel->get_implementation()->get_ptx_version().ver() <3.0){
+	   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( grid );
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index 9681a67..23e8d00 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -878,7 +878,7 @@ class core_t {
         void initilizeSIMTStack(unsigned warps, unsigned warpsSize);
         warp_inst_t getExecuteWarp(unsigned warpId);
         void get_pdom_stack_top_info( unsigned warpId, unsigned *pc, unsigned *rpc ) const;
-
+        kernel_info_t * get_kernel_info(){ return m_kernel;}
     protected:
         class gpgpu_sim *m_gpu;
         kernel_info_t *m_kernel;
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index de041fc..dd45a67 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -183,7 +183,7 @@ breakaddr  TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
 \.address_size TC; return ADDRESS_SIZE_DIRECTIVE;
 
 \.constptr TC; return CONSTPTR_DIRECTIVE; /* Ptx plus directive for pointer to constant memory */
-
+\.ptr TC; return PTR_DIRECTIVE; /* Added for new OpenCL genrated code */
 
 "%clock" TC; ptx_lval.int_value = CLOCK_REG; return SPECIAL_REGISTER;
 "%halfclock" TC; ptx_lval.int_value = HALFCLOCK_ID; return SPECIAL_REGISTER;
diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y
index 360be51..f8b29c4 100644
--- a/src/cuda-sim/ptx.y
+++ b/src/cuda-sim/ptx.y
@@ -44,6 +44,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 %token  CALLTARGETS_DIRECTIVE
 %token  <int_value> CONST_DIRECTIVE
 %token  CONSTPTR_DIRECTIVE
+%token  PTR_DIRECTIVE
 %token  ENTRY_DIRECTIVE
 %token  EXTERN_DIRECTIVE
 %token  FILE_DIRECTIVE
@@ -242,9 +243,20 @@ param_list: /*empty*/
 	| param_entry { add_directive(); }
 	| param_list COMMA {func_header_info(",");} param_entry { add_directive(); }
 
-param_entry: PARAM_DIRECTIVE { add_space_spec(param_space_unclassified,0); } variable_spec identifier_spec { add_function_arg(); }
+param_entry: PARAM_DIRECTIVE { add_space_spec(param_space_unclassified,0); } variable_spec ptr_spec identifier_spec { add_function_arg(); }
 	| REG_DIRECTIVE { add_space_spec(reg_space,0); } variable_spec identifier_spec { add_function_arg(); }
 
+ptr_spec: /*empty*/
+        | PTR_DIRECTIVE ptr_space_spec ptr_align_spec
+        | PTR_DIRECTIVE ptr_align_spec
+
+ptr_space_spec: GLOBAL_DIRECTIVE
+                                | LOCAL_DIRECTIVE
+                                | CONST_DIRECTIVE
+                                | SHARED_DIRECTIVE
+
+ptr_align_spec: ALIGN_DIRECTIVE INT_OPERAND
+
 statement_block: LEFT_BRACE statement_list RIGHT_BRACE 
 
 statement_list: directive_statement { add_directive(); }
diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc
index b0393fa..29351a0 100644
--- a/src/cuda-sim/ptx_sim.cc
+++ b/src/cuda-sim/ptx_sim.cc
@@ -207,7 +207,36 @@ unsigned ptx_thread_info::get_builtin( int builtin_id, unsigned dim_mod )
       if( dim_mod == 2 ) return m_ctaid.z;
       abort();
       break;
-   case ENVREG_REG: feature_not_implemented( "%envreg" ); return 0;
+   case ENVREG_REG:{
+	int index = builtin_id >> 16;
+	dim3 gdim = this->get_core()->get_kernel_info()->get_grid_dim();
+		switch(index){
+		case 0:
+		case 1:
+		case 2:
+		case 3:
+		case 4:
+		case 5:
+			return 0;
+			break;
+		case 6:
+			return gdim.x;
+		case 7:
+			return gdim.y;
+		case 8:
+			return gdim.z;
+		case 9:
+			if(gdim.z == 1 && gdim.y == 1)
+				return 1;
+			else if(gdim.z == 1)
+				return 2;
+			else
+				return 3;
+			break;
+		default:
+			break;
+		}
+   }
    case GRIDID_REG:
       return m_gridid;
    case LANEID_REG: feature_not_implemented( "%laneid" ); return 0;