From 690774e8f70624ef759f0f76f88d804f89db8f46 Mon Sep 17 00:00:00 2001 From: Roy Spliet Date: Thu, 25 Jul 2019 16:43:06 +0100 Subject: libopencl: Stub clGetProgramBuildInfo() Assumes everything is ok, even if it isn't. Allows for some Rodinia benchmarks to fail slightly later. v2: Use available case macros. Signed-off-by: Roy Spliet --- libopencl/opencl_runtime_api.cc | 29 +++++++++++++++++++++++++++++ 1 file changed, 29 insertions(+) (limited to 'libopencl/opencl_runtime_api.cc') diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 97a54d8..a21e11d 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -1255,6 +1255,35 @@ clGetProgramInfo(cl_program program, return CL_SUCCESS; } +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, -- cgit v1.3 From 6cdb7b18f70f8a1da200020f1b72f56fc33864d8 Mon Sep 17 00:00:00 2001 From: Roy Spliet Date: Thu, 25 Jul 2019 18:20:56 +0100 Subject: libopencl: Perform PDOM analysis Signed-off-by: Roy Spliet --- libopencl/opencl_runtime_api.cc | 11 +++++++++++ 1 file changed, 11 insertions(+) (limited to 'libopencl/opencl_runtime_api.cc') diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index a21e11d..23e52f6 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -950,6 +950,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 -- cgit v1.3 From 55a05c11b0af3159d81bce823c1fda4fb3cd04e3 Mon Sep 17 00:00:00 2001 From: Roy Spliet Date: Thu, 25 Jul 2019 18:34:44 +0100 Subject: libopencl: Calculate a valid offset in bind_args() v2: Adhere to alignment requirements. v3: Small style fix. Signed-off-by: Roy Spliet --- libopencl/opencl_runtime_api.cc | 14 +++++++++++++- 1 file changed, 13 insertions(+), 1 deletion(-) (limited to 'libopencl/opencl_runtime_api.cc') diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 23e52f6..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::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; -- cgit v1.3