summaryrefslogtreecommitdiff
path: root/libopencl/opencl_runtime_api.cc
diff options
context:
space:
mode:
authorTor Aamodt <[email protected]>2010-08-08 21:50:37 -0800
committerTor Aamodt <[email protected]>2010-08-08 21:50:37 -0800
commit250b9f3e84ccfe87557e36230ea081d9dea772e7 (patch)
treec26cc9e633dc089bfdb2a86545b1907ab99dcd34 /libopencl/opencl_runtime_api.cc
parentb400820185db543ff98fd823a15958444e12fee6 (diff)
oclSortingNetworks passing: update workgroup size calculation (now ignores
shared memory usage since I believe the assumption is at least a single CTA can run for this calculation) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7169]
Diffstat (limited to 'libopencl/opencl_runtime_api.cc')
-rw-r--r--libopencl/opencl_runtime_api.cc7
1 files changed, 0 insertions, 7 deletions
diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index 51ae840..5f2b1b7 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -327,18 +327,11 @@ extern unsigned int gpu_n_thread_per_shader;
size_t _cl_kernel::get_workgroup_size()
{
- unsigned smem = ptx_kernel_shmem_size( m_kernel_impl );
unsigned nregs = ptx_kernel_nregs( m_kernel_impl );
- unsigned result_shmem = (unsigned)-1;
unsigned result_regs = (unsigned)-1;
-
- if( smem > 0 )
- result_shmem = gpgpu_shmem_size / smem;
if( nregs > 0 )
result_regs = gpgpu_shader_registers / ((nregs+3)&~3);
-
unsigned result = gpu_n_thread_per_shader;
- result = min(result, result_shmem);
result = min(result, result_regs);
return (size_t)result;
}