diff options
Diffstat (limited to 'intern/cycles/kernel/device/gpu/parallel_sorted_index.h')
-rw-r--r-- | intern/cycles/kernel/device/gpu/parallel_sorted_index.h | 49 |
1 files changed, 49 insertions, 0 deletions
diff --git a/intern/cycles/kernel/device/gpu/parallel_sorted_index.h b/intern/cycles/kernel/device/gpu/parallel_sorted_index.h new file mode 100644 index 00000000000..99b35468517 --- /dev/null +++ b/intern/cycles/kernel/device/gpu/parallel_sorted_index.h @@ -0,0 +1,49 @@ +/* + * Copyright 2021 Blender Foundation + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +CCL_NAMESPACE_BEGIN + +/* Given an array of states, build an array of indices for which the states + * are active and sorted by a given key. The prefix sum of the number of active + * states per key must have already been computed. + * + * TODO: there may be ways to optimize this to avoid this many atomic ops? */ + +#include "util/util_atomic.h" + +#define GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE 512 +#define GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY (~0) + +template<uint blocksize, typename GetKeyOp> +__device__ void gpu_parallel_sorted_index_array(const uint num_states, + int *indices, + int *num_indices, + int *key_prefix_sum, + GetKeyOp get_key_op) +{ + const uint state_index = ccl_gpu_block_idx_x * blocksize + ccl_gpu_thread_idx_x; + const int key = (state_index < num_states) ? get_key_op(state_index) : + GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY; + + if (key != GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY) { + const uint index = atomic_fetch_and_add_uint32(&key_prefix_sum[key], 1); + indices[index] = state_index; + } +} + +CCL_NAMESPACE_END |