diff options
| -rw-r--r-- | libopencl/opencl_runtime_api.cc | 54 | ||||
| -rw-r--r-- | src/cuda-sim/ptx.y | 18 |
2 files changed, 71 insertions, 1 deletions
diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 97a54d8..752bfdf 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -263,15 +263,27 @@ void _cl_kernel::SetKernelArg( cl_int _cl_kernel::bind_args( gpgpu_ptx_sim_arg_list_t &arg_list ) { + size_t offset = 0; + assert( arg_list.empty() ); unsigned k=0; std::map<unsigned, arg_info>::iterator i; for( i = m_args.begin(); i!=m_args.end(); i++ ) { if( i->first != k ) return CL_INVALID_KERNEL_ARGS; + arg_info arg = i->second; - gpgpu_ptx_sim_arg param( arg.m_arg_value, arg.m_arg_size, 0); + const symbol *sym = m_kernel_impl->get_arg(i->first); + const type_info_key &t = sym->type()->get_key(); + + int align = (t.get_alignment_spec() == -1) ? arg.m_arg_size : t.get_alignment_spec(); + if( offset % align ) + offset += (align - (offset % align)); + + gpgpu_ptx_sim_arg param( arg.m_arg_value, arg.m_arg_size, offset ); arg_list.push_front( param ); + + offset += arg.m_arg_size; k++; } return CL_SUCCESS; @@ -950,6 +962,17 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, 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); + + //do dynamic PDOM analysis for performance simulation scenario + std::string kname = grid->name(); + function_info *kernel_func_info = grid->entry(); + if (kernel_func_info->is_pdom_set()) { + printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", kname.c_str() ); + } else { + printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", kname.c_str() ); + kernel_func_info->do_pdom(); + kernel_func_info->set_pdom(); + } if ( g_ptx_sim_mode ) gpgpu_opencl_ptx_sim_main_func( grid ); else @@ -1256,6 +1279,35 @@ clGetProgramInfo(cl_program program, } extern CL_API_ENTRY cl_int CL_API_CALL +clGetProgramBuildInfo (cl_program program, + cl_device_id device, + cl_program_build_info param_name, + size_t param_value_size, + void * param_value, + size_t * param_value_size_ret) CL_API_SUFFIX__VERSION_1_0 +{ + char *buf = (char*)param_value; + + switch( param_name ) { + case CL_PROGRAM_BUILD_STATUS: + CL_CASE( cl_build_status, CL_BUILD_SUCCESS ); + break; + case CL_PROGRAM_BUILD_OPTIONS: + case CL_PROGRAM_BUILD_LOG: + CL_STRING_CASE( "" ); + break; + case CL_PROGRAM_BINARY_TYPE: + CL_CASE( cl_program_binary_type, CL_PROGRAM_BINARY_TYPE_EXECUTABLE ); + break; + default: + return CL_INVALID_VALUE; + break; + } + + return CL_SUCCESS; +} + +extern CL_API_ENTRY cl_int CL_API_CALL clEnqueueCopyBuffer(cl_command_queue command_queue, cl_mem src_buffer, cl_mem dst_buffer, diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index 45392fb..260564f 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -287,6 +287,7 @@ ptr_align_spec: ALIGN_DIRECTIVE INT_OPERAND statement_block: LEFT_BRACE statement_list RIGHT_BRACE statement_list: directive_statement { add_directive(); } + | statement_list prototype_block {printf("Prototype statement detected. WARNING: this is not supported yet on GPGPU-SIM\n"); } | instruction_statement { add_instruction(); } | statement_list directive_statement { add_directive(); } | statement_list instruction_statement { add_instruction(); } @@ -403,6 +404,23 @@ initializer_list: LEFT_BRACE literal_list RIGHT_BRACE { add_array_initializer(); literal_list: literal_operand | literal_list COMMA literal_operand; +// TODO: This is currently hardcoded to handle and ignore one specific case +// that all prototype statements follow in the PTX from Pytorch. As a +// workaround, this parses and ignores both the prototype declaration +// and calling of the prototype (which conveniently comes right after the +// declaration for all cases.) This should be changed to handle both +// declaring the prototype, and actually calling it. +prototype_block: prototype_decl prototype_call + +prototype_decl: IDENTIFIER COLON CALLPROTOTYPE_DIRECTIVE LEFT_PAREN prototype_param RIGHT_PAREN IDENTIFIER LEFT_PAREN prototype_param RIGHT_PAREN SEMI_COLON + +prototype_call: OPCODE LEFT_PAREN IDENTIFIER RIGHT_PAREN COMMA operand COMMA LEFT_PAREN IDENTIFIER RIGHT_PAREN COMMA IDENTIFIER SEMI_COLON + | OPCODE IDENTIFIER COMMA LEFT_PAREN IDENTIFIER RIGHT_PAREN COMMA IDENTIFIER SEMI_COLON + +prototype_param: /* empty */ + | PARAM_DIRECTIVE B64_TYPE IDENTIFIER + | PARAM_DIRECTIVE B32_TYPE IDENTIFIER + instruction_statement: instruction SEMI_COLON | IDENTIFIER COLON { add_label($1); } | pred_spec instruction SEMI_COLON; |