25 #define BYTE_REVERSAL(initial_rng_state) \ 26 initial_rng_state = (initial_rng_state>>24)&0xff \ 27 | (initial_rng_state>> 8)&0xff00 \ 28 | (initial_rng_state<< 8)&0xff0000 \ 29 | (initial_rng_state<<24)&0xff000000; 41 #define INITIALIZE_RNG(initial_rng_state) \ 42 initial_rng_state = i+(j+seed_idx); \ 45 BYTE_REVERSAL(initial_rng_state); \ 46 lehmer_rand_uint(&initial_rng_state); \ 47 BYTE_REVERSAL(initial_rng_state); 49 #ifdef KERNEL_INTEGRATE_TRAJECTORY 89 __global
const bool *mask_array,
90 __global
const float2 *uv_array,
91 __global char2 *trajectories_array,
92 __global ushort *traj_nsteps_array,
93 __global
float *traj_length_array,
94 __global uint *slc_array,
95 __global uint *slt_array )
98 const uint global_id = get_global_id(0u)+get_global_id(1u)*get_global_size(0u),
101 const float2 current_seed_point_vec = seed_point_array[seed_idx];
102 __global char2 *trajectory_vec = &trajectories_array[trajectory_index];
103 __private uint i,j, initial_rng_state;
107 printf(
"On GPU/OpenCL device: #workitems=%d #workgroups=%d\n",
108 get_local_size(0u), get_num_groups(0u));
113 trajectory_vec, global_id, seed_idx,
114 seed_point_array[seed_idx]);
128 current_seed_point_vec + (float2)(
__kernel void integrate_trajectory(__global const float2 *seed_point_array, __global const bool *mask_array, __global const float2 *uv_array, __global char2 *trajectories_array, __global ushort *traj_nsteps_array, __global float *traj_length_array, __global uint *slc_array, __global uint *slt_array)
GPU kernel that drives streamline integration from seed positions given in seed_point_array, controlled by the 'flow' vector field given in uv_array, and either terminated at pixels masked in mask_array or because a streamline exceeds a threshold distance (length or number of integration points) given by parameters stored in info.
#define SUBPIXEL_SEED_STEP
#define SUBPIXEL_SEED_POINT_DENSITY
#define SUBPIXEL_SEED_HALFSPAN
static void trajectory_jittered(__global const float2 *uv_array, __global const bool *mask_array, __global uint *slc_array, __global uint *slt_array, const uint global_id, const uint seed_idx, const float2 current_seed_point_vec, const uint initial_rng_state)
Integrate a jittered flow path downstream or upstream.
#define SEEDS_CHUNK_OFFSET
#define INITIALIZE_RNG(initial_rng_state)
Initialize the Lehmer random number generator.
static void trajectory_record(__global const float2 *uv_array, __global const bool *mask_array, __global ushort *traj_nsteps_array, __global float *traj_length_array, __global char2 *trajectory_vec, const uint global_id, const uint seed_idx, const float2 current_seed_point_vec)
Integrate a streamline downstream or upstream; record the trajectory.