slm: OpenCL code base  0.1
connect.cl
Go to the documentation of this file.
1 
9 #ifdef KERNEL_CONNECT_CHANNELS
10 __kernel void connect_channels(
27  __global const float2 *seed_point_array,
28  __global const bool *mask_array,
29  __global const float2 *uv_array,
30  __global uint *mapping_array
31  )
32 {
33  // For every non-masked pixel...
34 
35  const uint global_id = get_global_id(0u)+get_global_id(1u)*get_global_size(0u);
36  const float2 current_seed_point_vec = seed_point_array[global_id];
37  __private uint idx, prev_idx, n_steps = 0u, step=0u;
38  __private float l_trajectory = 0.0f, dl = 0.0f, dt = DT_MAX;
39  __private float2 next_vec, uv1_vec, uv2_vec, dxy1_vec, dxy2_vec,
40  vec = current_seed_point_vec, prev_vec = vec;
41  __private char2 trajectory_vec[INTERCHANNEL_MAX_N_STEPS];
42 
43  // Remember here
44  prev_vec = vec;
45  idx = get_array_idx(vec);
46  prev_idx = idx;
47  // Integrate downstream one pixel
48  while (prev_idx==idx && !mask_array[idx] && n_steps!=MAX_N_STEPS) {
49  compute_step_vec(dt, uv_array, &dxy1_vec, &dxy2_vec, &uv1_vec, &uv2_vec,
50  vec, &next_vec, &idx);
51  if (!mask_array[idx]) {
52  if (connect_runge_kutta_step_record(&dt, &dl, &l_trajectory,
53  &dxy1_vec, &dxy2_vec, &vec, &prev_vec,
54  &next_vec, &n_steps, &idx, &prev_idx,
55  trajectory_vec))
56  continue;
57  }
58  }
59  // Integrate until we're back onto a channel pixel OR we reach a masked pixel
60  while ((mapping_array[idx] & IS_CHANNEL)==0 && !mask_array[idx]
61  && n_steps!=INTERCHANNEL_MAX_N_STEPS) {
62  compute_step_vec(dt, uv_array, &dxy1_vec, &dxy2_vec, &uv1_vec, &uv2_vec,
63  vec, &next_vec, &idx);
64  if (!mask_array[idx]) {
65  if (connect_runge_kutta_step_record(&dt, &dl, &l_trajectory,
66  &dxy1_vec, &dxy2_vec, &vec, &prev_vec,
67  &next_vec, &n_steps, &idx, &prev_idx,
68  trajectory_vec))
69  continue;
70  }
71  }
72  if (n_steps>2 && n_steps<INTERCHANNEL_MAX_N_STEPS) {
73  // At this point, we have either connected some channel pixels (type=1)
74  // or simply reached the trajectory tracking limit, or reached a masked pixel.
75  // Now we need to designate all intervening pixels since the last channel pixel
76  // as 'intervening channel' type=3.
77  vec = current_seed_point_vec;
78  idx = get_array_idx(vec);
79  step = 0u;
80  while (!mask_array[idx] && step<n_steps-1) {
81  // If this pixel was between channels, flag as both (1=channel; 2=between)
82  if (mapping_array[idx]==0u) {
83  atomic_or(&mapping_array[idx],IS_INTERCHANNEL);
84  }
85  // Increment along recorded trajectory, skipping first point
86  vec = vec + uncompress(trajectory_vec[step]);
87  idx = get_array_idx(vec);
88  step++;
89  }
90  }
91  return;
92 }
93 #endif
94 
95 #ifdef KERNEL_PUSH_TO_EXIT
96 
97 #define CHECK_TAIL(nbr_vec) { \
98  nbr_idx = get_array_idx(nbr_vec); \
99  if ( !mask_array[nbr_idx] && (mapping_array[nbr_idx]&IS_THINCHANNEL) ) { \
100  tail_flag |= 1; \
101  } \
102 }
103 #define CHECK_EXIT(nbr_vec) { \
104  nbr_idx = get_array_idx(nbr_vec); \
105  if ( mask_array[nbr_idx] ) { \
106  exit_flag |= 1; \
107  } \
108 }
109 // Check in all 8 pixel-nbr directions
110 #define CHECK_E_TAIL(vec) CHECK_TAIL((float2)(vec[0]+1.0f, vec[1] ))
111 #define CHECK_NE_TAIL(vec) CHECK_TAIL((float2)(vec[0]+1.0f, vec[1]+1.0f ))
112 #define CHECK_N_TAIL(vec) CHECK_TAIL((float2)(vec[0], vec[1]+1.0f ))
113 #define CHECK_NW_TAIL(vec) CHECK_TAIL((float2)(vec[0]-1.0f, vec[1]+1.0f ))
114 #define CHECK_W_TAIL(vec) CHECK_TAIL((float2)(vec[0]-1.0f, vec[1] ))
115 #define CHECK_SW_TAIL(vec) CHECK_TAIL((float2)(vec[0]-1.0f, vec[1]-1.0f ))
116 #define CHECK_S_TAIL(vec) CHECK_TAIL((float2)(vec[0], vec[1]-1.0f ))
117 #define CHECK_SE_TAIL(vec) CHECK_TAIL((float2)(vec[0]+1.0f, vec[1]-1.0f ))
118 // Check in all 8 pixel-nbr directions
119 #define CHECK_E_EXIT(vec) CHECK_EXIT((float2)(vec[0]+1.0f, vec[1] ))
120 #define CHECK_NE_EXIT(vec) CHECK_EXIT((float2)(vec[0]+1.0f, vec[1]+1.0f ))
121 #define CHECK_N_EXIT(vec) CHECK_EXIT((float2)(vec[0], vec[1]+1.0f ))
122 #define CHECK_NW_EXIT(vec) CHECK_EXIT((float2)(vec[0]-1.0f, vec[1]+1.0f ))
123 #define CHECK_W_EXIT(vec) CHECK_EXIT((float2)(vec[0]-1.0f, vec[1] ))
124 #define CHECK_SW_EXIT(vec) CHECK_EXIT((float2)(vec[0]-1.0f, vec[1]-1.0f ))
125 #define CHECK_S_EXIT(vec) CHECK_EXIT((float2)(vec[0], vec[1]-1.0f ))
126 #define CHECK_SE_EXIT(vec) CHECK_EXIT((float2)(vec[0]+1.0f, vec[1]-1.0f ))
127 
144 __kernel void push_to_exit(
145  __global const float2 *seed_point_array,
146  __global const bool *mask_array,
147  __global const float2 *uv_array,
148  __global uint *mapping_array
149  )
150 {
151  // For every channel but not thin-channel pixel...
152 
153  const uint global_id = get_global_id(0u)+get_global_id(1u)*get_global_size(0u);
154  __private uint idx, nbr_idx, i;
155  __private float2 vec = seed_point_array[global_id];
156  __private uchar tail_flag=0, exit_flag=0;
157 
158  idx = get_array_idx(vec);
159  CHECK_N_TAIL(vec);
160  CHECK_S_TAIL(vec);
161  CHECK_E_TAIL(vec);
162  CHECK_W_TAIL(vec);
163  CHECK_NE_TAIL(vec);
164  CHECK_SE_TAIL(vec);
165  CHECK_NW_TAIL(vec);
166  CHECK_SW_TAIL(vec);
167  CHECK_N_EXIT(vec);
168  CHECK_S_EXIT(vec);
169  CHECK_E_EXIT(vec);
170  CHECK_W_EXIT(vec);
171  CHECK_NE_EXIT(vec);
172  CHECK_SE_EXIT(vec);
173  CHECK_NW_EXIT(vec);
174  CHECK_SW_EXIT(vec);
175  if (tail_flag && exit_flag) {
176  atomic_or(&mapping_array[idx],IS_THINCHANNEL);
177  }
178  return;
179 }
180 #endif
__kernel void push_to_exit(__global const float2 *seed_point_array, __global const bool *mask_array, __global const float2 *uv_array, __global uint *mapping_array)
Connect up dangling channels to the masked grid boundary.
Definition: connect.cl:144
#define DT_MAX
Definition: info.h:48
#define CHECK_NE_EXIT(vec)
Definition: connect.cl:120
#define CHECK_W_TAIL(vec)
Definition: connect.cl:114
#define CHECK_SW_EXIT(vec)
Definition: connect.cl:124
#define CHECK_SE_TAIL(vec)
Definition: connect.cl:117
#define CHECK_N_TAIL(vec)
Definition: connect.cl:112
__kernel void connect_channels(__global const float2 *seed_point_array, __global const bool *mask_array, __global const float2 *uv_array, __global uint *mapping_array)
Connect up channel strands by designating intervening pixels as channel pixels.
Definition: connect.cl:26
static float2 uncompress(char2 compressed_vector)
Unsquish a byte vector back to a float vector.
Definition: essentials.cl:146
#define CHECK_E_TAIL(vec)
Definition: connect.cl:110
#define CHECK_SE_EXIT(vec)
Definition: connect.cl:126
#define IS_INTERCHANNEL
Definition: info.h:98
#define CHECK_S_EXIT(vec)
Definition: connect.cl:125
#define CHECK_NW_EXIT(vec)
Definition: connect.cl:122
#define CHECK_W_EXIT(vec)
Definition: connect.cl:123
static void compute_step_vec(const float dt, const __global float2 *uv_array, float2 *dxy1_vec, float2 *dxy2_vec, float2 *uv1_vec, float2 *uv2_vec, const float2 vec, float2 *next_vec, uint *idx)
Compute a 2nd-order Runge-Kutta integration step along a streamline.
Definition: computestep.cl:41
#define CHECK_E_EXIT(vec)
Definition: connect.cl:119
#define CHECK_NE_TAIL(vec)
Definition: connect.cl:111
static uint get_array_idx(float2 vec)
Compute the array index of the padded grid pixel pointed to by a float2 grid position vector (choice ...
Definition: essentials.cl:62
#define CHECK_S_TAIL(vec)
Definition: connect.cl:116
#define CHECK_SW_TAIL(vec)
Definition: connect.cl:115
#define IS_THINCHANNEL
Definition: info.h:97
#define MAX_N_STEPS
Definition: info.h:49
#define INTERCHANNEL_MAX_N_STEPS
Definition: info.h:51
#define IS_CHANNEL
Definition: info.h:96
#define CHECK_N_EXIT(vec)
Definition: connect.cl:121
#define CHECK_NW_TAIL(vec)
Definition: connect.cl:113
static bool connect_runge_kutta_step_record(float *dt, float *dl, float *l_trajectory, float2 *dxy1_vec, float2 *dxy2_vec, float2 *vec, float2 *prev_vec, float2 *next_vec, uint *n_steps, uint *idx, uint *prev_idx, __private char2 *trajectory_vec)
Compute a single step of 2nd-order Runge-Kutta numerical integration of a streamline given precompute...