Commit: 25e3eb87299b0a430c584ef304df75ccfacc57b3
Author: Xavier Hallade
Date: Wed Jun 29 10:37:37 2022 +0200
Branches: cycles_oneapi
https://developer.blender.org/rB25e3eb87299b0a430c584ef304df75ccfacc57b3
Cycles oneAPI: re-merge parallel_active_index.h
===================================================================
M intern/cycles/kernel/device/gpu/parallel_active_index.h
===================================================================
diff --git a/intern/cycles/kernel/device/gpu/parallel_active_index.h
b/intern/cycles/kernel/device/gpu/parallel_active_index.h
index cf53bf85067..c1df49c4f49 100644
--- a/intern/cycles/kernel/device/gpu/parallel_active_index.h
+++ b/intern/cycles/kernel/device/gpu/parallel_active_index.h
@@ -22,21 +22,7 @@ CCL_NAMESPACE_BEGIN
* ccl_gpu_thread_warp, ccl_gpu_warp_index, ccl_gpu_num_warps for all devices
* and keep device specific code in compat.h */
-#ifdef __KERNEL_METAL__
-void gpu_parallel_active_index_array_impl(const uint num_states,
- ccl_global int *indices,
- ccl_global int *num_indices,
- const uint is_active,
- const uint blocksize,
- const int thread_index,
- const uint state_index,
- const int ccl_gpu_warp_size,
- const int thread_warp,
- const int warp_index,
- const int num_warps,
- threadgroup int *warp_offset)
-{
-#elif defined(__KERNEL_ONEAPI__)
+#ifdef __KERNEL_ONEAPI__
# ifdef WITH_ONEAPI_SYCL_HOST_ENABLED
template<typename IsActiveOp>
void cpu_serial_active_index_array_impl(const uint num_states,
@@ -52,7 +38,8 @@ void cpu_serial_active_index_array_impl(const uint num_states,
*num_indices = write_index;
return;
}
-# endif
+# endif /* WITH_ONEAPI_SYCL_HOST_ENABLED */
+
template<typename IsActiveOp>
void gpu_parallel_active_index_array_impl(const uint num_states,
ccl_global int *ccl_restrict indices,
@@ -83,12 +70,28 @@ void gpu_parallel_active_index_array_impl(const uint
num_states,
/* Test if state corresponding to this thread is active. */
const uint is_active = (state_index < num_states) ?
is_active_op(state_index) : 0;
-#else
+#else /* !__KERNEL__ONEAPI__ */
+# ifndef __KERNEL_METAL__
template<uint blocksize, typename IsActiveOp>
-__device__ void gpu_parallel_active_index_array_impl(const uint num_states,
- ccl_global int *indices,
- ccl_global int
*num_indices,
- IsActiveOp is_active_op)
+__device__
+# endif
+ void
+ gpu_parallel_active_index_array_impl(const uint num_states,
+ ccl_global int *indices,
+ ccl_global int *num_indices,
+# ifdef __KERNEL_METAL__
+ const uint is_active,
+ const uint blocksize,
+ const int thread_index,
+ const uint state_index,
+ const int ccl_gpu_warp_size,
+ const int thread_warp,
+ const int warp_index,
+ const int num_warps,
+ threadgroup int *warp_offset)
+{
+# else
+ IsActiveOp is_active_op)
{
extern ccl_gpu_shared int warp_offset[];
@@ -102,8 +105,8 @@ __device__ void gpu_parallel_active_index_array_impl(const
uint num_states,
/* Test if state corresponding to this thread is active. */
const uint is_active = (state_index < num_states) ?
is_active_op(state_index) : 0;
-#endif
-
+# endif
+#endif /* !__KERNEL_ONEAPI__ */
/* For each thread within a warp compute how many other active states
precede it. */
#ifdef __KERNEL_ONEAPI__
const uint thread_offset = sycl::exclusive_scan_over_group(
_______________________________________________
Bf-blender-cvs mailing list
Bf-blender-cvs@blender.org
List details, subscription details or unsubscribe:
https://lists.blender.org/mailman/listinfo/bf-blender-cvs
