slm: OpenCL code base  0.1
Functions
Streamline integration

Kernels and functions used to integrate streamlines. More...

Functions

__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. More...
 
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. More...
 
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. More...
 

Detailed Description

Kernels and functions used to integrate streamlines.

Function Documentation

◆ integrate_trajectory()

__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.

Further integration parameters are provided in this struct.

The kernel acts on one seed point only. It chooses this seed point by computing a global id and using it to index the seed_point_array. UPDATE: now doing sub-pixel streamlines as a set per seed point... need to doc here

Each streamline trajectory is returned in the appropriate location in trajectories_array as a list of compressed-into-byte dx,dy values.

Compiled if KERNEL_INTEGRATE_TRAJECTORY is defined.

Parameters
[in]seed_point_arraylist of initial streamline point vectors, one allotted to each kernel instance
[in]mask_arraygrid pixel mask (padded), with true = masked, false = good
[in]uv_arrayflow unit velocity vector grid (padded)
[out]trajectories_arraylists of streamline trajectories, stored as compressed-into-byte dx,dy vector sequences; one list per seed_point_array vector
[out]traj_nsteps_arraylist of number of steps along each trajectory; one per seed_point_array vector
[out]traj_length_arraylist of lengths of each trajectory; one per seed_point_array vector
[out]slc_arraygrid recording accumulated count of streamline integration steps across each pixel (padded)
[out]slt_arraygrid recording accumulated count of streamline segment lengths crossing each pixel (padded)
Returns
void

Definition at line 88 of file integration.cl.

◆ trajectory_jittered()

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 
)
inlinestatic

Integrate a jittered flow path downstream or upstream.

Write the streamline count and lengths to slc, slt arrays. Don't record the trajectory itself.

Compiled if KERNEL_INTEGRATE_TRAJECTORY is defined.

Parameters
[in]uv_arrayflow unit velocity vector grid (padded)
[in]mask_arraygrid pixel mask (padded), with true = masked, false = good
[out]slc_arraygrid recording accumulated count of streamline integration steps across each pixel (padded)
[out]slt_arraygrid recording accumulated count of streamline segment lengths crossing each pixel (padded)
[in]global_idID of the kernel instance
[in]seed_idxindex of the seed vector in the list seed_point_array; if chunkified, the sequence of indexes is offset from global_id by SEEDS_CHUNK_OFFSET
[in]current_seed_point_vecvector (real, float2) for the current point along the streamline trajectory
[in]initial_rng_stateRNG state and integer variate
Returns
void

Definition at line 107 of file trajectory.cl.

◆ trajectory_record()

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 
)
inlinestatic

Integrate a streamline downstream or upstream; record the trajectory.

Compiled if KERNEL_INTEGRATE_TRAJECTORY is defined.

Parameters
[in]uv_arrayflow unit velocity vector grid (padded)
[in]mask_arraygrid pixel mask (padded), with true = masked, false = good
[out]traj_nsteps_arraylist of number of steps along each trajectory; one per seed_point_array vector
[out]traj_length_arraylist of lengths of each trajectory; one per seed_point_array vector
[out]trajectory_vecsequence of compressed-as-byte dx,dy integration steps along this streamline trajectory
[in]global_idID of the kernel instance
[in]seed_idxindex of the seed vector in the list seed_point_array; if chunkified, the sequence of indexes is offset from global_id by SEEDS_CHUNK_OFFSET
[in]current_seed_point_vecvector (real, float2) for the current point along the streamline trajectory
Returns
void

Definition at line 40 of file trajectory.cl.