Welcome to mirror list, hosted at ThFree Co, Russian Federation.

git.blender.org/blender.git - Unnamed repository; edit this file 'description' to name the repository.
summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorStefan Werner <stefan.werner@tangent-animation.com>2018-10-28 12:37:41 +0300
committerStefan Werner <stefan.werner@tangent-animation.com>2018-10-28 12:37:41 +0300
commite58c6cf0c678849cf9c348a8df5e0ec24a6abd4d (patch)
tree7c6dc887ad93e87ad56c3d0c627265a99c7f1e7d /intern/cycles/kernel/kernels/cuda
parentc0b3e3daebd36a483e659d32e6517f2fb9b0e277 (diff)
Cycles: Added Cryptomatte output.
This allows for extra output passes that encode automatic object and material masks for the entire scene. It is an implementation of the Cryptomatte standard as introduced by Psyop. A good future extension would be to add a manifest to the export and to do plenty of testing to ensure that it is fully compatible with other renderers and compositing programs that use Cryptomatte. Internally, it adds the ability for Cycles to have several passes of the same type that are distinguished by their name. Differential Revision: https://developer.blender.org/D3538
Diffstat (limited to 'intern/cycles/kernel/kernels/cuda')
-rw-r--r--intern/cycles/kernel/kernels/cuda/kernel.cu30
1 files changed, 22 insertions, 8 deletions
diff --git a/intern/cycles/kernel/kernels/cuda/kernel.cu b/intern/cycles/kernel/kernels/cuda/kernel.cu
index 8a180a509e8..af311027f78 100644
--- a/intern/cycles/kernel/kernels/cuda/kernel.cu
+++ b/intern/cycles/kernel/kernels/cuda/kernel.cu
@@ -40,14 +40,21 @@ CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
kernel_cuda_path_trace(WorkTile *tile, uint total_work_size)
{
int work_index = ccl_global_id(0);
-
- if(work_index < total_work_size) {
- uint x, y, sample;
+ bool thread_is_active = work_index < total_work_size;
+ uint x, y, sample;
+ KernelGlobals kg;
+ if(thread_is_active) {
get_work_pixel(tile, work_index, &x, &y, &sample);
- KernelGlobals kg;
kernel_path_trace(&kg, tile->buffer, sample, x, y, tile->offset, tile->stride);
}
+
+ if(kernel_data.film.cryptomatte_passes) {
+ __syncthreads();
+ if(thread_is_active) {
+ kernel_cryptomatte_post(&kg, tile->buffer, sample, x, y, tile->offset, tile->stride);
+ }
+ }
}
#ifdef __BRANCHED_PATH__
@@ -56,14 +63,21 @@ CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_BRANCHED_MAX_REGISTERS)
kernel_cuda_branched_path_trace(WorkTile *tile, uint total_work_size)
{
int work_index = ccl_global_id(0);
-
- if(work_index < total_work_size) {
- uint x, y, sample;
+ bool thread_is_active = work_index < total_work_size;
+ uint x, y, sample;
+ KernelGlobals kg;
+ if(thread_is_active) {
get_work_pixel(tile, work_index, &x, &y, &sample);
- KernelGlobals kg;
kernel_branched_path_trace(&kg, tile->buffer, sample, x, y, tile->offset, tile->stride);
}
+
+ if(kernel_data.film.cryptomatte_passes) {
+ __syncthreads();
+ if(thread_is_active) {
+ kernel_cryptomatte_post(&kg, tile->buffer, sample, x, y, tile->offset, tile->stride);
+ }
+ }
}
#endif