Page MenuHome
Differential D12888 Diff 43411 intern/cycles/kernel/device/gpu/kernel.h

Changeset View

Standalone View

View Options

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