Differential D13361 Diff 45286 intern/cycles/kernel/device/gpu/parallel_active_index.h

Changeset View

Standalone View

Show First 20 Lines • Show All 79 Lines • ▼ Show 20 Lines	__device__ void gpu_parallel_active_index_array(const uint num_states,

const uint state_index = ccl_gpu_block_idx_x * blocksize + thread_index;		const uint state_index = ccl_gpu_block_idx_x * blocksize + thread_index;
#endif		#endif

/* Test if state corresponding to this thread is active. */		/* Test if state corresponding to this thread is active. */
const uint is_active = (state_index < num_states) ? is_active_op(state_index) : 0;		const uint is_active = (state_index < num_states) ? is_active_op(state_index) : 0;

/* For each thread within a warp compute how many other active states precede it. */		/* For each thread within a warp compute how many other active states precede it. */
const uint thread_offset = ccl_gpu_popc(ccl_gpu_ballot(is_active) &		const uint thread_offset = popcount(ccl_gpu_ballot(is_active) &
ccl_gpu_thread_mask(thread_warp));		ccl_gpu_thread_mask(thread_warp));

/* Last thread in warp stores number of active states for each warp. */		/* Last thread in warp stores number of active states for each warp. */
if (thread_warp == ccl_gpu_warp_size - 1) {		if (thread_warp == ccl_gpu_warp_size - 1) {
warp_offset[warp_index] = thread_offset + is_active;		warp_offset[warp_index] = thread_offset + is_active;
}		}

ccl_gpu_syncthreads();		ccl_gpu_syncthreads();

▲ Show 20 Lines • Show All 42 Lines • Show Last 20 Lines