From 0b5efbc556621998ac42e96024163fae05beaf85 Mon Sep 17 00:00:00 2001 From: Hadi Jooybar Date: Mon, 5 Nov 2012 19:35:21 -0800 Subject: OpenCL newer diver fix [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 14554] --- libopencl/opencl_runtime_api.cc | 23 +++++++++++++++-------- 1 file changed, 15 insertions(+), 8 deletions(-) (limited to 'libopencl') 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 ); -- cgit v1.3