Changeset View
Changeset View
Standalone View
Standalone View
intern/cycles/kernel/device/gpu/kernel.h
| Show First 20 Lines • Show All 45 Lines • ▼ Show 20 Lines | |||||
| */ | */ | ||||
| ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS) | ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS) | ||||
| kernel_gpu_integrator_reset(int num_states) | kernel_gpu_integrator_reset(int num_states) | ||||
| { | { | ||||
| const int state = ccl_gpu_global_id_x(); | const int state = ccl_gpu_global_id_x(); | ||||
| if (state < num_states) { | if (state < num_states) { | ||||
| INTEGRATOR_STATE_WRITE(path, queued_kernel) = 0; | INTEGRATOR_STATE_WRITE(state, path, queued_kernel) = 0; | ||||
| INTEGRATOR_STATE_WRITE(shadow_path, queued_kernel) = 0; | INTEGRATOR_STATE_WRITE(state, shadow_path, queued_kernel) = 0; | ||||
| } | } | ||||
| } | } | ||||
| ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS) | ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS) | ||||
| kernel_gpu_integrator_init_from_camera(KernelWorkTile *tiles, | kernel_gpu_integrator_init_from_camera(KernelWorkTile *tiles, | ||||
| const int num_tiles, | const int num_tiles, | ||||
| float *render_buffer, | float *render_buffer, | ||||
| const int max_tile_work_size) | const int max_tile_work_size) | ||||
| ▲ Show 20 Lines • Show All 175 Lines • ▼ Show 20 Lines | |||||
| extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE) | extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE) | ||||
| kernel_gpu_integrator_queued_paths_array(int num_states, | kernel_gpu_integrator_queued_paths_array(int num_states, | ||||
| int *indices, | int *indices, | ||||
| int *num_indices, | int *num_indices, | ||||
| int kernel) | int kernel) | ||||
| { | { | ||||
| gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>( | gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>( | ||||
| num_states, indices, num_indices, [kernel](const int state) { | num_states, indices, num_indices, [kernel](const int state) { | ||||
| return (INTEGRATOR_STATE(path, queued_kernel) == kernel); | return (INTEGRATOR_STATE(state, path, queued_kernel) == kernel); | ||||
| }); | }); | ||||
| } | } | ||||
| extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE) | extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE) | ||||
| kernel_gpu_integrator_queued_shadow_paths_array(int num_states, | kernel_gpu_integrator_queued_shadow_paths_array(int num_states, | ||||
| int *indices, | int *indices, | ||||
| int *num_indices, | int *num_indices, | ||||
| int kernel) | int kernel) | ||||
| { | { | ||||
| gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>( | gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>( | ||||
| num_states, indices, num_indices, [kernel](const int state) { | num_states, indices, num_indices, [kernel](const int state) { | ||||
| return (INTEGRATOR_STATE(shadow_path, queued_kernel) == kernel); | return (INTEGRATOR_STATE(state, shadow_path, queued_kernel) == kernel); | ||||
| }); | }); | ||||
| } | } | ||||
| extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE) | extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE) | ||||
| kernel_gpu_integrator_active_paths_array(int num_states, int *indices, int *num_indices) | kernel_gpu_integrator_active_paths_array(int num_states, int *indices, int *num_indices) | ||||
| { | { | ||||
| gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>( | gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>( | ||||
| num_states, indices, num_indices, [](const int state) { | num_states, indices, num_indices, [](const int state) { | ||||
| return (INTEGRATOR_STATE(path, queued_kernel) != 0) || | return (INTEGRATOR_STATE(state, path, queued_kernel) != 0) || | ||||
| (INTEGRATOR_STATE(shadow_path, queued_kernel) != 0); | (INTEGRATOR_STATE(state, shadow_path, queued_kernel) != 0); | ||||
| }); | }); | ||||
| } | } | ||||
| extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE) | extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE) | ||||
| kernel_gpu_integrator_terminated_paths_array(int num_states, | kernel_gpu_integrator_terminated_paths_array(int num_states, | ||||
| int *indices, | int *indices, | ||||
| int *num_indices, | int *num_indices, | ||||
| int indices_offset) | int indices_offset) | ||||
| { | { | ||||
| gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>( | gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>( | ||||
| num_states, indices + indices_offset, num_indices, [](const int state) { | num_states, indices + indices_offset, num_indices, [](const int state) { | ||||
| return (INTEGRATOR_STATE(path, queued_kernel) == 0) && | return (INTEGRATOR_STATE(state, path, queued_kernel) == 0) && | ||||
| (INTEGRATOR_STATE(shadow_path, queued_kernel) == 0); | (INTEGRATOR_STATE(state, shadow_path, queued_kernel) == 0); | ||||
| }); | }); | ||||
| } | } | ||||
| extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE) | extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE) | ||||
| kernel_gpu_integrator_sorted_paths_array( | kernel_gpu_integrator_sorted_paths_array( | ||||
| int num_states, int *indices, int *num_indices, int *key_prefix_sum, int kernel) | int num_states, int *indices, int *num_indices, int *key_prefix_sum, int kernel) | ||||
| { | { | ||||
| gpu_parallel_sorted_index_array<GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE>( | gpu_parallel_sorted_index_array<GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE>( | ||||
| num_states, indices, num_indices, key_prefix_sum, [kernel](const int state) { | num_states, indices, num_indices, key_prefix_sum, [kernel](const int state) { | ||||
| return (INTEGRATOR_STATE(path, queued_kernel) == kernel) ? | return (INTEGRATOR_STATE(state, path, queued_kernel) == kernel) ? | ||||
| INTEGRATOR_STATE(path, shader_sort_key) : | INTEGRATOR_STATE(state, path, shader_sort_key) : | ||||
| GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY; | GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY; | ||||
| }); | }); | ||||
| } | } | ||||
| extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE) | extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE) | ||||
| kernel_gpu_integrator_compact_paths_array(int num_states, | kernel_gpu_integrator_compact_paths_array(int num_states, | ||||
| int *indices, | int *indices, | ||||
| int *num_indices, | int *num_indices, | ||||
| int num_active_paths) | int num_active_paths) | ||||
| { | { | ||||
| gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>( | gpu_parallel_active_index_array<GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE>( | ||||
| num_states, indices, num_indices, [num_active_paths](const int state) { | num_states, indices, num_indices, [num_active_paths](const int state) { | ||||
| return (state >= num_active_paths) && | return (state >= num_active_paths) && | ||||
| ((INTEGRATOR_STATE(path, queued_kernel) != 0) || | ((INTEGRATOR_STATE(state, path, queued_kernel) != 0) || | ||||
| (INTEGRATOR_STATE(shadow_path, queued_kernel) != 0)); | (INTEGRATOR_STATE(state, shadow_path, queued_kernel) != 0)); | ||||
| }); | }); | ||||
| } | } | ||||
| extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE) | extern "C" __global__ void __launch_bounds__(GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE) | ||||
| kernel_gpu_integrator_compact_states(const int *active_terminated_states, | kernel_gpu_integrator_compact_states(const int *active_terminated_states, | ||||
| const int active_states_offset, | const int active_states_offset, | ||||
| const int terminated_states_offset, | const int terminated_states_offset, | ||||
| const int work_size) | const int work_size) | ||||
| ▲ Show 20 Lines • Show All 531 Lines • Show Last 20 Lines | |||||