Hello!
First - thank you for your forum and support!
I decided to create a new thread as it will be better this way. Earlier, I asked about the total load of all cores and Mali-T628MP6 though not yet figured out until the end there, but there was another question during these experiments.
I do not understand uchar16 class structure in Kernel code (see below code). Where it is declared and it's header? I had a problem that the FPS of the camera drops sharply (by 2 times! From 30 to 15!) And I found it in different lighting conditions the camera! And this behavior is observed in 3 different cameras. Perhaps it is not connected with the GPU work and I try to clarify it, But if you can explain about this class uchar16, I would be very grateful. A code of waters where it is used (kernel OpenCL code - color_convert.cl):
#define SCALE 8 #define SHIFT 14 #define R_W ((ushort)( 4899)) #define G_W ((ushort)( 9617)) #define B_W ((ushort)( 1868)) __kernel void convert_BGR2GRAY(__global uchar *src, __global uchar *dst, uint src_width_1_16 ) { int x = get_global_id(0); int y = get_global_id(1); size_t offset = ( x + y * src_width_1_16 ); uint16 out; uchar16 in0 = vload16( offset * 3 + 0, src ); uchar16 in1 = vload16( offset * 3 + 1, src ); uchar16 in2 = vload16( offset * 3 + 2, src ); /* Reorganize the data so that: in0 contains all the red in1 contains all the green in2 contains all the blue */ { uchar8 r,g,b; b = (uchar8) (in0.s0369, in0.scf, in1.s25); g = (uchar8) (in0.s147a, in0.sd, in1.s036); r = (uchar8) (in0.s258b, in0.se, in1.s147); in0 = (uchar16) (r, b); in1.s01234567 = g; b = (uchar8) (in1.s8be, in2.s147a, in2.sd); g = (uchar8) (in1.s9cf, in2.s258b, in2.se); r = (uchar8) (in1.sad, in2.s0369, in2.scf); in2 = (uchar16) (in0.s89abcdef, b); in0.s89abcdef = r; in1.s89abcdef = g; } out = convert_uint16(in0) * R_W; out += convert_uint16(in1) * G_W; out += convert_uint16(in2) * B_W; out += (uint16)(1 << (SHIFT-1)); out >>= SHIFT; vstore16( convert_uchar16( out ) , offset, dst ); } __kernel void convert_YUYV2GRAY(__global uchar *src, __global uchar *dst, uint src_width_1_16 ) { int x = get_global_id(0); int y = get_global_id(1); size_t offset = ( x + y * src_width_1_16 ); uint16 out; uchar16 in0 = vload16( offset * 2 + 0, src ); uchar16 in1 = vload16( offset * 2 + 1, src ); /* Keep only the Y component */ vstore16( (uchar16)( in0.even, in1.even ) , offset, dst ); }
For example what is this - in0.s0369 in line #27? Where this 's0369' field is declared?? Where this uchar16 declared? I could not find its definition. On the off. site including. Found here, but there is no 's0369' field declaration. Thank you.
********* 20161810 Update 01
I attach a video of the effect [removed] This is a screenshot of working that example App ("Accessing a Camera Device under Linux"). Pink little left (screen capture software error) - In fact, all black and white but it does not matter . But the main question is about uchar16 . Thank you.
Hi aleksbak,
uchar16
This is a vector of 16 uchar elements.
s0369
This is a swizzle, extracting lanes 0, 3, 6, and 9 of the vector to form a new 4 element vector containing the given elements.
HTH,Pete