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.
Slightly modified the code to see the times of get of the frame from the camera and process it inside GPU. It turned out that the "blame" the camera (in this video I tried another camera sensor - similar behavior from it) - see the video below (again, I'm sorry that it pink - it's bug the software to capture images from the screen (gtk-recordmydesktop)). But still the question of the structure of the class uchar16 remains. Many thanks.
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