mirror of
https://github.com/blender/blender
synced 2026-09-29 04:37:17 +03:00
9433b655c3
Each thread in the launched block runs over a portion of the counter array. First to calculate a sum of each portion, then summing those across all threads to get the base offset of each portion, before each thread does the prefix sum of its local portion again including that base offset. This could be improved further, but it already solves negative performance impact when increasing the number of states in the following commit. Pull Request: https://projects.blender.org/blender/blender/pulls/163930
66 lines
1.6 KiB
C
66 lines
1.6 KiB
C
/* SPDX-FileCopyrightText: 2021-2022 Blender Foundation
|
|
*
|
|
* SPDX-License-Identifier: Apache-2.0 */
|
|
|
|
#pragma once
|
|
|
|
CCL_NAMESPACE_BEGIN
|
|
|
|
/* Parallel prefix sum. */
|
|
|
|
#include "util/atomic.h"
|
|
|
|
__device__ void gpu_parallel_prefix_sum(const int global_id,
|
|
ccl_global int *ccl_restrict counter,
|
|
ccl_global int *ccl_restrict prefix_sum,
|
|
const int num_values)
|
|
{
|
|
int partition_start = 0;
|
|
int partition_end = num_values;
|
|
int offset = 0;
|
|
|
|
#ifndef __KERNEL_METAL__
|
|
const ushort local_id = ccl_gpu_thread_idx_x;
|
|
const ushort local_size = ccl_gpu_block_dim_x;
|
|
|
|
if (num_values >= local_size) {
|
|
const int partition_size = divide_up(num_values, local_size);
|
|
partition_start = partition_size * local_id;
|
|
partition_end = min(num_values, partition_start + partition_size);
|
|
|
|
int total = 0;
|
|
for (int i = partition_start; i < partition_end; i++) {
|
|
total += counter[i];
|
|
}
|
|
prefix_sum[local_id] = total;
|
|
|
|
ccl_gpu_syncthreads();
|
|
|
|
if (local_id == 0) {
|
|
for (int i = 0; i < local_size; i++) {
|
|
const int new_offset = offset + prefix_sum[i];
|
|
prefix_sum[i] = offset;
|
|
offset = new_offset;
|
|
}
|
|
}
|
|
|
|
ccl_gpu_syncthreads();
|
|
offset = prefix_sum[local_id];
|
|
ccl_gpu_syncthreads();
|
|
}
|
|
else
|
|
#endif
|
|
if (global_id != 0)
|
|
{
|
|
return;
|
|
}
|
|
|
|
for (int i = partition_start; i < partition_end; i++) {
|
|
const int new_offset = offset + counter[i];
|
|
prefix_sum[i] = offset;
|
|
counter[i] = 0;
|
|
offset = new_offset;
|
|
}
|
|
}
|
|
|
|
CCL_NAMESPACE_END
|