diff options
| author | Hadi Jooybar <[email protected]> | 2012-11-05 19:35:21 -0800 |
|---|---|---|
| committer | Andrew Boktor <[email protected]> | 2014-08-14 13:49:21 -0700 |
| commit | 0b5efbc556621998ac42e96024163fae05beaf85 (patch) | |
| tree | ae17e55c4103b246597e6dd2ad1981cee90e0b7a /libopencl | |
| parent | 9af6f86d0f06c2ca1b117d358009b91a646b0e83 (diff) | |
OpenCL newer diver fix
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 14554]
Diffstat (limited to 'libopencl')
| -rw-r--r-- | libopencl/opencl_runtime_api.cc | 23 |
1 files changed, 15 insertions, 8 deletions
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 ); |