| 123 | // int n_per_thread = 2; // 2x2 per thread = 4 total elems per thread |
| 124 | template<typename T, int block_dim, int n_per_thread, bool full_conn> |
| 125 | __global__ static void update_equiv(arrayfire::cuda::Param<T> equiv_map, |
| 126 | const cudaTextureObject_t tex) { |
| 127 | // Basic coordinates |
| 128 | const int base_x = (blockIdx.x * blockDim.x * n_per_thread) + threadIdx.x; |
| 129 | const int base_y = (blockIdx.y * blockDim.y * n_per_thread) + threadIdx.y; |
| 130 | |
| 131 | const int width = equiv_map.dims[0]; |
| 132 | const int height = equiv_map.dims[1]; |
| 133 | |
| 134 | bool tid_changed = false; |
| 135 | |
| 136 | // Per element write flags and label, initially 0 |
| 137 | char write[n_per_thread * n_per_thread]; |
| 138 | T best_label[n_per_thread * n_per_thread]; |
| 139 | |
| 140 | #pragma unroll |
| 141 | for (int i = 0; i < n_per_thread * n_per_thread; ++i) { |
| 142 | write[i] = (char)0; |
| 143 | best_label[i] = (T)0; |
| 144 | } |
| 145 | |
| 146 | // Cached tile of the equivalency map |
| 147 | __shared__ T s_tile[n_per_thread * block_dim][(n_per_thread * block_dim)]; |
| 148 | |
| 149 | #pragma unroll |
| 150 | for (int xb = 0; xb < n_per_thread; ++xb) { |
| 151 | #pragma unroll |
| 152 | for (int yb = 0; yb < n_per_thread; ++yb) { |
| 153 | // Indexing variables |
| 154 | const int x = base_x + (xb * blockDim.x); |
| 155 | const int y = base_y + (yb * blockDim.y); |
| 156 | const int tx = threadIdx.x + (xb * blockDim.x); |
| 157 | const int ty = threadIdx.y + (yb * blockDim.y); |
| 158 | const int tid_i = xb * n_per_thread + yb; |
| 159 | const int n = y * width + x; |
| 160 | |
| 161 | // Get the label for this pixel if we're in bounds |
| 162 | const T orig_label = |
| 163 | (x < width && y < height) ? fetch<T>(n, equiv_map, tex) : (T)0; |
| 164 | s_tile[ty][tx] = orig_label; |
| 165 | |
| 166 | // Find the lowest label of the nearest valid pixel |
| 167 | // So far, all we know is that this pixel is valid. |
| 168 | best_label[tid_i] = orig_label; |
| 169 | |
| 170 | if (orig_label != (T)0) { |
| 171 | const int south_y = min(y, height - 2) + 1; |
| 172 | const int north_y = max(y, 1) - 1; |
| 173 | const int east_x = min(x, width - 2) + 1; |
| 174 | const int west_x = max(x, 1) - 1; |
| 175 | |
| 176 | // Check bottom |
| 177 | best_label[tid_i] = |
| 178 | relabel(best_label[tid_i], |
| 179 | fetch((south_y)*width + x, equiv_map, tex)); |
| 180 | |
| 181 | // Check right neighbor |
| 182 | best_label[tid_i] = |