bad AF state logging (#1365)
parent
e568d3cadc
commit
968e2585cc
8 changed files with 358 additions and 4 deletions
@ -0,0 +1,110 @@ |
|||||||
|
// const __constant float3 rgb_weights = (0.299, 0.587, 0.114); // opencv rgb2gray weights |
||||||
|
// const __constant float3 bgr_weights = (0.114, 0.587, 0.299); // bgr2gray weights |
||||||
|
|
||||||
|
// convert input rgb image to single channel then conv |
||||||
|
__kernel void rgb2gray_conv2d( |
||||||
|
const __global uchar * input, |
||||||
|
__global short * output, |
||||||
|
__constant short * filter, |
||||||
|
__local uchar3 * cached |
||||||
|
) |
||||||
|
{ |
||||||
|
const int rowOffset = get_global_id(1) * IMAGE_W; |
||||||
|
const int my = get_global_id(0) + rowOffset; |
||||||
|
|
||||||
|
const int localRowLen = TWICE_HALF_FILTER_SIZE + get_local_size(0); |
||||||
|
const int localRowOffset = ( get_local_id(1) + HALF_FILTER_SIZE ) * localRowLen; |
||||||
|
const int myLocal = localRowOffset + get_local_id(0) + HALF_FILTER_SIZE; |
||||||
|
|
||||||
|
// cache local pixels |
||||||
|
cached[ myLocal ].x = input[ my * 3 ]; // r |
||||||
|
cached[ myLocal ].y = input[ my * 3 + 1]; // g |
||||||
|
cached[ myLocal ].z = input[ my * 3 + 2]; // b |
||||||
|
|
||||||
|
// pad |
||||||
|
if ( |
||||||
|
get_global_id(0) < HALF_FILTER_SIZE || |
||||||
|
get_global_id(0) > IMAGE_W - HALF_FILTER_SIZE - 1 || |
||||||
|
get_global_id(1) < HALF_FILTER_SIZE || |
||||||
|
get_global_id(1) > IMAGE_H - HALF_FILTER_SIZE - 1 |
||||||
|
) |
||||||
|
{ |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
return; |
||||||
|
} |
||||||
|
else |
||||||
|
{ |
||||||
|
int localColOffset = -1; |
||||||
|
int globalColOffset = -1; |
||||||
|
|
||||||
|
// cache extra |
||||||
|
if ( get_local_id(0) < HALF_FILTER_SIZE ) |
||||||
|
{ |
||||||
|
localColOffset = get_local_id(0); |
||||||
|
globalColOffset = -HALF_FILTER_SIZE; |
||||||
|
|
||||||
|
cached[ localRowOffset + get_local_id(0) ].x = input[ my * 3 - HALF_FILTER_SIZE * 3 ]; |
||||||
|
cached[ localRowOffset + get_local_id(0) ].y = input[ my * 3 - HALF_FILTER_SIZE * 3 + 1]; |
||||||
|
cached[ localRowOffset + get_local_id(0) ].z = input[ my * 3 - HALF_FILTER_SIZE * 3 + 2]; |
||||||
|
} |
||||||
|
else if ( get_local_id(0) >= get_local_size(0) - HALF_FILTER_SIZE ) |
||||||
|
{ |
||||||
|
localColOffset = get_local_id(0) + TWICE_HALF_FILTER_SIZE; |
||||||
|
globalColOffset = HALF_FILTER_SIZE; |
||||||
|
|
||||||
|
cached[ myLocal + HALF_FILTER_SIZE ].x = input[ my * 3 + HALF_FILTER_SIZE * 3 ]; |
||||||
|
cached[ myLocal + HALF_FILTER_SIZE ].y = input[ my * 3 + HALF_FILTER_SIZE * 3 + 1]; |
||||||
|
cached[ myLocal + HALF_FILTER_SIZE ].z = input[ my * 3 + HALF_FILTER_SIZE * 3 + 2]; |
||||||
|
} |
||||||
|
|
||||||
|
|
||||||
|
if ( get_local_id(1) < HALF_FILTER_SIZE ) |
||||||
|
{ |
||||||
|
cached[ get_local_id(1) * localRowLen + get_local_id(0) + HALF_FILTER_SIZE ].x = input[ my * 3 - HALF_FILTER_SIZE_IMAGE_W * 3 ]; |
||||||
|
cached[ get_local_id(1) * localRowLen + get_local_id(0) + HALF_FILTER_SIZE ].y = input[ my * 3 - HALF_FILTER_SIZE_IMAGE_W * 3 + 1]; |
||||||
|
cached[ get_local_id(1) * localRowLen + get_local_id(0) + HALF_FILTER_SIZE ].z = input[ my * 3 - HALF_FILTER_SIZE_IMAGE_W * 3 + 2]; |
||||||
|
if (localColOffset > 0) |
||||||
|
{ |
||||||
|
cached[ get_local_id(1) * localRowLen + localColOffset ].x = input[ my * 3 - HALF_FILTER_SIZE_IMAGE_W * 3 + globalColOffset * 3]; |
||||||
|
cached[ get_local_id(1) * localRowLen + localColOffset ].y = input[ my * 3 - HALF_FILTER_SIZE_IMAGE_W * 3 + globalColOffset * 3 + 1]; |
||||||
|
cached[ get_local_id(1) * localRowLen + localColOffset ].z = input[ my * 3 - HALF_FILTER_SIZE_IMAGE_W * 3 + globalColOffset * 3 + 2]; |
||||||
|
} |
||||||
|
} |
||||||
|
else if ( get_local_id(1) >= get_local_size(1) -HALF_FILTER_SIZE ) |
||||||
|
{ |
||||||
|
int offset = ( get_local_id(1) + TWICE_HALF_FILTER_SIZE ) * localRowLen; |
||||||
|
cached[ offset + get_local_id(0) + HALF_FILTER_SIZE ].x = input[ my * 3 + HALF_FILTER_SIZE_IMAGE_W * 3 ]; |
||||||
|
cached[ offset + get_local_id(0) + HALF_FILTER_SIZE ].y = input[ my * 3 + HALF_FILTER_SIZE_IMAGE_W * 3 + 1]; |
||||||
|
cached[ offset + get_local_id(0) + HALF_FILTER_SIZE ].z = input[ my * 3 + HALF_FILTER_SIZE_IMAGE_W * 3 + 2]; |
||||||
|
if (localColOffset > 0) |
||||||
|
{ |
||||||
|
cached[ offset + localColOffset ].x = input[ my * 3 + HALF_FILTER_SIZE_IMAGE_W * 3 + globalColOffset * 3]; |
||||||
|
cached[ offset + localColOffset ].y = input[ my * 3 + HALF_FILTER_SIZE_IMAGE_W * 3 + globalColOffset * 3 + 1]; |
||||||
|
cached[ offset + localColOffset ].z = input[ my * 3 + HALF_FILTER_SIZE_IMAGE_W * 3 + globalColOffset * 3 + 2]; |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
// sync |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
// perform convolution |
||||||
|
int fIndex = 0; |
||||||
|
short sum = 0; |
||||||
|
|
||||||
|
for (int r = -HALF_FILTER_SIZE; r <= HALF_FILTER_SIZE; r++) |
||||||
|
{ |
||||||
|
int curRow = r * localRowLen; |
||||||
|
for (int c = -HALF_FILTER_SIZE; c <= HALF_FILTER_SIZE; c++, fIndex++) |
||||||
|
{ |
||||||
|
if (!FLIP_RB){ |
||||||
|
// sum += dot(rgb_weights, cached[ myLocal + curRow + c ]) * filter[ fIndex ]; |
||||||
|
sum += (cached[ myLocal + curRow + c ].x / 3 + cached[ myLocal + curRow + c ].y / 2 + cached[ myLocal + curRow + c ].z / 9) * filter[ fIndex ]; |
||||||
|
} else { |
||||||
|
// sum += dot(bgr_weights, cached[ myLocal + curRow + c ]) * filter[ fIndex ]; |
||||||
|
sum += (cached[ myLocal + curRow + c ].x / 9 + cached[ myLocal + curRow + c ].y / 2 + cached[ myLocal + curRow + c ].z / 3) * filter[ fIndex ]; |
||||||
|
} |
||||||
|
} |
||||||
|
} |
||||||
|
output[my] = sum; |
||||||
|
} |
||||||
|
} |
@ -0,0 +1,34 @@ |
|||||||
|
// calculate variance in each subregion |
||||||
|
__kernel void var_pool( |
||||||
|
const __global char * input, |
||||||
|
__global ushort * output // should not be larger than 128*128 so uint16 |
||||||
|
) |
||||||
|
{ |
||||||
|
const int xidx = get_global_id(0) + ROI_X_MIN; |
||||||
|
const int yidx = get_global_id(1) + ROI_Y_MIN; |
||||||
|
|
||||||
|
const int size = X_PITCH * Y_PITCH; |
||||||
|
|
||||||
|
float fsum = 0; |
||||||
|
char mean, max; |
||||||
|
|
||||||
|
for (int i = 0; i < size; i++) { |
||||||
|
int x_offset = i % X_PITCH; |
||||||
|
int y_offset = i / X_PITCH; |
||||||
|
fsum += input[xidx*X_PITCH + yidx*Y_PITCH*FULL_STRIDE_X + x_offset + y_offset*FULL_STRIDE_X]; |
||||||
|
max = input[xidx*X_PITCH + yidx*Y_PITCH*FULL_STRIDE_X + x_offset + y_offset*FULL_STRIDE_X]>max ? input[xidx*X_PITCH + yidx*Y_PITCH*FULL_STRIDE_X + x_offset + y_offset*FULL_STRIDE_X]:max; |
||||||
|
} |
||||||
|
|
||||||
|
mean = convert_char_rte(fsum / size); |
||||||
|
|
||||||
|
float fvar = 0; |
||||||
|
for (int i = 0; i < size; i++) { |
||||||
|
int x_offset = i % X_PITCH; |
||||||
|
int y_offset = i / X_PITCH; |
||||||
|
fvar += (input[xidx*X_PITCH + yidx*Y_PITCH*FULL_STRIDE_X + x_offset + y_offset*FULL_STRIDE_X] - mean) * (input[xidx*X_PITCH + yidx*Y_PITCH*FULL_STRIDE_X + x_offset + y_offset*FULL_STRIDE_X] - mean); |
||||||
|
} |
||||||
|
|
||||||
|
fvar = fvar / size; |
||||||
|
|
||||||
|
output[(xidx-ROI_X_MIN)+(yidx-ROI_Y_MIN)*(ROI_X_MAX-ROI_X_MIN+1)] = convert_ushort_rte(5 * fvar + convert_float_rte(max)); |
||||||
|
} |
@ -0,0 +1,44 @@ |
|||||||
|
#include "utils.h" |
||||||
|
#include <stdio.h> |
||||||
|
#include <algorithm> |
||||||
|
|
||||||
|
// calculate score based on laplacians in one area
|
||||||
|
void get_lapmap_one(int16_t *lap, uint16_t *res, int x_pitch, int y_pitch) { |
||||||
|
int size = x_pitch * y_pitch; |
||||||
|
// avg and max of roi
|
||||||
|
float fsum = 0; |
||||||
|
int16_t mean, max; |
||||||
|
max = 0; |
||||||
|
|
||||||
|
for (int i = 0; i < size; i++) { |
||||||
|
int x_offset = i % x_pitch; |
||||||
|
int y_offset = i / x_pitch; |
||||||
|
fsum += lap[x_offset + y_offset*x_pitch]; |
||||||
|
max = std::max(lap[x_offset + y_offset*x_pitch], max); |
||||||
|
} |
||||||
|
|
||||||
|
mean = fsum / size; |
||||||
|
|
||||||
|
// var of roi
|
||||||
|
float fvar = 0; |
||||||
|
for (int i = 0; i < size; i++) { |
||||||
|
int x_offset = i % x_pitch; |
||||||
|
int y_offset = i / x_pitch; |
||||||
|
fvar += (float)((lap[x_offset + y_offset*x_pitch] - mean) * (lap[x_offset + y_offset*x_pitch] - mean)); |
||||||
|
} |
||||||
|
|
||||||
|
fvar = fvar / size; |
||||||
|
|
||||||
|
*res = std::min(5 * fvar + max, (float)65535); |
||||||
|
} |
||||||
|
|
||||||
|
bool is_blur(uint16_t *lapmap) { |
||||||
|
int n_roi = (ROI_X_MAX - ROI_X_MIN + 1) * (ROI_Y_MAX - ROI_Y_MIN + 1); |
||||||
|
float bad_sum = 0; |
||||||
|
for (int i = 0; i < n_roi; i++) { |
||||||
|
if (*(lapmap + i) < LM_THRESH) { |
||||||
|
bad_sum += 1/(float)n_roi; |
||||||
|
} |
||||||
|
} |
||||||
|
return (bad_sum > LM_PREC_THRESH); |
||||||
|
} |
@ -0,0 +1,30 @@ |
|||||||
|
#ifndef IMGPROC_UTILS |
||||||
|
#define IMGPROC_UTILS |
||||||
|
|
||||||
|
#include <stdint.h> |
||||||
|
|
||||||
|
#define NUM_SEGMENTS_X 8 |
||||||
|
#define NUM_SEGMENTS_Y 6 |
||||||
|
|
||||||
|
#define ROI_X_MIN 1 |
||||||
|
#define ROI_X_MAX 6 |
||||||
|
#define ROI_Y_MIN 2 |
||||||
|
#define ROI_Y_MAX 3 |
||||||
|
|
||||||
|
#define LM_THRESH 222 |
||||||
|
#define LM_PREC_THRESH 0.9 |
||||||
|
|
||||||
|
// only apply to QCOM
|
||||||
|
#define FULL_STRIDE_X 1280 |
||||||
|
#define FULL_STRIDE_Y 896 |
||||||
|
|
||||||
|
#define CONV_LOCAL_WORKSIZE 16 |
||||||
|
|
||||||
|
const int16_t lapl_conv_krnl[9] = {0, 1, 0, |
||||||
|
1, -4, 1, |
||||||
|
0, 1, 0}; |
||||||
|
|
||||||
|
void get_lapmap_one(int16_t *lap, uint16_t *res, int x_pitch, int y_pitch); |
||||||
|
bool is_blur(uint16_t *lapmap); |
||||||
|
|
||||||
|
#endif |
Loading…
Reference in new issue