ilusha/blender
1
0
Fork 0
mirror of https://github.com/blender/blender synced 2026-09-29 04:37:17 +03:00
Code Issues Projects Releases Packages Wiki Activity Actions

Cycles: oneAPI: Use sycl::inclusive_scan_over_group instead of ballot

Browse source 
Use sycl::inclusive_scan_over_group instead of the group_ballot
extension in adaptive_sampling_convergence_check and
integrator_shadow_catcher_count_possible_splits.

The current DPC++ implementation of the extension only supports
devices with sub-group sizes up to 64. The inclusive  scan works
also on devices with larger sub-group sizes.

Fixes the behaviour of the two affected kernels on the Snapdragon X
Elite GPU.

Pull Request: https://projects.blender.org/blender/blender/pulls/155801
This commit is contained in:
Rafal Bielski 2026-04-08 12:58:17 +02:00 • committed by Sergey Sharybin
parent 5dc6305602
commit d454d502b2
2 changed files with 26 additions and 7 deletions
Download patch file Download diff file Expand all files Collapse all files

22
intern/cycles/kernel/device/gpu/kernel.h
View file

@ -726,6 +726,7 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
const int work_index = ccl_gpu_global_id_x();
const int y = work_index / sw;
const int x = work_index - y * sw;
const int lane_id = ccl_gpu_thread_idx_x % ccl_gpu_warp_size;
bool converged = true;
@ -734,12 +735,20 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
nullptr, render_buffer, sx + x, sy + y, threshold, reset, offset, stride));
}
#ifdef __KERNEL_ONEAPI__
const sycl::nd_item<1> &item_id = sycl::ext::oneapi::this_work_item::get_nd_item<1>();
const uint num_active_pixels_in_warp = sycl::inclusive_scan_over_group(
item_id.get_sub_group(), static_cast<uint>(!converged), std::plus<>());
if (lane_id == item_id.get_sub_group().get_local_range()[0] - 1) {
atomic_fetch_and_add_uint32(num_active_pixels, num_active_pixels_in_warp);
}
#else
/* NOTE: All threads specified in the mask must execute the intrinsic. */
const auto num_active_pixels_mask = ccl_gpu_ballot(!converged);
const int lane_id = ccl_gpu_thread_idx_x % ccl_gpu_warp_size;
if (lane_id == 0) {
atomic_fetch_and_add_uint32(num_active_pixels, popcount(num_active_pixels_mask));
}
#endif
}
ccl_gpu_kernel_postfix
@ -1285,6 +1294,7 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
ccl_global uint *num_possible_splits)
{
const int state = ccl_gpu_global_id_x();
const int lane_id = ccl_gpu_thread_idx_x % ccl_gpu_warp_size;
bool can_split = false;
@ -1292,12 +1302,20 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
can_split = ccl_gpu_kernel_call(kernel_shadow_catcher_path_can_split(state));
}
#ifdef __KERNEL_ONEAPI__
const sycl::nd_item<1> &item_id = sycl::ext::oneapi::this_work_item::get_nd_item<1>();
const uint num_possible_splits_in_warp = sycl::inclusive_scan_over_group(
item_id.get_sub_group(), static_cast<uint>(can_split), std::plus<>());
if (lane_id == item_id.get_sub_group().get_local_range()[0] - 1) {
atomic_fetch_and_add_uint32(num_possible_splits, num_possible_splits_in_warp);
}
#else
/* NOTE: All threads specified in the mask must execute the intrinsic. */
const auto can_split_mask = ccl_gpu_ballot(can_split);
const int lane_id = ccl_gpu_thread_idx_x % ccl_gpu_warp_size;
if (lane_id == 0) {
atomic_fetch_and_add_uint32(num_possible_splits, popcount(can_split_mask));
}
#endif
}
ccl_gpu_kernel_postfix

11
intern/cycles/kernel/device/oneapi/compat.h
View file

@ -110,11 +110,12 @@ void oneapi_kernel_##name(KernelGlobalsGPU *ccl_restrict kg, \
/* GPU warp synchronization */
# define ccl_gpu_syncthreads() sycl::ext::oneapi::this_work_item::get_nd_item<1>().barrier()
# define ccl_gpu_local_syncthreads() sycl::ext::oneapi::this_work_item::get_nd_item<1>().barrier(sycl::access::fence_space::local_space)
# ifdef __SYCL_DEVICE_ONLY__
# define ccl_gpu_ballot(predicate) (sycl::ext::oneapi::group_ballot(sycl::ext::oneapi::this_work_item::get_sub_group(), predicate).count())
# else
# define ccl_gpu_ballot(predicate) (predicate ? 1 : 0)
# endif
/* A ballot in SYCL is only available as an Intel extension and its DPC++ v6.3 implementation
* does not support devices with sub-group sizes above 64. Summing values (of any type) within
* sub-groups can be achieved with the SYCL core feature inclusive_scan_over_group, which has
* better support on non-Intel devices. */
# define ccl_gpu_ballot(predicate) 0; static_assert(false, "Use sycl::inclusive_scan_over_group on oneAPI device instead of ccl_gpu_ballot")
/* Debug defines */
#if defined(__SYCL_DEVICE_ONLY__)