summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--CHANGES3
-rw-r--r--libopencl/opencl_runtime_api.cc23
-rw-r--r--src/abstract_hardware_model.h2
-rw-r--r--src/cuda-sim/ptx.l2
-rw-r--r--src/cuda-sim/ptx.y14
-rw-r--r--src/cuda-sim/ptx_sim.cc31
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;
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 23:26:20 +0000