diff options
Diffstat (limited to 'intern/cycles/kernel/device/gpu/parallel_active_index.h')
| -rw-r--r-- | intern/cycles/kernel/device/gpu/parallel_active_index.h | 33 |
1 files changed, 5 insertions, 28 deletions
diff --git a/intern/cycles/kernel/device/gpu/parallel_active_index.h b/intern/cycles/kernel/device/gpu/parallel_active_index.h index c1df49c4f49..38cdcb572eb 100644 --- a/intern/cycles/kernel/device/gpu/parallel_active_index.h +++ b/intern/cycles/kernel/device/gpu/parallel_active_index.h @@ -23,22 +23,6 @@ CCL_NAMESPACE_BEGIN * and keep device specific code in compat.h */ #ifdef __KERNEL_ONEAPI__ -# ifdef WITH_ONEAPI_SYCL_HOST_ENABLED -template<typename IsActiveOp> -void cpu_serial_active_index_array_impl(const uint num_states, - ccl_global int *ccl_restrict indices, - ccl_global int *ccl_restrict num_indices, - IsActiveOp is_active_op) -{ - int write_index = 0; - for (int state_index = 0; state_index < num_states; state_index++) { - if (is_active_op(state_index)) - indices[write_index++] = state_index; - } - *num_indices = write_index; - return; -} -# endif /* WITH_ONEAPI_SYCL_HOST_ENABLED */ template<typename IsActiveOp> void gpu_parallel_active_index_array_impl(const uint num_states, @@ -182,18 +166,11 @@ __device__ num_simd_groups, \ simdgroup_offset) #elif defined(__KERNEL_ONEAPI__) -# ifdef WITH_ONEAPI_SYCL_HOST_ENABLED -# define gpu_parallel_active_index_array( \ - blocksize, num_states, indices, num_indices, is_active_op) \ - if (ccl_gpu_global_size_x() == 1) \ - cpu_serial_active_index_array_impl(num_states, indices, num_indices, is_active_op); \ - else \ - gpu_parallel_active_index_array_impl(num_states, indices, num_indices, is_active_op); -# else -# define gpu_parallel_active_index_array( \ - blocksize, num_states, indices, num_indices, is_active_op) \ - gpu_parallel_active_index_array_impl(num_states, indices, num_indices, is_active_op) -# endif + +# define gpu_parallel_active_index_array( \ + blocksize, num_states, indices, num_indices, is_active_op) \ + gpu_parallel_active_index_array_impl(num_states, indices, num_indices, is_active_op) + #else # define gpu_parallel_active_index_array( \ |