summaryrefslogtreecommitdiff
path: root/libopencl/opencl_runtime_api.cc
diff options
context:
space:
mode:
Diffstat (limited to 'libopencl/opencl_runtime_api.cc')
-rw-r--r--libopencl/opencl_runtime_api.cc95
1 files changed, 78 insertions, 17 deletions
diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index 9da895d..30a50fc 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -74,7 +74,9 @@
#define __CUDA_RUNTIME_API_H__
#include "host_defines.h"
#include "builtin_types.h"
+#if (CUDART_VERSION < 8000)
#include "__cudaFatFormat.h"
+#endif
#include "../src/abstract_hardware_model.h"
#include "../src/cuda-sim/cuda-sim.h"
#include "../src/cuda-sim/ptx_loader.h"
@@ -82,6 +84,7 @@
#include "../src/gpgpusim_entrypoint.h"
#include "../src/gpgpu-sim/gpu-sim.h"
#include "../src/gpgpu-sim/shader.h"
+#include "../libcuda/gpgpu_context.h"
//# define __my_func__ __PRETTY_FUNCTION__
# if defined __cplusplus ? __GNUC_PREREQ (2, 6) : __GNUC_PREREQ (2, 4)
@@ -261,15 +264,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;
@@ -422,6 +437,8 @@ void register_ptx_function( const char *name, function_info *impl )
void _cl_program::Build(const char *options)
{
+ gpgpu_context *ctx;
+ ctx = GPGPU_Context();
printf("GPGPU-Sim OpenCL API: compiling OpenCL kernels...\n");
std::map<cl_uint,pgm_info>::iterator i;
for( i = m_pgm.begin(); i!= m_pgm.end(); i++ ) {
@@ -532,7 +549,7 @@ void _cl_program::Build(const char *options)
exit(1);
}
}
- if( !g_keep_intermediate_files ) {
+ if( !ctx->ptxinfo->g_keep_intermediate_files ) {
// clean up files...
snprintf(commandline,1024,"rm -f %s", cl_fname );
int result = system(commandline);
@@ -574,8 +591,8 @@ void _cl_program::Build(const char *options)
}
}
info.m_asm = tmp;
- info.m_symtab = gpgpu_ptx_sim_load_ptx_from_string( tmp, source_num );
- gpgpu_ptxinfo_load_from_string( tmp, source_num );
+ info.m_symtab = ctx->gpgpu_ptx_sim_load_ptx_from_string( tmp, source_num );
+ ctx->gpgpu_ptxinfo_load_from_string( tmp, source_num );
free(tmp);
}
printf("GPGPU-Sim OpenCL API: finished compiling OpenCL kernels.\n");
@@ -642,11 +659,13 @@ unsigned _cl_kernel::sm_context_uid = 0;
class _cl_device_id *GPGPUSim_Init()
{
static _cl_device_id *the_device = NULL;
+ gpgpu_context *ctx;
+ ctx = GPGPU_Context();
if( !the_device ) {
- gpgpu_sim *the_gpu = gpgpu_ptx_sim_init_perf();
+ gpgpu_sim *the_gpu = ctx->gpgpu_ptx_sim_init_perf();
the_device = new _cl_device_id(the_gpu);
}
- start_sim_thread(2);
+ ctx->start_sim_thread(2);
return the_device;
}
@@ -870,14 +889,16 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
const cl_event * event_wait_list,
cl_event * event) CL_API_SUFFIX__VERSION_1_0
{
+ gpgpu_context *ctx;
+ ctx = GPGPU_Context();
int _global_size[3];
int zeros[3] = { 0, 0, 0};
printf("\n\n\n");
char *mode = getenv("PTX_SIM_MODE_FUNC");
if ( mode )
- sscanf(mode,"%u", &g_ptx_sim_mode);
+ sscanf(mode,"%u", &(ctx->func_sim->g_ptx_sim_mode));
printf("GPGPU-Sim OpenCL API: clEnqueueNDRangeKernel '%s' (mode=%s)\n", kernel->name().c_str(),
- g_ptx_sim_mode?"functional simulation":"performance simulation");
+ (ctx->func_sim->g_ptx_sim_mode)?"functional simulation":"performance simulation");
if ( !work_dim || work_dim > 3 ) return CL_INVALID_WORK_DIMENSION;
size_t _local_size[3];
if( local_work_size != NULL ) {
@@ -941,17 +962,28 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
gpgpu_t *gpu = command_queue->get_device()->the_device();
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 );
+ ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1, gpu );
+ ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1, gpu );
+ ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1, gpu );
+ ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
+ ctx->func_sim->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 );
+ kernel_info_t *grid = ctx->func_sim->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 ( ctx->func_sim->g_ptx_sim_mode )
+ ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid );
else
- gpgpu_opencl_ptx_sim_main_perf( grid );
+ ctx->gpgpu_opencl_ptx_sim_main_perf( grid );
return CL_SUCCESS;
}
@@ -1254,6 +1286,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,
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 04:22:41 +0000